Author: Steffen Larsen Date: 2026-02-09T08:12:34+01:00 New Revision: aa808967eb96995f58288bb3bc0f37516ca70cbe
URL: https://github.com/llvm/llvm-project/commit/aa808967eb96995f58288bb3bc0f37516ca70cbe DIFF: https://github.com/llvm/llvm-project/commit/aa808967eb96995f58288bb3bc0f37516ca70cbe.diff LOG: [Clang][HIP][CUDA] Validate that variable type fits in address spaces (#178909) 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]> Co-authored-by: Steffen Holst Larsen <[email protected]> Added: clang/test/SemaHIP/shared-variable-too-large.hip clang/test/SemaOpenCL/amdgpu-variables-too-large-for-address-space.cl Modified: clang/include/clang/AST/ASTContext.h clang/include/clang/Basic/DiagnosticSemaKinds.td clang/include/clang/Sema/Sema.h clang/lib/Sema/SemaDecl.cpp clang/lib/Sema/SemaDeclAttr.cpp clang/lib/Sema/SemaType.cpp Removed: ################################################################################ diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h index c8d6de1689512..6b819de2fb36d 100644 --- a/clang/include/clang/AST/ASTContext.h +++ b/clang/include/clang/AST/ASTContext.h @@ -2698,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() || 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 9171cb69d7829..96a1e3481b2ca 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 fe4616d89df89..9ae2fa52a441a 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -15350,6 +15350,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 3b2c93b9fe7b5..7af6ce62d08dd 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 c19ec95ee7b7f..5dbff18fff7a9 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<CUDAConstantAttr>()) { @@ -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<CUDADeviceAttr>()) { @@ -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<HIPManagedAttr>()) D->addAttr(::new (S.Context) HIPManagedAttr(S.Context, AL)); diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index 28d1d63ff7acf..348823ab2e9ca 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 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)}} +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
