From eaa10ea40fc6786c2bab138512d796857371dad4 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Sat, 15 Feb 2020 11:22:51 -0800 Subject: [PATCH 01/18] disallow zero length arrays in SYCL kernels Signed-off-by: Chris Perkins --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 2 ++ clang/lib/Sema/SemaType.cpp | 7 +++++++ clang/test/SemaSYCL/sycl-restrict.cpp | 3 +++ 3 files changed, 12 insertions(+) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index afa83e18935c7..b90606103c04e 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10429,6 +10429,8 @@ def err_conflicting_sycl_kernel_attributes : Error< "conflicting attributes applied to a SYCL kernel">; def err_conflicting_sycl_function_attributes : Error< "%0 attribute conflicts with '%1' attribute">; +def err_sycl_typecheck_zero_array_size : Error< + "zero-length arrays are not permitted in SYCL kernels">; def err_sycl_x_y_z_arguments_must_be_one : Error< "%0 X-, Y- and Z- sizes must be 1 when %1 attribute is used with value 0">; def err_intel_attribute_argument_is_not_in_range: Error< diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index 438956488ad5f..af83397df9418 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -2340,6 +2340,13 @@ QualType Sema::BuildArrayType(QualType T, ArrayType::ArraySizeModifier ASM, << ArraySize->getSourceRange(); ASM = ArrayType::Normal; } + + // SYCL kernels reject zero length arrays + if (getLangOpts().SYCLIsDevice) { + SYCLDiagIfDeviceCode(ArraySize->getBeginLoc(), + diag::err_sycl_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/sycl-restrict.cpp b/clang/test/SemaSYCL/sycl-restrict.cpp index 658395e92c3ff..61a7e499aa936 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 SYCL kernels}} + int BadArray[0]; } namespace ns { From e9dd04a4c84b542208016e823e7c048fc544ecad Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 20 Feb 2020 10:16:55 -0800 Subject: [PATCH 02/18] changing diagnostic code per feedback Signed-off-by: Chris Perkins --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 2 -- clang/lib/Sema/SemaType.cpp | 2 +- clang/test/SemaSYCL/sycl-restrict.cpp | 2 +- 3 files changed, 2 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index b90606103c04e..afa83e18935c7 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10429,8 +10429,6 @@ def err_conflicting_sycl_kernel_attributes : Error< "conflicting attributes applied to a SYCL kernel">; def err_conflicting_sycl_function_attributes : Error< "%0 attribute conflicts with '%1' attribute">; -def err_sycl_typecheck_zero_array_size : Error< - "zero-length arrays are not permitted in SYCL kernels">; def err_sycl_x_y_z_arguments_must_be_one : Error< "%0 X-, Y- and Z- sizes must be 1 when %1 attribute is used with value 0">; def err_intel_attribute_argument_is_not_in_range: Error< diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index af83397df9418..ddb0cd15e246d 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -2344,7 +2344,7 @@ QualType Sema::BuildArrayType(QualType T, ArrayType::ArraySizeModifier ASM, // SYCL kernels reject zero length arrays if (getLangOpts().SYCLIsDevice) { SYCLDiagIfDeviceCode(ArraySize->getBeginLoc(), - diag::err_sycl_typecheck_zero_array_size) + diag::err_typecheck_zero_array_size) << ArraySize->getSourceRange(); } } else if (!T->isDependentType() && !T->isVariablyModifiedType() && diff --git a/clang/test/SemaSYCL/sycl-restrict.cpp b/clang/test/SemaSYCL/sycl-restrict.cpp index 61a7e499aa936..53269924d060c 100644 --- a/clang/test/SemaSYCL/sycl-restrict.cpp +++ b/clang/test/SemaSYCL/sycl-restrict.cpp @@ -144,7 +144,7 @@ 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 SYCL kernels}} + // expected-error@+1 {{zero-length arrays are not permitted in C++}} int BadArray[0]; } From 765d43bc57c589b4eb66a83300735c07be3d523b Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 20 Feb 2020 11:34:04 -0800 Subject: [PATCH 03/18] remove braces, per PR request Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaType.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index ddb0cd15e246d..83e7cfc82f53b 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -2342,11 +2342,10 @@ QualType Sema::BuildArrayType(QualType T, ArrayType::ArraySizeModifier ASM, } // SYCL kernels reject zero length arrays - if (getLangOpts().SYCLIsDevice) { + 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? From fb3d2ace1c1b1a0679c87cd12fbae325ccfc1799 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 20 Feb 2020 14:27:31 -0800 Subject: [PATCH 04/18] adding test of deferred diagnostics, per PR request Signed-off-by: Chris Perkins --- .../SemaSYCL/deferred-diagnostics-emit.cpp | 72 +++++++++++++++++++ 1 file changed, 72 insertions(+) create mode 100644 clang/test/SemaSYCL/deferred-diagnostics-emit.cpp diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp new file mode 100644 index 0000000000000..afce62a86022c --- /dev/null +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -0,0 +1,72 @@ +// RUN: %clang_cc1 -fcxx-exceptions -triple spir64 -fsycl-is-device -Wno-return-type -verify -fsyntax-only -std=c++17 %s + +/* + ensuring that the SYCL diagnostics that are typically deferred, correctly emit +*/ + +struct S { + virtual void foo() {} +}; + +int calledFromKernel(int a){ + // expected-error@+1 {{zero-length arrays are not permitted in C++}} + int MalArray[0]; + + // expected-error@+1 {{__float128 is not supported on this target}} + __float128 malFloat = 40; + + S mal; + // expected-error@+1 {{SYCL kernel cannot call a virtual function}} + mal.foo(); + + return a + 20; +} + +inline namespace cl { +namespace sycl { +class queue { +public: + template void submit(T CGF) {} +}; + +template class id {}; + +template class range {}; + +class handler { +public: + template + __attribute__((sycl_kernel)) void kernel_parallel_for(KernelType kernelFunc) { + kernelFunc(id<1>{}); + } + template + void parallel_for(range NWI, KernelType kernelFunc) { + kernel_parallel_for(kernelFunc); + } +}; +} +} + +int main(int argc, char **argv) { + + cl::sycl::range<1> numOfItems; + cl::sycl::queue deviceQueue; + + deviceQueue.submit([&](cl::sycl::handler &cgh) { + cgh.parallel_for(numOfItems, [=](cl::sycl::id<1> wiID) { + // 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 + + S s; + // expected-error@+1 {{zero-length arrays are not permitted in C++}} + s.foo(); + + calledFromKernel(10); + }); + }); + + return 0; +} \ No newline at end of file From 9618333a582fa2e9df9e37ebc4212b169700e654 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 20 Feb 2020 15:10:37 -0800 Subject: [PATCH 05/18] improved deferred diagnostic testing Signed-off-by: Chris Perkins --- .../SemaSYCL/deferred-diagnostics-emit.cpp | 38 ++++++++++++++++++- 1 file changed, 37 insertions(+), 1 deletion(-) diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index afce62a86022c..56aca69a84a3d 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -15,6 +15,7 @@ int calledFromKernel(int a){ // expected-error@+1 {{__float128 is not supported on this target}} __float128 malFloat = 40; + // not sure if 'no virtual function' is a _deferred_ diagnostic, testing anyway S mal; // expected-error@+1 {{SYCL kernel cannot call a virtual function}} mal.foo(); @@ -47,8 +48,34 @@ class handler { } } +/* + template used to specialize a function that contains a lambda that should + result in a deferred diagnostic being emitted. + HOWEVER, this is not working presently. + TODO: re-test after new deferred diagnostic system is merged. + restore the "FIX!!" tests below +*/ +template +void setup_sycl_operation(const T VA[]) { + + cl::sycl::range<1> numOfItems; + cl::sycl::queue deviceQueue; + + deviceQueue.submit([&](cl::sycl::handler &cgh) { + cgh.parallel_for(numOfItems, [=](cl::sycl::id<1> wiID) { + // FIX!! expected-error@+1 {{zero-length arrays are not permitted in C++}} + int OverlookedBadArray[0]; + + // FIX!! expected-error@+1 {{__float128 is not supported on this target}} + __float128 overlookedBadFloat = 40; + + }); + }); +} + int main(int argc, char **argv) { + // --- direct lambda testing --- cl::sycl::range<1> numOfItems; cl::sycl::queue deviceQueue; @@ -60,13 +87,22 @@ int main(int argc, char **argv) { // expected-error@+1 {{__float128 is not supported on this target}} __float128 badFloat = 40; // this SHOULD trigger a diagnostic + // not sure if 'no virtual function' is a _deferred_ diagnostic, but testing anyway. S s; - // expected-error@+1 {{zero-length arrays are not permitted in C++}} + // expected-error@+1 {{SYCL kernel cannot call a virtual function}} s.foo(); 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; } \ No newline at end of file From cc9aa9fcacf451b5618fb35a75f314256c40bde5 Mon Sep 17 00:00:00 2001 From: cperkinsintel Date: Fri, 21 Feb 2020 10:06:38 -0800 Subject: [PATCH 06/18] Update clang/test/SemaSYCL/deferred-diagnostics-emit.cpp Signed-off-by: Chris Perkins Co-Authored-By: Mariya Podchishchaeva --- clang/test/SemaSYCL/deferred-diagnostics-emit.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index 56aca69a84a3d..50ec43a986a5d 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fcxx-exceptions -triple spir64 -fsycl-is-device -Wno-return-type -verify -fsyntax-only -std=c++17 %s +// RUN: %clang_cc1 -fcxx-exceptions -triple spir64 -fsycl-is-device -Wno-return-type -verify -fsyntax-only %s /* ensuring that the SYCL diagnostics that are typically deferred, correctly emit @@ -105,4 +105,4 @@ int main(int argc, char **argv) { setup_sycl_operation(A); return 0; -} \ No newline at end of file +} From 17a2f4467b8b0a6cd9c90f637cd8313528ce449a Mon Sep 17 00:00:00 2001 From: cperkinsintel Date: Fri, 21 Feb 2020 10:06:56 -0800 Subject: [PATCH 07/18] Update clang/lib/Sema/SemaType.cpp Signed-off-by: Chris Perkins Co-Authored-By: Mariya Podchishchaeva --- clang/lib/Sema/SemaType.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index 83e7cfc82f53b..e080375937bf1 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -2341,7 +2341,7 @@ QualType Sema::BuildArrayType(QualType T, ArrayType::ArraySizeModifier ASM, ASM = ArrayType::Normal; } - // SYCL kernels reject zero length arrays + // Zero length arrays are disallowed in SYCL device code. if (getLangOpts().SYCLIsDevice) SYCLDiagIfDeviceCode(ArraySize->getBeginLoc(), diag::err_typecheck_zero_array_size) From b8a67d31be9054efb11d43a21484b9a1dc491abd Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 21 Feb 2020 11:08:01 -0800 Subject: [PATCH 08/18] changing comment formatting per PR feedback Signed-off-by: Chris Perkins --- .../SemaSYCL/deferred-diagnostics-emit.cpp | 20 +++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index 50ec43a986a5d..2941e4c75f3ee 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -1,8 +1,8 @@ // RUN: %clang_cc1 -fcxx-exceptions -triple spir64 -fsycl-is-device -Wno-return-type -verify -fsyntax-only %s -/* - ensuring that the SYCL diagnostics that are typically deferred, correctly emit -*/ +// +// ensuring that the SYCL diagnostics that are typically deferred, correctly emit +// struct S { virtual void foo() {} @@ -48,13 +48,13 @@ class handler { } } -/* - template used to specialize a function that contains a lambda that should - result in a deferred diagnostic being emitted. - HOWEVER, this is not working presently. - TODO: re-test after new deferred diagnostic system is merged. - restore the "FIX!!" tests below -*/ + +// template used to specialize a function that contains a lambda that should +// result in a deferred diagnostic being emitted. +// HOWEVER, this is not working presently. +// TODO: re-test after new deferred diagnostic system is merged. +// restore the "FIX!!" tests below + template void setup_sycl_operation(const T VA[]) { From 8d92c4dac6950dfebe3cd3b049269b2c36c57040 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 21 Feb 2020 11:37:49 -0800 Subject: [PATCH 09/18] switched to common fake SYCL scaffolding Signed-off-by: Chris Perkins --- .../SemaSYCL/deferred-diagnostics-emit.cpp | 33 +++---------------- 1 file changed, 5 insertions(+), 28 deletions(-) diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index 2941e4c75f3ee..c1ed067e5ce45 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -1,9 +1,11 @@ -// RUN: %clang_cc1 -fcxx-exceptions -triple spir64 -fsycl-is-device -Wno-return-type -verify -fsyntax-only %s +// RUN: %clang_cc1 -I %S/Inputs -fcxx-exceptions -triple spir64 -fsycl-is-device -Wno-return-type -verify -fsyntax-only %s // // ensuring that the SYCL diagnostics that are typically deferred, correctly emit // +#include + struct S { virtual void foo() {} }; @@ -23,31 +25,6 @@ int calledFromKernel(int a){ return a + 20; } -inline namespace cl { -namespace sycl { -class queue { -public: - template void submit(T CGF) {} -}; - -template class id {}; - -template class range {}; - -class handler { -public: - template - __attribute__((sycl_kernel)) void kernel_parallel_for(KernelType kernelFunc) { - kernelFunc(id<1>{}); - } - template - void parallel_for(range NWI, KernelType kernelFunc) { - kernel_parallel_for(kernelFunc); - } -}; -} -} - // template used to specialize a function that contains a lambda that should // result in a deferred diagnostic being emitted. @@ -62,7 +39,7 @@ void setup_sycl_operation(const T VA[]) { cl::sycl::queue deviceQueue; deviceQueue.submit([&](cl::sycl::handler &cgh) { - cgh.parallel_for(numOfItems, [=](cl::sycl::id<1> wiID) { + cgh.single_task([=]() { // FIX!! expected-error@+1 {{zero-length arrays are not permitted in C++}} int OverlookedBadArray[0]; @@ -80,7 +57,7 @@ int main(int argc, char **argv) { cl::sycl::queue deviceQueue; deviceQueue.submit([&](cl::sycl::handler &cgh) { - cgh.parallel_for(numOfItems, [=](cl::sycl::id<1> wiID) { + cgh.single_task([=]() { // expected-error@+1 {{zero-length arrays are not permitted in C++}} int BadArray[0]; From e28c09e457d415277ffa562b7fcdb88bc7c74372 Mon Sep 17 00:00:00 2001 From: cperkinsintel Date: Fri, 21 Feb 2020 13:06:37 -0800 Subject: [PATCH 10/18] Update clang/test/SemaSYCL/deferred-diagnostics-emit.cpp Signed-off-by: Chris Perkins Co-Authored-By: Alexey Bader --- clang/test/SemaSYCL/deferred-diagnostics-emit.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index c1ed067e5ce45..b9abc3d9dfe2c 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -I %S/Inputs -fcxx-exceptions -triple spir64 -fsycl-is-device -Wno-return-type -verify -fsyntax-only %s +// RUN: %clang_cc1 -I %S/Inputs -triple spir64 -fsycl-is-device -verify -fsyntax-only %s // // ensuring that the SYCL diagnostics that are typically deferred, correctly emit From 8634c6e053045fb7d5199812f6d9aeb2a6e7b959 Mon Sep 17 00:00:00 2001 From: cperkinsintel Date: Tue, 25 Feb 2020 08:49:01 -0800 Subject: [PATCH 11/18] Update clang/test/SemaSYCL/deferred-diagnostics-emit.cpp Signed-off-by: Chris Perkins chris.perkins@intel.com Co-Authored-By: Mariya Podchishchaeva --- clang/test/SemaSYCL/deferred-diagnostics-emit.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index b9abc3d9dfe2c..4d55ee272e952 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -1,7 +1,7 @@ // RUN: %clang_cc1 -I %S/Inputs -triple spir64 -fsycl-is-device -verify -fsyntax-only %s // -// ensuring that the SYCL diagnostics that are typically deferred, correctly emit +// Ensure that the SYCL diagnostics that are typically deferred, correctly emitted. // #include From 0a208a182404bdd0532b8008ee1b75a84d4e41ce Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 25 Feb 2020 11:23:29 -0800 Subject: [PATCH 12/18] updating test. Having problem with 'note' expectation. Signed-off-by: Chris Perkins --- clang/test/SemaSYCL/deferred-diagnostics-emit.cpp | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index b9abc3d9dfe2c..049438f89dfbe 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -1,5 +1,4 @@ -// RUN: %clang_cc1 -I %S/Inputs -triple spir64 -fsycl-is-device -verify -fsyntax-only %s - +// RUN: %clang_cc1 -I %S/Inputs -fsycl -triple spir64 -fsycl-is-device -verify -fsyntax-only %s // // ensuring that the SYCL diagnostics that are typically deferred, correctly emit // @@ -40,10 +39,10 @@ void setup_sycl_operation(const T VA[]) { deviceQueue.submit([&](cl::sycl::handler &cgh) { cgh.single_task([=]() { - // FIX!! expected-error@+1 {{zero-length arrays are not permitted in C++}} + // FIX!! xpected-error@+1 {{zero-length arrays are not permitted in C++}} int OverlookedBadArray[0]; - // FIX!! expected-error@+1 {{__float128 is not supported on this target}} + // FIX!! xpected-error@+1 {{__float128 is not supported on this target}} __float128 overlookedBadFloat = 40; }); @@ -56,8 +55,11 @@ int main(int argc, char **argv) { cl::sycl::range<1> numOfItems; cl::sycl::queue deviceQueue; + deviceQueue.submit([&](cl::sycl::handler &cgh) { + cgh.single_task([=]() { + // expected-error@+1 {{zero-length arrays are not permitted in C++}} int BadArray[0]; @@ -69,10 +71,11 @@ int main(int argc, char **argv) { // expected-error@+1 {{SYCL kernel cannot call a virtual function}} s.foo(); + // expected-note@+1 {{called by 'operator()'}} calledFromKernel(10); }); }); - + // --- lambda in specialized function testing --- From a1c2eb30238bb8241f9a7491306b12702527b7b0 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 26 Feb 2020 11:10:31 -0800 Subject: [PATCH 13/18] llvm-lit test passing Signed-off-by: Chris Perkins --- .../SemaSYCL/deferred-diagnostics-emit.cpp | 53 +++++++++++++------ 1 file changed, 38 insertions(+), 15 deletions(-) diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index d19c927be1cdf..2b0622ed26914 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -1,13 +1,39 @@ -// RUN: %clang_cc1 -I %S/Inputs -fsycl -triple spir64 -fsycl-is-device -verify -fsyntax-only %s +// RUN: %clang_cc1 -fsycl -triple spir64 -fsycl-is-device -verify -fsyntax-only %s // // Ensure that the SYCL diagnostics that are typically deferred, correctly emitted. // -#include +// testing that the deferred diagnostics work in conjunction with the SYCL namespaces. +inline namespace cl { +namespace sycl { +class queue { +public: + template void submit(T CGF) {} +}; -struct S { - virtual void foo() {} +template class id {}; + +template class range {}; + +class handler { +public: + template + __attribute__((sycl_kernel)) void kernel_parallel_for(KernelType kernelFunc) { + // expected-note@+1 2{{called by 'kernel_parallel_for{}); + } + template + void parallel_for(range NWI, KernelType kernelFunc) { + kernel_parallel_for(kernelFunc); + } }; +} +} + +//variadic functions from SYCL kernels emit a deferred diagnostic +void variadic(int, ...); + + int calledFromKernel(int a){ // expected-error@+1 {{zero-length arrays are not permitted in C++}} @@ -16,10 +42,8 @@ int calledFromKernel(int a){ // expected-error@+1 {{__float128 is not supported on this target}} __float128 malFloat = 40; - // not sure if 'no virtual function' is a _deferred_ diagnostic, testing anyway - S mal; - // expected-error@+1 {{SYCL kernel cannot call a virtual function}} - mal.foo(); + //expected-error@+1 {{SYCL kernel cannot call a variadic function}} + variadic(5); return a + 20; } @@ -38,7 +62,7 @@ void setup_sycl_operation(const T VA[]) { cl::sycl::queue deviceQueue; deviceQueue.submit([&](cl::sycl::handler &cgh) { - cgh.single_task([=]() { + cgh.parallel_for(numOfItems, [=](cl::sycl::id<1> wiID) { // FIX!! xpected-error@+1 {{zero-length arrays are not permitted in C++}} int OverlookedBadArray[0]; @@ -57,8 +81,8 @@ int main(int argc, char **argv) { deviceQueue.submit([&](cl::sycl::handler &cgh) { - - cgh.single_task([=]() { + + cgh.parallel_for(numOfItems, [=](cl::sycl::id<1> wiID) { // expected-error@+1 {{zero-length arrays are not permitted in C++}} int BadArray[0]; @@ -66,11 +90,10 @@ int main(int argc, char **argv) { // expected-error@+1 {{__float128 is not supported on this target}} __float128 badFloat = 40; // this SHOULD trigger a diagnostic - // not sure if 'no virtual function' is a _deferred_ diagnostic, but testing anyway. - S s; - // expected-error@+1 {{SYCL kernel cannot call a virtual function}} - s.foo(); + //expected-error@+1 {{SYCL kernel cannot call a variadic function}} + variadic(5); + // expected-note@+1 {{called by 'operator()'}} calledFromKernel(10); }); From 11b2e941d5460ec934cd9d80e6912ccc4e8a5425 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 26 Feb 2020 11:30:17 -0800 Subject: [PATCH 14/18] clang-format Signed-off-by: Chris Perkins --- .../SemaSYCL/deferred-diagnostics-emit.cpp | 52 ++++++++----------- 1 file changed, 23 insertions(+), 29 deletions(-) diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index 2b0622ed26914..3411e2f3afa9a 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -3,17 +3,20 @@ // Ensure that the SYCL diagnostics that are typically deferred, correctly emitted. // -// testing that the deferred diagnostics work in conjunction with the SYCL namespaces. +// testing that the deferred diagnostics work in conjunction with the SYCL namespaces. inline namespace cl { namespace sycl { class queue { public: - template void submit(T CGF) {} + template + void submit(T CGF) {} }; -template class id {}; +template +class id {}; -template class range {}; +template +class range {}; class handler { public: @@ -27,48 +30,44 @@ class handler { kernel_parallel_for(kernelFunc); } }; -} -} +} // namespace sycl +} // namespace cl //variadic functions from SYCL kernels emit a deferred diagnostic void variadic(int, ...); - - -int calledFromKernel(int a){ +int calledFromKernel(int a) { // expected-error@+1 {{zero-length arrays are not permitted in C++}} - int MalArray[0]; + int MalArray[0]; // expected-error@+1 {{__float128 is not supported on this target}} - __float128 malFloat = 40; + __float128 malFloat = 40; //expected-error@+1 {{SYCL kernel cannot call a variadic function}} - variadic(5); + variadic(5); return a + 20; } - // template used to specialize a function that contains a lambda that should // result in a deferred diagnostic being emitted. -// HOWEVER, this is not working presently. -// TODO: re-test after new deferred diagnostic system is merged. +// HOWEVER, this is not working presently. +// TODO: re-test after new deferred diagnostic system is merged. // restore the "FIX!!" tests below template void setup_sycl_operation(const T VA[]) { - + cl::sycl::range<1> numOfItems; cl::sycl::queue deviceQueue; deviceQueue.submit([&](cl::sycl::handler &cgh) { cgh.parallel_for(numOfItems, [=](cl::sycl::id<1> wiID) { // FIX!! xpected-error@+1 {{zero-length arrays are not permitted in C++}} - int OverlookedBadArray[0]; - + int OverlookedBadArray[0]; + // FIX!! xpected-error@+1 {{__float128 is not supported on this target}} - __float128 overlookedBadFloat = 40; - + __float128 overlookedBadFloat = 40; }); }); } @@ -79,30 +78,25 @@ int main(int argc, char **argv) { cl::sycl::range<1> numOfItems; cl::sycl::queue deviceQueue; - deviceQueue.submit([&](cl::sycl::handler &cgh) { - cgh.parallel_for(numOfItems, [=](cl::sycl::id<1> wiID) { - // expected-error@+1 {{zero-length arrays are not permitted in C++}} - int BadArray[0]; + 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); + variadic(5); - - // expected-note@+1 {{called by 'operator()'}} + // expected-note@+1 {{called by 'operator()'}} calledFromKernel(10); }); }); - // --- lambda in specialized function testing --- - //array A is only used to feed the template + //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); From 7d7a2ec7074a0451dbd7177413e6b2b9d4ac4cc7 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 26 Feb 2020 11:40:27 -0800 Subject: [PATCH 15/18] quick comm Signed-off-by: Chris Perkins --- clang/test/SemaSYCL/deferred-diagnostics-emit.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index 3411e2f3afa9a..0d556c367974d 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -1,7 +1,6 @@ // RUN: %clang_cc1 -fsycl -triple spir64 -fsycl-is-device -verify -fsyntax-only %s // // Ensure that the SYCL diagnostics that are typically deferred, correctly emitted. -// // testing that the deferred diagnostics work in conjunction with the SYCL namespaces. inline namespace cl { From 77a5c8879c9d61c746359cdcf03da4726f09ed7a Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 26 Feb 2020 16:35:56 -0800 Subject: [PATCH 16/18] cf trouble Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaType.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index e080375937bf1..4024ebc00cc9e 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -2342,7 +2342,7 @@ QualType Sema::BuildArrayType(QualType T, ArrayType::ArraySizeModifier ASM, } // Zero length arrays are disallowed in SYCL device code. - if (getLangOpts().SYCLIsDevice) + if (getLangOpts().SYCLIsDevice) SYCLDiagIfDeviceCode(ArraySize->getBeginLoc(), diag::err_typecheck_zero_array_size) << ArraySize->getSourceRange(); From f63f410c6ee11dbcfafb799333128b3a66290cb7 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 27 Feb 2020 14:27:07 -0800 Subject: [PATCH 17/18] comment Signed-off-by: Chris Perkins --- clang/test/SemaSYCL/deferred-diagnostics-emit.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index 0d556c367974d..3272d78215b5a 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -1,6 +1,6 @@ // RUN: %clang_cc1 -fsycl -triple spir64 -fsycl-is-device -verify -fsyntax-only %s // -// Ensure that the SYCL diagnostics that are typically deferred, correctly emitted. +// 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 { From dd87de274ff803562868c6b81b8f35beff4c1fec Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 28 Feb 2020 14:55:28 -0800 Subject: [PATCH 18/18] lightening test Signed-off-by: Chris Perkins --- .../SemaSYCL/deferred-diagnostics-emit.cpp | 69 ++++++------------- 1 file changed, 21 insertions(+), 48 deletions(-) diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index 3272d78215b5a..030985832f3cb 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -5,30 +5,13 @@ // testing that the deferred diagnostics work in conjunction with the SYCL namespaces. inline namespace cl { namespace sycl { -class queue { -public: - template - void submit(T CGF) {} -}; - -template -class id {}; - -template -class range {}; - -class handler { -public: - template - __attribute__((sycl_kernel)) void kernel_parallel_for(KernelType kernelFunc) { - // expected-note@+1 2{{called by 'kernel_parallel_for{}); - } - template - void parallel_for(range NWI, KernelType kernelFunc) { - kernel_parallel_for(kernelFunc); - } -}; + +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::range<1> numOfItems; - cl::sycl::queue deviceQueue; + cl::sycl::kernel_single_task([]() { + // FIX!! xpected-error@+1 {{zero-length arrays are not permitted in C++}} + int OverlookedBadArray[0]; - deviceQueue.submit([&](cl::sycl::handler &cgh) { - cgh.parallel_for(numOfItems, [=](cl::sycl::id<1> wiID) { - // 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; - }); + // 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::range<1> numOfItems; - cl::sycl::queue deviceQueue; - - deviceQueue.submit([&](cl::sycl::handler &cgh) { - cgh.parallel_for(numOfItems, [=](cl::sycl::id<1> wiID) { - // expected-error@+1 {{zero-length arrays are not permitted in C++}} - int BadArray[0]; + 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 {{__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-error@+1 {{SYCL kernel cannot call a variadic function}} + variadic(5); - // expected-note@+1 {{called by 'operator()'}} - calledFromKernel(10); - }); + // expected-note@+1 {{called by 'operator()'}} + calledFromKernel(10); }); // --- lambda in specialized function testing ---