Skip to content

[SPIR-V] Map SPIR-V friendly work-item function to built-in variables #148567

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 4 commits into from
Jul 18, 2025
Merged

[SPIR-V] Map SPIR-V friendly work-item function to built-in variables #148567

Show file tree
Hide file tree
Changes from 3 commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
0d5ae36
[SPIR-V] Map SPIR-V friendly work-item function to built-in variables
wenju-he Jul 14, 2025
a4523a5
Variable -> GetQuery
wenju-he Jul 15, 2025
a55fd37
align SPIR-V function name in clang/lib/Headers/__clang_spirv_builtins.h
wenju-he Jul 16, 2025
79dd451
check storage class
wenju-he Jul 17, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
24 changes: 12 additions & 12 deletions clang/lib/Headers/__clang_spirv_builtins.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -52,30 +52,30 @@
// Builtin IDs and sizes

extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_num_workgroups) __size_t
__spirv_NumWorkgroups(int);
__spirv_BuiltInNumWorkgroups(int);
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_workgroup_size) __size_t
__spirv_WorkgroupSize(int);
__spirv_BuiltInWorkgroupSize(int);
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_workgroup_id) __size_t
__spirv_WorkgroupId(int);
__spirv_BuiltInWorkgroupId(int);
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_local_invocation_id) __size_t
__spirv_LocalInvocationId(int);
__spirv_BuiltInLocalInvocationId(int);
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_invocation_id) __size_t
__spirv_GlobalInvocationId(int);
__spirv_BuiltInGlobalInvocationId(int);

extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_size) __size_t
__spirv_GlobalSize(int);
__spirv_BuiltInGlobalSize(int);
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_offset) __size_t
__spirv_GlobalOffset(int);
__spirv_BuiltInGlobalOffset(int);
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_size) __uint32_t
__spirv_SubgroupSize();
__spirv_BuiltInSubgroupSize();
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_max_size) __uint32_t
__spirv_SubgroupMaxSize();
__spirv_BuiltInSubgroupMaxSize();
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_num_subgroups) __uint32_t
__spirv_NumSubgroups();
__spirv_BuiltInNumSubgroups();
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_id) __uint32_t
__spirv_SubgroupId();
__spirv_BuiltInSubgroupId();
extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_local_invocation_id)
__uint32_t __spirv_SubgroupLocalInvocationId();
__uint32_t __spirv_BuiltInSubgroupLocalInvocationId();

// OpGenericCastToPtrExplicit

Expand Down
104 changes: 52 additions & 52 deletions clang/test/Headers/spirv_ids.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,58 +53,58 @@
// CHECK: call i32 @llvm.spv.subgroup.id()
// CHECK: call i32 @llvm.spv.subgroup.local.invocation.id()

