Skip to content

Commit 64205ad

Browse files
wenju-hewenju-he
authored
[SPIR-V] Map SPIR-V friendly work-item function to built-in variables (#148567)
The mapping ensures the function is lowered to SPIR-V built-in variables in SPIR-V. This can fix pre-commit CI fail in #19359 Also add BuiltIn to SPIR-V Builtin function name in __clang_spirv_builtins.h to align with https://github.com/llvm/llvm-project/blob/main/llvm/docs/SPIRVUsage.rst#builtin-variables
1 parent c0294f4 commit 64205ad

File tree

4 files changed

+186
-69
lines changed

4 files changed

+186
-69
lines changed

‎clang/lib/Headers/__clang_spirv_builtins.h

Lines changed: 12 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -52,30 +52,30 @@
5252
// Builtin IDs and sizes
5353

5454
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_num_workgroups) __size_t
55-
__spirv_NumWorkgroups(int);
55+
__spirv_BuiltInNumWorkgroups(int);
5656
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_workgroup_size) __size_t
57-
__spirv_WorkgroupSize(int);
57+
__spirv_BuiltInWorkgroupSize(int);
5858
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_workgroup_id) __size_t
59-
__spirv_WorkgroupId(int);
59+
__spirv_BuiltInWorkgroupId(int);
6060
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_local_invocation_id) __size_t
61-
__spirv_LocalInvocationId(int);
61+
__spirv_BuiltInLocalInvocationId(int);
6262
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_invocation_id) __size_t
63-
__spirv_GlobalInvocationId(int);
63+
__spirv_BuiltInGlobalInvocationId(int);
6464

6565
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_size) __size_t
66-
__spirv_GlobalSize(int);
66+
__spirv_BuiltInGlobalSize(int);
6767
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_offset) __size_t
68-
__spirv_GlobalOffset(int);
68+
__spirv_BuiltInGlobalOffset(int);
6969
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_size) __uint32_t
70-
__spirv_SubgroupSize();
70+
__spirv_BuiltInSubgroupSize();
7171
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_max_size) __uint32_t
72-
__spirv_SubgroupMaxSize();
72+
__spirv_BuiltInSubgroupMaxSize();
7373
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_num_subgroups) __uint32_t
74-
__spirv_NumSubgroups();
74+
__spirv_BuiltInNumSubgroups();
7575
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_id) __uint32_t
76-
__spirv_SubgroupId();
76+
__spirv_BuiltInSubgroupId();
7777
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_local_invocation_id)
78-
__uint32_t __spirv_SubgroupLocalInvocationId();
78+
__uint32_t __spirv_BuiltInSubgroupLocalInvocationId();
7979

8080
// OpGenericCastToPtrExplicit
8181

‎clang/test/Headers/spirv_ids.cpp

Lines changed: 52 additions & 52 deletions
Original file line numberDiff line numberDiff line change
@@ -53,58 +53,58 @@
5353
// CHECK: call i32 @llvm.spv.subgroup.id()
5454
// CHECK: call i32 @llvm.spv.subgroup.local.invocation.id()
5555

