[clang] 88840f2 - [CIR][CUDA] Add support for scoped NVVM atomic builtins (#210863)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Jul 22 04:30:49 PDT 2026
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);
+}
More information about the cfe-commits
mailing list