From 694bbcbdb7e2a484f98b1f2bdb75dfa1f31c085d Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Tue, 4 Aug 2020 16:33:28 -0700
Subject: [PATCH 01/17] immutable functions redo. WIP.

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/include/clang/Basic/DiagnosticGroups.td |   4 +-
 .../clang/Basic/DiagnosticSemaKinds.td        |   9 +
 clang/lib/Frontend/CompilerInvocation.cpp     |   1 +
 clang/lib/Sema/SemaSYCL.cpp                   |  35 +++-
 clang/test/CodeGenSYCL/Inputs/sycl.hpp        |  12 +-
 .../CodeGenSYCL/address-space-cond-op.cpp     |   2 +-
 clang/test/CodeGenSYCL/address-space-new.cpp  |   4 +-
 .../CodeGenSYCL/address-space-of-returns.cpp  |   2 +-
 .../address-space-parameter-conversions.cpp   |   2 +-
 .../test/CodeGenSYCL/basic-kernel-wrapper.cpp |   2 +-
 clang/test/CodeGenSYCL/const-wg-init.cpp      |   2 +-
 .../CodeGenSYCL/debug-info-srcpos-kernel.cpp  |   2 +-
 clang/test/CodeGenSYCL/device-functions.cpp   |   2 +-
 clang/test/CodeGenSYCL/device-variables.cpp   |   2 +-
 .../emit-kernel-in-virtual-func.cpp           |   2 +-
 .../emit-kernel-in-virtual-func2.cpp          |   4 +-
 clang/test/CodeGenSYCL/esimd_metadata1.cpp    |   2 +-
 clang/test/CodeGenSYCL/fpga_pipes.cpp         |   2 +-
 clang/test/CodeGenSYCL/inheritance.cpp        |   2 +-
 clang/test/CodeGenSYCL/inline_asm.cpp         |   2 +-
 clang/test/CodeGenSYCL/inlining.cpp           |   2 +-
 clang/test/CodeGenSYCL/int_header1.cpp        |   2 +-
 .../test/CodeGenSYCL/int_header_inline_ns.cpp |   2 +-
 clang/test/CodeGenSYCL/integration_header.cpp |   2 +-
 .../CodeGenSYCL/intel-fpga-ivdep-array.cpp    |   2 +-
 .../intel-fpga-ivdep-embedded-loops.cpp       |   2 +-
 .../CodeGenSYCL/intel-fpga-ivdep-global.cpp   |   2 +-
 clang/test/CodeGenSYCL/intel-fpga-local.cpp   |   2 +-
 clang/test/CodeGenSYCL/intel-fpga-loops.cpp   |   2 +-
 .../CodeGenSYCL/intel-fpga-mem-builtin.cpp    |   2 +-
 .../intel-fpga-no-global-work-offset.cpp      |   4 +-
 clang/test/CodeGenSYCL/intel-fpga-reg.cpp     |   2 +-
 .../CodeGenSYCL/intel-max-global-work-dim.cpp |   4 +-
 .../CodeGenSYCL/intel-max-work-group-size.cpp |   4 +-
 clang/test/CodeGenSYCL/intel-restrict.cpp     |   2 +-
 .../test/CodeGenSYCL/kernel-by-reference.cpp  |  42 +++++
 clang/test/CodeGenSYCL/kernel-name.cpp        |   2 +-
 .../CodeGenSYCL/kernel-param-acc-array-ih.cpp |   2 +-
 .../CodeGenSYCL/kernel-param-acc-array.cpp    |   2 +-
 .../kernel-param-member-acc-array-ih.cpp      |   2 +-
 .../kernel-param-member-acc-array.cpp         |   2 +-
 .../CodeGenSYCL/kernel-param-pod-array-ih.cpp |   2 +-
 .../CodeGenSYCL/kernel-param-pod-array.cpp    |   2 +-
 clang/test/CodeGenSYCL/kernel_functor.cpp     |  63 +------
 .../CodeGenSYCL/kernel_name_with_typedefs.cpp |   4 +-
 clang/test/CodeGenSYCL/kernelname-enum.cpp    |  16 +-
 clang/test/CodeGenSYCL/loop_unroll.cpp        |   2 +-
 clang/test/CodeGenSYCL/module-id.cpp          |   2 +-
 clang/test/CodeGenSYCL/noexcept.cpp           |   2 +-
 .../test/CodeGenSYCL/num-simd-work-items.cpp  |   4 +-
 clang/test/CodeGenSYCL/remove-ur-inst.cpp     |   2 +-
 .../test/CodeGenSYCL/reqd-sub-group-size.cpp  |   8 +-
 .../test/CodeGenSYCL/reqd-work-group-size.cpp |   6 +-
 clang/test/CodeGenSYCL/sampler.cpp            |   2 +-
 clang/test/CodeGenSYCL/spir-calling-conv.cpp  |   2 +-
 clang/test/CodeGenSYCL/spir-enum.cpp          |   2 +-
 .../test/CodeGenSYCL/spir-opencl-version.cpp  |   2 +-
 .../CodeGenSYCL/sycl-device-static-init.cpp   |   2 +-
 .../CodeGenSYCL/sycl-multi-kernel-attr.cpp    |   5 +-
 clang/test/CodeGenSYCL/unique-stable-name.cpp |   2 +-
 clang/test/CodeGenSYCL/usm-int-header.cpp     |   2 +-
 clang/test/CodeGenSYCL/virtual-types.cpp      |   2 +-
 clang/test/CodeGenSYCL/wrapped-accessor.cpp   |   2 +-
 clang/test/SemaSYCL/Inputs/sycl.hpp           |   4 +-
 .../SemaSYCL/accessor-type-diagnostics.cpp    |   4 +-
 .../test/SemaSYCL/accessors-targets-image.cpp |   2 +-
 clang/test/SemaSYCL/accessors-targets.cpp     |   2 +-
 .../SemaSYCL/allow-constexpr-recursion.cpp    |   4 +-
 .../test/SemaSYCL/array-kernel-param-neg.cpp  |   6 +-
 clang/test/SemaSYCL/array-kernel-param.cpp    |   2 +-
 clang/test/SemaSYCL/basic-kernel-wrapper.cpp  |   2 +-
 .../SemaSYCL/built-in-type-kernel-arg.cpp     |   2 +-
 .../SemaSYCL/call-to-undefined-function.cpp   |   4 +-
 .../deferred-diagnostics-aux-builtin.cpp      |   4 +-
 .../SemaSYCL/deferred-diagnostics-emit.cpp    |   4 +-
 clang/test/SemaSYCL/fake-accessors.cpp        |   2 +-
 clang/test/SemaSYCL/float128.cpp              |   6 +-
 clang/test/SemaSYCL/forward-decl.cpp          |   4 +-
 clang/test/SemaSYCL/fpga_pipes.cpp            |   2 +-
 clang/test/SemaSYCL/implicit_kernel_type.cpp  |  50 ++++--
 clang/test/SemaSYCL/inheritance.cpp           |   2 +-
 clang/test/SemaSYCL/inline-asm.cpp            |   8 +-
 clang/test/SemaSYCL/intel-fpga-local.cpp      |   4 +-
 clang/test/SemaSYCL/intel-fpga-loops.cpp      |   4 +-
 .../test/SemaSYCL/intel-fpga-mem-builtin.cpp  |   2 +-
 clang/test/SemaSYCL/intel-fpga-reg.cpp        |   2 +-
 .../SemaSYCL/intel-max-global-work-dim.cpp    |  27 +--
 .../SemaSYCL/intel-max-work-group-size.cpp    |  19 ++-
 .../SemaSYCL/intel-reqd-work-group-size.cpp   |  28 ++--
 clang/test/SemaSYCL/intel-restrict.cpp        |   9 +-
 clang/test/SemaSYCL/kernel-function-type.cpp  |   4 +-
 clang/test/SemaSYCL/kernelname-enum.cpp       |  10 +-
 .../SemaSYCL/lambda_implicit_capture_this.cpp |   4 +-
 clang/test/SemaSYCL/loop_unroll.cpp           |   4 +-
 clang/test/SemaSYCL/mangle-kernel.cpp         |   2 +-
 .../SemaSYCL/markfunction-astconsumer.cpp     |   4 +-
 clang/test/SemaSYCL/no-vtables.cpp            |   5 +-
 clang/test/SemaSYCL/no-vtables2.cpp           |   5 +-
 clang/test/SemaSYCL/non-std-layout-param.cpp  |   9 +-
 .../non-trivially-copyable-kernel-param.cpp   |   4 +-
 clang/test/SemaSYCL/num_simd_work_items.cpp   |  16 +-
 clang/test/SemaSYCL/prohibit-thread-local.cpp |   4 +-
 .../test/SemaSYCL/reference-kernel-param.cpp  |   4 +-
 .../SemaSYCL/reqd-sub-group-size-device.cpp   |  14 +-
 .../SemaSYCL/reqd-work-group-size-device.cpp  |  18 +-
 clang/test/SemaSYCL/restrict-recursion.cpp    |  10 +-
 clang/test/SemaSYCL/restrict-recursion2.cpp   |   4 +-
 clang/test/SemaSYCL/restrict-recursion3.cpp   |   4 +-
 clang/test/SemaSYCL/restrict-recursion4.cpp   |   4 +-
 clang/test/SemaSYCL/sampler.cpp               |   2 +-
 .../SemaSYCL/spec_const_and_accesor_crash.cpp |   2 +-
 clang/test/SemaSYCL/sycl-callstack.cpp        |   7 +-
 clang/test/SemaSYCL/sycl-cconv.cpp            |   4 +-
 .../SemaSYCL/sycl-device-const-static.cpp     |   4 +-
 .../SemaSYCL/sycl-device-static-restrict.cpp  |   4 +-
 .../SemaSYCL/sycl-dllimport-dllexport.cpp     |  10 +-
 clang/test/SemaSYCL/sycl-esimd.cpp            |  12 +-
 clang/test/SemaSYCL/sycl-fptr-lambda.cpp      |   4 +-
 clang/test/SemaSYCL/sycl-pseudo-dtor.cpp      |   4 +-
 clang/test/SemaSYCL/sycl-restrict.cpp         |  10 +-
 clang/test/SemaSYCL/sycl-varargs-cconv.cpp    |   2 +-
 clang/test/SemaSYCL/tls_error.cpp             |   4 +-
 clang/test/SemaSYCL/unevaluated-function.cpp  |   4 +-
 clang/test/SemaSYCL/unnamed-kernel.cpp        |   4 +-
 clang/test/SemaSYCL/unsupported_math.cpp      |   4 +-
 clang/test/SemaSYCL/variadic-func-call.cpp    |   4 +-
 clang/test/SemaSYCL/wrapped-accessor.cpp      |   2 +-
 sycl/include/CL/sycl/handler.hpp              | 115 ++++++++++++-
 sycl/include/CL/sycl/queue.hpp                | 154 +++++++++++++++---
 sycl/test/functor/functor_inheritance.cpp     |   2 +-
 sycl/test/functor/kernel_functor.cpp          |   8 +-
 sycl/test/hier_par/hier_par_basic.cpp         |   8 +-
 sycl/test/separate-compile/same-kernel.cpp    |   4 +-
 sycl/test/sub_group/attributes.cpp            |   2 +-
 134 files changed, 641 insertions(+), 396 deletions(-)
 create mode 100644 clang/test/CodeGenSYCL/kernel-by-reference.cpp

diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td
index a55227673b65..117590d6c6ad 100644
--- a/clang/include/clang/Basic/DiagnosticGroups.td
+++ b/clang/include/clang/Basic/DiagnosticGroups.td
@@ -1128,7 +1128,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.
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 62a46831da8f..77017bdcfac0 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -10995,6 +10995,7 @@ 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">;
+<<<<<<< HEAD
 def err_sycl_compiletime_property_duplication : Error<
   "Can't apply %0 property twice to the same accessor">;
 def err_sycl_invalid_property_list_param_number : Error<
@@ -11005,6 +11006,14 @@ 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>;
+>>>>>>> immutable functions redo. WIP.
 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 "
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index 010383c1d5f5..023c31c53339 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -2586,6 +2586,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) {
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index f339c4a69b3b..969ad1c680e5 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -707,7 +707,39 @@ 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->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
@@ -2284,6 +2316,7 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
   // The first argument to the KernelCallerFunc is the lambda object.
   const CXXRecordDecl *KernelObj = getKernelObjectType(KernelCallerFunc);
   assert(KernelObj && "invalid kernel caller");
+  checkKernelAndCaller(*this, KernelCallerFunc, KernelObj);
 
   // Calculate both names, since Integration headers need both.
   std::string CalculatedName, StableName;
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index de3e8999961b..60d5d2bc2307 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/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);
diff --git a/clang/test/CodeGenSYCL/address-space-cond-op.cpp b/clang/test/CodeGenSYCL/address-space-cond-op.cpp
index b43ed4e380f4..b2428dca72df 100644
--- a/clang/test/CodeGenSYCL/address-space-cond-op.cpp
+++ b/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();
 }
 
diff --git a/clang/test/CodeGenSYCL/address-space-new.cpp b/clang/test/CodeGenSYCL/address-space-new.cpp
index 827215cec8ef..1caf5d49dd20 100644
--- a/clang/test/CodeGenSYCL/address-space-new.cpp
+++ b/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;
diff --git a/clang/test/CodeGenSYCL/address-space-of-returns.cpp b/clang/test/CodeGenSYCL/address-space-of-returns.cpp
index 69104c261ddf..24bd762bb28d 100644
--- a/clang/test/CodeGenSYCL/address-space-of-returns.cpp
+++ b/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();
 }
 
diff --git a/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp b/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp
index 632453486c2f..cbb645009bb3 100644
--- a/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp
+++ b/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() {
diff --git a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
index 31795fc73b77..7c89a0154ec2 100644
--- a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
+++ b/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();
 }
 
diff --git a/clang/test/CodeGenSYCL/const-wg-init.cpp b/clang/test/CodeGenSYCL/const-wg-init.cpp
index 07a2a746dffb..a31a008b1752 100644
--- a/clang/test/CodeGenSYCL/const-wg-init.cpp
+++ b/clang/test/CodeGenSYCL/const-wg-init.cpp
@@ -4,7 +4,7 @@
 
 template <typename KernelName, typename KernelType>
 __attribute__((sycl_kernel)) void
-kernel_parallel_for_work_group(KernelType KernelFunc) {
+kernel_parallel_for_work_group(const KernelType &KernelFunc) {
   cl::sycl::group<1> G;
   KernelFunc(G);
 }
diff --git a/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp b/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp
index fcff6aae4f76..5cfd9a939913 100644
--- a/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp
+++ b/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp
@@ -11,7 +11,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();
 }
 
diff --git a/clang/test/CodeGenSYCL/device-functions.cpp b/clang/test/CodeGenSYCL/device-functions.cpp
index d52ba4c13a7f..f2c180a80ccf 100644
--- a/clang/test/CodeGenSYCL/device-functions.cpp
+++ b/clang/test/CodeGenSYCL/device-functions.cpp
@@ -13,7 +13,7 @@ T bar(T arg) {
 }
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/device-variables.cpp b/clang/test/CodeGenSYCL/device-variables.cpp
index 83d6c4258995..e02ac0f6ae42 100644
--- a/clang/test/CodeGenSYCL/device-variables.cpp
+++ b/clang/test/CodeGenSYCL/device-variables.cpp
@@ -11,7 +11,7 @@ static constexpr int my_array[1] = {42};
 void foo(const test_type &) {}
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func.cpp b/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func.cpp
index 081d6b5d93c7..2c4734b509c6 100644
--- a/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func.cpp
+++ b/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func.cpp
@@ -1,7 +1,7 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func2.cpp b/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func2.cpp
index 2120176c5cd3..15d11b379fa2 100644
--- a/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func2.cpp
+++ b/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func2.cpp
@@ -1,8 +1,8 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
-    kernelFunc();
+__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
+  kernelFunc();
 }
 
 template <class TAR>
diff --git a/clang/test/CodeGenSYCL/esimd_metadata1.cpp b/clang/test/CodeGenSYCL/esimd_metadata1.cpp
index 554f508ef4f3..0f776e9b43dd 100644
--- a/clang/test/CodeGenSYCL/esimd_metadata1.cpp
+++ b/clang/test/CodeGenSYCL/esimd_metadata1.cpp
@@ -9,7 +9,7 @@
 // 3. Proper module !spirv.Source metadata is generated
 
 template <typename name, typename Func>
