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

VigneshwarJ · 2024-10-23T15:23:47Z

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

llvmbot · 2024-10-23T15:24:23Z

@llvm/pr-subscribers-clang

Author: Vigneshwar Jayakumar (VigneshwarJ)

Changes

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

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

3 Files Affected:

(modified) clang/include/clang/Basic/DiagnosticSemaKinds.td (+1-1)
(modified) clang/lib/Sema/SemaType.cpp (+11)
(added) clang/test/SemaHIP/zero-sized-device-array.hip (+25)

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

JonChesterfield

Error looks good. Might want to add a case for "dynamic shared" to the test file as the syntax is very close to the case being diagnosed - iirc it's things like

extern __shared__ float array[];

Some existing handling in C like languages conflates [] and [0], so might also want extern __shared__ float array2[0];

That's probably already covered by existing HIP test cases though, so if everything else is passing I think we're good. Thanks!

tbaederr · 2024-10-23T23:32:25Z

clang/lib/Sema/SemaType.cpp

+          auto *FD = dyn_cast_or_null<FunctionDecl>(CurContext);
+          if (FD && (FD->hasAttr<CUDADeviceAttr>() ||
+                     FD->hasAttr<CUDAGlobalAttr>())) {


Suggested change

auto *FD = dyn_cast_or_null<FunctionDecl>(CurContext);

if (FD && (FD->hasAttr<CUDADeviceAttr>() ||

FD->hasAttr<CUDAGlobalAttr>())) {

if (const auto FD = dyn_cast_if_present(...);FD && (FD->hasAttr<CUDADeviceAttr>() ||

FD->hasAttr<CUDAGlobalAttr>())) {

efriedma-quic · 2024-10-23T23:53:05Z

clang/lib/Sema/SemaType.cpp

            << 0 << ArraySize->getSourceRange();
+
+        // zero sized static arrays are not allowed in HIP device functions
+        if (LangOpts.HIP && LangOpts.CUDAIsDevice) {


I'm not sure you're diagnosing this in the right place; BuildArrayType gets called at the point we parse the type, not the variable, so you can easily avoid the diagnostic using something like the following:

typedef float floatarr[0]; __global__ void global_fun() { floatarr array; }

Thanks for pointing it out, moved the logic to SemaVarDeclaration, also checking for the typedef and pointers

moved code to SemaVarDecl also check the pointer types to figure out its within any typedefs or pointers.

efriedma-quic · 2024-10-25T20:55:33Z

clang/lib/Sema/SemaDecl.cpp

+      };
+      QualType NextTy = NewVD->getType();
+      while (NextTy->isAnyPointerType() || NextTy->isArrayType() ||
+             NextTy->isReferenceType()) {


I'm having trouble imagining how a pointer to a zero-size array could cause issues; there isn't any way to tell in IR that pointer points to a zero-length array.

Yes, you are right. I got confused. That would not cause any issue, The while logic is unnecessary, updated to simply check if its an array type.

efriedma-quic · 2024-10-30T23:57:37Z

clang/lib/Sema/SemaDecl.cpp

+      if (QualType NextTy = NewVD->getType(); NextTy->isArrayType()) {
+        if (const ConstantArrayType *ArrayT =
+                getASTContext().getAsConstantArrayType(NextTy);


Suggested change

if (QualType NextTy = NewVD->getType(); NextTy->isArrayType()) {

if (const ConstantArrayType *ArrayT =

getASTContext().getAsConstantArrayType(NextTy);

if (const ConstantArrayType *ArrayT =

getASTContext().getAsConstantArrayType(T);

clang/lib/Sema/SemaDecl.cpp

efriedma-quic · 2024-11-13T23:54:02Z

clang/lib/Sema/SemaDecl.cpp


+  // zero sized static arrays are not allowed in HIP device functions
+  if (getLangOpts().HIP &&
+      DeclAttrsMatchCUDAMode(getLangOpts(), getCurFunctionDecl())) {


Do you need to check LangOpts.CUDAIsDevice?

Yes and that is being checked in the function DeclAttrsMatchCUDAMode.

The check there is isDeviceSideDecl == LangOpts.CUDAIsDevice... but you don't want to diagnose this on host, do you? (The regression test only checks the -fcuda-is-device case.)

Yes, that's right, I dont want to diagnose this on host, but the device side decl are cuda device specific. So I thought there would not be any problem.

else it would be cleaner to go back to how I was checking previously as this is the only helper function I found that is closely related.

The code, in its current state, will diagnose on host. (The testcase doesn't show this because it only runs with -fcuda-is-device.)

…oject into SWDEV_449592

efriedma-quic

LGTM

clang/lib/Sema/SemaDecl.cpp

Revert: breaks build of rocFFT 854d730 [Clang/AMDGPU] Zero sized arrays not allowed in HIP device code. (llvm#113470) Change-Id: Idb4e97a583e9e595154b2da69ef68568b6d96eb5

…m#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@amd.com> (cherry picked from commit 854d730) Reapply a commit which was reverted due to an issue in rocFFT that needed to be fixed first.

…m#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@amd.com> (cherry picked from commit 854d730) Reapply a commit which was reverted due to an issue in rocFFT that needed to be fixed first. (cherry picked from commit 7cc4722)

…m#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@amd.com> (cherry picked from commit 854d730)

…m#113470) (llvm#4650)

[Clang/AMDGPU] Zero sized arrays not allowed in HIP device code.

0e2ee52

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

llvmbot added clang Clang issues not falling into any other category clang:frontend Language frontend issues, e.g. anything involving "Sema" labels Oct 23, 2024

bcahoon requested a review from yxsamliu October 23, 2024 15:49

JonChesterfield approved these changes Oct 23, 2024

View reviewed changes

tbaederr reviewed Oct 23, 2024

View reviewed changes

efriedma-quic reviewed Oct 23, 2024

View reviewed changes

fix review comments

fed2349

moved code to SemaVarDecl also check the pointer types to figure out its within any typedefs or pointers.

efriedma-quic reviewed Oct 25, 2024

View reviewed changes

review comment fixes

7ee4272

efriedma-quic reviewed Oct 30, 2024

View reviewed changes

VigneshwarJ and others added 2 commits November 7, 2024 08:38

review comments

a7065f4

Merge branch 'main' into SWDEV_449592

187615c

efriedma-quic reviewed Nov 13, 2024

View reviewed changes

VigneshwarJ added 2 commits November 20, 2024 12:28

review changes

d613d59

Merge branch 'SWDEV_449592' of https://github.com/VigneshwarJ/llvm-pr…

38a59ec

…oject into SWDEV_449592

efriedma-quic approved these changes Nov 20, 2024

View reviewed changes

clang/lib/Sema/SemaDecl.cpp Show resolved Hide resolved

bcahoon merged commit 854d730 into llvm:main Nov 27, 2024
8 checks passed

rahulc-gh pushed a commit to ROCm/llvm-project that referenced this pull request Nov 6, 2025

[Clang/AMDGPU] Zero sized arrays not allowed in HIP device code. (llv…

27682a1

…m#113470) (llvm#4650)

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

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

Uh oh!

Conversation

VigneshwarJ commented Oct 23, 2024

Uh oh!

llvmbot commented Oct 23, 2024

Uh oh!

JonChesterfield left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

efriedma-quic left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants