llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Steffen Larsen (steffenlarsen) <details> <summary>Changes</summary> 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. --- Full diff: https://github.com/llvm/llvm-project/pull/180756.diff 2 Files Affected: - (modified) clang/lib/Sema/SemaTemplateInstantiateDecl.cpp (+14) - (modified) clang/test/SemaHIP/shared-variable-too-large.hip (+7-3) ``````````diff 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}} +} `````````` </details> https://github.com/llvm/llvm-project/pull/180756 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
