diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index 438956488ad5f..4024ebc00cc9e 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -2340,6 +2340,12 @@ QualType Sema::BuildArrayType(QualType T, ArrayType::ArraySizeModifier ASM, << ArraySize->getSourceRange(); ASM = ArrayType::Normal; } + + // Zero length arrays are disallowed in SYCL device code. + if (getLangOpts().SYCLIsDevice) + SYCLDiagIfDeviceCode(ArraySize->getBeginLoc(), + diag::err_typecheck_zero_array_size) + << ArraySize->getSourceRange(); } else if (!T->isDependentType() && !T->isVariablyModifiedType() && !T->isIncompleteType() && !T->isUndeducedType()) { // Is the array too large? diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp new file mode 100644 index 0000000000000..030985832f3cb --- /dev/null +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -0,0 +1,77 @@ +// RUN: %clang_cc1 -fsycl -triple spir64 -fsycl-is-device -verify -fsyntax-only %s +// +// Ensure that the SYCL diagnostics that are typically deferred are correctly emitted. + +// testing that the deferred diagnostics work in conjunction with the SYCL namespaces. +inline namespace cl { +namespace sycl { + +template +__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { + // expected-note@+1 2{{called by 'kernel_single_task +void setup_sycl_operation(const T VA[]) { + + cl::sycl::kernel_single_task([]() { + // FIX!! xpected-error@+1 {{zero-length arrays are not permitted in C++}} + int OverlookedBadArray[0]; + + // FIX!! xpected-error@+1 {{__float128 is not supported on this target}} + __float128 overlookedBadFloat = 40; + }); +} + +int main(int argc, char **argv) { + + // --- direct lambda testing --- + cl::sycl::kernel_single_task([]() { + // expected-error@+1 {{zero-length arrays are not permitted in C++}} + int BadArray[0]; + + // expected-error@+1 {{__float128 is not supported on this target}} + __float128 badFloat = 40; // this SHOULD trigger a diagnostic + + //expected-error@+1 {{SYCL kernel cannot call a variadic function}} + variadic(5); + + // expected-note@+1 {{called by 'operator()'}} + calledFromKernel(10); + }); + + // --- lambda in specialized function testing --- + + //array A is only used to feed the template + const int array_size = 4; + int A[array_size] = {1, 2, 3, 4}; + setup_sycl_operation(A); + + return 0; +} diff --git a/clang/test/SemaSYCL/sycl-restrict.cpp b/clang/test/SemaSYCL/sycl-restrict.cpp index 658395e92c3ff..53269924d060c 100644 --- a/clang/test/SemaSYCL/sycl-restrict.cpp +++ b/clang/test/SemaSYCL/sycl-restrict.cpp @@ -143,6 +143,9 @@ void usage(myFuncDef functionPtr) { // expected-error@+1 {{__float128 is not supported on this target}} __float128 A; + + // expected-error@+1 {{zero-length arrays are not permitted in C++}} + int BadArray[0]; } namespace ns {