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

[Brendon Cahoon via cfe-commits]

================
@@ -9180,17 +9180,13 @@ 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;
-      }
-    }
+  // Zero-length __shared__ arrays are often meant as dynamic shared memory.
+  if (getLangOpts().CUDAIsDevice && NewVD->hasAttr<CUDASharedAttr>() &&
+      !inTemplateInstantiation()) {
----------------
bcahoon wrote:

Adding @VigneshwarJ.  I'm fine with changing this as appropriate. Though, I do think its useful to have warning message for the cases that result in run-time errors that are difficult to diagnose otherwise.

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