Author: Vigneshwar Jayakumar Date: 2024-11-27T14:20:07-06:00 New Revision: 854d7301f989dd1e3c838ef4f48cb57bb7d496e0
URL: https://github.com/llvm/llvm-project/commit/854d7301f989dd1e3c838ef4f48cb57bb7d496e0 DIFF: https://github.com/llvm/llvm-project/commit/854d7301f989dd1e3c838ef4f48cb57bb7d496e0.diff LOG: [Clang/AMDGPU] Zero sized arrays not allowed in HIP device code. (#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.jayaku...@amd.com> Added: clang/test/SemaHIP/zero-sized-device-array.hip Modified: clang/include/clang/Basic/DiagnosticSemaKinds.td clang/lib/Sema/SemaDecl.cpp Removed: ################################################################################ diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 1d777f670097b5..a6c5c5806c33f2 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -6270,7 +6270,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/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index c33701c0157714..ba574307055c6b 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -8762,6 +8762,19 @@ 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/test/SemaHIP/zero-sized-device-array.hip b/clang/test/SemaHIP/zero-sized-device-array.hip new file mode 100644 index 00000000000000..8b25ec38fb01ac --- /dev/null +++ b/clang/test/SemaHIP/zero-sized-device-array.hip @@ -0,0 +1,38 @@ +// 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}} +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