-void kernel(Func f) __attribute__((sycl_kernel)) {
+void kernel(const Func &f) __attribute__((sycl_kernel)) {
   f();
 }
 
diff --git a/clang/test/CodeGenSYCL/fpga_pipes.cpp b/clang/test/CodeGenSYCL/fpga_pipes.cpp
index bf1ed38afe24..237f9988f589 100644
--- a/clang/test/CodeGenSYCL/fpga_pipes.cpp
+++ b/clang/test/CodeGenSYCL/fpga_pipes.cpp
@@ -45,7 +45,7 @@ class pipe {
 };
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/inheritance.cpp b/clang/test/CodeGenSYCL/inheritance.cpp
index 4ac785336fb3..f648b244e552 100644
--- a/clang/test/CodeGenSYCL/inheritance.cpp
+++ b/clang/test/CodeGenSYCL/inheritance.cpp
@@ -24,7 +24,7 @@ struct base {
 struct derived : base, second_base {
   int a;
 
-  void operator()() {
+  void operator()() const {
   }
 };
 
diff --git a/clang/test/CodeGenSYCL/inline_asm.cpp b/clang/test/CodeGenSYCL/inline_asm.cpp
index db1956673a15..be63c64f07e7 100644
--- a/clang/test/CodeGenSYCL/inline_asm.cpp
+++ b/clang/test/CodeGenSYCL/inline_asm.cpp
@@ -3,7 +3,7 @@
 class kernel;
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
   // CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [100 x i32], align 4
   // CHECK: %[[IDX:.*]] = getelementptr inbounds [100 x i32], [100 x i32]* %[[ARRAY_A]], i64 0, i64 0
   int a[100], i = 0;
diff --git a/clang/test/CodeGenSYCL/inlining.cpp b/clang/test/CodeGenSYCL/inlining.cpp
index da01f2b70a8c..a816d16f88a0 100644
--- a/clang/test/CodeGenSYCL/inlining.cpp
+++ b/clang/test/CodeGenSYCL/inlining.cpp
@@ -1,7 +1,7 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice %s -S -emit-llvm -o - | FileCheck %s
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/int_header1.cpp b/clang/test/CodeGenSYCL/int_header1.cpp
index 65b1b6cdab05..76e37c4a45b5 100644
--- a/clang/test/CodeGenSYCL/int_header1.cpp
+++ b/clang/test/CodeGenSYCL/int_header1.cpp
@@ -22,7 +22,7 @@
 #include "sycl.hpp"
 
 template <typename KernelName, typename KernelType>
-__attribute__((sycl_kernel)) void kernel_single_task(KernelType kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/int_header_inline_ns.cpp b/clang/test/CodeGenSYCL/int_header_inline_ns.cpp
index dedeb0345f0c..efe4e12e9853 100644
--- a/clang/test/CodeGenSYCL/int_header_inline_ns.cpp
+++ b/clang/test/CodeGenSYCL/int_header_inline_ns.cpp
@@ -8,7 +8,7 @@
 #include "sycl.hpp"
 
 template <typename KernelName, typename KernelType>
-__attribute__((sycl_kernel)) void kernel_single_task(KernelType kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp
index e285bfce8f53..65c1e22be188 100644
--- a/clang/test/CodeGenSYCL/integration_header.cpp
+++ b/clang/test/CodeGenSYCL/integration_header.cpp
@@ -68,7 +68,7 @@
 #include "sycl.hpp"
 
 template <typename KernelName, typename KernelType>
-__attribute__((sycl_kernel)) void kernel_single_task(KernelType kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) {
   kernelFunc();
 }
 struct x {};
diff --git a/clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp b/clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp
index 707eeff1ea39..a59d4733e59c 100644
--- a/clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp
+++ b/clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp
@@ -178,7 +178,7 @@ void ivdep_struct() {
 }
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/intel-fpga-ivdep-embedded-loops.cpp b/clang/test/CodeGenSYCL/intel-fpga-ivdep-embedded-loops.cpp
index c90671c8f04a..72434fe3480a 100644
--- a/clang/test/CodeGenSYCL/intel-fpga-ivdep-embedded-loops.cpp
+++ b/clang/test/CodeGenSYCL/intel-fpga-ivdep-embedded-loops.cpp
@@ -148,7 +148,7 @@ void ivdep_embedded_multiple_dimensions() {
 }
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/intel-fpga-ivdep-global.cpp b/clang/test/CodeGenSYCL/intel-fpga-ivdep-global.cpp
index df745e91d63d..c6a029833613 100644
--- a/clang/test/CodeGenSYCL/intel-fpga-ivdep-global.cpp
+++ b/clang/test/CodeGenSYCL/intel-fpga-ivdep-global.cpp
@@ -80,7 +80,7 @@ void ivdep_conflicting_safelen() {
 }
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/intel-fpga-local.cpp b/clang/test/CodeGenSYCL/intel-fpga-local.cpp
index 41a483f33bab..ad02d32e2aca 100644
--- a/clang/test/CodeGenSYCL/intel-fpga-local.cpp
+++ b/clang/test/CodeGenSYCL/intel-fpga-local.cpp
@@ -256,7 +256,7 @@ void field_addrspace_cast() {
 }
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/intel-fpga-loops.cpp b/clang/test/CodeGenSYCL/intel-fpga-loops.cpp
index 0556feadcf7a..ce7fb46b9c66 100644
--- a/clang/test/CodeGenSYCL/intel-fpga-loops.cpp
+++ b/clang/test/CodeGenSYCL/intel-fpga-loops.cpp
@@ -95,7 +95,7 @@ void speculated_iterations() {
 }
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/intel-fpga-mem-builtin.cpp b/clang/test/CodeGenSYCL/intel-fpga-mem-builtin.cpp
index e330fffb8eb6..abc75cf10dd8 100644
--- a/clang/test/CodeGenSYCL/intel-fpga-mem-builtin.cpp
+++ b/clang/test/CodeGenSYCL/intel-fpga-mem-builtin.cpp
@@ -71,7 +71,7 @@ void foo(float *A, int *B, State *C, State &D) {
 // CHECK-DAG: attributes [[ATT]] = { readnone }
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp
index d1352b190fa9..3b2ffd1cae5e 100644
--- a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp
+++ b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp
@@ -2,11 +2,11 @@
 
 class Foo {
 public:
-  [[intelfpga::no_global_work_offset(1)]] void operator()() {}
+  [[intelfpga::no_global_work_offset(1)]] void operator()() const {}
 };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/intel-fpga-reg.cpp b/clang/test/CodeGenSYCL/intel-fpga-reg.cpp
index 9428243813f4..066ccc406493 100644
--- a/clang/test/CodeGenSYCL/intel-fpga-reg.cpp
+++ b/clang/test/CodeGenSYCL/intel-fpga-reg.cpp
@@ -143,7 +143,7 @@ void foo() {
 }
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp b/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp
index 5208db6ec390..0fcd1ab01486 100644
--- a/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp
+++ b/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp
@@ -2,11 +2,11 @@
 
 class Foo {
 public:
-  [[intelfpga::max_global_work_dim(1)]] void operator()() {}
+  [[intelfpga::max_global_work_dim(1)]] void operator()() const {}
 };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp b/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp
index 13bbb54f3419..e9b4360701e0 100644
--- a/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp
+++ b/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp
@@ -2,11 +2,11 @@
 
 class Foo {
 public:
-  [[intelfpga::max_work_group_size(1, 1, 1)]] void operator()() {}
+  [[intelfpga::max_work_group_size(1, 1, 1)]] void operator()() const {}
 };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/intel-restrict.cpp b/clang/test/CodeGenSYCL/intel-restrict.cpp
index ec4d9f968544..12b870ba625f 100644
--- a/clang/test/CodeGenSYCL/intel-restrict.cpp
+++ b/clang/test/CodeGenSYCL/intel-restrict.cpp
@@ -1,7 +1,7 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device %s -emit-llvm -triple spir64-unknown-unknown-sycldevice -o - | FileCheck %s
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/kernel-by-reference.cpp b/clang/test/CodeGenSYCL/kernel-by-reference.cpp
new file mode 100644
index 000000000000..575ed9af6dac
--- /dev/null
+++ b/clang/test/CodeGenSYCL/kernel-by-reference.cpp
@@ -0,0 +1,42 @@
+// 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
+
+// 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();
+}
+
+// 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();
+}
+
+int do_nothing(int i) {
+  return i + 1;
+}
+
+// ensure both compile.
+int main() {
+  sycl_2017_single_task<class sycl12>([]() {
+    do_nothing(10);
+  });
+
+  sycl_2020_single_task<class sycl2020>([]() {
+    do_nothing(11);
+  });
+
+  return 0;
+}
+#if defined(NODIAG)
+// expected-no-diagnostics
+#endif
\ No newline at end of file
diff --git a/clang/test/CodeGenSYCL/kernel-name.cpp b/clang/test/CodeGenSYCL/kernel-name.cpp
index 69298f9f8b16..167203fd266a 100644
--- a/clang/test/CodeGenSYCL/kernel-name.cpp
+++ b/clang/test/CodeGenSYCL/kernel-name.cpp
@@ -9,7 +9,7 @@ inline namespace cl {
 using namespace cl::sycl;
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp
index 902ad7ddc339..7bf277437a03 100644
--- a/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp
@@ -37,7 +37,7 @@
 using namespace cl::sycl;
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp
index ec8ac8bc01f5..cf17a9d7e3e8 100644
--- a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp
@@ -7,7 +7,7 @@
 using namespace cl::sycl;
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp
index ddfaca966485..242d3c2f5069 100644
--- a/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp
@@ -37,7 +37,7 @@
 using namespace cl::sycl;
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
index 1b1b25dcd3ff..b77eecfc85d6 100644
--- a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
@@ -7,7 +7,7 @@
 using namespace cl::sycl;
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp
index d4a5c8d5995a..b0dc2132fb06 100755
--- a/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp
@@ -39,7 +39,7 @@
 using namespace cl::sycl;
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp b/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp
index 3ff1669dc4d1..78fb63c71476 100644
--- a/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp
@@ -7,7 +7,7 @@
 using namespace cl::sycl;
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/kernel_functor.cpp b/clang/test/CodeGenSYCL/kernel_functor.cpp
index 8712895c2065..b75e4d465f88 100644
--- a/clang/test/CodeGenSYCL/kernel_functor.cpp
+++ b/clang/test/CodeGenSYCL/kernel_functor.cpp
@@ -9,32 +9,9 @@ constexpr auto sycl_read_write = cl::sycl::access::mode::read_write;
 constexpr auto sycl_global_buffer = cl::sycl::access::target::global_buffer;
 
 // Case 1:
-// - functor class is defined in an anonymous namespace
-// - the '()' operator:
-//   * does not have parameters (to be used in 'single_task').
-//   * has no 'const' qualifier
-namespace {
-  class Functor1 {
-  public:
-    Functor1(int X_, cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> &Acc_) :
-      X(X_), Acc(Acc_)
-    {}
-
-    void operator()() {
-      Acc.use(X);
-    }
-
-  private:
-    int X;
-    cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> Acc;
-  };
-}
-
-// Case 2:
 // - functor class is defined in a namespace
 // - the '()' operator:
 //   * does not have parameters (to be used in 'single_task').
-//   * has the 'const' qualifier
 namespace ns {
   class Functor2 {
   public:
@@ -52,31 +29,10 @@ namespace ns {
   };
 }
 
-// Case 3:
-// - functor class is templated and defined in the translation unit scope
-// - the '()' operator:
-//   * has a parameter of type cl::sycl::id<1> (to be used in 'parallel_for').
-//   * has no 'const' qualifier
-template <typename T> class TmplFunctor {
-public:
-  TmplFunctor(T X_, cl::sycl::accessor<T, 1, sycl_read_write, sycl_global_buffer> &Acc_) :
-    X(X_), Acc(Acc_)
-  {}
-
-  void operator()(cl::sycl::id<1> id) {
-    Acc.use(id, X);
-  }
-
-private:
-  T X;
-  cl::sycl::accessor<T, 1, sycl_read_write, sycl_global_buffer> Acc;
-};
-
-// Case 4:
+// Case 2:
 // - functor class is templated and defined in the translation unit scope
 // - the '()' operator:
 //   * has a parameter of type cl::sycl::id<1> (to be used in 'parallel_for').
-//   * has the 'const' qualifier
 template <typename T> class TmplConstFunctor {
 public:
   TmplConstFunctor(T X_, cl::sycl::accessor<T, 1, sycl_read_write, sycl_global_buffer> &Acc_) :
@@ -99,12 +55,6 @@ int foo(int X) {
     cl::sycl::queue Q;
     cl::sycl::buffer<int, 1> Buf(A, 1);
 
-    Q.submit([&](cl::sycl::handler& cgh) {
-      auto Acc = Buf.get_access<sycl_read_write, sycl_global_buffer>(cgh);
-      Functor1 F(X, Acc);
-
-      cgh.single_task(F);
-    });
     Q.submit([&](cl::sycl::handler& cgh) {
       auto Acc = Buf.get_access<sycl_read_write, sycl_global_buffer>(cgh);
       ns::Functor2 F(X, Acc);
@@ -129,13 +79,6 @@ template <typename T> T bar(T X) {
   {
     cl::sycl::queue Q;
     cl::sycl::buffer<T, 1> Buf(A, ARR_LEN(A));
-
-    Q.submit([&](cl::sycl::handler& cgh) {
-      auto Acc = Buf.template get_access<sycl_read_write, sycl_global_buffer>(cgh);
-      TmplFunctor<T> F(X, Acc);
-
-      cgh.parallel_for(cl::sycl::range<1>(ARR_LEN(A)), F);
-    });
     // Spice with lambdas to make sure functors and lambdas work together.
     Q.submit([&](cl::sycl::handler& cgh) {
       auto Acc = Buf.template get_access<sycl_read_write, sycl_global_buffer>(cgh);
@@ -165,12 +108,8 @@ int main() {
   const int Gold2 = 80;
 
 #ifndef __SYCL_DEVICE_ONLY__
-  cl::sycl::detail::KernelInfo<Functor1>::getName();
-  // CHECK: Functor1
   cl::sycl::detail::KernelInfo<ns::Functor2>::getName();
   // CHECK: ns::Functor2
-  cl::sycl::detail::KernelInfo<TmplFunctor<int>>::getName();
-  // CHECK: TmplFunctor<int>
   cl::sycl::detail::KernelInfo<TmplConstFunctor<int>>::getName();
   // CHECK: TmplConstFunctor<int>
 #endif // __SYCL_DEVICE_ONLY__
diff --git a/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp b/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp
index 7a1bb7bb447d..543b3ddfc3f5 100644
--- a/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp
+++ b/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp
@@ -4,12 +4,12 @@
 #include "sycl.hpp"
 
 template <typename KernelName, typename KernelType>
-__attribute__((sycl_kernel)) void single_task(KernelType kernelFunc) {
+__attribute__((sycl_kernel)) void single_task(const KernelType &kernelFunc) {
   kernelFunc();
 }
 
 struct dummy_functor {
-  void operator()() {}
+  void operator()() const {}
 };
 
 typedef int int_t;
diff --git a/clang/test/CodeGenSYCL/kernelname-enum.cpp b/clang/test/CodeGenSYCL/kernelname-enum.cpp
index 8580f1c53279..3529fa1534b0 100644
--- a/clang/test/CodeGenSYCL/kernelname-enum.cpp
+++ b/clang/test/CodeGenSYCL/kernelname-enum.cpp
@@ -40,43 +40,43 @@ enum class no_type_set {
 template <no_namespace_int EnumType>
 class dummy_functor_1 {
 public:
-  void operator()() {}
+  void operator()() const {}
 };
 
 template <no_namespace_short EnumType>
 class dummy_functor_2 {
 public:
-  void operator()() {}
+  void operator()() const {}
 };
 
 template <internal::namespace_short EnumType>
 class dummy_functor_3 {
 public:
-  void operator()() {}
+  void operator()() const {}
 };
 
 template <enum_in_anonNS EnumType>
 class dummy_functor_4 {
 public:
-  void operator()() {}
+  void operator()() const {}
 };
 
 template <no_type_set EnumType>
 class dummy_functor_5 {
 public:
-  void operator()() {}
+  void operator()() const {}
 };
 
 template <unscoped_enum EnumType>
 class dummy_functor_6 {
 public:
-  void operator()() {}
+  void operator()() const {}
 };
 
 template <typename EnumType>
 class dummy_functor_7 {
 public:
-  void operator()() {}
+  void operator()() const {}
 };
 
 namespace type_argument_template_enum {
@@ -105,7 +105,7 @@ class Baz;
 template <typename EnumTypeOut, template <EnumValueIn EnumValue, typename EnumTypeIn> class T>
 class dummy_functor_8 {
 public:
-  void operator()() {}
+  void operator()() const {}
 };
 
 int main() {
diff --git a/clang/test/CodeGenSYCL/loop_unroll.cpp b/clang/test/CodeGenSYCL/loop_unroll.cpp
index 7e80c92db523..3b9928873a6e 100644
--- a/clang/test/CodeGenSYCL/loop_unroll.cpp
+++ b/clang/test/CodeGenSYCL/loop_unroll.cpp
@@ -30,7 +30,7 @@ void disable() {
 }
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/module-id.cpp b/clang/test/CodeGenSYCL/module-id.cpp
index d120ee295c28..08ad6178ddbf 100644
--- a/clang/test/CodeGenSYCL/module-id.cpp
+++ b/clang/test/CodeGenSYCL/module-id.cpp
@@ -1,7 +1,7 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/noexcept.cpp b/clang/test/CodeGenSYCL/noexcept.cpp
index 816f2c43ebfe..c0f4e55400a9 100644
--- a/clang/test/CodeGenSYCL/noexcept.cpp
+++ b/clang/test/CodeGenSYCL/noexcept.cpp
@@ -48,7 +48,7 @@ void foo_cleanup() {
 }
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/num-simd-work-items.cpp b/clang/test/CodeGenSYCL/num-simd-work-items.cpp
index 8b8b8ba22d0d..5613969b88a9 100644
--- a/clang/test/CodeGenSYCL/num-simd-work-items.cpp
+++ b/clang/test/CodeGenSYCL/num-simd-work-items.cpp
@@ -2,11 +2,11 @@
 
 class Foo {
 public:
-  [[intelfpga::num_simd_work_items(1)]] void operator()() {}
+  [[intelfpga::num_simd_work_items(1)]] void operator()() const {}
 };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/remove-ur-inst.cpp b/clang/test/CodeGenSYCL/remove-ur-inst.cpp
index 77792438d481..4b799f0996ab 100644
--- a/clang/test/CodeGenSYCL/remove-ur-inst.cpp
+++ b/clang/test/CodeGenSYCL/remove-ur-inst.cpp
@@ -4,7 +4,7 @@
 SYCL_EXTERNAL void doesNotReturn() throw() __attribute__((__noreturn__));
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp
index d95b97180b3f..1b62cfa81052 100644
--- a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp
+++ b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp
@@ -2,14 +2,14 @@
 
 class Functor16 {
 public:
-  [[cl::intel_reqd_sub_group_size(16)]] void operator()() {}
+  [[cl::intel_reqd_sub_group_size(16)]] void operator()() const {}
 };
 
 [[cl::intel_reqd_sub_group_size(8)]] void foo() {}
 
 class Functor {
 public:
-  void operator()() {
+  void operator()() const {
     foo();
   }
 };
@@ -17,11 +17,11 @@ class Functor {
 template <int SIZE>
 class Functor5 {
 public:
-  [[cl::intel_reqd_sub_group_size(SIZE)]] void operator()() {}
+  [[cl::intel_reqd_sub_group_size(SIZE)]] void operator()() const {}
 };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp
index bfb08c7ce6c2..72ba3492d2ed 100644
--- a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp
+++ b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp
@@ -2,20 +2,20 @@
 
 class Functor32x16x16 {
 public:
-  [[cl::reqd_work_group_size(32, 16, 16)]] void operator()() {}
+  [[cl::reqd_work_group_size(32, 16, 16)]] void operator()() const {}
 };
 
 [[cl::reqd_work_group_size(8, 1, 1)]] void f8x1x1() {}
 
 class Functor {
 public:
-  void operator()() {
+  void operator()() const {
     f8x1x1();
   }
 };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/sampler.cpp b/clang/test/CodeGenSYCL/sampler.cpp
index 39755aa6b9ae..68c856f4532a 100644
--- a/clang/test/CodeGenSYCL/sampler.cpp
+++ b/clang/test/CodeGenSYCL/sampler.cpp
@@ -43,7 +43,7 @@ struct sampler_wrapper {
 };
 
 template <typename KernelName, typename KernelType>
-__attribute__((sycl_kernel)) void kernel_single_task(KernelType kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/spir-calling-conv.cpp b/clang/test/CodeGenSYCL/spir-calling-conv.cpp
index bed31dcb96e4..35f1b280881e 100644
--- a/clang/test/CodeGenSYCL/spir-calling-conv.cpp
+++ b/clang/test/CodeGenSYCL/spir-calling-conv.cpp
@@ -1,7 +1,7 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/spir-enum.cpp b/clang/test/CodeGenSYCL/spir-enum.cpp
index 738d1337e02a..185ef144c7dd 100644
--- a/clang/test/CodeGenSYCL/spir-enum.cpp
+++ b/clang/test/CodeGenSYCL/spir-enum.cpp
@@ -1,7 +1,7 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/spir-opencl-version.cpp b/clang/test/CodeGenSYCL/spir-opencl-version.cpp
index 43e187b46315..572902088356 100644
--- a/clang/test/CodeGenSYCL/spir-opencl-version.cpp
+++ b/clang/test/CodeGenSYCL/spir-opencl-version.cpp
@@ -1,7 +1,7 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/sycl-device-static-init.cpp b/clang/test/CodeGenSYCL/sycl-device-static-init.cpp
index b811b1061bfd..95b220388d89 100644
--- a/clang/test/CodeGenSYCL/sycl-device-static-init.cpp
+++ b/clang/test/CodeGenSYCL/sycl-device-static-init.cpp
@@ -28,7 +28,7 @@ template <class T>
 const int BaseInit<T>::var = 9;
 template struct BaseInit<TestBaseType>;
 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() {
diff --git a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp
index 3df365b3fef1..457766d4c5c1 100644
--- a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp
+++ b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp
@@ -2,12 +2,11 @@
 
 class Functor {
 public:
-  [[cl::intel_reqd_sub_group_size(4), cl::reqd_work_group_size(32, 16, 16)]] void operator()() {}
+  [[cl::intel_reqd_sub_group_size(4), cl::reqd_work_group_size(32, 16, 16)]] void operator()() const {}
 };
 
-
 template <typename Name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/CodeGenSYCL/unique-stable-name.cpp b/clang/test/CodeGenSYCL/unique-stable-name.cpp
index f24bef81bb8c..66ca499e6cda 100644
--- a/clang/test/CodeGenSYCL/unique-stable-name.cpp
+++ b/clang/test/CodeGenSYCL/unique-stable-name.cpp
@@ -33,7 +33,7 @@ void lambda_in_dependent_function() {
   {DEF_IN_MACRO();}{DEF_IN_MACRO();}
 
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/usm-int-header.cpp b/clang/test/CodeGenSYCL/usm-int-header.cpp
index 5b3ffa00050f..298a7464ac58 100644
--- a/clang/test/CodeGenSYCL/usm-int-header.cpp
+++ b/clang/test/CodeGenSYCL/usm-int-header.cpp
@@ -17,7 +17,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();
 }
 
diff --git a/clang/test/CodeGenSYCL/virtual-types.cpp b/clang/test/CodeGenSYCL/virtual-types.cpp
index b78f93712f9a..48bb1b17456f 100644
--- a/clang/test/CodeGenSYCL/virtual-types.cpp
+++ b/clang/test/CodeGenSYCL/virtual-types.cpp
@@ -1,6 +1,6 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-linux-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
 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();
 }
 
diff --git a/clang/test/CodeGenSYCL/wrapped-accessor.cpp b/clang/test/CodeGenSYCL/wrapped-accessor.cpp
index 0355a2a17b15..bf511f6ae9fa 100644
--- a/clang/test/CodeGenSYCL/wrapped-accessor.cpp
+++ b/clang/test/CodeGenSYCL/wrapped-accessor.cpp
@@ -34,7 +34,7 @@ template <typename Acc>
 struct AccWrapper { Acc accessor; };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp
index 1b3ab8b20763..5e3478abf517 100644
--- a/clang/test/SemaSYCL/Inputs/sycl.hpp
+++ b/clang/test/SemaSYCL/Inputs/sycl.hpp
@@ -202,13 +202,13 @@ struct get_kernel_name_t<auto_name, Type> {
 };
 #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();
 }
 class handler {
 public:
   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);
diff --git a/clang/test/SemaSYCL/accessor-type-diagnostics.cpp b/clang/test/SemaSYCL/accessor-type-diagnostics.cpp
index 3a814fbf7801..e6b958d130ca 100644
--- a/clang/test/SemaSYCL/accessor-type-diagnostics.cpp
+++ b/clang/test/SemaSYCL/accessor-type-diagnostics.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -I %S/Inputs -fsycl -triple spir64 -fsycl-is-device -verify -fsyntax-only  %s
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -triple spir64 -fsycl-is-device -verify -fsyntax-only -Wno-sycl-2017-compat  %s
 //
 // Ensure SYCL type restrictions are applied to accessors as well.
 
@@ -7,7 +7,7 @@
 using namespace cl::sycl;
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/SemaSYCL/accessors-targets-image.cpp b/clang/test/SemaSYCL/accessors-targets-image.cpp
index 18fac9940cb1..385b1d11bf21 100644
--- a/clang/test/SemaSYCL/accessors-targets-image.cpp
+++ b/clang/test/SemaSYCL/accessors-targets-image.cpp
@@ -8,7 +8,7 @@
 using namespace cl::sycl;
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/SemaSYCL/accessors-targets.cpp b/clang/test/SemaSYCL/accessors-targets.cpp
index dbaab2664e95..b589bb462c69 100644
--- a/clang/test/SemaSYCL/accessors-targets.cpp
+++ b/clang/test/SemaSYCL/accessors-targets.cpp
@@ -8,7 +8,7 @@
 using namespace cl::sycl;
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/SemaSYCL/allow-constexpr-recursion.cpp b/clang/test/SemaSYCL/allow-constexpr-recursion.cpp
index fed70a02f702..6b17af6351e2 100644
--- a/clang/test/SemaSYCL/allow-constexpr-recursion.cpp
+++ b/clang/test/SemaSYCL/allow-constexpr-recursion.cpp
@@ -1,7 +1,7 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -verify -fsyntax-only -std=c++20 -Werror=vla %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -Wno-sycl-2017-compat -verify -fsyntax-only -std=c++20 -Werror=vla %s
 
 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();
 }
 
diff --git a/clang/test/SemaSYCL/array-kernel-param-neg.cpp b/clang/test/SemaSYCL/array-kernel-param-neg.cpp
index 0618014c9fb1..5ad8569feebe 100755
--- a/clang/test/SemaSYCL/array-kernel-param-neg.cpp
+++ b/clang/test/SemaSYCL/array-kernel-param-neg.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -verify -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-sycl-2017-compat -verify -fsyntax-only %s
 
 // This test checks if compiler reports compilation error on an attempt to pass
 // an array of non-trivially copyable structs as SYCL kernel parameter or
@@ -20,11 +20,11 @@ class E {
   int i[];
 
 public:
-  int operator()() { return i[0]; }
+  int operator()() const { return i[0]; }
 };
 
 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();
 }
 
diff --git a/clang/test/SemaSYCL/array-kernel-param.cpp b/clang/test/SemaSYCL/array-kernel-param.cpp
index 548077be4f82..2308d6a3360d 100644
--- a/clang/test/SemaSYCL/array-kernel-param.cpp
+++ b/clang/test/SemaSYCL/array-kernel-param.cpp
@@ -8,7 +8,7 @@
 using namespace cl::sycl;
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp
index 54ae68ae5b8c..9d32dd144545 100644
--- a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp
+++ b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp
@@ -9,7 +9,7 @@ template <typename Acc>
 struct AccWrapper { Acc accessor; };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
index ef3751cd7655..f9a219f40752 100644
--- a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
+++ b/clang/test/SemaSYCL/built-in-type-kernel-arg.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();
 }
 
diff --git a/clang/test/SemaSYCL/call-to-undefined-function.cpp b/clang/test/SemaSYCL/call-to-undefined-function.cpp
index cc30af261936..8fd8a6fd7ca7 100644
--- a/clang/test/SemaSYCL/call-to-undefined-function.cpp
+++ b/clang/test/SemaSYCL/call-to-undefined-function.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-sycl-2017-compat -verify -fsyntax-only %s
 
 void defined() {
 }
@@ -9,7 +9,7 @@ void undefined();
 SYCL_EXTERNAL void undefinedExternal();
 
 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();
 }
 
diff --git a/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp b/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp
index 8da7b60bcdf6..b4df07772539 100644
--- a/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp
+++ b/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp
@@ -1,10 +1,10 @@
-// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fsycl -fsycl-is-device -aux-triple x86_64-unknown-linux-gnu -verify -fsyntax-only  %s
+// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fsycl -fsycl-is-device -aux-triple x86_64-unknown-linux-gnu -Wno-sycl-2017-compat -verify -fsyntax-only  %s
 
 inline namespace cl {
 namespace sycl {
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
   // expected-note@+1 {{called by 'kernel_single_task<AName, (lambda}}
   kernelFunc();
 }
diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp
index b637301fb5d2..c8369bf1d033 100644
--- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp
+++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1  -fsycl -triple spir64 -fsycl-is-device -verify -fsyntax-only  %s
+// RUN: %clang_cc1  -fsycl -triple spir64 -fsycl-is-device -Wno-sycl-2017-compat -verify -fsyntax-only  %s
 //
 // Ensure that the SYCL diagnostics that are typically deferred are correctly emitted.
 
@@ -12,7 +12,7 @@ inline namespace cl {
 namespace sycl {
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
   // expected-note@+1 3{{called by 'kernel_single_task<AName, (lambda}}
   kernelFunc();
 }
diff --git a/clang/test/SemaSYCL/fake-accessors.cpp b/clang/test/SemaSYCL/fake-accessors.cpp
index a1357083dd2c..b6faceef5826 100644
--- a/clang/test/SemaSYCL/fake-accessors.cpp
+++ b/clang/test/SemaSYCL/fake-accessors.cpp
@@ -26,7 +26,7 @@ using MyAccessorA = cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_writ
                                        cl::sycl::access::target::global_buffer>;
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/SemaSYCL/float128.cpp b/clang/test/SemaSYCL/float128.cpp
index 0904e740d508..b61787a153fd 100644
--- a/clang/test/SemaSYCL/float128.cpp
+++ b/clang/test/SemaSYCL/float128.cpp
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -verify -fsyntax-only %s
-// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsycl -fsycl-is-device -fsyntax-only %s
+// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -Wno-sycl-2017-compat -verify -fsyntax-only %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsycl -fsycl-is-device -Wno-sycl-2017-compat -fsyntax-only %s
 
 typedef __float128 BIGTY;
 
@@ -63,7 +63,7 @@ void foo2(){};
 __float128 foo(__float128 P) { return P; }
 
 template <typename Name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   // expected-note@+1 6{{called by 'kernel}}
   kernelFunc();
 }
diff --git a/clang/test/SemaSYCL/forward-decl.cpp b/clang/test/SemaSYCL/forward-decl.cpp
index 4a7a741991ed..7d252af5abc8 100644
--- a/clang/test/SemaSYCL/forward-decl.cpp
+++ b/clang/test/SemaSYCL/forward-decl.cpp
@@ -1,8 +1,8 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify -pedantic %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify -pedantic %s
 
 // expected-no-diagnostics
 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();
 }
 
diff --git a/clang/test/SemaSYCL/fpga_pipes.cpp b/clang/test/SemaSYCL/fpga_pipes.cpp
index 7f988c0ff49e..ef25438a0816 100644
--- a/clang/test/SemaSYCL/fpga_pipes.cpp
+++ b/clang/test/SemaSYCL/fpga_pipes.cpp
@@ -34,7 +34,7 @@ template <int N>
 pipe_storage Storage5 __attribute__((io_pipe_id(N)));
 
 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();
 }
 
diff --git a/clang/test/SemaSYCL/implicit_kernel_type.cpp b/clang/test/SemaSYCL/implicit_kernel_type.cpp
index 97a0ca90a2bc..022d34ad5922 100644
--- a/clang/test/SemaSYCL/implicit_kernel_type.cpp
+++ b/clang/test/SemaSYCL/implicit_kernel_type.cpp
@@ -1,13 +1,43 @@
 // RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s -Werror=sycl-strict -DERROR
-// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s -Wsycl-strict -DWARN
-// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -verify %s -Werror=sycl-strict
-#include <sycl.hpp>
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s  -Wsycl-strict -DWARN
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -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
 
 #ifdef __SYCL_UNNAMED_LAMBDA__
 // expected-no-diagnostics
 #endif
 
-using namespace cl::sycl;
+//using namespace cl::sycl;
 
 void function() {
 }
@@ -20,11 +50,11 @@ struct myWrapper {
 class myWrapper2;
 
 int main() {
-  cl::sycl::queue q;
+  queue q;
 #ifndef __SYCL_UNNAMED_LAMBDA__
   // expected-note@+1 {{InvalidKernelName1 declared here}}
   class InvalidKernelName1 {};
-  q.submit([&](cl::sycl::handler &h) {
+  q.submit([&](handler &h) {
     // expected-error@+1 {{kernel needs to have a globally-visible name}}
     h.single_task<InvalidKernelName1>([]() {});
   });
@@ -36,7 +66,7 @@ int main() {
   // 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}}
 #endif
-  cl::sycl::kernel_single_task<class fake_kernel>([]() { function(); });
+  sycl_121_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}}
@@ -44,10 +74,10 @@ int main() {
   // 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}}
 #endif
-  cl::sycl::kernel_single_task<class fake_kernel2>([]() {
+  sycl_121_single_task<class fake_kernel2>([]() {
     auto l = [](auto f) { f(); };
   });
-  cl::sycl::kernel_single_task<class myWrapper>([]() { function(); });
-  cl::sycl::kernel_single_task<class myWrapper2>([]() { function(); });
+  sycl_121_single_task<class myWrapper>([]() { function(); });
+  sycl_121_single_task<class myWrapper2>([]() { function(); });
   return 0;
 }
diff --git a/clang/test/SemaSYCL/inheritance.cpp b/clang/test/SemaSYCL/inheritance.cpp
index ff0b263449a3..b801c416cada 100644
--- a/clang/test/SemaSYCL/inheritance.cpp
+++ b/clang/test/SemaSYCL/inheritance.cpp
@@ -24,7 +24,7 @@ struct base {
 struct derived : base, second_base {
   int a;
 
-  void operator()() {
+  void operator()() const {
   }
 };
 
diff --git a/clang/test/SemaSYCL/inline-asm.cpp b/clang/test/SemaSYCL/inline-asm.cpp
index 36237ee3fd38..bb8a03b7a409 100644
--- a/clang/test/SemaSYCL/inline-asm.cpp
+++ b/clang/test/SemaSYCL/inline-asm.cpp
@@ -1,6 +1,6 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify %s -DLINUX_ASM
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify %s -DLINUX_ASM -DSPIR_CHECK -triple spir64-unknown-unknown-sycldevice
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify -triple x86_64-windows -fasm-blocks %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify %s -DLINUX_ASM
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify %s -DLINUX_ASM -DSPIR_CHECK -triple spir64-unknown-unknown-sycldevice
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify -triple x86_64-windows -fasm-blocks %s
 
 #ifndef SPIR_CHECK
 //expected-no-diagnostics
@@ -50,7 +50,7 @@ void bar() {
 }
 
 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();
 #ifdef LINUX_ASM
   __asm__("int3");
diff --git a/clang/test/SemaSYCL/intel-fpga-local.cpp b/clang/test/SemaSYCL/intel-fpga-local.cpp
index cb406b84f6e3..4d637a4a9213 100644
--- a/clang/test/SemaSYCL/intel-fpga-local.cpp
+++ b/clang/test/SemaSYCL/intel-fpga-local.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-return-type -fcxx-exceptions -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-return-type -fcxx-exceptions -fsyntax-only -ast-dump -Wno-sycl-2017-compat -verify -pedantic %s | FileCheck %s
 
 //CHECK: FunctionDecl{{.*}}check_ast
 void check_ast()
@@ -781,7 +781,7 @@ struct templ_st {
 };
 
 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();
 }
 
diff --git a/clang/test/SemaSYCL/intel-fpga-loops.cpp b/clang/test/SemaSYCL/intel-fpga-loops.cpp
index d32c3cfa4a0c..164d720eb156 100644
--- a/clang/test/SemaSYCL/intel-fpga-loops.cpp
+++ b/clang/test/SemaSYCL/intel-fpga-loops.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify -pedantic %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify -pedantic %s
 
 // Test for Intel FPGA loop attributes applied not to a loop
 void foo() {
@@ -371,7 +371,7 @@ void max_concurrency_dependent() {
 }
 
 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();
 }
 
diff --git a/clang/test/SemaSYCL/intel-fpga-mem-builtin.cpp b/clang/test/SemaSYCL/intel-fpga-mem-builtin.cpp
index fa3709b4eccb..bfc3d803ec44 100644
--- a/clang/test/SemaSYCL/intel-fpga-mem-builtin.cpp
+++ b/clang/test/SemaSYCL/intel-fpga-mem-builtin.cpp
@@ -56,7 +56,7 @@ void foo(float *A, int *B, State *C) {
 }
 
 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() {
diff --git a/clang/test/SemaSYCL/intel-fpga-reg.cpp b/clang/test/SemaSYCL/intel-fpga-reg.cpp
index 5d09258d3917..213038b44a81 100644
--- a/clang/test/SemaSYCL/intel-fpga-reg.cpp
+++ b/clang/test/SemaSYCL/intel-fpga-reg.cpp
@@ -57,7 +57,7 @@ void foo() {
 }
 
 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() {
diff --git a/clang/test/SemaSYCL/intel-max-global-work-dim.cpp b/clang/test/SemaSYCL/intel-max-global-work-dim.cpp
index d391924773f1..cfde3fd4d7c3 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim.cpp
@@ -1,15 +1,16 @@
-// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DTRIGGER_ERROR -verify
-// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 | FileCheck %s
-// RUN: %clang_cc1 -fsycl -fsycl-is-host -fsyntax-only -verify %s
+// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -Wno-sycl-2017-compat -triple spir64 -DTRIGGER_ERROR -verify
+// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -Wno-sycl-2017-compat -triple spir64 | FileCheck %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-host -Wno-sycl-2017-compat -fsyntax-only -verify %s
 
 #ifndef __SYCL_DEVICE_ONLY__
 struct FuncObj {
   [[intelfpga::max_global_work_dim(1)]] // expected-no-diagnostics
-  void operator()() {}
+  void
+  operator()() const {}
 };
 
 template <typename name, typename Func>
-void kernel(Func kernelFunc) {
+void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
@@ -23,22 +24,21 @@ void foo() {
 [[intelfpga::max_global_work_dim(2)]] void func_do_not_ignore() {}
 
 struct FuncObj {
-  [[intelfpga::max_global_work_dim(1)]]
-  void operator()() {}
+  [[intelfpga::max_global_work_dim(1)]] void operator()() const {}
 };
 
 struct TRIFuncObjGood1 {
   [[intelfpga::max_global_work_dim(0)]]
   [[intelfpga::max_work_group_size(1, 1, 1)]]
-  [[cl::reqd_work_group_size(1, 1, 1)]]
-  void operator()() {}
+  [[cl::reqd_work_group_size(1, 1, 1)]] void
+  operator()() const {}
 };
 
 struct TRIFuncObjGood2 {
   [[intelfpga::max_global_work_dim(3)]]
   [[intelfpga::max_work_group_size(8, 1, 1)]]
-  [[cl::reqd_work_group_size(4, 1, 1)]]
-  void operator()() {}
+  [[cl::reqd_work_group_size(4, 1, 1)]] void
+  operator()() const {}
 };
 
 #ifdef TRIGGER_ERROR
@@ -46,12 +46,13 @@ struct TRIFuncObjBad {
   [[intelfpga::max_global_work_dim(0)]]
   [[intelfpga::max_work_group_size(8, 8, 8)]] // expected-error{{'max_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
   [[cl::reqd_work_group_size(4, 4, 4)]] // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
-  void operator()() {}
+  void
+  operator()() const {}
 };
 #endif // TRIGGER_ERROR
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/SemaSYCL/intel-max-work-group-size.cpp b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
index f5b2a67599e8..210944f6b2fd 100644
--- a/clang/test/SemaSYCL/intel-max-work-group-size.cpp
+++ b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
@@ -1,15 +1,16 @@
-// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DTRIGGER_ERROR -verify
-// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 | FileCheck %s
-// RUN: %clang_cc1 -fsycl -fsycl-is-host -fsyntax-only -verify %s
+// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -Wno-sycl-2017-compat -triple spir64 -DTRIGGER_ERROR -verify
+// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -Wno-sycl-2017-compat -triple spir64 | FileCheck %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-host -fsyntax-only -Wno-sycl-2017-compat -verify %s
 
 #ifndef __SYCL_DEVICE_ONLY__
 struct FuncObj {
   [[intelfpga::max_work_group_size(1, 1, 1)]] // expected-no-diagnostics
-  void operator()() {}
+  void
+  operator()() const {}
 };
 
 template <typename name, typename Func>
-void kernel(Func kernelFunc) {
+void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
@@ -23,20 +24,20 @@ void foo() {
 [[intelfpga::max_work_group_size(2, 2, 2)]] void func_do_not_ignore() {}
 
 struct FuncObj {
-  [[intelfpga::max_work_group_size(4, 4, 4)]]
-  void operator()() {}
+  [[intelfpga::max_work_group_size(4, 4, 4)]] void operator()() const {}
 };
 
 #ifdef TRIGGER_ERROR
 struct DAFuncObj {
   [[intelfpga::max_work_group_size(4, 4, 4)]]
   [[cl::reqd_work_group_size(8, 8, 4)]] // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}
-  void operator()() {}
+  void
+  operator()() const {}
 };
 #endif // TRIGGER_ERROR
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp
index 361b3d083c0e..09fe0934029b 100644
--- a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp
+++ b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp
@@ -1,16 +1,16 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify -DTRIGGER_ERROR %s
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
-// RUN: %clang_cc1 -fsycl -fsycl-is-host -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-sycl-2017-compat -fsyntax-only -verify -DTRIGGER_ERROR %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-host -Wno-sycl-2017-compat -fsyntax-only -verify %s
 
 #ifndef __SYCL_DEVICE_ONLY__
 // expected-no-diagnostics
 class Functor {
 public:
-  [[intel::reqd_work_group_size(4)]] void operator()() {}
+  [[intel::reqd_work_group_size(4)]] void operator()() const {}
 };
 
 template <typename name, typename Func>
-void kernel(Func kernelFunc) {
+void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
@@ -32,50 +32,50 @@ void bar() {
 #ifdef TRIGGER_ERROR
 class Functor32 {
 public:
-  [[cl::reqd_work_group_size(32)]] void operator()() {} // expected-error {{'reqd_work_group_size' attribute requires exactly 3 arguments}}
+  [[cl::reqd_work_group_size(32)]] void operator()() const {} // expected-error {{'reqd_work_group_size' attribute requires exactly 3 arguments}}
 };
 class Functor33 {
 public:
-  [[intel::reqd_work_group_size(32, -4)]] void operator()() {} // expected-error {{'reqd_work_group_size' attribute requires a non-negative integral compile time constant expression}}
+  [[intel::reqd_work_group_size(32, -4)]] void operator()() const {} // expected-error {{'reqd_work_group_size' attribute requires a non-negative integral compile time constant expression}}
 };
 #endif // TRIGGER_ERROR
 
 class Functor16 {
 public:
-  [[intel::reqd_work_group_size(16)]] void operator()() {}
+  [[intel::reqd_work_group_size(16)]] void operator()() const {}
 };
 
 class Functor64 {
 public:
-  [[intel::reqd_work_group_size(64, 64)]] void operator()() {}
+  [[intel::reqd_work_group_size(64, 64)]] void operator()() const {}
 };
 
 class Functor16x16x16 {
 public:
-  [[intel::reqd_work_group_size(16, 16, 16)]] void operator()() {}
+  [[intel::reqd_work_group_size(16, 16, 16)]] void operator()() const {}
 };
 
 class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}}
 public:
-  [[intel::reqd_work_group_size(8)]] void operator()() { // expected-note {{conflicting attribute is here}}
+  [[intel::reqd_work_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}}
     f4x1x1();
   }
 };
 
 class Functor {
 public:
-  void operator()() {
+  void operator()() const {
     f4x1x1();
   }
 };
 
 class FunctorAttr {
 public:
-  __attribute__((reqd_work_group_size(128, 128, 128))) void operator()() {}
+  __attribute__((reqd_work_group_size(128, 128, 128))) void operator()() const {}
 };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/SemaSYCL/intel-restrict.cpp b/clang/test/SemaSYCL/intel-restrict.cpp
index b998b261f432..c878cec15be3 100644
--- a/clang/test/SemaSYCL/intel-restrict.cpp
+++ b/clang/test/SemaSYCL/intel-restrict.cpp
@@ -1,15 +1,14 @@
-// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DCHECKDIAG -verify
-// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 | FileCheck %s
+// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -Wno-sycl-2017-compat -triple spir64 -DCHECKDIAG -verify
+// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -Wno-sycl-2017-compat -triple spir64 | FileCheck %s
 
 [[intel::kernel_args_restrict]] void func_do_not_ignore() {}
 
 struct FuncObj {
-  [[intel::kernel_args_restrict]]
-  void operator()() {}
+  [[intel::kernel_args_restrict]] void operator()() const {}
 };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 #ifdef CHECKDIAG
   [[intel::kernel_args_restrict]] int invalid = 42; // expected-error{{'kernel_args_restrict' attribute only applies to functions}}
diff --git a/clang/test/SemaSYCL/kernel-function-type.cpp b/clang/test/SemaSYCL/kernel-function-type.cpp
index 7b33ec5ffd07..86b0399d79ed 100644
--- a/clang/test/SemaSYCL/kernel-function-type.cpp
+++ b/clang/test/SemaSYCL/kernel-function-type.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify %s
 // expected-no-diagnostics
 
 // The kernel_single_task call is emitted as an OpenCL kernel function. The call
@@ -29,7 +29,7 @@ void foo(j e, d k) {
 }
 
 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();
 }
 
diff --git a/clang/test/SemaSYCL/kernelname-enum.cpp b/clang/test/SemaSYCL/kernelname-enum.cpp
index 22a9f96acc50..8ed02e06f747 100644
--- a/clang/test/SemaSYCL/kernelname-enum.cpp
+++ b/clang/test/SemaSYCL/kernelname-enum.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -Wno-sycl-2017-compat -verify %s
 
 #include "sycl.hpp"
 
@@ -26,26 +26,26 @@ enum class scoped_enum_no_type_set {
 template <unscoped_enum_int EnumType>
 class dummy_functor_1 {
 public:
-  void operator()() {}
+  void operator()() const {}
 };
 
 // expected-error@+2 {{kernel name is invalid. Unscoped enum requires fixed underlying type}}
 template <unscoped_enum_no_type_set EnumType>
 class dummy_functor_2 {
 public:
-  void operator()() {}
+  void operator()() const {}
 };
 
 template <scoped_enum_int EnumType>
 class dummy_functor_3 {
 public:
-  void operator()() {}
+  void operator()() const {}
 };
 
 template <scoped_enum_no_type_set EnumType>
 class dummy_functor_4 {
 public:
-  void operator()() {}
+  void operator()() const {}
 };
 
 int main() {
diff --git a/clang/test/SemaSYCL/lambda_implicit_capture_this.cpp b/clang/test/SemaSYCL/lambda_implicit_capture_this.cpp
index 98212d3ca644..3868601f353f 100644
--- a/clang/test/SemaSYCL/lambda_implicit_capture_this.cpp
+++ b/clang/test/SemaSYCL/lambda_implicit_capture_this.cpp
@@ -1,7 +1,7 @@
-// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsyntax-only -verify %s
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify %s
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/SemaSYCL/loop_unroll.cpp b/clang/test/SemaSYCL/loop_unroll.cpp
index 2e7e205532fd..13ecae08d743 100644
--- a/clang/test/SemaSYCL/loop_unroll.cpp
+++ b/clang/test/SemaSYCL/loop_unroll.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify -pedantic %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify -pedantic %s
 
 template <int A>
 void bar() {
@@ -60,7 +60,7 @@ void foo() {
 }
 
 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();
 }
 
diff --git a/clang/test/SemaSYCL/mangle-kernel.cpp b/clang/test/SemaSYCL/mangle-kernel.cpp
index 8cc2950019b4..7b515dd50830 100644
--- a/clang/test/SemaSYCL/mangle-kernel.cpp
+++ b/clang/test/SemaSYCL/mangle-kernel.cpp
@@ -5,7 +5,7 @@
 #include <stdlib.h>
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/SemaSYCL/markfunction-astconsumer.cpp b/clang/test/SemaSYCL/markfunction-astconsumer.cpp
index c4883f64a8c3..8315c890fc35 100644
--- a/clang/test/SemaSYCL/markfunction-astconsumer.cpp
+++ b/clang/test/SemaSYCL/markfunction-astconsumer.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-return-type -verify -fsyntax-only -std=c++17 %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-return-type -verify -Wno-sycl-2017-compat -fsyntax-only -std=c++17 %s
 void bar();
 
 template <typename T>
@@ -7,7 +7,7 @@ void usage(T func) {
 }
 
 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();
 }
 
diff --git a/clang/test/SemaSYCL/no-vtables.cpp b/clang/test/SemaSYCL/no-vtables.cpp
index 1891ed278d44..68f9953bf678 100644
--- a/clang/test/SemaSYCL/no-vtables.cpp
+++ b/clang/test/SemaSYCL/no-vtables.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -verify -fsyntax-only -emit-llvm-only %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -verify -Wno-sycl-2017-compat -fsyntax-only -emit-llvm-only %s
 // expected-no-diagnostics
 // Should never fail, since the type is never used in kernel code.
 
@@ -17,9 +17,8 @@ void always_uses() {
 void usage() {
 }
 
-
 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() {
diff --git a/clang/test/SemaSYCL/no-vtables2.cpp b/clang/test/SemaSYCL/no-vtables2.cpp
index e2468121291f..92bde4db57b3 100644
--- a/clang/test/SemaSYCL/no-vtables2.cpp
+++ b/clang/test/SemaSYCL/no-vtables2.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-return-type -verify -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-return-type -verify -Wno-sycl-2017-compat -fsyntax-only %s
 
 struct Base {
   virtual void f() const {}
@@ -28,9 +28,8 @@ Inherit usage() {
   usage_child();
 }
 
-
 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() {
diff --git a/clang/test/SemaSYCL/non-std-layout-param.cpp b/clang/test/SemaSYCL/non-std-layout-param.cpp
index 0a8ffbf798f9..0698b3b7573a 100644
--- a/clang/test/SemaSYCL/non-std-layout-param.cpp
+++ b/clang/test/SemaSYCL/non-std-layout-param.cpp
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-std-layout-kernel-params -verify -fsyntax-only %s
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-std-layout-kernel-params -verify -Wno-sycl-2017-compat -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-sycl-2017-compat -fsyntax-only %s
 
 // This test checks if compiler reports compilation error on an attempt to pass
 // non-standard layout struct object as SYCL kernel parameter.
@@ -15,11 +15,10 @@ struct C : public Base {
 };
 
 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();
 }
 
-
 void test() {
   C C0;
   C0.Y=0;
@@ -30,7 +29,7 @@ void test() {
 }
 
 struct Kernel {
-  void operator()() {
+  void operator()() const {
     (void) c1;
     (void) c2;
     (void) p;
diff --git a/clang/test/SemaSYCL/non-trivially-copyable-kernel-param.cpp b/clang/test/SemaSYCL/non-trivially-copyable-kernel-param.cpp
index 4453e5273f9a..00d1ec0bc041 100644
--- a/clang/test/SemaSYCL/non-trivially-copyable-kernel-param.cpp
+++ b/clang/test/SemaSYCL/non-trivially-copyable-kernel-param.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -Wno-sycl-2017-compat -fsyntax-only %s
 
 // This test checks if compiler reports compilation error on an attempt to pass
 // a struct with non-trivially copyable type as SYCL kernel parameter.
@@ -22,7 +22,7 @@ struct D {
 };
 
 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();
 }
 
diff --git a/clang/test/SemaSYCL/num_simd_work_items.cpp b/clang/test/SemaSYCL/num_simd_work_items.cpp
index 638539917a28..d6528ec4e665 100644
--- a/clang/test/SemaSYCL/num_simd_work_items.cpp
+++ b/clang/test/SemaSYCL/num_simd_work_items.cpp
@@ -1,15 +1,16 @@
-// RUN: %clang_cc1 %s -fsycl -fsycl-is-device -triple spir64 -fsyntax-only -DTRIGGER_ERROR -verify
-// RUN: %clang_cc1 %s -fsycl -fsycl-is-device -triple spir64 -fsyntax-only -ast-dump | FileCheck %s
-// RUN: %clang_cc1 -fsycl -fsycl-is-host -fsyntax-only -verify %s
+// RUN: %clang_cc1 %s -fsycl -fsycl-is-device -triple spir64 -fsyntax-only -Wno-sycl-2017-compat -DTRIGGER_ERROR -verify
+// RUN: %clang_cc1 %s -fsycl -fsycl-is-device -triple spir64 -fsyntax-only -Wno-sycl-2017-compat -ast-dump | FileCheck %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-host -fsyntax-only -Wno-sycl-2017-compat -verify %s
 
 #ifndef __SYCL_DEVICE_ONLY__
 struct FuncObj {
   [[intelfpga::num_simd_work_items(42)]] // expected-no-diagnostics
-  void operator()() {}
+  void
+  operator()() const {}
 };
 
 template <typename name, typename Func>
-void kernel(Func kernelFunc) {
+void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
@@ -23,12 +24,11 @@ void foo() {
 [[intelfpga::num_simd_work_items(2)]] void func_do_not_ignore() {}
 
 struct FuncObj {
-  [[intelfpga::num_simd_work_items(42)]]
-  void operator()() {}
+  [[intelfpga::num_simd_work_items(42)]] void operator()() const {}
 };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/SemaSYCL/prohibit-thread-local.cpp b/clang/test/SemaSYCL/prohibit-thread-local.cpp
index e07353e86247..1cdfb4c5ea3b 100644
--- a/clang/test/SemaSYCL/prohibit-thread-local.cpp
+++ b/clang/test/SemaSYCL/prohibit-thread-local.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64 -verify -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64 -verify -Wno-sycl-2017-compat -fsyntax-only %s
 
 thread_local const int prohobit_ns_scope = 0;
 thread_local int prohobit_ns_scope2 = 0;
@@ -41,7 +41,7 @@ template <typename name, typename Func>
 __attribute__((sycl_kernel))
 // expected-note@+2 2{{called by}}
 void
-kernel_single_task(Func kernelFunc) { kernelFunc(); }
+kernel_single_task(const Func &kernelFunc) { kernelFunc(); }
 
 int main() {
   // expected-note@+1 2{{called by}}
diff --git a/clang/test/SemaSYCL/reference-kernel-param.cpp b/clang/test/SemaSYCL/reference-kernel-param.cpp
index 3bd548e23a6f..06809a8fcd1c 100644
--- a/clang/test/SemaSYCL/reference-kernel-param.cpp
+++ b/clang/test/SemaSYCL/reference-kernel-param.cpp
@@ -1,10 +1,10 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -Wno-sycl-2017-compat -fsyntax-only %s
 
 // This test checks if compiler reports compilation error on an attempt to pass
 // a reference as SYCL kernel parameter.
 
 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();
 }
 
diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp
index ffa4b176207a..56fca9002d43 100644
--- a/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp
+++ b/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify -DTRIGGER_ERROR %s
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify -DTRIGGER_ERROR %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s
 
 [[intel::reqd_sub_group_size(4)]] void foo() {} // expected-note {{conflicting attribute is here}}
 // expected-note@-1 {{conflicting attribute is here}}
@@ -9,30 +9,30 @@ class Functor16 {
 public:
   // expected-warning@+2 {{attribute 'intel_reqd_sub_group_size' is deprecated}}
   // expected-note@+1 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
-  [[cl::intel_reqd_sub_group_size(16)]] void operator()() {}
+  [[cl::intel_reqd_sub_group_size(16)]] void operator()() const {}
 };
 
 class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}}
 public:
-  [[intel::reqd_sub_group_size(8)]] void operator()() { // expected-note {{conflicting attribute is here}}
+  [[intel::reqd_sub_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}}
     foo();
   }
 };
 
 class Functor4 {
 public:
-  [[intel::reqd_sub_group_size(12)]] void operator()() {}
+  [[intel::reqd_sub_group_size(12)]] void operator()() const {}
 };
 
 class Functor {
 public:
-  void operator()() {
+  void operator()() const {
     foo();
   }
 };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp
index 7175322fb43c..65bb0b13474d 100644
--- a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp
+++ b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify -DTRIGGER_ERROR %s
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify -DTRIGGER_ERROR %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s
 
 [[cl::reqd_work_group_size(4, 1, 1)]] void f4x1x1() {} // expected-note {{conflicting attribute is here}}
 // expected-note@-1 {{conflicting attribute is here}}
@@ -13,7 +13,7 @@
 
 class Functor16 {
 public:
-  [[cl::reqd_work_group_size(16, 1, 1)]] [[cl::reqd_work_group_size(16, 1, 1)]] void operator()() {}
+  [[cl::reqd_work_group_size(16, 1, 1)]] [[cl::reqd_work_group_size(16, 1, 1)]] void operator()() const {}
 };
 
 #ifdef TRIGGER_ERROR
@@ -21,35 +21,35 @@ class Functor32 {
 public:
   //expected-warning@+2{{attribute 'reqd_work_group_size' is already applied with different parameters}}
   // expected-error@+1{{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
-  [[cl::reqd_work_group_size(32, 1, 1)]] [[cl::reqd_work_group_size(1, 1, 32)]] void operator()() {}
+  [[cl::reqd_work_group_size(32, 1, 1)]] [[cl::reqd_work_group_size(1, 1, 32)]] void operator()() const {}
 };
 #endif
 class Functor16x16x16 {
 public:
-  [[cl::reqd_work_group_size(16, 16, 16)]] void operator()() {}
+  [[cl::reqd_work_group_size(16, 16, 16)]] void operator()() const {}
 };
 
 class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}}
 public:
-  [[cl::reqd_work_group_size(1, 1, 8)]] void operator()() { // expected-note {{conflicting attribute is here}}
+  [[cl::reqd_work_group_size(1, 1, 8)]] void operator()() const { // expected-note {{conflicting attribute is here}}
     f4x1x1();
   }
 };
 
 class Functor {
 public:
-  void operator()() {
+  void operator()() const {
     f4x1x1();
   }
 };
 
 class FunctorAttr {
 public:
-  __attribute__((reqd_work_group_size(128, 128, 128))) void operator()() {}
+  __attribute__((reqd_work_group_size(128, 128, 128))) void operator()() const {}
 };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/SemaSYCL/restrict-recursion.cpp b/clang/test/SemaSYCL/restrict-recursion.cpp
index 3547e2f793e1..3e29271561fe 100644
--- a/clang/test/SemaSYCL/restrict-recursion.cpp
+++ b/clang/test/SemaSYCL/restrict-recursion.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -verify -fsyntax-only -std=c++17 %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -verify -Wno-sycl-2017-compat -fsyntax-only -std=c++17 %s
 
 // This recursive function is not called from sycl kernel,
 // so it should not be diagnosed.
@@ -57,7 +57,7 @@ bool isa_B(void) {
 }
 
 template <typename N, typename L>
-__attribute__((sycl_kernel)) void kernel(L l) {
+__attribute__((sycl_kernel)) void kernel(const L &l) {
   l();
 }
 
@@ -85,13 +85,13 @@ int addInt(int n, int m) {
 }
 
 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();
 }
 
 template <typename name, typename Func>
-  // expected-note@+1 2{{function implemented using recursion declared here}}
-__attribute__((sycl_kernel)) void kernel_single_task2(Func kernelFunc) {
+// expected-note@+1 2{{function implemented using recursion declared here}}
+__attribute__((sycl_kernel)) void kernel_single_task2(const Func &kernelFunc) {
   kernelFunc();
   // expected-error@+1 2{{SYCL kernel cannot call a recursive function}}
   kernel_single_task2<name, Func>(kernelFunc);
diff --git a/clang/test/SemaSYCL/restrict-recursion2.cpp b/clang/test/SemaSYCL/restrict-recursion2.cpp
index 51688fea7b1b..b9eaffa4149b 100644
--- a/clang/test/SemaSYCL/restrict-recursion2.cpp
+++ b/clang/test/SemaSYCL/restrict-recursion2.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -verify -fsyntax-only -std=c++17 %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -verify -Wno-sycl-2017-compat -fsyntax-only -std=c++17 %s
 
 // This recursive function is not called from sycl kernel,
 // so it should not be diagnosed.
@@ -71,7 +71,7 @@ int addInt(int n, int m) {
 }
 
 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();
 }
 
diff --git a/clang/test/SemaSYCL/restrict-recursion3.cpp b/clang/test/SemaSYCL/restrict-recursion3.cpp
index 1a4a67190416..f5289b3142e2 100644
--- a/clang/test/SemaSYCL/restrict-recursion3.cpp
+++ b/clang/test/SemaSYCL/restrict-recursion3.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -Wno-error=sycl-strict -verify -fsyntax-only -std=c++17 %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -Wno-sycl-2017-compat -Wno-error=sycl-strict -verify -fsyntax-only -std=c++17 %s
 
 // This recursive function is not called from sycl kernel,
 // so it should not be diagnosed.
@@ -32,7 +32,7 @@ int addInt(int n, int m) {
 
 template <typename name, typename Func>
 // expected-note@+1 2{{function implemented using recursion declared here}}
-__attribute__((sycl_kernel)) void kernel_single_task2(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task2(const Func &kernelFunc) {
   // expected-note@+1 {{called by 'kernel_single_task2}}
   kernelFunc();
   // expected-warning@+1 2{{SYCL kernel cannot call a recursive function}}
diff --git a/clang/test/SemaSYCL/restrict-recursion4.cpp b/clang/test/SemaSYCL/restrict-recursion4.cpp
index c755f0eaec63..dde093792cf2 100644
--- a/clang/test/SemaSYCL/restrict-recursion4.cpp
+++ b/clang/test/SemaSYCL/restrict-recursion4.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -Wno-error=sycl-strict -verify -fsyntax-only -std=c++17 %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -Wno-sycl-2017-compat -Wno-error=sycl-strict -verify -fsyntax-only -std=c++17 %s
 
 // This recursive function is not called from sycl kernel,
 // so it should not be diagnosed.
@@ -34,7 +34,7 @@ int addInt(int n, int m) {
 }
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
   // expected-note@+1 {{called by 'kernel_single_task}}
   kernelFunc();
 }
diff --git a/clang/test/SemaSYCL/sampler.cpp b/clang/test/SemaSYCL/sampler.cpp
index 2a3171ce06f0..a1fb423b109c 100644
--- a/clang/test/SemaSYCL/sampler.cpp
+++ b/clang/test/SemaSYCL/sampler.cpp
@@ -3,7 +3,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();
 }
 
diff --git a/clang/test/SemaSYCL/spec_const_and_accesor_crash.cpp b/clang/test/SemaSYCL/spec_const_and_accesor_crash.cpp
index cbfba6e4be32..54a6f4629026 100644
--- a/clang/test/SemaSYCL/spec_const_and_accesor_crash.cpp
+++ b/clang/test/SemaSYCL/spec_const_and_accesor_crash.cpp
@@ -5,7 +5,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();
 }
 
diff --git a/clang/test/SemaSYCL/sycl-callstack.cpp b/clang/test/SemaSYCL/sycl-callstack.cpp
index 15da9d68255f..8c8a53753743 100644
--- a/clang/test/SemaSYCL/sycl-callstack.cpp
+++ b/clang/test/SemaSYCL/sycl-callstack.cpp
@@ -1,10 +1,9 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -verify -fsyntax-only -std=c++17 %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -verify -Wno-sycl-2017-compat -fsyntax-only -std=c++17 %s
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel))
-void kernel_single_task(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
   // expected-note@+1 {{called by 'kernel_single_task}}
-    kernelFunc();
+  kernelFunc();
 }
 
 void foo() {
diff --git a/clang/test/SemaSYCL/sycl-cconv.cpp b/clang/test/SemaSYCL/sycl-cconv.cpp
index f905417027c4..a42d4e7622aa 100644
--- a/clang/test/SemaSYCL/sycl-cconv.cpp
+++ b/clang/test/SemaSYCL/sycl-cconv.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-windows-sycldevice -aux-triple x86_64-pc-windows-msvc -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-windows-sycldevice -aux-triple x86_64-pc-windows-msvc -fsyntax-only -Wno-sycl-2017-compat -verify %s
 
 // expected-no-warning@+1
 __inline __cdecl int printf(char const* const _Format, ...) { return 0; }
@@ -13,7 +13,7 @@ void bar() {
 
 template <typename name, typename Func>
 // expected-no-warning@+1
-__cdecl __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+__cdecl __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
   // expected-error@+1{{SYCL kernel cannot call a variadic function}}
   printf("cannot call from here\n");
   // expected-no-error@+1
diff --git a/clang/test/SemaSYCL/sycl-device-const-static.cpp b/clang/test/SemaSYCL/sycl-device-const-static.cpp
index 0e90057c73d6..02a97810b7e5 100644
--- a/clang/test/SemaSYCL/sycl-device-const-static.cpp
+++ b/clang/test/SemaSYCL/sycl-device-const-static.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -Wno-sycl-2017-compat -fsyntax-only %s
 
 struct Base {};
 struct S {
@@ -37,7 +37,7 @@ void usage() {
 }
 
 template <typename Name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
   // expected-error@+1{{SYCL kernel cannot use a non-const static data variable}}
   static int z;
   // expected-note-re@+3{{called by 'kernel_single_task<fake_kernel, (lambda at {{.*}})>}}
diff --git a/clang/test/SemaSYCL/sycl-device-static-restrict.cpp b/clang/test/SemaSYCL/sycl-device-static-restrict.cpp
index 03132b4530b8..b3517770e529 100644
--- a/clang/test/SemaSYCL/sycl-device-static-restrict.cpp
+++ b/clang/test/SemaSYCL/sycl-device-static-restrict.cpp
@@ -1,8 +1,8 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -Wno-sycl-2017-compat -fsyntax-only %s
 const int glob1 = 1;
 int glob2 = 2;
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
   // expected-note-re@+1{{called by 'kernel_single_task<fake_kernel, (lambda at {{.*}})>}}
   kernelFunc();
 }
diff --git a/clang/test/SemaSYCL/sycl-dllimport-dllexport.cpp b/clang/test/SemaSYCL/sycl-dllimport-dllexport.cpp
index b2c3fbf3aa1f..e03d749a09db 100644
--- a/clang/test/SemaSYCL/sycl-dllimport-dllexport.cpp
+++ b/clang/test/SemaSYCL/sycl-dllimport-dllexport.cpp
@@ -1,19 +1,19 @@
 // RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fms-extensions \
 // RUN: -aux-triple x86_64-unknown-linux-gnu -fsycl -fsycl-is-device \
-// RUN: -fsyntax-only -DWARNCHECK %s -o /dev/null 2>&1 | FileCheck %s
+// RUN: -fsyntax-only -Wno-sycl-2017-compat -DWARNCHECK %s -o /dev/null 2>&1 | FileCheck %s
 // check random triple aux-triple with sycl-device
 
-// RUN: %clang_cc1 -triple spir64-unknown-windows-sycldevice -fsyntax-only \
+// RUN: %clang_cc1 -triple spir64-unknown-windows-sycldevice -Wno-sycl-2017-compat -fsyntax-only \
 // RUN: -fms-extensions -DWARNCHECK %s -o /dev/null 2>&1 | FileCheck %s
 // check without -aux-triple but sycl-device
 
 // RUN: %clang_cc1 -triple spir64-unknown-windows-sycldevice -fsycl \
 // RUN: -fsycl-is-device -aux-triple x86_64-pc-windows-msvc -fms-extensions \
-// RUN: -fsyntax-only -DWARNCHECK %s -o /dev/null 2>&1 | \
+// RUN: -fsyntax-only -Wno-sycl-2017-compat -DWARNCHECK %s -o /dev/null 2>&1 | \
 // RUN: FileCheck %s --check-prefixes CHECKALL
 // check -aux-tripe without sycl-device
 
-// RUN: %clang_cc1 -triple spir64-unknown-windows-sycldevice -fsyntax-only \
+// RUN: %clang_cc1 -triple spir64-unknown-windows-sycldevice -Wno-sycl-2017-compat -fsyntax-only \
 // RUN: -aux-triple x86_64-pc-windows-msvc -fsycl -fsycl-is-device \
 // RUN: -fms-extensions -verify  %s
 // check error message when dllimport function gets called in sycl-kernel code
@@ -51,7 +51,7 @@ int foobar()  // expected-warning {{'foobar' redeclared without 'dllimport' attr
 }
 
 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();
 }
 
diff --git a/clang/test/SemaSYCL/sycl-esimd.cpp b/clang/test/SemaSYCL/sycl-esimd.cpp
index f235394b9971..68ef557818dd 100644
--- a/clang/test/SemaSYCL/sycl-esimd.cpp
+++ b/clang/test/SemaSYCL/sycl-esimd.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-explicit-simd -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-explicit-simd -fsyntax-only -Wno-sycl-2017-compat -verify %s
 
 // ----------- Negative tests
 
@@ -12,7 +12,7 @@ bar() {}
 // -- ESIMD kernel can't call functions with required subgroup size != 1
 
 template <typename ID, typename F>
-void kernel0(F f) __attribute__((sycl_kernel)) {
+void kernel0(const F &f) __attribute__((sycl_kernel)) {
   f();
 }
 
@@ -27,7 +27,7 @@ void test0() {
 
 // -- Usual kernel can't call ESIMD function
 template <typename ID, typename F>
-void kernel1(F f) __attribute__((sycl_kernel)) {
+void kernel1(const F &f) __attribute__((sycl_kernel)) {
   f();
 }
 
@@ -43,7 +43,7 @@ void test1() {
 
 // -- Kernel-function call, both have the attribute, lambda kernel.
 template <typename ID, typename F>
-void kernel2(F f) __attribute__((sycl_kernel)) {
+void kernel2(const F &f) __attribute__((sycl_kernel)) {
   f();
 }
 
@@ -64,12 +64,12 @@ class A {
 // --  Functor object kernel.
 
 template <typename F, typename ID = F>
-void kernel3(F f) __attribute__((sycl_kernel)) {
+void kernel3(const F &f) __attribute__((sycl_kernel)) {
   f();
 }
 
 struct Kernel3 {
-  void operator()() __attribute__((sycl_explicit_simd)) {}
+  void operator()() const __attribute__((sycl_explicit_simd)) {}
 };
 
 void bar3() {
diff --git a/clang/test/SemaSYCL/sycl-fptr-lambda.cpp b/clang/test/SemaSYCL/sycl-fptr-lambda.cpp
index ba44c7286ec2..5ec68265a2ee 100644
--- a/clang/test/SemaSYCL/sycl-fptr-lambda.cpp
+++ b/clang/test/SemaSYCL/sycl-fptr-lambda.cpp
@@ -1,8 +1,8 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -std=c++14 -verify -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -std=c++14 -verify -Wno-sycl-2017-compat -fsyntax-only %s
 // expected-no-diagnostics
 
 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();
 }
 
diff --git a/clang/test/SemaSYCL/sycl-pseudo-dtor.cpp b/clang/test/SemaSYCL/sycl-pseudo-dtor.cpp
index b4700cadd3f2..cc839a18446e 100644
--- a/clang/test/SemaSYCL/sycl-pseudo-dtor.cpp
+++ b/clang/test/SemaSYCL/sycl-pseudo-dtor.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -Wno-sycl-2017-compat -fsyntax-only %s
 
 template <typename functor_t>
 struct functor_wrapper{
@@ -13,7 +13,7 @@ struct S { virtual void foo(); };
 struct T { virtual ~T(); };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
   // expected-no-note@+1
   using DATA_I = int;
   using DATA_S = S;
diff --git a/clang/test/SemaSYCL/sycl-restrict.cpp b/clang/test/SemaSYCL/sycl-restrict.cpp
index 00b215d7d459..2c963f95ff76 100644
--- a/clang/test/SemaSYCL/sycl-restrict.cpp
+++ b/clang/test/SemaSYCL/sycl-restrict.cpp
@@ -1,6 +1,6 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -triple spir64 -Wno-return-type -verify -fsyntax-only -std=c++17 %s
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -triple spir64 -fno-sycl-allow-func-ptr -Wno-return-type -verify -fsyntax-only -std=c++17 %s
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -triple spir64 -DALLOW_FP=1 -fsycl-allow-func-ptr -Wno-return-type -verify -fsyntax-only -std=c++17 %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -triple spir64 -Wno-return-type -verify -Wno-sycl-2017-compat -fsyntax-only -std=c++17 %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -triple spir64 -fno-sycl-allow-func-ptr -Wno-return-type -verify -Wno-sycl-2017-compat -fsyntax-only -std=c++17 %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -triple spir64 -DALLOW_FP=1 -fsycl-allow-func-ptr -Wno-return-type -verify -Wno-sycl-2017-compat -fsyntax-only -std=c++17 %s
 
 namespace std {
 class type_info;
@@ -86,7 +86,7 @@ bool isa_B(A *a) {
 }
 
 template <typename N, typename L>
-__attribute__((sycl_kernel)) void kernel1(L l) {
+__attribute__((sycl_kernel)) void kernel1(const L &l) {
   l(); //#rtti_kernel  // expected-note 2{{called by 'kernel1<kernel_name, (lambda at }}
 }
 } // namespace Check_RTTI_Restriction
@@ -374,7 +374,7 @@ int use2(a_type ab, a_type *abp) {
 }
 
 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(); //#call_kernelFunc // expected-note 3{{called by 'kernel_single_task<fake_kernel, (lambda at}}
 }
 
diff --git a/clang/test/SemaSYCL/sycl-varargs-cconv.cpp b/clang/test/SemaSYCL/sycl-varargs-cconv.cpp
index 4f1866fdcacd..8e0f8e0dd535 100644
--- a/clang/test/SemaSYCL/sycl-varargs-cconv.cpp
+++ b/clang/test/SemaSYCL/sycl-varargs-cconv.cpp
@@ -43,7 +43,7 @@ void bar() {
 }
 
 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(); //expected-note 2+ {{called by 'kernel_single_task}}
 }
 
diff --git a/clang/test/SemaSYCL/tls_error.cpp b/clang/test/SemaSYCL/tls_error.cpp
index eb8961610595..bacecf782780 100644
--- a/clang/test/SemaSYCL/tls_error.cpp
+++ b/clang/test/SemaSYCL/tls_error.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64 -verify -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64 -verify -Wno-sycl-2017-compat -fsyntax-only %s
 
 extern __thread void* __once_callable;  // expected-no-error
 extern __thread void (*__once_call)();  // expected-no-error
@@ -14,7 +14,7 @@ void usage() {
 }
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
   //expected-note@+1{{called by}}
   kernelFunc();
 }
diff --git a/clang/test/SemaSYCL/unevaluated-function.cpp b/clang/test/SemaSYCL/unevaluated-function.cpp
index 6e4f3b0fa3f5..49ffe954ab7a 100644
--- a/clang/test/SemaSYCL/unevaluated-function.cpp
+++ b/clang/test/SemaSYCL/unevaluated-function.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -verify -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -verify -Wno-sycl-2017-compat -fsyntax-only %s
 
 // Check that a function used in an unevaluated context is not subject
 // to delayed device diagnostics.
@@ -24,7 +24,7 @@ bool foo3() {
 }
 
 template <typename Name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
   // expected-note@+1 1{{called by}}
   kernelFunc();
 }
diff --git a/clang/test/SemaSYCL/unnamed-kernel.cpp b/clang/test/SemaSYCL/unnamed-kernel.cpp
index 33e5656ac822..45954f325b4d 100644
--- a/clang/test/SemaSYCL/unnamed-kernel.cpp
+++ b/clang/test/SemaSYCL/unnamed-kernel.cpp
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s
-// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -verify %s
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -Wno-sycl-2017-compat -verify %s
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -Wno-sycl-2017-compat -verify %s
 #include <sycl.hpp>
 
 #ifdef __SYCL_UNNAMED_LAMBDA__
diff --git a/clang/test/SemaSYCL/unsupported_math.cpp b/clang/test/SemaSYCL/unsupported_math.cpp
index e93bf6128358..3c0de837dcd7 100644
--- a/clang/test/SemaSYCL/unsupported_math.cpp
+++ b/clang/test/SemaSYCL/unsupported_math.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify %s
 extern "C" float  sinf(float);
 extern "C" float  cosf(float);
 extern "C" float  logf(float);
@@ -6,7 +6,7 @@ extern "C" double sin(double);
 extern "C" double cos(double);
 extern "C" double log(double);
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/clang/test/SemaSYCL/variadic-func-call.cpp b/clang/test/SemaSYCL/variadic-func-call.cpp
index 896ba16e46b9..bdb7ace1690b 100644
--- a/clang/test/SemaSYCL/variadic-func-call.cpp
+++ b/clang/test/SemaSYCL/variadic-func-call.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown -fsyntax-only -Wno-sycl-2017-compat -verify %s
 
 void variadic(int, ...) {}
 namespace NS {
@@ -18,7 +18,7 @@ void foo() {
 void overloaded(int, int) {}
 void overloaded(int, ...) {}
 template <typename, typename Func>
-__attribute__((sycl_kernel)) void task(Func KF) {
+__attribute__((sycl_kernel)) void task(const Func &KF) {
   KF(); // expected-note 2 {{called by 'task}}
 }
 
diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp
index 9f9a3196024c..c6ebc3c48730 100644
--- a/clang/test/SemaSYCL/wrapped-accessor.cpp
+++ b/clang/test/SemaSYCL/wrapped-accessor.cpp
@@ -9,7 +9,7 @@ template <typename Acc>
 struct AccWrapper { Acc accessor; };
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index 83cc08f88cbb..737afb7d0327 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -33,6 +33,10 @@
 #include <memory>
 #include <type_traits>
 
+#if !CL_SYCL_LANGUAGE_VERSION || CL_SYCL_LANGUAGE_VERSION < 2020
+#define __SYCL_NONCONST_FUNCTOR__
+#endif
+
 template <typename DataT, int Dimensions, cl::sycl::access::mode AccessMode,
           cl::sycl::access::target AccessTarget,
           cl::sycl::access::placeholder IsPlaceholder>
@@ -719,14 +723,24 @@ class __SYCL_EXPORT handler {
   // NOTE: the name of this function - "kernel_single_task" - is used by the
   // Front End to determine kernel invocation kind.
   template <typename KernelName, typename KernelType>
-  __attribute__((sycl_kernel)) void kernel_single_task(KernelType KernelFunc) {
+  __attribute__((sycl_kernel)) void
+#ifdef __SYCL_NONCONST_FUNCTOR__
+  kernel_single_task(KernelType KernelFunc) {
+#else
+  kernel_single_task(const KernelType &KernelFunc) {
+#endif
     KernelFunc();
   }
 
   // NOTE: the name of these functions - "kernel_parallel_for" - are used by the
   // Front End to determine kernel invocation kind.
   template <typename KernelName, typename ElementType, typename KernelType>
-  __attribute__((sycl_kernel)) void kernel_parallel_for(KernelType KernelFunc) {
+  __attribute__((sycl_kernel)) void
+#ifdef __SYCL_NONCONST_FUNCTOR__
+  kernel_parallel_for(KernelType KernelFunc) {
+#else
+  kernel_parallel_for(const KernelType &KernelFunc) {
+#endif
     KernelFunc(
         detail::Builder::getElement(static_cast<ElementType *>(nullptr)));
   }
@@ -735,7 +749,11 @@ class __SYCL_EXPORT handler {
   // used by the Front End to determine kernel invocation kind.
   template <typename KernelName, typename ElementType, typename KernelType>
   __attribute__((sycl_kernel)) void
+#ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_parallel_for_work_group(KernelType KernelFunc) {
+#else
+  kernel_parallel_for_work_group(const KernelType &KernelFunc) {
+#endif
     KernelFunc(
         detail::Builder::getElement(static_cast<ElementType *>(nullptr)));
   }
@@ -839,7 +857,11 @@ class __SYCL_EXPORT handler {
   ///
   /// \param KernelFunc is a SYCL kernel function.
   template <typename KernelName = detail::auto_name, typename KernelType>
+#ifdef __SYCL_NONCONST_FUNCTOR__
   void single_task(KernelType KernelFunc) {
+#else
+  void single_task(const KernelType &KernelFunc) {
+#endif
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -856,17 +878,29 @@ class __SYCL_EXPORT handler {
   }
 
   template <typename KernelName = detail::auto_name, typename KernelType>
+#ifdef __SYCL_NONCONST_FUNCTOR__
   void parallel_for(range<1> NumWorkItems, KernelType KernelFunc) {
+#else
+  void parallel_for(range<1> NumWorkItems, const KernelType &KernelFunc) {
+#endif
     parallel_for_lambda_impl<KernelName>(NumWorkItems, std::move(KernelFunc));
   }
 
   template <typename KernelName = detail::auto_name, typename KernelType>
+#ifdef __SYCL_NONCONST_FUNCTOR__
   void parallel_for(range<2> NumWorkItems, KernelType KernelFunc) {
+#else
+  void parallel_for(range<2> NumWorkItems, const KernelType &KernelFunc) {
+#endif
     parallel_for_lambda_impl<KernelName>(NumWorkItems, std::move(KernelFunc));
   }
 
   template <typename KernelName = detail::auto_name, typename KernelType>
+#ifdef __SYCL_NONCONST_FUNCTOR__
   void parallel_for(range<3> NumWorkItems, KernelType KernelFunc) {
+#else
+  void parallel_for(range<3> NumWorkItems, const KernelType &KernelFunc) {
+#endif
     parallel_for_lambda_impl<KernelName>(NumWorkItems, std::move(KernelFunc));
   }
 
@@ -929,7 +963,11 @@ class __SYCL_EXPORT handler {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   void parallel_for(range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
+#ifdef __SYCL_NONCONST_FUNCTOR__
                     KernelType KernelFunc) {
+#else
+                    const KernelType &KernelFunc) {
+#endif
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -961,7 +999,12 @@ class __SYCL_EXPORT handler {
   /// \param KernelFunc is a SYCL kernel function.
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
-  void parallel_for(nd_range<Dims> ExecutionRange, KernelType KernelFunc) {
+  void parallel_for(nd_range<Dims> ExecutionRange,
+#ifdef __SYCL_NONCONST_FUNCTOR__
+                    KernelType KernelFunc) {
+#else
+                    const KernelType &KernelFunc) {
+#endif
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -988,7 +1031,12 @@ class __SYCL_EXPORT handler {
             int Dims, typename Reduction>
   detail::enable_if_t<Reduction::accessor_mode == access::mode::read_write &&
                       Reduction::has_fast_atomics && !Reduction::is_usm>
-  parallel_for(nd_range<Dims> Range, Reduction Redu, KernelType KernelFunc) {
+  parallel_for(nd_range<Dims> Range, Reduction Redu,
+#ifdef __SYCL_NONCONST_FUNCTOR__
+               KernelType KernelFunc) {
+#else
+               const KernelType &KernelFunc) {
+#endif
     intel::detail::reduCGFunc<KernelName>(*this, KernelFunc, Range, Redu,
                                           Redu.getUserAccessor());
   }
@@ -1001,7 +1049,12 @@ class __SYCL_EXPORT handler {
             int Dims, typename Reduction>
   detail::enable_if_t<Reduction::accessor_mode == access::mode::read_write &&
                       Reduction::has_fast_atomics && Reduction::is_usm>
-  parallel_for(nd_range<Dims> Range, Reduction Redu, KernelType KernelFunc) {
+  parallel_for(nd_range<Dims> Range, Reduction Redu,
+#ifdef __SYCL_NONCONST_FUNCTOR__
+               KernelType KernelFunc) {
+#else
+               const KernelType &KernelFunc) {
+#endif
     intel::detail::reduCGFunc<KernelName>(*this, KernelFunc, Range, Redu,
                                           Redu.getUSMPointer());
   }
@@ -1020,7 +1073,12 @@ class __SYCL_EXPORT handler {
             int Dims, typename Reduction>
   detail::enable_if_t<Reduction::accessor_mode == access::mode::discard_write &&
                       Reduction::has_fast_atomics>
-  parallel_for(nd_range<Dims> Range, Reduction Redu, KernelType KernelFunc) {
+  parallel_for(nd_range<Dims> Range, Reduction Redu,
+#ifdef __SYCL_NONCONST_FUNCTOR__
+               KernelType KernelFunc) {
+#else
+               const KernelType &KernelFunc) {
+#endif
     shared_ptr_class<detail::queue_impl> QueueCopy = MQueue;
     auto RWAcc = Redu.getReadWriteScalarAcc(*this);
     intel::detail::reduCGFunc<KernelName>(*this, KernelFunc, Range, Redu,
@@ -1056,7 +1114,12 @@ class __SYCL_EXPORT handler {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims, typename Reduction>
   detail::enable_if_t<!Reduction::has_fast_atomics>
-  parallel_for(nd_range<Dims> Range, Reduction Redu, KernelType KernelFunc) {
+  parallel_for(nd_range<Dims> Range, Reduction Redu,
+#ifdef __SYCL_NONCONST_FUNCTOR__
+               KernelType KernelFunc) {
+#else
+               const KernelType &KernelFunc) {
+#endif
     // This parallel_for() is lowered to the following sequence:
     // 1) Call a kernel that a) call user's lambda function and b) performs
     //    one iteration of reduction, storing the partial reductions/sums
@@ -1131,7 +1194,11 @@ class __SYCL_EXPORT handler {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   void parallel_for_work_group(range<Dims> NumWorkGroups,
+#ifdef __SYCL_NONCONST_FUNCTOR__
                                KernelType KernelFunc) {
+#else
+                               const KernelType &KernelFunc) {
+#endif
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -1164,7 +1231,11 @@ class __SYCL_EXPORT handler {
             int Dims>
   void parallel_for_work_group(range<Dims> NumWorkGroups,
                                range<Dims> WorkGroupSize,
+#ifdef __SYCL_NONCONST_FUNCTOR__
                                KernelType KernelFunc) {
+#else
+                               const KernelType &KernelFunc) {
+#endif
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -1266,7 +1337,12 @@ class __SYCL_EXPORT handler {
   /// \param KernelFunc is a lambda that is used if device, queue is bound to,
   /// is a host device.
   template <typename KernelName = detail::auto_name, typename KernelType>
-  void single_task(kernel Kernel, KernelType KernelFunc) {
+  void single_task(kernel Kernel,
+#ifdef __SYCL_NONCONST_FUNCTOR__
+                   KernelType KernelFunc) {
+#else
+                   const KernelType &KernelFunc) {
+#endif
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -1306,7 +1382,11 @@ class __SYCL_EXPORT handler {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   void parallel_for(kernel Kernel, range<Dims> NumWorkItems,
+#ifdef __SYCL_NONCONST_FUNCTOR__
                     KernelType KernelFunc) {
+#else
+                    const KernelType &KernelFunc) {
+#endif
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -1341,7 +1421,12 @@ class __SYCL_EXPORT handler {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   void parallel_for(kernel Kernel, range<Dims> NumWorkItems,
-                    id<Dims> WorkItemOffset, KernelType KernelFunc) {
+                    id<Dims> WorkItemOffset,
+#ifdef __SYCL_NONCONST_FUNCTOR__
+                    KernelType KernelFunc) {
+#else
+                    const KernelType &KernelFunc) {
+#endif
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -1378,7 +1463,11 @@ class __SYCL_EXPORT handler {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   void parallel_for(kernel Kernel, nd_range<Dims> NDRange,
+#ifdef __SYCL_NONCONST_FUNCTOR__
                     KernelType KernelFunc) {
+#else
+                    const KernelType &KernelFunc) {
+#endif
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -1420,7 +1509,11 @@ class __SYCL_EXPORT handler {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   void parallel_for_work_group(kernel Kernel, range<Dims> NumWorkGroups,
+#ifdef __SYCL_NONCONST_FUNCTOR__
                                KernelType KernelFunc) {
+#else
+                               const KernelType &KernelFunc) {
+#endif
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -1458,7 +1551,11 @@ class __SYCL_EXPORT handler {
             int Dims>
   void parallel_for_work_group(kernel Kernel, range<Dims> NumWorkGroups,
                                range<Dims> WorkGroupSize,
+#ifdef __SYCL_NONCONST_FUNCTOR__
                                KernelType KernelFunc) {
+#else
+                               const KernelType &KernelFunc) {
+#endif
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp
index 5338567f5eda..09c7da71a3dc 100644
--- a/sycl/include/CL/sycl/queue.hpp
+++ b/sycl/include/CL/sycl/queue.hpp
@@ -361,21 +361,47 @@ class __SYCL_EXPORT queue {
     return submit([=](handler &CGH) { CGH.prefetch(Ptr, Count); });
   }
 
+// having _TWO_ mid-param #ifdefs makes the functions very difficult to read.
+// Here we simplify the &CodeLoc declaration to be CODELOCPARAM(A,B,C,D,E) and
+// CODELOCARG(B,C,D) Similarly, the KernelFunc param is simplified to be
+// KERNELFUNCPARAM(B,F,G) Once the queue kernel functions are defined, these
+// macros are #undef immediately. Compare this first definition of single_task
+// here, with the second one below. replace CODELOCPARAM(A,B,C,D,E) with nothing
+// or :   , const detail::code_location &CodeLoc =
+// detail::code_location::current() replace CODELOCARG(B,C,D) with nothign or :
+// const detail::code_location &CodeLoc = {};
+#define A ,
+#define B const
+#define C detail::code_location
+#define D &CodeLoc
+#define E detail::code_location::current()
+#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
+#define CODELOCPARAM(a, b, c, d, e) a b c d = e
+
+#define CODELOCARG(b, c, d)
+#else
+#define CODELOCPARAM(a, b, c, d, e)
+
+#define CODELOCARG(b, c, d) b c d = {}
+#endif
+// replace KERNELFUNCPARAM(B,F,G) with   KernelType KernelFunc
+//                                  or   const KernelType &KernelFunc
+#define F KernelType
+#define G KernelFunc
+#ifdef __SYCL_NONCONST_FUNCTOR__
+#define KERNELFUNCPARAM(B, F, G) F G
+#else
+#define KERNELFUNCPARAM(B, F, G) B F &G
+#endif
+
   /// single_task version with a kernel represented as a lambda.
   ///
   /// \param KernelFunc is the Kernel functor or lambda
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
-  event single_task(
-      KernelType KernelFunc
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event single_task(KERNELFUNCPARAM(B, F, G) CODELOCPARAM(A, B, C, D, E)) {
+    CODELOCARG(B, C, D);
+
     return submit(
         [&](handler &CGH) {
           CGH.template single_task<KernelName, KernelType>(KernelFunc);
@@ -383,6 +409,18 @@ class __SYCL_EXPORT queue {
         CodeLoc);
   }
 
+// Clean up CODELOC and KERNELFUNC macros.
+#undef A
+#undef B
+#undef C
+#undef D
+#undef E
+#undef F
+#undef G
+#undef CODELOCPARAM
+#undef CODELOCARG
+#undef KERNELFUNCPARAM
+
   /// single_task version with a kernel represented as a lambda.
   ///
   /// \param DepEvent is an event that specifies the kernel dependencies
@@ -390,7 +428,13 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
   event single_task(
-      event DepEvent, KernelType KernelFunc
+      event DepEvent,
+#ifdef __SYCL_NONCONST_FUNCTOR__
+      KernelType KernelFunc
+#else
+      const KernelType &KernelFunc
+#endif
+
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
       ,
       const detail::code_location &CodeLoc = detail::code_location::current()
@@ -415,7 +459,13 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
   event single_task(
-      const vector_class<event> &DepEvents, KernelType KernelFunc
+      const vector_class<event> &DepEvents,
+#ifdef __SYCL_NONCONST_FUNCTOR__
+      KernelType KernelFunc
+#else
+      const KernelType &KernelFunc
+#endif
+
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
       ,
       const detail::code_location &CodeLoc = detail::code_location::current()
@@ -440,7 +490,13 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(
-      range<1> NumWorkItems, KernelType KernelFunc
+      range<1> NumWorkItems,
+#ifdef __SYCL_NONCONST_FUNCTOR__
+      KernelType KernelFunc
+#else
+      const KernelType &KernelFunc
+#endif
+
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
       ,
       const detail::code_location &CodeLoc = detail::code_location::current()
@@ -460,7 +516,13 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(
-      range<2> NumWorkItems, KernelType KernelFunc
+      range<2> NumWorkItems,
+#ifdef __SYCL_NONCONST_FUNCTOR__
+      KernelType KernelFunc
+#else
+      const KernelType &KernelFunc
+#endif
+
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
       ,
       const detail::code_location &CodeLoc = detail::code_location::current()
@@ -480,7 +542,13 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(
-      range<3> NumWorkItems, KernelType KernelFunc
+      range<3> NumWorkItems,
+#ifdef __SYCL_NONCONST_FUNCTOR__
+      KernelType KernelFunc
+#else
+      const KernelType &KernelFunc
+#endif
+
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
       ,
       const detail::code_location &CodeLoc = detail::code_location::current()
@@ -501,7 +569,13 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(
-      range<1> NumWorkItems, event DepEvent, KernelType KernelFunc
+      range<1> NumWorkItems, event DepEvent,
+#ifdef __SYCL_NONCONST_FUNCTOR__
+      KernelType KernelFunc
+#else
+      const KernelType &KernelFunc
+#endif
+
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
       ,
       const detail::code_location &CodeLoc = detail::code_location::current()
@@ -523,7 +597,13 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(
-      range<2> NumWorkItems, event DepEvent, KernelType KernelFunc
+      range<2> NumWorkItems, event DepEvent,
+#ifdef __SYCL_NONCONST_FUNCTOR__
+      KernelType KernelFunc
+#else
+      const KernelType &KernelFunc
+#endif
+
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
       ,
       const detail::code_location &CodeLoc = detail::code_location::current()
@@ -545,7 +625,13 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(
-      range<3> NumWorkItems, event DepEvent, KernelType KernelFunc
+      range<3> NumWorkItems, event DepEvent,
+#ifdef __SYCL_NONCONST_FUNCTOR__
+      KernelType KernelFunc
+#else
+      const KernelType &KernelFunc
+#endif
+
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
       ,
       const detail::code_location &CodeLoc = detail::code_location::current()
@@ -569,7 +655,7 @@ class __SYCL_EXPORT queue {
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(
       range<1> NumWorkItems, const vector_class<event> &DepEvents,
-      KernelType KernelFunc
+      const KernelType &KernelFunc
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
       ,
       const detail::code_location &CodeLoc = detail::code_location::current()
@@ -593,7 +679,7 @@ class __SYCL_EXPORT queue {
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(
       range<2> NumWorkItems, const vector_class<event> &DepEvents,
-      KernelType KernelFunc
+      const KernelType &KernelFunc
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
       ,
       const detail::code_location &CodeLoc = detail::code_location::current()
@@ -617,7 +703,7 @@ class __SYCL_EXPORT queue {
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(
       range<3> NumWorkItems, const vector_class<event> &DepEvents,
-      KernelType KernelFunc
+      const KernelType &KernelFunc
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
       ,
       const detail::code_location &CodeLoc = detail::code_location::current()
@@ -640,7 +726,8 @@ class __SYCL_EXPORT queue {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   event parallel_for(
-      range<Dims> NumWorkItems, id<Dims> WorkItemOffset, KernelType KernelFunc
+      range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
+      const KernelType &KernelFunc
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
       ,
       const detail::code_location &CodeLoc = detail::code_location::current()
@@ -669,7 +756,7 @@ class __SYCL_EXPORT queue {
             int Dims>
   event parallel_for(
       range<Dims> NumWorkItems, id<Dims> WorkItemOffset, event DepEvent,
-      KernelType KernelFunc
+      const KernelType &KernelFunc
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
       ,
       const detail::code_location &CodeLoc = detail::code_location::current()
@@ -700,7 +787,13 @@ class __SYCL_EXPORT queue {
             int Dims>
   event parallel_for(
       range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
-      const vector_class<event> &DepEvents, KernelType KernelFunc
+      const vector_class<event> &DepEvents,
+#ifdef __SYCL_NONCONST_FUNCTOR__
+      KernelType KernelFunc
+#else
+      const KernelType &KernelFunc
+#endif
+
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
       ,
       const detail::code_location &CodeLoc = detail::code_location::current()
@@ -728,7 +821,13 @@ class __SYCL_EXPORT queue {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   event parallel_for(
-      nd_range<Dims> ExecutionRange, KernelType KernelFunc
+      nd_range<Dims> ExecutionRange,
+#ifdef __SYCL_NONCONST_FUNCTOR__
+      KernelType KernelFunc
+#else
+      const KernelType &KernelFunc
+#endif
+
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
       ,
       const detail::code_location &CodeLoc = detail::code_location::current()
@@ -756,7 +855,8 @@ class __SYCL_EXPORT queue {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   event parallel_for(
-      nd_range<Dims> ExecutionRange, event DepEvent, KernelType KernelFunc
+      nd_range<Dims> ExecutionRange, event DepEvent,
+      const KernelType &KernelFunc
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
       ,
       const detail::code_location &CodeLoc = detail::code_location::current()
@@ -787,7 +887,7 @@ class __SYCL_EXPORT queue {
             int Dims>
   event parallel_for(
       nd_range<Dims> ExecutionRange, const vector_class<event> &DepEvents,
-      KernelType KernelFunc
+      const KernelType &KernelFunc
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
       ,
       const detail::code_location &CodeLoc = detail::code_location::current()
diff --git a/sycl/test/functor/functor_inheritance.cpp b/sycl/test/functor/functor_inheritance.cpp
index 3a3d5218ef6e..436e8ce74a7d 100644
--- a/sycl/test/functor/functor_inheritance.cpp
+++ b/sycl/test/functor/functor_inheritance.cpp
@@ -36,7 +36,7 @@ struct Derived : public Base, public SecondBase {
       int _A, int _B, int _C, int _D, int _E,
       cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> &_Acc)
       : A(_A), Acc(_Acc), /*Out(_Out),*/ Base(_B, _C, _D), SecondBase(_E) {}
-  void operator()() {
+  void operator()() const {
     Acc[0] = this->A + this->B + this->InnerObj.C + this->InnerObj.D + this->E;
   }
 
diff --git a/sycl/test/functor/kernel_functor.cpp b/sycl/test/functor/kernel_functor.cpp
index 9dd5e0f2fecd..9bed0b451ca2 100644
--- a/sycl/test/functor/kernel_functor.cpp
+++ b/sycl/test/functor/kernel_functor.cpp
@@ -23,7 +23,7 @@ constexpr auto sycl_global_buffer = cl::sycl::access::target::global_buffer;
 // - functor class is defined in an anonymous namespace
 // - the '()' operator:
 //   * does not have parameters (to be used in 'single_task').
-//   * has no 'const' qualifier
+//   * has the 'const' qualifier
 namespace {
 class Functor1 {
 public:
@@ -32,7 +32,7 @@ class Functor1 {
       cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> &Acc_)
       : X(X_), Acc(Acc_) {}
 
-  void operator()() { Acc[0] += X; }
+  void operator()() const { Acc[0] += X; }
 
 private:
   int X;
@@ -66,14 +66,14 @@ class Functor2 {
 // - functor class is templated and defined in the translation unit scope
 // - the '()' operator:
 //   * has a parameter of type cl::sycl::id<1> (to be used in 'parallel_for').
-//   * has no 'const' qualifier
+//   * has the 'const' qualifier
 template <typename T> class TmplFunctor {
 public:
   TmplFunctor(
       T X_, cl::sycl::accessor<T, 1, sycl_read_write, sycl_global_buffer> &Acc_)
       : X(X_), Acc(Acc_) {}
 
-  void operator()(cl::sycl::id<1> id) { Acc[id] += X; }
+  void operator()(cl::sycl::id<1> id) const { Acc[id] += X; }
 
 private:
   T X;
diff --git a/sycl/test/hier_par/hier_par_basic.cpp b/sycl/test/hier_par/hier_par_basic.cpp
index d1a94ea1a711..5670663e2e14 100644
--- a/sycl/test/hier_par/hier_par_basic.cpp
+++ b/sycl/test/hier_par/hier_par_basic.cpp
@@ -55,7 +55,7 @@ struct PFWIFunctor {
       : wg_chunk(wg_chunk), wg_size(wg_size), wg_offset(wg_offset),
         range_length(range_length), v(v), dev_ptr(dev_ptr) {}
 
-  void operator()(h_item<1> i) {
+  void operator()(h_item<1> i) const {
     // number of buf elements per work item:
     size_t wi_chunk = (wg_chunk + wg_size - 1) / wg_size;
     auto id = i.get_physical_local_id().get(0);
@@ -82,7 +82,7 @@ struct PFWGFunctor {
       : wg_chunk(wg_chunk), range_length(range_length), dev_ptr(dev_ptr),
         addend(addend), n_iter(n_iter) {}
 
-  void operator()(group<1> g) {
+  void operator()(group<1> g) const {
     int v = addend; // to check constant initializer works too
     size_t wg_offset = wg_chunk * g.get_id(0);
     size_t wg_size = g.get_local_range(0);
@@ -95,13 +95,13 @@ struct PFWGFunctor {
   }
   // Dummy operator '()' to make sure compiler can handle multiple '()'
   // operators/ and pick the right one for PFWG kernel code generation.
-  void operator()(int ind, int val) { dev_ptr[ind] += val; }
+  void operator()(int ind, int val) const { dev_ptr[ind] += val; }
 
   const size_t wg_chunk;
   const size_t range_length;
   const int n_iter;
   const int addend;
-  AccTy dev_ptr;
+  mutable AccTy dev_ptr;
 };
 
 int main() {
diff --git a/sycl/test/separate-compile/same-kernel.cpp b/sycl/test/separate-compile/same-kernel.cpp
index 66ca32780f3c..ad72decfa360 100644
--- a/sycl/test/separate-compile/same-kernel.cpp
+++ b/sycl/test/separate-compile/same-kernel.cpp
@@ -26,9 +26,7 @@ class TestFnObj {
   TestFnObj(buffer<int> &buf, handler &cgh) :
     data(buf.get_access<access::mode::write>(cgh)) {}
   accessor<int, 1, access::mode::write, access::target::global_buffer> data;
-  void operator()(id<1> item) {
-    data[item] = item[0];
-  }
+  void operator()(id<1> item) const { data[item] = item[0]; }
 };
 
 void kernel2();
diff --git a/sycl/test/sub_group/attributes.cpp b/sycl/test/sub_group/attributes.cpp
index d8173d2d1cf7..650ba23dd8c7 100644
--- a/sycl/test/sub_group/attributes.cpp
+++ b/sycl/test/sub_group/attributes.cpp
@@ -22,7 +22,7 @@
   class KernelFunctor##SIZE {                                                  \
   public:                                                                      \
     [[cl::intel_reqd_sub_group_size(SIZE)]] void                               \
-    operator()(cl::sycl::nd_item<1> Item) {                                    \
+    operator()(cl::sycl::nd_item<1> Item) const {                              \
       const auto GID = Item.get_global_id();                                   \
     }                                                                          \
   };

From d68ba917b35e7124272ab47c256a640bd26ada35 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Tue, 4 Aug 2020 18:24:41 -0700
Subject: [PATCH 02/17] hooked up -sycl-std flag to set SYCL version constant. 
 Needs test.

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/lib/Driver/ToolChains/Clang.cpp | 11 +++++++++++
 sycl/include/CL/sycl/queue.hpp        |  9 ++++++---
 2 files changed, 17 insertions(+), 3 deletions(-)

diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 61ee698ffbd9..283a458db2fc 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -4160,9 +4160,20 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
     if (SYCLStdArg) {
       SYCLStdArg->render(Args, CmdArgs);
       CmdArgs.push_back("-fsycl-std-layout-kernel-params");
+
+      unsigned version = llvm::StringSwitch<unsigned>(SYCLStdArg->getValue())
+                             .Cases("2017", "1.2.1", "121", "sycl-1.2.1", 2017)
+                             .Case("2020", 2020)
+                             .Default(0);
+      if (version == 2017)
+        CmdArgs.push_back("-DCL_SYCL_LANGUAGE_VERSION=121");
+      else if (version == 2020)
+        CmdArgs.push_back("-DCL_SYCL_LANGUAGE_VERSION=2020");
+
     } else {
       // Ensure the default version in SYCL mode is 1.2.1 (aka 2017)
       CmdArgs.push_back("-sycl-std=2017");
+      CmdArgs.push_back("-DCL_SYCL_LANGUAGE_VERSION=121");
     }
   }
 
diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp
index 09c7da71a3dc..bdd0823cd352 100644
--- a/sycl/include/CL/sycl/queue.hpp
+++ b/sycl/include/CL/sycl/queue.hpp
@@ -366,10 +366,13 @@ class __SYCL_EXPORT queue {
 // CODELOCARG(B,C,D) Similarly, the KernelFunc param is simplified to be
 // KERNELFUNCPARAM(B,F,G) Once the queue kernel functions are defined, these
 // macros are #undef immediately. Compare this first definition of single_task
-// here, with the second one below. replace CODELOCPARAM(A,B,C,D,E) with nothing
+// here, with the second one below.
+
+// replace CODELOCPARAM(A,B,C,D,E) with nothing
 // or :   , const detail::code_location &CodeLoc =
-// detail::code_location::current() replace CODELOCARG(B,C,D) with nothign or :
-// const detail::code_location &CodeLoc = {};
+// detail::code_location::current()
+// replace CODELOCARG(B,C,D) with nothing
+// or :  const detail::code_location &CodeLoc = {}
 #define A ,
 #define B const
 #define C detail::code_location

From 8390646aaee35f4eca8300a606a83ab11df03d29 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Tue, 4 Aug 2020 20:23:17 -0700
Subject: [PATCH 03/17] simplified macros - tricky re-use not worth it.

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 sycl/include/CL/sycl/queue.hpp | 60 ++++++++++++++--------------------
 1 file changed, 24 insertions(+), 36 deletions(-)

diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp
index bdd0823cd352..880e6f5bf652 100644
--- a/sycl/include/CL/sycl/queue.hpp
+++ b/sycl/include/CL/sycl/queue.hpp
@@ -361,40 +361,35 @@ class __SYCL_EXPORT queue {
     return submit([=](handler &CGH) { CGH.prefetch(Ptr, Count); });
   }
 
-// having _TWO_ mid-param #ifdefs makes the functions very difficult to read.
-// Here we simplify the &CodeLoc declaration to be CODELOCPARAM(A,B,C,D,E) and
-// CODELOCARG(B,C,D) Similarly, the KernelFunc param is simplified to be
-// KERNELFUNCPARAM(B,F,G) Once the queue kernel functions are defined, these
-// macros are #undef immediately. Compare this first definition of single_task
-// here, with the second one below.
-
-// replace CODELOCPARAM(A,B,C,D,E) with nothing
-// or :   , const detail::code_location &CodeLoc =
-// detail::code_location::current()
-// replace CODELOCARG(B,C,D) with nothing
-// or :  const detail::code_location &CodeLoc = {}
-#define A ,
-#define B const
-#define C detail::code_location
-#define D &CodeLoc
-#define E detail::code_location::current()
+  // having _TWO_ mid-param #ifdefs makes the functions very difficult to read.
+  // Here we simplify the &CodeLoc declaration to be CODELOCPARAM(&CodeLoc) and
+  // CODELOCARG(&CodeLoc) Similarly, the KernelFunc param is simplified to be
+  // KERNELFUNCPARAM(KernelFunc) Once the queue kernel functions are defined,
+  // these macros are #undef immediately. Compare this first definition of
+  // single_task here, with the second one below.
+
+  // replace CODELOCPARAM(&CodeLoc) with nothing
+  // or :   , const detail::code_location &CodeLoc =
+  // detail::code_location::current()
+  // replace CODELOCARG(&CodeLoc) with nothing
+  // or :  const detail::code_location &CodeLoc = {}
+
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-#define CODELOCPARAM(a, b, c, d, e) a b c d = e
+#define CODELOCPARAM(a)                                                        \
+  , const detail::code_location a = detail::code_location::current()
 
-#define CODELOCARG(b, c, d)
+#define CODELOCARG(a)
 #else
-#define CODELOCPARAM(a, b, c, d, e)
+#define CODELOCPARAM(a)
 
-#define CODELOCARG(b, c, d) b c d = {}
+#define CODELOCARG(a) const detail::code_location a = {}
 #endif
-// replace KERNELFUNCPARAM(B,F,G) with   KernelType KernelFunc
-//                                  or   const KernelType &KernelFunc
-#define F KernelType
-#define G KernelFunc
+// replace KERNELFUNCPARAM(KernelFunc) with   KernelType KernelFunc
+//                                     or     const KernelType &KernelFunc
 #ifdef __SYCL_NONCONST_FUNCTOR__
-#define KERNELFUNCPARAM(B, F, G) F G
+#define KERNELFUNCPARAM(a) KernelType a
 #else
-#define KERNELFUNCPARAM(B, F, G) B F &G
+#define KERNELFUNCPARAM(a) const KernelType &a
 #endif
 
   /// single_task version with a kernel represented as a lambda.
@@ -402,8 +397,8 @@ class __SYCL_EXPORT queue {
   /// \param KernelFunc is the Kernel functor or lambda
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
-  event single_task(KERNELFUNCPARAM(B, F, G) CODELOCPARAM(A, B, C, D, E)) {
-    CODELOCARG(B, C, D);
+  event single_task(KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
 
     return submit(
         [&](handler &CGH) {
@@ -413,13 +408,6 @@ class __SYCL_EXPORT queue {
   }
 
 // Clean up CODELOC and KERNELFUNC macros.
-#undef A
-#undef B
-#undef C
-#undef D
-#undef E
-#undef F
-#undef G
 #undef CODELOCPARAM
 #undef CODELOCARG
 #undef KERNELFUNCPARAM

From f5df1d581d9ea9c2d6a00f7e71939aba188086d5 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Tue, 4 Aug 2020 20:53:33 -0700
Subject: [PATCH 04/17] test added

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/test/CodeGenSYCL/sycl-std-flag.cpp | 26 ++++++++++++++++++++++++
 1 file changed, 26 insertions(+)
 create mode 100644 clang/test/CodeGenSYCL/sycl-std-flag.cpp

diff --git a/clang/test/CodeGenSYCL/sycl-std-flag.cpp b/clang/test/CodeGenSYCL/sycl-std-flag.cpp
new file mode 100644
index 000000000000..0411c1922785
--- /dev/null
+++ b/clang/test/CodeGenSYCL/sycl-std-flag.cpp
@@ -0,0 +1,26 @@
+// RUN: %clang++  -fsycl  -fsyntax-only -sycl-std=2017 -DSYCL2017 %s
+// RUN: %clang++  -fsycl  -fsyntax-only -sycl-std=121 -DSYCL2017 %s
+// RUN: %clang++  -fsycl  -fsyntax-only -sycl-std=1.2.1 -DSYCL2017 %s
+// RUN: %clang++  -fsycl  -fsyntax-only -sycl-std=sycl-1.2.1 -DSYCL2017 %s
+// RUN: %clang++  -fsycl  -fsyntax-only -sycl-std=2020 -DSYCL2020 %s
+
+
+#if defined(SYCL2017)
+    #if !CL_SYCL_LANGUAGE_VERSION
+        static_assert(false , "CL_SYCL_LANGUAGE_VERSION should be defined when -sycl-std flag is used");
+    #endif
+    #if  CL_SYCL_LANGUAGE_VERSION > 121
+        static_assert(false, "-sycl-std flag should have set CL_SYCL_LANGUAGE_VERSION to 121");
+    #endif
+#endif
+
+#if defined(SYCL2020)
+    #if !CL_SYCL_LANGUAGE_VERSION
+        static_assert(false , "CL_SYCL_LANGUAGE_VERSION should be defined when -sycl-std flag is used");
+    #endif
+    #if  CL_SYCL_LANGUAGE_VERSION < 2020
+        static_assert(false , "-sycl-std=2020 flag should set CL_SYCL_LANGUAGE_VERSION to 2020");
+    #endif
+#endif
+
+// expected-no-diagnostics
\ No newline at end of file

From 02dcc87d7454b990ff8693ed5de6f577e0cea932 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Wed, 5 Aug 2020 10:08:54 -0700
Subject: [PATCH 05/17] last changes, clang-formattery.

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/test/CodeGenSYCL/sycl-std-flag.cpp      |  35 +-
 .../SemaSYCL/intel-max-global-work-dim.cpp    |   2 +-
 sycl/include/CL/sycl/queue.hpp                | 307 ++++--------------
 3 files changed, 82 insertions(+), 262 deletions(-)

diff --git a/clang/test/CodeGenSYCL/sycl-std-flag.cpp b/clang/test/CodeGenSYCL/sycl-std-flag.cpp
index 0411c1922785..e2b4c352127e 100644
--- a/clang/test/CodeGenSYCL/sycl-std-flag.cpp
+++ b/clang/test/CodeGenSYCL/sycl-std-flag.cpp
@@ -1,26 +1,27 @@
-// RUN: %clang++  -fsycl  -fsyntax-only -sycl-std=2017 -DSYCL2017 %s
-// RUN: %clang++  -fsycl  -fsyntax-only -sycl-std=121 -DSYCL2017 %s
-// RUN: %clang++  -fsycl  -fsyntax-only -sycl-std=1.2.1 -DSYCL2017 %s
-// RUN: %clang++  -fsycl  -fsyntax-only -sycl-std=sycl-1.2.1 -DSYCL2017 %s
-// RUN: %clang++  -fsycl  -fsyntax-only -sycl-std=2020 -DSYCL2020 %s
+// REQUIRES: clang-driver
 
+// RUN: %clangxx  -fsycl  -fsyntax-only -sycl-std=2017 -DSYCL2017 %s
+// RUN: %clangxx  -fsycl  -fsyntax-only -sycl-std=121 -DSYCL2017 %s
+// RUN: %clangxx  -fsycl  -fsyntax-only -sycl-std=1.2.1 -DSYCL2017 %s
+// RUN: %clangxx  -fsycl  -fsyntax-only -sycl-std=sycl-1.2.1 -DSYCL2017 %s
+// RUN: %clangxx  -fsycl  -fsyntax-only -sycl-std=2020 -DSYCL2020 %s
 
 #if defined(SYCL2017)
-    #if !CL_SYCL_LANGUAGE_VERSION
-        static_assert(false , "CL_SYCL_LANGUAGE_VERSION should be defined when -sycl-std flag is used");
-    #endif
-    #if  CL_SYCL_LANGUAGE_VERSION > 121
-        static_assert(false, "-sycl-std flag should have set CL_SYCL_LANGUAGE_VERSION to 121");
-    #endif
+#if !CL_SYCL_LANGUAGE_VERSION
+static_assert(false, "CL_SYCL_LANGUAGE_VERSION should be defined when -sycl-std flag is used");
+#endif
+#if CL_SYCL_LANGUAGE_VERSION > 121
+static_assert(false, "-sycl-std flag should have set CL_SYCL_LANGUAGE_VERSION to 121");
+#endif
 #endif
 
 #if defined(SYCL2020)
-    #if !CL_SYCL_LANGUAGE_VERSION
-        static_assert(false , "CL_SYCL_LANGUAGE_VERSION should be defined when -sycl-std flag is used");
-    #endif
-    #if  CL_SYCL_LANGUAGE_VERSION < 2020
-        static_assert(false , "-sycl-std=2020 flag should set CL_SYCL_LANGUAGE_VERSION to 2020");
-    #endif
+#if !CL_SYCL_LANGUAGE_VERSION
+static_assert(false, "CL_SYCL_LANGUAGE_VERSION should be defined when -sycl-std flag is used");
+#endif
+#if CL_SYCL_LANGUAGE_VERSION < 2020
+static_assert(false, "-sycl-std=2020 flag should set CL_SYCL_LANGUAGE_VERSION to 2020");
+#endif
 #endif
 
 // expected-no-diagnostics
\ No newline at end of file
diff --git a/clang/test/SemaSYCL/intel-max-global-work-dim.cpp b/clang/test/SemaSYCL/intel-max-global-work-dim.cpp
index cfde3fd4d7c3..c039787b6d8c 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim.cpp
@@ -45,7 +45,7 @@ struct TRIFuncObjGood2 {
 struct TRIFuncObjBad {
   [[intelfpga::max_global_work_dim(0)]]
   [[intelfpga::max_work_group_size(8, 8, 8)]] // expected-error{{'max_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
-  [[cl::reqd_work_group_size(4, 4, 4)]] // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+  [[cl::reqd_work_group_size(4, 4, 4)]]       // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
   void
   operator()() const {}
 };
diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp
index 880e6f5bf652..8d6a01d07af0 100644
--- a/sycl/include/CL/sycl/queue.hpp
+++ b/sycl/include/CL/sycl/queue.hpp
@@ -365,8 +365,7 @@ class __SYCL_EXPORT queue {
   // Here we simplify the &CodeLoc declaration to be CODELOCPARAM(&CodeLoc) and
   // CODELOCARG(&CodeLoc) Similarly, the KernelFunc param is simplified to be
   // KERNELFUNCPARAM(KernelFunc) Once the queue kernel functions are defined,
-  // these macros are #undef immediately. Compare this first definition of
-  // single_task here, with the second one below.
+  // these macros are #undef immediately.
 
   // replace CODELOCPARAM(&CodeLoc) with nothing
   // or :   , const detail::code_location &CodeLoc =
@@ -407,10 +406,6 @@ class __SYCL_EXPORT queue {
         CodeLoc);
   }
 
-// Clean up CODELOC and KERNELFUNC macros.
-#undef CODELOCPARAM
-#undef CODELOCARG
-#undef KERNELFUNCPARAM
 
   /// single_task version with a kernel represented as a lambda.
   ///
@@ -418,22 +413,9 @@ class __SYCL_EXPORT queue {
   /// \param KernelFunc is the Kernel functor or lambda
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
-  event single_task(
-      event DepEvent,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-      KernelType KernelFunc
-#else
-      const KernelType &KernelFunc
-#endif
-
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event single_task(event DepEvent,
+                    KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
     return submit(
         [&](handler &CGH) {
           CGH.depends_on(DepEvent);
@@ -449,22 +431,9 @@ class __SYCL_EXPORT queue {
   /// \param KernelFunc is the Kernel functor or lambda
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
-  event single_task(
-      const vector_class<event> &DepEvents,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-      KernelType KernelFunc
-#else
-      const KernelType &KernelFunc
-#endif
-
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event single_task(const vector_class<event> &DepEvents,
+                    KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
     return submit(
         [&](handler &CGH) {
           CGH.depends_on(DepEvents);
@@ -480,22 +449,9 @@ class __SYCL_EXPORT queue {
   /// \param KernelFunc is the Kernel functor or lambda
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
-  event parallel_for(
-      range<1> NumWorkItems,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-      KernelType KernelFunc
-#else
-      const KernelType &KernelFunc
-#endif
-
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event parallel_for(range<1> NumWorkItems,
+                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, KernelFunc, CodeLoc);
   }
 
@@ -506,22 +462,9 @@ class __SYCL_EXPORT queue {
   /// \param KernelFunc is the Kernel functor or lambda
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
-  event parallel_for(
-      range<2> NumWorkItems,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-      KernelType KernelFunc
-#else
-      const KernelType &KernelFunc
-#endif
-
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event parallel_for(range<2> NumWorkItems,
+                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, KernelFunc, CodeLoc);
   }
 
@@ -532,22 +475,9 @@ class __SYCL_EXPORT queue {
   /// \param KernelFunc is the Kernel functor or lambda
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
-  event parallel_for(
-      range<3> NumWorkItems,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-      KernelType KernelFunc
-#else
-      const KernelType &KernelFunc
-#endif
-
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event parallel_for(range<3> NumWorkItems,
+                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, KernelFunc, CodeLoc);
   }
 
@@ -559,22 +489,9 @@ class __SYCL_EXPORT queue {
   /// \param KernelFunc is the Kernel functor or lambda
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
-  event parallel_for(
-      range<1> NumWorkItems, event DepEvent,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-      KernelType KernelFunc
-#else
-      const KernelType &KernelFunc
-#endif
-
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event parallel_for(range<1> NumWorkItems, event DepEvent,
+                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, DepEvent, KernelFunc,
                                          CodeLoc);
   }
@@ -587,22 +504,9 @@ class __SYCL_EXPORT queue {
   /// \param KernelFunc is the Kernel functor or lambda
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
-  event parallel_for(
-      range<2> NumWorkItems, event DepEvent,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-      KernelType KernelFunc
-#else
-      const KernelType &KernelFunc
-#endif
-
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event parallel_for(range<2> NumWorkItems, event DepEvent,
+                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, DepEvent, KernelFunc,
                                          CodeLoc);
   }
@@ -615,22 +519,9 @@ class __SYCL_EXPORT queue {
   /// \param KernelFunc is the Kernel functor or lambda
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
-  event parallel_for(
-      range<3> NumWorkItems, event DepEvent,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-      KernelType KernelFunc
-#else
-      const KernelType &KernelFunc
-#endif
-
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event parallel_for(range<3> NumWorkItems, event DepEvent,
+                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, DepEvent, KernelFunc,
                                          CodeLoc);
   }
@@ -644,17 +535,10 @@ class __SYCL_EXPORT queue {
   /// \param KernelFunc is the Kernel functor or lambda
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
-  event parallel_for(
-      range<1> NumWorkItems, const vector_class<event> &DepEvents,
-      const KernelType &KernelFunc
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event parallel_for(range<1> NumWorkItems,
+                     const vector_class<event> &DepEvents,
+                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, DepEvents, KernelFunc,
                                          CodeLoc);
   }
@@ -668,17 +552,10 @@ class __SYCL_EXPORT queue {
   /// \param KernelFunc is the Kernel functor or lambda
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
-  event parallel_for(
-      range<2> NumWorkItems, const vector_class<event> &DepEvents,
-      const KernelType &KernelFunc
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event parallel_for(range<2> NumWorkItems,
+                     const vector_class<event> &DepEvents,
+                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, DepEvents, KernelFunc,
                                          CodeLoc);
   }
@@ -692,17 +569,10 @@ class __SYCL_EXPORT queue {
   /// \param KernelFunc is the Kernel functor or lambda
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
-  event parallel_for(
-      range<3> NumWorkItems, const vector_class<event> &DepEvents,
-      const KernelType &KernelFunc
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event parallel_for(range<3> NumWorkItems,
+                     const vector_class<event> &DepEvents,
+                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, DepEvents, KernelFunc,
                                          CodeLoc);
   }
@@ -716,17 +586,9 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
-  event parallel_for(
-      range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
-      const KernelType &KernelFunc
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event parallel_for(range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
+                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
     return submit(
         [&](handler &CGH) {
           CGH.template parallel_for<KernelName, KernelType>(
@@ -745,17 +607,10 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
-  event parallel_for(
-      range<Dims> NumWorkItems, id<Dims> WorkItemOffset, event DepEvent,
-      const KernelType &KernelFunc
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event parallel_for(range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
+                     event DepEvent,
+                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
     return submit(
         [&](handler &CGH) {
           CGH.depends_on(DepEvent);
@@ -776,23 +631,10 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
-  event parallel_for(
-      range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
-      const vector_class<event> &DepEvents,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-      KernelType KernelFunc
-#else
-      const KernelType &KernelFunc
-#endif
-
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event parallel_for(range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
+                     const vector_class<event> &DepEvents,
+                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
     return submit(
         [&](handler &CGH) {
           CGH.depends_on(DepEvents);
@@ -811,22 +653,9 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
-  event parallel_for(
-      nd_range<Dims> ExecutionRange,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-      KernelType KernelFunc
-#else
-      const KernelType &KernelFunc
-#endif
-
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event parallel_for(nd_range<Dims> ExecutionRange,
+                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
     return submit(
         [&](handler &CGH) {
           CGH.template parallel_for<KernelName, KernelType>(ExecutionRange,
@@ -845,17 +674,9 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
-  event parallel_for(
-      nd_range<Dims> ExecutionRange, event DepEvent,
-      const KernelType &KernelFunc
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event parallel_for(nd_range<Dims> ExecutionRange, event DepEvent,
+                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
     return submit(
         [&](handler &CGH) {
           CGH.depends_on(DepEvent);
@@ -876,17 +697,10 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
-  event parallel_for(
-      nd_range<Dims> ExecutionRange, const vector_class<event> &DepEvents,
-      const KernelType &KernelFunc
-#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-      ,
-      const detail::code_location &CodeLoc = detail::code_location::current()
-#endif
-  ) {
-#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA
-    const detail::code_location &CodeLoc = {};
-#endif
+  event parallel_for(nd_range<Dims> ExecutionRange,
+                     const vector_class<event> &DepEvents,
+                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
+    CODELOCARG(&CodeLoc);
     return submit(
         [&](handler &CGH) {
           CGH.depends_on(DepEvents);
@@ -896,6 +710,11 @@ class __SYCL_EXPORT queue {
         CodeLoc);
   }
 
+// Clean up CODELOC and KERNELFUNC macros.
+#undef CODELOCPARAM
+#undef CODELOCARG
+#undef KERNELFUNCPARAM
+
   /// Returns whether the queue is in order or OoO
   ///
   /// Equivalent to has_property<property::queue::in_order>()

From b0cffd8199090d8c9c0e83768b6548dbf0488a3d Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Wed, 5 Aug 2020 10:18:10 -0700
Subject: [PATCH 06/17] clang-formatota

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 sycl/include/CL/sycl/queue.hpp | 1 -
 1 file changed, 1 deletion(-)

diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp
index 8d6a01d07af0..adbc33533514 100644
--- a/sycl/include/CL/sycl/queue.hpp
+++ b/sycl/include/CL/sycl/queue.hpp
@@ -406,7 +406,6 @@ class __SYCL_EXPORT queue {
         CodeLoc);
   }
 
-
   /// single_task version with a kernel represented as a lambda.
   ///
   /// \param DepEvent is an event that specifies the kernel dependencies

From 7d43180024a2346113bc4bbe0b819c0c8d8b68b4 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Wed, 5 Aug 2020 13:08:20 -0700
Subject: [PATCH 07/17] rebased to remove conflicts

---
 clang/include/clang/Basic/DiagnosticSemaKinds.td | 3 ---
 1 file changed, 3 deletions(-)

diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 77017bdcfac0..f3b49747761f 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -10995,7 +10995,6 @@ 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">;
-<<<<<<< HEAD
 def err_sycl_compiletime_property_duplication : Error<
   "Can't apply %0 property twice to the same accessor">;
 def err_sycl_invalid_property_list_param_number : Error<
@@ -11006,14 +11005,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>;
->>>>>>> immutable functions redo. WIP.
 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 "

From daf304c69ed402738619ea729b1ffb23190c3ab3 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Wed, 5 Aug 2020 14:59:22 -0700
Subject: [PATCH 08/17] new test needed flag

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/test/SemaSYCL/buffer_location.cpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/clang/test/SemaSYCL/buffer_location.cpp b/clang/test/SemaSYCL/buffer_location.cpp
index bd2daaa01ba4..65742af0cbe1 100644
--- a/clang/test/SemaSYCL/buffer_location.cpp
+++ b/clang/test/SemaSYCL/buffer_location.cpp
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
-// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -verify -pedantic -DTRIGGER_ERROR %s
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -Wno-sycl-2017-compat -verify -pedantic -DTRIGGER_ERROR %s
 
 #include "sycl.hpp"
 

From 388a6db8a3a3ae889266c1ae6a4f4eeef6c396c4 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Wed, 12 Aug 2020 14:39:46 -0700
Subject: [PATCH 09/17] rebased, incorporated latest PR feedback

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/lib/Driver/ToolChains/Clang.cpp        |   6 --
 clang/lib/Frontend/InitPreprocessor.cpp      |   3 +
 clang/test/CodeGenSYCL/sycl-std-flag.cpp     |  27 -----
 clang/test/Preprocessor/sycl-macro.cpp       |   5 +
 clang/test/SemaSYCL/implicit_kernel_type.cpp |   1 -
 sycl/include/CL/sycl/handler.hpp             |  97 +++++-------------
 sycl/include/CL/sycl/queue.hpp               | 102 +++++++++----------
 7 files changed, 84 insertions(+), 157 deletions(-)
 delete mode 100644 clang/test/CodeGenSYCL/sycl-std-flag.cpp

diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 283a458db2fc..4f100cb68eb3 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -4165,15 +4165,9 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
                              .Cases("2017", "1.2.1", "121", "sycl-1.2.1", 2017)
                              .Case("2020", 2020)
                              .Default(0);
-      if (version == 2017)
-        CmdArgs.push_back("-DCL_SYCL_LANGUAGE_VERSION=121");
-      else if (version == 2020)
-        CmdArgs.push_back("-DCL_SYCL_LANGUAGE_VERSION=2020");
-
     } else {
       // Ensure the default version in SYCL mode is 1.2.1 (aka 2017)
       CmdArgs.push_back("-sycl-std=2017");
-      CmdArgs.push_back("-DCL_SYCL_LANGUAGE_VERSION=121");
     }
   }
 
diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
index e0342209d781..a8c78a5846c5 100644
--- a/clang/lib/Frontend/InitPreprocessor.cpp
+++ b/clang/lib/Frontend/InitPreprocessor.cpp
@@ -465,6 +465,9 @@ static void InitializeStandardPredefinedMacros(const TargetInfo &TI,
     // SYCL Version is set to a value when building SYCL applications
     if (LangOpts.SYCLVersion == 2017)
       Builder.defineMacro("CL_SYCL_LANGUAGE_VERSION", "121");
+    else if (LangOpts.SYCLVersion == 2020)
+      Builder.defineMacro("CL_SYCL_LANGUAGE_VERSION", "202001");
+
     if (LangOpts.SYCLValueFitInMaxInt)
       Builder.defineMacro("__SYCL_ID_QUERIES_FIT_IN_INT__", "1");
   }
diff --git a/clang/test/CodeGenSYCL/sycl-std-flag.cpp b/clang/test/CodeGenSYCL/sycl-std-flag.cpp
deleted file mode 100644
index e2b4c352127e..000000000000
--- a/clang/test/CodeGenSYCL/sycl-std-flag.cpp
+++ /dev/null
@@ -1,27 +0,0 @@
-// REQUIRES: clang-driver
-
-// RUN: %clangxx  -fsycl  -fsyntax-only -sycl-std=2017 -DSYCL2017 %s
-// RUN: %clangxx  -fsycl  -fsyntax-only -sycl-std=121 -DSYCL2017 %s
-// RUN: %clangxx  -fsycl  -fsyntax-only -sycl-std=1.2.1 -DSYCL2017 %s
-// RUN: %clangxx  -fsycl  -fsyntax-only -sycl-std=sycl-1.2.1 -DSYCL2017 %s
-// RUN: %clangxx  -fsycl  -fsyntax-only -sycl-std=2020 -DSYCL2020 %s
-
-#if defined(SYCL2017)
-#if !CL_SYCL_LANGUAGE_VERSION
-static_assert(false, "CL_SYCL_LANGUAGE_VERSION should be defined when -sycl-std flag is used");
-#endif
-#if CL_SYCL_LANGUAGE_VERSION > 121
-static_assert(false, "-sycl-std flag should have set CL_SYCL_LANGUAGE_VERSION to 121");
-#endif
-#endif
-
-#if defined(SYCL2020)
-#if !CL_SYCL_LANGUAGE_VERSION
-static_assert(false, "CL_SYCL_LANGUAGE_VERSION should be defined when -sycl-std flag is used");
-#endif
-#if CL_SYCL_LANGUAGE_VERSION < 2020
-static_assert(false, "-sycl-std=2020 flag should set CL_SYCL_LANGUAGE_VERSION to 2020");
-#endif
-#endif
-
-// expected-no-diagnostics
\ No newline at end of file
diff --git a/clang/test/Preprocessor/sycl-macro.cpp b/clang/test/Preprocessor/sycl-macro.cpp
index 21f94e9fac2d..24145987637b 100644
--- a/clang/test/Preprocessor/sycl-macro.cpp
+++ b/clang/test/Preprocessor/sycl-macro.cpp
@@ -1,11 +1,14 @@
 // RUN: %clang_cc1 %s -E -dM | FileCheck %s
 // RUN: %clang_cc1 %s -fsycl -fsycl-id-queries-fit-in-int -fsycl-is-host -sycl-std=2017 -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD %s
 // RUN: %clang_cc1 %s -fsycl -fsycl-id-queries-fit-in-int -fsycl-is-device -sycl-std=2017 -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD %s
+// RUN: %clang_cc1 %s -fsycl -fsycl-id-queries-fit-in-int -fsycl-is-device -sycl-std=2020 -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD-2020 %s
 // RUN: %clang_cc1 %s -fsycl -fsycl-id-queries-fit-in-int -fsycl-is-device -sycl-std=1.2.1 -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD-DEVICE %s
 // RUNx: %clang_cc1 %s -fsycl -fsycl-id-queries-fit-in-int -fsycl-is-device -E -dM -fms-compatibility | FileCheck --check-prefix=CHECK-MSVC %s
 // RUN: %clang_cc1 -fno-sycl-id-queries-fit-in-int %s -E -dM | FileCheck \
 // RUN: --check-prefix=CHECK-NO-SYCL_FIT_IN_INT %s
 
+
+
 // CHECK-NOT:#define __SYCL_DEVICE_ONLY__ 1
 // CHECK-NOT:#define SYCL_EXTERNAL
 // CHECK-NOT:#define CL_SYCL_LANGUAGE_VERSION 121
@@ -14,6 +17,8 @@
 // CHECK-SYCL-STD:#define CL_SYCL_LANGUAGE_VERSION 121
 // CHECK-SYCL-STD:#define __SYCL_ID_QUERIES_FIT_IN_INT__ 1
 
+// CHECK-SYCL-STD-2020:#define CL_SYCL_LANGUAGE_VERSION 202001
+
 // CHECK-SYCL-STD-DEVICE:#define __SYCL_DEVICE_ONLY__ 1
 // CHECK-SYCL-STD-DEVICE:#define __SYCL_ID_QUERIES_FIT_IN_INT__ 1
 
diff --git a/clang/test/SemaSYCL/implicit_kernel_type.cpp b/clang/test/SemaSYCL/implicit_kernel_type.cpp
index 022d34ad5922..b80dbc14ce38 100644
--- a/clang/test/SemaSYCL/implicit_kernel_type.cpp
+++ b/clang/test/SemaSYCL/implicit_kernel_type.cpp
@@ -37,7 +37,6 @@ class handler {
 // expected-no-diagnostics
 #endif
 
-//using namespace cl::sycl;
 
 void function() {
 }
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index 737afb7d0327..2c9feb2b5eba 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -947,6 +947,14 @@ class __SYCL_EXPORT handler {
     MCGType = detail::CG::CODEPLAY_HOST_TASK;
   }
 
+// replace _KERNELFUNCPARAM(KernelFunc) with   KernelType KernelFunc
+//                                     or     const KernelType &KernelFunc
+#ifdef __SYCL_NONCONST_FUNCTOR__
+#define _KERNELFUNCPARAM(a) KernelType a
+#else
+#define _KERNELFUNCPARAM(a) const KernelType &a
+#endif
+
   /// Defines and invokes a SYCL kernel function for the specified range and
   /// offset.
   ///
@@ -963,11 +971,7 @@ class __SYCL_EXPORT handler {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   void parallel_for(range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-                    KernelType KernelFunc) {
-#else
-                    const KernelType &KernelFunc) {
-#endif
+                    _KERNELFUNCPARAM(KernelFunc)) {
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -1000,11 +1004,7 @@ class __SYCL_EXPORT handler {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   void parallel_for(nd_range<Dims> ExecutionRange,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-                    KernelType KernelFunc) {
-#else
-                    const KernelType &KernelFunc) {
-#endif
+                    _KERNELFUNCPARAM(KernelFunc)) {
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -1032,11 +1032,7 @@ class __SYCL_EXPORT handler {
   detail::enable_if_t<Reduction::accessor_mode == access::mode::read_write &&
                       Reduction::has_fast_atomics && !Reduction::is_usm>
   parallel_for(nd_range<Dims> Range, Reduction Redu,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-               KernelType KernelFunc) {
-#else
-               const KernelType &KernelFunc) {
-#endif
+               _KERNELFUNCPARAM(KernelFunc)) {
     intel::detail::reduCGFunc<KernelName>(*this, KernelFunc, Range, Redu,
                                           Redu.getUserAccessor());
   }
@@ -1050,11 +1046,7 @@ class __SYCL_EXPORT handler {
   detail::enable_if_t<Reduction::accessor_mode == access::mode::read_write &&
                       Reduction::has_fast_atomics && Reduction::is_usm>
   parallel_for(nd_range<Dims> Range, Reduction Redu,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-               KernelType KernelFunc) {
-#else
-               const KernelType &KernelFunc) {
-#endif
+               _KERNELFUNCPARAM(KernelFunc)) {
     intel::detail::reduCGFunc<KernelName>(*this, KernelFunc, Range, Redu,
                                           Redu.getUSMPointer());
   }
@@ -1074,11 +1066,7 @@ class __SYCL_EXPORT handler {
   detail::enable_if_t<Reduction::accessor_mode == access::mode::discard_write &&
                       Reduction::has_fast_atomics>
   parallel_for(nd_range<Dims> Range, Reduction Redu,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-               KernelType KernelFunc) {
-#else
-               const KernelType &KernelFunc) {
-#endif
+               _KERNELFUNCPARAM(KernelFunc)) {
     shared_ptr_class<detail::queue_impl> QueueCopy = MQueue;
     auto RWAcc = Redu.getReadWriteScalarAcc(*this);
     intel::detail::reduCGFunc<KernelName>(*this, KernelFunc, Range, Redu,
@@ -1115,11 +1103,7 @@ class __SYCL_EXPORT handler {
             int Dims, typename Reduction>
   detail::enable_if_t<!Reduction::has_fast_atomics>
   parallel_for(nd_range<Dims> Range, Reduction Redu,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-               KernelType KernelFunc) {
-#else
-               const KernelType &KernelFunc) {
-#endif
+               _KERNELFUNCPARAM(KernelFunc)) {
     // This parallel_for() is lowered to the following sequence:
     // 1) Call a kernel that a) call user's lambda function and b) performs
     //    one iteration of reduction, storing the partial reductions/sums
@@ -1194,11 +1178,7 @@ class __SYCL_EXPORT handler {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   void parallel_for_work_group(range<Dims> NumWorkGroups,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-                               KernelType KernelFunc) {
-#else
-                               const KernelType &KernelFunc) {
-#endif
+                               _KERNELFUNCPARAM(KernelFunc)) {
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -1231,11 +1211,7 @@ class __SYCL_EXPORT handler {
             int Dims>
   void parallel_for_work_group(range<Dims> NumWorkGroups,
                                range<Dims> WorkGroupSize,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-                               KernelType KernelFunc) {
-#else
-                               const KernelType &KernelFunc) {
-#endif
+                               _KERNELFUNCPARAM(KernelFunc)) {
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -1337,12 +1313,7 @@ class __SYCL_EXPORT handler {
   /// \param KernelFunc is a lambda that is used if device, queue is bound to,
   /// is a host device.
   template <typename KernelName = detail::auto_name, typename KernelType>
-  void single_task(kernel Kernel,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-                   KernelType KernelFunc) {
-#else
-                   const KernelType &KernelFunc) {
-#endif
+  void single_task(kernel Kernel, _KERNELFUNCPARAM(KernelFunc)) {
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -1382,11 +1353,7 @@ class __SYCL_EXPORT handler {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   void parallel_for(kernel Kernel, range<Dims> NumWorkItems,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-                    KernelType KernelFunc) {
-#else
-                    const KernelType &KernelFunc) {
-#endif
+                    _KERNELFUNCPARAM(KernelFunc)) {
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -1421,12 +1388,7 @@ class __SYCL_EXPORT handler {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   void parallel_for(kernel Kernel, range<Dims> NumWorkItems,
-                    id<Dims> WorkItemOffset,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-                    KernelType KernelFunc) {
-#else
-                    const KernelType &KernelFunc) {
-#endif
+                    id<Dims> WorkItemOffset, _KERNELFUNCPARAM(KernelFunc)) {
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -1463,11 +1425,7 @@ class __SYCL_EXPORT handler {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   void parallel_for(kernel Kernel, nd_range<Dims> NDRange,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-                    KernelType KernelFunc) {
-#else
-                    const KernelType &KernelFunc) {
-#endif
+                    _KERNELFUNCPARAM(KernelFunc)) {
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -1509,11 +1467,7 @@ class __SYCL_EXPORT handler {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   void parallel_for_work_group(kernel Kernel, range<Dims> NumWorkGroups,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-                               KernelType KernelFunc) {
-#else
-                               const KernelType &KernelFunc) {
-#endif
+                               _KERNELFUNCPARAM(KernelFunc)) {
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -1551,11 +1505,7 @@ class __SYCL_EXPORT handler {
             int Dims>
   void parallel_for_work_group(kernel Kernel, range<Dims> NumWorkGroups,
                                range<Dims> WorkGroupSize,
-#ifdef __SYCL_NONCONST_FUNCTOR__
-                               KernelType KernelFunc) {
-#else
-                               const KernelType &KernelFunc) {
-#endif
+                               _KERNELFUNCPARAM(KernelFunc)) {
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
@@ -1579,6 +1529,9 @@ class __SYCL_EXPORT handler {
 #endif // __SYCL_DEVICE_ONLY__
   }
 
+  // Clean up KERNELFUNC macro.
+#undef _KERNELFUNCPARAM
+
   // Explicit copy operations API
 
   /// Copies the content of memory object accessed by Src into the memory
diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp
index adbc33533514..2c6ffdfdfed7 100644
--- a/sycl/include/CL/sycl/queue.hpp
+++ b/sycl/include/CL/sycl/queue.hpp
@@ -362,33 +362,33 @@ class __SYCL_EXPORT queue {
   }
 
   // having _TWO_ mid-param #ifdefs makes the functions very difficult to read.
-  // Here we simplify the &CodeLoc declaration to be CODELOCPARAM(&CodeLoc) and
-  // CODELOCARG(&CodeLoc) Similarly, the KernelFunc param is simplified to be
-  // KERNELFUNCPARAM(KernelFunc) Once the queue kernel functions are defined,
+  // Here we simplify the &CodeLoc declaration to be _CODELOCPARAM(&CodeLoc) and
+  // _CODELOCARG(&CodeLoc) Similarly, the KernelFunc param is simplified to be
+  // _KERNELFUNCPARAM(KernelFunc) Once the queue kernel functions are defined,
   // these macros are #undef immediately.
 
-  // replace CODELOCPARAM(&CodeLoc) with nothing
+  // replace _CODELOCPARAM(&CodeLoc) with nothing
   // or :   , const detail::code_location &CodeLoc =
   // detail::code_location::current()
-  // replace CODELOCARG(&CodeLoc) with nothing
+  // replace _CODELOCARG(&CodeLoc) with nothing
   // or :  const detail::code_location &CodeLoc = {}
 
 #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA
-#define CODELOCPARAM(a)                                                        \
+#define _CODELOCPARAM(a)                                                       \
   , const detail::code_location a = detail::code_location::current()
 
-#define CODELOCARG(a)
+#define _CODELOCARG(a)
 #else
-#define CODELOCPARAM(a)
+#define _CODELOCPARAM(a)
 
-#define CODELOCARG(a) const detail::code_location a = {}
+#define _CODELOCARG(a) const detail::code_location a = {}
 #endif
-// replace KERNELFUNCPARAM(KernelFunc) with   KernelType KernelFunc
+// replace _KERNELFUNCPARAM(KernelFunc) with   KernelType KernelFunc
 //                                     or     const KernelType &KernelFunc
 #ifdef __SYCL_NONCONST_FUNCTOR__
-#define KERNELFUNCPARAM(a) KernelType a
+#define _KERNELFUNCPARAM(a) KernelType a
 #else
-#define KERNELFUNCPARAM(a) const KernelType &a
+#define _KERNELFUNCPARAM(a) const KernelType &a
 #endif
 
   /// single_task version with a kernel represented as a lambda.
@@ -396,8 +396,8 @@ class __SYCL_EXPORT queue {
   /// \param KernelFunc is the Kernel functor or lambda
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
-  event single_task(KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+  event single_task(_KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
 
     return submit(
         [&](handler &CGH) {
@@ -413,8 +413,8 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
   event single_task(event DepEvent,
-                    KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+                    _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
     return submit(
         [&](handler &CGH) {
           CGH.depends_on(DepEvent);
@@ -431,8 +431,8 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
   event single_task(const vector_class<event> &DepEvents,
-                    KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+                    _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
     return submit(
         [&](handler &CGH) {
           CGH.depends_on(DepEvents);
@@ -449,8 +449,8 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(range<1> NumWorkItems,
-                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, KernelFunc, CodeLoc);
   }
 
@@ -462,8 +462,8 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(range<2> NumWorkItems,
-                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, KernelFunc, CodeLoc);
   }
 
@@ -475,8 +475,8 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(range<3> NumWorkItems,
-                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, KernelFunc, CodeLoc);
   }
 
@@ -489,8 +489,8 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(range<1> NumWorkItems, event DepEvent,
-                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, DepEvent, KernelFunc,
                                          CodeLoc);
   }
@@ -504,8 +504,8 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(range<2> NumWorkItems, event DepEvent,
-                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, DepEvent, KernelFunc,
                                          CodeLoc);
   }
@@ -519,8 +519,8 @@ class __SYCL_EXPORT queue {
   /// \param CodeLoc contains the code location of user code
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(range<3> NumWorkItems, event DepEvent,
-                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, DepEvent, KernelFunc,
                                          CodeLoc);
   }
@@ -536,8 +536,8 @@ class __SYCL_EXPORT queue {
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(range<1> NumWorkItems,
                      const vector_class<event> &DepEvents,
-                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, DepEvents, KernelFunc,
                                          CodeLoc);
   }
@@ -553,8 +553,8 @@ class __SYCL_EXPORT queue {
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(range<2> NumWorkItems,
                      const vector_class<event> &DepEvents,
-                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, DepEvents, KernelFunc,
                                          CodeLoc);
   }
@@ -570,8 +570,8 @@ class __SYCL_EXPORT queue {
   template <typename KernelName = detail::auto_name, typename KernelType>
   event parallel_for(range<3> NumWorkItems,
                      const vector_class<event> &DepEvents,
-                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
     return parallel_for_impl<KernelName>(NumWorkItems, DepEvents, KernelFunc,
                                          CodeLoc);
   }
@@ -586,8 +586,8 @@ class __SYCL_EXPORT queue {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   event parallel_for(range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
-                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
     return submit(
         [&](handler &CGH) {
           CGH.template parallel_for<KernelName, KernelType>(
@@ -608,8 +608,8 @@ class __SYCL_EXPORT queue {
             int Dims>
   event parallel_for(range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
                      event DepEvent,
-                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
     return submit(
         [&](handler &CGH) {
           CGH.depends_on(DepEvent);
@@ -632,8 +632,8 @@ class __SYCL_EXPORT queue {
             int Dims>
   event parallel_for(range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
                      const vector_class<event> &DepEvents,
-                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
     return submit(
         [&](handler &CGH) {
           CGH.depends_on(DepEvents);
@@ -653,8 +653,8 @@ class __SYCL_EXPORT queue {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   event parallel_for(nd_range<Dims> ExecutionRange,
-                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
     return submit(
         [&](handler &CGH) {
           CGH.template parallel_for<KernelName, KernelType>(ExecutionRange,
@@ -674,8 +674,8 @@ class __SYCL_EXPORT queue {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   event parallel_for(nd_range<Dims> ExecutionRange, event DepEvent,
-                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
     return submit(
         [&](handler &CGH) {
           CGH.depends_on(DepEvent);
@@ -698,8 +698,8 @@ class __SYCL_EXPORT queue {
             int Dims>
   event parallel_for(nd_range<Dims> ExecutionRange,
                      const vector_class<event> &DepEvents,
-                     KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) {
-    CODELOCARG(&CodeLoc);
+                     _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
+    _CODELOCARG(&CodeLoc);
     return submit(
         [&](handler &CGH) {
           CGH.depends_on(DepEvents);
@@ -710,9 +710,9 @@ class __SYCL_EXPORT queue {
   }
 
 // Clean up CODELOC and KERNELFUNC macros.
-#undef CODELOCPARAM
-#undef CODELOCARG
-#undef KERNELFUNCPARAM
+#undef _CODELOCPARAM
+#undef _CODELOCARG
+#undef _KERNELFUNCPARAM
 
   /// Returns whether the queue is in order or OoO
   ///

From 6bace1ef56d5a84504614025233bf3bf65f219d5 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Wed, 12 Aug 2020 14:50:05 -0700
Subject: [PATCH 10/17] overlooked test.

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/test/SemaSYCL/accessor-type-diagnostics.cpp | 7 +------
 1 file changed, 1 insertion(+), 6 deletions(-)

diff --git a/clang/test/SemaSYCL/accessor-type-diagnostics.cpp b/clang/test/SemaSYCL/accessor-type-diagnostics.cpp
index e6b958d130ca..a409e35527e4 100644
--- a/clang/test/SemaSYCL/accessor-type-diagnostics.cpp
+++ b/clang/test/SemaSYCL/accessor-type-diagnostics.cpp
@@ -6,11 +6,6 @@
 
 using namespace cl::sycl;
 
-template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
-  kernelFunc();
-}
-
 //alias template
 template <typename...>
 using int128alias_t = __uint128_t;
@@ -42,7 +37,7 @@ int main() {
   // -- Accessor of struct that contains a prohibited type.
   accessor<Mesh, 1, access::mode::read_write> struct_acc;
 
-  kernel<class use_local>(
+  kernel_single_task<class use_local>(
       [=]() {
         ok_acc.use();
 

From dc5863d60866bf5622f7ead63722c3a2d730d0e0 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Wed, 12 Aug 2020 14:55:29 -0700
Subject: [PATCH 11/17] clang-format is grumpy

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/lib/Sema/SemaSYCL.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 969ad1c680e5..f200b9cc456c 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -707,7 +707,7 @@ getKernelInvocationKind(FunctionDecl *KernelCallerFunc) {
 
 static const CXXRecordDecl *getKernelObjectType(FunctionDecl *Caller) {
   assert(Caller->getNumParams() > 0 && "Insufficient kernel parameters");
-  
+
   QualType KernelParamTy = (*Caller->param_begin())->getType();
   // In SYCL 2020 kernels are now passed by reference.
   if (KernelParamTy->isReferenceType())

From bbf2f9528f25e3d17e5a170fb3199e9ee2937888 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Wed, 12 Aug 2020 16:59:34 -0700
Subject: [PATCH 12/17] move conformance check

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/lib/Sema/SemaSYCL.cpp | 39 +++++++++++++------------------------
 1 file changed, 13 insertions(+), 26 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index f200b9cc456c..ff7499ba6466 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -708,7 +708,7 @@ getKernelInvocationKind(FunctionDecl *KernelCallerFunc) {
 static const CXXRecordDecl *getKernelObjectType(FunctionDecl *Caller) {
   assert(Caller->getNumParams() > 0 && "Insufficient kernel parameters");
 
-  QualType KernelParamTy = (*Caller->param_begin())->getType();
+  QualType KernelParamTy = Caller->getParamDecl(0)->getType();
   // In SYCL 2020 kernels are now passed by reference.
   if (KernelParamTy->isReferenceType())
     return KernelParamTy->getPointeeCXXRecordDecl();
@@ -717,30 +717,6 @@ static const CXXRecordDecl *getKernelObjectType(FunctionDecl *Caller) {
   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
 /// \param Src  field declaration to construct name from
@@ -2280,6 +2256,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);
@@ -2316,7 +2304,6 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
   // The first argument to the KernelCallerFunc is the lambda object.
   const CXXRecordDecl *KernelObj = getKernelObjectType(KernelCallerFunc);
   assert(KernelObj && "invalid kernel caller");
-  checkKernelAndCaller(*this, KernelCallerFunc, KernelObj);
 
   // Calculate both names, since Integration headers need both.
   std::string CalculatedName, StableName;

From d84c1e9c0300fef58059fd6e64b49d15964f27d0 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Wed, 12 Aug 2020 17:04:06 -0700
Subject: [PATCH 13/17] hmm - git clang-format is skipping certain whitespace
 changes.

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/lib/Sema/SemaSYCL.cpp | 1 -
 1 file changed, 1 deletion(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index ff7499ba6466..ac0ea7b12518 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -717,7 +717,6 @@ static const CXXRecordDecl *getKernelObjectType(FunctionDecl *Caller) {
   return KernelParamTy->getAsCXXRecordDecl();
 }
 
-
 /// Creates a kernel parameter descriptor
 /// \param Src  field declaration to construct name from
 /// \param Ty   the desired parameter type

From 0f9fe01403c850ebe318849d9632f77ef84c4513 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Thu, 13 Aug 2020 08:52:44 -0700
Subject: [PATCH 14/17] removed unused var.  stray from earlier change

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/lib/Driver/ToolChains/Clang.cpp | 5 -----
 1 file changed, 5 deletions(-)

diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 4f100cb68eb3..61ee698ffbd9 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -4160,11 +4160,6 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
     if (SYCLStdArg) {
       SYCLStdArg->render(Args, CmdArgs);
       CmdArgs.push_back("-fsycl-std-layout-kernel-params");
-
-      unsigned version = llvm::StringSwitch<unsigned>(SYCLStdArg->getValue())
-                             .Cases("2017", "1.2.1", "121", "sycl-1.2.1", 2017)
-                             .Case("2020", 2020)
-                             .Default(0);
     } else {
       // Ensure the default version in SYCL mode is 1.2.1 (aka 2017)
       CmdArgs.push_back("-sycl-std=2017");

From b87d642072b5f9af22f3a2ec34205678a6bddb1e Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Thu, 13 Aug 2020 09:32:16 -0700
Subject: [PATCH 15/17] latest review requests

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 sycl/include/CL/sycl/handler.hpp | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index 2c9feb2b5eba..084f1a4b2641 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -33,7 +33,8 @@
 #include <memory>
 #include <type_traits>
 
-#if !CL_SYCL_LANGUAGE_VERSION || CL_SYCL_LANGUAGE_VERSION < 2020
+// CL_SYCL_LANGUAGE_VERSION is 4 digit year followed by 2 digit revision
+#if !CL_SYCL_LANGUAGE_VERSION || CL_SYCL_LANGUAGE_VERSION < 202001
 #define __SYCL_NONCONST_FUNCTOR__
 #endif
 

From 728dd4558c9921f20b514542b0ec6505afc427eb Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Thu, 13 Aug 2020 11:00:57 -0700
Subject: [PATCH 16/17] clang-format, of course

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/test/Preprocessor/sycl-macro.cpp | 2 --
 1 file changed, 2 deletions(-)

diff --git a/clang/test/Preprocessor/sycl-macro.cpp b/clang/test/Preprocessor/sycl-macro.cpp
index 24145987637b..e80c1fe1d3bf 100644
--- a/clang/test/Preprocessor/sycl-macro.cpp
+++ b/clang/test/Preprocessor/sycl-macro.cpp
@@ -7,8 +7,6 @@
 // RUN: %clang_cc1 -fno-sycl-id-queries-fit-in-int %s -E -dM | FileCheck \
 // RUN: --check-prefix=CHECK-NO-SYCL_FIT_IN_INT %s
 
-
-
 // CHECK-NOT:#define __SYCL_DEVICE_ONLY__ 1
 // CHECK-NOT:#define SYCL_EXTERNAL
 // CHECK-NOT:#define CL_SYCL_LANGUAGE_VERSION 121

From 391ba54e16dfa2afe589b2ab2043edcb0f53778d Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Fri, 14 Aug 2020 09:35:18 -0700
Subject: [PATCH 17/17] new SYCL_LANGUAGE_VERSION macro and etc.

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 clang/lib/Frontend/InitPreprocessor.cpp | 8 +++++---
 clang/test/Preprocessor/sycl-macro.cpp  | 3 ++-
 sycl/include/CL/sycl/handler.hpp        | 6 +++---
 sycl/test/basic_tests/macros.cpp        | 3 +--
 sycl/test/basic_tests/set_arg_error.cpp | 2 +-
 5 files changed, 12 insertions(+), 10 deletions(-)

diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
index a8c78a5846c5..b48917f7b1eb 100644
--- a/clang/lib/Frontend/InitPreprocessor.cpp
+++ b/clang/lib/Frontend/InitPreprocessor.cpp
@@ -463,10 +463,12 @@ 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");
-    else if (LangOpts.SYCLVersion == 2020)
-      Builder.defineMacro("CL_SYCL_LANGUAGE_VERSION", "202001");
+      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");
diff --git a/clang/test/Preprocessor/sycl-macro.cpp b/clang/test/Preprocessor/sycl-macro.cpp
index e80c1fe1d3bf..8a0be0841da4 100644
--- a/clang/test/Preprocessor/sycl-macro.cpp
+++ b/clang/test/Preprocessor/sycl-macro.cpp
@@ -13,9 +13,10 @@
 // CHECK-NOT:#define __SYCL_ID_QUERIES_FIT_IN_INT__ 1
 
 // CHECK-SYCL-STD:#define CL_SYCL_LANGUAGE_VERSION 121
+// CHECK-SYCL-STD:#define SYCL_LANGUAGE_VERSION 201707
 // CHECK-SYCL-STD:#define __SYCL_ID_QUERIES_FIT_IN_INT__ 1
 
-// CHECK-SYCL-STD-2020:#define CL_SYCL_LANGUAGE_VERSION 202001
+// CHECK-SYCL-STD-2020:#define SYCL_LANGUAGE_VERSION 202001
 
 // CHECK-SYCL-STD-DEVICE:#define __SYCL_DEVICE_ONLY__ 1
 // CHECK-SYCL-STD-DEVICE:#define __SYCL_ID_QUERIES_FIT_IN_INT__ 1
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index 084f1a4b2641..03b3d5eb9d74 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -33,8 +33,8 @@
 #include <memory>
 #include <type_traits>
 
-// CL_SYCL_LANGUAGE_VERSION is 4 digit year followed by 2 digit revision
-#if !CL_SYCL_LANGUAGE_VERSION || CL_SYCL_LANGUAGE_VERSION < 202001
+// SYCL_LANGUAGE_VERSION is 4 digit year followed by 2 digit revision
+#if !SYCL_LANGUAGE_VERSION || SYCL_LANGUAGE_VERSION < 202001
 #define __SYCL_NONCONST_FUNCTOR__
 #endif
 
@@ -812,7 +812,7 @@ class __SYCL_EXPORT handler {
   template <typename T> struct ShouldEnableSetArg {
     static constexpr bool value =
         std::is_trivially_copyable<detail::remove_reference_t<T>>::value
-#if CL_SYCL_LANGUAGE_VERSION && CL_SYCL_LANGUAGE_VERSION <= 121
+#if SYCL_LANGUAGE_VERSION && SYCL_LANGUAGE_VERSION <= 201707
             && std::is_standard_layout<detail::remove_reference_t<T>>::value
 #endif
         || is_same_type<sampler, T>::value // Sampler
diff --git a/sycl/test/basic_tests/macros.cpp b/sycl/test/basic_tests/macros.cpp
index 2be95b94964a..726f070406e2 100644
--- a/sycl/test/basic_tests/macros.cpp
+++ b/sycl/test/basic_tests/macros.cpp
@@ -11,8 +11,7 @@
 #include <iostream>
 
 int main() {
-  std::cout << "SYCL language version: " << CL_SYCL_LANGUAGE_VERSION
-            << std::endl;
+  std::cout << "SYCL language version: " << SYCL_LANGUAGE_VERSION << std::endl;
   std::cout << "SYCL compiler version: " << __SYCL_COMPILER_VERSION
             << std::endl;
   return 0;
diff --git a/sycl/test/basic_tests/set_arg_error.cpp b/sycl/test/basic_tests/set_arg_error.cpp
index 52df555902b4..d4215766c7a5 100644
--- a/sycl/test/basic_tests/set_arg_error.cpp
+++ b/sycl/test/basic_tests/set_arg_error.cpp
@@ -44,7 +44,7 @@ int main() {
         5, ntc);
     h.set_arg( // expected-error {{no matching member function for call to 'set_arg'}}
         4, NonTriviallyCopyable{});
-#if CL_SYCL_LANGUAGE_VERSION && CL_SYCL_LANGUAGE_VERSION <= 121
+#if SYCL_LANGUAGE_VERSION && SYCL_LANGUAGE_VERSION <= 201707
     NonStdLayout nstd;
     h.set_arg( // expected-error {{no matching member function for call to 'set_arg'}}
         6, nstd);