[SYCL] Disallow passing kernel by value for sycl_kernel attribute

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -13172,9 +13172,8 @@ def err_sycl_num_kernel_wrong_reqd_wg_size : Error<
 def err_sycl_invalid_aspect_argument : Error<
   "%0 attribute argument is invalid; argument must be device aspect of type sycl::aspect">;
 
-def warn_sycl_pass_by_value_deprecated 
-    : Warning<"passing kernel functions by value is deprecated in SYCL 2020">,
-      InGroup<Sycl2020Compat>, ShowInSystemHeader;
+def err_sycl_kernel_pass_by_value
+    : Error<"SYCL kernel function must be passed by reference">;
 def warn_sycl_potentially_invalid_as_cast : Warning<
   "explicit cast from %0 to %1 potentially leads to an invalid address space"
   " cast in the resulting code">, InGroup<SyclStrict>,

clang/lib/Sema/SemaSYCL.cpp

@@ -1153,17 +1153,15 @@ getKernelInvocationKind(FunctionDecl *KernelCallerFunc) {
 
 // The SYCL kernel's 'object type' used for diagnostics and naming/mangling is
 // the first parameter to a function template using the sycl_kernel
-// attribute. In SYCL 1.2.1, this was passed by value,
-// and in SYCL 2020, it is passed by reference.
+// attribute.
 static QualType GetSYCLKernelObjectType(const FunctionDecl *KernelCaller) {
   assert(KernelCaller->getNumParams() > 0 && "Insufficient kernel parameters");
   QualType KernelParamTy = KernelCaller->getParamDecl(0)->getType();
 
-  // SYCL 2020 kernels are passed by reference.
+  // The kernel parameter is supposed to be a reference, but diagnostics are
+  // done at a later stage, so we accept a by-value kernel function type here.
   if (KernelParamTy->isReferenceType())
     KernelParamTy = KernelParamTy->getPointeeType();
-
-  // SYCL 1.2.1
   return KernelParamTy.getUnqualifiedType();
 }
 
@@ -5243,13 +5241,9 @@ void SemaSYCL::CheckSYCLKernelCall(FunctionDecl *KernelFunc,
       }
   }
 
-  // check that calling kernel conforms to spec
-  QualType KernelParamTy = KernelFunc->getParamDecl(0)->getType();
-  if (not KernelParamTy->isReferenceType()) {
-    // passing by value.  emit warning if using SYCL 2020 or greater
-    if (SemaRef.LangOpts.getSYCLVersion() >= LangOptions::SYCL_2020)
-      Diag(KernelFunc->getLocation(), diag::warn_sycl_pass_by_value_deprecated);
-  }
+  // SYCL only supports passing kernel functions by reference.
+  if (!KernelFunc->getParamDecl(0)->getType()->isReferenceType())
+    Diag(KernelFunc->getLocation(), diag::err_sycl_kernel_pass_by_value);
 
   // Do not visit invalid kernel object.
   if (KernelObj->isInvalidDecl())
@@ -5396,8 +5390,8 @@ void SemaSYCL::SetSYCLKernelNames() {
 //
 // Example of kernel caller function:
 //   template <typename KernelName, typename KernelType/*, ...*/>
-//   __attribute__((sycl_kernel)) void kernel_caller_function(KernelType
-//                                                            KernelFuncObj) {
+//   __attribute__((sycl_kernel))
+//   void kernel_caller_function(const KernelType &KernelFuncObj) {
 //     KernelFuncObj();
 //   }
 //

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -558,11 +558,6 @@ ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc, kernel_ha
   kernelFunc(kh);
 }
 
