181 lines · plain
1; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK2; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}3 4target datalayout = "e-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 = "spirv64-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: [[I64:%[0-9]*]] = OpTypeInt 64 021; CHECK: [[I32PTR:%[0-9]*]] = OpTypePointer Input [[I32]]22; CHECK: [[I64V3:%[0-9]*]] = OpTypeVector [[I64]] 323; CHECK: [[I64V3PTR:%[0-9]*]] = OpTypePointer Input [[I64V3]]24; CHECK: [[NumWorkgroups]] = OpVariable [[I64V3PTR]] Input25; CHECK: [[WorkgroupSize]] = OpVariable [[I64V3PTR]] Input26; CHECK: [[WorkgroupId]] = OpVariable [[I64V3PTR]] Input27; CHECK: [[LocalInvocationId]] = OpVariable [[I64V3PTR]] Input28; CHECK: [[GlobalInvocationId]] = OpVariable [[I64V3PTR]] Input29; CHECK: [[GlobalSize]] = OpVariable [[I64V3PTR]] Input30; CHECK: [[GlobalOffset]] = OpVariable [[I64V3PTR]] Input31; CHECK: [[SubgroupSize]] = OpVariable [[I32PTR]] Input32; CHECK: [[SubgroupMaxSize]] = OpVariable [[I32PTR]] Input33; CHECK: [[NumSubgroups]] = OpVariable [[I32PTR]] Input34; CHECK: [[SubgroupId]] = OpVariable [[I32PTR]] Input35; CHECK: [[SubgroupLocalInvocationId]] = OpVariable [[I32PTR]] Input36 37@G_spv_num_workgroups_0 = global i64 038@G_spv_num_workgroups_1 = global i64 039@G_spv_num_workgroups_2 = global i64 040@G_spv_workgroup_size_0 = global i64 041@G_spv_workgroup_size_1 = global i64 042@G_spv_workgroup_size_2 = global i64 043@G_spv_group_id_0 = global i64 044@G_spv_group_id_1 = global i64 045@G_spv_group_id_2 = global i64 046@G_spv_thread_id_in_group_0 = global i64 047@G_spv_thread_id_in_group_1 = global i64 048@G_spv_thread_id_in_group_2 = global i64 049@G_spv_thread_id_0 = global i64 050@G_spv_thread_id_1 = global i64 051@G_spv_thread_id_2 = global i64 052@G_spv_global_size_0 = global i64 053@G_spv_global_size_1 = global i64 054@G_spv_global_size_2 = global i64 055@G_spv_global_offset_0 = global i64 056@G_spv_global_offset_1 = global i64 057@G_spv_global_offset_2 = global i64 058 59; Function Attrs: convergent noinline norecurse nounwind optnone60define spir_func void @test_id_and_range() {61entry:62 %ssize = alloca i32, align 463 %smax = alloca i32, align 464 %snum = alloca i32, align 465 %sid = alloca i32, align 466 %sinvocid = alloca i32, align 467; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[NumWorkgroups]]68; CHECK: OpCompositeExtract [[I64]] [[LD]] 069 %spv.num.workgroups = call i64 @llvm.spv.num.workgroups.i64(i32 0)70 store i64 %spv.num.workgroups, i64* @G_spv_num_workgroups_071; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[NumWorkgroups]]72; CHECK: OpCompositeExtract [[I64]] [[LD]] 173 %spv.num.workgroups1 = call i64 @llvm.spv.num.workgroups.i64(i32 1)74 store i64 %spv.num.workgroups1, i64* @G_spv_num_workgroups_175; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[NumWorkgroups]]76; CHECK: OpCompositeExtract [[I64]] [[LD]] 277 %spv.num.workgroups2 = call i64 @llvm.spv.num.workgroups.i64(i32 2)78 store i64 %spv.num.workgroups2, i64* @G_spv_num_workgroups_279; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupSize]]80; CHECK: OpCompositeExtract [[I64]] [[LD]] 081 %spv.workgroup.size = call i64 @llvm.spv.workgroup.size.i64(i32 0)82 store i64 %spv.workgroup.size, i64* @G_spv_workgroup_size_083; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupSize]]84; CHECK: OpCompositeExtract [[I64]] [[LD]] 185 %spv.workgroup.size3 = call i64 @llvm.spv.workgroup.size.i64(i32 1)86 store i64 %spv.workgroup.size3, i64* @G_spv_workgroup_size_187; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupSize]]88; CHECK: OpCompositeExtract [[I64]] [[LD]] 289 %spv.workgroup.size4 = call i64 @llvm.spv.workgroup.size.i64(i32 2)90 store i64 %spv.workgroup.size4, i64* @G_spv_workgroup_size_291; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupId]]92; CHECK: OpCompositeExtract [[I64]] [[LD]] 093 %spv.group.id = call i64 @llvm.spv.group.id.i64(i32 0)94 store i64 %spv.group.id, i64* @G_spv_group_id_095; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupId]]96; CHECK: OpCompositeExtract [[I64]] [[LD]] 197 %spv.group.id5 = call i64 @llvm.spv.group.id.i64(i32 1)98 store i64 %spv.group.id5, i64* @G_spv_group_id_199; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupId]]100; CHECK: OpCompositeExtract [[I64]] [[LD]] 2101 %spv.group.id6 = call i64 @llvm.spv.group.id.i64(i32 2)102 store i64 %spv.group.id6, i64* @G_spv_group_id_2103; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[LocalInvocationId]]104; CHECK: OpCompositeExtract [[I64]] [[LD]] 0105 %spv.thread.id.in.group = call i64 @llvm.spv.thread.id.in.group.i64(i32 0)106 store i64 %spv.thread.id.in.group, i64* @G_spv_thread_id_in_group_0107; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[LocalInvocationId]]108; CHECK: OpCompositeExtract [[I64]] [[LD]] 1109 %spv.thread.id.in.group7 = call i64 @llvm.spv.thread.id.in.group.i64(i32 1)110 store i64 %spv.thread.id.in.group7, i64* @G_spv_thread_id_in_group_1111; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[LocalInvocationId]]112; CHECK: OpCompositeExtract [[I64]] [[LD]] 2113 %spv.thread.id.in.group8 = call i64 @llvm.spv.thread.id.in.group.i64(i32 2)114 store i64 %spv.thread.id.in.group8, i64* @G_spv_thread_id_in_group_2115; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalInvocationId]]116; CHECK: OpCompositeExtract [[I64]] [[LD]] 0117 %spv.thread.id = call i64 @llvm.spv.thread.id.i64(i32 0)118 store i64 %spv.thread.id, i64* @G_spv_thread_id_0119; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalInvocationId]]120; CHECK: OpCompositeExtract [[I64]] [[LD]] 1121 %spv.thread.id9 = call i64 @llvm.spv.thread.id.i64(i32 1)122 store i64 %spv.thread.id9, i64* @G_spv_thread_id_1123; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalInvocationId]]124; CHECK: OpCompositeExtract [[I64]] [[LD]] 2125 %spv.thread.id10 = call i64 @llvm.spv.thread.id.i64(i32 2)126 store i64 %spv.thread.id10, i64* @G_spv_thread_id_2127; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalSize]]128; CHECK: OpCompositeExtract [[I64]] [[LD]] 0129 %spv.num.workgroups11 = call i64 @llvm.spv.global.size.i64(i32 0)130 store i64 %spv.num.workgroups11, i64* @G_spv_global_size_0131; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalSize]]132; CHECK: OpCompositeExtract [[I64]] [[LD]] 1133 %spv.num.workgroups12 = call i64 @llvm.spv.global.size.i64(i32 1)134 store i64 %spv.num.workgroups12, i64* @G_spv_global_size_1135; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalSize]]136; CHECK: OpCompositeExtract [[I64]] [[LD]] 2137 %spv.num.workgroups13 = call i64 @llvm.spv.global.size.i64(i32 2)138 store i64 %spv.num.workgroups13, i64* @G_spv_global_size_2139; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalOffset]]140; CHECK: OpCompositeExtract [[I64]] [[LD]] 0141 %spv.global.offset = call i64 @llvm.spv.global.offset.i64(i32 0)142 store i64 %spv.global.offset, i64* @G_spv_global_offset_0143; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalOffset]]144; CHECK: OpCompositeExtract [[I64]] [[LD]] 1145 %spv.global.offset14 = call i64 @llvm.spv.global.offset.i64(i32 1)146 store i64 %spv.global.offset14, i64* @G_spv_global_offset_1147; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalOffset]]148; CHECK: OpCompositeExtract [[I64]] [[LD]] 2149 %spv.global.offset15 = call i64 @llvm.spv.global.offset.i64(i32 2)150 store i64 %spv.global.offset15, i64* @G_spv_global_offset_2151; CHECK: OpLoad %5 [[SubgroupSize]]152 %0 = call i32 @llvm.spv.subgroup.size()153 store i32 %0, ptr %ssize, align 4154; CHECK: OpLoad %5 [[SubgroupMaxSize]]155 %1 = call i32 @llvm.spv.subgroup.max.size()156 store i32 %1, ptr %smax, align 4157; CHECK: OpLoad %5 [[NumSubgroups]]158 %2 = call i32 @llvm.spv.num.subgroups()159 store i32 %2, ptr %snum, align 4160; CHECK: OpLoad %5 [[SubgroupId]]161 %3 = call i32 @llvm.spv.subgroup.id()162 store i32 %3, ptr %sid, align 4163; CHECK: OpLoad %5 [[SubgroupLocalInvocationId]]164 %4 = call i32 @llvm.spv.subgroup.local.invocation.id()165 store i32 %4, ptr %sinvocid, align 4166 ret void167}168 169declare i64 @llvm.spv.num.workgroups.i64(i32)170declare i64 @llvm.spv.workgroup.size.i64(i32)171declare i64 @llvm.spv.group.id.i64(i32)172declare i64 @llvm.spv.thread.id.in.group.i64(i32)173declare i64 @llvm.spv.thread.id.i64(i32)174declare i64 @llvm.spv.global.size.i64(i32)175declare i64 @llvm.spv.global.offset.i64(i32)176declare noundef i32 @llvm.spv.subgroup.size()177declare noundef i32 @llvm.spv.subgroup.max.size()178declare noundef i32 @llvm.spv.num.subgroups()179declare noundef i32 @llvm.spv.subgroup.id()180declare noundef i32 @llvm.spv.subgroup.local.invocation.id()181