https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/210534
>From 341eb81eb1ef09fee1130ae7ec7a1e1ae97f3f6a Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Sat, 18 Jul 2026 19:20:58 +0000 Subject: [PATCH 1/4] add support for nvvm atomic add --- clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp | 15 ++++---- clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp | 8 ++-- clang/lib/CIR/CodeGen/CIRGenFunction.h | 6 +++ .../CIR/CodeGenCUDA/builtins-nvvm-atomic.cu | 37 +++++++++++++++++++ 4 files changed, 54 insertions(+), 12 deletions(-) create mode 100644 clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp index 3c48b8b67d9ee..9673b38fc9f5c 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp @@ -163,11 +163,10 @@ static Address checkAtomicAlignment(CIRGenFunction &cgf, const CallExpr *e) { /// Utility to insert an atomic instruction based on Intrinsic::ID /// and the expression node. -static mlir::Value makeBinaryAtomicValue( - CIRGenFunction &cgf, cir::AtomicFetchKind kind, const CallExpr *expr, - mlir::Type *originalArgType = nullptr, - mlir::Value *emittedArgValue = nullptr, - cir::MemOrder ordering = cir::MemOrder::SequentiallyConsistent) { +mlir::Value CIRGenFunction::makeBinaryAtomicValue( + cir::AtomicFetchKind kind, const CallExpr *expr, mlir::Type *originalArgType, + mlir::Value *emittedArgValue, cir::MemOrder ordering) { + CIRGenFunction &cgf = *this; QualType type = expr->getType(); QualType ptrType = expr->getArg(0)->getType(); @@ -224,7 +223,7 @@ static mlir::Value makeBinaryAtomicValue( static RValue emitBinaryAtomic(CIRGenFunction &cgf, cir::AtomicFetchKind atomicOpkind, const CallExpr *e) { - return RValue::get(makeBinaryAtomicValue(cgf, atomicOpkind, e)); + return RValue::get(cgf.makeBinaryAtomicValue(atomicOpkind, e)); } template <typename BinOp> @@ -234,8 +233,8 @@ static RValue emitBinaryAtomicPost(CIRGenFunction &cgf, mlir::Value emittedArgValue; mlir::Type originalArgType; clang::QualType typ = e->getType(); - mlir::Value result = makeBinaryAtomicValue( - cgf, atomicOpkind, e, &originalArgType, &emittedArgValue); + mlir::Value result = cgf.makeBinaryAtomicValue( + atomicOpkind, e, &originalArgType, &emittedArgValue); clang::CIRGen::CIRGenBuilderTy &builder = cgf.getBuilder(); result = BinOp::create(builder, result.getLoc(), result, emittedArgValue); diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp index 4db2d7259c6ba..c8469d3a575b4 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp @@ -39,10 +39,10 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { case NVPTX::BI__nvvm_atom_add_gen_i: case NVPTX::BI__nvvm_atom_add_gen_l: case NVPTX::BI__nvvm_atom_add_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeBinaryAtomicValue(cir::AtomicFetchKind::Add, expr, + /*originalArgType=*/nullptr, + /*emittedArgValue=*/nullptr, + cir::MemOrder::Relaxed); case NVPTX::BI__nvvm_atom_sub_gen_i: case NVPTX::BI__nvvm_atom_sub_gen_l: case NVPTX::BI__nvvm_atom_sub_gen_ll: diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h index 322355fde3957..fa926ef54f7e2 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.h +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h @@ -1648,6 +1648,12 @@ class CIRGenFunction : public CIRGenTypeCache { const Expr *memOrder, bool isStore, bool isLoad, bool isFence, llvm::function_ref<void(cir::MemOrder)> emitAtomicOp); + mlir::Value makeBinaryAtomicValue( + cir::AtomicFetchKind kind, const clang::CallExpr *expr, + mlir::Type *originalArgType = nullptr, + mlir::Value *emittedArgValue = nullptr, + cir::MemOrder ordering = cir::MemOrder::SequentiallyConsistent); + mlir::LogicalResult emitAttributedStmt(const AttributedStmt &s); AutoVarEmission emitAutoVarAlloca(const clang::VarDecl &d, diff --git a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu new file mode 100644 index 0000000000000..f532b3fedf46c --- /dev/null +++ b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu @@ -0,0 +1,37 @@ +#include "Inputs/cuda.h" + +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_80 -x cuda \ +// RUN: -fcuda-is-device -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_80 -x cuda \ +// RUN: -fcuda-is-device -fclangir -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s + +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_80 -x cuda \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// CIR-LABEL: @_Z19test_atom_add_gen_iPii +// CIR: cir.atomic.fetch add relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z19test_atom_add_gen_iPii +// LLVM: atomicrmw add ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_add_gen_i(int *p, int val) { + __nvvm_atom_add_gen_i(p, val); +} + +// CIR-LABEL: @_Z19test_atom_add_gen_lPll +// CIR: cir.atomic.fetch add relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z19test_atom_add_gen_lPll +// LLVM: atomicrmw add ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_add_gen_l(long *p, long val) { + __nvvm_atom_add_gen_l(p, val); +} + +// CIR-LABEL: @_Z20test_atom_add_gen_llPxx +// CIR: cir.atomic.fetch add relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z20test_atom_add_gen_llPxx +// LLVM: atomicrmw add ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_add_gen_ll(long long *p, long long val) { + __nvvm_atom_add_gen_ll(p, val); +} >From ed266e0075ef8d9736f10aa035f28c58093d1fa9 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Sat, 18 Jul 2026 19:28:25 +0000 Subject: [PATCH 2/4] fix style --- clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp index 9673b38fc9f5c..965ac2fa45832 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp @@ -163,9 +163,11 @@ static Address checkAtomicAlignment(CIRGenFunction &cgf, const CallExpr *e) { /// Utility to insert an atomic instruction based on Intrinsic::ID /// and the expression node. -mlir::Value CIRGenFunction::makeBinaryAtomicValue( - cir::AtomicFetchKind kind, const CallExpr *expr, mlir::Type *originalArgType, - mlir::Value *emittedArgValue, cir::MemOrder ordering) { +mlir::Value CIRGenFunction::makeBinaryAtomicValue(cir::AtomicFetchKind kind, + const CallExpr *expr, + mlir::Type *originalArgType, + mlir::Value *emittedArgValue, + cir::MemOrder ordering) { CIRGenFunction &cgf = *this; QualType type = expr->getType(); >From 323e7c758b4ed006791d6c4d1986a1223c0d35a4 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Sat, 18 Jul 2026 20:32:46 +0000 Subject: [PATCH 3/4] add more binary builtins --- clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp | 82 ++++---- .../CIR/CodeGenCUDA/builtins-nvvm-atomic.cu | 192 ++++++++++++++++++ 2 files changed, 234 insertions(+), 40 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp index c8469d3a575b4..7bccede65e221 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp @@ -46,31 +46,31 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { case NVPTX::BI__nvvm_atom_sub_gen_i: case NVPTX::BI__nvvm_atom_sub_gen_l: case NVPTX::BI__nvvm_atom_sub_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeBinaryAtomicValue(cir::AtomicFetchKind::Sub, expr, + /*originalArgType=*/nullptr, + /*emittedArgValue=*/nullptr, + cir::MemOrder::Relaxed); case NVPTX::BI__nvvm_atom_and_gen_i: case NVPTX::BI__nvvm_atom_and_gen_l: case NVPTX::BI__nvvm_atom_and_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeBinaryAtomicValue(cir::AtomicFetchKind::And, expr, + /*originalArgType=*/nullptr, + /*emittedArgValue=*/nullptr, + cir::MemOrder::Relaxed); case NVPTX::BI__nvvm_atom_or_gen_i: case NVPTX::BI__nvvm_atom_or_gen_l: case NVPTX::BI__nvvm_atom_or_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeBinaryAtomicValue(cir::AtomicFetchKind::Or, expr, + /*originalArgType=*/nullptr, + /*emittedArgValue=*/nullptr, + cir::MemOrder::Relaxed); case NVPTX::BI__nvvm_atom_xor_gen_i: case NVPTX::BI__nvvm_atom_xor_gen_l: case NVPTX::BI__nvvm_atom_xor_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeBinaryAtomicValue(cir::AtomicFetchKind::Xor, expr, + /*originalArgType=*/nullptr, + /*emittedArgValue=*/nullptr, + cir::MemOrder::Relaxed); case NVPTX::BI__nvvm_atom_xchg_gen_i: case NVPTX::BI__nvvm_atom_xchg_gen_l: case NVPTX::BI__nvvm_atom_xchg_gen_ll: @@ -81,31 +81,33 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { case NVPTX::BI__nvvm_atom_max_gen_i: case NVPTX::BI__nvvm_atom_max_gen_l: case NVPTX::BI__nvvm_atom_max_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeBinaryAtomicValue(cir::AtomicFetchKind::Max, expr, + /*originalArgType=*/nullptr, + /*emittedArgValue=*/nullptr, + cir::MemOrder::Relaxed); case NVPTX::BI__nvvm_atom_max_gen_ui: case NVPTX::BI__nvvm_atom_max_gen_ul: case NVPTX::BI__nvvm_atom_max_gen_ull: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + // The unsigned operand type makes this lower to `umax`. + return makeBinaryAtomicValue(cir::AtomicFetchKind::Max, expr, + /*originalArgType=*/nullptr, + /*emittedArgValue=*/nullptr, + cir::MemOrder::Relaxed); case NVPTX::BI__nvvm_atom_min_gen_i: case NVPTX::BI__nvvm_atom_min_gen_l: case NVPTX::BI__nvvm_atom_min_gen_ll: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeBinaryAtomicValue(cir::AtomicFetchKind::Min, expr, + /*originalArgType=*/nullptr, + /*emittedArgValue=*/nullptr, + cir::MemOrder::Relaxed); case NVPTX::BI__nvvm_atom_min_gen_ui: case NVPTX::BI__nvvm_atom_min_gen_ul: case NVPTX::BI__nvvm_atom_min_gen_ull: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + // The unsigned operand type makes this lower to `umin`. + return makeBinaryAtomicValue(cir::AtomicFetchKind::Min, expr, + /*originalArgType=*/nullptr, + /*emittedArgValue=*/nullptr, + cir::MemOrder::Relaxed); case NVPTX::BI__nvvm_atom_cas_gen_us: case NVPTX::BI__nvvm_atom_cas_gen_i: case NVPTX::BI__nvvm_atom_cas_gen_l: @@ -122,15 +124,15 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { getContext().BuiltinInfo.getName(builtinId)); return mlir::Value{}; case NVPTX::BI__nvvm_atom_inc_gen_ui: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeBinaryAtomicValue(cir::AtomicFetchKind::UIncWrap, expr, + /*originalArgType=*/nullptr, + /*emittedArgValue=*/nullptr, + cir::MemOrder::Relaxed); case NVPTX::BI__nvvm_atom_dec_gen_ui: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeBinaryAtomicValue(cir::AtomicFetchKind::UDecWrap, expr, + /*originalArgType=*/nullptr, + /*emittedArgValue=*/nullptr, + cir::MemOrder::Relaxed); case NVPTX::BI__nvvm_ldg_c: case NVPTX::BI__nvvm_ldg_sc: case NVPTX::BI__nvvm_ldg_c2: diff --git a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu index f532b3fedf46c..ac9eb60110905 100644 --- a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu +++ b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-atomic.cu @@ -35,3 +35,195 @@ __device__ void test_atom_add_gen_l(long *p, long val) { __device__ void test_atom_add_gen_ll(long long *p, long long val) { __nvvm_atom_add_gen_ll(p, val); } + +// CIR-LABEL: @_Z19test_atom_sub_gen_iPii +// CIR: cir.atomic.fetch sub relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z19test_atom_sub_gen_iPii +// LLVM: atomicrmw sub ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_sub_gen_i(int *p, int val) { + __nvvm_atom_sub_gen_i(p, val); +} + +// CIR-LABEL: @_Z19test_atom_sub_gen_lPll +// CIR: cir.atomic.fetch sub relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z19test_atom_sub_gen_lPll +// LLVM: atomicrmw sub ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sub_gen_l(long *p, long val) { + __nvvm_atom_sub_gen_l(p, val); +} + +// CIR-LABEL: @_Z20test_atom_sub_gen_llPxx +// CIR: cir.atomic.fetch sub relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z20test_atom_sub_gen_llPxx +// LLVM: atomicrmw sub ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_sub_gen_ll(long long *p, long long val) { + __nvvm_atom_sub_gen_ll(p, val); +} + +// CIR-LABEL: @_Z19test_atom_and_gen_iPii +// CIR: cir.atomic.fetch and relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z19test_atom_and_gen_iPii +// LLVM: atomicrmw and ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_and_gen_i(int *p, int val) { + __nvvm_atom_and_gen_i(p, val); +} + +// CIR-LABEL: @_Z19test_atom_and_gen_lPll +// CIR: cir.atomic.fetch and relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z19test_atom_and_gen_lPll +// LLVM: atomicrmw and ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_and_gen_l(long *p, long val) { + __nvvm_atom_and_gen_l(p, val); +} + +// CIR-LABEL: @_Z20test_atom_and_gen_llPxx +// CIR: cir.atomic.fetch and relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z20test_atom_and_gen_llPxx +// LLVM: atomicrmw and ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_and_gen_ll(long long *p, long long val) { + __nvvm_atom_and_gen_ll(p, val); +} + +// CIR-LABEL: @_Z18test_atom_or_gen_iPii +// CIR: cir.atomic.fetch or relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z18test_atom_or_gen_iPii +// LLVM: atomicrmw or ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_or_gen_i(int *p, int val) { + __nvvm_atom_or_gen_i(p, val); +} + +// CIR-LABEL: @_Z18test_atom_or_gen_lPll +// CIR: cir.atomic.fetch or relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z18test_atom_or_gen_lPll +// LLVM: atomicrmw or ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_or_gen_l(long *p, long val) { + __nvvm_atom_or_gen_l(p, val); +} + +// CIR-LABEL: @_Z19test_atom_or_gen_llPxx +// CIR: cir.atomic.fetch or relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z19test_atom_or_gen_llPxx +// LLVM: atomicrmw or ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_or_gen_ll(long long *p, long long val) { + __nvvm_atom_or_gen_ll(p, val); +} + +// CIR-LABEL: @_Z19test_atom_xor_gen_iPii +// CIR: cir.atomic.fetch xor relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z19test_atom_xor_gen_iPii +// LLVM: atomicrmw xor ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_xor_gen_i(int *p, int val) { + __nvvm_atom_xor_gen_i(p, val); +} + +// CIR-LABEL: @_Z19test_atom_xor_gen_lPll +// CIR: cir.atomic.fetch xor relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z19test_atom_xor_gen_lPll +// LLVM: atomicrmw xor ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_xor_gen_l(long *p, long val) { + __nvvm_atom_xor_gen_l(p, val); +} + +// CIR-LABEL: @_Z20test_atom_xor_gen_llPxx +// CIR: cir.atomic.fetch xor relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z20test_atom_xor_gen_llPxx +// LLVM: atomicrmw xor ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_xor_gen_ll(long long *p, long long val) { + __nvvm_atom_xor_gen_ll(p, val); +} + +// CIR-LABEL: @_Z19test_atom_max_gen_iPii +// CIR: cir.atomic.fetch max relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z19test_atom_max_gen_iPii +// LLVM: atomicrmw max ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_max_gen_i(int *p, int val) { + __nvvm_atom_max_gen_i(p, val); +} + +// CIR-LABEL: @_Z19test_atom_max_gen_lPll +// CIR: cir.atomic.fetch max relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z19test_atom_max_gen_lPll +// LLVM: atomicrmw max ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_max_gen_l(long *p, long val) { + __nvvm_atom_max_gen_l(p, val); +} + +// CIR-LABEL: @_Z20test_atom_max_gen_llPxx +// CIR: cir.atomic.fetch max relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z20test_atom_max_gen_llPxx +// LLVM: atomicrmw max ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_max_gen_ll(long long *p, long long val) { + __nvvm_atom_max_gen_ll(p, val); +} + +// CIR-LABEL: @_Z20test_atom_max_gen_uiPjj +// CIR: cir.atomic.fetch max relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u32i>, !u32i) -> !u32i +// LLVM-LABEL: @_Z20test_atom_max_gen_uiPjj +// LLVM: atomicrmw umax ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_max_gen_ui(unsigned *p, unsigned val) { + __nvvm_atom_max_gen_ui(p, val); +} + +// CIR-LABEL: @_Z20test_atom_max_gen_ulPmm +// CIR: cir.atomic.fetch max relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u64i>, !u64i) -> !u64i +// LLVM-LABEL: @_Z20test_atom_max_gen_ulPmm +// LLVM: atomicrmw umax ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_max_gen_ul(unsigned long *p, unsigned long val) { + __nvvm_atom_max_gen_ul(p, val); +} + +// CIR-LABEL: @_Z21test_atom_max_gen_ullPyy +// CIR: cir.atomic.fetch max relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u64i>, !u64i) -> !u64i +// LLVM-LABEL: @_Z21test_atom_max_gen_ullPyy +// LLVM: atomicrmw umax ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_max_gen_ull(unsigned long long *p, unsigned long long val) { + __nvvm_atom_max_gen_ull(p, val); +} + +// CIR-LABEL: @_Z19test_atom_min_gen_iPii +// CIR: cir.atomic.fetch min relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i +// LLVM-LABEL: @_Z19test_atom_min_gen_iPii +// LLVM: atomicrmw min ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_min_gen_i(int *p, int val) { + __nvvm_atom_min_gen_i(p, val); +} + +// CIR-LABEL: @_Z19test_atom_min_gen_lPll +// CIR: cir.atomic.fetch min relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z19test_atom_min_gen_lPll +// LLVM: atomicrmw min ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_min_gen_l(long *p, long val) { + __nvvm_atom_min_gen_l(p, val); +} + +// CIR-LABEL: @_Z20test_atom_min_gen_llPxx +// CIR: cir.atomic.fetch min relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!s64i>, !s64i) -> !s64i +// LLVM-LABEL: @_Z20test_atom_min_gen_llPxx +// LLVM: atomicrmw min ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_min_gen_ll(long long *p, long long val) { + __nvvm_atom_min_gen_ll(p, val); +} + +// CIR-LABEL: @_Z20test_atom_min_gen_uiPjj +// CIR: cir.atomic.fetch min relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u32i>, !u32i) -> !u32i +// LLVM-LABEL: @_Z20test_atom_min_gen_uiPjj +// LLVM: atomicrmw umin ptr %{{.*}}, i32 %{{.*}} monotonic, align 4 +__device__ void test_atom_min_gen_ui(unsigned *p, unsigned val) { + __nvvm_atom_min_gen_ui(p, val); +} + +// CIR-LABEL: @_Z20test_atom_min_gen_ulPmm +// CIR: cir.atomic.fetch min relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u64i>, !u64i) -> !u64i +// LLVM-LABEL: @_Z20test_atom_min_gen_ulPmm +// LLVM: atomicrmw umin ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_min_gen_ul(unsigned long *p, unsigned long val) { + __nvvm_atom_min_gen_ul(p, val); +} + +// CIR-LABEL: @_Z21test_atom_min_gen_ullPyy +// CIR: cir.atomic.fetch min relaxed syncscope(system) fetch_first %{{.*}}, %{{.*}} : (!cir.ptr<!u64i>, !u64i) -> !u64i +// LLVM-LABEL: @_Z21test_atom_min_gen_ullPyy +// LLVM: atomicrmw umin ptr %{{.*}}, i64 %{{.*}} monotonic, align 8 +__device__ void test_atom_min_gen_ull(unsigned long long *p, unsigned long long val) { + __nvvm_atom_min_gen_ull(p, val); +} >From 58e2dca482b5db47f233c649dd13a08d799a1e51 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Sat, 18 Jul 2026 20:35:37 +0000 Subject: [PATCH 4/4] remove uneccessary 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 7bccede65e221..86a7bbd5026f6 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp @@ -88,7 +88,6 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { case NVPTX::BI__nvvm_atom_max_gen_ui: case NVPTX::BI__nvvm_atom_max_gen_ul: case NVPTX::BI__nvvm_atom_max_gen_ull: - // The unsigned operand type makes this lower to `umax`. return makeBinaryAtomicValue(cir::AtomicFetchKind::Max, expr, /*originalArgType=*/nullptr, /*emittedArgValue=*/nullptr, @@ -103,7 +102,6 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { case NVPTX::BI__nvvm_atom_min_gen_ui: case NVPTX::BI__nvvm_atom_min_gen_ul: case NVPTX::BI__nvvm_atom_min_gen_ull: - // The unsigned operand type makes this lower to `umin`. return makeBinaryAtomicValue(cir::AtomicFetchKind::Min, expr, /*originalArgType=*/nullptr, /*emittedArgValue=*/nullptr, _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