// NV: call noundef i64 @_Z21__spirv_NumWorkgroupsi(i32 noundef 0) #2
// NV: call noundef i64 @_Z21__spirv_NumWorkgroupsi(i32 noundef 1) #2
// NV: call noundef i64 @_Z21__spirv_NumWorkgroupsi(i32 noundef 2) #2
// NV: call noundef i64 @_Z21__spirv_WorkgroupSizei(i32 noundef 0) #2
// NV: call noundef i64 @_Z21__spirv_WorkgroupSizei(i32 noundef 1) #2
// NV: call noundef i64 @_Z21__spirv_WorkgroupSizei(i32 noundef 2) #2
// NV: call noundef i64 @_Z19__spirv_WorkgroupIdi(i32 noundef 0) #2
// NV: call noundef i64 @_Z19__spirv_WorkgroupIdi(i32 noundef 1) #2
// NV: call noundef i64 @_Z19__spirv_WorkgroupIdi(i32 noundef 2) #2
// NV: call noundef i64 @_Z25__spirv_LocalInvocationIdi(i32 noundef 0) #2
// NV: call noundef i64 @_Z25__spirv_LocalInvocationIdi(i32 noundef 1) #2
// NV: call noundef i64 @_Z25__spirv_LocalInvocationIdi(i32 noundef 2) #2
// NV: call noundef i64 @_Z26__spirv_GlobalInvocationIdi(i32 noundef 0) #2
// NV: call noundef i64 @_Z26__spirv_GlobalInvocationIdi(i32 noundef 1) #2
// NV: call noundef i64 @_Z26__spirv_GlobalInvocationIdi(i32 noundef 2) #2
// NV: call noundef i64 @_Z18__spirv_GlobalSizei(i32 noundef 0) #2
// NV: call noundef i64 @_Z18__spirv_GlobalSizei(i32 noundef 1) #2
// NV: call noundef i64 @_Z18__spirv_GlobalSizei(i32 noundef 2) #2
// NV: call noundef i64 @_Z20__spirv_GlobalOffseti(i32 noundef 0) #2
// NV: call noundef i64 @_Z20__spirv_GlobalOffseti(i32 noundef 1) #2
// NV: call noundef i64 @_Z20__spirv_GlobalOffseti(i32 noundef 2) #2
// NV: call noundef i32 @_Z20__spirv_SubgroupSizev() #2
// NV: call noundef i32 @_Z23__spirv_SubgroupMaxSizev() #2
// NV: call noundef i32 @_Z20__spirv_NumSubgroupsv() #2
// NV: call noundef i32 @_Z18__spirv_SubgroupIdv() #2
// NV: call noundef i32 @_Z33__spirv_SubgroupLocalInvocationIdv() #2
// NV: call noundef i64 @_Z28__spirv_BuiltInNumWorkgroupsi(i32 noundef 0) #2
// NV: call noundef i64 @_Z28__spirv_BuiltInNumWorkgroupsi(i32 noundef 1) #2
// NV: call noundef i64 @_Z28__spirv_BuiltInNumWorkgroupsi(i32 noundef 2) #2
// NV: call noundef i64 @_Z28__spirv_BuiltInWorkgroupSizei(i32 noundef 0) #2
// NV: call noundef i64 @_Z28__spirv_BuiltInWorkgroupSizei(i32 noundef 1) #2
// NV: call noundef i64 @_Z28__spirv_BuiltInWorkgroupSizei(i32 noundef 2) #2
// NV: call noundef i64 @_Z26__spirv_BuiltInWorkgroupIdi(i32 noundef 0) #2
// NV: call noundef i64 @_Z26__spirv_BuiltInWorkgroupIdi(i32 noundef 1) #2
// NV: call noundef i64 @_Z26__spirv_BuiltInWorkgroupIdi(i32 noundef 2) #2
// NV: call noundef i64 @_Z32__spirv_BuiltInLocalInvocationIdi(i32 noundef 0) #2
// NV: call noundef i64 @_Z32__spirv_BuiltInLocalInvocationIdi(i32 noundef 1) #2
// NV: call noundef i64 @_Z32__spirv_BuiltInLocalInvocationIdi(i32 noundef 2) #2
// NV: call noundef i64 @_Z33__spirv_BuiltInGlobalInvocationIdi(i32 noundef 0) #2
// NV: call noundef i64 @_Z33__spirv_BuiltInGlobalInvocationIdi(i32 noundef 1) #2
// NV: call noundef i64 @_Z33__spirv_BuiltInGlobalInvocationIdi(i32 noundef 2) #2
// NV: call noundef i64 @_Z25__spirv_BuiltInGlobalSizei(i32 noundef 0) #2
// NV: call noundef i64 @_Z25__spirv_BuiltInGlobalSizei(i32 noundef 1) #2
// NV: call noundef i64 @_Z25__spirv_BuiltInGlobalSizei(i32 noundef 2) #2
// NV: call noundef i64 @_Z27__spirv_BuiltInGlobalOffseti(i32 noundef 0) #2
// NV: call noundef i64 @_Z27__spirv_BuiltInGlobalOffseti(i32 noundef 1) #2
// NV: call noundef i64 @_Z27__spirv_BuiltInGlobalOffseti(i32 noundef 2) #2
// NV: call noundef i32 @_Z27__spirv_BuiltInSubgroupSizev() #2
// NV: call noundef i32 @_Z30__spirv_BuiltInSubgroupMaxSizev() #2
// NV: call noundef i32 @_Z27__spirv_BuiltInNumSubgroupsv() #2
// NV: call noundef i32 @_Z25__spirv_BuiltInSubgroupIdv() #2
// NV: call noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #2

