[SYCL] Emit suppressed warnings from SYCL headers (#2575)

This updates FE tests related to the following warnings:

- Passing kernel functions by value is deprecated in SYCL 2020.
- Passing of kernel functions by reference is a SYCL 2020 extension.
- SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable.

The updated tests include the mock <sycl.hpp> header file and call the kernel invocation functions defined in the mock header.

Author: srividya-sundaram

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<Sycl2020Compat>;
+      InGroup<Sycl2020Compat>, ShowInSystemHeader;
 def warn_sycl_pass_by_reference_future 
     : Warning<"Passing of kernel functions by reference is a SYCL 2020 extension">,
-      InGroup<Sycl2017Compat>;
+      InGroup<Sycl2017Compat>, 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<SyclStrict>, DefaultIgnore;
+      InGroup<SyclStrict>, ShowInSystemHeader, DefaultIgnore;
 def warn_sycl_restrict_recursion
     : Warning<"SYCL kernel cannot call a recursive function">,
       InGroup<SyclStrict>, DefaultError;

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -281,6 +281,11 @@ ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) {
   kernelFunc();
 }
 
+template <typename KernelName = auto_name, typename KernelType>
+ATTR_SYCL_KERNEL void kernel_single_task_2017(KernelType kernelFunc) {
+  kernelFunc();
+}
+
 template <typename KernelName, typename KernelType, int Dims>
 ATTR_SYCL_KERNEL void
 kernel_parallel_for(const KernelType &KernelFunc) {
@@ -323,6 +328,16 @@ class handler {
     kernel_single_task<NameT>(kernelFunc);
 #else
     kernelFunc();
+#endif
+  }
+
+  template <typename KernelName = auto_name, typename KernelType>
+  void single_task_2017(KernelType kernelFunc) {
+    using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
+#ifdef __SYCL_DEVICE_ONLY__
+    kernel_single_task_2017<NameT>(kernelFunc);
+#else
+    kernelFunc();
 #endif
   }
 };

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 <typename name, typename Func>
-#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 <typename name, typename Func>
-#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<class sycl12>([]() {
-    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<class sycl2017>([]() { simple_add(10); });
   });
 
-  sycl_2020_single_task<class sycl2020>([]() {
-    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<class sycl2020>([]() { simple_add(11); });
   });
 
   return 0;
 }
 #if defined(NODIAG)
 // expected-no-diagnostics
-#endif
+#endif

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 <typename Name, typename F>
-__attribute__((sycl_kernel)) void kernel(F KernelFunc) {
-  KernelFunc();
-}
+using namespace cl::sycl;
 
-template <typename Name, typename F>
-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<Name>(KernelFunc);
-}
+queue q;
 
 using Accessor =
-    cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer>;
+    accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer>;
+
+// 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<Foo}}
-  parallel_for<Foo>(L);
-}
+  q.submit([&](handler &h) {
+    // expected-note@+1 {{in instantiation of function template specialization}}
+    h.single_task<class Foo>(L);
+  });
+}

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 <typename name, typename Func>
-__attribute__((sycl_kernel)) void sycl_121_single_task(Func kernelFunc) {
-  kernelFunc();
-}
-
-class event {};
-class queue {
-public:
-  template <typename T>
-  event submit(T cgf) { return event{}; }
-};
-class auto_name {};
-template <typename Name, typename Type>
-struct get_kernel_name_t {
-  using name = Name;
-};
-class handler {
-public:
-  template <typename KernelName = auto_name, typename KernelType>
-  void single_task(KernelType kernelFunc) {
-    using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
-#ifdef __SYCL_DEVICE_ONLY__
-    sycl_121_single_task<NameT>(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<InvalidKernelName1>([]() {});
   });
-#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<class fake_kernel>([]() { 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<class fake_kernel>([]() { 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<class fake_kernel2>([]() {
-    auto l = [](auto f) { f(); };
+    h.single_task<class fake_kernel2>([]() {
+      auto l = [](auto f) { f(); };
+    });
+  });
+
+  q.submit([&](handler &h) {
+    h.single_task<class myWrapper>([]() { function(); });
+  });
+
+  q.submit([&](handler &h) {
+    h.single_task<class myWrapper2>([]() { function(); });
   });
-  sycl_121_single_task<class myWrapper>([]() { function(); });
-  sycl_121_single_task<class myWrapper2>([]() { function(); });
   return 0;
 }