Skip to content

Commit 03fd104

Browse files
wenju-hewenju-he
authored
[SYCL] Add __attribute__((const)) to SPIR-V BuiltIn function declaration (#19674)
For SPIRV target, llvm-spirv adds memory(none) attribute to work-item function when translating back to LLVM IR. This PR adds the attribute for targets, that bypass SPIRV, as well.
1 parent 735b688 commit 03fd104

File tree

8 files changed

+629
-547
lines changed

8 files changed

+629
-547
lines changed

‎sycl/include/sycl/__spirv/spirv_vars.hpp

Lines changed: 24 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -17,19 +17,31 @@
1717

1818
// SPIR-V built-in variables mapped to function call.
1919

20-
__DPCPP_SYCL_EXTERNAL size_t __spirv_BuiltInGlobalInvocationId(int);
21-
__DPCPP_SYCL_EXTERNAL size_t __spirv_BuiltInGlobalSize(int);
22-
__DPCPP_SYCL_EXTERNAL size_t __spirv_BuiltInGlobalOffset(int);
23-
__DPCPP_SYCL_EXTERNAL size_t __spirv_BuiltInNumWorkgroups(int);
24-
__DPCPP_SYCL_EXTERNAL size_t __spirv_BuiltInWorkgroupSize(int);
25-
__DPCPP_SYCL_EXTERNAL size_t __spirv_BuiltInWorkgroupId(int);
26-
__DPCPP_SYCL_EXTERNAL size_t __spirv_BuiltInLocalInvocationId(int);
20+
__DPCPP_SYCL_EXTERNAL __attribute__((const)) size_t
21+
__spirv_BuiltInGlobalInvocationId(int);
22+
__DPCPP_SYCL_EXTERNAL __attribute__((const)) size_t
23+
__spirv_BuiltInGlobalSize(int);
24+
__DPCPP_SYCL_EXTERNAL __attribute__((const)) size_t
25+
__spirv_BuiltInGlobalOffset(int);
26+
__DPCPP_SYCL_EXTERNAL __attribute__((const)) size_t
27+
__spirv_BuiltInNumWorkgroups(int);
28+
__DPCPP_SYCL_EXTERNAL __attribute__((const)) size_t
29+
__spirv_BuiltInWorkgroupSize(int);
30+
__DPCPP_SYCL_EXTERNAL __attribute__((const)) size_t
31+
__spirv_BuiltInWorkgroupId(int);
32+
__DPCPP_SYCL_EXTERNAL __attribute__((const)) size_t
33+
__spirv_BuiltInLocalInvocationId(int);
2734

28-
__DPCPP_SYCL_EXTERNAL uint32_t __spirv_BuiltInSubgroupSize();
29-
__DPCPP_SYCL_EXTERNAL uint32_t __spirv_BuiltInSubgroupMaxSize();
30-
__DPCPP_SYCL_EXTERNAL uint32_t __spirv_BuiltInNumSubgroups();
31-
__DPCPP_SYCL_EXTERNAL uint32_t __spirv_BuiltInSubgroupId();
32-
__DPCPP_SYCL_EXTERNAL uint32_t __spirv_BuiltInSubgroupLocalInvocationId();
35+
__DPCPP_SYCL_EXTERNAL __attribute__((const)) uint32_t
36+
__spirv_BuiltInSubgroupSize();
37+
__DPCPP_SYCL_EXTERNAL __attribute__((const)) uint32_t
38+
__spirv_BuiltInSubgroupMaxSize();
39+
__DPCPP_SYCL_EXTERNAL __attribute__((const)) uint32_t
40+
__spirv_BuiltInNumSubgroups();
41+
__DPCPP_SYCL_EXTERNAL __attribute__((const)) uint32_t
42+
__spirv_BuiltInSubgroupId();
43+
__DPCPP_SYCL_EXTERNAL __attribute__((const)) uint32_t
44+
__spirv_BuiltInSubgroupLocalInvocationId();
3345

3446
namespace __spirv {
3547

‎sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp

Lines changed: 20 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -25,11 +25,11 @@ int main() {
2525
auto Redu1 = sycl::reduction<int>(nullptr, sycl::plus<int>());
2626
auto Redu2 = sycl::reduction<float>(nullptr, sycl::multiplies<float>());
2727

28-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel0(){{.*}} #[[SGSizeAttr1:[0-9]+]]
28+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel0(){{.*}} #[[SGSizeAttr0:[0-9]+]]
2929
Q.single_task<class SGSizeKernel0>(Props, []() {});
30-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel1(){{.*}} #[[SGSizeAttr1]]
30+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel1(){{.*}} #[[SGSizeAttr0]]
3131
Q.single_task<class SGSizeKernel1>(Ev, Props, []() {});
32-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel2(){{.*}} #[[SGSizeAttr1]]
32+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel2(){{.*}} #[[SGSizeAttr0]]
3333
Q.single_task<class SGSizeKernel2>({Ev}, Props, []() {});
3434

3535
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel9(){{.*}} #[[SGSizeAttr2:[0-9]+]]
@@ -79,28 +79,28 @@ int main() {
7979
Q.parallel_for<class SGSizeKernel26>(R3, {Ev}, Props, Redu1,
8080
[](sycl::id<3>, auto &) {});
8181

82-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel27(){{.*}} #[[SGSizeAttr2]]
82+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel27(){{.*}} #[[SGSizeAttr6:[0-9]+]]
8383
Q.parallel_for<class SGSizeKernel27>(NDR1, Props, [](sycl::nd_item<1>) {});
84-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel28(){{.*}} #[[SGSizeAttr2]]
84+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel28(){{.*}} #[[SGSizeAttr6]]
8585
Q.parallel_for<class SGSizeKernel28>(NDR1, Ev, Props,
8686
[](sycl::nd_item<1>) {});
87-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel29(){{.*}} #[[SGSizeAttr2]]
87+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel29(){{.*}} #[[SGSizeAttr6]]
8888
Q.parallel_for<class SGSizeKernel29>(NDR1, {Ev}, Props,
8989
[](sycl::nd_item<1>) {});
90-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel30(){{.*}} #[[SGSizeAttr2]]
90+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel30(){{.*}} #[[SGSizeAttr6]]
9191
Q.parallel_for<class SGSizeKernel30>(NDR2, Props, [](sycl::nd_item<2>) {});
92-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel31(){{.*}} #[[SGSizeAttr2]]
92+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel31(){{.*}} #[[SGSizeAttr6]]
9393
Q.parallel_for<class SGSizeKernel31>(NDR2, Ev, Props,
9494
[](sycl::nd_item<2>) {});
95-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel32(){{.*}} #[[SGSizeAttr2]]
95+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel32(){{.*}} #[[SGSizeAttr6]]
9696
Q.parallel_for<class SGSizeKernel32>(NDR2, {Ev}, Props,
9797
[](sycl::nd_item<2>) {});
98-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel33(){{.*}} #[[SGSizeAttr2]]
98+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel33(){{.*}} #[[SGSizeAttr6]]
9999
Q.parallel_for<class SGSizeKernel33>(NDR3, Props, [](sycl::nd_item<3>) {});
100-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel34(){{.*}} #[[SGSizeAttr2]]
100+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel34(){{.*}} #[[SGSizeAttr6]]
101101
Q.parallel_for<class SGSizeKernel34>(NDR3, Ev, Props,
102102
[](sycl::nd_item<3>) {});
103-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel35(){{.*}} #[[SGSizeAttr2]]
103+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel35(){{.*}} #[[SGSizeAttr6]]
104104
Q.parallel_for<class SGSizeKernel35>(NDR3, {Ev}, Props,
105105
[](sycl::nd_item<3>) {});
106106

@@ -160,15 +160,15 @@ int main() {
160160
Q.parallel_for<class SGSizeKernel53>(NDR3, {Ev}, Props, Redu1, Redu2,
161161
[](sycl::nd_item<3>, auto &, auto &) {});
162162

163-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel54(){{.*}} #[[SGSizeAttr1]]
163+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel54(){{.*}} #[[SGSizeAttr0]]
164164
Q.submit([&](sycl::handler &CGH) {
165165
CGH.single_task<class SGSizeKernel54>(Props, []() {});
166166
});
167-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel55(){{.*}} #[[SGSizeAttr1]]
167+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel55(){{.*}} #[[SGSizeAttr0]]
168168
Q.submit([&](sycl::handler &CGH) {
169169
CGH.single_task<class SGSizeKernel55>(Props, []() {});
170170
});
171-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel56(){{.*}} #[[SGSizeAttr1]]
171+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel56(){{.*}} #[[SGSizeAttr0]]
172172
Q.submit([&](sycl::handler &CGH) {
173173
CGH.single_task<class SGSizeKernel56>(Props, []() {});
174174
});
@@ -202,17 +202,17 @@ int main() {
202202
[](sycl::id<3>, auto &) {});
203203
});
204204

205-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel63(){{.*}} #[[SGSizeAttr2]]
205+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel63(){{.*}} #[[SGSizeAttr6]]
206206
Q.submit([&](sycl::handler &CGH) {
207207
CGH.parallel_for<class SGSizeKernel63>(NDR1, Props,
208208
[](sycl::nd_item<1>) {});
209209
});
210-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel64(){{.*}} #[[SGSizeAttr2]]
210+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel64(){{.*}} #[[SGSizeAttr6]]
211211
Q.submit([&](sycl::handler &CGH) {
212212
CGH.parallel_for<class SGSizeKernel64>(NDR2, Props,
213213
[](sycl::nd_item<2>) {});
214214
});
215-
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel65(){{.*}} #[[SGSizeAttr2]]
215+
// CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel65(){{.*}} #[[SGSizeAttr6]]
216216
Q.submit([&](sycl::handler &CGH) {
217217
CGH.parallel_for<class SGSizeKernel65>(NDR3, Props,
218218
[](sycl::nd_item<3>) {});
@@ -275,5 +275,6 @@ int main() {
275275
return 0;
276276
}
277277

278-
// CHECK-IR: attributes #[[SGSizeAttr1]] = { {{.*}}"sycl-sub-group-size"="1"
278+
// CHECK-IR: attributes #[[SGSizeAttr0]] = { {{.*}}"sycl-sub-group-size"="1"
279279
// CHECK-IR: attributes #[[SGSizeAttr2]] = { {{.*}}"sycl-sub-group-size"="1"
280+
// CHECK-IR: attributes #[[SGSizeAttr6]] = { {{.*}}"sycl-sub-group-size"="1"

0 commit comments

Comments
 (0)