Skip to content

Commit 334cc23

Browse files
smanna12smanna12
committed
address @elizabeth's review comments
Signed-off-by: Soumi Manna <[email protected]>
1 parent def5f8f commit 334cc23

File tree

3 files changed

+90
-33
lines changed

3 files changed

+90
-33
lines changed

‎clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp

+46-30
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
// 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
22

33
// Tests for IR of [[intel::scheduler_target_fmax_mhz()]], [[intel::num_simd_work_items()]],
4-
// [[intel::no_global_work_offset()]], [[intel::no_global_work_offset()]], [[sycl::reqd_sub_group_size()]],
4+
// [[intel::no_global_work_offset()]], [[intel::max_global_work_dim()]], [[sycl::reqd_sub_group_size()]],
55
// [[sycl::reqd_work_group_size()]], [[intel::kernel_args_restrict]], [[intel::max_work_group_size()]],
66
// and [[intel::sycl_explicit_simd]] function attributes in SYCL 2020.
77

@@ -73,7 +73,7 @@ class Foo4 {
7373
class Functor4 {
7474
public:
7575
void operator()() const {
76-
foo();
76+
foo4();
7777
}
7878
};
7979

@@ -146,134 +146,148 @@ class Functor10 {
146146

147147
int main() {
148148
q.submit([&](handler &h) {
149-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM1:[0-9]+]]
149+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 !kernel_arg_buffer_location ![[NUM:[0-9]+]] !scheduler_target_fmax_mhz ![[NUM1:[0-9]+]]
150150
Foo boo;
151151
h.single_task<class kernel_name1>(boo);
152152

153-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM42:[0-9]+]]
153+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 !kernel_arg_buffer_location ![[NUM]] !scheduler_target_fmax_mhz ![[NUM42:[0-9]+]]
154154
h.single_task<class kernel_name2>(
155155
[]() [[intel::scheduler_target_fmax_mhz(42)]]{});
156156

157-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM2:[0-9]+]]
157+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 !kernel_arg_buffer_location ![[NUM]] !scheduler_target_fmax_mhz ![[NUM2:[0-9]+]]
158158
Functor<2> f;
159159
h.single_task<class kernel_name3>(f);
160160

161161
// Test attribute is not propagated.
162-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4()
162+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 !kernel_arg_buffer_location ![[NUM]]
163163
// CHECK-NOT: !scheduler_target_fmax_mhz
164+
// CHECK-SAME: {
165+
// CHECK: define dso_local spir_func void @_Z3foov()
164166
h.single_task<class kernel_name4>(
165167
[]() { foo(); });
166168

167-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !num_simd_work_items ![[NUM1]]
169+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 !kernel_arg_buffer_location ![[NUM]] !num_simd_work_items ![[NUM1]]
168170
Foo1 boo1;
169171
h.single_task<class kernel_name5>(boo1);
170172

171-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6() #0 {{.*}} !num_simd_work_items ![[NUM42]]
173+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6() #0 !kernel_arg_buffer_location ![[NUM]] !num_simd_work_items ![[NUM42]]
172174
h.single_task<class kernel_name6>(
173175
[]() [[intel::num_simd_work_items(42)]]{});
174176

175-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name7() #0 {{.*}} !num_simd_work_items ![[NUM2]]
177+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name7() #0 !kernel_arg_buffer_location ![[NUM]] !num_simd_work_items ![[NUM2]]
176178
Functor1<2> f1;
177179
h.single_task<class kernel_name7>(f1);
178180

179181
// Test attribute is not propagated.
180-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name8()
182+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name8() #0 !kernel_arg_buffer_location ![[NUM]]
181183
// CHECK-NOT: !num_simd_work_items
184+
// CHECK-SAME: {
185+
// CHECK: define dso_local spir_func void @_Z4foo1v()
182186
h.single_task<class kernel_name8>(
183187
[]() { foo1(); });
184188

185-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name9() #0 {{.*}} !no_global_work_offset ![[NUM:[0-9]+]]
189+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name9() #0 !kernel_arg_buffer_location ![[NUM]] !no_global_work_offset ![[NUM:[0-9]+]]
186190
Foo2 boo2;
187191
h.single_task<class kernel_name9>(boo2);
188192

189193
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name10() #0 {{.*}} ![[NUM0:[0-9]+]]
190194
h.single_task<class kernel_name10>(
191195
[]() [[intel::no_global_work_offset(0)]]{});
192196

193-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name11() #0 {{.*}} !no_global_work_offset ![[NUM]]
197+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name11() #0 !kernel_arg_buffer_location ![[NUM]] !no_global_work_offset ![[NUM]]
194198
Functor2<1> f2;
195199
h.single_task<class kernel_name11>(f2);
196200

197201
// Test attribute is not propagated.
198-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name12()
202+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name12() #0 !kernel_arg_buffer_location ![[NUM]]
199203
// CHECK-NOT: !no_global_work_offset
204+
// CHECK-SAME: {
205+
// CHECK: define dso_local spir_func void @_Z4foo2v()
200206
h.single_task<class kernel_name12>(
201207
[]() { foo2(); });
202208

