https://github.com/steffenlarsen created https://github.com/llvm/llvm-project/pull/180756
This commit adds a check for the size of variables in address spaces after template instantiation. This ensures that diagnostics are also issued when the size of a variable is not known prior to template instantiation for dependent variables. This is a follow-up from https://github.com/llvm/llvm-project/pull/178909. >From 6ce0f51817f7694137d5e472e1557708c1005a2c Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Tue, 10 Feb 2026 09:27:27 -0600 Subject: [PATCH] [Clang][HIP][CUDA] Validate size for variable in address space after instantiation This commit adds a check for the size of variables in address spaces after template instantiation. This ensures that diagnostics are also issued when the size of a variable is not known prior to template instantiation for dependent variables. This is a follow-up from https://github.com/llvm/llvm-project/pull/178909. Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/lib/Sema/SemaTemplateInstantiateDecl.cpp | 14 ++++++++++++++ clang/test/SemaHIP/shared-variable-too-large.hip | 10 +++++++--- 2 files changed, 21 insertions(+), 3 deletions(-) diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index e74c41517ecbf..760e0c2c33a1c 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -1031,6 +1031,20 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs, continue; } + // Check address space sizes for CUDA/HIP variable attributes now that we + // know the type of the instantiation. + if (auto *NewVar = dyn_cast<VarDecl>(New)) { + if ((isa<CUDADeviceAttr>(TmplAttr) || isa<HIPManagedAttr>(TmplAttr)) && + !CheckVarDeclSizeAddressSpace(NewVar, LangAS::cuda_device)) + return; + if (isa<CUDASharedAttr>(TmplAttr) && + !CheckVarDeclSizeAddressSpace(NewVar, LangAS::cuda_shared)) + return; + if (isa<CUDAConstantAttr>(TmplAttr) && + !CheckVarDeclSizeAddressSpace(NewVar, LangAS::cuda_constant)) + return; + } + assert(!TmplAttr->isPackExpansion()); if (TmplAttr->isLateParsed() && LateAttrs) { // Late parsed attributes must be instantiated and attached after the diff --git a/clang/test/SemaHIP/shared-variable-too-large.hip b/clang/test/SemaHIP/shared-variable-too-large.hip index eff5f8f6a7900..c877f435f4d33 100644 --- a/clang/test/SemaHIP/shared-variable-too-large.hip +++ b/clang/test/SemaHIP/shared-variable-too-large.hip @@ -14,10 +14,14 @@ __device__ void func() { __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)}} } +template <long long N> __device__ void dep_func() { + __shared__ char max_size_arr[N]; + __shared__ char too_large_arr[N+1]; // expected-error {{'char[4294967296]' 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. + dep_func<4294967295>(); // expected-note {{in instantiation of function template specialization 'dep_func<4294967295LL>' requested here}} +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
