[clang] Cir cuda builtin vars (PR #195539)
via cfe-commits
cfe-commits at lists.llvm.org
Sun May 3 10:39:17 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clangir
Author: AbdallahRashed (AbdallahRashed)
<details>
<summary>Changes</summary>
---
Full diff: https://github.com/llvm/llvm-project/pull/195539.diff
7 Files Affected:
- (modified) clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp (+4)
- (added) clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp (+30)
- (modified) clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp (+1-2)
- (modified) clang/lib/CIR/CodeGen/CIRGenFunction.cpp (+8-17)
- (modified) clang/lib/CIR/CodeGen/CIRGenFunction.h (+6)
- (modified) clang/lib/CIR/CodeGen/CMakeLists.txt (+1)
- (added) clang/test/CIR/CodeGenCUDA/cuda-builtin-vars.cu (+38)
``````````diff
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/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..fdaf720d728dd 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..20f5145eccc5f 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,
@@ -2018,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"
``````````
</details>
https://github.com/llvm/llvm-project/pull/195539
More information about the cfe-commits
mailing list