203-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name13() #0 {{.*}} !max_global_work_dim ![[NUM1]]
209+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name13() #0 !kernel_arg_buffer_location ![[NUM]] !max_global_work_dim ![[NUM1]]
204210
Foo3 boo3;
205211
h.single_task<class kernel_name13>(boo3);
206212

207-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name14() #0 {{.*}} !max_global_work_dim ![[NUM1]]
213+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name14() #0 !kernel_arg_buffer_location ![[NUM]] !max_global_work_dim ![[NUM1]]
208214
h.single_task<class kernel_name14>(
209215
[]() [[intel::max_global_work_dim(1)]]{});
210216

211-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name15() #0 {{.*}} !max_global_work_dim ![[NUM2]]
217+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name15() #0 !kernel_arg_buffer_location ![[NUM]] !max_global_work_dim ![[NUM2]]
212218
Functor3<2> f3;
213219
h.single_task<class kernel_name15>(f3);
214220

215221
// Test attribute is not propagated.
216-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name16()
222+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name16() #0 !kernel_arg_buffer_location ![[NUM]]
217223
// CHECK-NOT: !max_global_work_dim
224+
// CHECK-SAME: {
225+
// CHECK: define dso_local spir_func void @_Z4foo3v()
218226
h.single_task<class kernel_name16>(
219227
[]() { foo3(); });
220228

221-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name17() #0 {{.*}} !intel_reqd_sub_group_size ![[NUM16:[0-9]+]]
229+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name17() #0 !kernel_arg_buffer_location ![[NUM]] !intel_reqd_sub_group_size ![[NUM16:[0-9]+]]
222230
Foo4 boo4;
223231
h.single_task<class kernel_name17>(boo4);
224232

225-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name18() #0 {{.*}} !intel_reqd_sub_group_size ![[NUM1]]
233+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name18() #0 !kernel_arg_buffer_location ![[NUM]] !intel_reqd_sub_group_size ![[NUM1]]
226234
h.single_task<class kernel_name18>(
227235
[]() [[sycl::reqd_sub_group_size(1)]]{});
228236

229-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name19() #0 {{.*}} !intel_reqd_sub_group_size ![[NUM2]]
237+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name19() #0 !kernel_arg_buffer_location ![[NUM]] !intel_reqd_sub_group_size ![[NUM2]]
230238
Functor5<2> f5;
231239
h.single_task<class kernel_name19>(f5);
232240

233241
// Test attribute is not propagated.
234-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name20()
242+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name20() #0 !kernel_arg_buffer_location ![[NUM]]
235243
// CHECK-NOT: !reqd_sub_group_size
244+
// CHECK-SAME: {
245+
// CHECK: define dso_local spir_func void @_Z4foo4v()
236246
Functor4 f4;
237247
h.single_task<class kernel_name20>(f4);
238248

239-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name21() #0 {{.*}} !reqd_work_group_size ![[NUM32:[0-9]+]]
249+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name21() #0 !kernel_arg_buffer_location ![[NUM]] !reqd_work_group_size ![[NUM32:[0-9]+]]
240250
Foo5 boo5;
241251
h.single_task<class kernel_name21>(boo5);
242252

243-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name22() #0 {{.*}} !reqd_work_group_size ![[NUM88:[0-9]+]]
253+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name22() #0 !kernel_arg_buffer_location ![[NUM]] !reqd_work_group_size ![[NUM88:[0-9]+]]
244254
h.single_task<class kernel_name22>(
245255
[]() [[sycl::reqd_work_group_size(8, 8, 8)]]{});
246256

247-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name23() #0 {{.*}} !reqd_work_group_size ![[NUM22:[0-9]+]]
257+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name23() #0 !kernel_arg_buffer_location ![[NUM]] !reqd_work_group_size ![[NUM22:[0-9]+]]
248258
Functor7<2, 2, 2> f7;
249259
h.single_task<class kernel_name23>(f7);
250260

251261
// Test attribute is not propagated.
252-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name24()
262+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name24() #0 !kernel_arg_buffer_location ![[NUM]]
253263
// CHECK-NOT: !reqd_work_group_size
264+
// CHECK-SAME: {
265+
// CHECK: define dso_local spir_func void @_Z4foo5v()
254266
Functor6 f6;
255267
h.single_task<class kernel_name24>(f6);
256268

257-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name25() #0 {{.*}} !max_work_group_size ![[NUM32]]
269+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name25() #0 !kernel_arg_buffer_location ![[NUM]] !max_work_group_size ![[NUM32]]
258270
Foo6 boo6;
259271
h.single_task<class kernel_name25>(boo6);
260272

