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 <[email protected]> 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 <[email protected]> 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>(); } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
