[clang] [CIR][CUDA] Add support for `__nvvm_atom_add_gen` builtins (PR #210534)
Ayokunle Amodu via cfe-commits
cfe-commits at lists.llvm.org
Sat Jul 18 12:28:41 PDT 2026
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 <ayokunle321 at gmail.com>
Date: Sat, 18 Jul 2026 19:20:58 +0000
Subject: [PATCH 1/2] 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 <ayokunle321 at gmail.com>
Date: Sat, 18 Jul 2026 19:28:25 +0000
Subject: [PATCH 2/2] 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();
More information about the cfe-commits
mailing list