diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 9320e56e69c6f..f2ed088e268a0 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -415,6 +415,10 @@ def SYCLSimdDocs : Documentation {
     The compiler may decide to compile such functions using different optimization
     and code generation pipeline. Also, this attribute is used to distinguish
     ESIMD private globals from regular SYCL global variables.
+
+    In SYCL 1.2.1 mode, the ``intel::sycl_explicit_simd`` attribute is propagated
+    from the function it is applied to onto the kernel which calls the function.
+    In SYCL 2020 mode, the attribute is not propagated to the kernel.
   }];
 }
 
@@ -2443,8 +2447,9 @@ lambda capture, or function object member, of the callable to which the
 attribute was applied. This effect is equivalent to annotating restrict on
 **all** kernel pointer arguments in an OpenCL or SPIR-V kernel.
 
-If ``intel::kernel_args_restrict`` is applied to a function called from a device
-kernel, the attribute is not ignored and it is propagated to the kernel.
+In SYCL 1.2.1 mode, the ``intel::kernel_args_restrict`` attribute is propagated
+from the function it is applied to onto the kernel which calls the function.
+In SYCL 2020 mode, the attribute is not propagated to the kernel.
 
 The attribute forms an unchecked assertion, in that implementations
 do not need to check/confirm the pre-condition in any way. If a user applies
