Author: Ayokunle Amodu Date: 2026-07-29T19:48:08-04:00 New Revision: 3ca8b6e0afef0370c189721f71cb9092ea8009fd
URL: https://github.com/llvm/llvm-project/commit/3ca8b6e0afef0370c189721f71cb9092ea8009fd DIFF: https://github.com/llvm/llvm-project/commit/3ca8b6e0afef0370c189721f71cb9092ea8009fd.diff LOG: [CIR][CUDA] Add support for NVVM xchg builtins (#211815) Adds codegen support for the scoped and unscoped NVVM atomic exchange builtins: `atom_xchg,` `atom_cta_xchg,` and `atom_sys_xchg.` These are lowered to the corresponding CIR `cir.atomic.xchg` operations and subsequently lowered to LLVM `atomicrmw xchg` instructions. 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 11c4abc0b3ba5..839e768c7fc88 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp @@ -64,6 +64,19 @@ static mlir::Value makeScopedAtomicRMW(CIRGenFunction &cgf, return rmwi->getResult(0); } +static mlir::Value makeScopedAtomicXchg(CIRGenFunction &cgf, + const CallExpr *expr, + cir::SyncScopeKind scope) { + auto &builder = cgf.getBuilder(); + Address destAddr = cgf.emitPointerWithAlignment(expr->getArg(0)); + mlir::Value destValue = destAddr.emitRawPointer(); + mlir::Value val = cgf.emitScalarExpr(expr->getArg(1)); + auto xchg = cir::AtomicXchgOp::create( + builder, cgf.getLoc(expr->getSourceRange()), destValue, val, + cir::MemOrder::Relaxed, scope, /*is_volatile=*/false); + return xchg.getResult(); +} + std::optional<mlir::Value> CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { switch (builtinId) { @@ -105,10 +118,7 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { case NVPTX::BI__nvvm_atom_xchg_gen_i: case NVPTX::BI__nvvm_atom_xchg_gen_l: case NVPTX::BI__nvvm_atom_xchg_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicXchg(*this, expr, cir::SyncScopeKind::System); case NVPTX::BI__nvvm_atom_max_gen_i: case NVPTX::BI__nvvm_atom_max_gen_l: case NVPTX::BI__nvvm_atom_max_gen_ll: @@ -257,17 +267,11 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { case NVPTX::BI__nvvm_atom_cta_xchg_gen_i: case NVPTX::BI__nvvm_atom_cta_xchg_gen_l: case NVPTX::BI__nvvm_atom_cta_xchg_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicXchg(*this, expr, cir::SyncScopeKind::Workgroup); case NVPTX::BI__nvvm_atom_sys_xchg_gen_i: case NVPTX::BI__nvvm_atom_sys_xchg_gen_l: case NVPTX::BI__nvvm_atom_sys_xchg_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicXchg(*this, expr, cir::SyncScopeKind::System); case NVPTX::BI__nvvm_atom_cta_max_gen_i: case NVPTX::BI__nvvm_atom_cta_max_gen_ui: case NVPTX::BI__nvvm_atom_cta_max_gen_l: diff --git a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu index a5ed80759d2ad..1ca877017beb2 100644 --- a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu +++ b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu @@ -228,6 +228,30 @@ __device__ void test_atom_min_gen_ull(unsigned long long *p, unsigned long long __nvvm_atom_min_gen_ull(p, val); } +// CIR-LABEL: @_Z20test_atom_xchg_gen_iPii +// CIR: cir.atomic.xchg relaxed syncscope(system) %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z20test_atom_xchg_gen_iPii +// LLVM: atomicrmw xchg ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_xchg_gen_i(int *p, int val) { + __nvvm_atom_xchg_gen_i(p, val); +} + +// CIR-LABEL: @_Z20test_atom_xchg_gen_lPll +// CIR: cir.atomic.xchg relaxed syncscope(system) %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z20test_atom_xchg_gen_lPll +// LLVM: atomicrmw xchg ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_xchg_gen_l(long *p, long val) { + __nvvm_atom_xchg_gen_l(p, val); +} + +// CIR-LABEL: @_Z21test_atom_xchg_gen_llPxx +// CIR: cir.atomic.xchg relaxed syncscope(system) %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z21test_atom_xchg_gen_llPxx +// LLVM: atomicrmw xchg ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_xchg_gen_ll(long long *p, long long val) { + __nvvm_atom_xchg_gen_ll(p, val); +} + // CIR-LABEL: @_Z23test_atom_cta_add_gen_iPii // CIR: cir.atomic.fetch add relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i // LLVM-LABEL: @_Z23test_atom_cta_add_gen_iPii @@ -675,3 +699,51 @@ __device__ void test_atom_cta_xor_gen_ll(long long *p, long long val) { __device__ void test_atom_sys_xor_gen_ll(long long *p, long long val) { __nvvm_atom_sys_xor_gen_ll(p, val); } + +// CIR-LABEL: @_Z24test_atom_cta_xchg_gen_iPii +// CIR: cir.atomic.xchg relaxed syncscope(workgroup) %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z24test_atom_cta_xchg_gen_iPii +// LLVM: atomicrmw xchg ptr %{{.*}}, i32 %{{.*}} syncscope("block") monotonic, align 4 +__device__ void test_atom_cta_xchg_gen_i(int *p, int val) { + __nvvm_atom_cta_xchg_gen_i(p, val); +} + +// CIR-LABEL: @_Z24test_atom_sys_xchg_gen_iPii +// CIR: cir.atomic.xchg relaxed syncscope(system) %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z24test_atom_sys_xchg_gen_iPii +// LLVM: atomicrmw xchg ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_sys_xchg_gen_i(int *p, int val) { + __nvvm_atom_sys_xchg_gen_i(p, val); +} + +// CIR-LABEL: @_Z24test_atom_cta_xchg_gen_lPll +// CIR: cir.atomic.xchg relaxed syncscope(workgroup) %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z24test_atom_cta_xchg_gen_lPll +// LLVM: atomicrmw xchg ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_xchg_gen_l(long *p, long val) { + __nvvm_atom_cta_xchg_gen_l(p, val); +} + +// CIR-LABEL: @_Z24test_atom_sys_xchg_gen_lPll +// CIR: cir.atomic.xchg relaxed syncscope(system) %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z24test_atom_sys_xchg_gen_lPll +// LLVM: atomicrmw xchg ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_xchg_gen_l(long *p, long val) { + __nvvm_atom_sys_xchg_gen_l(p, val); +} + +// CIR-LABEL: @_Z25test_atom_cta_xchg_gen_llPxx +// CIR: cir.atomic.xchg relaxed syncscope(workgroup) %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z25test_atom_cta_xchg_gen_llPxx +// LLVM: atomicrmw xchg ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_xchg_gen_ll(long long *p, long long val) { + __nvvm_atom_cta_xchg_gen_ll(p, val); +} + +// CIR-LABEL: @_Z25test_atom_sys_xchg_gen_llPxx +// CIR: cir.atomic.xchg relaxed syncscope(system) %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z25test_atom_sys_xchg_gen_llPxx +// LLVM: atomicrmw xchg ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_xchg_gen_ll(long long *p, long long val) { + __nvvm_atom_sys_xchg_gen_ll(p, val); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
