Author: Ayokunle Amodu Date: 2026-07-22T07:30:44-04:00 New Revision: 88840f264de6e0e6548219429ff62b3354db5210
URL: https://github.com/llvm/llvm-project/commit/88840f264de6e0e6548219429ff62b3354db5210 DIFF: https://github.com/llvm/llvm-project/commit/88840f264de6e0e6548219429ff62b3354db5210.diff LOG: [CIR][CUDA] Add support for scoped NVVM atomic builtins (#210863) Adds codegen support for the block and system-scoped NVVM atomic builtins: add, and, or, xor, min, max, inc and dec. These are lowered to the corresponding CIR `cir.atomic.fetch` operations and subsequently lowered to LLVM `atomicrmw` instructions. Added: Modified: clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.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 86a7bbd5026f6..e1e829fc2ab22 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp @@ -33,6 +33,21 @@ static mlir::Value emitUnaryNVVMIntrinsic(CIRGenFunction &cgf, .getResult(); } +static mlir::Value makeScopedAtomicRMW(CIRGenFunction &cgf, + const CallExpr *expr, + cir::AtomicFetchKind kind, + 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 rmwi = cir::AtomicFetchOp::create( + builder, cgf.getLoc(expr->getSourceRange()), destValue, val, kind, + cir::MemOrder::Relaxed, scope, /*is_volatile=*/false, + /*fetch_first=*/true); + return rmwi->getResult(0); +} + std::optional<mlir::Value> CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { switch (builtinId) { @@ -214,29 +229,21 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { case NVPTX::BI__nvvm_atom_cta_add_gen_i: case NVPTX::BI__nvvm_atom_cta_add_gen_l: case NVPTX::BI__nvvm_atom_cta_add_gen_ll: - 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::Workgroup); case NVPTX::BI__nvvm_atom_sys_add_gen_i: case NVPTX::BI__nvvm_atom_sys_add_gen_l: case NVPTX::BI__nvvm_atom_sys_add_gen_ll: - 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_cta_add_gen_f: case NVPTX::BI__nvvm_atom_cta_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::Workgroup); case NVPTX::BI__nvvm_atom_sys_add_gen_f: case NVPTX::BI__nvvm_atom_sys_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_cta_xchg_gen_i: case NVPTX::BI__nvvm_atom_cta_xchg_gen_l: case NVPTX::BI__nvvm_atom_cta_xchg_gen_ll: @@ -257,102 +264,74 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { case NVPTX::BI__nvvm_atom_cta_max_gen_ul: case NVPTX::BI__nvvm_atom_cta_max_gen_ll: case NVPTX::BI__nvvm_atom_cta_max_gen_ull: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicRMW(*this, expr, cir::AtomicFetchKind::Max, + cir::SyncScopeKind::Workgroup); case NVPTX::BI__nvvm_atom_sys_max_gen_i: case NVPTX::BI__nvvm_atom_sys_max_gen_ui: case NVPTX::BI__nvvm_atom_sys_max_gen_l: case NVPTX::BI__nvvm_atom_sys_max_gen_ul: case NVPTX::BI__nvvm_atom_sys_max_gen_ll: case NVPTX::BI__nvvm_atom_sys_max_gen_ull: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicRMW(*this, expr, cir::AtomicFetchKind::Max, + cir::SyncScopeKind::System); case NVPTX::BI__nvvm_atom_cta_min_gen_i: case NVPTX::BI__nvvm_atom_cta_min_gen_ui: case NVPTX::BI__nvvm_atom_cta_min_gen_l: case NVPTX::BI__nvvm_atom_cta_min_gen_ul: case NVPTX::BI__nvvm_atom_cta_min_gen_ll: case NVPTX::BI__nvvm_atom_cta_min_gen_ull: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicRMW(*this, expr, cir::AtomicFetchKind::Min, + cir::SyncScopeKind::Workgroup); case NVPTX::BI__nvvm_atom_sys_min_gen_i: case NVPTX::BI__nvvm_atom_sys_min_gen_ui: case NVPTX::BI__nvvm_atom_sys_min_gen_l: case NVPTX::BI__nvvm_atom_sys_min_gen_ul: case NVPTX::BI__nvvm_atom_sys_min_gen_ll: case NVPTX::BI__nvvm_atom_sys_min_gen_ull: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicRMW(*this, expr, cir::AtomicFetchKind::Min, + cir::SyncScopeKind::System); case NVPTX::BI__nvvm_atom_cta_inc_gen_ui: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicRMW(*this, expr, cir::AtomicFetchKind::UIncWrap, + cir::SyncScopeKind::Workgroup); case NVPTX::BI__nvvm_atom_cta_dec_gen_ui: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicRMW(*this, expr, cir::AtomicFetchKind::UDecWrap, + cir::SyncScopeKind::Workgroup); case NVPTX::BI__nvvm_atom_sys_inc_gen_ui: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicRMW(*this, expr, cir::AtomicFetchKind::UIncWrap, + cir::SyncScopeKind::System); case NVPTX::BI__nvvm_atom_sys_dec_gen_ui: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicRMW(*this, expr, cir::AtomicFetchKind::UDecWrap, + cir::SyncScopeKind::System); case NVPTX::BI__nvvm_atom_cta_and_gen_i: case NVPTX::BI__nvvm_atom_cta_and_gen_l: case NVPTX::BI__nvvm_atom_cta_and_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicRMW(*this, expr, cir::AtomicFetchKind::And, + cir::SyncScopeKind::Workgroup); case NVPTX::BI__nvvm_atom_sys_and_gen_i: case NVPTX::BI__nvvm_atom_sys_and_gen_l: case NVPTX::BI__nvvm_atom_sys_and_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicRMW(*this, expr, cir::AtomicFetchKind::And, + cir::SyncScopeKind::System); case NVPTX::BI__nvvm_atom_cta_or_gen_i: case NVPTX::BI__nvvm_atom_cta_or_gen_l: case NVPTX::BI__nvvm_atom_cta_or_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicRMW(*this, expr, cir::AtomicFetchKind::Or, + cir::SyncScopeKind::Workgroup); case NVPTX::BI__nvvm_atom_sys_or_gen_i: case NVPTX::BI__nvvm_atom_sys_or_gen_l: case NVPTX::BI__nvvm_atom_sys_or_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicRMW(*this, expr, cir::AtomicFetchKind::Or, + cir::SyncScopeKind::System); case NVPTX::BI__nvvm_atom_cta_xor_gen_i: case NVPTX::BI__nvvm_atom_cta_xor_gen_l: case NVPTX::BI__nvvm_atom_cta_xor_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicRMW(*this, expr, cir::AtomicFetchKind::Xor, + cir::SyncScopeKind::Workgroup); case NVPTX::BI__nvvm_atom_sys_xor_gen_i: case NVPTX::BI__nvvm_atom_sys_xor_gen_l: case NVPTX::BI__nvvm_atom_sys_xor_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeScopedAtomicRMW(*this, expr, cir::AtomicFetchKind::Xor, + cir::SyncScopeKind::System); case NVPTX::BI__nvvm_atom_cta_cas_gen_us: case NVPTX::BI__nvvm_atom_cta_cas_gen_i: case NVPTX::BI__nvvm_atom_cta_cas_gen_l: diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp index 806e3235b6a8e..94c5a01842613 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp @@ -30,6 +30,13 @@ class NVPTXTargetLoweringInfo : public TargetLoweringInfo { "Unknown CIR address space for NVPTX target"); return NVPTXAddrSpaceMap[idx]; } + + cir::SyncScopeKind + convertSyncScope(cir::SyncScopeKind syncScope) const override { + if (syncScope == cir::SyncScopeKind::Workgroup) + return cir::SyncScopeKind::Workgroup; + return cir::SyncScopeKind::System; + } }; } // namespace diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index edbf9e3883399..f3f77a9151c35 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -1017,7 +1017,14 @@ getLLVMMemOrder(std::optional<cir::MemOrder> memorder) { } static llvm::StringRef getLLVMSyncScope(cir::SyncScopeKind syncScope) { - return syncScope == cir::SyncScopeKind::SingleThread ? "singlethread" : ""; + switch (syncScope) { + case cir::SyncScopeKind::SingleThread: + return "singlethread"; + case cir::SyncScopeKind::Workgroup: + return "block"; + default: + return ""; + } } static std::optional<llvm::StringRef> diff --git a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu index ac9eb60110905..a5ed80759d2ad 100644 --- a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu +++ b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu @@ -227,3 +227,451 @@ __device__ void test_atom_min_gen_ul(unsigned long *p, unsigned long val) { __device__ void test_atom_min_gen_ull(unsigned long long *p, unsigned long long val) { __nvvm_atom_min_gen_ull(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 +// LLVM: atomicrmw add ptr %{{.*}}, i32 %{{.*}} syncscope("block") monotonic, align 4 +__device__ void test_atom_cta_add_gen_i(int *p, int val) { + __nvvm_atom_cta_add_gen_i(p, val); +} + +// CIR-LABEL: @_Z23test_atom_sys_add_gen_iPii +// CIR: cir.atomic.fetch add relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z23test_atom_sys_add_gen_iPii +// LLVM: atomicrmw add ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_sys_add_gen_i(int *p, int val) { + __nvvm_atom_sys_add_gen_i(p, val); +} + +// CIR-LABEL: @_Z23test_atom_cta_add_gen_lPll +// CIR: cir.atomic.fetch add relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z23test_atom_cta_add_gen_lPll +// LLVM: atomicrmw add ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_add_gen_l(long *p, long val) { + __nvvm_atom_cta_add_gen_l(p, val); +} + +// CIR-LABEL: @_Z23test_atom_sys_add_gen_lPll +// CIR: cir.atomic.fetch add relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z23test_atom_sys_add_gen_lPll +// LLVM: atomicrmw add ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_add_gen_l(long *p, long val) { + __nvvm_atom_sys_add_gen_l(p, val); +} + +// CIR-LABEL: @_Z24test_atom_cta_add_gen_llPxx +// CIR: cir.atomic.fetch add relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z24test_atom_cta_add_gen_llPxx +// LLVM: atomicrmw add ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_add_gen_ll(long long *p, long long val) { + __nvvm_atom_cta_add_gen_ll(p, val); +} + +// CIR-LABEL: @_Z24test_atom_sys_add_gen_llPxx +// CIR: cir.atomic.fetch add relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z24test_atom_sys_add_gen_llPxx +// LLVM: atomicrmw add ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_add_gen_ll(long long *p, long long val) { + __nvvm_atom_sys_add_gen_ll(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 +// LLVM: atomicrmw fadd ptr %{{.*}}, float %{{.*}} syncscope("block") monotonic, align 4 +__device__ void test_atom_cta_add_gen_f(float *p, float val) { + __nvvm_atom_cta_add_gen_f(p, val); +} + +// CIR-LABEL: @_Z23test_atom_sys_add_gen_fPff +// CIR: cir.atomic.fetch add relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!cir.float>, !cir.float) -> !cir.float +// LLVM-LABEL: @_Z23test_atom_sys_add_gen_fPff +// LLVM: atomicrmw fadd ptr %{{.*}}, float %{{.*}} monotonic, align 4 +__device__ void test_atom_sys_add_gen_f(float *p, float val) { + __nvvm_atom_sys_add_gen_f(p, val); +} + +// CIR-LABEL: @_Z23test_atom_cta_add_gen_dPdd +// CIR: cir.atomic.fetch add relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!cir.double>, !cir.double) -> !cir.double +// LLVM-LABEL: @_Z23test_atom_cta_add_gen_dPdd +// LLVM: atomicrmw fadd ptr %{{.*}}, double %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_add_gen_d(double *p, double val) { + __nvvm_atom_cta_add_gen_d(p, val); +} + +// CIR-LABEL: @_Z23test_atom_sys_add_gen_dPdd +// CIR: cir.atomic.fetch add relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!cir.double>, !cir.double) -> !cir.double +// LLVM-LABEL: @_Z23test_atom_sys_add_gen_dPdd +// LLVM: atomicrmw fadd ptr %{{.*}}, double %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_add_gen_d(double *p, double val) { + __nvvm_atom_sys_add_gen_d(p, val); +} + +// CIR-LABEL: @_Z23test_atom_cta_max_gen_iPii +// CIR: cir.atomic.fetch max relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z23test_atom_cta_max_gen_iPii +// LLVM: atomicrmw max ptr %{{.*}}, i32 %{{.*}} syncscope("block") monotonic, align 4 +__device__ void test_atom_cta_max_gen_i(int *p, int val) { + __nvvm_atom_cta_max_gen_i(p, val); +} + +// CIR-LABEL: @_Z23test_atom_sys_max_gen_iPii +// CIR: cir.atomic.fetch max relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z23test_atom_sys_max_gen_iPii +// LLVM: atomicrmw max ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_sys_max_gen_i(int *p, int val) { + __nvvm_atom_sys_max_gen_i(p, val); +} + +// CIR-LABEL: @_Z24test_atom_cta_max_gen_uiPjj +// CIR: cir.atomic.fetch max relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u32i>, !u32i) -> !u32i +// LLVM-LABEL: @_Z24test_atom_cta_max_gen_uiPjj +// LLVM: atomicrmw umax ptr %{{.*}}, i32 %{{.*}} syncscope("block") monotonic, align 4 +__device__ void test_atom_cta_max_gen_ui(unsigned *p, unsigned val) { + __nvvm_atom_cta_max_gen_ui(p, val); +} + +// CIR-LABEL: @_Z24test_atom_sys_max_gen_uiPjj +// CIR: cir.atomic.fetch max relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u32i>, !u32i) -> !u32i +// LLVM-LABEL: @_Z24test_atom_sys_max_gen_uiPjj +// LLVM: atomicrmw umax ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_sys_max_gen_ui(unsigned *p, unsigned val) { + __nvvm_atom_sys_max_gen_ui(p, val); +} + +// CIR-LABEL: @_Z23test_atom_cta_max_gen_lPll +// CIR: cir.atomic.fetch max relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z23test_atom_cta_max_gen_lPll +// LLVM: atomicrmw max ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_max_gen_l(long *p, long val) { + __nvvm_atom_cta_max_gen_l(p, val); +} + +// CIR-LABEL: @_Z23test_atom_sys_max_gen_lPll +// CIR: cir.atomic.fetch max relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z23test_atom_sys_max_gen_lPll +// LLVM: atomicrmw max ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_max_gen_l(long *p, long val) { + __nvvm_atom_sys_max_gen_l(p, val); +} + +// CIR-LABEL: @_Z24test_atom_cta_max_gen_ulPmm +// CIR: cir.atomic.fetch max relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u64i>, !u64i) -> !u64i +// LLVM-LABEL: @_Z24test_atom_cta_max_gen_ulPmm +// LLVM: atomicrmw umax ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_max_gen_ul(unsigned long *p, unsigned long val) { + __nvvm_atom_cta_max_gen_ul(p, val); +} + +// CIR-LABEL: @_Z24test_atom_sys_max_gen_ulPmm +// CIR: cir.atomic.fetch max relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u64i>, !u64i) -> !u64i +// LLVM-LABEL: @_Z24test_atom_sys_max_gen_ulPmm +// LLVM: atomicrmw umax ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_max_gen_ul(unsigned long *p, unsigned long val) { + __nvvm_atom_sys_max_gen_ul(p, val); +} + +// CIR-LABEL: @_Z24test_atom_cta_max_gen_llPxx +// CIR: cir.atomic.fetch max relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z24test_atom_cta_max_gen_llPxx +// LLVM: atomicrmw max ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_max_gen_ll(long long *p, long long val) { + __nvvm_atom_cta_max_gen_ll(p, val); +} + +// CIR-LABEL: @_Z24test_atom_sys_max_gen_llPxx +// CIR: cir.atomic.fetch max relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z24test_atom_sys_max_gen_llPxx +// LLVM: atomicrmw max ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_max_gen_ll(long long *p, long long val) { + __nvvm_atom_sys_max_gen_ll(p, val); +} + +// CIR-LABEL: @_Z25test_atom_cta_max_gen_ullPyy +// CIR: cir.atomic.fetch max relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u64i>, !u64i) -> !u64i +// LLVM-LABEL: @_Z25test_atom_cta_max_gen_ullPyy +// LLVM: atomicrmw umax ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_max_gen_ull(unsigned long long *p, unsigned long long val) { + __nvvm_atom_cta_max_gen_ull(p, val); +} + +// CIR-LABEL: @_Z25test_atom_sys_max_gen_ullPyy +// CIR: cir.atomic.fetch max relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u64i>, !u64i) -> !u64i +// LLVM-LABEL: @_Z25test_atom_sys_max_gen_ullPyy +// LLVM: atomicrmw umax ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_max_gen_ull(unsigned long long *p, unsigned long long val) { + __nvvm_atom_sys_max_gen_ull(p, val); +} + +// CIR-LABEL: @_Z23test_atom_cta_min_gen_iPii +// CIR: cir.atomic.fetch min relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z23test_atom_cta_min_gen_iPii +// LLVM: atomicrmw min ptr %{{.*}}, i32 %{{.*}} syncscope("block") monotonic, align 4 +__device__ void test_atom_cta_min_gen_i(int *p, int val) { + __nvvm_atom_cta_min_gen_i(p, val); +} + +// CIR-LABEL: @_Z23test_atom_sys_min_gen_iPii +// CIR: cir.atomic.fetch min relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z23test_atom_sys_min_gen_iPii +// LLVM: atomicrmw min ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_sys_min_gen_i(int *p, int val) { + __nvvm_atom_sys_min_gen_i(p, val); +} + +// CIR-LABEL: @_Z24test_atom_cta_min_gen_uiPjj +// CIR: cir.atomic.fetch min relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u32i>, !u32i) -> !u32i +// LLVM-LABEL: @_Z24test_atom_cta_min_gen_uiPjj +// LLVM: atomicrmw umin ptr %{{.*}}, i32 %{{.*}} syncscope("block") monotonic, align 4 +__device__ void test_atom_cta_min_gen_ui(unsigned *p, unsigned val) { + __nvvm_atom_cta_min_gen_ui(p, val); +} + +// CIR-LABEL: @_Z24test_atom_sys_min_gen_uiPjj +// CIR: cir.atomic.fetch min relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u32i>, !u32i) -> !u32i +// LLVM-LABEL: @_Z24test_atom_sys_min_gen_uiPjj +// LLVM: atomicrmw umin ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_sys_min_gen_ui(unsigned *p, unsigned val) { + __nvvm_atom_sys_min_gen_ui(p, val); +} + +// CIR-LABEL: @_Z23test_atom_cta_min_gen_lPll +// CIR: cir.atomic.fetch min relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z23test_atom_cta_min_gen_lPll +// LLVM: atomicrmw min ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_min_gen_l(long *p, long val) { + __nvvm_atom_cta_min_gen_l(p, val); +} + +// CIR-LABEL: @_Z23test_atom_sys_min_gen_lPll +// CIR: cir.atomic.fetch min relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z23test_atom_sys_min_gen_lPll +// LLVM: atomicrmw min ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_min_gen_l(long *p, long val) { + __nvvm_atom_sys_min_gen_l(p, val); +} + +// CIR-LABEL: @_Z24test_atom_cta_min_gen_ulPmm +// CIR: cir.atomic.fetch min relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u64i>, !u64i) -> !u64i +// LLVM-LABEL: @_Z24test_atom_cta_min_gen_ulPmm +// LLVM: atomicrmw umin ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_min_gen_ul(unsigned long *p, unsigned long val) { + __nvvm_atom_cta_min_gen_ul(p, val); +} + +// CIR-LABEL: @_Z24test_atom_sys_min_gen_ulPmm +// CIR: cir.atomic.fetch min relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u64i>, !u64i) -> !u64i +// LLVM-LABEL: @_Z24test_atom_sys_min_gen_ulPmm +// LLVM: atomicrmw umin ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_min_gen_ul(unsigned long *p, unsigned long val) { + __nvvm_atom_sys_min_gen_ul(p, val); +} + +// CIR-LABEL: @_Z24test_atom_cta_min_gen_llPxx +// CIR: cir.atomic.fetch min relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z24test_atom_cta_min_gen_llPxx +// LLVM: atomicrmw min ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_min_gen_ll(long long *p, long long val) { + __nvvm_atom_cta_min_gen_ll(p, val); +} + +// CIR-LABEL: @_Z24test_atom_sys_min_gen_llPxx +// CIR: cir.atomic.fetch min relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z24test_atom_sys_min_gen_llPxx +// LLVM: atomicrmw min ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_min_gen_ll(long long *p, long long val) { + __nvvm_atom_sys_min_gen_ll(p, val); +} + +// CIR-LABEL: @_Z25test_atom_cta_min_gen_ullPyy +// CIR: cir.atomic.fetch min relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u64i>, !u64i) -> !u64i +// LLVM-LABEL: @_Z25test_atom_cta_min_gen_ullPyy +// LLVM: atomicrmw umin ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_min_gen_ull(unsigned long long *p, unsigned long long val) { + __nvvm_atom_cta_min_gen_ull(p, val); +} + +// CIR-LABEL: @_Z25test_atom_sys_min_gen_ullPyy +// CIR: cir.atomic.fetch min relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u64i>, !u64i) -> !u64i +// LLVM-LABEL: @_Z25test_atom_sys_min_gen_ullPyy +// LLVM: atomicrmw umin ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_min_gen_ull(unsigned long long *p, unsigned long long val) { + __nvvm_atom_sys_min_gen_ull(p, val); +} + +// CIR-LABEL: @_Z24test_atom_cta_inc_gen_uiPjj +// CIR: cir.atomic.fetch uinc_wrap relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u32i>, !u32i) -> !u32i +// LLVM-LABEL: @_Z24test_atom_cta_inc_gen_uiPjj +// LLVM: atomicrmw uinc_wrap ptr %{{.*}}, i32 %{{.*}} syncscope("block") monotonic, align 4 +__device__ void test_atom_cta_inc_gen_ui(unsigned *p, unsigned val) { + __nvvm_atom_cta_inc_gen_ui(p, val); +} + +// CIR-LABEL: @_Z24test_atom_sys_inc_gen_uiPjj +// CIR: cir.atomic.fetch uinc_wrap relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u32i>, !u32i) -> !u32i +// LLVM-LABEL: @_Z24test_atom_sys_inc_gen_uiPjj +// LLVM: atomicrmw uinc_wrap ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_sys_inc_gen_ui(unsigned *p, unsigned val) { + __nvvm_atom_sys_inc_gen_ui(p, val); +} + +// CIR-LABEL: @_Z24test_atom_cta_dec_gen_uiPjj +// CIR: cir.atomic.fetch udec_wrap relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u32i>, !u32i) -> !u32i +// LLVM-LABEL: @_Z24test_atom_cta_dec_gen_uiPjj +// LLVM: atomicrmw udec_wrap ptr %{{.*}}, i32 %{{.*}} syncscope("block") monotonic, align 4 +__device__ void test_atom_cta_dec_gen_ui(unsigned *p, unsigned val) { + __nvvm_atom_cta_dec_gen_ui(p, val); +} + +// CIR-LABEL: @_Z24test_atom_sys_dec_gen_uiPjj +// CIR: cir.atomic.fetch udec_wrap relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u32i>, !u32i) -> !u32i +// LLVM-LABEL: @_Z24test_atom_sys_dec_gen_uiPjj +// LLVM: atomicrmw udec_wrap ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_sys_dec_gen_ui(unsigned *p, unsigned val) { + __nvvm_atom_sys_dec_gen_ui(p, val); +} + +// CIR-LABEL: @_Z23test_atom_cta_and_gen_iPii +// CIR: cir.atomic.fetch and relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z23test_atom_cta_and_gen_iPii +// LLVM: atomicrmw and ptr %{{.*}}, i32 %{{.*}} syncscope("block") monotonic, align 4 +__device__ void test_atom_cta_and_gen_i(int *p, int val) { + __nvvm_atom_cta_and_gen_i(p, val); +} + +// CIR-LABEL: @_Z23test_atom_sys_and_gen_iPii +// CIR: cir.atomic.fetch and relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z23test_atom_sys_and_gen_iPii +// LLVM: atomicrmw and ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_sys_and_gen_i(int *p, int val) { + __nvvm_atom_sys_and_gen_i(p, val); +} + +// CIR-LABEL: @_Z23test_atom_cta_and_gen_lPll +// CIR: cir.atomic.fetch and relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z23test_atom_cta_and_gen_lPll +// LLVM: atomicrmw and ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_and_gen_l(long *p, long val) { + __nvvm_atom_cta_and_gen_l(p, val); +} + +// CIR-LABEL: @_Z23test_atom_sys_and_gen_lPll +// CIR: cir.atomic.fetch and relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z23test_atom_sys_and_gen_lPll +// LLVM: atomicrmw and ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_and_gen_l(long *p, long val) { + __nvvm_atom_sys_and_gen_l(p, val); +} + +// CIR-LABEL: @_Z24test_atom_cta_and_gen_llPxx +// CIR: cir.atomic.fetch and relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z24test_atom_cta_and_gen_llPxx +// LLVM: atomicrmw and ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_and_gen_ll(long long *p, long long val) { + __nvvm_atom_cta_and_gen_ll(p, val); +} + +// CIR-LABEL: @_Z24test_atom_sys_and_gen_llPxx +// CIR: cir.atomic.fetch and relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z24test_atom_sys_and_gen_llPxx +// LLVM: atomicrmw and ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_and_gen_ll(long long *p, long long val) { + __nvvm_atom_sys_and_gen_ll(p, val); +} + +// CIR-LABEL: @_Z22test_atom_cta_or_gen_iPii +// CIR: cir.atomic.fetch or relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z22test_atom_cta_or_gen_iPii +// LLVM: atomicrmw or ptr %{{.*}}, i32 %{{.*}} syncscope("block") monotonic, align 4 +__device__ void test_atom_cta_or_gen_i(int *p, int val) { + __nvvm_atom_cta_or_gen_i(p, val); +} + +// CIR-LABEL: @_Z22test_atom_sys_or_gen_iPii +// CIR: cir.atomic.fetch or relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z22test_atom_sys_or_gen_iPii +// LLVM: atomicrmw or ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_sys_or_gen_i(int *p, int val) { + __nvvm_atom_sys_or_gen_i(p, val); +} + +// CIR-LABEL: @_Z22test_atom_cta_or_gen_lPll +// CIR: cir.atomic.fetch or relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z22test_atom_cta_or_gen_lPll +// LLVM: atomicrmw or ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_or_gen_l(long *p, long val) { + __nvvm_atom_cta_or_gen_l(p, val); +} + +// CIR-LABEL: @_Z22test_atom_sys_or_gen_lPll +// CIR: cir.atomic.fetch or relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z22test_atom_sys_or_gen_lPll +// LLVM: atomicrmw or ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_or_gen_l(long *p, long val) { + __nvvm_atom_sys_or_gen_l(p, val); +} + +// CIR-LABEL: @_Z23test_atom_cta_or_gen_llPxx +// CIR: cir.atomic.fetch or relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z23test_atom_cta_or_gen_llPxx +// LLVM: atomicrmw or ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_or_gen_ll(long long *p, long long val) { + __nvvm_atom_cta_or_gen_ll(p, val); +} + +// CIR-LABEL: @_Z23test_atom_sys_or_gen_llPxx +// CIR: cir.atomic.fetch or relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z23test_atom_sys_or_gen_llPxx +// LLVM: atomicrmw or ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_or_gen_ll(long long *p, long long val) { + __nvvm_atom_sys_or_gen_ll(p, val); +} + +// CIR-LABEL: @_Z23test_atom_cta_xor_gen_iPii +// CIR: cir.atomic.fetch xor relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z23test_atom_cta_xor_gen_iPii +// LLVM: atomicrmw xor ptr %{{.*}}, i32 %{{.*}} syncscope("block") monotonic, align 4 +__device__ void test_atom_cta_xor_gen_i(int *p, int val) { + __nvvm_atom_cta_xor_gen_i(p, val); +} + +// CIR-LABEL: @_Z23test_atom_sys_xor_gen_iPii +// CIR: cir.atomic.fetch xor relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z23test_atom_sys_xor_gen_iPii +// LLVM: atomicrmw xor ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_sys_xor_gen_i(int *p, int val) { + __nvvm_atom_sys_xor_gen_i(p, val); +} + +// CIR-LABEL: @_Z23test_atom_cta_xor_gen_lPll +// CIR: cir.atomic.fetch xor relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z23test_atom_cta_xor_gen_lPll +// LLVM: atomicrmw xor ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_xor_gen_l(long *p, long val) { + __nvvm_atom_cta_xor_gen_l(p, val); +} + +// CIR-LABEL: @_Z23test_atom_sys_xor_gen_lPll +// CIR: cir.atomic.fetch xor relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z23test_atom_sys_xor_gen_lPll +// LLVM: atomicrmw xor ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_xor_gen_l(long *p, long val) { + __nvvm_atom_sys_xor_gen_l(p, val); +} + +// CIR-LABEL: @_Z24test_atom_cta_xor_gen_llPxx +// CIR: cir.atomic.fetch xor relaxed syncscope(workgroup) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z24test_atom_cta_xor_gen_llPxx +// LLVM: atomicrmw xor ptr %{{.*}}, i64 %{{.*}} syncscope("block") monotonic, align 8 +__device__ void test_atom_cta_xor_gen_ll(long long *p, long long val) { + __nvvm_atom_cta_xor_gen_ll(p, val); +} + +// CIR-LABEL: @_Z24test_atom_sys_xor_gen_llPxx +// CIR: cir.atomic.fetch xor relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z24test_atom_sys_xor_gen_llPxx +// LLVM: atomicrmw xor ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sys_xor_gen_ll(long long *p, long long val) { + __nvvm_atom_sys_xor_gen_ll(p, val); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
