[clang] [CIR][CUDA] Add support for NVVM ldu builtins (PR #211843)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 24 09:27:34 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Ayokunle Amodu (ayokunle321)
<details>
<summary>Changes</summary>
Adds CIR codegen support for the NVVM global uniform load builtins `nvvm_ldu_*` ,including the integer, floating-point, and vector variants.
These are lowered to the corresponding `llvm.nvvm.ldu.global.*` intrinsics based on the pointee type.
---
Full diff: https://github.com/llvm/llvm-project/pull/211843.diff
2 Files Affected:
- (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp (+19-12)
- (added) clang/test/CIR/CodeGenCUDA/builtins-nvvm-ldu.cu (+155)
``````````diff
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
index e1e829fc2ab22..11c4abc0b3ba5 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
@@ -20,6 +20,22 @@
using namespace clang;
using namespace clang::CIRGen;
+static mlir::Value makeLdu(CIRGenFunction &cgf, const CallExpr *expr,
+ llvm::StringRef intrinsicName) {
+ auto &builder = cgf.getBuilder();
+ Address ptr = cgf.emitPointerWithAlignment(expr->getArg(0));
+ QualType argType = expr->getArg(0)->getType();
+ mlir::Type elemTy = cgf.convertTypeForMem(argType->getPointeeType());
+ clang::CharUnits align = ptr.getAlignment();
+ mlir::Location loc = cgf.getLoc(expr->getExprLoc());
+ mlir::Value alignVal =
+ builder.getConstantInt(loc, builder.getSInt32Ty(), align.getQuantity());
+ return cir::LLVMIntrinsicCallOp::create(
+ builder, loc, builder.getStringAttr(intrinsicName), elemTy,
+ {ptr.emitRawPointer(), alignVal})
+ .getResult();
+}
+
/// Emit a CIR LLVMIntrinsicCallOp for a unary NVVM intrinsic.
/// The result type is inferred from the single argument.
static mlir::Value emitUnaryNVVMIntrinsic(CIRGenFunction &cgf,
@@ -213,19 +229,13 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) {
case NVPTX::BI__nvvm_ldu_ul2:
case NVPTX::BI__nvvm_ldu_ull:
case NVPTX::BI__nvvm_ldu_ull2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return makeLdu(*this, expr, "nvvm.ldu.global.i");
case NVPTX::BI__nvvm_ldu_f:
case NVPTX::BI__nvvm_ldu_f2:
case NVPTX::BI__nvvm_ldu_f4:
case NVPTX::BI__nvvm_ldu_d:
case NVPTX::BI__nvvm_ldu_d2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return makeLdu(*this, expr, "nvvm.ldu.global.f");
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:
@@ -788,10 +798,7 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) {
return mlir::Value{};
case NVPTX::BI__nvvm_ldu_h:
case NVPTX::BI__nvvm_ldu_h2:
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented NVPTX builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ return makeLdu(*this, expr, "nvvm.ldu.global.f");
case NVPTX::BI__nvvm_cp_async_ca_shared_global_4:
cgm.errorNYI(expr->getSourceRange(),
std::string("unimplemented NVPTX builtin call: ") +
diff --git a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-ldu.cu b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-ldu.cu
new file mode 100644
index 0000000000000..cee0038d770da
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-ldu.cu
@@ -0,0 +1,155 @@
+#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
+
+// FIXME: CIR doesn't propagate the 'contract' fast-math flag to LLVM IR calls
+// yet, so the floating-point LLVM check lines use {{.*}} to tolerate the
+// difference between CIR (no flags) and classic codegen ('contract').
+
+typedef char char2 __attribute__((ext_vector_type(2)));
+typedef unsigned char uchar2 __attribute__((ext_vector_type(2)));
+typedef signed char schar2 __attribute__((ext_vector_type(2)));
+typedef char char4 __attribute__((ext_vector_type(4)));
+typedef unsigned char uchar4 __attribute__((ext_vector_type(4)));
+typedef signed char schar4 __attribute__((ext_vector_type(4)));
+typedef short short2 __attribute__((ext_vector_type(2)));
+typedef unsigned short ushort2 __attribute__((ext_vector_type(2)));
+typedef short short4 __attribute__((ext_vector_type(4)));
+typedef unsigned short ushort4 __attribute__((ext_vector_type(4)));
+typedef int int2 __attribute__((ext_vector_type(2)));
+typedef unsigned int uint2 __attribute__((ext_vector_type(2)));
+typedef int int4 __attribute__((ext_vector_type(4)));
+typedef unsigned int uint4 __attribute__((ext_vector_type(4)));
+typedef long long2 __attribute__((ext_vector_type(2)));
+typedef unsigned long ulong2 __attribute__((ext_vector_type(2)));
+typedef long long longlong2 __attribute__((ext_vector_type(2)));
+typedef unsigned long long ulonglong2 __attribute__((ext_vector_type(2)));
+typedef float float2 __attribute__((ext_vector_type(2)));
+typedef float float4 __attribute__((ext_vector_type(4)));
+typedef double double2 __attribute__((ext_vector_type(2)));
+
+// CIR-LABEL: @_Z8nvvm_lduPKv
+// LLVM-LABEL: @_Z8nvvm_lduPKv
+__device__ void nvvm_ldu(const void *p) {
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!s8i>, !s32i) -> !s8i
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!u8i>, !s32i) -> !u8i
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!s8i>, !s32i) -> !s8i
+ // LLVM: call i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr {{%[0-9]+}}, i32 1)
+ // LLVM: call i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr {{%[0-9]+}}, i32 1)
+ // LLVM: call i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr {{%[0-9]+}}, i32 1)
+ __nvvm_ldu_c((const char *)p);
+ __nvvm_ldu_uc((const unsigned char *)p);
+ __nvvm_ldu_sc((const signed char *)p);
+
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!s16i>, !s32i) -> !s16i
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!u16i>, !s32i) -> !u16i
+ // LLVM: call i16 @llvm.nvvm.ldu.global.i.i16.p0(ptr {{%[0-9]+}}, i32 2)
+ // LLVM: call i16 @llvm.nvvm.ldu.global.i.i16.p0(ptr {{%[0-9]+}}, i32 2)
+ __nvvm_ldu_s((const short *)p);
+ __nvvm_ldu_us((const unsigned short *)p);
+
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!u32i>, !s32i) -> !u32i
+ // LLVM: call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr {{%[0-9]+}}, i32 4)
+ // LLVM: call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr {{%[0-9]+}}, i32 4)
+ __nvvm_ldu_i((const int *)p);
+ __nvvm_ldu_ui((const unsigned int *)p);
+
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!s64i>, !s32i) -> !s64i
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!u64i>, !s32i) -> !u64i
+ // LLVM: call i64 @llvm.nvvm.ldu.global.i.i64.p0(ptr {{%[0-9]+}}, i32 8)
+ // LLVM: call i64 @llvm.nvvm.ldu.global.i.i64.p0(ptr {{%[0-9]+}}, i32 8)
+ __nvvm_ldu_l((const long *)p);
+ __nvvm_ldu_ul((const unsigned long *)p);
+
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.f" {{.*}} : (!cir.ptr<!cir.float>, !s32i) -> !cir.float
+ // LLVM: call {{.*}}float @llvm.nvvm.ldu.global.f.f32.p0(ptr {{%[0-9]+}}, i32 4)
+ __nvvm_ldu_f((const float *)p);
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.f" {{.*}} : (!cir.ptr<!cir.double>, !s32i) -> !cir.double
+ // LLVM: call {{.*}}double @llvm.nvvm.ldu.global.f.f64.p0(ptr {{%[0-9]+}}, i32 8)
+ __nvvm_ldu_d((const double *)p);
+
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s8i>>, !s32i) -> !cir.vector<2 x !s8i>
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !u8i>>, !s32i) -> !cir.vector<2 x !u8i>
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s8i>>, !s32i) -> !cir.vector<2 x !s8i>
+ // LLVM: call <2 x i8> @llvm.nvvm.ldu.global.i.v2i8.p0(ptr {{%[0-9]+}}, i32 2)
+ // LLVM: call <2 x i8> @llvm.nvvm.ldu.global.i.v2i8.p0(ptr {{%[0-9]+}}, i32 2)
+ // LLVM: call <2 x i8> @llvm.nvvm.ldu.global.i.v2i8.p0(ptr {{%[0-9]+}}, i32 2)
+ __nvvm_ldu_c2((const char2 *)p);
+ __nvvm_ldu_uc2((const uchar2 *)p);
+ __nvvm_ldu_sc2((const schar2 *)p);
+
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !s8i>>, !s32i) -> !cir.vector<4 x !s8i>
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !u8i>>, !s32i) -> !cir.vector<4 x !u8i>
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !s8i>>, !s32i) -> !cir.vector<4 x !s8i>
+ // LLVM: call <4 x i8> @llvm.nvvm.ldu.global.i.v4i8.p0(ptr {{%[0-9]+}}, i32 4)
+ // LLVM: call <4 x i8> @llvm.nvvm.ldu.global.i.v4i8.p0(ptr {{%[0-9]+}}, i32 4)
+ // LLVM: call <4 x i8> @llvm.nvvm.ldu.global.i.v4i8.p0(ptr {{%[0-9]+}}, i32 4)
+ __nvvm_ldu_c4((const char4 *)p);
+ __nvvm_ldu_uc4((const uchar4 *)p);
+ __nvvm_ldu_sc4((const schar4 *)p);
+
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s16i>>, !s32i) -> !cir.vector<2 x !s16i>
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !u16i>>, !s32i) -> !cir.vector<2 x !u16i>
+ // LLVM: call <2 x i16> @llvm.nvvm.ldu.global.i.v2i16.p0(ptr {{%[0-9]+}}, i32 4)
+ // LLVM: call <2 x i16> @llvm.nvvm.ldu.global.i.v2i16.p0(ptr {{%[0-9]+}}, i32 4)
+ __nvvm_ldu_s2((const short2 *)p);
+ __nvvm_ldu_us2((const ushort2 *)p);
+
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !s16i>>, !s32i) -> !cir.vector<4 x !s16i>
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !u16i>>, !s32i) -> !cir.vector<4 x !u16i>
+ // LLVM: call <4 x i16> @llvm.nvvm.ldu.global.i.v4i16.p0(ptr {{%[0-9]+}}, i32 8)
+ // LLVM: call <4 x i16> @llvm.nvvm.ldu.global.i.v4i16.p0(ptr {{%[0-9]+}}, i32 8)
+ __nvvm_ldu_s4((const short4 *)p);
+ __nvvm_ldu_us4((const ushort4 *)p);
+
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s32i>>, !s32i) -> !cir.vector<2 x !s32i>
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !u32i>>, !s32i) -> !cir.vector<2 x !u32i>
+ // LLVM: call <2 x i32> @llvm.nvvm.ldu.global.i.v2i32.p0(ptr {{%[0-9]+}}, i32 8)
+ // LLVM: call <2 x i32> @llvm.nvvm.ldu.global.i.v2i32.p0(ptr {{%[0-9]+}}, i32 8)
+ __nvvm_ldu_i2((const int2 *)p);
+ __nvvm_ldu_ui2((const uint2 *)p);
+
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !s32i>>, !s32i) -> !cir.vector<4 x !s32i>
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !u32i>>, !s32i) -> !cir.vector<4 x !u32i>
+ // LLVM: call <4 x i32> @llvm.nvvm.ldu.global.i.v4i32.p0(ptr {{%[0-9]+}}, i32 16)
+ // LLVM: call <4 x i32> @llvm.nvvm.ldu.global.i.v4i32.p0(ptr {{%[0-9]+}}, i32 16)
+ __nvvm_ldu_i4((const int4 *)p);
+ __nvvm_ldu_ui4((const uint4 *)p);
+
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s64i>>, !s32i) -> !cir.vector<2 x !s64i>
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !u64i>>, !s32i) -> !cir.vector<2 x !u64i>
+ // LLVM: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16)
+ // LLVM: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16)
+ __nvvm_ldu_l2((const long2 *)p);
+ __nvvm_ldu_ul2((const ulong2 *)p);
+
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s64i>>, !s32i) -> !cir.vector<2 x !s64i>
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !u64i>>, !s32i) -> !cir.vector<2 x !u64i>
+ // LLVM: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16)
+ // LLVM: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16)
+ __nvvm_ldu_ll2((const longlong2 *)p);
+ __nvvm_ldu_ull2((const ulonglong2 *)p);
+
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.f" {{.*}} : (!cir.ptr<!cir.vector<2 x !cir.float>>, !s32i) -> !cir.vector<2 x !cir.float>
+ // LLVM: call {{.*}}<2 x float> @llvm.nvvm.ldu.global.f.v2f32.p0(ptr {{%[0-9]+}}, i32 8)
+ __nvvm_ldu_f2((const float2 *)p);
+
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.f" {{.*}} : (!cir.ptr<!cir.vector<4 x !cir.float>>, !s32i) -> !cir.vector<4 x !cir.float>
+ // LLVM: call {{.*}}<4 x float> @llvm.nvvm.ldu.global.f.v4f32.p0(ptr {{%[0-9]+}}, i32 16)
+ __nvvm_ldu_f4((const float4 *)p);
+
+ // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.f" {{.*}} : (!cir.ptr<!cir.vector<2 x !cir.double>>, !s32i) -> !cir.vector<2 x !cir.double>
+ // LLVM: call {{.*}}<2 x double> @llvm.nvvm.ldu.global.f.v2f64.p0(ptr {{%[0-9]+}}, i32 16)
+ __nvvm_ldu_d2((const double2 *)p);
+}
``````````
</details>
https://github.com/llvm/llvm-project/pull/211843
More information about the cfe-commits
mailing list