diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp
deleted file mode 100644
index 67ae007429213..0000000000000
--- a/clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp
+++ /dev/null
@@ -1,110 +0,0 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s
-
-// The test checks AST of [[intel::reqd_sub_group_size()]] attribute.
-
-#include "sycl.hpp"
-
-using namespace cl::sycl;
-queue q;
-
-[[intel::reqd_sub_group_size(4)]] void foo() {}
-
-class Functor16 {
-public:
-  [[intel::reqd_sub_group_size(16)]] void operator()() const {}
-};
-
-class Functor {
-public:
-  void operator()() const {
-    foo();
-  }
-};
-
-// Test that checks template parameter support on member function of class template.
-template <int SIZE>
-class KernelFunctor {
-public:
-  [[intel::reqd_sub_group_size(SIZE)]] void operator()() const {}
-};
-
-// Test that checks template parameter support on function.
-// CHECK: FunctionTemplateDecl {{.*}} func
-// CHECK: FunctionDecl {{.*}} func 'void ()'
-// CHECK-NEXT: CompoundStmt
-// CHECK_NEXT: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
-// CHECK_NEXT: DeclRefExpr {{.*}} 'int' NonTypeTemplateParm {{.*}} 'N' 'int'
-// CHECK: FunctionDecl {{.*}} func 'void ()'
-// CHECK-NEXT: TemplateArgument integral 12
-// CHECK-NEXT: CompoundStmt
-// CHECK-NEXT: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
-// CHECK-NEXT: ConstantExpr{{.*}}'int'
-// CHECK-NEXT: value: Int 12
-// CHECK-NEXT: SubstNonTypeTemplateParmExpr
-// CHECK-NEXT: NonTypeTemplateParmDecl
-// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 12
-template <int N>
-[[intel::reqd_sub_group_size(N)]] void func() {}
-
-int main() {
-  q.submit([&](handler &h) {
-    // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1
-    // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
-    // CHECK-NEXT: ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT: value: Int 16
-    // CHECK-NEXT: IntegerLiteral{{.*}}16{{$}}
-    Functor16 f16;
-    h.single_task<class kernel_name1>(f16);
-
-    // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2
-    // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
-    // CHECK-NEXT: ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT: value: Int 4
-    // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
-    Functor f;
-    h.single_task<class kernel_name2>(f);
-
-    // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name3
-    // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
-    // CHECK-NEXT: ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT: value: Int 2
-    // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
-    h.single_task<class kernel_name3>([]() [[intel::reqd_sub_group_size(2)]] {});
-
-    // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name4
-    // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
-    // CHECK-NEXT: ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT: value: Int 4
-    // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
-    h.single_task<class kernel_name4>([]() [[intel::reqd_sub_group_size(4)]] { foo(); });
-    // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5
-    // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
-    // CHECK-NEXT: ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT: value: Int 6
-    // CHECK-NEXT: IntegerLiteral{{.*}}6{{$}}
-    h.single_task<class kernel_name5>([]() [[intel::reqd_sub_group_size(6)]] {});
-
-    // CHECK: FunctionDecl {{.*}}kernel_name_6
-    // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
-    // CHECK-NEXT: ConstantExpr{{.*}}'int'
-    // CHECK-NEXT: value: Int 10
-    // CHECK-NEXT: SubstNonTypeTemplateParmExpr
-    // CHECK-NEXT: NonTypeTemplateParmDecl
-    // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 10
-    KernelFunctor<10> f2;
-    h.single_task<class kernel_name_6>(f2);
-
-    // Ignore duplicate attribute.
-    h.single_task<class kernel_name_7>(
-        // CHECK: FunctionDecl {{.*}}kernel_name_7
-        // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
-        // CHECK-NEXT: ConstantExpr {{.*}} 'int'
-        // CHECK-NEXT: value: Int 8
-        // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}
-        // CHECK-NOT: IntelReqdSubGroupSizeAttr
-        []() [[intel::reqd_sub_group_size(8),
-               intel::reqd_sub_group_size(8)]] {});
-  });
-  func<12>();
-  return 0;
-}
diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp
new file mode 100644
index 0000000000000..4569ca5284e83
--- /dev/null
+++ b/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp
@@ -0,0 +1,125 @@
+// 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"
+
+using namespace cl::sycl;
+queue q;
+
+[[intel::reqd_sub_group_size(4)]] void foo() {} // expected-note {{conflicting attribute is here}}
+// expected-note@-1 {{conflicting attribute is here}}
+[[intel::reqd_sub_group_size(32)]] void baz() {} // expected-note {{conflicting attribute is here}}
+
+// No diagnostic is emitted because the arguments match.
+[[intel::reqd_sub_group_size(12)]] void bar();
+[[intel::reqd_sub_group_size(12)]] void bar() {} // OK
+
+// No diagnostic because the attributes are synonyms with identical behavior.
+[[sycl::reqd_sub_group_size(12)]] void bar(); // OK
+
+// Diagnostic is emitted because the arguments mismatch.
+[[intel::reqd_sub_group_size(12)]] void quux(); // expected-note {{previous attribute is here}}
+[[intel::reqd_sub_group_size(100)]] void quux(); // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}} expected-note {{previous attribute is here}}
+[[sycl::reqd_sub_group_size(200)]] void quux(); // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}}
+
+#ifdef TRIGGER_ERROR
+// Make sure there's at least one argument passed.
+[[sycl::reqd_sub_group_size]] void quibble(); // expected-error {{'reqd_sub_group_size' attribute takes one argument}}
+#endif // TRIGGER_ERROR
+
+class Functor16 {
+public:
+  [[intel::reqd_sub_group_size(16)]] void operator()() const {}
+};
+
+class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}}
+public:
+  [[intel::reqd_sub_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}}
+    foo();
+  }
+};
+
+class Functor4 {
+public:
+  [[intel::reqd_sub_group_size(12)]] void operator()() const {}
+};
+
+class Functor {
+public:
+  void operator()() const {
+    foo();
+  }
+};
+
+int main() {
+  q.submit([&](handler &h) {
+    Functor16 f16;
+    h.single_task<class kernel_name1>(f16);
+
+    Functor f;
+    h.single_task<class kernel_name2>(f);
+
+#ifdef TRIGGER_ERROR
+    Functor8 f8;
+    h.single_task<class kernel_name3>(f8);
+
+    h.single_task<class kernel_name4>([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}}
+      foo();
+      baz();
+    });
+#endif
+
+    h.single_task<class kernel_name5>([]() [[intel::reqd_sub_group_size(2)]]{});
+    h.single_task<class kernel_name6>([]() [[intel::reqd_sub_group_size(4)]] { foo(); });
+    h.single_task<class kernel_name7>([]() [[intel::reqd_sub_group_size(6)]]{});
+
+    Functor4 f4;
+    h.single_task<class kernel_name8>(f4);
+  });
+  return 0;
+}
+
+[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B();
+[[intel::reqd_sub_group_size(16)]] void A() {
+}
+
+[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B() {
+  A();
+}
+
+#ifdef TRIGGER_ERROR
+// expected-note@+1 {{conflicting attribute is here}}
+[[intel::reqd_sub_group_size(2)]] void sg_size2() {}
+
+// expected-note@+2 {{conflicting attribute is here}}
+// expected-error@+1 {{conflicting attributes applied to a SYCL kernel}}
+[[intel::reqd_sub_group_size(4)]] __attribute__((sycl_device)) void sg_size4() {
+  sg_size2();
+}
+#endif
+
+// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1
+// CHECK: IntelReqdSubGroupSizeAttr {{.*}}
+// CHECK-NEXT: ConstantExpr {{.*}} 'int'
+// CHECK-NEXT: value: Int 16
+// CHECK-NEXT: IntegerLiteral{{.*}}16{{$}}
+// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2
+// CHECK: IntelReqdSubGroupSizeAttr {{.*}}
+// CHECK-NEXT: ConstantExpr {{.*}} 'int'
+// CHECK-NEXT: value: Int 4
+// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
+// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5
+// CHECK: IntelReqdSubGroupSizeAttr {{.*}}
+// CHECK-NEXT: ConstantExpr {{.*}} 'int'
+// CHECK-NEXT: value: Int 2
+// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
+// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name7
+// CHECK: IntelReqdSubGroupSizeAttr {{.*}}
+// CHECK-NEXT: ConstantExpr {{.*}} 'int'
+// CHECK-NEXT: value: Int 6
+// CHECK-NEXT: IntegerLiteral{{.*}}6{{$}}
+// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name8
+// CHECK: IntelReqdSubGroupSizeAttr {{.*}}
+// CHECK-NEXT: ConstantExpr {{.*}} 'int'
+// CHECK-NEXT: value: Int 12
+// CHECK-NEXT: IntegerLiteral{{.*}}12{{$}}
diff --git a/clang/test/SemaSYCL/reqd-sub-group-size.cpp b/clang/test/SemaSYCL/reqd-sub-group-size.cpp
deleted file mode 100644
index 6b6a51fcbc360..0000000000000
--- a/clang/test/SemaSYCL/reqd-sub-group-size.cpp
+++ /dev/null
@@ -1,134 +0,0 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify -pedantic %s
-
-// The test checks functionality of [[intel::reqd_sub_group_size()]] attribute on SYCL kernel.
-
-#include "sycl.hpp"
-
-using namespace cl::sycl;
-queue q;
-
-[[intel::reqd_sub_group_size(4)]] void foo() {} // expected-note {{conflicting attribute is here}}
-// expected-note@-1 {{conflicting attribute is here}}
-[[intel::reqd_sub_group_size(32)]] void baz() {} // expected-note {{conflicting attribute is here}}
-
-class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}}
-public:
-  [[intel::reqd_sub_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}}
-    foo();
-  }
-};
-
-int main() {
-  q.submit([&](handler &h) {
-    Functor8 f8;
-    h.single_task<class kernel_name1>(f8);
-
-    h.single_task<class kernel_name2>([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}}
-      foo();
-      baz();
-    });
-  });
-  return 0;
-}
-[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B();
-[[intel::reqd_sub_group_size(16)]] void A() {
-}
-
-[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B() {
-  A();
-}
-// expected-note@+1 {{conflicting attribute is here}}
-[[intel::reqd_sub_group_size(2)]] void sg_size2() {}
-
-// expected-note@+2 {{conflicting attribute is here}}
-// expected-error@+1 {{conflicting attributes applied to a SYCL kernel}}
-[[intel::reqd_sub_group_size(4)]] __attribute__((sycl_device)) void sg_size4() {
-  sg_size2();
-}
-
-// Test that checks support and functionality of reqd_sub_group_size attribute support on function.
-
-// Tests for incorrect argument values for Intel reqd_sub_group_size attribute.
-[[intel::reqd_sub_group_size]] void one() {}         // expected-error {{'reqd_sub_group_size' attribute takes one argument}}
-[[intel::reqd_sub_group_size(5)]] int a;             // expected-error{{'reqd_sub_group_size' attribute only applies to functions}}
-[[intel::reqd_sub_group_size("foo")]] void func() {} // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'const char[4]'}}
-[[intel::reqd_sub_group_size(-1)]] void func1() {}   // expected-error{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
-[[intel::reqd_sub_group_size(0, 1)]] void arg() {}   // expected-error{{'reqd_sub_group_size' attribute takes one argument}}
-
-// Diagnostic is emitted because the arguments mismatch.
-[[intel::reqd_sub_group_size(12)]] void quux();  // expected-note {{previous attribute is here}}
-[[intel::reqd_sub_group_size(100)]] void quux(); // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}} expected-note {{previous attribute is here}}
-[[sycl::reqd_sub_group_size(200)]] void quux();  // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}}
-
-// Make sure there's at least one argument passed.
-[[sycl::reqd_sub_group_size]] void quibble(); // expected-error {{'reqd_sub_group_size' attribute takes one argument}}
-
-// No diagnostic is emitted because the arguments match.
-[[intel::reqd_sub_group_size(12)]] void same();
-[[intel::reqd_sub_group_size(12)]] void same() {} // OK
-
-// No diagnostic because the attributes are synonyms with identical behavior.
-[[sycl::reqd_sub_group_size(12)]] void same(); // OK
-
-// Test that checks wrong function template instantiation and ensures that the type
-// is checked properly when instantiating from the template definition.
-template <typename Ty>
-// expected-error@+3{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
-// expected-error@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}}
-// expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}}
-[[intel::reqd_sub_group_size(Ty{})]] void func() {}
-
-struct S {};
-void test() {
-  // expected-note@+1{{in instantiation of function template specialization 'func<S>' requested here}}
-  func<S>();
-  // expected-note@+1{{in instantiation of function template specialization 'func<float>' requested here}}
-  func<float>();
-  // expected-note@+1{{in instantiation of function template specialization 'func<int>' requested here}}
-  func<int>();
-}
-
-// Test that checks expression is not a constant expression.
-// expected-note@+1{{declared here}}
-int foo1();
-// expected-error@+2{{expression is not an integral constant expression}}
-// expected-note@+1{{non-constexpr function 'foo1' cannot be used in a constant expression}}
-[[intel::reqd_sub_group_size(foo1() + 12)]] void func1();
-
-// Test that checks expression is a constant expression.
-constexpr int bar1() { return 0; }
-[[intel::reqd_sub_group_size(bar1() + 12)]] void func2(); // OK
-
-// Test that checks template parameter support on member function of class template.
-template <int SIZE>
-class KernelFunctor {
-public:
-  // expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
-  [[intel::reqd_sub_group_size(SIZE)]] void operator()() {}
-};
-
-int check() {
-  // expected-note@+1{{in instantiation of template class 'KernelFunctor<-1>' requested here}}
-  KernelFunctor<-1>();
-  return 0;
-}
-
-// Test that checks template parameter support on function.
-template <int N>
-// expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
-[[intel::reqd_sub_group_size(N)]] void func3() {}
-
-template <int N>
-[[intel::reqd_sub_group_size(4)]] void func4(); // expected-note {{previous attribute is here}}
-
-template <int N>
-[[intel::reqd_sub_group_size(N)]] void func4() {} // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}}
-
-int check1() {
-  // no error expected
-  func3<12>();
-  // expected-note@+1{{in instantiation of function template specialization 'func3<-1>' requested here}}
-  func3<-1>();
-  func4<6>(); // expected-note {{in instantiation of function template specialization 'func4<6>' requested here}}
-  return 0;
-}
diff --git a/clang/test/SemaSYCL/sycl-device-reqd-sub-group-size-template.cpp b/clang/test/SemaSYCL/sycl-device-reqd-sub-group-size-template.cpp
new file mode 100644
index 0000000000000..90908b01ce2ad
--- /dev/null
+++ b/clang/test/SemaSYCL/sycl-device-reqd-sub-group-size-template.cpp
@@ -0,0 +1,94 @@
+// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s
+
+// Test that checkes template parameter support for 'reqd_sub_group_size' attribute on sycl device.
+
+// Test that checks wrong function template instantiation and ensures that the type
+// is checked properly when instantiating from the template definition.
+template <typename Ty>
+// expected-error@+3{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
+// expected-error@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}}
+// expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}}
+[[intel::reqd_sub_group_size(Ty{})]] void func() {}
+
+struct S {};
+void test() {
+  //expected-note@+1{{in instantiation of function template specialization 'func<S>' requested here}}
+  func<S>();
+  //expected-note@+1{{in instantiation of function template specialization 'func<float>' requested here}}
+  func<float>();
+  //expected-note@+1{{in instantiation of function template specialization 'func<int>' requested here}}
+  func<int>();
+}
+
+// Test that checks expression is not a constant expression.
+// expected-note@+1{{declared here}}
+int foo();
+// expected-error@+2{{expression is not an integral constant expression}}
+// expected-note@+1{{non-constexpr function 'foo' cannot be used in a constant expression}}
+[[intel::reqd_sub_group_size(foo() + 12)]] void func1();
+
+// Test that checks expression is a constant expression.
+constexpr int bar() { return 0; }
+[[intel::reqd_sub_group_size(bar() + 12)]] void func2(); // OK
+
+// Test that checks template parameter support on member function of class template.
+template <int SIZE>
+class KernelFunctor {
+public:
+  // expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
+  [[intel::reqd_sub_group_size(SIZE)]] void operator()() {}
+};
+
+int main() {
+  //expected-note@+1{{in instantiation of template class 'KernelFunctor<-1>' requested here}}
+  KernelFunctor<-1>();
+  // no error expected
+  KernelFunctor<10>();
+  return 0;
+}
+
+// CHECK: ClassTemplateDecl {{.*}} {{.*}} KernelFunctor
+// CHECK: ClassTemplateSpecializationDecl {{.*}} {{.*}} class KernelFunctor definition
+// CHECK: CXXRecordDecl {{.*}} {{.*}} implicit class KernelFunctor
+// CHECK: IntelReqdSubGroupSizeAttr {{.*}}
+// CHECK: SubstNonTypeTemplateParmExpr {{.*}}
+// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
+// CHECK-NEXT: IntegerLiteral{{.*}}10{{$}}
+
+// Test that checks template parameter support on function.
+template <int N>
+// expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
+[[intel::reqd_sub_group_size(N)]] void func3() {}
+
+template <int N>
+[[intel::reqd_sub_group_size(4)]] void func4(); // expected-note {{previous attribute is here}}
+
+template <int N>
+[[intel::reqd_sub_group_size(N)]] void func4() {} // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}}
+
+int check() {
+  // no error expected
+  func3<12>();
+  //expected-note@+1{{in instantiation of function template specialization 'func3<-1>' requested here}}
+  func3<-1>();
+  func4<6>(); //expected-note {{in instantiation of function template specialization 'func4<6>' requested here}}
+  return 0;
+}
+
+// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored.
+[[intel::reqd_sub_group_size(8)]]
+[[intel::reqd_sub_group_size(8)]] void func5() {}
+
+// CHECK: FunctionTemplateDecl {{.*}} {{.*}} func3
+// CHECK: NonTypeTemplateParmDecl {{.*}} {{.*}} referenced 'int' depth 0 index 0 N
+// CHECK: FunctionDecl {{.*}} {{.*}} func3 'void ()'
+// CHECK: IntelReqdSubGroupSizeAttr {{.*}}
+// CHECK: SubstNonTypeTemplateParmExpr {{.*}}
+// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
+// CHECK-NEXT: IntegerLiteral{{.*}}12{{$}}
+
+// CHECK: FunctionDecl {{.*}} {{.*}} func5 'void ()'
+// CHECK: IntelReqdSubGroupSizeAttr {{.*}}
+// CHECK-NEXT: ConstantExpr {{.*}} 'int'
+// CHECK-NEXT: value: Int 8
+// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}