56-
// NV: call noundef i64 @_Z21__spirv_NumWorkgroupsi(i32 noundef 0) #2
57-
// NV: call noundef i64 @_Z21__spirv_NumWorkgroupsi(i32 noundef 1) #2
58-
// NV: call noundef i64 @_Z21__spirv_NumWorkgroupsi(i32 noundef 2) #2
59-
// NV: call noundef i64 @_Z21__spirv_WorkgroupSizei(i32 noundef 0) #2
60-
// NV: call noundef i64 @_Z21__spirv_WorkgroupSizei(i32 noundef 1) #2
61-
// NV: call noundef i64 @_Z21__spirv_WorkgroupSizei(i32 noundef 2) #2
62-
// NV: call noundef i64 @_Z19__spirv_WorkgroupIdi(i32 noundef 0) #2
63-
// NV: call noundef i64 @_Z19__spirv_WorkgroupIdi(i32 noundef 1) #2
64-
// NV: call noundef i64 @_Z19__spirv_WorkgroupIdi(i32 noundef 2) #2
65-
// NV: call noundef i64 @_Z25__spirv_LocalInvocationIdi(i32 noundef 0) #2
66-
// NV: call noundef i64 @_Z25__spirv_LocalInvocationIdi(i32 noundef 1) #2
67-
// NV: call noundef i64 @_Z25__spirv_LocalInvocationIdi(i32 noundef 2) #2
68-
// NV: call noundef i64 @_Z26__spirv_GlobalInvocationIdi(i32 noundef 0) #2
69-
// NV: call noundef i64 @_Z26__spirv_GlobalInvocationIdi(i32 noundef 1) #2
70-
// NV: call noundef i64 @_Z26__spirv_GlobalInvocationIdi(i32 noundef 2) #2
71-
// NV: call noundef i64 @_Z18__spirv_GlobalSizei(i32 noundef 0) #2
72-
// NV: call noundef i64 @_Z18__spirv_GlobalSizei(i32 noundef 1) #2
73-
// NV: call noundef i64 @_Z18__spirv_GlobalSizei(i32 noundef 2) #2
74-
// NV: call noundef i64 @_Z20__spirv_GlobalOffseti(i32 noundef 0) #2
75-
// NV: call noundef i64 @_Z20__spirv_GlobalOffseti(i32 noundef 1) #2
76-
// NV: call noundef i64 @_Z20__spirv_GlobalOffseti(i32 noundef 2) #2
77-
// NV: call noundef i32 @_Z20__spirv_SubgroupSizev() #2
78-
// NV: call noundef i32 @_Z23__spirv_SubgroupMaxSizev() #2
79-
// NV: call noundef i32 @_Z20__spirv_NumSubgroupsv() #2
80-
// NV: call noundef i32 @_Z18__spirv_SubgroupIdv() #2
81-
// NV: call noundef i32 @_Z33__spirv_SubgroupLocalInvocationIdv() #2
56+
// NV: call noundef i64 @_Z28__spirv_BuiltInNumWorkgroupsi(i32 noundef 0) #2
57+
// NV: call noundef i64 @_Z28__spirv_BuiltInNumWorkgroupsi(i32 noundef 1) #2
58+
// NV: call noundef i64 @_Z28__spirv_BuiltInNumWorkgroupsi(i32 noundef 2) #2
59+
// NV: call noundef i64 @_Z28__spirv_BuiltInWorkgroupSizei(i32 noundef 0) #2
60+
// NV: call noundef i64 @_Z28__spirv_BuiltInWorkgroupSizei(i32 noundef 1) #2
61+
// NV: call noundef i64 @_Z28__spirv_BuiltInWorkgroupSizei(i32 noundef 2) #2
62+
// NV: call noundef i64 @_Z26__spirv_BuiltInWorkgroupIdi(i32 noundef 0) #2
63+
// NV: call noundef i64 @_Z26__spirv_BuiltInWorkgroupIdi(i32 noundef 1) #2
64+
// NV: call noundef i64 @_Z26__spirv_BuiltInWorkgroupIdi(i32 noundef 2) #2
65+
// NV: call noundef i64 @_Z32__spirv_BuiltInLocalInvocationIdi(i32 noundef 0) #2
66+
// NV: call noundef i64 @_Z32__spirv_BuiltInLocalInvocationIdi(i32 noundef 1) #2
67+
// NV: call noundef i64 @_Z32__spirv_BuiltInLocalInvocationIdi(i32 noundef 2) #2
68+
// NV: call noundef i64 @_Z33__spirv_BuiltInGlobalInvocationIdi(i32 noundef 0) #2
69+
// NV: call noundef i64 @_Z33__spirv_BuiltInGlobalInvocationIdi(i32 noundef 1) #2
70+
// NV: call noundef i64 @_Z33__spirv_BuiltInGlobalInvocationIdi(i32 noundef 2) #2
71+
// NV: call noundef i64 @_Z25__spirv_BuiltInGlobalSizei(i32 noundef 0) #2
72+
// NV: call noundef i64 @_Z25__spirv_BuiltInGlobalSizei(i32 noundef 1) #2
73+
// NV: call noundef i64 @_Z25__spirv_BuiltInGlobalSizei(i32 noundef 2) #2
74+
// NV: call noundef i64 @_Z27__spirv_BuiltInGlobalOffseti(i32 noundef 0) #2
75+
// NV: call noundef i64 @_Z27__spirv_BuiltInGlobalOffseti(i32 noundef 1) #2
76+
// NV: call noundef i64 @_Z27__spirv_BuiltInGlobalOffseti(i32 noundef 2) #2
77+
// NV: call noundef i32 @_Z27__spirv_BuiltInSubgroupSizev() #2
78+
// NV: call noundef i32 @_Z30__spirv_BuiltInSubgroupMaxSizev() #2
79+
// NV: call noundef i32 @_Z27__spirv_BuiltInNumSubgroupsv() #2
80+
// NV: call noundef i32 @_Z25__spirv_BuiltInSubgroupIdv() #2
81+
// NV: call noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #2
8282

