[clang] [Clang/AMDGPU] Allow zero sized arrays in HIP device code (PR #195716)

[via cfe-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clang

Author: Joseph Huber (jhuber6)

<details>
<summary>Changes</summary>

Summary:
I tested the original reproducer and was unable to reproduce the
original issue. There are already tests that check for zero-sized LDS
usage so I believe this should be fine. This should be a better solution
than https://github.com/llvm/llvm-project/pull/195700.

This is really confusing to me because there seemed to be a backend fix
for this problem in 7c7704c946ab but we then landed a Sema fix for it
half a year later? Maybe someone else can confirm this, but I really
don't see any particular reason to split host / device compilation on
this.

This reverts commit 854d7301f989dd1e3c838ef4f48cb57bb7d496e0.


---
Full diff: https://github.com/llvm/llvm-project/pull/195716.diff


4 Files Affected:

- (modified) clang/include/clang/Basic/DiagnosticSemaKinds.td (+1-1) 
- (modified) clang/lib/Sema/SemaDecl.cpp (-13) 
- (modified) clang/lib/Sema/SemaType.cpp (+1-1) 
- (removed) clang/test/SemaHIP/zero-sized-device-array.hip (-38) 


``````````diff
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index e059260778631..d1553fd924c28 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6694,7 +6694,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|HIP device code|OpenCL}0">;
+  "zero-length arrays are not permitted in %select{C++|SYCL device code|OpenCL}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 eb5b6d65b4d58..fe97433567a7c 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -9058,19 +9058,6 @@ 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/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index a5fceb065fb88..1b6aa61383a98 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -2290,7 +2290,7 @@ QualType Sema::BuildArrayType(QualType T, ArraySizeModifier ASM,
       if (ConstVal == 0 && !T.isWebAssemblyReferenceType()) {
         if (getLangOpts().OpenCL) {
           Diag(ArraySize->getBeginLoc(), diag::err_typecheck_zero_array_size)
-              << 3 << ArraySize->getSourceRange();
+              << 2 << ArraySize->getSourceRange();
           return QualType();
         }
 
diff --git a/clang/test/SemaHIP/zero-sized-device-array.hip b/clang/test/SemaHIP/zero-sized-device-array.hip
deleted file mode 100644
index 8b25ec38fb01a..0000000000000
--- a/clang/test/SemaHIP/zero-sized-device-array.hip
+++ /dev/null
@@ -1,38 +0,0 @@
-// 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}} 
-}

``````````

</details>


https://github.com/llvm/llvm-project/pull/195716