https://github.com/RiverDave created https://github.com/llvm/llvm-project/pull/221261
Blocker from CUB/Thrust and MiniFE (atomicCAS). >From e606c060bf585101d062dc482387bccf1944606b Mon Sep 17 00:00:00 2001 From: David Rivera <[email protected]> Date: Fri, 4 Sep 2026 10:58:55 -0400 Subject: [PATCH 1/2] [CIR][NVPTX] Lower __nvvm_atom_cas_gen_* builtins Blocked CUB/Thrust and MiniFE (atomicCAS). --- clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp | 33 +++++---- .../CIR/CodeGenCUDA/builtins-nvvm-atomic.cu | 70 +++++++++++++++++++ 2 files changed, 90 insertions(+), 13 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp index ae994005c588a..aedc34b836d4d 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp @@ -97,6 +97,23 @@ static mlir::Value makeScopedAtomicXchg(CIRGenFunction &cgf, return xchg.getResult(); } +/// Lower __nvvm_atom{,_cta,_sys}_cas_gen_* like classic codegen: +/// cmpxchg monotonic monotonic, return the old value (not the success flag). +static mlir::Value makeScopedAtomicCAS(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 expected = cgf.emitScalarExpr(expr->getArg(1)); + mlir::Value desired = cgf.emitScalarExpr(expr->getArg(2)); + auto cmpxchg = cir::AtomicCmpXchgOp::create( + builder, cgf.getLoc(expr->getSourceRange()), destValue, expected, desired, + cir::MemOrder::Relaxed, cir::MemOrder::Relaxed, scope, + /*alignment=*/nullptr, /*weak=*/false, /*is_volatile=*/false); + return cmpxchg.getOld(); +} + std::optional<mlir::Value> CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { switch (builtinId) { @@ -171,11 +188,7 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { case NVPTX::BI__nvvm_atom_cas_gen_i: case NVPTX::BI__nvvm_atom_cas_gen_l: case NVPTX::BI__nvvm_atom_cas_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - // success flag. + return makeScopedAtomicCAS(*this, expr, cir::SyncScopeKind::System); case NVPTX::BI__nvvm_atom_add_gen_f: case NVPTX::BI__nvvm_atom_add_gen_d: cgm.errorNYI(expr->getSourceRange(), @@ -367,18 +380,12 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { case NVPTX::BI__nvvm_atom_cta_cas_gen_i: case NVPTX::BI__nvvm_atom_cta_cas_gen_l: case NVPTX::BI__nvvm_atom_cta_cas_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicCAS(*this, expr, cir::SyncScopeKind::Workgroup); case NVPTX::BI__nvvm_atom_sys_cas_gen_us: case NVPTX::BI__nvvm_atom_sys_cas_gen_i: case NVPTX::BI__nvvm_atom_sys_cas_gen_l: case NVPTX::BI__nvvm_atom_sys_cas_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicCAS(*this, expr, cir::SyncScopeKind::System); case NVPTX::BI__nvvm_match_all_sync_i32p: case NVPTX::BI__nvvm_match_all_sync_i64p: cgm.errorNYI(expr->getSourceRange(), diff --git a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu index 1ca877017beb2..6695a94345e07 100644 --- a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu +++ b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu @@ -252,6 +252,76 @@ __device__ void test_atom_xchg_gen_ll(long long *p, long long val) { __nvvm_atom_xchg_gen_ll(p, val); } +// CIR-LABEL: @_Z19test_atom_cas_gen_iPiii +// CIR: cir.atomic.cmpxchg success(relaxed) failure(relaxed) syncscope(system) %{{.*}}, %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i, !s32i) -> (!s32i, !cir.bool) +// LLVM-LABEL: @_Z19test_atom_cas_gen_iPiii +// LLVM: %[[CAS:.*]] = cmpxchg ptr %{{.*}}, i32 %{{.*}}, i32 %{{.*}} monotonic monotonic, align 4 +// LLVM: extractvalue { i32, i1 } %[[CAS]], 0 +__device__ int test_atom_cas_gen_i(int *p, int cmp, int val) { + return __nvvm_atom_cas_gen_i(p, cmp, val); +} + +// CIR-LABEL: @_Z20test_atom_cas_gen_llPxxx +// CIR: cir.atomic.cmpxchg success(relaxed) failure(relaxed) syncscope(system) %{{.*}}, %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i, !s64i) -> (!s64i, !cir.bool) +// LLVM-LABEL: @_Z20test_atom_cas_gen_llPxxx +// LLVM: %[[CAS:.*]] = cmpxchg ptr %{{.*}}, i64 %{{.*}}, i64 %{{.*}} monotonic monotonic, align 8 +// LLVM: extractvalue { i64, i1 } %[[CAS]], 0 +__device__ long long test_atom_cas_gen_ll(long long *p, long long cmp, + long long val) { + return __nvvm_atom_cas_gen_ll(p, cmp, val); +} + +// CIR-LABEL: @_Z20test_atom_cas_gen_usPttt +// CIR: cir.atomic.cmpxchg success(relaxed) failure(relaxed) syncscope(system) %{{.*}}, %{{.*}}, %{{.*}} : (!cir.ptr<!u16i>, !u16i, !u16i) -> (!u16i, !cir.bool) +// LLVM-LABEL: @_Z20test_atom_cas_gen_usPttt +// LLVM: %[[CAS:.*]] = cmpxchg ptr %{{.*}}, i16 %{{.*}}, i16 %{{.*}} monotonic monotonic, align 2 +// LLVM: extractvalue { i16, i1 } %[[CAS]], 0 +__device__ unsigned short test_atom_cas_gen_us(unsigned short *p, + unsigned short cmp, + unsigned short val) { + return __nvvm_atom_cas_gen_us(p, cmp, val); +} + +// CIR-LABEL: @_Z23test_atom_cta_cas_gen_iPiii +// CIR: cir.atomic.cmpxchg success(relaxed) failure(relaxed) syncscope(workgroup) %{{.*}}, %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i, !s32i) -> (!s32i, !cir.bool) +// LLVM-LABEL: @_Z23test_atom_cta_cas_gen_iPiii +// LLVM: %[[CAS:.*]] = cmpxchg ptr %{{.*}}, i32 %{{.*}}, i32 %{{.*}} syncscope("block") monotonic monotonic, align 4 +// LLVM: extractvalue { i32, i1 } %[[CAS]], 0 +__device__ int test_atom_cta_cas_gen_i(int *p, int cmp, int val) { + return __nvvm_atom_cta_cas_gen_i(p, cmp, val); +} + +// CIR-LABEL: @_Z23test_atom_sys_cas_gen_iPiii +// CIR: cir.atomic.cmpxchg success(relaxed) failure(relaxed) syncscope(system) %{{.*}}, %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i, !s32i) -> (!s32i, !cir.bool) +// LLVM-LABEL: @_Z23test_atom_sys_cas_gen_iPiii +// LLVM: %[[CAS:.*]] = cmpxchg ptr %{{.*}}, i32 %{{.*}}, i32 %{{.*}} monotonic monotonic, align 4 +// LLVM: extractvalue { i32, i1 } %[[CAS]], 0 +__device__ int test_atom_sys_cas_gen_i(int *p, int cmp, int val) { + return __nvvm_atom_sys_cas_gen_i(p, cmp, val); +} + +// CIR-LABEL: @_Z24test_atom_cta_cas_gen_usPttt +// CIR: cir.atomic.cmpxchg success(relaxed) failure(relaxed) syncscope(workgroup) %{{.*}}, %{{.*}}, %{{.*}} : (!cir.ptr<!u16i>, !u16i, !u16i) -> (!u16i, !cir.bool) +// LLVM-LABEL: @_Z24test_atom_cta_cas_gen_usPttt +// LLVM: %[[CAS:.*]] = cmpxchg ptr %{{.*}}, i16 %{{.*}}, i16 %{{.*}} syncscope("block") monotonic monotonic, align 2 +// LLVM: extractvalue { i16, i1 } %[[CAS]], 0 +__device__ unsigned short test_atom_cta_cas_gen_us(unsigned short *p, + unsigned short cmp, + unsigned short val) { + return __nvvm_atom_cta_cas_gen_us(p, cmp, val); +} + +// CIR-LABEL: @_Z24test_atom_sys_cas_gen_usPttt +// CIR: cir.atomic.cmpxchg success(relaxed) failure(relaxed) syncscope(system) %{{.*}}, %{{.*}}, %{{.*}} : (!cir.ptr<!u16i>, !u16i, !u16i) -> (!u16i, !cir.bool) +// LLVM-LABEL: @_Z24test_atom_sys_cas_gen_usPttt +// LLVM: %[[CAS:.*]] = cmpxchg ptr %{{.*}}, i16 %{{.*}}, i16 %{{.*}} monotonic monotonic, align 2 +// LLVM: extractvalue { i16, i1 } %[[CAS]], 0 +__device__ unsigned short test_atom_sys_cas_gen_us(unsigned short *p, + unsigned short cmp, + unsigned short val) { + return __nvvm_atom_sys_cas_gen_us(p, cmp, 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 >From a2caa55f292f826d4ea090d4a09aa0b2ae666a00 Mon Sep 17 00:00:00 2001 From: David Rivera <[email protected]> Date: Fri, 4 Sep 2026 11:37:54 -0400 Subject: [PATCH 2/2] [CIR][NVPTX] Drop classic-codegen comments --- clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp index aedc34b836d4d..adabbd682daea 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp @@ -97,8 +97,6 @@ static mlir::Value makeScopedAtomicXchg(CIRGenFunction &cgf, return xchg.getResult(); } -/// Lower __nvvm_atom{,_cta,_sys}_cas_gen_* like classic codegen: -/// cmpxchg monotonic monotonic, return the old value (not the success flag). static mlir::Value makeScopedAtomicCAS(CIRGenFunction &cgf, const CallExpr *expr, cir::SyncScopeKind scope) { _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
