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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
