Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[CIR][CUDA] initial support for __constant__ variables #1436

Closed
wants to merge 3 commits into from
Closed

[CIR][CUDA] initial support for __constant__ variables #1436

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
e440258
emit const attr for cuda
anominos Feb 25, 2025
73fabd6
emitGlobal: move shouldEmitCUDAGlobalVar to first check
anominos Feb 28, 2025
76901bb
tests: add constant to global-vars.cu
anominos Mar 3, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
29 changes: 12 additions & 17 deletions clang/lib/CIR/CodeGen/CIRGenModule.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -517,15 +517,11 @@ bool CIRGenModule::shouldEmitCUDAGlobalVar(const VarDecl *global) const {
// device-side variables because the CUDA runtime needs their
// size and host-side address in order to provide access to
// their device-side incarnations.

if (global->hasAttr<CUDAConstantAttr>() ||
global->hasAttr<CUDASharedAttr>() ||
global->getType()->isCUDADeviceBuiltinSurfaceType() ||
global->getType()->isCUDADeviceBuiltinTextureType()) {
llvm_unreachable("NYI");
}

return !langOpts.CUDAIsDevice || global->hasAttr<CUDADeviceAttr>();
return !langOpts.CUDAIsDevice || global->hasAttr<CUDADeviceAttr>() ||
global->hasAttr<CUDAConstantAttr>() ||
global->hasAttr<CUDASharedAttr>() ||
global->getType()->isCUDADeviceBuiltinSurfaceType() ||
global->getType()->isCUDADeviceBuiltinTextureType();
}

void CIRGenModule::emitGlobal(GlobalDecl gd) {
Expand All @@ -549,7 +545,10 @@ void CIRGenModule::emitGlobal(GlobalDecl gd) {

if (langOpts.CUDA || langOpts.HIP) {
// clang uses the same flag when building HIP code
if (langOpts.CUDAIsDevice) {
if (const auto *vd = dyn_cast<VarDecl>(global)) {
if (!shouldEmitCUDAGlobalVar(vd))
return;
} else if (langOpts.CUDAIsDevice) {
// This will implicitly mark templates and their
// specializations as __host__ __device__.
if (langOpts.OffloadImplicitHostDeviceTemplates)
Expand All @@ -571,11 +570,6 @@ void CIRGenModule::emitGlobal(GlobalDecl gd) {
return;
}
}

if (const auto *vd = dyn_cast<VarDecl>(global)) {
if (!shouldEmitCUDAGlobalVar(vd))
return;
}
}

if (langOpts.OpenMP) {
Expand Down Expand Up @@ -1452,8 +1446,9 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *d,
emitter->finalize(gv);

// TODO(cir): If it is safe to mark the global 'constant', do so now.
gv.setConstant(!needsGlobalCtor && !needsGlobalDtor &&
isTypeConstant(d->getType(), true, true));
gv.setConstant((d->hasAttr<CUDAConstantAttr>() && langOpts.CUDAIsDevice) ||
(!needsGlobalCtor && !needsGlobalDtor &&
isTypeConstant(d->getType(), true, true)));

// If it is in a read-only section, mark it 'constant'.
if (const SectionAttr *sa = d->getAttr<SectionAttr>())
Expand Down
2 changes: 1 addition & 1 deletion clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -652,6 +652,7 @@ AddressSpaceAttr::getValueFromLangAS(clang::LangAS langAS) {
case LangAS::cuda_device:
return Kind::offload_global;
case LangAS::opencl_constant:
case LangAS::cuda_constant:
return Kind::offload_constant;
case LangAS::opencl_private:
return Kind::offload_private;
Expand All @@ -660,7 +661,6 @@ AddressSpaceAttr::getValueFromLangAS(clang::LangAS langAS) {

case LangAS::opencl_global_device:
case LangAS::opencl_global_host:
case LangAS::cuda_constant:
case LangAS::sycl_global:
case LangAS::sycl_global_device:
case LangAS::sycl_global_host:
Expand Down
6 changes: 5 additions & 1 deletion clang/test/CIR/CodeGen/CUDA/global-vars.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,4 +8,8 @@

__device__ int a;

// CIR-DEVICE: cir.global external addrspace(offload_global) @a = #cir.int<0> : !s32i {alignment = 4 : i64} loc(#loc3)
// CIR-DEVICE: cir.global external addrspace(offload_global) @a ={{.*}}

__constant__ int b;

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Could you please add a checker for the corresponding LLVM IR?

// CIR-DEVICE: cir.global constant external addrspace(offload_constant) @b ={{.*}}
Loading