[clang] [Clang/AMDGPU] Zero sized arrays not allowed in HIP device code. (PR #113470)

Vigneshwar Jayakumar via cfe-commits cfe-commits at lists.llvm.org
Wed Nov 20 10:31:38 PST 2024


https://github.com/VigneshwarJ updated https://github.com/llvm/llvm-project/pull/113470

>From 0e2ee524f5b5c19169e446c55a386a00cfb0f6bc Mon Sep 17 00:00:00 2001
From: vigneshwar jayakumar <vigneshwar.jayakumar at amd.com>
Date: Wed, 23 Oct 2024 09:20:16 -0500
Subject: [PATCH 1/5] [Clang/AMDGPU] Zero sized arrays not allowed in HIP
 device code.

Added diagnosis to throw error when zero sized arrays are used in
the HIP device code. SWDEV-449592
---
 .../clang/Basic/DiagnosticSemaKinds.td        |  2 +-
 clang/lib/Sema/SemaType.cpp                   | 11 ++++++++
 .../test/SemaHIP/zero-sized-device-array.hip  | 25 +++++++++++++++++++
 3 files changed, 37 insertions(+), 1 deletion(-)
 create mode 100644 clang/test/SemaHIP/zero-sized-device-array.hip

diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 8e4718008ece72..b5fad40294c368 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6251,7 +6251,7 @@ def err_typecheck_invalid_restrict_invalid_pointee : Error<
 def ext_typecheck_zero_array_size : Extension<
   "zero size arrays are an extension">, InGroup<ZeroLengthArray>;
 def err_typecheck_zero_array_size : Error<
-  "zero-length arrays are not permitted in %select{C++|SYCL device code}0">;
+  "zero-length arrays are not permitted in %select{C++|SYCL device code|HIP device code}0">;
 def err_array_size_non_int : Error<"size of array has non-integer type %0">;
 def err_init_element_not_constant : Error<
   "initializer element is not a compile-time constant">;
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index 6387fe9f1129ba..3f940102da51d2 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -2259,6 +2259,17 @@ QualType Sema::BuildArrayType(QualType T, ArraySizeModifier ASM,
              isSFINAEContext() ? diag::err_typecheck_zero_array_size
                                : diag::ext_typecheck_zero_array_size)
             << 0 << ArraySize->getSourceRange();
+
+        // zero sized static arrays are not allowed in HIP device functions
+        if (LangOpts.HIP && LangOpts.CUDAIsDevice) {
+          auto *FD = dyn_cast_or_null<FunctionDecl>(CurContext);
+          if (FD && (FD->hasAttr<CUDADeviceAttr>() ||
+                     FD->hasAttr<CUDAGlobalAttr>())) {
+            Diag(ArraySize->getBeginLoc(), diag::err_typecheck_zero_array_size)
+                << 2 << ArraySize->getSourceRange();
+            return QualType();
+          }
+        }
       }
 
       // Is the array too large?
diff --git a/clang/test/SemaHIP/zero-sized-device-array.hip b/clang/test/SemaHIP/zero-sized-device-array.hip
new file mode 100644
index 00000000000000..31fc943f5ae75b
--- /dev/null
+++ b/clang/test/SemaHIP/zero-sized-device-array.hip
@@ -0,0 +1,25 @@
+// REQUIRES: amdgpu-registered-target
+// RUN:  %clang_cc1 -fsyntax-only -x hip -fcuda-is-device -verify -triple amdgcn %s
+#define __device__ __attribute__((device))
+#define __host__ __attribute__((host))
+#define __global__ __attribute__((global))
+#define __shared__ __attribute__((shared))
+
+__global__ void global_fun() {
+    float array[0];  // expected-error {{zero-length arrays are not permitted in HIP device code}}
+}
+
+// should not throw error for host side code.
+__host__ void host_fun() {
+    float array[0];
+}
+
+__host__ __device__ void host_dev_fun()
+{
+    float array[0]; // expected-error {{zero-length arrays are not permitted in HIP device code}}
+}
+
+__device__ void device_fun()
+{
+    __shared__ float array[0]; // expected-error {{zero-length arrays are not permitted in HIP device code}}
+}

