[SYCL][NFC] Refactor [[intel::reqd_sub_group_size()]] attribute tests

clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp

@@ -0,0 +1,110 @@
+// 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;
+}

clang/test/SemaSYCL/reqd-sub-group-size.cpp

@@ -0,0 +1,134 @@
+// 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;
+}