llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clangir @llvm/pr-subscribers-clang Author: David Rivera (RiverDave) <details> <summary>Changes</summary> Blocker from CUB/Thrust and MiniFE (atomicCAS). --- Full diff: https://github.com/llvm/llvm-project/pull/221261.diff 2 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp (+18-13) - (modified) clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu (+70) ``````````diff diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp index ae994005c588a..adabbd682daea 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp @@ -97,6 +97,21 @@ 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) { @@ -171,11 +186,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 +378,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 `````````` </details> https://github.com/llvm/llvm-project/pull/221261 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
