[clang] Cir cuda builtin vars (PR #195539)
via cfe-commits
cfe-commits at lists.llvm.org
Sun May 3 10:46:29 PDT 2026
https://github.com/AbdallahRashed updated https://github.com/llvm/llvm-project/pull/195539
>From ff33ae555b73dcdf30c1dd7588988a7ef43db665 Mon Sep 17 00:00:00 2001
From: AbdallahRashed <abdallah.mrashed at gmail.com>
Date: Sun, 3 May 2026 19:44:55 +0200
Subject: [PATCH 1/2] [CIR] Support PseudoObjectExpr scalar/RValue emission
Implement emitPseudoObjectRValue and fix VisitPseudoObjectExpr in the
scalar emitter to call it instead of errorNYI. Also remove the
errorNYI guard for unique OpaqueValueExprs in the PseudoObjectExpr
emission loop, matching the behavior of classic CodeGen (just assert
and continue).
This unblocks CUDA builtin variable access (e.g. threadIdx.x) which
goes through PseudoObjectExpr in the AST.
---
clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp | 3 +--
clang/lib/CIR/CodeGen/CIRGenFunction.cpp | 25 +++++++---------------
clang/lib/CIR/CodeGen/CIRGenFunction.h | 2 ++
3 files changed, 11 insertions(+), 19 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp
index 92b7156f3a3a8..231039ec5da29 100644
--- a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp
@@ -272,8 +272,7 @@ class ScalarExprEmitter : public StmtVisitor<ScalarExprEmitter, mlir::Value> {
convertType(e->getType()), e->getPackLength());
}
mlir::Value VisitPseudoObjectExpr(PseudoObjectExpr *e) {
- cgf.cgm.errorNYI(e->getSourceRange(), "ScalarExprEmitter: pseudo object");
- return {};
+ return cgf.emitPseudoObjectRValue(e).getValue();
}
mlir::Value VisitSYCLUniqueStableNameExpr(SYCLUniqueStableNameExpr *e) {
cgf.cgm.errorNYI(e->getSourceRange(),
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp
index 32b4881a93095..5119b0cddd87e 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp
@@ -1075,10 +1075,6 @@ emitPseudoObjectExpr(CIRGenFunction &cgf, const PseudoObjectExpr *e,
// Skip unique OVEs.
if (ov->isUnique()) {
- // FIXME: This doesn't really affect anything, but I cannot find a test
- // for this, so leave an ErrorNYI here until we can find one.
- cgf.cgm.errorNYI(e->getSourceRange(),
- "emitPseudoObjectExpr skipped for uniqueness");
assert(ov != resultExpr &&
"A unique OVE cannot be used as the result expression");
continue;
@@ -1096,15 +1092,10 @@ emitPseudoObjectExpr(CIRGenFunction &cgf, const PseudoObjectExpr *e,
// If this is the result, also evaluate the result now.
if (ov == resultExpr) {
- // FIXME: This doesn't really affect anything, but I cannot find a
- // test for this, so leave an ErrorNYI here until we can find one.
- cgf.cgm.errorNYI(e->getSourceRange(),
- "emitPseudoObjectExpr as result");
if (forLValue)
result = cgf.emitLValue(ov);
else
- cgf.cgm.errorNYI(e->getSourceRange(),
- "emitPseudoObjectExpr as an RValue");
+ result = cgf.emitAnyExpr(ov, slot);
}
}
opaques.push_back(opaqueData);
@@ -1114,14 +1105,8 @@ emitPseudoObjectExpr(CIRGenFunction &cgf, const PseudoObjectExpr *e,
if (forLValue)
result = cgf.emitLValue(semantic);
else
- cgf.cgm.errorNYI(
- e->getSourceRange(),
- "emitPseudoObjectExpr as an RValue, when semantic is result");
+ result = cgf.emitAnyExpr(semantic, slot);
} else {
- // FIXME: best I can tell, this is only reachable as an r-value, so this
- // isn't properly tested.
- cgf.cgm.errorNYI(e->getSourceRange(),
- "emitPseudoObjectExpr as an ignored value");
// Otherwise, evaluate the expression in an ignored context.
cgf.emitIgnoredExpr(semantic);
}
@@ -1130,6 +1115,12 @@ emitPseudoObjectExpr(CIRGenFunction &cgf, const PseudoObjectExpr *e,
return result;
}
+RValue CIRGenFunction::emitPseudoObjectRValue(const PseudoObjectExpr *e,
+ AggValueSlot slot) {
+ return std::get<RValue>(
+ emitPseudoObjectExpr(*this, e, /*forLValue=*/false, slot));
+}
+
LValue CIRGenFunction::emitPseudoObjectLValue(const PseudoObjectExpr *e) {
return std::get<LValue>(emitPseudoObjectExpr(*this, e, /*forLValue=*/true,
AggValueSlot::ignored()));
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index 3905c154e472c..5c3c1e887ae01 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -1563,6 +1563,8 @@ class CIRGenFunction : public CIRGenTypeCache {
AutoVarEmission emitAutoVarAlloca(const clang::VarDecl &d,
mlir::OpBuilder::InsertPoint ip = {});
+ RValue emitPseudoObjectRValue(const PseudoObjectExpr *e,
+ AggValueSlot slot = AggValueSlot::ignored());
LValue emitPseudoObjectLValue(const PseudoObjectExpr *E);
/// Emit code and set up symbol table for a variable declaration with auto,
>From 8e890851ad2745b539f7f94cb88dc773312db5fe Mon Sep 17 00:00:00 2001
From: AbdallahRashed <abdallah.mrashed at gmail.com>
Date: Sun, 3 May 2026 19:45:13 +0200
Subject: [PATCH 2/2] [CIR][CUDA] Support CUDA builtin variables (threadIdx,
blockIdx, etc.)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
Add CIRGenBuiltinNVPTX.cpp as infrastructure for NVPTX-specific
builtin emission in ClangIR, and wire it into emitTargetArchBuiltinExpr.
The actual PTX special-register reads (__nvvm_read_ptx_sreg_tid_x,
etc.) are already handled by the generic intrinsic path via
Intrinsic::getIntrinsicForClangBuiltin, which maps them to
cir.call_llvm_intrinsic "nvvm.read.ptx.sreg.*" operations. The
NVPTX dispatch file serves as the entry point for future builtins
that need special codegen (e.g., atomics, MMA/tensor operations).
The CUDA builtin variables (threadIdx, blockIdx, blockDim, gridDim)
access these intrinsics through PseudoObjectExpr → inline
__fetch_builtin_{x,y,z}() methods → __nvvm_read_ptx_sreg_* builtins.
Partially addresses llvm/llvm-project#179278.
---
clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp | 4 ++
clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp | 30 +++++++++++++++
clang/lib/CIR/CodeGen/CIRGenFunction.h | 4 ++
clang/lib/CIR/CodeGen/CMakeLists.txt | 1 +
.../test/CIR/CodeGenCUDA/cuda-builtin-vars.cu | 38 +++++++++++++++++++
5 files changed, 77 insertions(+)
create mode 100644 clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
create mode 100644 clang/test/CIR/CodeGenCUDA/cuda-builtin-vars.cu
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
index fa0d02f9e4eef..09a58036c1883 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
@@ -2607,8 +2607,12 @@ emitTargetArchBuiltinExpr(CIRGenFunction *cgf, unsigned builtinID,
case llvm::Triple::amdgcn:
return cgf->emitAMDGPUBuiltinExpr(builtinID, e);
case llvm::Triple::systemz:
+ // These are actually NYI, but that will be reported by emitBuiltinExpr.
+ // At this point, we don't even know that the builtin is target-specific.
+ return std::nullopt;
case llvm::Triple::nvptx:
case llvm::Triple::nvptx64:
+ return cgf->emitNVPTXBuiltinExpr(builtinID, e);
case llvm::Triple::wasm32:
case llvm::Triple::wasm64:
case llvm::Triple::hexagon:
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
new file mode 100644
index 0000000000000..8577808a42660
--- /dev/null
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
@@ -0,0 +1,30 @@
+//===--- CIRGenBuiltinNVPTX.cpp - Emit CIR for NVPTX builtins -------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This contains code to emit NVPTX builtins.
+//
+//===----------------------------------------------------------------------===//
+
+#include "CIRGenFunction.h"
+
+#include "clang/Basic/TargetBuiltins.h"
+
+using namespace clang;
+using namespace clang::CIRGen;
+
+std::optional<mlir::Value>
+CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinID, const CallExpr *e) {
+ // Most NVPTX builtins (including the PTX special-register reads like
+ // __nvvm_read_ptx_sreg_tid_x) are handled by the generic intrinsic path
+ // in emitBuiltinExpr via Intrinsic::getIntrinsicForClangBuiltin.
+ // This function handles builtins that require special codegen.
+ switch (builtinID) {
+ default:
+ return std::nullopt;
+ }
+}
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index 5c3c1e887ae01..20f5145eccc5f 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -2020,6 +2020,10 @@ class CIRGenFunction : public CIRGenTypeCache {
std::optional<mlir::Value> emitAMDGPUBuiltinExpr(unsigned builtinID,
const CallExpr *expr);
+ /// Emit a call to an NVPTX builtin function.
+ std::optional<mlir::Value> emitNVPTXBuiltinExpr(unsigned builtinID,
+ const CallExpr *e);
+
LValue emitOpaqueValueLValue(const OpaqueValueExpr *e);
LValue emitConditionalOperatorLValue(const AbstractConditionalOperator *expr);
diff --git a/clang/lib/CIR/CodeGen/CMakeLists.txt b/clang/lib/CIR/CodeGen/CMakeLists.txt
index b59addbfd6beb..2b6bc036a5d54 100644
--- a/clang/lib/CIR/CodeGen/CMakeLists.txt
+++ b/clang/lib/CIR/CodeGen/CMakeLists.txt
@@ -15,6 +15,7 @@ add_clang_library(clangCIR
CIRGenBuiltinAArch64.cpp
CIRGenBuiltinAMDGPU.cpp
CIRGenAMDGPU.cpp
+ CIRGenBuiltinNVPTX.cpp
CIRGenBuiltinRISCV.cpp
CIRGenBuiltinX86.cpp
CIRGenCall.cpp
diff --git a/clang/test/CIR/CodeGenCUDA/cuda-builtin-vars.cu b/clang/test/CIR/CodeGenCUDA/cuda-builtin-vars.cu
new file mode 100644
index 0000000000000..9ffc2c700b2f5
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/cuda-builtin-vars.cu
@@ -0,0 +1,38 @@
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda \
+// RUN: -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --input-file=%t.cir %s
+
+#include "__clang_cuda_builtin_vars.h"
+
+__attribute__((global))
+void kernel(int *out) {
+ int i = 0;
+ out[i++] = threadIdx.x;
+ out[i++] = threadIdx.y;
+ out[i++] = threadIdx.z;
+
+ out[i++] = blockIdx.x;
+ out[i++] = blockIdx.y;
+ out[i++] = blockIdx.z;
+
+ out[i++] = blockDim.x;
+ out[i++] = blockDim.y;
+ out[i++] = blockDim.z;
+
+ out[i++] = gridDim.x;
+ out[i++] = gridDim.y;
+ out[i++] = gridDim.z;
+}
+
+// CHECK-DAG: cir.call_llvm_intrinsic "nvvm.read.ptx.sreg.tid.x"
+// CHECK-DAG: cir.call_llvm_intrinsic "nvvm.read.ptx.sreg.tid.y"
+// CHECK-DAG: cir.call_llvm_intrinsic "nvvm.read.ptx.sreg.tid.z"
+// CHECK-DAG: cir.call_llvm_intrinsic "nvvm.read.ptx.sreg.ctaid.x"
+// CHECK-DAG: cir.call_llvm_intrinsic "nvvm.read.ptx.sreg.ctaid.y"
+// CHECK-DAG: cir.call_llvm_intrinsic "nvvm.read.ptx.sreg.ctaid.z"
+// CHECK-DAG: cir.call_llvm_intrinsic "nvvm.read.ptx.sreg.ntid.x"
+// CHECK-DAG: cir.call_llvm_intrinsic "nvvm.read.ptx.sreg.ntid.y"
+// CHECK-DAG: cir.call_llvm_intrinsic "nvvm.read.ptx.sreg.ntid.z"
+// CHECK-DAG: cir.call_llvm_intrinsic "nvvm.read.ptx.sreg.nctaid.x"
+// CHECK-DAG: cir.call_llvm_intrinsic "nvvm.read.ptx.sreg.nctaid.y"
+// CHECK-DAG: cir.call_llvm_intrinsic "nvvm.read.ptx.sreg.nctaid.z"
More information about the cfe-commits
mailing list