diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 8dcce853e13e..6122a6833fcc 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -11060,10 +11060,10 @@ def err_sycl_invalid_accessor_property_list_template_param : Error< "%select{parameter pack|type|non-negative integer}1">; def warn_sycl_pass_by_value_deprecated : Warning<"Passing kernel functions by value is deprecated in SYCL 2020">, - InGroup; + InGroup, ShowInSystemHeader; def warn_sycl_pass_by_reference_future : Warning<"Passing of kernel functions by reference is a SYCL 2020 extension">, - InGroup; + InGroup, ShowInSystemHeader; def warn_sycl_attibute_function_raw_ptr : Warning<"SYCL 1.2.1 specification does not allow %0 attribute applied " "to a function with a raw pointer " @@ -11073,7 +11073,7 @@ def warn_sycl_implicit_decl : Warning<"SYCL 1.2.1 specification requires an explicit forward " "declaration for a kernel type name; your program may not " "be portable">, - InGroup, DefaultIgnore; + InGroup, ShowInSystemHeader, DefaultIgnore; def warn_sycl_restrict_recursion : Warning<"SYCL kernel cannot call a recursive function">, InGroup, DefaultError; diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index e266ae3d3bea..402663eafb59 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -281,6 +281,11 @@ ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { kernelFunc(); } +template +ATTR_SYCL_KERNEL void kernel_single_task_2017(KernelType kernelFunc) { + kernelFunc(); +} + template ATTR_SYCL_KERNEL void kernel_parallel_for(const KernelType &KernelFunc) { @@ -323,6 +328,16 @@ class handler { kernel_single_task(kernelFunc); #else kernelFunc(); +#endif + } + + template + void single_task_2017(KernelType kernelFunc) { + using NameT = typename get_kernel_name_t::name; +#ifdef __SYCL_DEVICE_ONLY__ + kernel_single_task_2017(kernelFunc); +#else + kernelFunc(); #endif } }; diff --git a/clang/test/CodeGenSYCL/kernel-by-reference.cpp b/clang/test/CodeGenSYCL/kernel-by-reference.cpp index 575ed9af6dac..6645c52e33de 100644 --- a/clang/test/CodeGenSYCL/kernel-by-reference.cpp +++ b/clang/test/CodeGenSYCL/kernel-by-reference.cpp @@ -1,42 +1,37 @@ -// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -verify -fsyntax-only -sycl-std=2017 -DSYCL2017 %s -// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -verify -fsyntax-only -sycl-std=2020 -DSYCL2020 %s -// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -verify -fsyntax-only -Wno-sycl-strict -DNODIAG %s -// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -verify -fsyntax-only -sycl-std=2020 -Wno-sycl-strict -DNODIAG %s +// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -sycl-std=2017 -DSYCL2017 %s +// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -sycl-std=2020 -DSYCL2020 %s +// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -Wno-sycl-strict -DNODIAG %s +// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -sycl-std=2020 -Wno-sycl-strict -DNODIAG %s -// SYCL 1.2/2017 - kernel functions passed directly. (Also no const requirement, though mutable lambdas never supported) -template -#if defined(SYCL2020) -// expected-warning@+2 {{Passing kernel functions by value is deprecated in SYCL 2020}} -#endif -__attribute__((sycl_kernel)) void sycl_2017_single_task(Func kernelFunc) { - kernelFunc(); -} +#include "sycl.hpp" -// SYCL 2020 - kernel functions are passed by reference. -template -#if defined(SYCL2017) -// expected-warning@+2 {{Passing of kernel functions by reference is a SYCL 2020 extension}} -#endif -__attribute__((sycl_kernel)) void sycl_2020_single_task(const Func &kernelFunc) { - kernelFunc(); -} +using namespace cl::sycl; -int do_nothing(int i) { +int simple_add(int i) { return i + 1; } // ensure both compile. int main() { - sycl_2017_single_task([]() { - do_nothing(10); + queue q; +#if defined(SYCL2020) + // expected-warning@Inputs/sycl.hpp:285 {{Passing kernel functions by value is deprecated in SYCL 2020}} + // expected-note@+3 {{in instantiation of function template specialization}} +#endif + q.submit([&](handler &h) { + h.single_task_2017([]() { simple_add(10); }); }); - sycl_2020_single_task([]() { - do_nothing(11); +#if defined(SYCL2017) + // expected-warning@Inputs/sycl.hpp:280 {{Passing of kernel functions by reference is a SYCL 2020 extension}} + // expected-note@+3 {{in instantiation of function template specialization}} +#endif + q.submit([&](handler &h) { + h.single_task([]() { simple_add(11); }); }); return 0; } #if defined(NODIAG) // expected-no-diagnostics -#endif \ No newline at end of file +#endif diff --git a/clang/test/SemaSYCL/args-size-overflow.cpp b/clang/test/SemaSYCL/args-size-overflow.cpp index d312b8f72316..fb3bcace406a 100644 --- a/clang/test/SemaSYCL/args-size-overflow.cpp +++ b/clang/test/SemaSYCL/args-size-overflow.cpp @@ -1,26 +1,16 @@ -// RUN: %clang_cc1 -fsycl -triple spir64 -fsycl-is-device -fsyntax-only -verify %s -// RUN: %clang_cc1 -fsycl -triple spir64 -Werror=sycl-strict -DERROR -fsycl-is-device -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -Wsycl-strict -sycl-std=2020 -verify %s -#include "Inputs/sycl.hpp" +#include "sycl.hpp" class Foo; -template -__attribute__((sycl_kernel)) void kernel(F KernelFunc) { - KernelFunc(); -} +using namespace cl::sycl; -template -void parallel_for(F KernelFunc) { -#ifdef ERROR - // expected-error@+4 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}} -#else - // expected-warning@+2 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}} -#endif - kernel(KernelFunc); -} +queue q; using Accessor = - cl::sycl::accessor; + accessor; + +// expected-warning@Inputs/sycl.hpp:220 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}} void use() { struct S { @@ -31,6 +21,8 @@ void use() { int Array[1991]; } Args; auto L = [=]() { (void)Args; }; - // expected-note@+1 {{in instantiation of function template specialization 'parallel_for(L); -} + q.submit([&](handler &h) { + // expected-note@+1 {{in instantiation of function template specialization}} + h.single_task(L); + }); +} \ No newline at end of file diff --git a/clang/test/SemaSYCL/implicit_kernel_type.cpp b/clang/test/SemaSYCL/implicit_kernel_type.cpp index 4083c82987c7..d014f333b983 100644 --- a/clang/test/SemaSYCL/implicit_kernel_type.cpp +++ b/clang/test/SemaSYCL/implicit_kernel_type.cpp @@ -1,47 +1,20 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s -Werror=sycl-strict -DERROR -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s -Wsycl-strict -DWARN -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -verify %s -Werror=sycl-strict +// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsycl-int-header=%t.h -fsyntax-only -sycl-std=2020 -verify %s -Werror=sycl-strict -DERROR +// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsycl-int-header=%t.h -fsyntax-only -sycl-std=2020 -verify %s -Wsycl-strict -DWARN +// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -sycl-std=2020 -verify %s -Werror=sycl-strict -// SYCL 1.2 Definitions -template -__attribute__((sycl_kernel)) void sycl_121_single_task(Func kernelFunc) { - kernelFunc(); -} - -class event {}; -class queue { -public: - template - event submit(T cgf) { return event{}; } -}; -class auto_name {}; -template -struct get_kernel_name_t { - using name = Name; -}; -class handler { -public: - template - void single_task(KernelType kernelFunc) { - using NameT = typename get_kernel_name_t::name; -#ifdef __SYCL_DEVICE_ONLY__ - sycl_121_single_task(kernelFunc); -#else - kernelFunc(); -#endif - } -}; -// -- /Definitions +#include "sycl.hpp" #ifdef __SYCL_UNNAMED_LAMBDA__ // expected-no-diagnostics #endif +using namespace cl::sycl; +// user-defined function void function() { } -// user-defined class +// user-defined struct struct myWrapper { }; @@ -50,34 +23,57 @@ class myWrapper2; int main() { queue q; -#ifndef __SYCL_UNNAMED_LAMBDA__ + +#if defined(WARN) + // expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}} + // expected-note@+7 {{InvalidKernelName1 declared here}} + // expected-note@+8 {{in instantiation of function template specialization}} +#elif defined(ERROR) + // expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}} // expected-note@+3 {{InvalidKernelName1 declared here}} - // expected-note@+4{{in instantiation of function template specialization}} - // expected-error@28 {{kernel needs to have a globally-visible name}} + // expected-note@+4 {{in instantiation of function template specialization}} +#endif class InvalidKernelName1 {}; q.submit([&](handler &h) { h.single_task([]() {}); }); -#endif + #if defined(WARN) - // expected-warning@+6 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} - // expected-note@+5 {{fake_kernel declared here}} + // expected-warning@Inputs/sycl.hpp:220 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #elif defined(ERROR) - // expected-error@+3 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} - // expected-note@+2 {{fake_kernel declared here}} + // expected-error@Inputs/sycl.hpp:220 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #endif - sycl_121_single_task([]() { function(); }); + + q.submit([&](handler &h) { +#ifndef __SYCL_UNNAMED_LAMBDA__ + // expected-note@+3 {{fake_kernel declared here}} + // expected-note@+2 {{in instantiation of function template specialization}} +#endif + h.single_task([]() { function(); }); + }); + #if defined(WARN) - // expected-warning@+6 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} - // expected-note@+5 {{fake_kernel2 declared here}} + // expected-warning@Inputs/sycl.hpp:220 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #elif defined(ERROR) - // expected-error@+3 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} - // expected-note@+2 {{fake_kernel2 declared here}} + // expected-error@Inputs/sycl.hpp:220 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} +#endif + + q.submit([&](handler &h) { +#ifndef __SYCL_UNNAMED_LAMBDA__ + // expected-note@+3 {{fake_kernel2 declared here}} + // expected-note@+2 {{in instantiation of function template specialization}} #endif - sycl_121_single_task([]() { - auto l = [](auto f) { f(); }; + h.single_task([]() { + auto l = [](auto f) { f(); }; + }); + }); + + q.submit([&](handler &h) { + h.single_task([]() { function(); }); + }); + + q.submit([&](handler &h) { + h.single_task([]() { function(); }); }); - sycl_121_single_task([]() { function(); }); - sycl_121_single_task([]() { function(); }); return 0; } diff --git a/sycl/test/basic_tests/args-size-overflow.cpp b/sycl/test/basic_tests/args-size-overflow.cpp deleted file mode 100644 index 66b24c912e45..000000000000 --- a/sycl/test/basic_tests/args-size-overflow.cpp +++ /dev/null @@ -1,21 +0,0 @@ -// RUN: %clangxx -fsycl -fsycl-device-only -Xclang -verify -Wsycl-strict %s -fsyntax-only - -#include - -using namespace cl::sycl; - -// expected-warning@../../include/sycl/CL/sycl/handler.hpp:919 {{size of kernel arguments (8068 bytes) may exceed the supported maximum of 2048 bytes on some devices}} - -int main() { - struct S { - int A; - int B; - int Array[2015]; - } Args; - queue myQueue; - myQueue.submit([&](handler &cgh) { - // expected-note@+1 {{in instantiation of function template specialization 'cl::sycl::handler::single_task}} - cgh.single_task([=]() { (void)Args; }); - }); - return 0; -}