8383
void test_id_and_range() {
84-
__spirv_NumWorkgroups(0);
85-
__spirv_NumWorkgroups(1);
86-
__spirv_NumWorkgroups(2);
87-
__spirv_WorkgroupSize(0);
88-
__spirv_WorkgroupSize(1);
89-
__spirv_WorkgroupSize(2);
90-
__spirv_WorkgroupId(0);
91-
__spirv_WorkgroupId(1);
92-
__spirv_WorkgroupId(2);
93-
__spirv_LocalInvocationId(0);
94-
__spirv_LocalInvocationId(1);
95-
__spirv_LocalInvocationId(2);
96-
__spirv_GlobalInvocationId(0);
97-
__spirv_GlobalInvocationId(1);
98-
__spirv_GlobalInvocationId(2);
99-
__spirv_GlobalSize(0);
100-
__spirv_GlobalSize(1);
101-
__spirv_GlobalSize(2);
102-
__spirv_GlobalOffset(0);
103-
__spirv_GlobalOffset(1);
104-
__spirv_GlobalOffset(2);
105-
unsigned int ssize = __spirv_SubgroupSize();
106-
unsigned int smax = __spirv_SubgroupMaxSize();
107-
unsigned int snum = __spirv_NumSubgroups();
108-
unsigned int sid = __spirv_SubgroupId();
109-
unsigned int sinvocid = __spirv_SubgroupLocalInvocationId();
84+
__spirv_BuiltInNumWorkgroups(0);
85+
__spirv_BuiltInNumWorkgroups(1);
86+
__spirv_BuiltInNumWorkgroups(2);
87+
__spirv_BuiltInWorkgroupSize(0);
88+
__spirv_BuiltInWorkgroupSize(1);
89+
__spirv_BuiltInWorkgroupSize(2);
90+
__spirv_BuiltInWorkgroupId(0);
91+
__spirv_BuiltInWorkgroupId(1);
92+
__spirv_BuiltInWorkgroupId(2);
93+
__spirv_BuiltInLocalInvocationId(0);
94+
__spirv_BuiltInLocalInvocationId(1);
95+
__spirv_BuiltInLocalInvocationId(2);
96+
__spirv_BuiltInGlobalInvocationId(0);
97+
__spirv_BuiltInGlobalInvocationId(1);
98+
__spirv_BuiltInGlobalInvocationId(2);
99+
__spirv_BuiltInGlobalSize(0);
100+
__spirv_BuiltInGlobalSize(1);
101+
__spirv_BuiltInGlobalSize(2);
102+
__spirv_BuiltInGlobalOffset(0);
103+
__spirv_BuiltInGlobalOffset(1);
104+
__spirv_BuiltInGlobalOffset(2);
105+
unsigned int ssize = __spirv_BuiltInSubgroupSize();
106+
unsigned int smax = __spirv_BuiltInSubgroupMaxSize();
107+
unsigned int snum = __spirv_BuiltInNumSubgroups();
108+
unsigned int sid = __spirv_BuiltInSubgroupId();
109+
unsigned int sinvocid = __spirv_BuiltInSubgroupLocalInvocationId();
110110
}

‎llvm/lib/Target/SPIRV/SPIRVBuiltins.td