@@ -2482,8 +2487,10 @@ def SYCLIntelNumSimdWorkItemsAttrDocs : Documentation {
   let Content = [{
 Applies to a device function/lambda function. Indicates the number of work
 items that should be processed in parallel. Valid values are positive integers.
-If ``intel::num_simd_work_items`` is applied to a function called from a
-device kernel, the attribute is not ignored and it is propagated to the kernel.
+
+In SYCL 1.2.1 mode, the ``intel::num_simd_work_items`` attribute is propagated
+from the function it is applied to onto the kernel which calls the function.
+In SYCL 2020 mode, the attribute is not propagated to the kernel.
 
 .. code-block:: c++
 
@@ -2656,6 +2663,11 @@ allows the Y and Z arguments to be optional. If not provided by the user, the
 value of Y and Z defaults to 1. See section 5.8.1 Kernel Attributes for more
 details.
 
+In SYCL 1.2.1 mode, the ``intel::reqd_work_group_size``,
+``cl::reqd_work_group_size``, and ``sycl::reqd_work_group_size`` attributes are
+propagated from the function they are applied to onto the kernel which calls the
+function. In SYCL 2020 mode, the attributes are not propagated to the kernel.
+
 .. code-block:: c++
 
   [[sycl::reqd_work_group_size(4, 4, 4)]] void foo() {}
@@ -2800,8 +2812,10 @@ Applies to a device function/lambda function. Indicates the maximum dimensions
 of a work group. Values must be positive integers. This is similar to
 reqd_work_group_size, but allows work groups that are smaller or equal to the
 specified sizes.
-If ``intel::max_work_group_size`` is applied to a function called from a
-device kernel, the attribute is not ignored and it is propagated to the kernel.
+
+In SYCL 1.2.1 mode, the ``intel::max_work_group_size`` attribute is propagated
+from the function it is applied to onto the kernel which calls the function.
+In SYCL 2020 mode, the attribute is not propagated to the kernel.
 
 .. code-block:: c++
 
@@ -2832,8 +2846,10 @@ Applies to a device function/lambda function or function call operator (of a
 function object). Indicates the largest valid global work dimension that will be
 accepted when running the kernel on a device. Valid values are integers in a
 range of [0, 3].
-If ``intel::max_global_work_dim`` is applied to a function called from a
-device kernel, the attribute is not ignored and it is propagated to the kernel.
+
+In SYCL 1.2.1 mode, the ``intel::max_global_work_dim`` attribute is propagated
+from the function it is applied to onto the kernel which calls the function.
+In SYCL 2020 mode, the attribute is not propagated to the kernel.
 
 .. code-block:: c++
 
@@ -2890,6 +2906,10 @@ device operation, guiding the FPGA backend to insert the appropriate number of
 registers to break-up the combinational logic circuit, and thereby controlling
 the length of the longest combinational path.
 
+In SYCL 1.2.1 mode, the ``intel::scheduler_target_fmax_mhz`` attribute is
+propagated from the function it is applied to onto the kernel which calls the
+function. In SYCL 2020 mode, the attribute is not propagated to the kernel.
+
 .. code-block:: c++
 
   [[intel::scheduler_target_fmax_mhz(4)]] void foo() {}
@@ -2920,6 +2940,10 @@ function object). If 1, compiler doesn't use the global work offset values for
 the device function. Valid values are 0 and 1. If used without argument, value
 of 1 is set implicitly.
 
+In SYCL 1.2.1 mode, the ``intel::no_global_work_offset`` attribute is
+propagated from the function it is applied to onto the kernel which calls the
+function. In SYCL 2020 mode, the attribute is not propagated to the kernel.
+
 .. code-block:: c++
 
   [[intel::no_global_work_offset]]
@@ -4607,6 +4631,10 @@ the ``[[intel::named_sub_group_size(NAME)]]`` documentation for clarification.
 This attribute is mutually exclusive with ``[[intel::named_sub_group_size(NAME)]]``
 and ``[[intel::sycl_explicit_simd]]``.
 
+In SYCL 1.2.1 mode, the ``intel::reqd_sub_group_size`` attribute is propagated
+from the function it is applied to onto the kernel which calls the function.
+In SYCL 2020 mode, the attribute is not propagated to the kernel.
+
 In addition to device functions, the required sub-group size attribute may also
 be specified in the definition of a named functor object and lambda functions,
 as in the examples below:
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 5126475f8aa4c..acaefa9f91a36 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -558,15 +558,21 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD,
   if (!FD->hasAttrs())
     return;
 
-  llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) {
-    // FIXME: Make this list self-adapt as new SYCL attributes are added.
-    return isa<IntelReqdSubGroupSizeAttr, IntelNamedSubGroupSizeAttr,
-               ReqdWorkGroupSizeAttr, SYCLIntelKernelArgsRestrictAttr,
-               SYCLIntelNumSimdWorkItemsAttr,
-               SYCLIntelSchedulerTargetFmaxMhzAttr,
-               SYCLIntelMaxWorkGroupSizeAttr, SYCLIntelMaxGlobalWorkDimAttr,
-               SYCLIntelNoGlobalWorkOffsetAttr, SYCLSimdAttr>(A);
-  });
+  // In SYCL 1.2.1 mode, the attributes are propagated from the function they
+  // are applied to onto the kernel which calls the function.
+  // In SYCL 2020 mode, the attributes are not propagated to the kernel.
+  if (DirectlyCalled || S.getASTContext().getLangOpts().getSYCLVersion() <
+                            LangOptions::SYCL_2020) {
+    llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) {
+      // FIXME: Make this list self-adapt as new SYCL attributes are added.
+      return isa<IntelReqdSubGroupSizeAttr, IntelNamedSubGroupSizeAttr,
+                 ReqdWorkGroupSizeAttr, SYCLIntelKernelArgsRestrictAttr,
+                 SYCLIntelNumSimdWorkItemsAttr,
+                 SYCLIntelSchedulerTargetFmaxMhzAttr,
+                 SYCLIntelMaxWorkGroupSizeAttr, SYCLIntelMaxGlobalWorkDimAttr,
+                 SYCLIntelNoGlobalWorkOffsetAttr, SYCLSimdAttr>(A);
+    });
+  }
 
   // Attributes that should not be propagated from device functions to a kernel.
   if (DirectlyCalled) {
diff --git a/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp b/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp
new file mode 100644
index 0000000000000..a04610d556360
--- /dev/null
+++ b/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp
@@ -0,0 +1,335 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s
+
+// Tests for IR of [[intel::scheduler_target_fmax_mhz()]], [[intel::num_simd_work_items()]],
+// [[intel::no_global_work_offset()]], [[intel::max_global_work_dim()]], [[sycl::reqd_sub_group_size()]],
+// [[sycl::reqd_work_group_size()]], [[intel::kernel_args_restrict]], [[intel::max_work_group_size()]],
+// and [[intel::sycl_explicit_simd]] function attributes in SYCL 2020.
+
+#include "sycl.hpp"
+
+using namespace cl::sycl;
+queue q;
+
+class Foo {
+public:
+  [[intel::scheduler_target_fmax_mhz(1)]] void operator()() const {}
+};
+
+template <int SIZE>
+class Functor {
+public:
+  [[intel::scheduler_target_fmax_mhz(SIZE)]] void operator()() const {}
+};
+
+[[intel::scheduler_target_fmax_mhz(5)]] void foo() {}
+
+class Foo1 {
+public:
+  [[intel::num_simd_work_items(1)]] void operator()() const {}
+};
+
+template <int SIZE>
+class Functor1 {
+public:
+  [[intel::num_simd_work_items(SIZE)]] void operator()() const {}
+};
+
+[[intel::num_simd_work_items(5)]] void foo1() {}
+
+class Foo2 {
+public:
+  [[intel::no_global_work_offset(1)]] void operator()() const {}
+};
+
+template <int SIZE>
+class Functor2 {
+public:
+  [[intel::no_global_work_offset(SIZE)]] void operator()() const {}
+};
+
+[[intel::no_global_work_offset(0)]] void foo2() {}
+
+class Foo3 {
+public:
+  [[intel::max_global_work_dim(1)]] void operator()() const {}
+};
+
+template <int SIZE>
+class Functor3 {
+public:
+  [[intel::max_global_work_dim(SIZE)]] void operator()() const {}
+};
+
+[[intel::max_global_work_dim(1)]] void foo3() {}
+
+
+class Foo4 {
+public:
+  [[sycl::reqd_sub_group_size(16)]] void operator()() const {}
+};
+
+[[sycl::reqd_sub_group_size(8)]] void foo4() {}
+
+class Functor4 {
+public:
+  void operator()() const {
+    foo4();
+  }
+};
+
+template <int SIZE>
+class Functor5 {
+public:
+  [[sycl::reqd_sub_group_size(SIZE)]] void operator()() const {}
+};
+
+class Foo5 {
+public:
+  [[sycl::reqd_work_group_size(32, 16, 16)]] void operator()() const {}
+};
+
+[[sycl::reqd_work_group_size(8, 1, 1)]] void foo5() {}
+
+class Functor6 {
+public:
+  void operator()() const {
+    foo5();
+  }
+};
+
+template <int SIZE, int SIZE1, int SIZE2>
+class Functor7 {
+public:
+  [[sycl::reqd_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() const {}
+};
+
+class Foo6 {
+public:
+  [[intel::max_work_group_size(32, 16, 16)]] void operator()() const {}
+};
+
+[[intel::max_work_group_size(8, 1, 1)]] void foo6() {}
+
+class Functor8 {
+public:
+  void operator()() const {
+    foo6();
+  }
+};
+
+template <int SIZE, int SIZE1, int SIZE2>
+class Functor9 {
+public:
+  [[intel::max_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() const {}
+};
+
+class Foo7 {
+public:
+  [[intel::sycl_explicit_simd]] void operator()() const {}
+};
+
+[[intel::sycl_explicit_simd]] void foo7() {}
+
+class Foo8 {
+public:
+  [[intel::kernel_args_restrict]] void operator()() const {}
+};
+
+[[intel::kernel_args_restrict]] void foo8() {}
+
+class Functor10 {
+public:
+  void operator()() const {
+    foo8();
+  }
+};
+
+int main() {
+  q.submit([&](handler &h) {
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 !kernel_arg_buffer_location ![[NUM:[0-9]+]] !scheduler_target_fmax_mhz ![[NUM1:[0-9]+]]
+    Foo boo;
+    h.single_task<class kernel_name1>(boo);
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 !kernel_arg_buffer_location ![[NUM]] !scheduler_target_fmax_mhz ![[NUM42:[0-9]+]]
+    h.single_task<class kernel_name2>(
+        []() [[intel::scheduler_target_fmax_mhz(42)]]{});
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 !kernel_arg_buffer_location ![[NUM]] !scheduler_target_fmax_mhz ![[NUM2:[0-9]+]]
+    Functor<2> f;
+    h.single_task<class kernel_name3>(f);
+
+    // Test attribute is not propagated.
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 !kernel_arg_buffer_location ![[NUM]]
+    // CHECK-NOT: !scheduler_target_fmax_mhz
+    // CHECK-SAME: {
+    // CHECK: define dso_local spir_func void @_Z3foov()
+    h.single_task<class kernel_name4>(
+        []() { foo(); });
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 !kernel_arg_buffer_location ![[NUM]] !num_simd_work_items ![[NUM1]]
+    Foo1 boo1;
+    h.single_task<class kernel_name5>(boo1);
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6() #0 !kernel_arg_buffer_location ![[NUM]] !num_simd_work_items ![[NUM42]]
+    h.single_task<class kernel_name6>(
+        []() [[intel::num_simd_work_items(42)]]{});
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name7() #0 !kernel_arg_buffer_location ![[NUM]] !num_simd_work_items ![[NUM2]]
+    Functor1<2> f1;
+    h.single_task<class kernel_name7>(f1);
+
+    // Test attribute is not propagated.
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name8() #0 !kernel_arg_buffer_location ![[NUM]]
+    // CHECK-NOT: !num_simd_work_items
+    // CHECK-SAME: {
+    // CHECK: define dso_local spir_func void @_Z4foo1v()
+    h.single_task<class kernel_name8>(
+        []() { foo1(); });
+    
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name9() #0 !kernel_arg_buffer_location ![[NUM]] !no_global_work_offset ![[NUM:[0-9]+]]
+    Foo2 boo2;
+    h.single_task<class kernel_name9>(boo2);
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name10() #0 {{.*}} ![[NUM0:[0-9]+]]
+    h.single_task<class kernel_name10>(
+        []() [[intel::no_global_work_offset(0)]]{});
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name11() #0 !kernel_arg_buffer_location ![[NUM]] !no_global_work_offset ![[NUM]]
+    Functor2<1> f2;
+    h.single_task<class kernel_name11>(f2);
+
+    // Test attribute is not propagated.
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name12() #0 !kernel_arg_buffer_location ![[NUM]]
+    // CHECK-NOT: !no_global_work_offset
+    // CHECK-SAME: {
+    // CHECK: define dso_local spir_func void @_Z4foo2v()
+    h.single_task<class kernel_name12>(
+        []() { foo2(); });
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name13() #0 !kernel_arg_buffer_location ![[NUM]] !max_global_work_dim ![[NUM1]]
+    Foo3 boo3;
+    h.single_task<class kernel_name13>(boo3);
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name14() #0 !kernel_arg_buffer_location ![[NUM]] !max_global_work_dim ![[NUM1]]
+    h.single_task<class kernel_name14>(
+        []() [[intel::max_global_work_dim(1)]]{});
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name15() #0 !kernel_arg_buffer_location ![[NUM]] !max_global_work_dim ![[NUM2]]
+    Functor3<2> f3;
+    h.single_task<class kernel_name15>(f3);
+
+    // Test attribute is not propagated.
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name16() #0 !kernel_arg_buffer_location ![[NUM]]
+    // CHECK-NOT: !max_global_work_dim
+    // CHECK-SAME: {
+    // CHECK: define dso_local spir_func void @_Z4foo3v()
+    h.single_task<class kernel_name16>(
+        []() { foo3(); });
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name17() #0 !kernel_arg_buffer_location ![[NUM]] !intel_reqd_sub_group_size ![[NUM16:[0-9]+]]
+    Foo4 boo4;
+    h.single_task<class kernel_name17>(boo4);
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name18() #0 !kernel_arg_buffer_location ![[NUM]] !intel_reqd_sub_group_size ![[NUM1]]
+    h.single_task<class kernel_name18>(
+        []() [[sycl::reqd_sub_group_size(1)]]{});
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name19() #0 !kernel_arg_buffer_location ![[NUM]] !intel_reqd_sub_group_size ![[NUM2]]
+    Functor5<2> f5;
+    h.single_task<class kernel_name19>(f5);
+
+    // Test attribute is not propagated.
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name20() #0 !kernel_arg_buffer_location ![[NUM]]
+    // CHECK-NOT: !reqd_sub_group_size
+    // CHECK-SAME: {
+    // CHECK: define dso_local spir_func void @_Z4foo4v()
+    Functor4 f4;
+    h.single_task<class kernel_name20>(f4);
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name21() #0 !kernel_arg_buffer_location ![[NUM]] !reqd_work_group_size ![[NUM32:[0-9]+]]
+    Foo5 boo5;
+    h.single_task<class kernel_name21>(boo5);
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name22() #0 !kernel_arg_buffer_location ![[NUM]] !reqd_work_group_size ![[NUM88:[0-9]+]]
+    h.single_task<class kernel_name22>(
+        []() [[sycl::reqd_work_group_size(8, 8, 8)]]{});
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name23() #0 !kernel_arg_buffer_location ![[NUM]] !reqd_work_group_size ![[NUM22:[0-9]+]]
+    Functor7<2, 2, 2> f7;
+    h.single_task<class kernel_name23>(f7);
+
+    // Test attribute is not propagated.
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name24() #0 !kernel_arg_buffer_location ![[NUM]]
+    // CHECK-NOT: !reqd_work_group_size
+    // CHECK-SAME: {
+    // CHECK: define dso_local spir_func void @_Z4foo5v()
+    Functor6 f6;
+    h.single_task<class kernel_name24>(f6);
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name25() #0 !kernel_arg_buffer_location ![[NUM]] !max_work_group_size ![[NUM32]]
+    Foo6 boo6;
+    h.single_task<class kernel_name25>(boo6);
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name26() #0 !kernel_arg_buffer_location ![[NUM]] !max_work_group_size ![[NUM88]]
+    h.single_task<class kernel_name26>(
+        []() [[intel::max_work_group_size(8, 8, 8)]]{});
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name27() #0 !kernel_arg_buffer_location ![[NUM]] !max_work_group_size ![[NUM22]]
+    Functor9<2, 2, 2> f9;
+    h.single_task<class kernel_name27>(f9);
+
+    // Test attribute is not propagated.
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name28() #0 !kernel_arg_buffer_location ![[NUM]]
+    // CHECK-NOT: !max_work_group_size
+    // CHECK-SAME: {
+    // CHECK: define dso_local spir_func void @_Z4foo6v()
+    Functor8 f8;
+    h.single_task<class kernel_name28>(f8);
+
+    // Test attribute is not propagated.
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name29() #0 !kernel_arg_buffer_location ![[NUM]]
+    // CHECK-NOT: !sycl_explicit_simd
+    // CHECK-SAME: {
+    // CHECK: define {{.*}}spir_func void @{{.*}}foo7{{.*}} !sycl_explicit_simd ![[NUM]]
+    h.single_task<class kernel_name29>(
+        []() { foo7(); });
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name30() #0 !intel_reqd_sub_group_size ![[NUM1]] !sycl_explicit_simd ![[NUM]]
+    Foo7 boo7;
+    h.single_task<class kernel_name30>(boo7);
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name31() #0 !intel_reqd_sub_group_size ![[NUM1]] !sycl_explicit_simd ![[NUM]]
+    h.single_task<class kernel_name31>(
+        []() [[intel::sycl_explicit_simd]]{});
+
+    // Test attribute is not propagated.
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name32() #0 !kernel_arg_buffer_location ![[NUM]]
+    // CHECK: define {{.*}}spir_func void @{{.*}}Functor10{{.*}}(%class._ZTS9Functor10.Functor10 addrspace(4)* align 1 dereferenceable_or_null(1) %this) #3 comdat align 2
+    // CHECK-NOT: noalias
+    // CHECK-SAME: {
+    // CHECK: define dso_local spir_func void @_Z4foo8v()
+    Functor10 f10;
+    h.single_task<class kernel_name32>(f10);
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name33() #0 !kernel_arg_buffer_location ![[NUM]]
+    // CHECK: define {{.*}}spir_func void @{{.*}}Foo8{{.*}}(%class.{{.*}}Foo8.Foo8 addrspace(4)* noalias align 1 dereferenceable_or_null(1) %this) #3 comdat align 2
+    Foo8 boo8;
+    h.single_task<class kernel_name33>(boo8);
+
+    // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name34() #0 !kernel_arg_buffer_location ![[NUM]]
+    // CHECK: define {{.*}}spir_func void @{{.*}}(%class.{{.*}}.anon addrspace(4)* noalias align 1 dereferenceable_or_null(1) %this) #4 align 2
+    h.single_task<class kernel_name34>(
+        []() [[intel::kernel_args_restrict]]{});
+  });
+  return 0;
+}
+
+// CHECK: ![[NUM]] = !{}
+// CHECK: ![[NUM1]] = !{i32 1}
+// CHECK: ![[NUM42]] = !{i32 42}
+// CHECK: ![[NUM2]] = !{i32 2}
+// CHECK-NOT: ![[NUM0]]  = !{i32 0}
+// CHECK: ![[NUM16]] = !{i32 16}
+// CHECK: ![[NUM32]] = !{i32 16, i32 16, i32 32}
+// CHECK: ![[NUM88]] = !{i32 8, i32 8, i32 8}
+// CHECK: ![[NUM22]] = !{i32 2, i32 2, i32 2}
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 d85295bc53178..d63d6b076fd81 100644
--- a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp
+++ b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s
 
 #include "sycl.hpp"
 
diff --git a/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp b/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp
index 53ab36ec9ab92..e3be51dd70b96 100644
--- a/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp
+++ b/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s
 
 #include "sycl.hpp"
 
diff --git a/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp b/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp
index 28c0bae228871..254a9175db744 100644
--- a/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp
+++ b/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s
 
 #include "sycl.hpp"
 
diff --git a/clang/test/CodeGenSYCL/intel-restrict.cpp b/clang/test/CodeGenSYCL/intel-restrict.cpp
index 20aa089e35f26..343616ccc2d54 100644
--- a/clang/test/CodeGenSYCL/intel-restrict.cpp
+++ b/clang/test/CodeGenSYCL/intel-restrict.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device %s -emit-llvm -triple spir64-unknown-unknown-sycldevice -o - | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device %s -emit-llvm -sycl-std=2017 -triple spir64-unknown-unknown-sycldevice -o - | FileCheck %s
 
 template <typename name, typename Func>
 __attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
diff --git a/clang/test/CodeGenSYCL/num-simd-work-items.cpp b/clang/test/CodeGenSYCL/num-simd-work-items.cpp
index eebce1f408de4..56c7bd511e346 100644
--- a/clang/test/CodeGenSYCL/num-simd-work-items.cpp
+++ b/clang/test/CodeGenSYCL/num-simd-work-items.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s
 
 #include "sycl.hpp"
 
diff --git a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp
index 128a69052118e..01957696603e3 100644
--- a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp
+++ b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -triple spir64-unknown-unknown-sycldevice -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -sycl-std=2017 -triple spir64-unknown-unknown-sycldevice -emit-llvm -o - %s | FileCheck %s
 
 #include "sycl.hpp"
 
diff --git a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp
index 0986e89738514..d9e79bd4c31af 100644
--- a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp
+++ b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s
 
 #include "sycl.hpp"
 
diff --git a/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp b/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp
index 55cf1f176e9fe..1357347f2208b 100644
--- a/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp
+++ b/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
 
 #include "sycl.hpp"
 
diff --git a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp
index fd4e014d5825a..c316e611a18a8 100644
--- a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp
+++ b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
 
 #include "sycl.hpp"
 
diff --git a/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp b/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp
new file mode 100644
index 0000000000000..45fe936643b33
--- /dev/null
+++ b/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp
@@ -0,0 +1,290 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump %s | FileCheck %s
+
+// Tests for AST of [[intel::scheduler_target_fmax_mhz()]], [[intel::num_simd_work_items()]],
+// [[intel::no_global_work_offset()]], [[intel::max_global_work_dim()]], [[intel::sycl_explicit_simd]],
+// [[sycl::reqd_sub_group_size()]], [[sycl::reqd_work_group_size()]], [[intel::kernel_args_restrict]], and
+// [[intel::max_work_group_size()]] function attributes in SYCL 2020.
+
+#include "sycl.hpp"
+
+sycl::queue deviceQueue;
+
+struct FuncObj {
+  [[intel::sycl_explicit_simd]] void operator()() const {}
+};
+
+struct FuncObj1 {
+  [[intel::no_global_work_offset(1)]] void operator()() const {}
+};
+
+struct FuncObj2 {
+  [[intel::scheduler_target_fmax_mhz(10)]] void operator()() const {}
+};
+
+struct FuncObj3 {
+  [[intel::max_work_group_size(2, 2, 2)]] void operator()() const {}
+};
+
+struct FuncObj4 {
+  [[sycl::reqd_work_group_size(2, 2, 2)]] void operator()() const {}
+};
+
+struct FuncObj5 {
+  [[intel::num_simd_work_items(8)]] void operator()() const {}
+};
+
+struct FuncObj6 {
+  [[intel::kernel_args_restrict]] void operator()() const {}
+};
+
+struct FuncObj7 {
+  [[intel::max_global_work_dim(1)]] void operator()() const {}
+};
+
+[[intel::sycl_explicit_simd]] void func() {}
+
+[[intel::no_global_work_offset(1)]] void func1() {}
+
+[[intel::scheduler_target_fmax_mhz(2)]] void func2() {}
+
+[[intel::max_work_group_size(1, 1, 1)]] void func3() {}
+
+[[sycl::reqd_work_group_size(1, 1, 1)]] void func4() {}
+
+[[intel::num_simd_work_items(5)]] void func5() {}
+
+[[intel::kernel_args_restrict]] void func6() {}
+
+[[intel::max_global_work_dim(0)]] void func7() {}
+
+[[sycl::reqd_sub_group_size(4)]] void func8() {}
+
+class Functor {
+public:
+  void operator()() const {
+    func8();
+  }
+};
+
+class Functor1 {
+public:
+  [[sycl::reqd_sub_group_size(12)]] void operator()() const {}
+};
+
+int main() {
+  deviceQueue.submit([&](sycl::handler &h) {
+    // CHECK:       FunctionDecl {{.*}}test_kernel1
+    // CHECK:       SYCLSimdAttr {{.*}} Implicit
+    // CHECK-NEXT:  SYCLKernelAttr {{.*}} Implicit
+    // CHECK-NEXT:  SYCLSimdAttr
+    h.single_task<class test_kernel1>(
+        FuncObj());
+    // CHECK:       FunctionDecl {{.*}}test_kernel2
+    // CHECK:       SYCLSimdAttr {{.*}} Implicit
+    // CHECK-NEXT:  SYCLKernelAttr {{.*}} Implicit
+    // CHECK-NEXT:  SYCLSimdAttr
+    h.single_task<class test_kernel2>(
+        []() [[intel::sycl_explicit_simd]]{});
+
+    // Test attribute is not propagated.
+    // CHECK:      FunctionDecl {{.*}}test_kernel3
+    // CHECK:      SYCLSimdAttr {{.*}} Implicit
+    // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit
+    // CHECK-NEXT: SYCLSimdAttr
+    // CHECK-NOT:  SYCLSimdAttr
+    h.single_task<class test_kernel3>(
+        []() [[intel::sycl_explicit_simd]] { func(); });
+
+    // Test attribute is not propagated.
+    // CHECK:      FunctionDecl {{.*}}test_kernel4
+    // CHECK-NOT:  SYCLIntelNoGlobalWorkOffsetAttr
+    h.single_task<class test_kernel4>(
+        []() { func1(); });
+
+    // CHECK:      FunctionDecl {{.*}}test_kernel5
+    // CHECK:      SYCLIntelNoGlobalWorkOffsetAttr
+    // CHECK-NEXT: ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT: value: Int 1
+    // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1
+    h.single_task<class test_kernel5>(
+        FuncObj1());
+
+    // CHECK:       FunctionDecl {{.*}}test_kernel6
+    // CHECK:       SYCLIntelNoGlobalWorkOffsetAttr
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 1
+    h.single_task<class test_kernel6>(
+        []() [[intel::no_global_work_offset]]{});
+
+    // CHECK:       FunctionDecl {{.*}}test_kernel7
+    // CHECK:       SYCLIntelSchedulerTargetFmaxMhzAttr
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 10
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 10
+    h.single_task<class test_kernel7>(
+        FuncObj2());
+
+    // CHECK:       FunctionDecl {{.*}}test_kernel8
+    // CHECK:       SYCLIntelSchedulerTargetFmaxMhzAttr
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 20
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 20
+    h.single_task<class test_kernel8>(
+        []() [[intel::scheduler_target_fmax_mhz(20)]]{});
+
+    // Test attribute is not propagated.
+    // CHECK:      FunctionDecl {{.*}}test_kernel9
+    // CHECK-NOT:  SYCLIntelSchedulerTargetFmaxMhzAttr
+    h.single_task<class test_kernel9>(
+        []() { func2(); });
+
+    // CHECK:       FunctionDecl {{.*}}test_kernel10
+    // CHECK:       SYCLIntelMaxWorkGroupSizeAttr
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 2
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 2
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 2
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 2
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 2
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 2
+    h.single_task<class test_kernel10>(
+        FuncObj3());
+
+    // Test attribute is not propagated.
+    // CHECK:     FunctionDecl {{.*}}test_kernel11
+    // CHECK-NOT: SYCLIntelMaxWorkGroupSizeAttr
+    h.single_task<class test_kernel11>(
+        []() { func3(); });
+
+    // CHECK:       FunctionDecl {{.*}}test_kernel12
+    // CHECK:       SYCLIntelMaxWorkGroupSizeAttr
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 8
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 8
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 8
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 8
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 8
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 8
+    h.single_task<class test_kernel12>(
+        []() [[intel::max_work_group_size(8, 8, 8)]]{});
+
+    // CHECK:       FunctionDecl {{.*}}test_kernel13
+    // CHECK:       ReqdWorkGroupSizeAttr
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 2
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 2
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 2
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 2
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 2
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 2
+    h.single_task<class test_kernel13>(
+        FuncObj4());
+
+    // Test attribute is not propagated.
+    // CHECK:      FunctionDecl {{.*}}test_kernel14
+    // CHECK-NOT:  ReqdWorkGroupSizeAttr
+    h.single_task<class test_kernel14>(
+        []() { func4(); });
+
+    // CHECK:       FunctionDecl {{.*}}test_kernel15
+    // CHECK:       ReqdWorkGroupSizeAttr
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 8
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 8
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 8
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 8
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 8
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 8
+    h.single_task<class test_kernel15>(
+        []() [[sycl::reqd_work_group_size(8, 8, 8)]]{});
+
+    // CHECK:       FunctionDecl {{.*}}test_kernel16
+    // CHECK:       SYCLIntelNumSimdWorkItemsAttr
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 8
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 8
+    h.single_task<class test_kernel16>(
+        FuncObj5());
+
+    // CHECK:       FunctionDecl {{.*}}test_kernel17
+    // CHECK:       SYCLIntelNumSimdWorkItemsAttr
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 20
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 20
+    h.single_task<class test_kernel17>(
+        []() [[intel::num_simd_work_items(20)]]{});
+
+    // Test attribute is not propagated.
+    // CHECK:     FunctionDecl {{.*}}test_kernel18
+    // CHECK-NOT: SYCLIntelNumSimdWorkItemsAttr
+    h.single_task<class test_kernel18>(
+        []() { func5(); });
+
+    // CHECK: FunctionDecl {{.*}}test_kernel19
+    // CHECK: SYCLIntelKernelArgsRestrictAttr
+    h.single_task<class test_kernel19>(
+        FuncObj6());
+
+    // CHECK: FunctionDecl {{.*}}test_kernel20
+    // CHECK: SYCLIntelKernelArgsRestrictAttr
+    h.single_task<class test_kernel20>(
+        []() [[intel::kernel_args_restrict]]{});
+
+    // Test attribute is not propagated.
+    // CHECK:     FunctionDecl {{.*}}test_kernel21
+    // CHECK-NOT: SYCLIntelKernelArgsRestrictAttr
+    h.single_task<class test_kernel21>(
+        []() { func6(); });
+
+    // CHECK: FunctionDecl {{.*}}test_kernel22
+    // CHECK: SYCLIntelMaxGlobalWorkDimAttr
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 1
+    h.single_task<class test_kernel22>(
+        FuncObj7());
+
+    // CHECK: FunctionDecl {{.*}}test_kernel23
+    // CHECK: SYCLIntelMaxGlobalWorkDimAttr
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 0
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 0
+    h.single_task<class test_kernel23>(
+        []() [[intel::max_global_work_dim(0)]]{});
+
+    // Test attribute is not propagated.
+    // CHECK:     FunctionDecl {{.*}}test_kernel24
+    // CHECK-NOT: SYCLIntelMaxGlobalWorkDimAttr
+    h.single_task<class test_kernel24>(
+        []() { func7(); });
+
+    // Test attribute is not propagated.
+    // CHECK:     FunctionDecl {{.*}}test_kernel25
+    // CHECK-NOT: IntelReqdSubGroupSizeAttr
+    Functor f;
+    h.single_task<class test_kernel25>(f);
+
+    // CHECK:     FunctionDecl {{.*}}test_kernel26
+    // CHECK:     IntelReqdSubGroupSizeAttr
+    Functor1 f1;
+    h.single_task<class test_kernel26>(f1);
+
+    // CHECK:       FunctionDecl {{.*}}test_kernel27
+    // CHECK:       IntelReqdSubGroupSizeAttr
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 8
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 8
+    h.single_task<class test_kernel27>(
+        []() [[sycl::reqd_sub_group_size(8)]]{});
+  });
+  return 0;
+}
diff --git a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp b/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp
index a7bab5a2f43c1..bdabd1f40d52d 100644
--- a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp
+++ b/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp
@@ -1,6 +1,6 @@
-// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -triple spir64 -verify
-// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -triple spir64 -DTRIGGER_ERROR -verify
-// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -triple spir64 | FileCheck %s
+// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 -verify
+// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 -DTRIGGER_ERROR -verify
+// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -sycl-std=2017 -triple spir64 | FileCheck %s
 
 #ifndef TRIGGER_ERROR
 [[intel::no_global_work_offset]] void not_direct_one() {} // expected-no-diagnostics
diff --git a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp
index b3c2d4b8e40c2..d85ca9ddc1f96 100644
--- a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp
+++ b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-return-type -Wno-sycl-2017-compat -fcxx-exceptions -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-return-type -sycl-std=2017 -Wno-sycl-2017-compat -fcxx-exceptions -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s
 
 #include "sycl.hpp"
 
diff --git a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
index a93096acec29e..507001a889432 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -triple spir64 -DTRIGGER_ERROR -verify
-// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -triple spir64 | FileCheck %s
+// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -Wno-sycl-2017-compat -triple spir64 -DTRIGGER_ERROR -verify
+// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 | FileCheck %s
 
 #include "sycl.hpp"
 
diff --git a/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp b/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp
index f190e5920ee38..906d35c3eaac6 100644
--- a/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -triple spir64 -DTRIGGER_ERROR -verify
-// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -triple spir64 | FileCheck %s
+// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -triple spir64 -DTRIGGER_ERROR -verify
+// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple spir64 | FileCheck %s
 
 #include "sycl.hpp"
 
diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp
new file mode 100644
index 0000000000000..8909245dd2b19
--- /dev/null
+++ b/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp
@@ -0,0 +1,145 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -fsyntax-only -verify -DTRIGGER_ERROR %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump %s | FileCheck %s
+
+// Test for AST of reqd_work_group_size kernel attribute in SYCL 2020.
+
+#include "sycl.hpp"
+
+using namespace cl::sycl;
+queue q;
+
+#ifndef __SYCL_DEVICE_ONLY__
+// expected-no-diagnostics
+class Functor {
+public:
+  [[sycl::reqd_work_group_size(4)]] void operator()() const {}
+};
+
+void bar() {
+  q.submit([&](handler &h) {
+    Functor f;
+    h.single_task<class kernel_name>(f);
+  });
+}
+
+#else
+#ifdef TRIGGER_ERROR
+class Functor32 {
+public:
+  [[cl::reqd_work_group_size(32)]] void operator()() const {} // expected-error {{'reqd_work_group_size' attribute requires exactly 3 arguments}} \
+                                                              // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \
+                                                              // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}}
+};
+#endif // TRIGGER_ERROR
+
+class Functor33 {
+public:
+  // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}}
+  [[sycl::reqd_work_group_size(32, -4)]] void operator()() const {}
+};
+
+class Functor30 {
+public:
+  // expected-warning@+1 2{{implicit conversion changes signedness: 'int' to 'unsigned long long'}}
+  [[sycl::reqd_work_group_size(30, -30, -30)]] void operator()() const {}
+};
+
+class Functor16 {
+public:
+  [[sycl::reqd_work_group_size(16)]] void operator()() const {}
+};
+
+class Functor64 {
+public:
+  [[sycl::reqd_work_group_size(64, 64)]] void operator()() const {}
+};
+
+class Functor16x16x16 {
+public:
+  [[intel::reqd_work_group_size(16, 16, 16)]] void operator()() const {} // expected-warning {{attribute 'intel::reqd_work_group_size' is deprecated}} \
+                                                                         // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}}
+};
+
+class FunctorAttr {
+public:
+  __attribute__((reqd_work_group_size(128, 128, 128))) void operator()() const {} // expected-warning {{attribute 'reqd_work_group_size' is deprecated}} \
+                                                                                  // expected-note {{did you mean to use '[[sycl::reqd_work_group_size]]' instead?}}
+};
+
+int main() {
+  q.submit([&](handler &h) {
+    // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1
+    // CHECK: ReqdWorkGroupSizeAttr
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 16
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 16
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 1
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 1
+    Functor16 f16;
+    h.single_task<class kernel_name1>(f16);
+
+    // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2
+    // CHECK: ReqdWorkGroupSizeAttr
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 16
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 16
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 16
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 16
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 16
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 16
+    Functor16x16x16 f16x16x16;
+    h.single_task<class kernel_name2>(f16x16x16);
+
+    // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name3
+    // CHECK: ReqdWorkGroupSizeAttr
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 128
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 128
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 128
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 128
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 128
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 128
+    FunctorAttr fattr;
+    h.single_task<class kernel_name3>(fattr);
+
+    // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name4
+    // CHECK: ReqdWorkGroupSizeAttr
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 32
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 32
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int -4
+    // CHECK-NEXT:  UnaryOperator{{.*}} 'int' prefix '-'
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 4
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 1
+    Functor33 f33;
+    h.single_task<class kernel_name4>(f33);
+
+    // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5
+    // CHECK: ReqdWorkGroupSizeAttr
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 30
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 30
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int -30
+    // CHECK-NEXT:  UnaryOperator{{.*}} 'int' prefix '-'
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 30
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int -30
+    // CHECK-NEXT:  UnaryOperator{{.*}} 'int' prefix '-'
+    Functor30 f30;
+    h.single_task<class kernel_name5>(f30);
+  });
+  return 0;
+}
+#endif // __SYCL_DEVICE_ONLY__
diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp
index 52c9cdfe9a6f5..68d299d4d3a4f 100644
--- a/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp
+++ b/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp
@@ -1,5 +1,7 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -fsyntax-only -verify -DTRIGGER_ERROR %s
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -fsyntax-only -verify -DTRIGGER_ERROR %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s
+
+// Test for AST of reqd_work_group_size kernel attribute in SYCL 1.2.1.
 
 #include "sycl.hpp"
 
@@ -30,15 +32,6 @@ void bar() {
 [[sycl::reqd_work_group_size(32, 32)]] void f32x32x1() {}      // expected-note {{conflicting attribute is here}}
 [[sycl::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}}
 
-#ifdef TRIGGER_ERROR
-class Functor32 {
-public:
-  [[cl::reqd_work_group_size(32)]] void operator()() const {} // expected-error {{'reqd_work_group_size' attribute requires exactly 3 arguments}} \
-                                                              // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \
-                                                              // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}}
-};
-#endif // TRIGGER_ERROR
-
 class Functor33 {
 public:
   // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}}
@@ -63,8 +56,7 @@ class Functor64 {
 
 class Functor16x16x16 {
 public:
-  [[intel::reqd_work_group_size(16, 16, 16)]] void operator()() const {} // expected-warning {{attribute 'intel::reqd_work_group_size' is deprecated}} \
-                                                                         // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}}
+  [[sycl::reqd_work_group_size(16, 16, 16)]] void operator()() const {}
 };
 
 class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}}
@@ -83,8 +75,7 @@ class Functor {
 
 class FunctorAttr {
 public:
-  __attribute__((reqd_work_group_size(128, 128, 128))) void operator()() const {} // expected-warning {{attribute 'reqd_work_group_size' is deprecated}} \
-                                                                                  // expected-note {{did you mean to use '[[sycl::reqd_work_group_size]]' instead?}}
+  [[sycl::reqd_work_group_size(128, 128, 128)]] void operator()() const {}
 };
 
 int main() {
diff --git a/clang/test/SemaSYCL/intel-restrict.cpp b/clang/test/SemaSYCL/intel-restrict.cpp
index bd51eb6b5c055..f8e7a670fddd6 100644
--- a/clang/test/SemaSYCL/intel-restrict.cpp
+++ b/clang/test/SemaSYCL/intel-restrict.cpp
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -Wno-sycl-2017-compat -triple spir64 -DCHECKDIAG -verify
-// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -Wno-sycl-2017-compat -triple spir64 | FileCheck %s
+// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 -DCHECKDIAG -verify
+// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 | FileCheck %s
 
 [[intel::kernel_args_restrict]] void func_do_not_ignore() {}
 
diff --git a/clang/test/SemaSYCL/named_sub_group_size.cpp b/clang/test/SemaSYCL/named_sub_group_size.cpp
new file mode 100644
index 0000000000000..d207e73453fa4
--- /dev/null
+++ b/clang/test/SemaSYCL/named_sub_group_size.cpp
@@ -0,0 +1,46 @@
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -sycl-std=2020 -fsyntax-only -ast-dump -verify %s | FileCheck %s
+
+// Validate the semantic analysis checks for the named_sub_group_size attribute in SYCL 2020 mode.
+
+#include "Inputs/sycl.hpp"
+
+struct Functor {
+  [[intel::named_sub_group_size(automatic)]] void operator()() const {
+  }
+};
+
+struct Functor1 {
+  [[intel::named_sub_group_size(primary)]] void operator()() const {
+  }
+};
+
+// Test attribute gets propagated to the kernel.
+void calls_kernel_1() {
+  // CHECK: FunctionDecl {{.*}}Kernel1
+  // CHECK: IntelNamedSubGroupSizeAttr {{.*}} Automatic
+  sycl::kernel_single_task<class Kernel1>([]() [[intel::named_sub_group_size(automatic)]] {
+  });
+}
+
+void calls_kernel_2() {
+  Functor F;
+  // CHECK: FunctionDecl {{.*}}Kernel2
+  // CHECK: IntelNamedSubGroupSizeAttr {{.*}} Automatic
+  sycl::kernel_single_task<class Kernel2>(F);
+
+  Functor1 F1;
+  // CHECK: FunctionDecl {{.*}}Kernel3
+  // CHECK: IntelNamedSubGroupSizeAttr {{.*}} Primary
+  sycl::kernel_single_task<class Kernel3>(F1);
+}
+
+[[intel::named_sub_group_size(primary)]] void AttrFunc() {} // expected-error{{kernel-called function must have a sub group size that matches the size specified for the kernel}}
+
+// Test attribute does not get propagated to the kernel.
+void calls_kernel_3() {
+  // CHECK:     FunctionDecl {{.*}}Kernel4
+  // CHECK-NOT: IntelNamedSubGroupSizeAttr {{.*}}
+  sycl::kernel_single_task<class Kernel4>([]() { // expected-note{{kernel declared here}}
+    AttrFunc();
+  });
+}
diff --git a/clang/test/SemaSYCL/num_simd_work_items_device.cpp b/clang/test/SemaSYCL/num_simd_work_items_device.cpp
index 0841662f9b7bb..858fe713e412b 100644
--- a/clang/test/SemaSYCL/num_simd_work_items_device.cpp
+++ b/clang/test/SemaSYCL/num_simd_work_items_device.cpp
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 %s -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -fsyntax-only -Wno-sycl-2017-compat -DTRIGGER_ERROR -verify
-// RUN: %clang_cc1 %s -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -fsyntax-only -Wno-sycl-2017-compat -ast-dump | FileCheck %s
+// RUN: %clang_cc1 %s -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -fsyntax-only -sycl-std=2020 -Wno-sycl-2017-compat -DTRIGGER_ERROR -verify
+// RUN: %clang_cc1 %s -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -ast-dump | FileCheck %s
 
 #include "sycl.hpp"
 
diff --git a/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp b/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp
index dfc39bc84a6d2..6be07189d4d1d 100755
--- a/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp
+++ b/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -triple spir64 | FileCheck %s
+// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -sycl-std=2017 -fsycl-is-device -triple spir64 | FileCheck %s
 
 #include "Inputs/sycl.hpp"
 
diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
index 448c36d5fdc98..56b70a64b3b7d 100644
--- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
+++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
@@ -1,6 +1,6 @@
-// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -Wno-sycl-2017-compat -verify
-// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -DTRIGGER_ERROR -Wno-sycl-2017-compat -verify
-// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -Wno-sycl-2017-compat | FileCheck %s
+// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -sycl-std=2017 -Wno-sycl-2017-compat -verify
+// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -sycl-std=2017 -DTRIGGER_ERROR -Wno-sycl-2017-compat -verify
+// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple spir64 -Wno-sycl-2017-compat | FileCheck %s
 
 #include "sycl.hpp"
 
diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp
index d4694411f0fb8..4569ca5284e83 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-is-device -internal-isystem %S/Inputs -fsyntax-only -Wno-sycl-2017-compat -verify -DTRIGGER_ERROR %s
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify -DTRIGGER_ERROR %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s
 
 #include "sycl.hpp"
 
diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp
new file mode 100644
index 0000000000000..88b8ab6a3a3a1
--- /dev/null
+++ b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp
@@ -0,0 +1,77 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2020 -verify -DTRIGGER_ERROR %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump %s | FileCheck %s
+
+// Test for AST of reqd_work_group_size kernel attribute in SYCL 2020.
+
+#include "sycl.hpp"
+
+using namespace cl::sycl;
+queue q;
+
+// The GNU and [[cl::reqd_work_group_size]] spellings are deprecated in SYCL
+// mode, and still requires all three arguments.
+__attribute__((reqd_work_group_size(4, 4, 4))) void four_once_more(); // expected-warning {{attribute 'reqd_work_group_size' is deprecated}} \
+                                                                      // expected-note {{did you mean to use '[[sycl::reqd_work_group_size]]' instead?}}
+[[cl::reqd_work_group_size(4, 4, 4)]] void four_with_feeling(); // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \
+                                                                // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}}
+
+#ifdef TRIGGER_ERROR
+__attribute__((reqd_work_group_size(4))) void four_yet_again(); // expected-error {{'reqd_work_group_size' attribute requires exactly 3 arguments}} \
+                                                                // expected-warning {{attribute 'reqd_work_group_size' is deprecated}} \
+                                                                // expected-note {{did you mean to use '[[sycl::reqd_work_group_size]]' instead?}}
+
+[[cl::reqd_work_group_size(4)]] void four_with_more_feeling(); // expected-error {{'reqd_work_group_size' attribute requires exactly 3 arguments}} \
+                                                               // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \
+                                                               // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}}
+class Functor32 {
+public:
+  // expected-note@+3{{conflicting attribute is here}}
+  // expected-warning@+2{{attribute 'reqd_work_group_size' is already applied with different arguments}}
+  // expected-error@+1{{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
+  [[sycl::reqd_work_group_size(32, 1, 1)]] [[sycl::reqd_work_group_size(1, 1, 32)]] void operator()() const {}
+};
+#endif // TRIGGER_ERROR
+
+class Functor16x16x16 {
+public:
+  [[sycl::reqd_work_group_size(16, 16, 16)]] void operator()() const {}
+};
+
+class FunctorAttr {
+public:
+  __attribute__((reqd_work_group_size(128, 128, 128))) void operator()() const {} // expected-warning {{attribute 'reqd_work_group_size' is deprecated}} \
+                                                                                  // expected-note {{did you mean to use '[[sycl::reqd_work_group_size]]' instead?}}
+};
+
+int main() {
+  q.submit([&](handler &h) {
+    // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1
+    // CHECK: ReqdWorkGroupSizeAttr
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 16
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 16
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 16
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 16
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 16
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 16
+    Functor16x16x16 f16x16x16;
+    h.single_task<class kernel_name1>(f16x16x16);
+
+    // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2
+    // CHECK: ReqdWorkGroupSizeAttr
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 128
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 128
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 128
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 128
+    // CHECK-NEXT:  ConstantExpr{{.*}}'int'
+    // CHECK-NEXT:  value: Int 128
+    // CHECK-NEXT:  IntegerLiteral {{.*}} 'int' 128
+    FunctorAttr fattr;
+    h.single_task<class kernel_name2>(fattr);
+  });
+  return 0;
+}
diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp
index cda6a368be58c..700f861c7fce9 100644
--- a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp
+++ b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp
@@ -1,6 +1,7 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -Wno-sycl-2017-compat -verify -DTRIGGER_ERROR %s
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify -DTRIGGER_ERROR %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s
 
+// Test for AST of reqd_work_group_size kernel attribute in SYCL 1.2.1.
 #include "sycl.hpp"
 
 using namespace cl::sycl;
@@ -34,23 +35,8 @@ queue q;
 [[sycl::reqd_work_group_size(4, 1, 1)]] void four_again(); // OK
 [[sycl::reqd_work_group_size(4, 1, 1)]] void four_again(); // OK
 
-// The GNU and [[cl::reqd_work_group_size]] spellings are deprecated in SYCL
-// mode, and still requires all three arguments.
-__attribute__((reqd_work_group_size(4, 4, 4))) void four_once_more(); // expected-warning {{attribute 'reqd_work_group_size' is deprecated}} \
-                                                                      // expected-note {{did you mean to use '[[sycl::reqd_work_group_size]]' instead?}}
-[[cl::reqd_work_group_size(4, 4, 4)]] void four_with_feeling(); // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \
-                                                                // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}}
-
-#ifdef TRIGGER_ERROR
-__attribute__((reqd_work_group_size(4))) void four_yet_again(); // expected-error {{'reqd_work_group_size' attribute requires exactly 3 arguments}} \
-                                                                // expected-warning {{attribute 'reqd_work_group_size' is deprecated}} \
-                                                                // expected-note {{did you mean to use '[[sycl::reqd_work_group_size]]' instead?}}
-
-[[cl::reqd_work_group_size(4)]] void four_with_more_feeling(); // expected-error {{'reqd_work_group_size' attribute requires exactly 3 arguments}} \
-                                                               // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \
-                                                               // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}}
-
 // Make sure there's at least one argument passed for the SYCL spelling.
+#ifdef TRIGGER_ERROR
 [[sycl::reqd_work_group_size]] void four_no_more(); // expected-error {{'reqd_work_group_size' attribute takes at least 1 argument}}
 #endif // TRIGGER_ERROR
 
@@ -87,12 +73,6 @@ class Functor {
   }
 };
 
-class FunctorAttr {
-public:
-  __attribute__((reqd_work_group_size(128, 128, 128))) void operator()() const {} // expected-warning {{attribute 'reqd_work_group_size' is deprecated}} \
-                                                                                  // expected-note {{did you mean to use '[[sycl::reqd_work_group_size]]' instead?}}
-};
-
 int main() {
   q.submit([&](handler &h) {
     Functor16 f16;
@@ -104,9 +84,6 @@ int main() {
     Functor16x16x16 f16x16x16;
     h.single_task<class kernel_name3>(f16x16x16);
 
-    FunctorAttr fattr;
-    h.single_task<class kernel_name4>(fattr);
-
     h.single_task<class kernel_name5>([]() [[sycl::reqd_work_group_size(32, 32, 32), sycl::reqd_work_group_size(32, 32, 32)]] {
       f32x32x32();
     });
@@ -176,17 +153,6 @@ int main() {
 // CHECK-NEXT:  ConstantExpr{{.*}}'int'
 // CHECK-NEXT:  value: Int 16
 // CHECK-NEXT:  IntegerLiteral{{.*}}16{{$}}
-// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name4
-// CHECK: ReqdWorkGroupSizeAttr {{.*}}
-// CHECK-NEXT:  ConstantExpr{{.*}}'int'
-// CHECK-NEXT:  value: Int 128
-// CHECK-NEXT:  IntegerLiteral{{.*}}128{{$}}
-// CHECK-NEXT:  ConstantExpr{{.*}}'int'
-// CHECK-NEXT:  value: Int 128
-// CHECK-NEXT:  IntegerLiteral{{.*}}128{{$}}
-// CHECK-NEXT:  ConstantExpr{{.*}}'int'
-// CHECK-NEXT:  value: Int 128
-// CHECK-NEXT:  IntegerLiteral{{.*}}128{{$}}
 // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5
 // CHECK: ReqdWorkGroupSizeAttr {{.*}}
 // CHECK-NEXT:  ConstantExpr{{.*}}'int'
diff --git a/clang/test/SemaSYCL/sycl-esimd-ast.cpp b/clang/test/SemaSYCL/sycl-esimd-ast.cpp
new file mode 100644
index 0000000000000..a0eae761582dc
--- /dev/null
+++ b/clang/test/SemaSYCL/sycl-esimd-ast.cpp
@@ -0,0 +1,41 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -ast-dump %s | FileCheck %s
+
+// Tests for AST of sycl_explicit_simd function attribute in SYCL 1.2.1.
+
+#include "sycl.hpp"
+
+sycl::queue deviceQueue;
+
+struct FuncObj {
+  [[intel::sycl_explicit_simd]] void operator()() const {}
+};
+
+[[intel::sycl_explicit_simd]] void func() {}
+
+int main() {
+  deviceQueue.submit([&](sycl::handler &h) {
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel1
+    // CHECK:       SYCLSimdAttr {{.*}} Implicit
+    // CHECK-NEXT:  SYCLKernelAttr {{.*}} Implicit
+    // CHECK-NEXT:  SYCLSimdAttr {{.*}}
+    h.single_task<class test_kernel1>(
+        FuncObj());
+
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel2
+    // CHECK:       SYCLSimdAttr {{.*}} Implicit
+    // CHECK-NEXT:  SYCLKernelAttr {{.*}} Implicit
+    // CHECK-NEXT:  SYCLSimdAttr {{.*}}
+    h.single_task<class test_kernel2>(
+        []() [[intel::sycl_explicit_simd]]{});
+
+    // Test attribute is propagated.
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
+    // CHECK:       SYCLSimdAttr {{.*}} Implicit
+    // CHECK-NEXT:  SYCLKernelAttr {{.*}} Implicit
+    // CHECK-NEXT:  SYCLSimdAttr {{.*}}
+    // CHECK-NEXT:  SYCLSimdAttr {{.*}}
+    h.single_task<class test_kernel3>(
+        []() [[intel::sycl_explicit_simd]] { func(); });
+  });
+  return 0;
+}
diff --git a/clang/test/SemaSYCL/sycl-esimd.cpp b/clang/test/SemaSYCL/sycl-esimd.cpp
index 4c8d9da02e91d..5153aa1ea3b48 100644
--- a/clang/test/SemaSYCL/sycl-esimd.cpp
+++ b/clang/test/SemaSYCL/sycl-esimd.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify %s
+// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify %s
 
 // This test checks specifics of semantic analysis of ESIMD kernels.