>From fed2349c62763592866291d35b970d024858d8b1 Mon Sep 17 00:00:00 2001
From: vigneshwar jayakumar <vigneshwar.jayakumar at amd.com>
Date: Fri, 25 Oct 2024 14:52:39 -0500
Subject: [PATCH 2/5] fix review comments

moved code to SemaVarDecl also check the pointer types to figure out
its within any typedefs or pointers.
---
 clang/lib/Sema/SemaDecl.cpp                   | 25 +++++++++++++++++++
 clang/lib/Sema/SemaType.cpp                   | 11 --------
 .../test/SemaHIP/zero-sized-device-array.hip  | 12 +++++++++
 3 files changed, 37 insertions(+), 11 deletions(-)

diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 229c9080d558ec..55e5209d370745 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -8714,6 +8714,31 @@ void Sema::CheckVariableDeclarationType(VarDecl *NewVD) {
     }
   }
 
+  // zero sized static arrays are not allowed in HIP device functions
+  if (LangOpts.CUDAIsDevice && LangOpts.HIP) {
+    if (FunctionDecl *FD = getCurFunctionDecl();
+        FD &&
+        (FD->hasAttr<CUDADeviceAttr>() || FD->hasAttr<CUDAGlobalAttr>())) {
+
+      auto Check = [&](QualType TypeToCheck, const VarDecl *VD) {
+        if (const ConstantArrayType *ArrayT =
+                getASTContext().getAsConstantArrayType(TypeToCheck);
+            ArrayT && ArrayT->isZeroSize()) {
+          Diag(VD->getLocation(), diag::err_typecheck_zero_array_size) << 2;
+        }
+      };
+      QualType NextTy = NewVD->getType();
+      while (NextTy->isAnyPointerType() || NextTy->isArrayType() ||
+             NextTy->isReferenceType()) {
+        if (NextTy->isArrayType()) {
+          Check(NextTy, NewVD);
+          break;
+        } else
+          NextTy = NextTy->getPointeeType();
+      }
+    }
+  }
+
   bool isVM = T->isVariablyModifiedType();
   if (isVM || NewVD->hasAttr<CleanupAttr>() ||
       NewVD->hasAttr<BlocksAttr>())
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index 3f940102da51d2..6387fe9f1129ba 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -2259,17 +2259,6 @@ QualType Sema::BuildArrayType(QualType T, ArraySizeModifier ASM,
              isSFINAEContext() ? diag::err_typecheck_zero_array_size
                                : diag::ext_typecheck_zero_array_size)
             << 0 << ArraySize->getSourceRange();
-
-        // zero sized static arrays are not allowed in HIP device functions
-        if (LangOpts.HIP && LangOpts.CUDAIsDevice) {
-          auto *FD = dyn_cast_or_null<FunctionDecl>(CurContext);
-          if (FD && (FD->hasAttr<CUDADeviceAttr>() ||
-                     FD->hasAttr<CUDAGlobalAttr>())) {
-            Diag(ArraySize->getBeginLoc(), diag::err_typecheck_zero_array_size)
-                << 2 << ArraySize->getSourceRange();
-            return QualType();
-          }
-        }
       }
 
       // Is the array too large?
diff --git a/clang/test/SemaHIP/zero-sized-device-array.hip b/clang/test/SemaHIP/zero-sized-device-array.hip
index 31fc943f5ae75b..d9c74e6fbe536e 100644
--- a/clang/test/SemaHIP/zero-sized-device-array.hip
+++ b/clang/test/SemaHIP/zero-sized-device-array.hip
@@ -5,8 +5,13 @@
 #define __global__ __attribute__((global))
 #define __shared__ __attribute__((shared))
 