void test_id_and_range() {
__spirv_NumWorkgroups(0);
__spirv_NumWorkgroups(1);
__spirv_NumWorkgroups(2);
__spirv_WorkgroupSize(0);
__spirv_WorkgroupSize(1);
__spirv_WorkgroupSize(2);
__spirv_WorkgroupId(0);
__spirv_WorkgroupId(1);
__spirv_WorkgroupId(2);
__spirv_LocalInvocationId(0);
__spirv_LocalInvocationId(1);
__spirv_LocalInvocationId(2);
__spirv_GlobalInvocationId(0);
__spirv_GlobalInvocationId(1);
__spirv_GlobalInvocationId(2);
__spirv_GlobalSize(0);
__spirv_GlobalSize(1);
__spirv_GlobalSize(2);
__spirv_GlobalOffset(0);
__spirv_GlobalOffset(1);
__spirv_GlobalOffset(2);
unsigned int ssize = __spirv_SubgroupSize();
unsigned int smax = __spirv_SubgroupMaxSize();
unsigned int snum = __spirv_NumSubgroups();
unsigned int sid = __spirv_SubgroupId();
unsigned int sinvocid = __spirv_SubgroupLocalInvocationId();
__spirv_BuiltInNumWorkgroups(0);
__spirv_BuiltInNumWorkgroups(1);
__spirv_BuiltInNumWorkgroups(2);
__spirv_BuiltInWorkgroupSize(0);
__spirv_BuiltInWorkgroupSize(1);
__spirv_BuiltInWorkgroupSize(2);
__spirv_BuiltInWorkgroupId(0);
__spirv_BuiltInWorkgroupId(1);
__spirv_BuiltInWorkgroupId(2);
__spirv_BuiltInLocalInvocationId(0);
__spirv_BuiltInLocalInvocationId(1);
__spirv_BuiltInLocalInvocationId(2);
__spirv_BuiltInGlobalInvocationId(0);
__spirv_BuiltInGlobalInvocationId(1);
__spirv_BuiltInGlobalInvocationId(2);
__spirv_BuiltInGlobalSize(0);
__spirv_BuiltInGlobalSize(1);
__spirv_BuiltInGlobalSize(2);
__spirv_BuiltInGlobalOffset(0);
__spirv_BuiltInGlobalOffset(1);
__spirv_BuiltInGlobalOffset(2);
unsigned int ssize = __spirv_BuiltInSubgroupSize();
unsigned int smax = __spirv_BuiltInSubgroupMaxSize();
unsigned int snum = __spirv_BuiltInNumSubgroups();
unsigned int sid = __spirv_BuiltInSubgroupId();
unsigned int sinvocid = __spirv_BuiltInSubgroupLocalInvocationId();
}
27 changes: 26 additions & 1 deletion llvm/lib/Target/SPIRV/SPIRVBuiltins.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1364,7 +1364,24 @@ defm : DemangledGetBuiltin<"get_sub_group_gt_mask", OpenCL_std, Variable, Subgro
defm : DemangledGetBuiltin<"get_sub_group_le_mask", OpenCL_std, Variable, SubgroupLeMask>;
defm : DemangledGetBuiltin<"get_sub_group_lt_mask", OpenCL_std, Variable, SubgroupLtMask>;
defm : DemangledGetBuiltin<"__spirv_BuiltInGlobalLinearId", OpenCL_std, Variable, GlobalLinearId>;
defm : DemangledGetBuiltin<"__spirv_BuiltInGlobalInvocationId", OpenCL_std, Variable, GlobalInvocationId>;
defm : DemangledGetBuiltin<"__spirv_BuiltInLocalInvocationIndex", OpenCL_std, Variable, LocalInvocationIndex>;
defm : DemangledGetBuiltin<"__spirv_BuiltInWorkDim", OpenCL_std, Variable, WorkDim>;
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupSize", OpenCL_std, Variable, SubgroupSize>;
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupMaxSize", OpenCL_std, Variable, SubgroupMaxSize>;
defm : DemangledGetBuiltin<"__spirv_BuiltInNumSubgroups", OpenCL_std, Variable, NumSubgroups>;
defm : DemangledGetBuiltin<"__spirv_BuiltInNumEnqueuedSubgroups", OpenCL_std, Variable, NumEnqueuedSubgroups>;
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupId", OpenCL_std, Variable, SubgroupId>;
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupLocalInvocationId", OpenCL_std, Variable, SubgroupLocalInvocationId>;
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupEqMask", OpenCL_std, Variable, SubgroupEqMask>;
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupEqMaskKHR", OpenCL_std, Variable, SubgroupEqMask>;
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupGeMask", OpenCL_std, Variable, SubgroupGeMask>;
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupGeMaskKHR", OpenCL_std, Variable, SubgroupGeMask>;
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupGtMask", OpenCL_std, Variable, SubgroupGtMask>;
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupGtMaskKHR", OpenCL_std, Variable, SubgroupGtMask>;
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupLeMask", OpenCL_std, Variable, SubgroupLeMask>;
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupLeMaskKHR", OpenCL_std, Variable, SubgroupLeMask>;
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupLtMask", OpenCL_std, Variable, SubgroupLtMask>;
defm : DemangledGetBuiltin<"__spirv_BuiltInSubgroupLtMaskKHR", OpenCL_std, Variable, SubgroupLtMask>;

