diff --git a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
index 9b9ea1dfe26fd..f8504f12d4854 100644
--- a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
+++ b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
@@ -1,14 +1,11 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump %s | FileCheck %s
 
 // This test checks that compiler generates correct initialization for arguments
 // that have struct or built-in type inside the OpenCL kernel
 
-#include "Inputs/sycl.hpp"
+#include "sycl.hpp"
 
-template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
-  kernelFunc();
-}
+sycl::queue deviceQueue;
 
 struct test_struct {
   int data;
@@ -18,10 +15,12 @@ struct test_struct {
 };
 
 void test(const int some_const) {
-  kernel<class kernel_const>(
-      [=]() {
-        int a = some_const;
-      });
+  deviceQueue.submit([&](sycl::handler &h) {
+    h.single_task<class kernel_const>(
+        [=]() {
+          int a = some_const;
+        });
+  });
 }
 
 int main() {
@@ -31,20 +30,29 @@ int main() {
   int *ptr_array[2];
   test_struct s;
   s.data = data;
-  kernel<class kernel_int>(
-      [=]() {
-        int kernel_data = data;
-      });
-  kernel<class kernel_struct>(
-      [=]() {
-        test_struct k_s;
-        k_s = s;
-      });
-  kernel<class kernel_pointer>(
-      [=]() {
-        new_data_addr[0] = data_addr[0];
-        int *local = ptr_array[1];
-      });
+
+  deviceQueue.submit([&](sycl::handler &h) {
+    h.single_task<class kernel_int>(
+        [=]() {
+          int kernel_data = data;
+        });
+  });
+
+  deviceQueue.submit([&](sycl::handler &h) {
+    h.single_task<class kernel_struct>(
+        [=]() {
+          test_struct k_s;
+          k_s = s;
+        });
+  });
+
+  deviceQueue.submit([&](sycl::handler &h) {
+    h.single_task<class kernel_pointer>(
+        [=]() {
+          new_data_addr[0] = data_addr[0];
+          int *local = ptr_array[1];
+        });
+  });
 
   const int some_const = 10;
   test(some_const);
diff --git a/clang/test/SemaSYCL/call-to-undefined-function.cpp b/clang/test/SemaSYCL/call-to-undefined-function.cpp
index 8fd8a6fd7ca73..210ccaa8aa121 100644
--- a/clang/test/SemaSYCL/call-to-undefined-function.cpp
+++ b/clang/test/SemaSYCL/call-to-undefined-function.cpp
@@ -1,4 +1,8 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-sycl-2017-compat -verify -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -verify -fsyntax-only %s
+
+#include "sycl.hpp"
+
+sycl::queue deviceQueue;
 
 void defined() {
 }
@@ -8,11 +12,6 @@ void undefined();
 
 SYCL_EXTERNAL void undefinedExternal();
 
-template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
-  kernelFunc();
-}
-
 template <typename T>
 void definedTpl() {
 }
@@ -95,76 +94,78 @@ int main() {
   // No problems in host code
   undefined();
 
-  kernel_single_task<class CallToUndefinedFnTester>([]() {
-    // expected-note@-1 {{called by 'operator()'}}
-    // expected-note@-2 {{called by 'operator()'}}
-
-    // simple functions
-    defined();
-    undefinedExternal();
-    undefined();
-    // expected-error@-1 {{SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute}}
-
-    // templated functions
-    definedTpl<int>();
-    undefinedExternalTpl<int>();
-    undefinedTpl<int>();
-    // expected-error@-1 {{SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute}}
-
-    // partially specialized template function
-    definedPartialTpl<int, false>();
-    definedPartialTpl<int, true>();
-    definedPartialTpl<char, false>();
-    definedPartialTpl<char, true>();
-
-    // template class with specialization
-    {
-      Tpl<int, false> tpl;
-      tpl.defined();
-    }
-
-    {
-      Tpl<int, true> tpl;
-      tpl.defined();
-    }
-
-    // template class with template method, both have specializations.
-    {
-      TplWithTplMethod<int, false> tpl;
-      tpl.defined<char, false>();
-      tpl.defined<char, true>();
-      tpl.defined<int, false>();
-      tpl.defined<int, true>();
-    }
-
-    {
-      TplWithTplMethod<int, true> tpl;
-      tpl.defined<char, false>();
-      tpl.defined<char, true>();
-      tpl.defined<int, false>();
-      tpl.defined<int, true>();
-    }
-
-    {
-      TplWithTplMethod2<int, false> tpl;
-      tpl.defined<char, false>();
-      tpl.defined<char, true>();
-      tpl.defined<int, false>();
-      tpl.defined<int, true>();
-    }
-
-    {
-      TplWithTplMethod2<int, true> tpl;
-      tpl.defined<char, false>();
-      tpl.defined<char, true>();
-      tpl.defined<int, false>();
-      tpl.defined<int, true>();
-    }
-
-    // forward-declared function
-    useFwDeclFn();
-    forwardDeclFn();
-    forwardDeclFn2();
+  deviceQueue.submit([&](sycl::handler &h) {
+    h.single_task<class CallToUndefinedFnTester>([]() {
+      // expected-note@-1 {{called by 'operator()'}}
+      // expected-note@-2 {{called by 'operator()'}}
+
+      // simple functions
+      defined();
+      undefinedExternal();
+      undefined();
+      // expected-error@-1 {{SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute}}
+
+      // templated functions
+      definedTpl<int>();
+      undefinedExternalTpl<int>();
+      undefinedTpl<int>();
+      // expected-error@-1 {{SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute}}
+
+      // partially specialized template function
+      definedPartialTpl<int, false>();
+      definedPartialTpl<int, true>();
+      definedPartialTpl<char, false>();
+      definedPartialTpl<char, true>();
+
+      // template class with specialization
+      {
+        Tpl<int, false> tpl;
+        tpl.defined();
+      }
+
+      {
+        Tpl<int, true> tpl;
+        tpl.defined();
+      }
+
+      // template class with template method, both have specializations.
+      {
+        TplWithTplMethod<int, false> tpl;
+        tpl.defined<char, false>();
+        tpl.defined<char, true>();
+        tpl.defined<int, false>();
+        tpl.defined<int, true>();
+      }
+
+      {
+        TplWithTplMethod<int, true> tpl;
+        tpl.defined<char, false>();
+        tpl.defined<char, true>();
+        tpl.defined<int, false>();
+        tpl.defined<int, true>();
+      }
+
+      {
+        TplWithTplMethod2<int, false> tpl;
+        tpl.defined<char, false>();
+        tpl.defined<char, true>();
+        tpl.defined<int, false>();
+        tpl.defined<int, true>();
+      }
+
+      {
+        TplWithTplMethod2<int, true> tpl;
+        tpl.defined<char, false>();
+        tpl.defined<char, true>();
+        tpl.defined<int, false>();
+        tpl.defined<int, true>();
+      }
+
+      // forward-declared function
+      useFwDeclFn();
+      forwardDeclFn();
+      forwardDeclFn2();
+    });
   });
 }
 
diff --git a/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp b/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp
index b4df077725393..904b28bda7c87 100644
--- a/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp
+++ b/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp
@@ -1,16 +1,8 @@
-// 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
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -aux-triple x86_64-unknown-linux-gnu -Wno-sycl-2017-compat -verify -fsyntax-only  %s
 
-inline namespace cl {
-namespace sycl {
+#include "sycl.hpp"
 
-template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
-  // expected-note@+1 {{called by 'kernel_single_task<AName, (lambda}}
-  kernelFunc();
-}
-
-} // namespace sycl
-} // namespace cl
+sycl::queue deviceQueue;
 
 int main(int argc, char **argv) {
   //_mm_prefetch is an x86-64 specific builtin where the second integer parameter is required to be a constant
@@ -19,9 +11,13 @@ int main(int argc, char **argv) {
 
   _mm_prefetch("test", 8); // expected-error {{argument value 8 is outside the valid range [0, 7]}}
 
-  cl::sycl::kernel_single_task<class AName>([]() {
-    _mm_prefetch("test", 4); // expected-error {{builtin is not supported on this target}}
-    _mm_prefetch("test", 8); // expected-error {{argument value 8 is outside the valid range [0, 7]}} expected-error {{builtin is not supported on this target}}
+  deviceQueue.submit([&](sycl::handler &h) {
+    // expected-note@Inputs/sycl.hpp:212 {{called by 'kernel_single_task<AName, (lambda}}
+    h.single_task<class AName>([]() {
+      _mm_prefetch("test", 4); // expected-error {{builtin is not supported on this target}}
+      _mm_prefetch("test", 8); // expected-error {{argument value 8 is outside the valid range [0, 7]}} expected-error {{builtin is not supported on this target}}
+    });
   });
+
   return 0;
-}
+}
\ No newline at end of file
diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp
index 90a05ab3b69e4..f634e7bc6933c 100644
--- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp
+++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp
@@ -1,30 +1,21 @@
-// RUN: %clang_cc1  -fsycl -triple spir64 -fsycl-is-device          \
-// RUN:  -aux-triple x86_64-unknown-linux-gnu -Wno-sycl-2017-compat \
+// RUN: %clang_cc1  -fsycl -internal-isystem %S/Inputs -sycl-std=2020 -triple spir64 -fsycl-is-device \
+// RUN:  -aux-triple x86_64-unknown-linux-gnu \
 // RUN:  -verify -fsyntax-only  %s
-// RUN: %clang_cc1  -fsycl -triple spir64 -fsycl-is-device          \
-// RUN:  -aux-triple x86_64-pc-windows-msvc -Wno-sycl-2017-compat   \
+// RUN: %clang_cc1  -fsycl -internal-isystem %S/Inputs -sycl-std=2020 -triple spir64 -fsycl-is-device \
+// RUN:  -aux-triple x86_64-pc-windows-msvc   \
 // RUN:  -verify -fsyntax-only  %s
 //
 // Ensure that the SYCL diagnostics that are typically deferred are correctly emitted.
 
+#include "sycl.hpp"
+
+sycl::queue deviceQueue;
+
 namespace std {
 class type_info;
 typedef __typeof__(sizeof(int)) size_t;
 } // namespace std
 
-// testing that the deferred diagnostics work in conjunction with the SYCL namespaces.
-inline namespace cl {
-namespace sycl {
-
-template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
-  // expected-note@+1 3{{called by 'kernel_single_task<AName, (lambda}}
-  kernelFunc();
-}
-
-} // namespace sycl
-} // namespace cl
-
 //variadic functions from SYCL kernels emit a deferred diagnostic
 void variadic(int, ...) {}
 
@@ -72,107 +63,113 @@ void foo(){};
 template <typename T>
 void setup_sycl_operation(const T VA[]) {
 
-  cl::sycl::kernel_single_task<class AName>([]() {
-    // ======= Zero Length Arrays Not Allowed in Kernel ==========
-    // expected-error@+1 {{zero-length arrays are not permitted in C++}}
-    int MalArray[0];
-    // expected-error@+1 {{zero-length arrays are not permitted in C++}}
-    intDef MalArrayDef[0];
-    // ---- false positive tests. These should not generate any errors.
-    foo<int[0]>();
-    std::size_t arrSz = sizeof(int[0]);
-
-    // ======= Float128 Not Allowed in Kernel ==========
-    // expected-note@+2 {{'malFloat' defined here}}
-    // expected-error@+1 {{'__float128' is not supported on this target}}
-    __float128 malFloat = 40;
-    // expected-error@+1 {{'__float128' is not supported on this target}}
-    trickyFloatType malFloatTrick = 41;
-    // expected-error@+1 {{'__float128' is not supported on this target}}
-    floatDef malFloatDef = 44;
-    // expected-error@+2 {{'malFloat' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
-    // expected-error@+1 {{'__float128' is not supported on this target}}
-    auto whatFloat = malFloat;
-    // expected-error@+2 {{'bar<__float128>' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
-    // expected-error@+1 {{'__float128' is not supported on this target}}
-    auto malAutoTemp5 = bar<__float128>();
-    // expected-error@+2 {{'bar<const __float128>' requires 128 bit size 'const __float128' type support, but device 'spir64' does not support it}}
-    // expected-error@+1 {{'__float128' is not supported on this target}}
-    auto malAutoTemp6 = bar<trickyFloatType>();
-    // expected-error@+1 {{'__float128' is not supported on this target}}
-    decltype(malFloat) malDeclFloat = 42;
-    // ---- false positive tests
-    std::size_t someSz = sizeof(__float128);
-    foo<__float128>();
-
-    // ======= __int128 Not Allowed in Kernel ==========
-    // expected-note@+2 {{'malIntent' defined here}}
-    // expected-error@+1 {{'__int128' is not supported on this target}}
-    __int128 malIntent = 2;
-    // expected-error@+1 {{'__int128' is not supported on this target}}
-    tricky128Type mal128Trick = 2;
-    // expected-error@+1 {{'__int128' is not supported on this target}}
-    int128Def malIntDef = 9;
-    // expected-error@+2 {{'malIntent' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
-    // expected-error@+1 {{'__int128' is not supported on this target}}
-    auto whatInt128 = malIntent;
-    // expected-error@+2 {{'bar<__int128>' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
-    // expected-error@+1 {{'__int128' is not supported on this target}}
-    auto malAutoTemp = bar<__int128>();
-    // expected-error@+2 {{'bar<const __int128>' requires 128 bit size 'const __int128' type support, but device 'spir64' does not support it}}
-    // expected-error@+1 {{'__int128' is not supported on this target}}
-    auto malAutoTemp2 = bar<tricky128Type>();
-    // expected-error@+1 {{'__int128' is not supported on this target}}
-    decltype(malIntent) malDeclInt = 2;
-
-    // expected-error@+1 {{'__int128' is not supported on this target}}
-    __int128_t malInt128 = 2;
-    // expected-note@+2 {{'malUInt128' defined here}}
-    // expected-error@+1 {{'unsigned __int128' is not supported on this target}}
-    __uint128_t malUInt128 = 3;
-    // expected-error@+1 {{'unsigned __int128' is not supported on this target}}
-    megeType malTypeDefTrick = 4;
-    // expected-error@+1 {{'__int128' is not supported on this target}}
-    int128tDef malInt2Def = 6;
-    // expected-error@+2 {{'malUInt128' requires 128 bit size '__uint128_t' (aka 'unsigned __int128') type support, but device 'spir64' does not support it}}
-    // expected-error@+1 {{'unsigned __int128' is not supported on this target}}
-    auto whatUInt = malUInt128;
-    // expected-error@+2 {{'bar<__int128>' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
-    // expected-error@+1 {{'__int128' is not supported on this target}}
-    auto malAutoTemp3 = bar<__int128_t>();
-    // expected-error@+2 {{'bar<const unsigned __int128>' requires 128 bit size 'const unsigned __int128' type support, but device 'spir64' does not support it}}
-    // expected-error@+1 {{'unsigned __int128' is not supported on this target}}
-    auto malAutoTemp4 = bar<megeType>();
-    // expected-error@+1 {{'__int128' is not supported on this target}}
-    decltype(malInt128) malDeclInt128 = 5;
-
-    // ---- false positive tests These should not generate any errors.
-    std::size_t i128Sz = sizeof(__int128);
-    foo<__int128>();
-    std::size_t u128Sz = sizeof(__uint128_t);
-    foo<__int128_t>();
-
-    // ========= variadic
-    //expected-error@+1 {{SYCL kernel cannot call a variadic function}}
-    variadic(5);
+  deviceQueue.submit([&](sycl::handler &h) {
+    // expected-note@Inputs/sycl.hpp:212 {{called by 'kernel_single_task<AName, (lambda}}
+    h.single_task<class AName>([]() {
+      // ======= Zero Length Arrays Not Allowed in Kernel ==========
+      // expected-error@+1 {{zero-length arrays are not permitted in C++}}
+      int MalArray[0];
+      // expected-error@+1 {{zero-length arrays are not permitted in C++}}
+      intDef MalArrayDef[0];
+      // ---- false positive tests. These should not generate any errors.
+      foo<int[0]>();
+      std::size_t arrSz = sizeof(int[0]);
+
+      // ======= Float128 Not Allowed in Kernel ==========
+      // expected-note@+2 {{'malFloat' defined here}}
+      // expected-error@+1 {{'__float128' is not supported on this target}}
+      __float128 malFloat = 40;
+      // expected-error@+1 {{'__float128' is not supported on this target}}
+      trickyFloatType malFloatTrick = 41;
+      // expected-error@+1 {{'__float128' is not supported on this target}}
+      floatDef malFloatDef = 44;
+      // expected-error@+2 {{'malFloat' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+      // expected-error@+1 {{'__float128' is not supported on this target}}
+      auto whatFloat = malFloat;
+      // expected-error@+2 {{'bar<__float128>' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+      // expected-error@+1 {{'__float128' is not supported on this target}}
+      auto malAutoTemp5 = bar<__float128>();
+      // expected-error@+2 {{'bar<const __float128>' requires 128 bit size 'const __float128' type support, but device 'spir64' does not support it}}
+      // expected-error@+1 {{'__float128' is not supported on this target}}
+      auto malAutoTemp6 = bar<trickyFloatType>();
+      // expected-error@+1 {{'__float128' is not supported on this target}}
+      decltype(malFloat) malDeclFloat = 42;
+      // ---- false positive tests
+      std::size_t someSz = sizeof(__float128);
+      foo<__float128>();
+
+      // ======= __int128 Not Allowed in Kernel ==========
+      // expected-note@+2 {{'malIntent' defined here}}
+      // expected-error@+1 {{'__int128' is not supported on this target}}
+      __int128 malIntent = 2;
+      // expected-error@+1 {{'__int128' is not supported on this target}}
+      tricky128Type mal128Trick = 2;
+      // expected-error@+1 {{'__int128' is not supported on this target}}
+      int128Def malIntDef = 9;
+      // expected-error@+2 {{'malIntent' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
+      // expected-error@+1 {{'__int128' is not supported on this target}}
+      auto whatInt128 = malIntent;
+      // expected-error@+2 {{'bar<__int128>' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
+      // expected-error@+1 {{'__int128' is not supported on this target}}
+      auto malAutoTemp = bar<__int128>();
+      // expected-error@+2 {{'bar<const __int128>' requires 128 bit size 'const __int128' type support, but device 'spir64' does not support it}}
+      // expected-error@+1 {{'__int128' is not supported on this target}}
+      auto malAutoTemp2 = bar<tricky128Type>();
+      // expected-error@+1 {{'__int128' is not supported on this target}}
+      decltype(malIntent) malDeclInt = 2;
+
+      // expected-error@+1 {{'__int128' is not supported on this target}}
+      __int128_t malInt128 = 2;
+      // expected-note@+2 {{'malUInt128' defined here}}
+      // expected-error@+1 {{'unsigned __int128' is not supported on this target}}
+      __uint128_t malUInt128 = 3;
+      // expected-error@+1 {{'unsigned __int128' is not supported on this target}}
+      megeType malTypeDefTrick = 4;
+      // expected-error@+1 {{'__int128' is not supported on this target}}
+      int128tDef malInt2Def = 6;
+      // expected-error@+2 {{'malUInt128' requires 128 bit size '__uint128_t' (aka 'unsigned __int128') type support, but device 'spir64' does not support it}}
+      // expected-error@+1 {{'unsigned __int128' is not supported on this target}}
+      auto whatUInt = malUInt128;
+      // expected-error@+2 {{'bar<__int128>' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
+      // expected-error@+1 {{'__int128' is not supported on this target}}
+      auto malAutoTemp3 = bar<__int128_t>();
+      // expected-error@+2 {{'bar<const unsigned __int128>' requires 128 bit size 'const unsigned __int128' type support, but device 'spir64' does not support it}}
+      // expected-error@+1 {{'unsigned __int128' is not supported on this target}}
+      auto malAutoTemp4 = bar<megeType>();
+      // expected-error@+1 {{'__int128' is not supported on this target}}
+      decltype(malInt128) malDeclInt128 = 5;
+
+      // ---- false positive tests These should not generate any errors.
+      std::size_t i128Sz = sizeof(__int128);
+      foo<__int128>();
+      std::size_t u128Sz = sizeof(__uint128_t);
+      foo<__int128_t>();
+
+      // ========= variadic
+      //expected-error@+1 {{SYCL kernel cannot call a variadic function}}
+      variadic(5);
+    });
   });
 }
 
 int main(int argc, char **argv) {
 
   // --- direct lambda testing ---
-  cl::sycl::kernel_single_task<class AName>([]() {
-    // expected-error@+1 {{zero-length arrays are not permitted in C++}}
-    int BadArray[0];
+  deviceQueue.submit([&](sycl::handler &h) {
+    // expected-note@Inputs/sycl.hpp:212 2 {{called by 'kernel_single_task<AName, (lambda}}
+    h.single_task<class AName>([]() {
+      // expected-error@+1 {{zero-length arrays are not permitted in C++}}
+      int BadArray[0];
 
-    // expected-error@+1 {{'__float128' is not supported on this target}}
-    __float128 badFloat = 40; // this SHOULD  trigger a diagnostic
+      // expected-error@+1 {{'__float128' is not supported on this target}}
+      __float128 badFloat = 40; // this SHOULD  trigger a diagnostic
 
-    //expected-error@+1 {{SYCL kernel cannot call a variadic function}}
-    variadic(5);
+      //expected-error@+1 {{SYCL kernel cannot call a variadic function}}
+      variadic(5);
 
-    // expected-note@+1 {{called by 'operator()'}}
-    calledFromKernel(10);
+      // expected-note@+1 {{called by 'operator()'}}
+      calledFromKernel(10);
+    });
   });
 
   // --- lambda in specialized function testing ---
diff --git a/clang/test/SemaSYCL/float128.cpp b/clang/test/SemaSYCL/float128.cpp
index b61787a153fd0..0f9d6180d5652 100644
--- a/clang/test/SemaSYCL/float128.cpp
+++ b/clang/test/SemaSYCL/float128.cpp
@@ -1,5 +1,9 @@
-// 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
+// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -verify -fsyntax-only %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsycl -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -fsyntax-only %s
+
+#include "sycl.hpp"
+
+sycl::queue deviceQueue;
 
 typedef __float128 BIGTY;
 
@@ -62,48 +66,51 @@ void foo2(){};
 // expected-note@+1 2{{'foo' defined here}}
 __float128 foo(__float128 P) { return P; }
 
-template <typename Name, typename Func>
-__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
-  // expected-note@+1 6{{called by 'kernel}}
-  kernelFunc();
-}
-
 int main() {
   // expected-note@+1 {{'CapturedToDevice' defined here}}
   __float128 CapturedToDevice = 1;
   host_ok();
-  kernel<class variables>([=]() {
-    // expected-error@+1 {{'__float128' is not supported on this target}}
-    decltype(CapturedToDevice) D;
-    // expected-error@+2 {{'CapturedToDevice' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
-    // expected-error@+1 {{'__float128' is not supported on this target}}
-    auto C = CapturedToDevice;
-    // expected-note@+1 3{{used here}}
-    Z<__float128> S;
-    // expected-error@+1 {{'field1' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
-    S.field1 += 1;
-    // expected-error@+1 {{'field' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
-    S.field = 1;
+  deviceQueue.submit([&](sycl::handler &h) {
+    // expected-note@Inputs/sycl.hpp:212 {{called by 'kernel_single_task<variables, (lambda}}
+    h.single_task<class variables>([=]() {
+      // expected-error@+1 {{'__float128' is not supported on this target}}
+      decltype(CapturedToDevice) D;
+      // expected-error@+2 {{'CapturedToDevice' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+      // expected-error@+1 {{'__float128' is not supported on this target}}
+      auto C = CapturedToDevice;
+      // expected-note@+1 3{{used here}}
+      Z<__float128> S;
+      // expected-error@+1 {{'field1' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+      S.field1 += 1;
+      // expected-error@+1 {{'field' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+      S.field = 1;
+    });
   });
 
-  kernel<class functions>([=]() {
-    // expected-note@+1 2{{called by 'operator()'}}
-    usage();
-    // expected-note@+2 {{'BBBB' defined here}}
-    // expected-error@+1 {{'__float128' is not supported on this target}}
-    BIGTY BBBB;
-    // expected-error@+4 {{'__float128' is not supported on this target}}
-    // expected-note@+3 {{called by 'operator()'}}
-    // expected-error@+2 2{{'foo' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
-    // expected-error@+1 {{'BBBB' requires 128 bit size 'BIGTY' (aka '__float128') type support, but device 'spir64' does not support it}}
-    auto A = foo(BBBB);
+  deviceQueue.submit([&](sycl::handler &h) {
+    // expected-note@Inputs/sycl.hpp:212 4{{called by 'kernel_single_task<functions, (lambda}}
+    h.single_task<class functions>([=]() {
+      // expected-note@+1 2{{called by 'operator()'}}
+      usage();
+      // expected-note@+2 {{'BBBB' defined here}}
+      // expected-error@+1 {{'__float128' is not supported on this target}}
+      BIGTY BBBB;
+      // expected-error@+4 {{'__float128' is not supported on this target}}
+      // expected-note@+3 {{called by 'operator()'}}
+      // expected-error@+2 2{{'foo' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+      // expected-error@+1 {{'BBBB' requires 128 bit size 'BIGTY' (aka '__float128') type support, but device 'spir64' does not support it}}
+      auto A = foo(BBBB);
+    });
   });
 
-  kernel<class ok>([=]() {
-    // expected-note@+1 3{{used here}}
-    Z<__float128> S;
-    foo2<__float128>();
-    auto A = sizeof(CapturedToDevice);
+  deviceQueue.submit([&](sycl::handler &h) {
+    // expected-note@Inputs/sycl.hpp:212 {{called by 'kernel_single_task<ok, (lambda}}
+    h.single_task<class ok>([=]() {
+      // expected-note@+1 3{{used here}}
+      Z<__float128> S;
+      foo2<__float128>();
+      auto A = sizeof(CapturedToDevice);
+    });
   });
 
   return 0;
diff --git a/clang/test/SemaSYCL/forward-decl.cpp b/clang/test/SemaSYCL/forward-decl.cpp
index 7d252af5abc89..b0187a4821ab6 100644
--- a/clang/test/SemaSYCL/forward-decl.cpp
+++ b/clang/test/SemaSYCL/forward-decl.cpp
@@ -1,22 +1,27 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify -pedantic %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -fsyntax-only -verify -pedantic %s
 
 // expected-no-diagnostics
-template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
-  kernelFunc();
-}
+
+#include "sycl.hpp"
+
+sycl::queue deviceQueue;
 
 int main() {
-  kernel_single_task<class kernel_function>([]() {
+  deviceQueue.submit([&](sycl::handler &h) {
+    h.single_task<class kernel_function>([]() {
       class Foo *F;
       class Boo {
       public:
         virtual int getBoo() { return 42; }
       };
+    });
   });
 
-  kernel_single_task<class kernel_function_2>([]() {
+  deviceQueue.submit([&](sycl::handler &h) {
+    h.single_task<class kernel_function_2>([]() {
       class Boo *B;
+    });
   });
+
   return 0;
 }
diff --git a/clang/test/SemaSYCL/fpga_pipes.cpp b/clang/test/SemaSYCL/fpga_pipes.cpp
index b78b4baae903c..0e37af3adf60f 100644
--- a/clang/test/SemaSYCL/fpga_pipes.cpp
+++ b/clang/test/SemaSYCL/fpga_pipes.cpp
@@ -33,11 +33,6 @@ int Storage4 __attribute__((io_pipe_id(5)));
 template <int N>
 pipe_storage Storage5 __attribute__((io_pipe_id(N)));
 
-template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
-  kernelFunc();
-}
-
 void foo(pipe_storage PS) {}
 
 int main() {
diff --git a/clang/test/SemaSYCL/int128.cpp b/clang/test/SemaSYCL/int128.cpp
index 95c93537edefe..2eea7b8734715 100644
--- a/clang/test/SemaSYCL/int128.cpp
+++ b/clang/test/SemaSYCL/int128.cpp
@@ -1,5 +1,9 @@
 // RUN: %clang_cc1 -triple spir64 -aux-triple x86_64-unknown-linux-gnu \
-// RUN:    -fsycl -fsycl-is-device -verify -fsyntax-only %s
+// RUN:    -fsycl -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -verify -fsyntax-only %s
+
+#include "sycl.hpp"
+
+sycl::queue deviceQueue;
 
 typedef __uint128_t BIGTY;
 
@@ -71,50 +75,53 @@ int foobar() {
   return a;
 }
 
-template <typename Name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
-  // expected-note@+1 7{{called by 'kernel}}
-  kernelFunc();
-}
-
 int main() {
   // expected-note@+1 {{'CapturedToDevice' defined here}}
   __int128 CapturedToDevice = 1;
   host_ok();
-  kernel<class variables>([=]() {
-    // expected-error@+1 {{'__int128' is not supported on this target}}
-    decltype(CapturedToDevice) D;
-    // expected-error@+2 {{'__int128' is not supported on this target}}
-    // expected-error@+1 {{'CapturedToDevice' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
-    auto C = CapturedToDevice;
-    // expected-note@+1 3{{used here}}
-    Z<__int128> S;
-    // expected-error@+1 {{'field1' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
-    S.field1 += 1;
-    // expected-error@+1 {{'field' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
-    S.field = 1;
+  deviceQueue.submit([&](sycl::handler &h) {
+    // expected-note@Inputs/sycl.hpp:212 {{called by 'kernel_single_task<variables, (lambda}}
+    h.single_task<class variables>([=]() {
+      // expected-error@+1 {{'__int128' is not supported on this target}}
+      decltype(CapturedToDevice) D;
+      // expected-error@+2 {{'__int128' is not supported on this target}}
+      // expected-error@+1 {{'CapturedToDevice' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
+      auto C = CapturedToDevice;
+      // expected-note@+1 3{{used here}}
+      Z<__int128> S;
+      // expected-error@+1 {{'field1' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
+      S.field1 += 1;
+      // expected-error@+1 {{'field' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
+      S.field = 1;
+    });
   });
 
-  kernel<class functions>([=]() {
-    // expected-note@+1 2{{called by 'operator()'}}
-    usage();
-    // expected-error@+2 {{'unsigned __int128' is not supported on this target}}
-    // expected-note@+1 {{'BBBB' defined here}}
-    BIGTY BBBB;
-    // expected-error@+4 {{'__int128' is not supported on this target}}
-    // expected-error@+3 {{'BBBB' requires 128 bit size 'BIGTY' (aka 'unsigned __int128') type support, but device 'spir64' does not support it}}
-    // expected-error@+2 2{{'foo' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
-    // expected-note@+1 1{{called by 'operator()'}}
-    auto A = foo(BBBB);
-    // expected-note@+1 {{called by 'operator()'}}
-    auto i = foobar();
+  deviceQueue.submit([&](sycl::handler &h) {
+    // expected-note@Inputs/sycl.hpp:212 5{{called by 'kernel_single_task<functions, (lambda}}
+    h.single_task<class functions>([=]() {
+      // expected-note@+1 2{{called by 'operator()'}}
+      usage();
+      // expected-error@+2 {{'unsigned __int128' is not supported on this target}}
+      // expected-note@+1 {{'BBBB' defined here}}
+      BIGTY BBBB;
+      // expected-error@+4 {{'__int128' is not supported on this target}}
+      // expected-error@+3 {{'BBBB' requires 128 bit size 'BIGTY' (aka 'unsigned __int128') type support, but device 'spir64' does not support it}}
+      // expected-error@+2 2{{'foo' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
+      // expected-note@+1 1{{called by 'operator()'}}
+      auto A = foo(BBBB);
+      // expected-note@+1 {{called by 'operator()'}}
+      auto i = foobar();
+    });
   });
 
-  kernel<class ok>([=]() {
-    // expected-note@+1 3{{used here}}
-    Z<__int128> S;
-    foo2<__int128>();
-    auto A = sizeof(CapturedToDevice);
+  deviceQueue.submit([&](sycl::handler &h) {
+    // expected-note@Inputs/sycl.hpp:212 {{called by 'kernel_single_task<ok, (lambda}}
+    h.single_task<class ok>([=]() {
+      // expected-note@+1 3{{used here}}
+      Z<__int128> S;
+      foo2<__int128>();
+      auto A = sizeof(CapturedToDevice);
+    });
   });
 
   return 0;
diff --git a/clang/test/SemaSYCL/intel-fpga-local.cpp b/clang/test/SemaSYCL/intel-fpga-local.cpp
index 08f914f06a999..858fc1f0afdab 100644
--- a/clang/test/SemaSYCL/intel-fpga-local.cpp
+++ b/clang/test/SemaSYCL/intel-fpga-local.cpp
@@ -1,4 +1,8 @@
-// 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
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -Wno-return-type -fcxx-exceptions -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s
+
+#include "sycl.hpp"
+
+sycl::queue deviceQueue;
 
 //CHECK: FunctionDecl{{.*}}check_ast
 void check_ast()
@@ -847,19 +851,17 @@ struct templ_st {
   [[intel::force_pow2_depth(A)]] unsigned int templ_force_p2d_field[64];
 };
 
-template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
-  kernelFunc();
-}
-
 int main() {
-  kernel_single_task<class kernel_function>([]() {
-    check_ast();
-    diagnostics();
-    check_gnu_style();
-    //expected-note@+1{{in instantiation of function template specialization}}
-    check_template_parameters<2, 4, 8, -1, 1>();
-    struct templ_st<0> ts {};
+  deviceQueue.submit([&](sycl::handler &h) {
+    h.single_task<class kernel_function>([]() {
+      check_ast();
+      diagnostics();
+      check_gnu_style();
+      //expected-note@+1{{in instantiation of function template specialization}}
+      check_template_parameters<2, 4, 8, -1, 1>();
+      struct templ_st<0> ts {};
+    });
   });
+
   return 0;
 }
diff --git a/clang/test/SemaSYCL/intel-fpga-loops.cpp b/clang/test/SemaSYCL/intel-fpga-loops.cpp
index 0c98e1e8a8226..be182ef5158b8 100644
--- a/clang/test/SemaSYCL/intel-fpga-loops.cpp
+++ b/clang/test/SemaSYCL/intel-fpga-loops.cpp
@@ -1,4 +1,8 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify -pedantic %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -fsyntax-only -verify -pedantic %s
+
+#include "sycl.hpp"
+
+sycl::queue deviceQueue;
 
 // Test for Intel FPGA loop attributes applied not to a loop
 void foo() {
@@ -393,28 +397,26 @@ void max_concurrency_dependent() {
       a[i] = 0;
 }
 
-template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
-  kernelFunc();
-}
-
 int main() {
-  kernel_single_task<class kernel_function>([]() {
-    foo();
-    foo_deprecated();
-    boo();
-    goo();
-    zoo();
-    loop_attrs_compatibility();
-    ivdep_dependent<4, 2, 1>();
-    //expected-note@-1 +{{in instantiation of function template specialization}}
-    ivdep_dependent<2, 4, -1>();
-    //expected-note@-1 +{{in instantiation of function template specialization}}
-    ii_dependent<2, 4, -1>();
-    //expected-note@-1 +{{in instantiation of function template specialization}}
-    max_concurrency_dependent<1, 4, -2>();
-    //expected-note@-1 +{{in instantiation of function template specialization}}
+  deviceQueue.submit([&](sycl::handler &h) {
+    h.single_task<class kernel_function>([]() {
+      foo();
+      foo_deprecated();
+      boo();
+      goo();
+      zoo();
+      loop_attrs_compatibility();
+      ivdep_dependent<4, 2, 1>();
+      //expected-note@-1 +{{in instantiation of function template specialization}}
+      ivdep_dependent<2, 4, -1>();
+      //expected-note@-1 +{{in instantiation of function template specialization}}
+      ii_dependent<2, 4, -1>();
+      //expected-note@-1 +{{in instantiation of function template specialization}}
+      max_concurrency_dependent<1, 4, -2>();
+      //expected-note@-1 +{{in instantiation of function template specialization}}
+    });
   });
+
   return 0;
 }