[clang] [Clang][HIP][CUDA] Validate size for variable in address space after instantiation (PR #180756)

[via cfe-commits]

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