// GetQuery builtin records:
defm : DemangledGetBuiltin<"get_local_id", OpenCL_std, GetQuery, LocalInvocationId>;
Expand All @@ -1375,6 +1392,14 @@ defm : DemangledGetBuiltin<"get_group_id", OpenCL_std, GetQuery, WorkgroupId>;
defm : DemangledGetBuiltin<"get_enqueued_local_size", OpenCL_std, GetQuery, EnqueuedWorkgroupSize>;
defm : DemangledGetBuiltin<"get_num_groups", OpenCL_std, GetQuery, NumWorkgroups>;
defm : DemangledGetBuiltin<"get_global_offset", OpenCL_std, GetQuery, GlobalOffset>;
defm : DemangledGetBuiltin<"__spirv_BuiltInLocalInvocationId", OpenCL_std, GetQuery, LocalInvocationId>;
defm : DemangledGetBuiltin<"__spirv_BuiltInGlobalInvocationId", OpenCL_std, GetQuery, GlobalInvocationId>;
defm : DemangledGetBuiltin<"__spirv_BuiltInWorkgroupSize", OpenCL_std, GetQuery, WorkgroupSize>;
defm : DemangledGetBuiltin<"__spirv_BuiltInGlobalSize", OpenCL_std, GetQuery, GlobalSize>;
defm : DemangledGetBuiltin<"__spirv_BuiltInWorkgroupId", OpenCL_std, GetQuery, WorkgroupId>;
defm : DemangledGetBuiltin<"__spirv_BuiltInEnqueuedWorkgroupSize", OpenCL_std, GetQuery, EnqueuedWorkgroupSize>;
defm : DemangledGetBuiltin<"__spirv_BuiltInNumWorkgroups", OpenCL_std, GetQuery, NumWorkgroups>;
defm : DemangledGetBuiltin<"__spirv_BuiltInGlobalOffset", OpenCL_std, GetQuery, GlobalOffset>;
defm : DemangledGetBuiltin<"__hlsl_wave_get_lane_index", GLSL_std_450, Wave, SubgroupLocalInvocationId>;

