[CIR][CUDA] Handle local, __device__, __shared__, and __constant__ variables - #184248
Merged
Merged
Conversation
Signed-off-by: ZakyHermawan <zaky.hermawan9615@gmail.com>
ZakyHermawan
requested review from
andykaylor,
bcardosolopes,
lanza and
xlauko
as code owners
March 2, 2026 22:22
llvmbot
added
clang
Member
|
@llvm/pr-subscribers-clang @llvm/pr-subscribers-clangir Author: Zaky Hermawan (ZakyHermawan) ChangesSupport device and shared variables References: #179278, llvm/clangir#1368, llvm/clangir#1394 Full diff: https://github.com/llvm/llvm-project/pull/184248.diff 7 Files Affected:
diff --git a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
index bb3117dfb2c98..b19e48d0f51d4 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
@@ -433,12 +433,15 @@ CIRGenModule::getOrCreateStaticVarDecl(const VarDecl &d,
mlir::Type lty = getTypes().convertTypeForMem(ty);
assert(!cir::MissingFeatures::addressSpace());
- if (d.hasAttr<LoaderUninitializedAttr>() || d.hasAttr<CUDASharedAttr>())
+ mlir::Attribute init = nullptr;
+ if (d.hasAttr<LoaderUninitializedAttr>())
errorNYI(d.getSourceRange(),
"getOrCreateStaticVarDecl: LoaderUninitializedAttr");
- assert(!cir::MissingFeatures::addressSpace());
+ else if (ty.getAddressSpace() != LangAS::opencl_local &&
+ !d.hasAttr<CUDASharedAttr>())
+ init = builder.getZeroInitAttr(convertType(ty));
- mlir::Attribute init = builder.getZeroInitAttr(convertType(ty));
+ assert(!cir::MissingFeatures::addressSpace());
cir::GlobalOp gv = builder.createVersionedGlobal(
getModule(), getLoc(d.getLocation()), name, lty, false, linkage);
@@ -665,11 +668,6 @@ void CIRGenFunction::emitStaticVarDecl(const VarDecl &d,
var.setAlignment(alignment.getAsAlign().value());
- // There are a lot of attributes that need to be handled here. Until
- // we start to support them, we just report an error if there are any.
- if (d.hasAttrs())
- cgm.errorNYI(d.getSourceRange(), "static var with attrs");
-
if (cgm.getCodeGenOpts().KeepPersistentStorageVariables)
cgm.errorNYI(d.getSourceRange(), "static var keep persistent storage");
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 223b53731359a..1517058af8782 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -797,6 +797,22 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty,
"external const declaration with initializer");
}
+ // TODO(cir): if this method is used to handle functions we must have
+ // something closer to GlobalValue::isDeclaration instead of checking for
+ // initializer.
+ if (gv.isDeclaration()) {
+ // TODO(cir): set target attributes
+
+ // External HIP managed variables needed to be recorded for transformation
+ // in both device and host compilations.
+ // External HIP managed variables needed to be recorded for transformation
+ // in both device and host compilations.
+ if (getLangOpts().CUDA && d && d->hasAttr<HIPManagedAttr>() &&
+ d->hasExternalStorage())
+ llvm_unreachable("NYI");
+ }
+
+ // TODO(cir): address space cast when needed for DAddrSpace.
return gv;
}
@@ -947,10 +963,6 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *vd,
errorNYI(vd->getSourceRange(), "annotate global variable");
}
- if (langOpts.CUDA) {
- errorNYI(vd->getSourceRange(), "CUDA global variable");
- }
-
// Set initializer and finalize emission
CIRGenModule::setInitializer(gv, init);
if (emitter)
@@ -1563,6 +1575,39 @@ CIRGenModule::getAddrOfConstantStringFromLiteral(const StringLiteral *s,
return builder.getGlobalViewAttr(ptrTy, gv);
}
+LangAS CIRGenModule::getGlobalVarAddressSpace(const VarDecl *d) {
+ if (langOpts.OpenCL) {
+ LangAS as = d ? d->getType().getAddressSpace() : LangAS::opencl_global;
+ assert(as == LangAS::opencl_global || as == LangAS::opencl_global_device ||
+ as == LangAS::opencl_global_host || as == LangAS::opencl_constant ||
+ as == LangAS::opencl_local || as >= LangAS::FirstTargetAddressSpace);
+ return as;
+ }
+
+ if (langOpts.SYCLIsDevice &&
+ (!d || d->getType().getAddressSpace() == LangAS::Default))
+ llvm_unreachable("NYI");
+
+ if (langOpts.CUDA && langOpts.CUDAIsDevice) {
+ if (d) {
+ if (d->hasAttr<CUDAConstantAttr>())
+ return LangAS::cuda_constant;
+ if (d->hasAttr<CUDASharedAttr>())
+ return LangAS::cuda_shared;
+ if (d->hasAttr<CUDADeviceAttr>())
+ return LangAS::cuda_device;
+ if (d->getType().isConstQualified())
+ return LangAS::cuda_constant;
+ }
+ return LangAS::cuda_device;
+ }
+
+ if (langOpts.OpenMP)
+ llvm_unreachable("NYI");
+
+ return getTargetCIRGenInfo().getGlobalVarAddressSpace(*this, d);
+}
+
// TODO(cir): this could be a common AST helper for both CIR and LLVM codegen.
LangAS CIRGenModule::getLangTempAllocaAddressSpace() const {
if (getLangOpts().OpenCL)
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index 52464a8bc30c4..d9173234868ee 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -359,6 +359,16 @@ class CIRGenModule : public CIRGenTypeCache {
getAddrOfConstantStringFromLiteral(const StringLiteral *s,
llvm::StringRef name = ".str");
+ /// Return the AST address space of the underlying global variable for D, as
+ /// determined by its declaration. Normally this is the same as the address
+ /// space of D's type, but in CUDA, address spaces are associated with
+ /// declarations, not types. If D is nullptr, return the default address
+ /// space for global variable.
+ ///
+ /// For languages without explicit address spaces, if D has default address
+ /// space, target-specific global or constant address space may be returned.
+ LangAS getGlobalVarAddressSpace(const VarDecl *d);
+
/// Returns the address space for temporary allocations in the language. This
/// ensures that the allocated variable's address space matches the
/// expectations of the AST, rather than using the target's allocation address
diff --git a/clang/lib/CIR/CodeGen/TargetInfo.cpp b/clang/lib/CIR/CodeGen/TargetInfo.cpp
index 2f3824d3d47a7..70ffb46050ea1 100644
--- a/clang/lib/CIR/CodeGen/TargetInfo.cpp
+++ b/clang/lib/CIR/CodeGen/TargetInfo.cpp
@@ -91,3 +91,12 @@ bool TargetCIRGenInfo::isNoProtoCallVariadic(
// For everything else, we just prefer false unless we opt out.
return false;
}
+
+clang::LangAS
+TargetCIRGenInfo::getGlobalVarAddressSpace(CIRGenModule &cgm,
+ const clang::VarDecl *d) const {
+ assert(!cgm.getLangOpts().OpenCL &&
+ !(cgm.getLangOpts().CUDA && cgm.getLangOpts().CUDAIsDevice) &&
+ "Address space agnostic languages only");
+ return d ? d->getType().getAddressSpace() : LangAS::Default;
+}
diff --git a/clang/lib/CIR/CodeGen/TargetInfo.h b/clang/lib/CIR/CodeGen/TargetInfo.h
index f4792d5309e36..8db2cbbce5d23 100644
--- a/clang/lib/CIR/CodeGen/TargetInfo.h
+++ b/clang/lib/CIR/CodeGen/TargetInfo.h
@@ -49,6 +49,13 @@ class TargetCIRGenInfo {
/// Returns ABI info helper for the target.
const ABIInfo &getABIInfo() const { return *info; }
+ /// Get target favored AST address space of a global variable for languages
+ /// other than OpenCL and CUDA.
+ /// If \p d is nullptr, returns the default target favored address space
+ /// for global variable.
+ virtual clang::LangAS getGlobalVarAddressSpace(CIRGenModule &cgm,
+ const clang::VarDecl *d) const;
+
/// Get the address space for alloca.
virtual mlir::ptr::MemorySpaceAttrInterface getCIRAllocaAddressSpace() const {
return cir::LangAddressSpaceAttr::get(&info->cgt.getMLIRContext(),
diff --git a/clang/test/CIR/CodeGenCUDA/address-spaces.cu b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
new file mode 100644
index 0000000000000..68905a6616ca7
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
@@ -0,0 +1,78 @@
+#include "Inputs/cuda.h"
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir \
+// RUN: -fcuda-is-device -emit-cir -target-sdk-version=12.3 \
+// RUN: -I%S/Inputs/ %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR-DEVICE --input-file=%t.cir %s
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \
+// RUN: -x cuda -emit-cir -target-sdk-version=12.3 \
+// RUN: -I%S/Inputs/ %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR-HOST --input-file=%t.cir %s
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir \
+// RUN: -fcuda-is-device -emit-llvm -target-sdk-version=12.3 \
+// RUN: -I%S/Inputs/ %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM-DEVICE --input-file=%t.ll %s
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \
+// RUN: -x cuda -emit-llvm -target-sdk-version=12.3 \
+// RUN: -I%S/Inputs/ %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM-HOST --input-file=%t.ll %s
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
+// RUN: -x cuda -emit-llvm -target-sdk-version=12.3 \
+// RUN: -I%S/Inputs/ %s -o %t.ll
+// RUN: FileCheck --check-prefix=OGCG-HOST --input-file=%t.ll %s
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda \
+// RUN: -fcuda-is-device -emit-llvm -target-sdk-version=12.3 \
+// RUN: -I%S/Inputs/ %s -o %t.ll
+// RUN: FileCheck --check-prefix=OGCG-DEVICE --input-file=%t.ll %s
+
+__global__ void fn() {
+ int i = 0;
+ __shared__ int j;
+ j = i;
+}
+
+// CIR-DEVICE: cir.global "private" internal dso_local @_ZZ2fnvE1j : !s32i
+// CIR-DEVICE: cir.func {{.*}}@_Z2fnv() {{.*}} {
+// CIR-DEVICE: %[[I:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["i", init]
+// CIR-DEVICE: %[[ZERO:.*]] = cir.const #cir.int<0> : !s32i
+// CIR-DEVICE: cir.store {{.*}}%[[ZERO]], %[[I]] : !s32i, !cir.ptr<!s32i>
+// CIR-DEVICE: %[[J:.*]] = cir.get_global @_ZZ2fnvE1j : !cir.ptr<!s32i>
+// CIR-DEVICE: %[[VAL:.*]] = cir.load {{.*}}%[[I]] : !cir.ptr<!s32i>, !s32i
+// CIR-DEVICE: cir.store {{.*}}%[[VAL]], %[[J]] : !s32i, !cir.ptr<!s32i>
+// CIR-DEVICE: cir.return
+
+// CIR-HOST: cir.func private dso_local @__cudaPopCallConfiguration
+// CIR-HOST: cir.func private dso_local @cudaLaunchKernel
+// CIR-HOST: cir.func {{.*}}@_Z17__device_stub__fnv()
+
+// LLVM-DEVICE: @_ZZ2fnvE1j = internal global i32 undef, align 4
+// LLVM-DEVICE: define dso_local void @_Z2fnv()
+// LLVM-DEVICE: %[[ALLOCA:.*]] = alloca i32, i64 1, align 4
+// LLVM-DEVICE: store i32 0, ptr %[[ALLOCA]], align 4
+// LLVM-DEVICE: %[[VAL:.*]] = load i32, ptr %[[ALLOCA]], align 4
+// LLVM-DEVICE: store i32 %[[VAL]], ptr @_ZZ2fnvE1j, align 4
+// LLVM-DEVICE: ret void
+
+// LLVM-HOST: %struct.dim3 = type { i32, i32, i32 }
+// LLVM-HOST: declare {{.*}}i32 @__cudaPopCallConfiguration(ptr, ptr, ptr, ptr)
+// LLVM-HOST: declare {{.*}}i32 @cudaLaunchKernel(ptr, %struct.dim3, %struct.dim3, ptr, i64, ptr)
+// LLVM-HOST: define dso_local void @_Z17__device_stub__fnv()
+
+// OGCG-HOST: define dso_local void @_Z17__device_stub__fnv()
+// OGCG-HOST: entry:
+// OGCG-HOST: call i32 @__cudaPopCallConfiguration
+// OGCG-HOST: call {{.*}}i32 @cudaLaunchKernel
+
+// OGCG-DEVICE: @_ZZ2fnvE1j = internal addrspace(3) global i32 undef, align 4
+// OGCG-DEVICE: define dso_local ptx_kernel void @_Z2fnv()
+// OGCG-DEVICE: entry:
+// OGCG-DEVICE: %[[I:.*]] = alloca i32, align 4
+// OGCG-DEVICE: store i32 0, ptr %[[I]], align 4
+// OGCG-DEVICE: %[[VAL:.*]] = load i32, ptr %[[I]], align 4
+// OGCG-DEVICE: store i32 %[[VAL]], ptr addrspacecast (ptr addrspace(3) @_ZZ2fnvE1j to ptr), align 4
+// OGCG-DEVICE: ret void
diff --git a/clang/test/CIR/CodeGenCUDA/global-vars.cu b/clang/test/CIR/CodeGenCUDA/global-vars.cu
new file mode 100644
index 0000000000000..f497d0e7f5f64
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/global-vars.cu
@@ -0,0 +1,47 @@
+#include "Inputs/cuda.h"
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir \
+// RUN: -fcuda-is-device -emit-cir -target-sdk-version=12.3 \
+// RUN: -I%S/Inputs/ %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR-DEVICE --input-file=%t.cir %s
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \
+// RUN: -x cuda -emit-cir -target-sdk-version=12.3 \
+// RUN: -I%S/Inputs/ %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR-HOST --input-file=%t.cir %s
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir \
+// RUN: -fcuda-is-device -emit-llvm -target-sdk-version=12.3 \
+// RUN: -I%S/Inputs/ %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM-DEVICE --input-file=%t.ll %s
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \
+// RUN: -x cuda -emit-llvm -target-sdk-version=12.3 \
+// RUN: -I%S/Inputs/ %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM-HOST --input-file=%t.ll %s
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda \
+// RUN: -fcuda-is-device -emit-llvm -target-sdk-version=12.3 \
+// RUN: -I%S/Inputs/ %s -o %t.ll
+// RUN: FileCheck --check-prefix=OGCG-DEVICE --input-file=%t.ll %s
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
+// RUN: -x cuda -emit-llvm -target-sdk-version=12.3 \
+// RUN: -I%S/Inputs/ %s -o %t.ll
+// RUN: FileCheck --check-prefix=OGCG-HOST --input-file=%t.ll %s
+
+__shared__ int a;
+// CIR-DEVICE: cir.global external [[SHARED:@.*]] = #cir.int<0> : !s32i {alignment = 4 : i64}
+// CIR-HOST: cir.global external [[SHARED_HOST:@.*]] = #cir.int<0> : !s32i {alignment = 4 : i64}
+// LLVM-DEVICE: @[[SHARED_LL:.*]] = global i32 0, align 4
+// LLVM-HOST: @[[SHARED_LH:.*]] = global i32 0, align 4
+// OGCG-DEVICE: @[[SHARED_OD:.*]] = addrspace(3) global i32 undef, align 4
+// OGCG-HOST: @[[SHARED_OH:.*]] = internal global i32 undef, align 4
+
+__device__ int b;
+// CIR-DEVICE: cir.global external [[DEV:@.*]] = #cir.int<0> : !s32i {alignment = 4 : i64}
+// CIR-HOST: cir.global external [[DEV_HOST:@.*]] = #cir.int<0> : !s32i {alignment = 4 : i64}
+// LLVM-DEVICE: @[[DEV_LD:.*]] = global i32 0, align 4
+// LLVM-HOST: @[[DEV_LH:.*]] = global i32 0, align 4
+// OGCG-HOST: @[[DEV_OH:.*]] = internal global i32 undef, align 4
+// OGCG-DEVICE: @[[DEV_OD:.*]] = addrspace(1) externally_initialized global i32 0, align 4
|
ZakyHermawan
changed the title
ZakyHermawan
changed the title
andykaylor
reviewed
Mar 4, 2026
Remove CIR-HOST LLVM-HOST and OGCG-HOST from global-vars.cu because shadow variables did not handled properly, yet Make few changes to handle __device__, __shared__, and __constant__ global variables using reference from OGCG Create and call a hook (setTargetAttributes) if the variable is global and declaration only. Signed-off-by: ZakyHermawan <zaky.hermawan9615@gmail.com>
ZakyHermawan
changed the title
Contributor
Author
|
I am adding Will create separate PR for shadow variables. |
Signed-off-by: ZakyHermawan <zaky.hermawan9615@gmail.com>
…riables Signed-off-by: ZakyHermawan <zaky.hermawan9615@gmail.com>
Contributor
Author
|
Can I get another round of review on this one ? |
andykaylor
reviewed
Mar 10, 2026
| init = builder.getZeroInitAttr(convertType(ty)); | ||
|
|
||
| mlir::Attribute init = builder.getZeroInitAttr(convertType(ty)); | ||
| assert(!cir::MissingFeatures::addressSpace()); |
Contributor
There was a problem hiding this comment.
What's missing here? Is it the AS argument to createVersionedGlobal()?
Contributor
Author
There was a problem hiding this comment.
createVersionedGlobal did not requrie AddressSpace as argument.
Do the assert statement still necessary ?
Contributor
There was a problem hiding this comment.
That is being added in #179082, so it will depend on the order in which these patches are merged. I don't think the assertion is needed since that's already in progress, but it's also fine to leave it.
…agnostics Set initial value for opencl local variables to undef. Set initial value for variable with attributes CUDASharedAttr and LoaderUninitializedAttr to undef. Remove unused LangAS getGlobalVarAddressSpace(const VarDecl *d); Improve diagnostics for NYI: print function name where the diagnostics are being emitted. Signed-off-by: ZakyHermawan <zaky.hermawan9615@gmail.com>
Contributor
Author
|
I am moving tests for global variables in address-spaces.cu, mirroring OGCG. |
ayasin-a
pushed a commit
to ayasin-a/llvm-project
that referenced
this pull request
Apr 19, 2026
…riables (llvm#184248) Support local, `__device__`, `__shared__`, and `__constant__` variables. Mark device variables as `externally_initialized`. References: llvm#175871, llvm#179278, llvm/clangir#1368, llvm/clangir#1394 --------- Signed-off-by: ZakyHermawan <zaky.hermawan9615@gmail.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Support local,
__device__,__shared__, and__constant__variables.Mark device variables as
externally_initialized.References: #175871, #179278, llvm/clangir#1368, llvm/clangir#1394