180 lines · plain
1; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK2; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}3 4target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1"5target triple = "spirv32-unknown-unknown"6 7; CHECK: OpDecorate [[NumWorkgroups:%[0-9]*]] BuiltIn NumWorkgroups8; CHECK: OpDecorate [[WorkgroupSize:%[0-9]*]] BuiltIn WorkgroupSize9; CHECK: OpDecorate [[WorkgroupId:%[0-9]*]] BuiltIn WorkgroupId10; CHECK: OpDecorate [[LocalInvocationId:%[0-9]*]] BuiltIn LocalInvocationId11; CHECK: OpDecorate [[GlobalInvocationId:%[0-9]*]] BuiltIn GlobalInvocationId12; CHECK: OpDecorate [[GlobalSize:%[0-9]*]] BuiltIn GlobalSize13; CHECK: OpDecorate [[GlobalOffset:%[0-9]*]] BuiltIn GlobalOffset14; CHECK: OpDecorate [[SubgroupSize:%[0-9]*]] BuiltIn SubgroupSize15; CHECK: OpDecorate [[SubgroupMaxSize:%[0-9]*]] BuiltIn SubgroupMaxSize16; CHECK: OpDecorate [[NumSubgroups:%[0-9]*]] BuiltIn NumSubgroups17; CHECK: OpDecorate [[SubgroupId:%[0-9]*]] BuiltIn SubgroupId18; CHECK: OpDecorate [[SubgroupLocalInvocationId:%[0-9]*]] BuiltIn SubgroupLocalInvocationId19; CHECK: [[I32:%[0-9]*]] = OpTypeInt 32 020; CHECK: [[I32PTR:%[0-9]*]] = OpTypePointer Input [[I32]]21; CHECK: [[I32V3:%[0-9]*]] = OpTypeVector [[I32]] 322; CHECK: [[I32V3PTR:%[0-9]*]] = OpTypePointer Input [[I32V3]]23; CHECK: [[NumWorkgroups]] = OpVariable [[I32V3PTR]] Input24; CHECK: [[WorkgroupSize]] = OpVariable [[I32V3PTR]] Input25; CHECK: [[WorkgroupId]] = OpVariable [[I32V3PTR]] Input26; CHECK: [[LocalInvocationId]] = OpVariable [[I32V3PTR]] Input27; CHECK: [[GlobalInvocationId]] = OpVariable [[I32V3PTR]] Input28; CHECK: [[GlobalSize]] = OpVariable [[I32V3PTR]] Input29; CHECK: [[GlobalOffset]] = OpVariable [[I32V3PTR]] Input30; CHECK: [[SubgroupSize]] = OpVariable [[I32PTR]] Input31; CHECK: [[SubgroupMaxSize]] = OpVariable [[I32PTR]] Input32; CHECK: [[NumSubgroups]] = OpVariable [[I32PTR]] Input33; CHECK: [[SubgroupId]] = OpVariable [[I32PTR]] Input34; CHECK: [[SubgroupLocalInvocationId]] = OpVariable [[I32PTR]] Input35 36@G_spv_num_workgroups_0 = global i32 037@G_spv_num_workgroups_1 = global i32 038@G_spv_num_workgroups_2 = global i32 039@G_spv_workgroup_size_0 = global i32 040@G_spv_workgroup_size_1 = global i32 041@G_spv_workgroup_size_2 = global i32 042@G_spv_group_id_0 = global i32 043@G_spv_group_id_1 = global i32 044@G_spv_group_id_2 = global i32 045@G_spv_thread_id_in_group_0 = global i32 046@G_spv_thread_id_in_group_1 = global i32 047@G_spv_thread_id_in_group_2 = global i32 048@G_spv_thread_id_0 = global i32 049@G_spv_thread_id_1 = global i32 050@G_spv_thread_id_2 = global i32 051@G_spv_global_size_0 = global i32 052@G_spv_global_size_1 = global i32 053@G_spv_global_size_2 = global i32 054@G_spv_global_offset_0 = global i32 055@G_spv_global_offset_1 = global i32 056@G_spv_global_offset_2 = global i32 057 58; Function Attrs: convergent noinline norecurse nounwind optnone59define spir_func void @test_id_and_range() {60entry:61 %ssize = alloca i32, align 462 %smax = alloca i32, align 463 %snum = alloca i32, align 464 %sid = alloca i32, align 465 %sinvocid = alloca i32, align 466; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[NumWorkgroups]]67; CHECK: OpCompositeExtract [[I32]] [[LD]] 068 %spv.num.workgroups = call i32 @llvm.spv.num.workgroups.i32(i32 0)69 store i32 %spv.num.workgroups, i32* @G_spv_num_workgroups_070; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[NumWorkgroups]]71; CHECK: OpCompositeExtract [[I32]] [[LD]] 172 %spv.num.workgroups1 = call i32 @llvm.spv.num.workgroups.i32(i32 1)73 store i32 %spv.num.workgroups1, i32* @G_spv_num_workgroups_174; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[NumWorkgroups]]75; CHECK: OpCompositeExtract [[I32]] [[LD]] 276 %spv.num.workgroups2 = call i32 @llvm.spv.num.workgroups.i32(i32 2)77 store i32 %spv.num.workgroups2, i32* @G_spv_num_workgroups_278; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupSize]]79; CHECK: OpCompositeExtract [[I32]] [[LD]] 080 %spv.workgroup.size = call i32 @llvm.spv.workgroup.size.i32(i32 0)81 store i32 %spv.workgroup.size, i32* @G_spv_workgroup_size_082; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupSize]]83; CHECK: OpCompositeExtract [[I32]] [[LD]] 184 %spv.workgroup.size3 = call i32 @llvm.spv.workgroup.size.i32(i32 1)85 store i32 %spv.workgroup.size3, i32* @G_spv_workgroup_size_186; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupSize]]87; CHECK: OpCompositeExtract [[I32]] [[LD]] 288 %spv.workgroup.size4 = call i32 @llvm.spv.workgroup.size.i32(i32 2)89 store i32 %spv.workgroup.size4, i32* @G_spv_workgroup_size_290; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupId]]91; CHECK: OpCompositeExtract [[I32]] [[LD]] 092 %spv.group.id = call i32 @llvm.spv.group.id.i32(i32 0)93 store i32 %spv.group.id, i32* @G_spv_group_id_094; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupId]]95; CHECK: OpCompositeExtract [[I32]] [[LD]] 196 %spv.group.id5 = call i32 @llvm.spv.group.id.i32(i32 1)97 store i32 %spv.group.id5, i32* @G_spv_group_id_198; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupId]]99; CHECK: OpCompositeExtract [[I32]] [[LD]] 2100 %spv.group.id6 = call i32 @llvm.spv.group.id.i32(i32 2)101 store i32 %spv.group.id6, i32* @G_spv_group_id_2102; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[LocalInvocationId]]103; CHECK: OpCompositeExtract [[I32]] [[LD]] 0104 %spv.thread.id.in.group = call i32 @llvm.spv.thread.id.in.group.i32(i32 0)105 store i32 %spv.thread.id.in.group, i32* @G_spv_thread_id_in_group_0106; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[LocalInvocationId]]107; CHECK: OpCompositeExtract [[I32]] [[LD]] 1108 %spv.thread.id.in.group7 = call i32 @llvm.spv.thread.id.in.group.i32(i32 1)109 store i32 %spv.thread.id.in.group7, i32* @G_spv_thread_id_in_group_1110; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[LocalInvocationId]]111; CHECK: OpCompositeExtract [[I32]] [[LD]] 2112 %spv.thread.id.in.group8 = call i32 @llvm.spv.thread.id.in.group.i32(i32 2)113 store i32 %spv.thread.id.in.group8, i32* @G_spv_thread_id_in_group_2114; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalInvocationId]]115; CHECK: OpCompositeExtract [[I32]] [[LD]] 0116 %spv.thread.id = call i32 @llvm.spv.thread.id.i32(i32 0)117 store i32 %spv.thread.id, i32* @G_spv_thread_id_0118; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalInvocationId]]119; CHECK: OpCompositeExtract [[I32]] [[LD]] 1120 %spv.thread.id9 = call i32 @llvm.spv.thread.id.i32(i32 1)121 store i32 %spv.thread.id9, i32* @G_spv_thread_id_1122; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalInvocationId]]123; CHECK: OpCompositeExtract [[I32]] [[LD]] 2124 %spv.thread.id10 = call i32 @llvm.spv.thread.id.i32(i32 2)125 store i32 %spv.thread.id10, i32* @G_spv_thread_id_2126; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalSize]]127; CHECK: OpCompositeExtract [[I32]] [[LD]] 0128 %spv.num.workgroups11 = call i32 @llvm.spv.global.size.i32(i32 0)129 store i32 %spv.num.workgroups11, i32* @G_spv_global_size_0130; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalSize]]131; CHECK: OpCompositeExtract [[I32]] [[LD]] 1132 %spv.num.workgroups12 = call i32 @llvm.spv.global.size.i32(i32 1)133 store i32 %spv.num.workgroups12, i32* @G_spv_global_size_1134; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalSize]]135; CHECK: OpCompositeExtract [[I32]] [[LD]] 2136 %spv.num.workgroups13 = call i32 @llvm.spv.global.size.i32(i32 2)137 store i32 %spv.num.workgroups13, i32* @G_spv_global_size_2138; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalOffset]]139; CHECK: OpCompositeExtract [[I32]] [[LD]] 0140 %spv.global.offset = call i32 @llvm.spv.global.offset.i32(i32 0)141 store i32 %spv.global.offset, i32* @G_spv_global_offset_0142; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalOffset]]143; CHECK: OpCompositeExtract [[I32]] [[LD]] 1144 %spv.global.offset14 = call i32 @llvm.spv.global.offset.i32(i32 1)145 store i32 %spv.global.offset14, i32* @G_spv_global_offset_1146; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalOffset]]147; CHECK: OpCompositeExtract [[I32]] [[LD]] 2148 %spv.global.offset15 = call i32 @llvm.spv.global.offset.i32(i32 2)149 store i32 %spv.global.offset15, i32* @G_spv_global_offset_2150; CHECK: OpLoad %5 [[SubgroupSize]]151 %0 = call i32 @llvm.spv.subgroup.size()152 store i32 %0, ptr %ssize, align 4153; CHECK: OpLoad %5 [[SubgroupMaxSize]]154 %1 = call i32 @llvm.spv.subgroup.max.size()155 store i32 %1, ptr %smax, align 4156; CHECK: OpLoad %5 [[NumSubgroups]]157 %2 = call i32 @llvm.spv.num.subgroups()158 store i32 %2, ptr %snum, align 4159; CHECK: OpLoad %5 [[SubgroupId]]160 %3 = call i32 @llvm.spv.subgroup.id()161 store i32 %3, ptr %sid, align 4162; CHECK: OpLoad %5 [[SubgroupLocalInvocationId]]163 %4 = call i32 @llvm.spv.subgroup.local.invocation.id()164 store i32 %4, ptr %sinvocid, align 4165 ret void166}167 168declare i32 @llvm.spv.num.workgroups.i32(i32)169declare i32 @llvm.spv.workgroup.size.i32(i32)170declare i32 @llvm.spv.group.id.i32(i32)171declare i32 @llvm.spv.thread.id.in.group.i32(i32)172declare i32 @llvm.spv.thread.id.i32(i32)173declare i32 @llvm.spv.global.size.i32(i32)174declare i32 @llvm.spv.global.offset.i32(i32)175declare noundef i32 @llvm.spv.subgroup.size()176declare noundef i32 @llvm.spv.subgroup.max.size()177declare noundef i32 @llvm.spv.num.subgroups()178declare noundef i32 @llvm.spv.subgroup.id()179declare noundef i32 @llvm.spv.subgroup.local.invocation.id()180