Make error and fix test arguments

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12519,9 +12519,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

@@ -1127,8 +1127,8 @@ 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.
-  assert(KernelParamTy->isReferenceType() && "Since SYCL 2020 kernels must be passed by reference.");
+  assert(KernelParamTy->isReferenceType() &&
+         "Kernel function must be passed by reference.");
   return KernelParamTy = KernelParamTy->getPointeeType();
 }
 
@@ -5032,13 +5032,9 @@ void SemaSYCL::CheckSYCLKernelCall(FunctionDecl *KernelFunc,
       }
   }
 
-  // check that calling kernel conforms to spec
-  QualType KernelParamTy = KernelFunc->getParamDecl(0)->getType();
-  if (! 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())
@@ -5163,8 +5159,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

@@ -542,11 +542,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) {
@@ -608,16 +603,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

@@ -20,7 +20,7 @@ class ESIMDFunctor {
 };
 
 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,17 +61,17 @@ 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();
 }
 

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
   }
 };