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.

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
Show file tree
Hide file tree
Changes from all commits
Commits
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
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
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
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
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 %[[#]] Input
; CHECK-SPIRV: %[[#Id1:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id2:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id3:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id4:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id5:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id6:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id7:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id8:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id9:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id10:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id11:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id12:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id13:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id14:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id15:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id16:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id17:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id18:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id19:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id20:]] = OpVariable %[[#]] Input
; CHECK-SPIRV: %[[#Id21:]] = OpVariable %[[#]] Input

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)