[clang] [Clang/AMDGPU] Allow zero sized arrays in HIP device code (PR #195716)
Joseph Huber via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 1 08:49:52 PDT 2026
https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/195716
>From c687ad90837ba4fc86526cd5b651e91dd04e6e56 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Mon, 4 May 2026 13:41:24 -0500
Subject: [PATCH 1/2] [Clang/AMDGPU] Allow zero sized arrays in HIP device code
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.
---
.../clang/Basic/DiagnosticSemaKinds.td | 2 +-
clang/lib/Sema/SemaDecl.cpp | 13 -------
clang/lib/Sema/SemaType.cpp | 2 +-
.../test/SemaHIP/zero-sized-device-array.hip | 34 ++++++++++---------
4 files changed, 20 insertions(+), 31 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index d293a9798da6a..2c3fbe47f3b51 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6761,7 +6761,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 db5e66cb96c3e..6d2d46f6b7ab2 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -9180,19 +9180,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 037583575d19f..34ba19223c3d3 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -2299,7 +2299,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
index 612ebeac7767c..bab5de7bba614 100644
--- a/clang/test/SemaHIP/zero-sized-device-array.hip
+++ b/clang/test/SemaHIP/zero-sized-device-array.hip
@@ -1,5 +1,9 @@
// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -fsyntax-only -x hip -fcuda-is-device -verify -triple amdgpu %s
+// RUN: %clang_cc1 -fsyntax-only -x hip -fcuda-is-device -verify -triple amdgpu %s
+// RUN: %clang_cc1 -fsyntax-only -x hip -verify -triple x86_64-unknown-linux-gnu %s
+
+// expected-no-diagnostics
+
#define __device__ __attribute__((device))
#define __host__ __attribute__((host))
#define __global__ __attribute__((global))
@@ -8,31 +12,29 @@
typedef float ZEROARR[0];
float global_array[0];
+__device__ float device_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}}
+ extern __shared__ float externArray[];
+ ZEROARR TypeDef;
+ float array[0];
}
-// should not throw error for host side code.
__host__ void host_fun() {
- float array[0];
+ 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}}
+__device__ void templated() {
+ Ty arr[Size];
+ __shared__ Ty shared_arr[Size];
}
-__host__ __device__ void host_dev_fun()
-{
- float array[0]; // expected-error {{zero-length arrays are not permitted in HIP device code}}
+__host__ __device__ void host_dev_fun() {
+ float array[0];
}
-__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}}
+__device__ void device_fun() {
+ __shared__ float array[0];
+ templated<int, 0>();
}
>From 32097c3d632ba4d19c638f4022cd31169776a25d Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Thu, 1 Oct 2026 10:43:28 -0500
Subject: [PATCH 2/2] [Clang] Warn on zero-length __shared__ arrays
Summary:
Zero-length `__shared__` arrays are now accepted, but they are easily
confused with dynamic shared memory declared as `extern __shared__ T x[]`.
Emit a warning in device compilation suggesting the latter.
---
clang/docs/ReleaseNotes.md | 2 ++
.../include/clang/Basic/DiagnosticSemaKinds.td | 4 ++++
clang/lib/Sema/SemaDecl.cpp | 9 +++++++++
clang/test/SemaCUDA/extern-shared.cu | 8 +++++---
clang/test/SemaHIP/zero-sized-device-array.hip | 18 ++++++++++++++----
5 files changed, 34 insertions(+), 7 deletions(-)
diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md
index 81e4d1fc8360f..cd2001745d0ab 100644
--- a/clang/docs/ReleaseNotes.md
+++ b/clang/docs/ReleaseNotes.md
@@ -922,6 +922,8 @@ features cannot lower the translation-unit ABI level;
libhipcxx headers to be included with paths such as `<cuda/std/atomic>`.
The `-nogpuinc` option disables this path together with the other HIP include
paths.
+- Zero-length arrays are now allowed in HIP device code. Zero-length
+ `__shared__` arrays are diagnosed by `-Wzero-length-shared-array`.
#### CUDA Support
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 2c3fbe47f3b51..3a1997d04c95f 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -9819,6 +9819,10 @@ def err_cuda_vla : Error<
"cannot use variable-length arrays in "
"%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
def err_cuda_extern_shared : Error<"__shared__ variable %0 cannot be 'extern'">;
+def warn_cuda_zero_length_shared_array : Warning<
+ "zero-length __shared__ array %0 has no storage; did you mean "
+ "'extern __shared__' with an unspecified size for dynamic shared memory?">,
+ InGroup<DiagGroup<"zero-length-shared-array">>;
def err_cuda_host_shared : Error<
"__shared__ local variables not allowed in "
"%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 6d2d46f6b7ab2..23b3d51973caf 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -9180,6 +9180,15 @@ void Sema::CheckVariableDeclarationType(VarDecl *NewVD) {
}
}
+ // Zero-length __shared__ arrays are often meant as dynamic shared memory.
+ if (getLangOpts().CUDAIsDevice && NewVD->hasAttr<CUDASharedAttr>() &&
+ !inTemplateInstantiation()) {
+ if (const ConstantArrayType *ArrayT = Context.getAsConstantArrayType(T);
+ ArrayT && ArrayT->isZeroSize())
+ Diag(NewVD->getLocation(), diag::warn_cuda_zero_length_shared_array)
+ << NewVD;
+ }
+
bool isVM = T->isVariablyModifiedType();
if (isVM || NewVD->hasAttr<CleanupAttr>() ||
NewVD->hasAttr<BlocksAttr>())
diff --git a/clang/test/SemaCUDA/extern-shared.cu b/clang/test/SemaCUDA/extern-shared.cu
index bc4f10d3082c0..7627d7eadda46 100644
--- a/clang/test/SemaCUDA/extern-shared.cu
+++ b/clang/test/SemaCUDA/extern-shared.cu
@@ -2,7 +2,7 @@
// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -fcuda-is-device -verify %s
// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -fgpu-rdc -verify=rdc %s
-// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -fcuda-is-device -fgpu-rdc -verify=rdc %s
+// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -fcuda-is-device -fgpu-rdc -verify=rdc,rdc-dev %s
// Most of these declarations are fine in separate compilation mode.
@@ -11,14 +11,16 @@
__device__ void foo() {
extern __shared__ int x; // expected-error {{__shared__ variable 'x' cannot be 'extern'}}
extern __shared__ int arr[]; // ok
- extern __shared__ int arr0[0]; // expected-error {{__shared__ variable 'arr0' cannot be 'extern'}}
+ extern __shared__ int arr0[0]; // expected-error {{__shared__ variable 'arr0' cannot be 'extern'}} \
+ // rdc-dev-warning {{zero-length __shared__ array 'arr0' has no storage}}
extern __shared__ int arr1[1]; // expected-error {{__shared__ variable 'arr1' cannot be 'extern'}}
extern __shared__ int* ptr ; // expected-error {{__shared__ variable 'ptr' cannot be 'extern'}}
}
__host__ __device__ void bar() {
extern __shared__ int arr[]; // ok
- extern __shared__ int arr0[0]; // expected-error {{__shared__ variable 'arr0' cannot be 'extern'}}
+ extern __shared__ int arr0[0]; // expected-error {{__shared__ variable 'arr0' cannot be 'extern'}} \
+ // rdc-dev-warning {{zero-length __shared__ array 'arr0' has no storage}}
extern __shared__ int arr1[1]; // expected-error {{__shared__ variable 'arr1' cannot be 'extern'}}
extern __shared__ int* ptr ; // expected-error {{__shared__ variable 'ptr' cannot be 'extern'}}
}
diff --git a/clang/test/SemaHIP/zero-sized-device-array.hip b/clang/test/SemaHIP/zero-sized-device-array.hip
index bab5de7bba614..34270567438f3 100644
--- a/clang/test/SemaHIP/zero-sized-device-array.hip
+++ b/clang/test/SemaHIP/zero-sized-device-array.hip
@@ -1,8 +1,10 @@
// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -fsyntax-only -x hip -fcuda-is-device -verify -triple amdgpu %s
-// RUN: %clang_cc1 -fsyntax-only -x hip -verify -triple x86_64-unknown-linux-gnu %s
+// RUN: %clang_cc1 -fsyntax-only -x hip -fcuda-is-device -verify=device -triple amdgpu %s
+// RUN: %clang_cc1 -fsyntax-only -x cuda -fcuda-is-device -verify=device -triple nvptx64-nvidia-cuda %s
+// RUN: %clang_cc1 -fsyntax-only -x hip -verify=host -triple x86_64-unknown-linux-gnu %s
+// RUN: %clang_cc1 -fsyntax-only -x hip -fcuda-is-device -Wno-zero-length-shared-array -verify=host -triple amdgpu %s
-// expected-no-diagnostics
+// host-no-diagnostics
#define __device__ __attribute__((device))
#define __host__ __attribute__((host))
@@ -13,6 +15,7 @@ typedef float ZEROARR[0];
float global_array[0];
__device__ float device_array[0];
+__shared__ float shared_global_array[0]; // device-warning {{zero-length __shared__ array 'shared_global_array' has no storage; did you mean 'extern __shared__' with an unspecified size for dynamic shared memory?}}
__global__ void global_fun() {
extern __shared__ float externArray[];
@@ -30,11 +33,18 @@ __device__ void templated() {
__shared__ Ty shared_arr[Size];
}
+template <typename Ty>
+__device__ void dependent_element() {
+ __shared__ Ty shared_arr[0]; // device-warning {{zero-length __shared__ array 'shared_arr' has no storage}}
+}
+
__host__ __device__ void host_dev_fun() {
float array[0];
}
__device__ void device_fun() {
- __shared__ float array[0];
+ __shared__ float array[0]; // device-warning {{zero-length __shared__ array 'array' has no storage}}
+ __shared__ ZEROARR typedef_array; // device-warning {{zero-length __shared__ array 'typedef_array' has no storage}}
templated<int, 0>();
+ dependent_element<int>();
}
More information about the cfe-commits
mailing list