llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clangir Author: David Rivera (RiverDave) <details> <summary>Changes</summary> Blocker from MiniFE (unscoped float atomicAdd). --- Full diff: https://github.com/llvm/llvm-project/pull/221262.diff 2 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp (+2-4) - (modified) clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu (+16) ``````````diff 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 `````````` </details> https://github.com/llvm/llvm-project/pull/221262 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