261-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name26() #0 {{.*}} !max_work_group_size ![[NUM88]]
273+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name26() #0 !kernel_arg_buffer_location ![[NUM]] !max_work_group_size ![[NUM88]]
262274
h.single_task<class kernel_name26>(
263275
[]() [[intel::max_work_group_size(8, 8, 8)]]{});
264276

265-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name27() #0 {{.*}} !max_work_group_size ![[NUM22]]
277+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name27() #0 !kernel_arg_buffer_location ![[NUM]] !max_work_group_size ![[NUM22]]
266278
Functor9<2, 2, 2> f9;
267279
h.single_task<class kernel_name27>(f9);
268280

269281
// Test attribute is not propagated.
270-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name28()
282+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name28() #0 !kernel_arg_buffer_location ![[NUM]]
271283
// CHECK-NOT: !max_work_group_size
284+
// CHECK-SAME: {
285+
// CHECK: define dso_local spir_func void @_Z4foo6v()
272286
Functor8 f8;
273287
h.single_task<class kernel_name28>(f8);
274288

275289
// Test attribute is not propagated.
276-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name29()
290+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name29() #0 !kernel_arg_buffer_location ![[NUM]]
277291
// CHECK-NOT: !sycl_explicit_simd
278292
// CHECK-SAME: {
279293
// CHECK: define {{.*}}spir_func void @{{.*}}foo7{{.*}} !sycl_explicit_simd ![[NUM]]
@@ -292,6 +306,8 @@ int main() {
292306
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name32() #0 !kernel_arg_buffer_location ![[NUM]]
293307
// CHECK: define {{.*}}spir_func void @{{.*}}Functor10{{.*}}(%class._ZTS9Functor10.Functor10 addrspace(4)* align 1 dereferenceable_or_null(1) %this) #3 comdat align 2
294308
// CHECK-NOT: noalias
309+
// CHECK-SAME: {
310+
// CHECK: define dso_local spir_func void @_Z4foo8v()
295311
Functor10 f10;
296312
h.single_task<class kernel_name32>(f10);
297313

‎clang/test/SemaSYCL/check-direct-attribute-propagation.cpp

+3-3
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -ast-dump %s | FileCheck %s
22

33
// Tests for AST of [[intel::scheduler_target_fmax_mhz()]], [[intel::num_simd_work_items()]],
4-
// [[intel::no_global_work_offset()]], [[intel::no_global_work_offset()]], [[intel::sycl_explicit_simd]],
4+
// [[intel::no_global_work_offset()]], [[intel::max_global_work_dim()]], [[intel::sycl_explicit_simd]],
55
// [[sycl::reqd_sub_group_size()]], [[sycl::reqd_work_group_size()]], [[intel::kernel_args_restrict]], and
66
// [[intel::max_work_group_size()]] function attributes in SYCL 2020.
77

@@ -109,12 +109,12 @@ int main() {
109109
h.single_task<class test_kernel5>(
110110
FuncObj1());
111111

112-
// CHECK: FunctionDecl {{.*}}test_kerne6
112+
// CHECK: FunctionDecl {{.*}}test_kernel6
113113
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr
114114
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
115115
// CHECK-NEXT: value: Int 1
116116
// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}}
117-
h.single_task<class test_kerne6>(
117+
h.single_task<class test_kernel6>(
118118
[]() [[intel::no_global_work_offset]]{});
119119

120120
// CHECK: FunctionDecl {{.*}}test_kernel7

‎clang/test/SemaSYCL/sycl-esimd-ast.cpp

+41
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,41 @@
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+
// Tests for AST of sycl_explicit_simd function attribute in SYCL 1.2.1.
4+
5+
#include "sycl.hpp"
6+
7+
sycl::queue deviceQueue;
8+
9+
struct FuncObj {
10+
[[intel::sycl_explicit_simd]] void operator()() const {}
11+
};
12+
13+
[[intel::sycl_explicit_simd]] void func() {}
14+
15+
int main() {
16+
deviceQueue.submit([&](sycl::handler &h) {
17+
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel1
18+
// CHECK: SYCLSimdAttr {{.*}} Implicit
19+
// CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit
20+
// CHECK-NEXT: SYCLSimdAttr {{.*}}
21+
h.single_task<class test_kernel1>(
22+
FuncObj());
23+
24+
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel2
25+
// CHECK: SYCLSimdAttr {{.*}} Implicit
26+
// CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit
27+
// CHECK-NEXT: SYCLSimdAttr {{.*}}
28+
h.single_task<class test_kernel2>(
29+
[]() [[intel::sycl_explicit_simd]]{});
30+
31+
// Test attribute is propagated.
32+
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
33+
// CHECK: SYCLSimdAttr {{.*}} Implicit
34+
// CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit
35+
// CHECK-NEXT: SYCLSimdAttr {{.*}}
36+
// CHECK-NEXT: SYCLSimdAttr {{.*}}
37+
h.single_task<class test_kernel3>(
38+
[]() [[intel::sycl_explicit_simd]] { func(); });
39+
});
40+
return 0;
41+
}

0 commit comments

Comments
 (0)