Skip to content

Commit c057d4d

Browse files
smanna12smanna12
authored
[FPGA][NFC] Refactor [[intel::max_work_group_size()]] attribute tests (#6240)
This patch updates test bases for [[intel::max_work_group_size()]] attribute: separates AST and diagnostics to make it easier to read/follow-up. removes duplicate test cases. no compiler changes. Signed-off-by: Soumi Manna [email protected]
1 parent 099af9c commit c057d4d

File tree

4 files changed

+170
-234
lines changed

4 files changed

+170
-234
lines changed

‎clang/test/SemaSYCL/intel-max-work-group-size-ast.cpp

Lines changed: 131 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,131 @@
1+
// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple spir64 | FileCheck %s
2+
3+
// The test checks support and functionality of [[intel:::max_work_group_size()]] attribute.
4+
#include "sycl.hpp"
5+
6+
using namespace cl::sycl;
7+
queue q;
8+
9+
[[intel::max_work_group_size(2, 2, 2)]] void func_do_not_ignore() {}
10+
11+
struct FuncObj {
12+
[[intel::max_work_group_size(4, 4, 4)]] void operator()() const {}
13+
};
14+
15+
// Test that checks template parameter support on member function of class template.
16+
template <int SIZE, int SIZE1, int SIZE2>
17+
class KernelFunctor {
18+
public:
19+
[[intel::max_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() {}
20+
};
21+
22+
// Test that checks template parameter support on function.
23+
template <int N, int N1, int N2>
24+
[[intel::max_work_group_size(N, N1, N2)]] void func() {}
25+
26+
int check() {
27+
// CHECK: ClassTemplateDecl {{.*}} {{.*}} KernelFunctor
28+
// CHECK: ClassTemplateSpecializationDecl {{.*}} {{.*}} class KernelFunctor definition
29+
// CHECK: CXXRecordDecl {{.*}} {{.*}} implicit class KernelFunctor
30+
// CHECK: SYCLIntelMaxWorkGroupSizeAttr
31+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
32+
// CHECK-NEXT: value: Int 4
33+
// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}}
34+
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
35+
// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
36+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
37+
// CHECK-NEXT: value: Int 4
38+
// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}}
39+
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
40+
// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
41+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
42+
// CHECK-NEXT: value: Int 4
43+
// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}}
44+
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
45+
// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
46+
KernelFunctor<4, 4, 4>();
47+
48+
// CHECK: FunctionTemplateDecl {{.*}} {{.*}} func
49+
// CHECK: FunctionDecl {{.*}} {{.*}} used func 'void ()'
50+
// CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}}
51+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
52+
// CHECK-NEXT: value: Int 8
53+
// CHECK: SubstNonTypeTemplateParmExpr {{.*}}
54+
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
55+
// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}
56+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
57+
// CHECK-NEXT: value: Int 8
58+
// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}}
59+
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
60+
// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}
61+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
62+
// CHECK-NEXT: value: Int 8
63+
// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}}
64+
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
65+
// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}
66+
func<8, 8, 8>();
67+
return 0;
68+
}
69+
70+
int main() {
71+
q.submit([&](handler &h) {
72+
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel1
73+
// CHECK: SYCLIntelMaxWorkGroupSizeAttr
74+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
75+
// CHECK-NEXT: value: Int 4
76+
// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
77+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
78+
// CHECK-NEXT: value: Int 4
79+
// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
80+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
81+
// CHECK-NEXT: value: Int 4
82+
// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
83+
h.single_task<class test_kernel1>(FuncObj());
84+
85+
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel2
86+
// CHECK: SYCLIntelMaxWorkGroupSizeAttr
87+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
88+
// CHECK-NEXT: value: Int 8
89+
// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}
90+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
91+
// CHECK-NEXT: value: Int 8
92+
// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}
93+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
94+
// CHECK-NEXT: value: Int 8
95+
// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}
96+
h.single_task<class test_kernel2>(
97+
[]() [[intel::max_work_group_size(8, 8, 8)]] {});
98+
99+
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
100+
// CHECK: SYCLIntelMaxWorkGroupSizeAttr
101+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
102+
// CHECK-NEXT: value: Int 2
103+
// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
104+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
105+
// CHECK-NEXT: value: Int 2
106+
// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
107+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
108+
// CHECK-NEXT: value: Int 2
109+
// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
110+
h.single_task<class test_kernel3>(
111+
[]() { func_do_not_ignore(); });
112+
113+
// Ignore duplicate attribute.
114+
h.single_task<class test_kernel10>(
115+
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel10
116+
// CHECK: SYCLIntelMaxWorkGroupSizeAttr
117+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
118+
// CHECK-NEXT: value: Int 2
119+
// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
120+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
121+
// CHECK-NEXT: value: Int 2
122+
// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
123+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
124+
// CHECK-NEXT: value: Int 2
125+
// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
126+
// CHECK-NOT: SYCLIntelMaxWorkGroupSizeAttr
127+
[]() [[intel::max_work_group_size(2, 2, 2),
128+
intel::max_work_group_size(2, 2, 2)]] {});
129+
});
130+
return 0;
131+
}

‎clang/test/SemaSYCL/intel-max-work-group-size-device.cpp

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

‎clang/test/SemaSYCL/intel-max-work-group-size.cpp

Lines changed: 39 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,7 @@
55
[[intel::max_work_group_size(12, 12, 12, 12)]] void f0(); // expected-error {{'max_work_group_size' attribute requires exactly 3 arguments}}
66
[[intel::max_work_group_size("derp", 1, 2)]] void f1(); // expected-error {{integral constant expression must have integral or unscoped enumeration type, not 'const char[5]'}}
77
[[intel::max_work_group_size(1, 1, 1)]] int i; // expected-error {{'max_work_group_size' attribute only applies to functions}}
8+
[[intel::max_work_group_size(-8, 8, -8)]] void neg(); // expected-error 2{{'max_work_group_size' attribute requires a positive integral compile time constant expression}}
89

910
// Tests for Intel FPGA 'max_work_group_size' attribute duplication.
1011
// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored.
@@ -105,3 +106,41 @@ f18();
105106

106107
[[intel::max_work_group_size(16, 16, 1)]] void f19();
107108
[[sycl::reqd_work_group_size(16, 16)]] void f19(); // OK
109+
110+
// Test that checks wrong function template instantiation and ensures that the type
111+
// is checked properly when instantiating from the template definition.
112+
113+
template <typename Ty, typename Ty1, typename Ty2>
114+
// expected-error@+1 3{{integral constant expression must have integral or unscoped enumeration type, not 'S'}}
115+
[[intel::max_work_group_size(Ty{}, Ty1{}, Ty2{})]] void f20() {}
116+
117+
struct S {};
118+
void var() {
119+
// expected-note@+1 {{in instantiation of function template specialization 'f20<S, S, S>' requested here}}
120+
f20<S, S, S>();
121+
}
122+
123+
// Test that checks expression is not a constant expression.
124+
// expected-note@+1 3{{declared here}}
125+
int foo();
126+
// expected-error@+2 3{{expression is not an integral constant expression}}
127+
// expected-note@+1 3{{non-constexpr function 'foo' cannot be used in a constant expression}}
128+
[[intel::max_work_group_size(foo() + 12, foo() + 12, foo() + 12)]] void f21();
129+
130+
// Test that checks expression is a constant expression.
131+
constexpr int bar() { return 0; }
132+
[[intel::max_work_group_size(bar() + 12, bar() + 12, bar() + 12)]] void f22(); // OK
133+
134+
struct DAFuncObj {
135+
[[intel::max_work_group_size(4, 4, 4)]] // expected-note {{conflicting attribute is here}}
136+
[[cl::reqd_work_group_size(8, 8, 4)]] // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} \
137+
// expected-warning{{attribute 'cl::reqd_work_group_size' is deprecated}} \
138+
// expected-note{{did you mean to use 'sycl::reqd_work_group_size' instead?}}
139+
void
140+
operator()() const {}
141+
};
142+
143+
struct Func {
144+
// expected-warning@+1 {{unknown attribute 'max_work_group_size' ignored}}
145+
[[intelfpga::max_work_group_size(1, 1, 1)]] void operator()() const {}
146+
};

0 commit comments

Comments
 (0)