[clang] [Clang][HIP][CUDA] Validate that variable type fits in address spaces (PR #178909)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Jan 30 07:56:10 PST 2026
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Steffen Larsen (steffenlarsen)
<details>
<summary>Changes</summary>
Currently, Clang only checks arrays and structures for size at a top-level view, that is it does not consider whether they will fit in the address space when applying the address space attribute. This can lead to situations where a variable is declared in an address space but its type is too large to fit in that address space, leading to potentially invalid modules.
This patch proposes a fix for this by checking the size of the type against the maximum size that can be addressed in the given address space when applying the address space attribute.
This does not currently handle instantiations of dependent variables, as the attributes are not re-processesd at that time. This is planned for further investigation and a follow-up patch.
---
Full diff: https://github.com/llvm/llvm-project/pull/178909.diff
5 Files Affected:
- (modified) clang/include/clang/AST/ASTContext.h (+11)
- (modified) clang/include/clang/Basic/DiagnosticSemaKinds.td (+3)
- (modified) clang/lib/Sema/SemaDeclAttr.cpp (+44)
- (added) clang/test/SemaHIP/shared-variable-too-large.hip (+18)
- (added) clang/test/SemaOpenCL/amdgpu-variables-too-large-for-address-space.cl (+10)
``````````diff
diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h
index 68205dd1c1fd9..c9745962674b7 100644
--- a/clang/include/clang/AST/ASTContext.h
+++ b/clang/include/clang/AST/ASTContext.h
@@ -2681,6 +2681,17 @@ class ASTContext : public RefCountedBase<ASTContext> {
uint64_t getTypeSize(QualType T) const { return getTypeInfo(T).Width; }
uint64_t getTypeSize(const Type *T) const { return getTypeInfo(T).Width; }
+ std::optional<uint64_t> getTypeSizeIfKnown(QualType Ty) const {
+ if (Ty->isIncompleteType() || Ty->isDependentType() ||
+ Ty->isUndeducedType())
+ return std::nullopt;
+ return getTypeSize(Ty);
+ }
+
+ std::optional<uint64_t> getTypeSizeIfKnown(const Type *Ty) const {
+ return getTypeSizeIfKnown(QualType(Ty, 0));
+ }
+
/// Return the size of the character type, in bits.
uint64_t getCharWidth() const {
return getTypeSize(CharTy);
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 807440c107897..cc57ea19c1743 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6558,6 +6558,9 @@ def err_vm_func_decl : Error<
def err_array_too_large : Error<
"array is too large (%0 elements)">;
+def err_type_too_large_for_address_space : Error<
+ "%0 is too large for the address space (maximum allowed size of %1 bytes)">;
+
def err_typecheck_negative_array_size : Error<"array size is negative">;
def warn_typecheck_function_qualifiers_ignored : Warning<
"'%0' qualifier on function type %1 has no effect">,
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index bee42cce09aca..77d4762c927ed 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -5134,12 +5134,34 @@ static void handleOptimizeNoneAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
D->addAttr(Optnone);
}
+static bool checkCommonVarDeclAddressSpaceAttr(Sema &S, const VarDecl *VD,
+ LangAS AS,
+ const ParsedAttr &AL) {
+ const ASTContext &Context = S.getASTContext();
+ QualType T = VD->getType();
+
+ // Check that the variable's type can fit in the specified address space. This
+ // is determined by how far a pointer in that address space can reach.
+ llvm::APInt MaxSizeForAddrSpace =
+ llvm::APInt::getMaxValue(Context.getTargetInfo().getPointerWidth(AS));
+ std::optional<uint64_t> TSizeInChars = Context.getTypeSizeIfKnown(T);
+ if (TSizeInChars && *TSizeInChars > MaxSizeForAddrSpace.getZExtValue()) {
+ S.Diag(AL.getLoc(), diag::err_type_too_large_for_address_space)
+ << T << MaxSizeForAddrSpace;
+ return false;
+ }
+
+ return true;
+}
+
static void handleConstantAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
const auto *VD = cast<VarDecl>(D);
if (VD->hasLocalStorage()) {
S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev);
return;
}
+ if (!checkCommonVarDeclAddressSpaceAttr(S, VD, LangAS::cuda_constant, AL))
+ return;
// constexpr variable may already get an implicit constant attr, which should
// be replaced by the explicit constant attr.
if (auto *A = D->getAttr<CUDAConstantAttr>()) {
@@ -5159,6 +5181,8 @@ static void handleSharedAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
S.Diag(AL.getLoc(), diag::err_cuda_extern_shared) << VD;
return;
}
+ if (!checkCommonVarDeclAddressSpaceAttr(S, VD, LangAS::cuda_shared, AL))
+ return;
if (S.getLangOpts().CUDA && VD->hasLocalStorage() &&
S.CUDA().DiagIfHostCode(AL.getLoc(), diag::err_cuda_host_shared)
<< S.CUDA().CurrentTarget())
@@ -5208,6 +5232,8 @@ static void handleDeviceAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev);
return;
}
+ if (!checkCommonVarDeclAddressSpaceAttr(S, VD, LangAS::cuda_device, AL))
+ return;
}
if (auto *A = D->getAttr<CUDADeviceAttr>()) {
@@ -5224,6 +5250,8 @@ static void handleManagedAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev);
return;
}
+ if (!checkCommonVarDeclAddressSpaceAttr(S, VD, LangAS::cuda_device, AL))
+ return;
}
if (!D->hasAttr<HIPManagedAttr>())
D->addAttr(::new (S.Context) HIPManagedAttr(S.Context, AL));
@@ -8135,6 +8163,22 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL,
case ParsedAttr::AT_GCCStruct:
handleGCCStructAttr(S, D, AL);
break;
+
+ case ParsedAttr::AT_OpenCLConstantAddressSpace:
+ case ParsedAttr::AT_OpenCLGlobalAddressSpace:
+ case ParsedAttr::AT_OpenCLGlobalDeviceAddressSpace:
+ case ParsedAttr::AT_OpenCLGlobalHostAddressSpace:
+ case ParsedAttr::AT_OpenCLLocalAddressSpace:
+ case ParsedAttr::AT_OpenCLPrivateAddressSpace:
+ case ParsedAttr::AT_OpenCLGenericAddressSpace: {
+ // OpenCL address space attributes are mainly checked during type
+ // checking. However, we need to do some common address space checking.
+ if (auto *VD = dyn_cast<VarDecl>(D)) {
+ LangAS AS = S.getLangOpts().SYCLIsDevice ? AL.asSYCLLangAS()
+ : AL.asOpenCLLangAS();
+ checkCommonVarDeclAddressSpaceAttr(S, VD, AS, AL);
+ }
+ }
}
}
diff --git a/clang/test/SemaHIP/shared-variable-too-large.hip b/clang/test/SemaHIP/shared-variable-too-large.hip
new file mode 100644
index 0000000000000..e04797d6c9418
--- /dev/null
+++ b/clang/test/SemaHIP/shared-variable-too-large.hip
@@ -0,0 +1,18 @@
+// RUN: %clang_cc1 -fsyntax-only -triple amdgcn -target-cpu gfx90a -verify %s -fcuda-is-device
+
+#define __global__ __attribute__((global))
+#define __device__ __attribute__((device))
+#define __shared__ __attribute__((shared))
+
+__shared__ short global_arr[2147483647]; // expected-error {{'short[2147483647]' is too large for the address space (maximum allowed size of 4'294'967'295 bytes)}}
+
+__device__ void func() {
+ __shared__ int arr[1073741823]; // expected-error {{'int[1073741823]' is too large for the address space (maximum allowed size of 4'294'967'295 bytes)}}
+}
+
+__global__ void kernel() {
+ __shared__ char arr[4294967295]; // expected-error {{'char[4294967295]' is too large for the address space (maximum allowed size of 4'294'967'295 bytes)}}
+}
+
+// TODO: The implementation of the __shared__ attribute doesn't check the
+// instantiation of dependent variables.
diff --git a/clang/test/SemaOpenCL/amdgpu-variables-too-large-for-address-space.cl b/clang/test/SemaOpenCL/amdgpu-variables-too-large-for-address-space.cl
new file mode 100644
index 0000000000000..2a4a60f181024
--- /dev/null
+++ b/clang/test/SemaOpenCL/amdgpu-variables-too-large-for-address-space.cl
@@ -0,0 +1,10 @@
+// RUN: %clang_cc1 -triple amdgcn-- -verify -fsyntax-only %s
+
+void func() {
+ __private char private_arr[4294967295]; // expected-error {{'__private char[4294967295]' is too large for the address space (maximum allowed size of 4'294'967'295 bytes)}}
+}
+
+void kernel kernel_func() {
+ __private int private_arr[1073741823]; // expected-error {{'__private int[1073741823]' is too large for the address space (maximum allowed size of 4'294'967'295 bytes)}}
+ __local long local_arr[536870911]; // expected-error {{'__local long[536870911]' is too large for the address space (maximum allowed size of 4'294'967'295 bytes)}}
+}
``````````
</details>
https://github.com/llvm/llvm-project/pull/178909
More information about the cfe-commits
mailing list