https://github.com/steffenlarsen updated https://github.com/llvm/llvm-project/pull/178909
>From 2649cb1197e7f13485770d2db25d5bb464e4485f Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Fri, 30 Jan 2026 04:24:30 -0600 Subject: [PATCH 1/3] [Clang][HIP][CUDA] Validate that variable type fits in address spaces 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. Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/include/clang/AST/ASTContext.h | 11 +++++ .../clang/Basic/DiagnosticSemaKinds.td | 3 ++ clang/lib/Sema/SemaDeclAttr.cpp | 44 +++++++++++++++++++ .../SemaHIP/shared-variable-too-large.hip | 18 ++++++++ ...u-variables-too-large-for-address-space.cl | 10 +++++ 5 files changed, 86 insertions(+) create mode 100644 clang/test/SemaHIP/shared-variable-too-large.hip create mode 100644 clang/test/SemaOpenCL/amdgpu-variables-too-large-for-address-space.cl 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)}} +} >From 16f6b0ef3c9753214724ac094a02e51f8edc7088 Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Mon, 2 Feb 2026 00:45:00 -0600 Subject: [PATCH 2/3] Check bytes instead of bits Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/include/clang/AST/ASTContext.h | 14 ++------------ clang/lib/Sema/SemaDeclAttr.cpp | 5 +++-- clang/test/SemaHIP/shared-variable-too-large.hip | 6 +++--- ...amdgpu-variables-too-large-for-address-space.cl | 6 +++--- 4 files changed, 11 insertions(+), 20 deletions(-) diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h index c9745962674b7..8c57f1497e274 100644 --- a/clang/include/clang/AST/ASTContext.h +++ b/clang/include/clang/AST/ASTContext.h @@ -2681,17 +2681,6 @@ 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); @@ -2709,7 +2698,8 @@ class ASTContext : public RefCountedBase<ASTContext> { CharUnits getTypeSizeInChars(const Type *T) const; std::optional<CharUnits> getTypeSizeInCharsIfKnown(QualType Ty) const { - if (Ty->isIncompleteType() || Ty->isDependentType()) + if (Ty->isIncompleteType() || Ty->isDependentType() || + Ty->isUndeducedType()) return std::nullopt; return getTypeSizeInChars(Ty); } diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 77d4762c927ed..623ea696a1e96 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5144,8 +5144,9 @@ static bool checkCommonVarDeclAddressSpaceAttr(Sema &S, const VarDecl *VD, // 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()) { + std::optional<CharUnits> TSizeInChars = Context.getTypeSizeInCharsIfKnown(T); + if (TSizeInChars && static_cast<uint64_t>(TSizeInChars->getQuantity()) > + MaxSizeForAddrSpace.getZExtValue()) { S.Diag(AL.getLoc(), diag::err_type_too_large_for_address_space) << T << MaxSizeForAddrSpace; return false; diff --git a/clang/test/SemaHIP/shared-variable-too-large.hip b/clang/test/SemaHIP/shared-variable-too-large.hip index e04797d6c9418..4d5040acfeb95 100644 --- a/clang/test/SemaHIP/shared-variable-too-large.hip +++ b/clang/test/SemaHIP/shared-variable-too-large.hip @@ -4,14 +4,14 @@ #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)}} +__shared__ short global_arr[2147483648]; // expected-error {{'short[2147483648]' 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)}} + __shared__ int arr[1073741824]; // expected-error {{'int[1073741824]' 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)}} + __shared__ char arr[4294967296]; // expected-error {{'char[4294967296]' 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 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 index 2a4a60f181024..5aff6729e6183 100644 --- a/clang/test/SemaOpenCL/amdgpu-variables-too-large-for-address-space.cl +++ b/clang/test/SemaOpenCL/amdgpu-variables-too-large-for-address-space.cl @@ -1,10 +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)}} + __private char private_arr[4294967296]; // expected-error {{'__private char[4294967296]' 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)}} + __private int private_arr[1073741824]; // expected-error {{'__private int[1073741824]' is too large for the address space (maximum allowed size of 4'294'967'295 bytes)}} + __local long local_arr[536870912]; // expected-error {{'__local long[536870912]' is too large for the address space (maximum allowed size of 4'294'967'295 bytes)}} } >From d2328a90cdcc71c4c5d1579399a5a430f7fc145d Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Mon, 2 Feb 2026 10:43:11 -0600 Subject: [PATCH 3/3] Move general address space checks Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/include/clang/AST/ASTContext.h | 2 +- clang/include/clang/Sema/Sema.h | 11 +++++ clang/lib/Sema/SemaDecl.cpp | 6 +++ clang/lib/Sema/SemaDeclAttr.cpp | 45 ++----------------- clang/lib/Sema/SemaType.cpp | 18 ++++++++ .../SemaHIP/shared-variable-too-large.hip | 3 +- 6 files changed, 42 insertions(+), 43 deletions(-) diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h index 8c57f1497e274..8dd99b71fb99c 100644 --- a/clang/include/clang/AST/ASTContext.h +++ b/clang/include/clang/AST/ASTContext.h @@ -2699,7 +2699,7 @@ class ASTContext : public RefCountedBase<ASTContext> { std::optional<CharUnits> getTypeSizeInCharsIfKnown(QualType Ty) const { if (Ty->isIncompleteType() || Ty->isDependentType() || - Ty->isUndeducedType()) + Ty->isUndeducedType() || Ty->isSizelessType()) return std::nullopt; return getTypeSizeInChars(Ty); } diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 7b3479bbc3677..3813f3b289edf 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -15340,6 +15340,17 @@ class Sema final : public SemaBase { bool AllowArrayTypes, bool OverrideExisting); + /// Check whether the given variable declaration has a size that fits within + /// the address space it is declared in. This issues a diagnostic if not. + /// + /// \param VD The variable declaration to check the size of. + /// + /// \param AS The address space to check the size of \p VD against. + /// + /// \returns true if the variable's size fits within the address space, false + /// otherwise. + bool CheckVarDeclSizeAddressSpace(const VarDecl *VD, LangAS AS); + /// Get the type of expression E, triggering instantiation to complete the /// type if necessary -- that is, if the expression refers to a templated /// static data member of incomplete array type. diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 907b7b367f19b..687e4db9ca00a 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -9194,6 +9194,12 @@ void Sema::CheckVariableDeclarationType(VarDecl *NewVD) { RISCV().checkRVVTypeSupport(T, NewVD->getLocation(), cast<Decl>(CurContext), CallerFeatureMap); } + + if (T.hasAddressSpace() && + !CheckVarDeclSizeAddressSpace(NewVD, T.getAddressSpace())) { + NewVD->setInvalidDecl(); + return; + } } bool Sema::CheckVariableDeclaration(VarDecl *NewVD, LookupResult &Previous) { diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 623ea696a1e96..f03b2aecacabf 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5134,34 +5134,13 @@ 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<CharUnits> TSizeInChars = Context.getTypeSizeInCharsIfKnown(T); - if (TSizeInChars && static_cast<uint64_t>(TSizeInChars->getQuantity()) > - 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)) + if (!S.CheckVarDeclSizeAddressSpace(VD, LangAS::cuda_constant)) return; // constexpr variable may already get an implicit constant attr, which should // be replaced by the explicit constant attr. @@ -5182,7 +5161,7 @@ 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)) + if (!S.CheckVarDeclSizeAddressSpace(VD, LangAS::cuda_shared)) return; if (S.getLangOpts().CUDA && VD->hasLocalStorage() && S.CUDA().DiagIfHostCode(AL.getLoc(), diag::err_cuda_host_shared) @@ -5233,7 +5212,7 @@ 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)) + if (!S.CheckVarDeclSizeAddressSpace(VD, LangAS::cuda_device)) return; } @@ -5251,7 +5230,7 @@ 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)) + if (!S.CheckVarDeclSizeAddressSpace(VD, LangAS::cuda_device)) return; } if (!D->hasAttr<HIPManagedAttr>()) @@ -8164,22 +8143,6 @@ 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/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index fa4dcdd9e1422..c58e0344a9895 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -7424,6 +7424,24 @@ bool Sema::CheckImplicitNullabilityTypeSpecifier(QualType &Type, /*isContextSensitive*/ false, AllowArrayTypes, OverrideExisting); } +bool Sema::CheckVarDeclSizeAddressSpace(const VarDecl *VD, LangAS AS) { + 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<CharUnits> TSizeInChars = Context.getTypeSizeInCharsIfKnown(T); + if (TSizeInChars && static_cast<uint64_t>(TSizeInChars->getQuantity()) > + MaxSizeForAddrSpace.getZExtValue()) { + Diag(VD->getLocation(), diag::err_type_too_large_for_address_space) + << T << MaxSizeForAddrSpace; + return false; + } + + return true; +} + /// Check the application of the Objective-C '__kindof' qualifier to /// the given type. static bool checkObjCKindOfType(TypeProcessingState &state, QualType &type, diff --git a/clang/test/SemaHIP/shared-variable-too-large.hip b/clang/test/SemaHIP/shared-variable-too-large.hip index 4d5040acfeb95..40b1acef04158 100644 --- a/clang/test/SemaHIP/shared-variable-too-large.hip +++ b/clang/test/SemaHIP/shared-variable-too-large.hip @@ -4,7 +4,8 @@ #define __device__ __attribute__((device)) #define __shared__ __attribute__((shared)) -__shared__ short global_arr[2147483648]; // expected-error {{'short[2147483648]' is too large for the address space (maximum allowed size of 4'294'967'295 bytes)}} +__shared__ short global_arr1[2147483648]; // expected-error {{'short[2147483648]' is too large for the address space (maximum allowed size of 4'294'967'295 bytes)}} +[[clang::loader_uninitialized]] short [[clang::address_space(3)]] global_arr1[2147483648]; // expected-error {{'__attribute__((address_space(3))) short[2147483648]' is too large for the address space (maximum allowed size of 4'294'967'295 bytes)}} __device__ void func() { __shared__ int arr[1073741824]; // expected-error {{'int[1073741824]' is too large for the address space (maximum allowed size of 4'294'967'295 bytes)}} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
