diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h index 68205dd1c1fd9..8dd99b71fb99c 100644 --- a/clang/include/clang/AST/ASTContext.h +++ b/clang/include/clang/AST/ASTContext.h @@ -2698,7 +2698,8 @@ class ASTContext : public RefCountedBase { CharUnits getTypeSizeInChars(const Type *T) const; std::optional getTypeSizeInCharsIfKnown(QualType Ty) const { - if (Ty->isIncompleteType() || Ty->isDependentType()) + if (Ty->isIncompleteType() || Ty->isDependentType() || + Ty->isUndeducedType() || Ty->isSizelessType()) return std::nullopt; return getTypeSizeInChars(Ty); } 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/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(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 bee42cce09aca..f03b2aecacabf 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5140,6 +5140,8 @@ static void handleConstantAttr(Sema &S, Decl *D, const ParsedAttr &AL) { S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev); return; } + 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. if (auto *A = D->getAttr()) { @@ -5159,6 +5161,8 @@ static void handleSharedAttr(Sema &S, Decl *D, const ParsedAttr &AL) { S.Diag(AL.getLoc(), diag::err_cuda_extern_shared) << VD; return; } + if (!S.CheckVarDeclSizeAddressSpace(VD, LangAS::cuda_shared)) + return; if (S.getLangOpts().CUDA && VD->hasLocalStorage() && S.CUDA().DiagIfHostCode(AL.getLoc(), diag::err_cuda_host_shared) << S.CUDA().CurrentTarget()) @@ -5208,6 +5212,8 @@ static void handleDeviceAttr(Sema &S, Decl *D, const ParsedAttr &AL) { S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev); return; } + if (!S.CheckVarDeclSizeAddressSpace(VD, LangAS::cuda_device)) + return; } if (auto *A = D->getAttr()) { @@ -5224,6 +5230,8 @@ static void handleManagedAttr(Sema &S, Decl *D, const ParsedAttr &AL) { S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev); return; } + if (!S.CheckVarDeclSizeAddressSpace(VD, LangAS::cuda_device)) + return; } if (!D->hasAttr()) D->addAttr(::new (S.Context) HIPManagedAttr(S.Context, 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 TSizeInChars = Context.getTypeSizeInCharsIfKnown(T); + if (TSizeInChars && static_cast(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 new file mode 100644 index 0000000000000..eff5f8f6a7900 --- /dev/null +++ b/clang/test/SemaHIP/shared-variable-too-large.hip @@ -0,0 +1,23 @@ +// 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 max_size_global_arr1[2147483647]; +[[clang::loader_uninitialized]] short [[clang::address_space(3)]] max_size_global_arr2[2147483647]; +__shared__ short too_large_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)]] too_large_global_arr2[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 max_size_arr[1073741823]; + __shared__ int too_large_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 max_size_arr[4294967295]; + __shared__ char too_large_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 +// 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..a0c2b8838761b --- /dev/null +++ b/clang/test/SemaOpenCL/amdgpu-variables-too-large-for-address-space.cl @@ -0,0 +1,13 @@ +// RUN: %clang_cc1 -triple amdgcn-- -verify -fsyntax-only %s + +void func() { + __private char max_size_private_arr[4294967295]; + __private char too_large_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 max_size_private_arr[1073741823]; + __local long max_size_local_arr[536870911]; + __private int too_large_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 too_large_local_arr[536870912]; // expected-error {{'__local long[536870912]' is too large for the address space (maximum allowed size of 4'294'967'295 bytes)}} +}