//===----------------------------------------------------------------------===//
Expand Down
100 changes: 96 additions & 4 deletions llvm/test/CodeGen/SPIRV/transcoding/builtin_calls.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,17 +1,109 @@
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}

; CHECK-SPIRV-DAG: OpDecorate %[[#Id:]] BuiltIn GlobalInvocationId
; CHECK-SPIRV-DAG: OpDecorate %[[#Id:]] BuiltIn GlobalLinearId
; CHECK-SPIRV: %[[#Id:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id:]] = OpVariable %[[#]]
; CHECK-SPIRV-DAG: OpDecorate %[[#Id0:]] BuiltIn GlobalLinearId
; CHECK-SPIRV-DAG: OpDecorate %[[#Id1:]] BuiltIn GlobalInvocationId
; CHECK-SPIRV-DAG: OpDecorate %[[#Id2:]] BuiltIn LocalInvocationIndex
; CHECK-SPIRV-DAG: OpDecorate %[[#Id3:]] BuiltIn WorkDim
; CHECK-SPIRV-DAG: OpDecorate %[[#Id4:]] BuiltIn SubgroupSize
; CHECK-SPIRV-DAG: OpDecorate %[[#Id5:]] BuiltIn SubgroupMaxSize
; CHECK-SPIRV-DAG: OpDecorate %[[#Id6:]] BuiltIn NumSubgroups
; CHECK-SPIRV-DAG: OpDecorate %[[#Id7:]] BuiltIn NumEnqueuedSubgroups
; CHECK-SPIRV-DAG: OpDecorate %[[#Id8:]] BuiltIn SubgroupId
; CHECK-SPIRV-DAG: OpDecorate %[[#Id9:]] BuiltIn SubgroupLocalInvocationId
; CHECK-SPIRV-DAG: OpDecorate %[[#Id10:]] BuiltIn SubgroupEqMask
; CHECK-SPIRV-DAG: OpDecorate %[[#Id11:]] BuiltIn SubgroupGeMask
; CHECK-SPIRV-DAG: OpDecorate %[[#Id12:]] BuiltIn SubgroupGtMask
; CHECK-SPIRV-DAG: OpDecorate %[[#Id13:]] BuiltIn SubgroupLeMask
; CHECK-SPIRV-DAG: OpDecorate %[[#Id14:]] BuiltIn SubgroupLtMask
; CHECK-SPIRV-DAG: OpDecorate %[[#Id15:]] BuiltIn LocalInvocationId
; CHECK-SPIRV-DAG: OpDecorate %[[#Id16:]] BuiltIn WorkgroupSize
; CHECK-SPIRV-DAG: OpDecorate %[[#Id17:]] BuiltIn GlobalSize
; CHECK-SPIRV-DAG: OpDecorate %[[#Id18:]] BuiltIn WorkgroupId
; CHECK-SPIRV-DAG: OpDecorate %[[#Id19:]] BuiltIn EnqueuedWorkgroupSize
; CHECK-SPIRV-DAG: OpDecorate %[[#Id20:]] BuiltIn NumWorkgroups
; CHECK-SPIRV-DAG: OpDecorate %[[#Id21:]] BuiltIn GlobalOffset

; CHECK-SPIRV: %[[#Id0:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id1:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id2:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id3:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id4:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id5:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id6:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id7:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id8:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id9:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id10:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id11:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id12:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id13:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id14:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id15:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id16:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id17:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id18:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id19:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id20:]] = OpVariable %[[#]]
; CHECK-SPIRV: %[[#Id21:]] = OpVariable %[[#]]

define spir_kernel void @f() {
entry:
%0 = call spir_func i32 @_Z29__spirv_BuiltInGlobalLinearIdv()
%1 = call spir_func i64 @_Z33__spirv_BuiltInGlobalInvocationIdi(i32 1)
%2 = call spir_func i64 @_Z35__spirv_BuiltInLocalInvocationIndexv()
%3 = call spir_func i32 @_Z22__spirv_BuiltInWorkDimv()
%4 = call spir_func i32 @_Z27__spirv_BuiltInSubgroupSizev()
%5 = call spir_func i32 @_Z30__spirv_BuiltInSubgroupMaxSizev()
%6 = call spir_func i32 @_Z27__spirv_BuiltInNumSubgroupsv()
%7 = call spir_func i32 @_Z35__spirv_BuiltInNumEnqueuedSubgroupsv()
%8 = call spir_func i32 @_Z25__spirv_BuiltInSubgroupIdv()
%9 = call spir_func i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv()
%10 = call spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupEqMaskv()
%11 = call spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupEqMaskKHRv()
%12 = call spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupGeMaskv()
%13 = call spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupGeMaskKHRv()
%14 = call spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupGtMaskv()
%15 = call spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupGtMaskKHRv()
%16 = call spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupLeMaskv()
%17 = call spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupLeMaskKHRv()
%18 = call spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupLtMaskv()
%19 = call spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupLtMaskKHRv()
%20 = call spir_func i64 @_Z32__spirv_BuiltInLocalInvocationIdi(i32 0)
%21 = call spir_func i64 @_Z28__spirv_BuiltInWorkgroupSizei(i32 0)
%22 = call spir_func i64 @_Z25__spirv_BuiltInGlobalSizei(i32 0)
%23 = call spir_func i64 @_Z26__spirv_BuiltInWorkgroupIdi(i32 0)
%24 = call spir_func i64 @_Z36__spirv_BuiltInEnqueuedWorkgroupSizei(i32 0)
%25 = call spir_func i64 @_Z28__spirv_BuiltInNumWorkgroupsi(i32 0)
%26 = call spir_func i64 @_Z27__spirv_BuiltInGlobalOffseti(i32 0)

ret void
}

declare spir_func i32 @_Z29__spirv_BuiltInGlobalLinearIdv()
declare spir_func i64 @_Z33__spirv_BuiltInGlobalInvocationIdi(i32)
declare spir_func i64 @_Z35__spirv_BuiltInLocalInvocationIndexv()
declare spir_func i32 @_Z22__spirv_BuiltInWorkDimv()
declare spir_func i32 @_Z27__spirv_BuiltInSubgroupSizev()
declare spir_func i32 @_Z30__spirv_BuiltInSubgroupMaxSizev()
declare spir_func i32 @_Z27__spirv_BuiltInNumSubgroupsv()
declare spir_func i32 @_Z35__spirv_BuiltInNumEnqueuedSubgroupsv()
declare spir_func i32 @_Z25__spirv_BuiltInSubgroupIdv()
declare spir_func i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv()
declare spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupEqMaskv()
declare spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupEqMaskKHRv()
declare spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupGeMaskv()
declare spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupGeMaskKHRv()
declare spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupGtMaskv()
declare spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupGtMaskKHRv()
declare spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupLeMaskv()
declare spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupLeMaskKHRv()
declare spir_func <4 x i32> @_Z29__spirv_BuiltInSubgroupLtMaskv()
declare spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupLtMaskKHRv()
declare spir_func i64 @_Z32__spirv_BuiltInLocalInvocationIdi(i32)
declare spir_func i64 @_Z28__spirv_BuiltInWorkgroupSizei(i32)
declare spir_func i64 @_Z25__spirv_BuiltInGlobalSizei(i32)
declare spir_func i64 @_Z26__spirv_BuiltInWorkgroupIdi(i32)
declare spir_func i64 @_Z36__spirv_BuiltInEnqueuedWorkgroupSizei(i32)
declare spir_func i64 @_Z28__spirv_BuiltInNumWorkgroupsi(i32)
declare spir_func i64 @_Z27__spirv_BuiltInGlobalOffseti(i32)