[llvm-branch-commits] [clang] [CIR][CUDA] Add NVPTX target info and CUDA/HIP global emission filtering (PR #177827)
via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Sat Jan 24 21:10:37 PST 2026
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clangir
Author: David Rivera (RiverDave)
<details>
<summary>Changes</summary>
This patch adds foundational infra for device-side CUDA/HIP compilation by introducing NVPTX target info and implementing the global emission filtering logic.
NVPTX Target Info to allows us to compile against that triple:
- Add NVPTXABIInfo and NVPTXTargetCIRGenInfo classes
- Wire up nvptx and nvptx64 triples in getTargetCIRGenInfo()
- Add createNVPTXTargetCIRGenInfo() factory function
CUDA/HIP Global Emission Filtering (most of this is boilerplate from the AST) This basically narrows down to:
- Skip host-only functions (no `__device__` attribute) when `-fcuda-is-device`
- Skip device-only functions (device without host) on host side
- Always emit ` __global__` kernels and `__host__` `__device__` functions on both sides
- Add `shouldEmitCUDAGlobalVar()` to handle variable emission (device/constant/shared variables)
- Handle special cases: implicit host/device templates, lambda call operators
---
Full diff: https://github.com/llvm/llvm-project/pull/177827.diff
6 Files Affected:
- (modified) clang/lib/CIR/CodeGen/CIRGenModule.cpp (+65)
- (modified) clang/lib/CIR/CodeGen/CIRGenModule.h (+4)
- (modified) clang/lib/CIR/CodeGen/TargetInfo.cpp (+19)
- (modified) clang/lib/CIR/CodeGen/TargetInfo.h (+2)
- (added) clang/test/CIR/CodeGen/CUDA/filter-decl.cu (+37)
- (added) clang/test/CIR/CodeGen/CUDA/nvptx-basic.cu (+30)
``````````diff
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index b535eab913a5d..ffe8d1e55056f 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -16,6 +16,7 @@
#include "CIRGenFunction.h"
#include "clang/AST/ASTContext.h"
+#include "clang/AST/ASTLambda.h"
#include "clang/AST/DeclBase.h"
#include "clang/AST/DeclOpenACC.h"
#include "clang/AST/GlobalDecl.h"
@@ -28,6 +29,7 @@
#include "clang/CIR/MissingFeatures.h"
#include "CIRGenFunctionInfo.h"
+#include "TargetInfo.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/Location.h"
#include "mlir/IR/MLIRContext.h"
@@ -242,6 +244,10 @@ const TargetCIRGenInfo &CIRGenModule::getTargetCIRGenInfo() {
return *theTargetCIRGenInfo;
}
}
+ case llvm::Triple::nvptx:
+ case llvm::Triple::nvptx64:
+ theTargetCIRGenInfo = createNVPTXTargetCIRGenInfo(genTypes);
+ return *theTargetCIRGenInfo;
}
}
@@ -358,6 +364,36 @@ void CIRGenModule::emitDeferred() {
}
}
+template <typename AttrT> static bool hasImplicitAttr(const ValueDecl *decl) {
+ if (!decl)
+ return false;
+ if (auto *attr = decl->getAttr<AttrT>())
+ return attr->isImplicit();
+ return decl->isImplicit();
+}
+
+// This function returns true if M is a specialization, a template,
+// or a non-generic lambda call operator.
+inline bool isLambdaCallOperator(const CXXMethodDecl *MD) {
+ const CXXRecordDecl *LambdaClass = MD->getParent();
+ if (!LambdaClass || !LambdaClass->isLambda())
+ return false;
+ return MD->getOverloadedOperator() == OO_Call;
+}
+
+bool CIRGenModule::shouldEmitCUDAGlobalVar(const VarDecl *global) const {
+ assert(langOpts.CUDA && "Should not be called by non-CUDA languages");
+ // We need to emit host-side 'shadows' for all global
+ // device-side variables because the CUDA runtime needs their
+ // size and host-side address in order to provide access to
+ // their device-side incarnations.
+ return !langOpts.CUDAIsDevice || global->hasAttr<CUDADeviceAttr>() ||
+ global->hasAttr<CUDAConstantAttr>() ||
+ global->hasAttr<CUDASharedAttr>() ||
+ global->getType()->isCUDADeviceBuiltinSurfaceType() ||
+ global->getType()->isCUDADeviceBuiltinTextureType();
+}
+
void CIRGenModule::emitGlobal(clang::GlobalDecl gd) {
if (const auto *cd = dyn_cast<clang::OpenACCConstructDecl>(gd.getDecl())) {
emitGlobalOpenACCDecl(cd);
@@ -372,6 +408,35 @@ void CIRGenModule::emitGlobal(clang::GlobalDecl gd) {
const auto *global = cast<ValueDecl>(gd.getDecl());
+ // If this is CUDA, be selective about which declarations we emit.
+ // Non-constexpr non-lambda implicit host device functions are not emitted
+ // unless they are used on device side.
+ if (langOpts.CUDA) {
+ assert((isa<FunctionDecl>(global) || isa<VarDecl>(global)) &&
+ "Expected Variable or Function");
+ if (const auto *varDecl = dyn_cast<VarDecl>(global)) {
+ if (!shouldEmitCUDAGlobalVar(varDecl))
+ return;
+ } else if (langOpts.CUDAIsDevice) {
+ const auto *functionDecl = dyn_cast<FunctionDecl>(global);
+ if ((!global->hasAttr<CUDADeviceAttr>() ||
+ (langOpts.OffloadImplicitHostDeviceTemplates &&
+ hasImplicitAttr<CUDAHostAttr>(functionDecl) &&
+ hasImplicitAttr<CUDADeviceAttr>(functionDecl) &&
+ !functionDecl->isConstexpr() &&
+ !isLambdaCallOperator(functionDecl) &&
+ !getASTContext().CUDAImplicitHostDeviceFunUsedByDevice.count(
+ functionDecl))) &&
+ !global->hasAttr<CUDAGlobalAttr>() &&
+ !(langOpts.HIPStdPar && isa<FunctionDecl>(global) &&
+ !global->hasAttr<CUDAHostAttr>()))
+ return;
+ // Device-only functions are the only things we skip.
+ } else if (!global->hasAttr<CUDAHostAttr>() &&
+ global->hasAttr<CUDADeviceAttr>())
+ return;
+ }
+
if (const auto *fd = dyn_cast<FunctionDecl>(global)) {
// Update deferred annotations with the latest declaration if the function
// was already used or defined.
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index 3c4f35bacc4f9..9503f993f56d3 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -557,6 +557,10 @@ class CIRGenModule : public CIRGenTypeCache {
static void setInitializer(cir::GlobalOp &op, mlir::Attribute value);
+ // Whether a global variable should be emitted by CUDA/HIP host/device
+ // related attributes.
+ bool shouldEmitCUDAGlobalVar(const VarDecl *global) const;
+
void replaceUsesOfNonProtoTypeWithRealFunction(mlir::Operation *old,
cir::FuncOp newFn);
diff --git a/clang/lib/CIR/CodeGen/TargetInfo.cpp b/clang/lib/CIR/CodeGen/TargetInfo.cpp
index 377c532e492d9..dc29dc0204c19 100644
--- a/clang/lib/CIR/CodeGen/TargetInfo.cpp
+++ b/clang/lib/CIR/CodeGen/TargetInfo.cpp
@@ -56,6 +56,25 @@ class X8664TargetCIRGenInfo : public TargetCIRGenInfo {
} // namespace
+namespace {
+
+class NVPTXABIInfo : public ABIInfo {
+public:
+ NVPTXABIInfo(CIRGenTypes &cgt) : ABIInfo(cgt) {}
+};
+
+class NVPTXTargetCIRGenInfo : public TargetCIRGenInfo {
+public:
+ NVPTXTargetCIRGenInfo(CIRGenTypes &cgt)
+ : TargetCIRGenInfo(std::make_unique<NVPTXABIInfo>(cgt)) {}
+};
+} // namespace
+
+std::unique_ptr<TargetCIRGenInfo>
+clang::CIRGen::createNVPTXTargetCIRGenInfo(CIRGenTypes &cgt) {
+ return std::make_unique<NVPTXTargetCIRGenInfo>(cgt);
+}
+
std::unique_ptr<TargetCIRGenInfo>
clang::CIRGen::createX8664TargetCIRGenInfo(CIRGenTypes &cgt) {
return std::make_unique<X8664TargetCIRGenInfo>(cgt);
diff --git a/clang/lib/CIR/CodeGen/TargetInfo.h b/clang/lib/CIR/CodeGen/TargetInfo.h
index 9535ba94fb08b..bab838692e215 100644
--- a/clang/lib/CIR/CodeGen/TargetInfo.h
+++ b/clang/lib/CIR/CodeGen/TargetInfo.h
@@ -124,6 +124,8 @@ class TargetCIRGenInfo {
std::unique_ptr<TargetCIRGenInfo> createX8664TargetCIRGenInfo(CIRGenTypes &cgt);
+std::unique_ptr<TargetCIRGenInfo> createNVPTXTargetCIRGenInfo(CIRGenTypes &cgt);
+
} // namespace clang::CIRGen
#endif // LLVM_CLANG_LIB_CIR_TARGETINFO_H
diff --git a/clang/test/CIR/CodeGen/CUDA/filter-decl.cu b/clang/test/CIR/CodeGen/CUDA/filter-decl.cu
new file mode 100644
index 0000000000000..ac1e7aeb4f1e1
--- /dev/null
+++ b/clang/test/CIR/CodeGen/CUDA/filter-decl.cu
@@ -0,0 +1,37 @@
+// Based on clang/test/CodeGenCUDA/filter-decl.cu tailored for CIR current capabilities.
+// Tests that host/device functions are emitted only on the appropriate side.
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x cuda \
+// RUN: -I%S/../inputs -emit-cir %s -o %t.host.cir
+// RUN: FileCheck --input-file=%t.host.cir %s --check-prefix=CHECK-HOST
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda \
+// RUN: -I%S/../inputs -fcuda-is-device -emit-cir %s -o %t.device.cir
+// RUN: FileCheck --input-file=%t.device.cir %s --check-prefix=CHECK-DEVICE
+
+#include "cuda.h"
+
+// Implicit host function (no attribute) — host only
+// CHECK-HOST: cir.func {{.*}} @_Z20implicithostonlyfuncv()
+// CHECK-DEVICE-NOT: @_Z20implicithostonlyfuncv
+void implicithostonlyfunc(void) {}
+
+// Explicit __host__ function — host only
+// CHECK-HOST: cir.func {{.*}} @_Z20explicithostonlyfuncv()
+// CHECK-DEVICE-NOT: @_Z20explicithostonlyfuncv
+__host__ void explicithostonlyfunc(void) {}
+
+// __device__ function — device only
+// CHECK-HOST-NOT: @_Z14deviceonlyfuncv
+// CHECK-DEVICE: cir.func {{.*}} @_Z14deviceonlyfuncv()
+__device__ void deviceonlyfunc(void) {}
+
+// __host__ __device__ function — both sides
+// CHECK-HOST: cir.func {{.*}} @_Z14hostdevicefuncv()
+// CHECK-DEVICE: cir.func {{.*}} @_Z14hostdevicefuncv()
+__host__ __device__ void hostdevicefunc(void) {}
+
+// __global__ kernel — both sides (stub on host, kernel on device)
+// CHECK-HOST: cir.func {{.*}} @__device_stub__globalfunc()
+// CHECK-DEVICE: cir.func {{.*}} @_Z10globalfuncv()
+__global__ void globalfunc(void) {}
diff --git a/clang/test/CIR/CodeGen/CUDA/nvptx-basic.cu b/clang/test/CIR/CodeGen/CUDA/nvptx-basic.cu
new file mode 100644
index 0000000000000..fe2233de4d10c
--- /dev/null
+++ b/clang/test/CIR/CodeGen/CUDA/nvptx-basic.cu
@@ -0,0 +1,30 @@
+// Based on clang/test/CodeGenCUDA/ptx-kernels.cu tailored for CIR current capabilities.
+// Tests basic device-side compilation with NVPTX target.
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda \
+// RUN: -I%S/../inputs -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --input-file=%t.cir %s
+
+#include "cuda.h"
+
+// CHECK: cir.func {{.*}} @device_function()
+extern "C"
+__device__ void device_function() {}
+
+// CHECK: cir.func {{.*}} @global_function()
+// CHECK: cir.call @device_function()
+extern "C"
+__global__ void global_function() {
+ device_function();
+}
+
+// Template kernel with explicit instantiation
+template <typename T> __global__ void templated_kernel(T param) {}
+template __global__ void templated_kernel<int>(int);
+// CHECK: cir.func {{.*}} @_Z16templated_kernelIiEvT_
+
+// Anonymous namespace kernel
+namespace {
+__global__ void anonymous_ns_kernel() {}
+// CHECK: cir.func {{.*}} @_ZN12_GLOBAL__N_119anonymous_ns_kernelEv
+}
``````````
</details>
https://github.com/llvm/llvm-project/pull/177827
More information about the llvm-branch-commits
mailing list