llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Vigneshwar Jayakumar (VigneshwarJ) <details> <summary>Changes</summary> 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 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}} +} `````````` </details> https://github.com/llvm/llvm-project/pull/113470 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
