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

[Vigneshwar Jayakumar via cfe-commits]

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

Added diagnosis to throw error when zero sized arrays are used in the HIP device code. SWDEV-449592

>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] [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}}
+}