Lines changed: 26 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1364,7 +1364,24 @@ defm : DemangledGetBuiltin<"get_sub_group_gt_mask", OpenCL_std, Variable, Subgro
13641364
defm : DemangledGetBuiltin<"get_sub_group_le_mask", OpenCL_std, Variable, SubgroupLeMask>;
13651365
defm : DemangledGetBuiltin<"get_sub_group_lt_mask", OpenCL_std, Variable, SubgroupLtMask>;
13661366
defm : DemangledGetBuiltin<"__spirv_BuiltInGlobalLinearId", OpenCL_std, Variable, GlobalLinearId>;
1367-
defm : DemangledGetBuiltin<"__spirv_BuiltInGlobalInvocationId", OpenCL_std, Variable, GlobalInvocationId>;
1367+
defm : DemangledGetBuiltin<"__spirv_BuiltInLocalInvocationIndex", OpenCL_std, Variable, LocalInvocationIndex>;
1368+
defm : DemangledGetBuiltin<"__spirv_BuiltInWorkDim", OpenCL_std, Variable, WorkDim>;
1369+
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupSize", OpenCL_std, Variable, SubgroupSize>;
1370+
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupMaxSize", OpenCL_std, Variable, SubgroupMaxSize>;
1371+
defm : DemangledGetBuiltin<"__spirv_BuiltInNumSubgroups", OpenCL_std, Variable, NumSubgroups>;
1372+
defm : DemangledGetBuiltin<"__spirv_BuiltInNumEnqueuedSubgroups", OpenCL_std, Variable, NumEnqueuedSubgroups>;
1373+
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupId", OpenCL_std, Variable, SubgroupId>;
1374+
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupLocalInvocationId", OpenCL_std, Variable, SubgroupLocalInvocationId>;
1375+
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupEqMask", OpenCL_std, Variable, SubgroupEqMask>;
1376+
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupEqMaskKHR", OpenCL_std, Variable, SubgroupEqMask>;
1377+
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupGeMask", OpenCL_std, Variable, SubgroupGeMask>;
1378+
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupGeMaskKHR", OpenCL_std, Variable, SubgroupGeMask>;
1379+
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupGtMask", OpenCL_std, Variable, SubgroupGtMask>;
1380+
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupGtMaskKHR", OpenCL_std, Variable, SubgroupGtMask>;
1381+
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupLeMask", OpenCL_std, Variable, SubgroupLeMask>;
1382+
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupLeMaskKHR", OpenCL_std, Variable, SubgroupLeMask>;
1383+
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupLtMask", OpenCL_std, Variable, SubgroupLtMask>;
1384+
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupLtMaskKHR", OpenCL_std, Variable, SubgroupLtMask>;
13681385

13691386
// GetQuery builtin records:
13701387
defm : DemangledGetBuiltin<"get_local_id", OpenCL_std, GetQuery, LocalInvocationId>;
@@ -1375,6 +1392,14 @@ defm : DemangledGetBuiltin<"get_group_id", OpenCL_std, GetQuery, WorkgroupId>;
13751392
defm : DemangledGetBuiltin<"get_enqueued_local_size", OpenCL_std, GetQuery, EnqueuedWorkgroupSize>;
13761393
defm : DemangledGetBuiltin<"get_num_groups", OpenCL_std, GetQuery, NumWorkgroups>;
13771394
defm : DemangledGetBuiltin<"get_global_offset", OpenCL_std, GetQuery, GlobalOffset>;
1395+
defm : DemangledGetBuiltin<"__spirv_BuiltInLocalInvocationId", OpenCL_std, GetQuery, LocalInvocationId>;
1396+
defm : DemangledGetBuiltin<"__spirv_BuiltInGlobalInvocationId", OpenCL_std, GetQuery, GlobalInvocationId>;
1397+
defm : DemangledGetBuiltin<"__spirv_BuiltInWorkgroupSize", OpenCL_std, GetQuery, WorkgroupSize>;
1398+
defm : DemangledGetBuiltin<"__spirv_BuiltInGlobalSize", OpenCL_std, GetQuery, GlobalSize>;
1399+
defm : DemangledGetBuiltin<"__spirv_BuiltInWorkgroupId", OpenCL_std, GetQuery, WorkgroupId>;
1400+
defm : DemangledGetBuiltin<"__spirv_BuiltInEnqueuedWorkgroupSize", OpenCL_std, GetQuery, EnqueuedWorkgroupSize>;
1401+
defm : DemangledGetBuiltin<"__spirv_BuiltInNumWorkgroups", OpenCL_std, GetQuery, NumWorkgroups>;
1402+
defm : DemangledGetBuiltin<"__spirv_BuiltInGlobalOffset", OpenCL_std, GetQuery, GlobalOffset>;
13781403
defm : DemangledGetBuiltin<"__hlsl_wave_get_lane_index", GLSL_std_450, Wave, SubgroupLocalInvocationId>;
13791404

13801405
//===----------------------------------------------------------------------===//

‎llvm/test/CodeGen/SPIRV/transcoding/builtin_calls.ll

