Skip to content

Commit 4071659

Browse files
smanna12smanna12
authored
[SYCL][NFC] Refactor [[intel::reqd_sub_group_size()]] attribute tests (#6243)
This patch updates tests for [[intel::reqd_sub_group_size()]] attribute: * separates AST and diagnostics to make it easier to read/follow-up. * removes duplicate test cases. * adds missing tests. * no compiler changes. Signed-off-by: Soumi Manna [email protected]
1 parent 5de8d94 commit 4071659

File tree

4 files changed

+244
-219
lines changed

4 files changed

+244
-219
lines changed

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

Lines changed: 110 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,110 @@
1+
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s
2+
3+
// The test checks AST of [[intel::reqd_sub_group_size()]] attribute.
4+
5+
#include "sycl.hpp"
6+
7+
using namespace cl::sycl;
8+
queue q;
9+
10+
[[intel::reqd_sub_group_size(4)]] void foo() {}
11+
12+
class Functor16 {
13+
public:
14+
[[intel::reqd_sub_group_size(16)]] void operator()() const {}
15+
};
16+
17+
class Functor {
18+
public:
19+
void operator()() const {
20+
foo();
21+
}
22+
};
23+
24+
// Test that checks template parameter support on member function of class template.
25+
template <int SIZE>
26+
class KernelFunctor {
27+
public:
28+
[[intel::reqd_sub_group_size(SIZE)]] void operator()() const {}
29+
};
30+
31+
// Test that checks template parameter support on function.
32+
// CHECK: FunctionTemplateDecl {{.*}} func
33+
// CHECK: FunctionDecl {{.*}} func 'void ()'
34+
// CHECK-NEXT: CompoundStmt
35+
// CHECK_NEXT: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
36+
// CHECK_NEXT: DeclRefExpr {{.*}} 'int' NonTypeTemplateParm {{.*}} 'N' 'int'
37+
// CHECK: FunctionDecl {{.*}} func 'void ()'
38+
// CHECK-NEXT: TemplateArgument integral 12
39+
// CHECK-NEXT: CompoundStmt
40+
// CHECK-NEXT: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
41+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
42+
// CHECK-NEXT: value: Int 12
43+
// CHECK-NEXT: SubstNonTypeTemplateParmExpr
44+
// CHECK-NEXT: NonTypeTemplateParmDecl
45+
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 12
46+
template <int N>
47+
[[intel::reqd_sub_group_size(N)]] void func() {}
48+
49+
int main() {
50+
q.submit([&](handler &h) {
51+
// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1
52+
// CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
53+
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
54+
// CHECK-NEXT: value: Int 16
55+
// CHECK-NEXT: IntegerLiteral{{.*}}16{{$}}
56+
Functor16 f16;
57+
h.single_task<class kernel_name1>(f16);
58+
59+
// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2
60+
// CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
61+
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
62+
// CHECK-NEXT: value: Int 4
63+
// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
64+
Functor f;
65+
h.single_task<class kernel_name2>(f);
66+
67+
// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name3
68+
// CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
69+
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
70+
// CHECK-NEXT: value: Int 2
71+
// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
72+
h.single_task<class kernel_name3>([]() [[intel::reqd_sub_group_size(2)]] {});
73+
74+
// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name4
75+
// CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
76+
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
77+
// CHECK-NEXT: value: Int 4
78+
// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
79+
h.single_task<class kernel_name4>([]() [[intel::reqd_sub_group_size(4)]] { foo(); });
80+
// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5
81+
// CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
82+
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
83+
// CHECK-NEXT: value: Int 6
84+
// CHECK-NEXT: IntegerLiteral{{.*}}6{{$}}
85+
h.single_task<class kernel_name5>([]() [[intel::reqd_sub_group_size(6)]] {});
86+
87+
// CHECK: FunctionDecl {{.*}}kernel_name_6
88+
// CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
89+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
90+
// CHECK-NEXT: value: Int 10
91+
// CHECK-NEXT: SubstNonTypeTemplateParmExpr
92+
// CHECK-NEXT: NonTypeTemplateParmDecl
93+
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 10
94+
KernelFunctor<10> f2;
95+
h.single_task<class kernel_name_6>(f2);
96+
97+
// Ignore duplicate attribute.
98+
h.single_task<class kernel_name_7>(
99+
// CHECK: FunctionDecl {{.*}}kernel_name_7
100+
// CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
101+
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
102+
// CHECK-NEXT: value: Int 8
103+
// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}
104+
// CHECK-NOT: IntelReqdSubGroupSizeAttr
105+
[]() [[intel::reqd_sub_group_size(8),
106+
intel::reqd_sub_group_size(8)]] {});
107+
});
108+
func<12>();
109+
return 0;
110+
}

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

Lines changed: 0 additions & 125 deletions
This file was deleted.

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

Lines changed: 134 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,134 @@
1+
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify -pedantic %s
2+
3+
// The test checks functionality of [[intel::reqd_sub_group_size()]] attribute on SYCL kernel.
4+
5+
#include "sycl.hpp"
6+
7+
using namespace cl::sycl;
8+
queue q;
9+
10+
[[intel::reqd_sub_group_size(4)]] void foo() {} // expected-note {{conflicting attribute is here}}
11+
// expected-note@-1 {{conflicting attribute is here}}
12+
[[intel::reqd_sub_group_size(32)]] void baz() {} // expected-note {{conflicting attribute is here}}
13+
14+
class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}}
15+
public:
16+
[[intel::reqd_sub_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}}
17+
foo();
18+
}
19+
};
20+
21+
int main() {
22+
q.submit([&](handler &h) {
23+
Functor8 f8;
24+
h.single_task<class kernel_name1>(f8);
25+
26+
h.single_task<class kernel_name2>([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}}
27+
foo();
28+
baz();
29+
});
30+
});
31+
return 0;
32+
}
33+
[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B();
34+
[[intel::reqd_sub_group_size(16)]] void A() {
35+
}
36+
37+
[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B() {
38+
A();
39+
}
40+
// expected-note@+1 {{conflicting attribute is here}}
41+
[[intel::reqd_sub_group_size(2)]] void sg_size2() {}
42+
43+
// expected-note@+2 {{conflicting attribute is here}}
44+
// expected-error@+1 {{conflicting attributes applied to a SYCL kernel}}
45+
[[intel::reqd_sub_group_size(4)]] __attribute__((sycl_device)) void sg_size4() {
46+
sg_size2();
47+
}
48+
49+
// Test that checks support and functionality of reqd_sub_group_size attribute support on function.
50+
51+
// Tests for incorrect argument values for Intel reqd_sub_group_size attribute.
52+
[[intel::reqd_sub_group_size]] void one() {} // expected-error {{'reqd_sub_group_size' attribute takes one argument}}
53+
[[intel::reqd_sub_group_size(5)]] int a; // expected-error{{'reqd_sub_group_size' attribute only applies to functions}}
54+
[[intel::reqd_sub_group_size("foo")]] void func() {} // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'const char[4]'}}
55+
[[intel::reqd_sub_group_size(-1)]] void func1() {} // expected-error{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
56+
[[intel::reqd_sub_group_size(0, 1)]] void arg() {} // expected-error{{'reqd_sub_group_size' attribute takes one argument}}
57+
58+
// Diagnostic is emitted because the arguments mismatch.
59+
[[intel::reqd_sub_group_size(12)]] void quux(); // expected-note {{previous attribute is here}}
60+
[[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}}
61+
[[sycl::reqd_sub_group_size(200)]] void quux(); // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}}
62+
63+
// Make sure there's at least one argument passed.
64+
[[sycl::reqd_sub_group_size]] void quibble(); // expected-error {{'reqd_sub_group_size' attribute takes one argument}}
65+
66+
// No diagnostic is emitted because the arguments match.
67+
[[intel::reqd_sub_group_size(12)]] void same();
68+
[[intel::reqd_sub_group_size(12)]] void same() {} // OK
69+
70+
// No diagnostic because the attributes are synonyms with identical behavior.
71+
[[sycl::reqd_sub_group_size(12)]] void same(); // OK
72+
73+
// Test that checks wrong function template instantiation and ensures that the type
74+
// is checked properly when instantiating from the template definition.
75+
template <typename Ty>
76+
// expected-error@+3{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
77+
// expected-error@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}}
78+
// expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}}
79+
[[intel::reqd_sub_group_size(Ty{})]] void func() {}
80+
81+
struct S {};
82+
void test() {
83+
// expected-note@+1{{in instantiation of function template specialization 'func<S>' requested here}}
84+
func<S>();
85+
// expected-note@+1{{in instantiation of function template specialization 'func<float>' requested here}}
86+
func<float>();
87+
// expected-note@+1{{in instantiation of function template specialization 'func<int>' requested here}}
88+
func<int>();
89+
}
90+
91+
// Test that checks expression is not a constant expression.
92+
// expected-note@+1{{declared here}}
93+
int foo1();
94+
// expected-error@+2{{expression is not an integral constant expression}}
95+
// expected-note@+1{{non-constexpr function 'foo1' cannot be used in a constant expression}}
96+
[[intel::reqd_sub_group_size(foo1() + 12)]] void func1();
97+
98+
// Test that checks expression is a constant expression.
99+
constexpr int bar1() { return 0; }
100+
[[intel::reqd_sub_group_size(bar1() + 12)]] void func2(); // OK
101+
102+
// Test that checks template parameter support on member function of class template.
103+
template <int SIZE>
104+
class KernelFunctor {
105+
public:
106+
// expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
107+
[[intel::reqd_sub_group_size(SIZE)]] void operator()() {}
108+
};
109+
110+
int check() {
111+
// expected-note@+1{{in instantiation of template class 'KernelFunctor<-1>' requested here}}
112+
KernelFunctor<-1>();
113+
return 0;
114+
}
115+
116+
// Test that checks template parameter support on function.
117+
template <int N>
118+
// expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
119+
[[intel::reqd_sub_group_size(N)]] void func3() {}
120+
121+
template <int N>
122+
[[intel::reqd_sub_group_size(4)]] void func4(); // expected-note {{previous attribute is here}}
123+
124+
template <int N>
125+
[[intel::reqd_sub_group_size(N)]] void func4() {} // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}}
126+
127+
int check1() {
128+
// no error expected
129+
func3<12>();
130+
// expected-note@+1{{in instantiation of function template specialization 'func3<-1>' requested here}}
131+
func3<-1>();
132+
func4<6>(); // expected-note {{in instantiation of function template specialization 'func4<6>' requested here}}
133+
return 0;
134+
}

0 commit comments

Comments
 (0)