diff --git a/clang/lib/Headers/__clang_spirv_builtins.h b/clang/lib/Headers/__clang_spirv_builtins.h index 9915cdfcae7cd..9c7215f506508 100644 --- a/clang/lib/Headers/__clang_spirv_builtins.h +++ b/clang/lib/Headers/__clang_spirv_builtins.h @@ -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 diff --git a/clang/test/Headers/spirv_ids.cpp b/clang/test/Headers/spirv_ids.cpp index 0cd74dbca53aa..466be5deee87a 100644 --- a/clang/test/Headers/spirv_ids.cpp +++ b/clang/test/Headers/spirv_ids.cpp @@ -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(); } diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.td b/llvm/lib/Target/SPIRV/SPIRVBuiltins.td index 6897865eb4e15..ea78dcd135267 100644 --- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.td +++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.td @@ -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>; @@ -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>; //===----------------------------------------------------------------------===// diff --git a/llvm/test/CodeGen/SPIRV/transcoding/builtin_calls.ll b/llvm/test/CodeGen/SPIRV/transcoding/builtin_calls.ll index 0a02a8bf56ace..b179732371d97 100644 --- a/llvm/test/CodeGen/SPIRV/transcoding/builtin_calls.ll +++ b/llvm/test/CodeGen/SPIRV/transcoding/builtin_calls.ll @@ -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)