https://github.com/RiverDave updated https://github.com/llvm/llvm-project/pull/221261
>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/3] [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/3] [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) { >From fda9d402fac3410e06dbd8885c8476342f58e5e9 Mon Sep 17 00:00:00 2001 From: David Rivera <[email protected]> Date: Fri, 4 Sep 2026 15:45:50 -0400 Subject: [PATCH 3/3] [CIR] Promote emitAtomicCmpXchg and reuse it for NVPTX CAS --- clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp | 36 +++++++++----------- clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp | 27 +++++---------- clang/lib/CIR/CodeGen/CIRGenFunction.h | 8 +++++ 3 files changed, 33 insertions(+), 38 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp index 61cb04828e272..09320480db133 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp @@ -397,29 +397,25 @@ static RValue emitBinaryAtomicPost(CIRGenFunction &cgf, return RValue::get(result); } -/// Emit a `cir.atomic.cmpxchg` for __sync_val_compare_and_swap_N and -/// __sync_bool_compare_and_swap_N. Returns the old value when `returnBool` is -/// false, otherwise returns a boolean success flag. -static RValue emitAtomicCmpXchg(CIRGenFunction &cgf, const CallExpr *e, - bool returnBool) { - Address destAddr = checkAtomicAlignment(cgf, e); - CIRGenBuilderTy &builder = cgf.getBuilder(); +mlir::Value CIRGenFunction::emitAtomicCmpXchg(const CallExpr *e, + bool returnBool, + cir::MemOrder successOrder, + cir::MemOrder failureOrder, + cir::SyncScopeKind scope) { + Address destAddr = checkAtomicAlignment(*this, e); + CIRGenBuilderTy &builder = getBuilder(); mlir::Value destValue = destAddr.emitRawPointer(); - mlir::Value expected = cgf.emitScalarExpr(e->getArg(1)); - mlir::Value desired = cgf.emitScalarExpr(e->getArg(2)); + mlir::Value expected = emitScalarExpr(e->getArg(1)); + mlir::Value desired = emitScalarExpr(e->getArg(2)); auto cmpxchg = cir::AtomicCmpXchgOp::create( - builder, cgf.getLoc(e->getSourceRange()), destValue, expected, desired, - cir::MemOrder::SequentiallyConsistent, - cir::MemOrder::SequentiallyConsistent, cir::SyncScopeKind::System, + builder, getLoc(e->getSourceRange()), destValue, expected, desired, + successOrder, failureOrder, scope, /*alignment=*/nullptr, /*weak=*/false, /*is_volatile=*/false); - if (returnBool) { - // cir.atomic.cmpxchg already returns (old, success). Use the success flag - // directly instead of re-emitting the expected argument and comparing. - return RValue::get(cmpxchg.getSuccess()); - } - return RValue::get(cmpxchg.getOld()); + if (returnBool) + return cmpxchg.getSuccess(); + return cmpxchg.getOld(); } /// Emit a `cir.atomic.xchg` for __sync_swap_N and __sync_lock_test_and_set_N. @@ -2498,12 +2494,12 @@ RValue CIRGenFunction::emitBuiltinExpr(const GlobalDecl &gd, unsigned builtinID, case Builtin::BI__sync_val_compare_and_swap_2: case Builtin::BI__sync_val_compare_and_swap_4: case Builtin::BI__sync_val_compare_and_swap_8: - return emitAtomicCmpXchg(*this, e, /*returnBool=*/false); + return RValue::get(emitAtomicCmpXchg(e, /*returnBool=*/false)); case Builtin::BI__sync_bool_compare_and_swap_1: case Builtin::BI__sync_bool_compare_and_swap_2: case Builtin::BI__sync_bool_compare_and_swap_4: case Builtin::BI__sync_bool_compare_and_swap_8: - return emitAtomicCmpXchg(*this, e, /*returnBool=*/true); + return RValue::get(emitAtomicCmpXchg(e, /*returnBool=*/true)); case Builtin::BI__sync_swap_1: case Builtin::BI__sync_swap_2: case Builtin::BI__sync_swap_4: diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp index adabbd682daea..0ea7a8ebc293b 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp @@ -97,21 +97,6 @@ static mlir::Value makeScopedAtomicXchg(CIRGenFunction &cgf, return xchg.getResult(); } -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) { @@ -186,7 +171,9 @@ 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: - return makeScopedAtomicCAS(*this, expr, cir::SyncScopeKind::System); + return emitAtomicCmpXchg(expr, /*returnBool=*/false, cir::MemOrder::Relaxed, + cir::MemOrder::Relaxed, + cir::SyncScopeKind::System); case NVPTX::BI__nvvm_atom_add_gen_f: case NVPTX::BI__nvvm_atom_add_gen_d: cgm.errorNYI(expr->getSourceRange(), @@ -378,12 +365,16 @@ 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: - return makeScopedAtomicCAS(*this, expr, cir::SyncScopeKind::Workgroup); + return emitAtomicCmpXchg(expr, /*returnBool=*/false, cir::MemOrder::Relaxed, + cir::MemOrder::Relaxed, + 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: - return makeScopedAtomicCAS(*this, expr, cir::SyncScopeKind::System); + return emitAtomicCmpXchg(expr, /*returnBool=*/false, cir::MemOrder::Relaxed, + cir::MemOrder::Relaxed, + 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/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h index e738c3b1fb72d..068bf4e92019b 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.h +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h @@ -1750,6 +1750,14 @@ class CIRGenFunction : public CIRGenTypeCache { mlir::Value *emittedArgValue = nullptr, cir::MemOrder ordering = cir::MemOrder::SequentiallyConsistent); + /// Emit `cir.atomic.cmpxchg`. Returns the old value, or the success flag + /// when `returnBool` is true. + mlir::Value emitAtomicCmpXchg( + const clang::CallExpr *expr, bool returnBool, + cir::MemOrder successOrder = cir::MemOrder::SequentiallyConsistent, + cir::MemOrder failureOrder = cir::MemOrder::SequentiallyConsistent, + cir::SyncScopeKind scope = cir::SyncScopeKind::System); + mlir::LogicalResult emitAttributedStmt(const AttributedStmt &s); AutoVarEmission emitAutoVarAlloca(const clang::VarDecl &d, _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
