atomic fences cause errors and atomic instructions are emitted without fences on `nvptx64-nvidia-cuda`

[usamoi]

I tried this code:

#![no_std]
#![allow(internal_features)]
#![feature(core_intrinsics)]

use core::sync::atomic::AtomicU32;

#[panic_handler]
fn panic_handler(_: &core::panic::PanicInfo<'_>) -> ! {
    loop {}
}

static COUNTER: AtomicU32 = AtomicU32::new(0);

// use intrinsics here, since inliner not work for `AtomicU32::fetch_add`

#[no_mangle]
fn atomic_xadd_seqcst() {
    unsafe {
        core::intrinsics::atomic_xadd_seqcst(COUNTER.as_ptr(), 1);
    }
}

#[no_mangle]
fn atomic_xadd_acqrel() {
    unsafe {
        core::intrinsics::atomic_xadd_acqrel(COUNTER.as_ptr(), 1);
    }
}

#[no_mangle]
fn atomic_xadd_relaxed() {
    unsafe {
        core::intrinsics::atomic_xadd_relaxed(COUNTER.as_ptr(), 1);
    }
}

[unstable]
build-std = ["core"]

[build]
target = "nvptx64-nvidia-cuda"
rustflags = [
    "-Clinker=llvm-bitcode-linker",
    "-Clinker-flavor=llbc",
    "-Zunstable-options",
    "-Clink-arg=-O3",
    "-Ctarget-cpu=sm_90",
]

Same instructions are generated for these 3 functions.

//
// Generated by LLVM NVPTX Back-End
//

.version 7.8
.target sm_90
.address_size 64

	// .globl	atomic_xadd_seqcst      // -- Begin function atomic_xadd_seqcst
.global .align 4 .b8 _ZN2nv7COUNTER17h4c11a16573192bf7E[4];
                                        // @atomic_xadd_seqcst
.visible .func atomic_xadd_seqcst()
{
	.reg .b32 	%r<2>;
	.reg .b64 	%rd<2>;

// %bb.0:
	mov.u64 	%rd1, _ZN2nv7COUNTER17h4c11a16573192bf7E;
	atom.global.add.u32 	%r1, [%rd1], 1;
	ret;
                                        // -- End function
}
	// .globl	atomic_xadd_acqrel      // -- Begin function atomic_xadd_acqrel
.visible .func atomic_xadd_acqrel()     // @atomic_xadd_acqrel
{
	.reg .b32 	%r<2>;
	.reg .b64 	%rd<2>;

// %bb.0:
	mov.u64 	%rd1, _ZN2nv7COUNTER17h4c11a16573192bf7E;
	atom.global.add.u32 	%r1, [%rd1], 1;
	ret;
                                        // -- End function
}
	// .globl	atomic_xadd_relaxed     // -- Begin function atomic_xadd_relaxed
.visible .func atomic_xadd_relaxed()    // @atomic_xadd_relaxed
{
	.reg .b32 	%r<2>;
	.reg .b64 	%rd<2>;

// %bb.0:
	mov.u64 	%rd1, _ZN2nv7COUNTER17h4c11a16573192bf7E;
	atom.global.add.u32 	%r1, [%rd1], 1;
	ret;
                                        // -- End function
}

Different instruction should be emitted.

Link for PTX document:

Parallel Synchronization and Communication Instructions: atom

Meta

rustc --version --verbose:

rustc 1.86.0-nightly (8239a37f9 2025-02-01)
binary: rustc
commit-hash: 8239a37f9c0951a037cfc51763ea52a20e71e6bd
commit-date: 2025-02-01
host: x86_64-unknown-linux-gnu
release: 1.86.0-nightly
LLVM version: 19.1.7

[usamoi]

This issue could be partially solved if llvm is updated.

If main branch of llvm is used, In latest nightly,

1. Fences don't cause errors any more.
2. rustc still emits atom.global.add.u32, instead of atom.global.add.acq_rel.u32.

[Noratrieb]

@RDambrosio016 @kjetilkjeka as nvptx maintainers

[kjetilkjeka]

Oh, that's not ideal. Good that it's partially fixed by llvm and good to have the problem documented.

Do you by any chance also have the exact llvm-ir that was produced by rustc?

[usamoi]

Do you by any chance also have the exact llvm-ir that was produced by rustc?

It's llvm-ir produced by nightly-2025-03-08.

; ModuleID = 'ru.5313cc6bb0664a09-cgu.0'
source_filename = "ru.5313cc6bb0664a09-cgu.0"
target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
target triple = "nvptx64-nvidia-cuda"

@_ZN2ru7COUNTER17he0d3afc3442f5f98E = internal global <{ [4 x i8] }> zeroinitializer, align 4

; Function Attrs: nofree norecurse noreturn nosync nounwind memory(none)
define void @rust_begin_unwind(ptr noalias nocapture noundef readonly align 8 dereferenceable(24) %_1) unnamed_addr #0 {
start:
  br label %bb1

bb1:                                              ; preds = %bb1, %start
  br label %bb1
}

; Function Attrs: mustprogress nofree norecurse nounwind willreturn
define void @foo() unnamed_addr #1 {
start:
  fence acquire
  ret void
}

; Function Attrs: mustprogress nofree norecurse nounwind willreturn memory(readwrite, argmem: none, inaccessiblemem: none)
define void @atomic_xadd_seqcst() unnamed_addr #2 {
start:
  %0 = atomicrmw add ptr @_ZN2ru7COUNTER17he0d3afc3442f5f98E, i32 1 seq_cst, align 4
  ret void
}

; Function Attrs: mustprogress nofree norecurse nounwind willreturn memory(readwrite, argmem: none, inaccessiblemem: none)
define void @atomic_xadd_acqrel() unnamed_addr #2 {
start:
  %0 = atomicrmw add ptr @_ZN2ru7COUNTER17he0d3afc3442f5f98E, i32 1 acq_rel, align 4
  ret void
}

; Function Attrs: mustprogress nofree norecurse nounwind willreturn memory(readwrite, argmem: none, inaccessiblemem: none)
define void @atomic_xadd_relaxed() unnamed_addr #2 {
start:
  %0 = atomicrmw add ptr @_ZN2ru7COUNTER17he0d3afc3442f5f98E, i32 1 monotonic, align 4
  ret void
}

attributes #0 = { nofree norecurse noreturn nosync nounwind memory(none) "target-cpu"="sm_30" }
attributes #1 = { mustprogress nofree norecurse nounwind willreturn "target-cpu"="sm_30" }
attributes #2 = { mustprogress nofree norecurse nounwind willreturn memory(readwrite, argmem: none, inaccessiblemem: none) "target-cpu"="sm_30" }

!llvm.module.flags = !{!0}
!llvm.ident = !{!1}

!0 = !{i32 8, !"PIC Level", i32 2}
!1 = !{!"rustc version 1.87.0-nightly (f5a1ef712 2025-03-07)"}