-template <typename KernelName, typename KernelType>
-ATTR_SYCL_KERNEL void kernel_single_task_2017(KernelType kernelFunc) { // #KernelSingleTask2017
-  kernelFunc();
-}
-
 template <typename KernelName, typename KernelType, int Dims>
 ATTR_SYCL_KERNEL void
 kernel_parallel_for(const KernelType &KernelFunc) {
@@ -624,16 +619,6 @@ class handler {
     kernel_single_task<NameT>(kernelFunc, kh);
 #else
     kernelFunc(kh);
-#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/esimd-private-global.cpp

@@ -7,7 +7,7 @@ __attribute__((opencl_private)) __attribute__((sycl_explicit_simd)) __attribute_
 // CHECK: @vc = {{.+}} i32 0, align 4 #[[ATTR:[0-9]+]]
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 

clang/test/CodeGenSYCL/esimd_metadata2.cpp

@@ -10,17 +10,17 @@ __attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func
 // CHECK-ESIMD-DAG: define {{.*}}spir_kernel void @{{.*}}kernel_cm() #{{[0-9]+}}{{.*}} !sycl_explicit_simd !{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[SGSIZE1:[0-9]+]] {{.*}}{
 // CHECK-ESIMD-DAG: define {{.*}}spir_func void @{{.*}}esimd_funcv() #{{[0-9]+}}{{.*}} !sycl_explicit_simd !{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[SGSIZE1]] {{.*}}{
 // CHECK-ESIMD-DAG: define {{.*}}spir_func void @{{.*}}shared_funcv() #{{[0-9]+}}{{.*}} {
-// CHECK-ESIMD-DAG: define linkonce_odr spir_func void @_ZN12ESIMDFunctorclEv({{.*}}) #{{[0-9]+}}{{.*}} !sycl_explicit_simd !{{[0-9]+}} {{.*}}{
+// CHECK-ESIMD-DAG: define linkonce_odr spir_func void @_ZNK12ESIMDFunctorclEv({{.*}}) #{{[0-9]+}}{{.*}} !sycl_explicit_simd !{{[0-9]+}} {{.*}}{
 
 class ESIMDFunctor {
 public:
-  void operator()() __attribute__((sycl_explicit_simd)) {
+  void operator()() const __attribute__((sycl_explicit_simd)) {
     esimd_func();
   }
 };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 

clang/test/CodeGenSYCL/union-kernel-param-ih.cpp

@@ -35,7 +35,7 @@ union MyUnion {
 };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 

clang/test/CodeGenSYCL/union-kernel-param.cpp

@@ -9,7 +9,7 @@ union MyUnion {
 };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 

clang/test/CodeGenSYCL/unique_stable_name.cpp

@@ -61,29 +61,29 @@ struct Derp {
 };
 
 template <typename KernelName, typename KernelType>
-[[clang::sycl_kernel]] void kernel_single_task(KernelType kernelFunc) {
+[[clang::sycl_kernel]] void kernel_single_task(const KernelType &kernelFunc) {
   kernelFunc();
 }
 
 template<typename KernelType>
-void unnamed_kernel_single_task(KernelType kernelFunc) {
+void unnamed_kernel_single_task(const KernelType &kernelFunc) {
   kernel_single_task<KernelType>(kernelFunc);
 }
 
 template <typename KernelName, typename KernelType>
-void not_kernel_single_task(KernelType kernelFunc) {
+void not_kernel_single_task(const KernelType &kernelFunc) {
   kernelFunc();
 }
 
 int main() {
   not_kernel_single_task<class kernel2>(func<Derp>);
-  // CHECK: call void @_Z22not_kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_(ptr noundef @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv)
+  // CHECK: call void @_Z22not_kernel_single_taskIZ4mainE7kernel2FPKcvEEvRKT0_(ptr noundef nonnull @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv)
 
   auto l1 = []() { return 1; };
   auto l2 = [](decltype(l1) *l = nullptr) { return 2; };
   kernel_single_task<decltype(l2)>(l2);
   puts(__builtin_sycl_unique_stable_name(decltype(l2)));
-  // CHECK: call void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvT0_
+  // CHECK: call void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvRKT0_
   // CHECK: call void @puts(ptr noundef @[[LAMBDA_KERNEL3]])
 
   constexpr const char str[] = "lalala";
@@ -98,10 +98,10 @@ int main() {
   puts(__builtin_sycl_unique_stable_name(int[++j]));
   // CHECK: call void @puts(ptr noundef @[[STRING]])
 
-  // CHECK: define internal void @_Z22not_kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_
+  // CHECK: define internal void @_Z22not_kernel_single_taskIZ4mainE7kernel2FPKcvEEvRKT0_
   // CHECK: declare noundef ptr @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv
-  // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvT0_
-  // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlvE0_S0_EvT0_
+  // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvRKT0_
+  // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlvE0_S0_EvRKT0_
 
   unnamed_kernel_single_task(
       []() {

clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp

@@ -3,32 +3,32 @@
 
 
 template<typename KN, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func F){
+__attribute__((sycl_kernel)) void kernel(const Func &F){
   F();
 }
 
 template<typename Func>
-void kernel_wrapper(Func F) {
+void kernel_wrapper(const Func &F) {
   kernel<Func>(F);
 }
 
 template<typename KN, typename Func>
-__attribute__((sycl_kernel)) void kernel2(Func F){
+__attribute__((sycl_kernel)) void kernel2(const Func &F){
   F(1);
 }
 
 template<typename Func>
-void kernel2_wrapper(Func F) {
+void kernel2_wrapper(const Func &F) {
   kernel2<Func>(F);
 }
 
 template<typename KN, typename Func>
-__attribute__((sycl_kernel)) void kernel3(Func F){
+__attribute__((sycl_kernel)) void kernel3(const Func &F){
   F(1.1);
 }
 
 template<typename Func>
-void kernel3_wrapper(Func F) {
+void kernel3_wrapper(const Func &F) {
   kernel3<Func>(F);
 }
 

clang/test/Frontend/Inputs/sycl.hpp

@@ -400,11 +400,6 @@ ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc, kernel_ha
   kernelFunc(kh);
 }
 
-template <typename KernelName, typename KernelType>
-ATTR_SYCL_KERNEL void kernel_single_task_2017(KernelType kernelFunc) { // #KernelSingleTask2017
-  kernelFunc();
-}
-
 template <typename KernelName, typename KernelType, int Dims>
 ATTR_SYCL_KERNEL void
 kernel_parallel_for(const KernelType &KernelFunc) {
@@ -466,16 +461,6 @@ class handler {
     kernel_single_task<NameT>(kernelFunc, kh);
 #else
     kernelFunc(kh);
-#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/SemaSYCL/kernel-attribute-on-non-sycl.cpp

@@ -9,6 +9,6 @@
 #endif
 
 template <typename T, typename A, int B>
-__attribute__((sycl_kernel)) void foo(T P);
+__attribute__((sycl_kernel)) void foo(const T& P);
 template <typename T, typename A, int B>
-[[clang::sycl_kernel]] void foo1(T P);
+[[clang::sycl_kernel]] void foo1(const T& P);

clang/test/SemaSYCL/kernel-attribute.cpp

@@ -11,43 +11,43 @@ __attribute__((sycl_kernel)) void foo(); // expected-warning {{'sycl_kernel' att
 
 // Attribute takes no arguments
 template <typename T, typename A>
-__attribute__((sycl_kernel(1))) void foo(T P); // expected-error {{'sycl_kernel' attribute takes no arguments}}
+__attribute__((sycl_kernel(1))) void foo(const T &P); // expected-error {{'sycl_kernel' attribute takes no arguments}}
 template <typename T, typename A, int I>
-[[clang::sycl_kernel(1)]] void foo1(T P);// expected-error {{'sycl_kernel' attribute takes no arguments}}
+[[clang::sycl_kernel(1)]] void foo1(const T &P);// expected-error {{'sycl_kernel' attribute takes no arguments}}
 
 // At least two template parameters
 template <typename T>
-__attribute__((sycl_kernel)) void foo(T P); // expected-warning {{'sycl_kernel' attribute only applies to a function template with at least two template parameters}}
+__attribute__((sycl_kernel)) void foo(const T &P); // expected-warning {{'sycl_kernel' attribute only applies to a function template with at least two template parameters}}
 template <typename T>
-[[clang::sycl_kernel]] void foo1(T P); // expected-warning {{'sycl_kernel' attribute only applies to a function template with at least two template parameters}}
+[[clang::sycl_kernel]] void foo1(const T &P); // expected-warning {{'sycl_kernel' attribute only applies to a function template with at least two template parameters}}
 
 // First two template parameters cannot be non-type template parameters
 template <typename T, int A>
-__attribute__((sycl_kernel)) void foo(T P); // expected-warning {{template parameter of a function template with the 'sycl_kernel' attribute cannot be a non-type template parameter}}
+__attribute__((sycl_kernel)) void foo(const T &P); // expected-warning {{template parameter of a function template with the 'sycl_kernel' attribute cannot be a non-type template parameter}}
 template <int A, typename T>
-[[clang::sycl_kernel]] void foo1(T P); // expected-warning {{template parameter of a function template with the 'sycl_kernel' attribute cannot be a non-type template parameter}}
+[[clang::sycl_kernel]] void foo1(const T &P); // expected-warning {{template parameter of a function template with the 'sycl_kernel' attribute cannot be a non-type template parameter}}
 
 // Must return void
 template <typename T, typename A>
-__attribute__((sycl_kernel)) int foo(T P); // expected-warning {{function template with 'sycl_kernel' attribute must have a 'void' return type}}
+__attribute__((sycl_kernel)) int foo(const T &P); // expected-warning {{function template with 'sycl_kernel' attribute must have a 'void' return type}}
 template <typename T, typename A>
-[[clang::sycl_kernel]] int foo1(T P); // expected-warning {{function template with 'sycl_kernel' attribute must have a 'void' return type}}
+[[clang::sycl_kernel]] int foo1(const T &P); // expected-warning {{function template with 'sycl_kernel' attribute must have a 'void' return type}}
 
 // Must take at least one argument
 template <typename T, typename A>
 __attribute__((sycl_kernel)) void foo(); // expected-warning {{function template with 'sycl_kernel' attribute must have at least one parameter}}
 template <typename T, typename A>
-[[clang::sycl_kernel]] void foo1(T t, A a); // no diagnostics
+[[clang::sycl_kernel]] void foo1(const T &t, A a); // no diagnostics
 
 // No diagnostics
 template <typename T, typename A>
-__attribute__((sycl_kernel)) void foo(T P);
+__attribute__((sycl_kernel)) void foo(const T &P);
 template <typename T, typename A, int I>
-[[clang::sycl_kernel]] void foo1(T P);
+[[clang::sycl_kernel]] void foo1(const T &P);
 
 #else
 
 // expected-no-diagnostics
 template <typename T, typename A>
-__attribute__((sycl_kernel)) void foo(T P);
+__attribute__((sycl_kernel)) void foo(const T &P);
 #endif

clang/test/SemaSYCL/kernel-by-val-error.cpp

@@ -0,0 +1,23 @@
+
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fsycl-is-device -verify %s
+// RUN: %clang_cc1 -fsycl-is-host -DHOST -fsyntax-only -verify %s
+
+// Kernel function argument must be passed by reference
+
+#ifndef HOST
+template <typename Name, typename T>
+__attribute__((sycl_kernel)) void foo(T P) {} // expected-error {{SYCL kernel function must be passed by reference}}
+template <typename Name, typename T>
+[[clang::sycl_kernel]] void bar(T P) {} // expected-error {{SYCL kernel function must be passed by reference}}
+#else
+// expected-no-diagnostics
+template <typename Name, typename T>
+__attribute__((sycl_kernel)) void foo(T P) {}
+template <typename Name, typename T>
+[[clang::sycl_kernel]] void bar(T P) {}
+#endif
+
+void F() {
+    foo<class FooKernel>([](){});
+    bar<class BarKernel>([](){});
+}

clang/test/SemaSYCL/kernel-not-functor.cpp

@@ -3,12 +3,12 @@
 // RUNX: %clang_cc1 -fsycl-is-host -fsyntax-only -verify %s
 
 template <typename Name, typename F>
-__attribute__((sycl_kernel)) void kernel(F kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const F &kernelFunc) {
   kernelFunc();
 }
 
 template <typename Name, typename F>
-void uses_kernel(F kernelFunc) {
+void uses_kernel(const F &kernelFunc) {
   // expected-error@+1{{kernel parameter must be a lambda or function object}}
   kernel<Name>(kernelFunc);
 }

clang/test/SemaSYCL/union-kernel-param.cpp

@@ -10,7 +10,7 @@ union MyUnion {
 };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 

clang/test/SemaSYCL/union-kernel-param1.cpp

@@ -4,7 +4,7 @@
 // unions containing Arrays.
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 

clang/test/SemaSYCL/unique-stable-name-multiple-target-crash.cpp

@@ -6,7 +6,8 @@
 //
 // expected-no-diagnostics
 class Empty {};
-template <typename, typename F> __attribute__((sycl_kernel)) void kernel(F) {
+template <typename, typename F> __attribute__((sycl_kernel))
+void kernel(const F&) {
     __builtin_sycl_unique_stable_name(F);
 }
 

clang/test/SemaSYCL/unique_stable_name.cpp

@@ -2,7 +2,7 @@
 // RUN: %clang_cc1 %s -std=c++17 -triple x86_64-linux-gnu -Wno-sycl-2020-compat -fsycl-is-device -verify -fsyntax-only -Wno-unused
 
 template <typename KernelName, typename KernelType>
-[[clang::sycl_kernel]] void kernel_single_task(KernelType kernelFunc) { // #kernelSingleTask
+[[clang::sycl_kernel]] void kernel_single_task(const KernelType &kernelFunc) { // #kernelSingleTask
   kernelFunc();
 }