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

[via cfe-commits]

Author: Vigneshwar Jayakumar
Date: 2024-11-27T14:20:07-06:00
New Revision: 854d7301f989dd1e3c838ef4f48cb57bb7d496e0

URL: https://github.com/llvm/llvm-project/commit/854d7301f989dd1e3c838ef4f48cb57bb7d496e0
DIFF: https://github.com/llvm/llvm-project/commit/854d7301f989dd1e3c838ef4f48cb57bb7d496e0.diff

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

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

---------

Co-authored-by: vigneshwar jayakumar <vigneshwar.jayakumar at amd.com>

Added: 
    clang/test/SemaHIP/zero-sized-device-array.hip

Modified: 
    clang/include/clang/Basic/DiagnosticSemaKinds.td
    clang/lib/Sema/SemaDecl.cpp

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 1d777f670097b5..a6c5c5806c33f2 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6270,7 +6270,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/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index c33701c0157714..ba574307055c6b 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -8762,6 +8762,19 @@ void Sema::CheckVariableDeclarationType(VarDecl *NewVD) {
     }
   }
 
+  // zero sized static arrays are not allowed in HIP device functions
+  if (getLangOpts().HIP && LangOpts.CUDAIsDevice) {
+    if (FunctionDecl *FD = getCurFunctionDecl();
+        FD &&
+        (FD->hasAttr<CUDADeviceAttr>() || FD->hasAttr<CUDAGlobalAttr>())) {
+      if (const ConstantArrayType *ArrayT =
+              getASTContext().getAsConstantArrayType(T);
+          ArrayT && ArrayT->isZeroSize()) {
+        Diag(NewVD->getLocation(), diag::err_typecheck_zero_array_size) << 2;
+      }
+    }
+  }
+
   bool isVM = T->isVariablyModifiedType();
   if (isVM || NewVD->hasAttr<CleanupAttr>() ||
       NewVD->hasAttr<BlocksAttr>())

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..8b25ec38fb01ac
--- /dev/null
+++ b/clang/test/SemaHIP/zero-sized-device-array.hip
@@ -0,0 +1,38 @@
+// 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))
+
+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}}
+    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];
+}
+
+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}}
+}
+
+__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}} 
+}