https://github.com/RiverDave created https://github.com/llvm/llvm-project/pull/221262
Blocker from MiniFE (unscoped float atomicAdd). >From 6729b77da177783bacb8c4a34b214020befaba75 Mon Sep 17 00:00:00 2001 From: David Rivera <[email protected]> Date: Fri, 4 Sep 2026 11:00:05 -0400 Subject: [PATCH] [CIR][NVPTX] Lower unscoped __nvvm_atom_add_gen_{f,d} Blocked MiniFE (unscoped float atomicAdd). --- clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp | 6 ++---- .../test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu | 16 ++++++++++++++++ 2 files changed, 18 insertions(+), 4 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp index ae994005c588a..e125f642f66cf 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp @@ -178,10 +178,8 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { // success flag. 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 1ca877017beb2..b0c6f394302e2 100644 --- a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu +++ b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu @@ -300,6 +300,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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
