[SYCL] Disallow mutable lambdas (#1785)

[SYCL] Disallow mutable lambdas (#1785) 

With the latest changes in the SYCL spec kernel lambdas (and functors) must
be const references. Have updated the SYCL headers and tests with this
change. Also had to make a minor adjustment to the kernel generation code
in Clang and update many of the SemaSYCL tests, because their mock kernel
defs needed to be updated as well. Clang also emits a diagnostic if version
 flag (-std-sycl) doesn't match kernel passing (new by const reference or old
 by value) and enabled -std-sycl=2020  
 
Signed-off-by: Chris Perkins [email protected]

Author: cperkinsintel

clang/include/clang/Basic/DiagnosticGroups.td

@@ -1127,7 +1127,9 @@ def OpenMP : DiagGroup<"openmp", [
   ]>;
 
 // SYCL warnings
-def SyclStrict : DiagGroup<"sycl-strict">;
+def Sycl2017Compat : DiagGroup<"sycl-2017-compat">;
+def Sycl2020Compat : DiagGroup<"sycl-2020-compat">;
+def SyclStrict : DiagGroup<"sycl-strict", [ Sycl2017Compat, Sycl2020Compat]>;
 def SyclTarget : DiagGroup<"sycl-target">;
 
 // Backend warnings.

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10974,6 +10974,12 @@ def warn_boolean_attribute_argument_is_not_valid: Warning<
 def err_sycl_attibute_cannot_be_applied_here
     : Error<"%0 attribute cannot be applied to a "
             "static function or function in an anonymous namespace">;
+def warn_sycl_pass_by_value_deprecated 
+    : Warning<"Passing kernel functions by value is deprecated in SYCL 2020">,
+      InGroup<Sycl2020Compat>;
+def warn_sycl_pass_by_reference_future 
+    : Warning<"Passing of kernel functions by reference is a SYCL 2020 extension">,
+      InGroup<Sycl2017Compat>;
 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 "

clang/lib/Frontend/CompilerInvocation.cpp

@@ -2587,6 +2587,7 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK,
     if (const Arg *A = Args.getLastArg(OPT_sycl_std_EQ)) {
       Opts.SYCLVersion = llvm::StringSwitch<unsigned>(A->getValue())
                              .Cases("2017", "1.2.1", "121", "sycl-1.2.1", 2017)
+                             .Case("2020", 2020)
                              .Default(0U);
 
       if (Opts.SYCLVersion == 0U) {

clang/lib/Sema/SemaSYCL.cpp

@@ -698,7 +698,38 @@ getKernelInvocationKind(FunctionDecl *KernelCallerFunc) {
 }
 
 static const CXXRecordDecl *getKernelObjectType(FunctionDecl *Caller) {
-  return (*Caller->param_begin())->getType()->getAsCXXRecordDecl();
+  QualType KernelParamTy = (*Caller->param_begin())->getType();
+  // In SYCL 2020 kernels are now passed by reference.
+  if (KernelParamTy->isReferenceType())
+    return KernelParamTy->getPointeeCXXRecordDecl();
+
+  // SYCL 1.2.1
+  return KernelParamTy->getAsCXXRecordDecl();
+}
+
+static void checkKernelAndCaller(Sema &SemaRef, FunctionDecl *Caller,
+                                 const CXXRecordDecl *KernelObj) {
+  // check captures
+  if (KernelObj->isLambda()) {
+    for (const LambdaCapture &LC : KernelObj->captures())
+      if (LC.capturesThis() && LC.isImplicit())
+        SemaRef.Diag(LC.getLocation(), diag::err_implicit_this_capture);
+  }
+
+  // check that calling kernel conforms to spec
+  assert(Caller->param_size() >= 1 && "missing kernel function argument.");
+  QualType KernelParamTy = (*Caller->param_begin())->getType();
+  if (KernelParamTy->isReferenceType()) {
+    // passing by reference, so emit warning if not using SYCL 2020
+    if (SemaRef.LangOpts.SYCLVersion < 2020)
+      SemaRef.Diag(Caller->getLocation(),
+                   diag::warn_sycl_pass_by_reference_future);
+  } else {
+    // passing by value.  emit warning if using SYCL 2020 or greater
+    if (SemaRef.LangOpts.SYCLVersion > 2017)
+      SemaRef.Diag(Caller->getLocation(),
+                   diag::warn_sycl_pass_by_value_deprecated);
+  }
 }
 
 /// Creates a kernel parameter descriptor
@@ -1913,11 +1944,8 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
       constructKernelName(*this, KernelCallerFunc, MC);
   StringRef KernelName(getLangOpts().SYCLUnnamedLambda ? StableName
                                                        : CalculatedName);
-  if (KernelObj->isLambda()) {
-    for (const LambdaCapture &LC : KernelObj->captures())
-      if (LC.capturesThis() && LC.isImplicit())
-        Diag(LC.getLocation(), diag::err_implicit_this_capture);
-  }
+
+  checkKernelAndCaller(*this, KernelCallerFunc, KernelObj);
   SyclKernelFieldChecker checker(*this);
   SyclKernelDeclCreator kernel_decl(
       *this, checker, KernelName, KernelObj->getLocation(),

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -260,26 +260,26 @@ class spec_constant {
 
 #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
 template <typename KernelName = auto_name, typename KernelType>
-ATTR_SYCL_KERNEL void kernel_single_task(KernelType kernelFunc) {
+ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) {
   kernelFunc();
 }
 
 template <typename KernelName, typename KernelType, int Dims>
 ATTR_SYCL_KERNEL void
-kernel_parallel_for(KernelType KernelFunc) {
+kernel_parallel_for(const KernelType &KernelFunc) {
   KernelFunc(id<Dims>());
 }
 
 template <typename KernelName, typename KernelType, int Dims>
 ATTR_SYCL_KERNEL void
-kernel_parallel_for_work_group(KernelType KernelFunc) {
+kernel_parallel_for_work_group(const KernelType &KernelFunc) {
   KernelFunc(group<Dims>());
 }
 
 class handler {
 public:
   template <typename KernelName = auto_name, typename KernelType, int Dims>
-  void parallel_for(range<Dims> numWorkItems, KernelType kernelFunc) {
+  void parallel_for(range<Dims> numWorkItems, const KernelType &kernelFunc) {
     using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
     kernel_parallel_for<NameT, KernelType, Dims>(kernelFunc);
@@ -289,7 +289,7 @@ class handler {
   }
 
   template <typename KernelName = auto_name, typename KernelType, int Dims>
-  void parallel_for_work_group(range<Dims> numWorkGroups, range<Dims> WorkGroupSize, KernelType kernelFunc) {
+  void parallel_for_work_group(range<Dims> numWorkGroups, range<Dims> WorkGroupSize, const KernelType &kernelFunc) {
     using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
     kernel_parallel_for_work_group<NameT, KernelType, Dims>(kernelFunc);
@@ -300,7 +300,7 @@ class handler {
   }
 
   template <typename KernelName = auto_name, typename KernelType>
-  void single_task(KernelType kernelFunc) {
+  void single_task(const KernelType &kernelFunc) {
     using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
     kernel_single_task<NameT>(kernelFunc);

clang/test/CodeGenSYCL/address-space-cond-op.cpp

@@ -24,7 +24,7 @@ S foo(bool cond, S &lhs, S rhs) {
 }
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 

clang/test/CodeGenSYCL/address-space-new.cpp

@@ -110,13 +110,11 @@ void test() {
   // CHECK: call spir_func void @{{.*}}bar{{.*}}(%struct.{{.*}}.HasX addrspace(4)* align 4 dereferenceable(4) %[[SECOND]])
 }
 
-
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
   kernelFunc();
 }
 
-
 int main() {
   kernel_single_task<class fake_kernel>([]() { test(); });
   return 0;

clang/test/CodeGenSYCL/address-space-of-returns.cpp

@@ -29,7 +29,7 @@ A ret_agg() {
 // CHECK: define spir_func void @{{.*}}ret_agg{{.*}}(%struct.{{.*}}.A addrspace(4)* {{.*}} %agg.result)
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
   kernelFunc();
 }
 

clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp

@@ -195,7 +195,7 @@ void usages2() {
 }
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
   kernelFunc();
 }
 int main() {

clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp

@@ -6,7 +6,7 @@
 #include "sycl.hpp"
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }