[SYCL] Immutable kernel functions redo. (#2259)

This is the same immutable kernel functions PR that was merged (and reverted) last week. The SYCL conformance tests have been updated, so hopefully this should be mergeable.   I've added the `__SYCL_NONCONST_FUNCTOR__` as requested.

I have also hooked up the `-sycl-std=` compiler driver flag so that it will pass `DCL_SYCL_LANGUAGE_VERSION=121` if SYCL 1.2.1 is chosen or `DCL_SYCL_LANGUAGE_VERSION=2020` if SYCL 2020 is chosen.

In SYCL 2020 kernel functions are now const and passed by reference.  In this PR the existing API constant CL_SYCL_LANGUAGE_VERSION now determine another constant ( __SYCL_NONCONST_FUNCTOR__ ) that makes the kernel functions be constant references (or not). 

In the Clang Front end, both by-reference and by-value passing of the kernel functions is supported, and a diagnostic emitted if that is mismatched.  The `sycl-std` flag now supports `2020` as a valid option.  And the `sycl-std` flag will pass `-DCL_SYCL_LANGUAGE_VERSION=` to match its election.    Tests have been added to confirm that both kernel passings are supported and that the sycl-std flag is working correctly.   Additionally, there are a lot of FE tests that employ mock kernels, these have been updated to the SYCL 2020 convention avoiding the deferred maintenance.  

Signed-off-by: Chris Perkins <[email protected]>

Author: cperkinsintel

clang/include/clang/Basic/DiagnosticGroups.td

@@ -1133,7 +1133,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

@@ -11026,6 +11026,12 @@ def err_sycl_invalid_property_list_template_param : Error<
   "%select{property_list|property_list pack argument|buffer_location}0 "
   "template parameter must be a "
   "%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>;
+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/Frontend/InitPreprocessor.cpp

@@ -463,8 +463,13 @@ static void InitializeStandardPredefinedMacros(const TargetInfo &TI,
 
   if (LangOpts.SYCL) {
     // SYCL Version is set to a value when building SYCL applications
-    if (LangOpts.SYCLVersion == 2017)
+    if (LangOpts.SYCLVersion == 2017) {
       Builder.defineMacro("CL_SYCL_LANGUAGE_VERSION", "121");
+      Builder.defineMacro("SYCL_LANGUAGE_VERSION", "201707");
+    } else if (LangOpts.SYCLVersion == 2020) {
+      Builder.defineMacro("SYCL_LANGUAGE_VERSION", "202001");
+    }
+
     if (LangOpts.SYCLValueFitInMaxInt)
       Builder.defineMacro("__SYCL_ID_QUERIES_FIT_IN_INT__", "1");
   }

clang/lib/Sema/SemaSYCL.cpp

@@ -707,7 +707,14 @@ getKernelInvocationKind(FunctionDecl *KernelCallerFunc) {
 
 static const CXXRecordDecl *getKernelObjectType(FunctionDecl *Caller) {
   assert(Caller->getNumParams() > 0 && "Insufficient kernel parameters");
-  return Caller->getParamDecl(0)->getType()->getAsCXXRecordDecl();
+
+  QualType KernelParamTy = Caller->getParamDecl(0)->getType();
+  // In SYCL 2020 kernels are now passed by reference.
+  if (KernelParamTy->isReferenceType())
+    return KernelParamTy->getPointeeCXXRecordDecl();
+
+  // SYCL 1.2.1
+  return KernelParamTy->getAsCXXRecordDecl();
 }
 
 /// Creates a kernel parameter descriptor
@@ -2248,6 +2255,18 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc,
 
   SyclKernelFieldChecker FieldChecker(*this);
   SyclKernelUnionChecker UnionChecker(*this);
+  // check that calling kernel conforms to spec
+  QualType KernelParamTy = KernelFunc->getParamDecl(0)->getType();
+  if (KernelParamTy->isReferenceType()) {
+    // passing by reference, so emit warning if not using SYCL 2020
+    if (LangOpts.SYCLVersion < 2020)
+      Diag(KernelFunc->getLocation(), diag::warn_sycl_pass_by_reference_future);
+  } else {
+    // passing by value.  emit warning if using SYCL 2020 or greater
+    if (LangOpts.SYCLVersion > 2017)
+      Diag(KernelFunc->getLocation(), diag::warn_sycl_pass_by_value_deprecated);
+  }
+
   KernelObjVisitor Visitor{*this};
   DiagnosingSYCLKernel = true;
   Visitor.VisitRecordBases(KernelObj, FieldChecker, UnionChecker);

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -271,26 +271,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);
@@ -300,7 +300,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);
@@ -311,7 +311,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() {