Lines changed: 96 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,17 +1,109 @@
11
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
22
; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
33

4-
; CHECK-SPIRV-DAG: OpDecorate %[[#Id:]] BuiltIn GlobalInvocationId
5-
; CHECK-SPIRV-DAG: OpDecorate %[[#Id:]] BuiltIn GlobalLinearId
6-
; CHECK-SPIRV: %[[#Id:]] = OpVariable %[[#]]
7-
; CHECK-SPIRV: %[[#Id:]] = OpVariable %[[#]]
4+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id0:]] BuiltIn GlobalLinearId
5+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id1:]] BuiltIn GlobalInvocationId
6+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id2:]] BuiltIn LocalInvocationIndex
7+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id3:]] BuiltIn WorkDim
8+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id4:]] BuiltIn SubgroupSize
9+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id5:]] BuiltIn SubgroupMaxSize
10+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id6:]] BuiltIn NumSubgroups
11+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id7:]] BuiltIn NumEnqueuedSubgroups
12+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id8:]] BuiltIn SubgroupId
13+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id9:]] BuiltIn SubgroupLocalInvocationId
14+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id10:]] BuiltIn SubgroupEqMask
15+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id11:]] BuiltIn SubgroupGeMask
16+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id12:]] BuiltIn SubgroupGtMask
17+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id13:]] BuiltIn SubgroupLeMask
18+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id14:]] BuiltIn SubgroupLtMask
19+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id15:]] BuiltIn LocalInvocationId
20+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id16:]] BuiltIn WorkgroupSize
21+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id17:]] BuiltIn GlobalSize
22+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id18:]] BuiltIn WorkgroupId
23+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id19:]] BuiltIn EnqueuedWorkgroupSize
24+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id20:]] BuiltIn NumWorkgroups
25+
; CHECK-SPIRV-DAG: OpDecorate %[[#Id21:]] BuiltIn GlobalOffset
26+
27+
; CHECK-SPIRV: %[[#Id0:]] = OpVariable %[[#]] Input
28+
; CHECK-SPIRV: %[[#Id1:]] = OpVariable %[[#]] Input
29+
; CHECK-SPIRV: %[[#Id2:]] = OpVariable %[[#]] Input
30+
; CHECK-SPIRV: %[[#Id3:]] = OpVariable %[[#]] Input
31+
; CHECK-SPIRV: %[[#Id4:]] = OpVariable %[[#]] Input
32+
; CHECK-SPIRV: %[[#Id5:]] = OpVariable %[[#]] Input
33+
; CHECK-SPIRV: %[[#Id6:]] = OpVariable %[[#]] Input
34+
; CHECK-SPIRV: %[[#Id7:]] = OpVariable %[[#]] Input
35+
; CHECK-SPIRV: %[[#Id8:]] = OpVariable %[[#]] Input
36+
; CHECK-SPIRV: %[[#Id9:]] = OpVariable %[[#]] Input
37+
; CHECK-SPIRV: %[[#Id10:]] = OpVariable %[[#]] Input
38+
; CHECK-SPIRV: %[[#Id11:]] = OpVariable %[[#]] Input
39+
; CHECK-SPIRV: %[[#Id12:]] = OpVariable %[[#]] Input
40+
; CHECK-SPIRV: %[[#Id13:]] = OpVariable %[[#]] Input
41+
; CHECK-SPIRV: %[[#Id14:]] = OpVariable %[[#]] Input
42+
; CHECK-SPIRV: %[[#Id15:]] = OpVariable %[[#]] Input
43+
; CHECK-SPIRV: %[[#Id16:]] = OpVariable %[[#]] Input
44+
; CHECK-SPIRV: %[[#Id17:]] = OpVariable %[[#]] Input
45+
; CHECK-SPIRV: %[[#Id18:]] = OpVariable %[[#]] Input
46+
; CHECK-SPIRV: %[[#Id19:]] = OpVariable %[[#]] Input
47+
; CHECK-SPIRV: %[[#Id20:]] = OpVariable %[[#]] Input
48+
; CHECK-SPIRV: %[[#Id21:]] = OpVariable %[[#]] Input
849

950
define spir_kernel void @f() {
1051
entry:
1152
%0 = call spir_func i32 @_Z29__spirv_BuiltInGlobalLinearIdv()
1253
%1 = call spir_func i64 @_Z33__spirv_BuiltInGlobalInvocationIdi(i32 1)
54+
%2 = call spir_func i64 @_Z35__spirv_BuiltInLocalInvocationIndexv()
55+
%3 = call spir_func i32 @_Z22__spirv_BuiltInWorkDimv()
56+
%4 = call spir_func i32 @_Z27__spirv_BuiltInSubgroupSizev()
57+
%5 = call spir_func i32 @_Z30__spirv_BuiltInSubgroupMaxSizev()
58+
%6 = call spir_func i32 @_Z27__spirv_BuiltInNumSubgroupsv()
59+
%7 = call spir_func i32 @_Z35__spirv_BuiltInNumEnqueuedSubgroupsv()
60+
%8 = call spir_func i32 @_Z25__spirv_BuiltInSubgroupIdv()
61+
%9 = call spir_func i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv()
62+
%10 = call spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupEqMaskv()
63+
%11 = call spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupEqMaskKHRv()
64+
%12 = call spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupGeMaskv()
65+
%13 = call spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupGeMaskKHRv()
66+
%14 = call spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupGtMaskv()
67+
%15 = call spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupGtMaskKHRv()
68+
%16 = call spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupLeMaskv()
69+
%17 = call spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupLeMaskKHRv()
70+
%18 = call spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupLtMaskv()
71+
%19 = call spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupLtMaskKHRv()
72+
%20 = call spir_func i64 @_Z32__spirv_BuiltInLocalInvocationIdi(i32 0)
73+
%21 = call spir_func i64 @_Z28__spirv_BuiltInWorkgroupSizei(i32 0)
74+
%22 = call spir_func i64 @_Z25__spirv_BuiltInGlobalSizei(i32 0)
75+
%23 = call spir_func i64 @_Z26__spirv_BuiltInWorkgroupIdi(i32 0)
76+
%24 = call spir_func i64 @_Z36__spirv_BuiltInEnqueuedWorkgroupSizei(i32 0)
77+
%25 = call spir_func i64 @_Z28__spirv_BuiltInNumWorkgroupsi(i32 0)
78+
%26 = call spir_func i64 @_Z27__spirv_BuiltInGlobalOffseti(i32 0)
79+
1380
ret void
1481
}
1582

1683
declare spir_func i32 @_Z29__spirv_BuiltInGlobalLinearIdv()
1784
declare spir_func i64 @_Z33__spirv_BuiltInGlobalInvocationIdi(i32)
85+
declare spir_func i64 @_Z35__spirv_BuiltInLocalInvocationIndexv()
86+
declare spir_func i32 @_Z22__spirv_BuiltInWorkDimv()
87+
declare spir_func i32 @_Z27__spirv_BuiltInSubgroupSizev()
88+
declare spir_func i32 @_Z30__spirv_BuiltInSubgroupMaxSizev()
89+
declare spir_func i32 @_Z27__spirv_BuiltInNumSubgroupsv()
90+
declare spir_func i32 @_Z35__spirv_BuiltInNumEnqueuedSubgroupsv()
91+
declare spir_func i32 @_Z25__spirv_BuiltInSubgroupIdv()
92+
declare spir_func i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv()
93+
declare spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupEqMaskv()
94+
declare spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupEqMaskKHRv()
95+
declare spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupGeMaskv()
96+
declare spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupGeMaskKHRv()
97+
declare spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupGtMaskv()
98+
declare spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupGtMaskKHRv()
99+
declare spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupLeMaskv()
100+
declare spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupLeMaskKHRv()
101+
declare spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupLtMaskv()
102+
declare spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupLtMaskKHRv()
103+
declare spir_func i64 @_Z32__spirv_BuiltInLocalInvocationIdi(i32)
104+
declare spir_func i64 @_Z28__spirv_BuiltInWorkgroupSizei(i32)
105+
declare spir_func i64 @_Z25__spirv_BuiltInGlobalSizei(i32)
106+
declare spir_func i64 @_Z26__spirv_BuiltInWorkgroupIdi(i32)
107+
declare spir_func i64 @_Z36__spirv_BuiltInEnqueuedWorkgroupSizei(i32)
108+
declare spir_func i64 @_Z28__spirv_BuiltInNumWorkgroupsi(i32)
109+
declare spir_func i64 @_Z27__spirv_BuiltInGlobalOffseti(i32)

0 commit comments

Comments
 (0)