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

[Steffen Larsen via cfe-commits]

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 <HolstLarsen.Steffen at amd.com>
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 <HolstLarsen.Steffen at amd.com>
---
 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}}
+}