[clang] c510577 - [CIR][NVPTX] Lower unscoped __nvvm_atom_add_gen_{f,d} (#221262)

[via cfe-commits]

Author: David Rivera
Date: 2026-09-05T02:06:32-04:00
New Revision: c510577ad11d4c965076935b8fabf1dfe8c5703b

URL: https://github.com/llvm/llvm-project/commit/c510577ad11d4c965076935b8fabf1dfe8c5703b
DIFF: https://github.com/llvm/llvm-project/commit/c510577ad11d4c965076935b8fabf1dfe8c5703b.diff

LOG: [CIR][NVPTX] Lower unscoped __nvvm_atom_add_gen_{f,d} (#221262)

Added: 
    

Modified: 
    clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
    clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu

Removed: 
    


################################################################################
diff  --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
index 8b33d85d49fe1..5f284c11c3272 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
@@ -193,10 +193,8 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) {
                              cir::SyncScopeKind::System);
   case NVPTX::BI__nvvm_atom_add_gen_f:
   case NVPTX::BI__nvvm_atom_add_gen_d:
-    cgm.errorNYI(expr->getSourceRange(),
-                 std::string("unimplemented NVPTX builtin call: ") +
-                     getContext().BuiltinInfo.getName(builtinId));
-    return mlir::Value{};
+    return makeScopedAtomicRMW(*this, expr, cir::AtomicFetchKind::Add,
+                               cir::SyncScopeKind::System);
   case NVPTX::BI__nvvm_atom_inc_gen_ui:
     return makeBinaryAtomicValue(cir::AtomicFetchKind::UIncWrap, expr,
                                  /*originalArgType=*/nullptr,

diff  --git a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu
index 6695a94345e07..466aeb0f4f847 100644
--- a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu
+++ b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu
@@ -370,6 +370,22 @@ __device__ void test_atom_sys_add_gen_ll(long long *p, long long val) {
   __nvvm_atom_sys_add_gen_ll(p, val);
 }
 
+// CIR-LABEL: @_Z19test_atom_add_gen_fPff
+// CIR: cir.atomic.fetch add relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!cir.float>, !cir.float) -> !cir.float
+// LLVM-LABEL: @_Z19test_atom_add_gen_fPff
+// LLVM: atomicrmw fadd ptr %{{.*}}, float %{{.*}} monotonic, align 4
+__device__ void test_atom_add_gen_f(float *p, float val) {
+  __nvvm_atom_add_gen_f(p, val);
+}
+
+// CIR-LABEL: @_Z19test_atom_add_gen_dPdd
+// CIR: cir.atomic.fetch add relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!cir.double>, !cir.double) -> !cir.double
+// LLVM-LABEL: @_Z19test_atom_add_gen_dPdd
+// LLVM: atomicrmw fadd ptr %{{.*}}, double %{{.*}} monotonic, align 8
+__device__ void test_atom_add_gen_d(double *p, double val) {
+  __nvvm_atom_add_gen_d(p, val);
+}
+
 // CIR-LABEL: @_Z23test_atom_cta_add_gen_fPff
 // CIR: cir.atomic.fetch add relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!cir.float>, !cir.float) -> !cir.float
 // LLVM-LABEL: @_Z23test_atom_cta_add_gen_fPff