+typedef float ZEROARR[0];
+
 __global__ void global_fun() {
+    extern __shared__ float externArray[];
+    ZEROARR TypeDef; // expected-error {{zero-length arrays are not permitted in HIP device code}}
     float array[0];  // expected-error {{zero-length arrays are not permitted in HIP device code}}
+    ZEROARR *Ptr;    // expected-error {{zero-length arrays are not permitted in HIP device code}}
 }
 
 // should not throw error for host side code.
@@ -14,6 +19,12 @@ __host__ void host_fun() {
     float array[0];
 }
 
+template <typename Ty, unsigned Size>
+__device__ void templated()
+{
+   Ty arr[Size];  // expected-error {{zero-length arrays are not permitted in HIP device code}}
+}
+
 __host__ __device__ void host_dev_fun()
 {
     float array[0]; // expected-error {{zero-length arrays are not permitted in HIP device code}}
@@ -22,4 +33,5 @@ __host__ __device__ void host_dev_fun()
 __device__ void device_fun()
 {
     __shared__ float array[0]; // expected-error {{zero-length arrays are not permitted in HIP device code}}
+    templated<int,0>(); // expected-note {{in instantiation of function template specialization 'templated<int, 0U>' requested here}} 
 }

>From 7ee4272ea27bef09762943c64c48c7287f99b19b Mon Sep 17 00:00:00 2001
From: vigneshwar jayakumar <vigneshwar.jayakumar at amd.com>
Date: Mon, 28 Oct 2024 10:05:32 -0500
Subject: [PATCH 3/5] review comment fixes

---
 clang/lib/Sema/SemaDecl.cpp                    | 16 +++-------------
 clang/test/SemaHIP/zero-sized-device-array.hip |  1 -
 2 files changed, 3 insertions(+), 14 deletions(-)

diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 55e5209d370745..623ff2f7750f7f 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -8719,22 +8719,12 @@ void Sema::CheckVariableDeclarationType(VarDecl *NewVD) {
     if (FunctionDecl *FD = getCurFunctionDecl();
         FD &&
         (FD->hasAttr<CUDADeviceAttr>() || FD->hasAttr<CUDAGlobalAttr>())) {
-
-      auto Check = [&](QualType TypeToCheck, const VarDecl *VD) {
+      if (QualType NextTy = NewVD->getType(); NextTy->isArrayType()) {
         if (const ConstantArrayType *ArrayT =
-                getASTContext().getAsConstantArrayType(TypeToCheck);
+                getASTContext().getAsConstantArrayType(NextTy);
             ArrayT && ArrayT->isZeroSize()) {
-          Diag(VD->getLocation(), diag::err_typecheck_zero_array_size) << 2;
+          Diag(NewVD->getLocation(), diag::err_typecheck_zero_array_size) << 2;
         }
-      };
-      QualType NextTy = NewVD->getType();
-      while (NextTy->isAnyPointerType() || NextTy->isArrayType() ||
-             NextTy->isReferenceType()) {
-        if (NextTy->isArrayType()) {
-          Check(NextTy, NewVD);
-          break;
-        } else
-          NextTy = NextTy->getPointeeType();
       }
     }
   }
diff --git a/clang/test/SemaHIP/zero-sized-device-array.hip b/clang/test/SemaHIP/zero-sized-device-array.hip
index d9c74e6fbe536e..4b4fa04403a722 100644
--- a/clang/test/SemaHIP/zero-sized-device-array.hip
+++ b/clang/test/SemaHIP/zero-sized-device-array.hip
@@ -11,7 +11,6 @@ __global__ void global_fun() {
     extern __shared__ float externArray[];
     ZEROARR TypeDef; // expected-error {{zero-length arrays are not permitted in HIP device code}}
     float array[0];  // expected-error {{zero-length arrays are not permitted in HIP device code}}
-    ZEROARR *Ptr;    // expected-error {{zero-length arrays are not permitted in HIP device code}}
 }
 
 // should not throw error for host side code.

>From a7065f43fa9b30cb82d5a3b28e90d9bb2fde197b Mon Sep 17 00:00:00 2001
From: vigneshwar jayakumar <vigneshwar.jayakumar at amd.com>
Date: Thu, 7 Nov 2024 08:38:49 -0600
Subject: [PATCH 4/5] review comments

---
 clang/lib/Sema/SemaDecl.cpp | 17 +++++++----------
 1 file changed, 7 insertions(+), 10 deletions(-)

diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 623ff2f7750f7f..82da28a4dbbc8f 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -8715,16 +8715,13 @@ void Sema::CheckVariableDeclarationType(VarDecl *NewVD) {
   }
 
   // zero sized static arrays are not allowed in HIP device functions
-  if (LangOpts.CUDAIsDevice && LangOpts.HIP) {
-    if (FunctionDecl *FD = getCurFunctionDecl();
-        FD &&
-        (FD->hasAttr<CUDADeviceAttr>() || FD->hasAttr<CUDAGlobalAttr>())) {
-      if (QualType NextTy = NewVD->getType(); NextTy->isArrayType()) {
-        if (const ConstantArrayType *ArrayT =
-                getASTContext().getAsConstantArrayType(NextTy);
-            ArrayT && ArrayT->isZeroSize()) {
-          Diag(NewVD->getLocation(), diag::err_typecheck_zero_array_size) << 2;
-        }
+  if (getLangOpts().HIP &&
+      DeclAttrsMatchCUDAMode(getLangOpts(), getCurFunctionDecl())) {
+    if (QualType NextTy = NewVD->getType(); NextTy->isArrayType()) {
+      if (const ConstantArrayType *ArrayT =
+              getASTContext().getAsConstantArrayType(NextTy);
+          ArrayT && ArrayT->isZeroSize()) {
+        Diag(NewVD->getLocation(), diag::err_typecheck_zero_array_size) << 2;
       }
     }
   }

>From d613d592d80bf16dac62ea64fbbe085187c14b6c Mon Sep 17 00:00:00 2001
From: vigneshwar jayakumar <vigneshwar.jayakumar at amd.com>
Date: Wed, 20 Nov 2024 12:28:28 -0600
Subject: [PATCH 5/5] review changes

---
 clang/lib/Sema/SemaDecl.cpp                    | 9 +++++----
 clang/test/SemaHIP/zero-sized-device-array.hip | 2 ++
 2 files changed, 7 insertions(+), 4 deletions(-)

diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 82da28a4dbbc8f..8bfc0c9035b554 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -8715,11 +8715,12 @@ void Sema::CheckVariableDeclarationType(VarDecl *NewVD) {
   }
 
   // zero sized static arrays are not allowed in HIP device functions
-  if (getLangOpts().HIP &&
-      DeclAttrsMatchCUDAMode(getLangOpts(), getCurFunctionDecl())) {
-    if (QualType NextTy = NewVD->getType(); NextTy->isArrayType()) {
+  if (getLangOpts().HIP && LangOpts.CUDAIsDevice) {
+    if (FunctionDecl *FD = getCurFunctionDecl();
+        FD &&
+        (FD->hasAttr<CUDADeviceAttr>() || FD->hasAttr<CUDAGlobalAttr>())) {
       if (const ConstantArrayType *ArrayT =
-              getASTContext().getAsConstantArrayType(NextTy);
+              getASTContext().getAsConstantArrayType(T);
           ArrayT && ArrayT->isZeroSize()) {
         Diag(NewVD->getLocation(), diag::err_typecheck_zero_array_size) << 2;
       }
diff --git a/clang/test/SemaHIP/zero-sized-device-array.hip b/clang/test/SemaHIP/zero-sized-device-array.hip
index 4b4fa04403a722..8b25ec38fb01ac 100644
--- a/clang/test/SemaHIP/zero-sized-device-array.hip
+++ b/clang/test/SemaHIP/zero-sized-device-array.hip
@@ -7,6 +7,8 @@
 
 typedef float ZEROARR[0];
 
+float global_array[0];
+
 __global__ void global_fun() {
     extern __shared__ float externArray[];
     ZEROARR TypeDef; // expected-error {{zero-length arrays are not permitted in HIP device code}}



More information about the cfe-commits mailing list