Skip to content

[CIR][CUDA] implement cuda constant variables #1444

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

Merged
merged 14 commits into from
Mar 11, 2025
Merged

[CIR][CUDA] implement cuda constant variables #1444

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
14 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
2617e0f
implement const
anominos Mar 6, 2025
a250e36
fix address space for offload_constant
anominos Mar 6, 2025
8a3c8f1
Merge branch 'main' into anom/constant-var
anominos Mar 6, 2025
d514fc4
Merge branch 'anom/const-var' into anom/constant-var
anominos Mar 6, 2025
33c353f
Fix address space values for NVPTX
anominos Mar 6, 2025
92cc36e
add testcase
anominos Mar 7, 2025
6e7f8ab
Merge branch 'anom/fix-nvptx-addrspace' into anom/constant-var
anominos Mar 7, 2025
07ac4f5
test: enable const test in addrspace-lowering.cu
anominos Mar 7, 2025
3431265
Merge branch 'main' into anom/constant-var
anominos Mar 8, 2025
bf4c995
remove duplicated case
anominos Mar 8, 2025
0a26ede
Merge branch 'main' into anom/constant-var
anominos Mar 11, 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
11 changes: 6 additions & 5 deletions clang/lib/CIR/CodeGen/CIRGenModule.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -569,12 +569,12 @@ bool CIRGenModule::shouldEmitCUDAGlobalVar(const VarDecl *global) const {
// size and host-side address in order to provide access to
// their device-side incarnations.

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

return !langOpts.CUDAIsDevice || global->hasAttr<CUDADeviceAttr>() ||
global->hasAttr<CUDAConstantAttr>() ||
global->hasAttr<CUDASharedAttr>() ||
global->getType()->isCUDADeviceBuiltinSurfaceType();
}
Expand Down Expand Up @@ -1492,7 +1492,7 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *d,
// __shared__ variables is not marked as externally initialized,
// because they must not be initialized.
if (linkage != cir::GlobalLinkageKind::InternalLinkage &&
(d->hasAttr<CUDADeviceAttr>() ||
(d->hasAttr<CUDADeviceAttr>() || d->hasAttr<CUDAConstantAttr>() ||
d->getType()->isCUDADeviceBuiltinSurfaceType())) {
gv->setAttr(CUDAExternallyInitializedAttr::getMnemonic(),
CUDAExternallyInitializedAttr::get(&getMLIRContext()));
Expand All @@ -1505,8 +1505,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
3 changes: 1 addition & 2 deletions 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,13 +652,12 @@ 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;
case LangAS::opencl_generic:
return Kind::offload_generic;
case LangAS::cuda_constant:
return Kind::offload_constant;
case LangAS::opencl_global_device:
case LangAS::opencl_global_host:
case LangAS::sycl_global:
Expand Down
4 changes: 2 additions & 2 deletions clang/test/CIR/CodeGen/CUDA/addrspace-lowering.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,6 +14,6 @@ __device__ int b;

// LLVM-DEVICE: @b = addrspace(1) {{.*}}

// __constant__ int c;
__constant__ int c;

// XFAIL-LLVM-DEVICE: @c = addrspace(4) {{.*}}
// LLVM-DEVICE: @c = addrspace(4) {{.*}}
4 changes: 4 additions & 0 deletions 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 @@ -17,3 +17,7 @@ __device__ int a;
__shared__ int shared;
// CIR-DEVICE: cir.global external addrspace(offload_local) @shared = #cir.undef
// LLVM-DEVICE: @shared = addrspace(3) global i32 undef, align 4

__constant__ int b;
// CIR-DEVICE: cir.global constant external addrspace(offload_constant) @b = #cir.int<0> : !s32i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized}
// LLVM-DEVICE: @b = addrspace(4) externally_initialized constant i32 0, align 4