207 lines · plain
1; RUN: sed 's/CODE_OBJECT_VERSION/400/g' %s | opt -S -mtriple=amdgcn-amd-amdhsa -passes=amdgpu-attributor -o %t.v4.ll2; RUN: sed 's/CODE_OBJECT_VERSION/600/g' %s | opt -S -mtriple=amdgcn-amd-amdhsa -passes=amdgpu-attributor -o %t.v6.ll3; RUN: llc -global-isel -new-reg-bank-select -mtriple=amdgcn-unknown-amdhsa < %t.v4.ll | FileCheck --check-prefixes=ALL,HSA,UNPACKED %s4; RUN: llc -global-isel -new-reg-bank-select -mtriple=amdgcn-unknown-amdhsa < %t.v4.ll | FileCheck --check-prefixes=ALL,HSA,UNPACKED %s5; RUN: llc -global-isel -new-reg-bank-select -mtriple=amdgcn-- -mcpu=hawaii -mattr=+flat-for-global < %t.v4.ll | FileCheck --check-prefixes=ALL,MESA,UNPACKED %s6; RUN: llc -global-isel -new-reg-bank-select -mtriple=amdgcn-- -mcpu=tonga -mattr=+flat-for-global < %t.v4.ll | FileCheck --check-prefixes=ALL,MESA,UNPACKED %s7; RUN: llc -global-isel -new-reg-bank-select -mtriple=amdgcn-unknown-mesa3d -mattr=+flat-for-global -mcpu=hawaii < %t.v4.ll | FileCheck -check-prefixes=ALL,MESA3D,UNPACKED %s8; RUN: llc -global-isel -new-reg-bank-select -mtriple=amdgcn-unknown-mesa3d -mcpu=tonga < %t.v4.ll | FileCheck -check-prefixes=ALL,MESA3D,UNPACKED %s9; RUN: llc -global-isel -new-reg-bank-select -mtriple=amdgcn-unknown-amdhsa -mcpu=gfx90a < %t.v4.ll | FileCheck -check-prefixes=ALL,PACKED-TID %s10; RUN: llc -global-isel -new-reg-bank-select -mtriple=amdgcn-unknown-amdhsa -mcpu=gfx1100 -amdgpu-enable-vopd=0 < %t.v4.ll | FileCheck -check-prefixes=ALL,PACKED-TID %s11; RUN: llc -global-isel -new-reg-bank-select -mtriple=amdgcn-unknown-amdhsa --amdhsa-code-object-version=6 -mcpu=gfx11-generic -amdgpu-enable-vopd=0 < %t.v6.ll | FileCheck -check-prefixes=ALL,PACKED-TID %s12 13declare i32 @llvm.amdgcn.workitem.id.x() #014declare i32 @llvm.amdgcn.workitem.id.y() #015declare i32 @llvm.amdgcn.workitem.id.z() #016 17; MESA: .section .AMDGPU.config18; MESA: .long 4718019; MESA-NEXT: .long 132{{$}}20 21; ALL-LABEL: {{^}}test_workitem_id_x:22; MESA3D: enable_vgpr_workitem_id = 023 24; ALL-NOT: v025; ALL: {{buffer|flat|global}}_store_{{dword|b32}} {{.*}}v026 27; PACKED-TID: .amdhsa_system_vgpr_workitem_id 028define amdgpu_kernel void @test_workitem_id_x(ptr addrspace(1) %out) #1 {29 %id = call i32 @llvm.amdgcn.workitem.id.x()30 store i32 %id, ptr addrspace(1) %out31 ret void32}33 34; MESA: .section .AMDGPU.config35; MESA: .long 4718036; MESA-NEXT: .long 2180{{$}}37 38; ALL-LABEL: {{^}}test_workitem_id_y:39; MESA3D: enable_vgpr_workitem_id = 140; MESA3D-NOT: v141; MESA3D: {{buffer|flat}}_store_dword {{.*}}v142 43; PACKED-TID: v_bfe_u32 [[ID:v[0-9]+]], v0, 10, 1044; PACKED-TID: {{buffer|flat|global}}_store_{{dword|b32}} {{.*}}[[ID]]45; PACKED-TID: .amdhsa_system_vgpr_workitem_id 146define amdgpu_kernel void @test_workitem_id_y(ptr addrspace(1) %out) #1 {47 %id = call i32 @llvm.amdgcn.workitem.id.y()48 store i32 %id, ptr addrspace(1) %out49 ret void50}51 52; MESA: .section .AMDGPU.config53; MESA: .long 4718054; MESA-NEXT: .long 4228{{$}}55 56; ALL-LABEL: {{^}}test_workitem_id_z:57; MESA3D: enable_vgpr_workitem_id = 258; MESA3D-NOT: v259; MESA3D: {{buffer|flat}}_store_dword {{.*}}v260 61; PACKED-TID: v_bfe_u32 [[ID:v[0-9]+]], v0, 20, 1062; PACKED-TID: {{buffer|flat|global}}_store_{{dword|b32}} {{.*}}[[ID]]63; PACKED-TID: .amdhsa_system_vgpr_workitem_id 264define amdgpu_kernel void @test_workitem_id_z(ptr addrspace(1) %out) #1 {65 %id = call i32 @llvm.amdgcn.workitem.id.z()66 store i32 %id, ptr addrspace(1) %out67 ret void68}69 70; ALL-LABEL: {{^}}test_workitem_id_x_usex2:71; ALL-NOT: v072; ALL: {{flat|global}}_store_{{dword|b32}} v{{.*}}, v073; ALL-NOT: v074; ALL: {{flat|global}}_store_{{dword|b32}} v{{.*}}, v075define amdgpu_kernel void @test_workitem_id_x_usex2(ptr addrspace(1) %out) #1 {76 %id0 = call i32 @llvm.amdgcn.workitem.id.x()77 store volatile i32 %id0, ptr addrspace(1) %out78 79 %id1 = call i32 @llvm.amdgcn.workitem.id.x()80 store volatile i32 %id1, ptr addrspace(1) %out81 ret void82}83 84; ALL-LABEL: {{^}}test_workitem_id_x_use_outside_entry:85; ALL-NOT: v086; ALL: {{flat|global}}_store_{{dword|b32}}87; ALL-NOT: v088; ALL: {{flat|global}}_store_{{dword|b32}} v{{.*}}, v089define amdgpu_kernel void @test_workitem_id_x_use_outside_entry(ptr addrspace(1) %out, i32 %arg) #1 {90bb0:91 store volatile i32 0, ptr addrspace(1) %out92 %cond = icmp eq i32 %arg, 093 br i1 %cond, label %bb1, label %bb294 95bb1:96 %id = call i32 @llvm.amdgcn.workitem.id.x()97 store volatile i32 %id, ptr addrspace(1) %out98 br label %bb299 100bb2:101 ret void102}103 104; ALL-LABEL: {{^}}test_workitem_id_x_func:105; ALL: s_waitcnt106; HSA-NEXT: v_and_b32_e32 v2, 0x3ff, v31107; MESA-NEXT: v_and_b32_e32 v2, 0x3ff, v31108define void @test_workitem_id_x_func(ptr addrspace(1) %out) #1 {109 %id = call i32 @llvm.amdgcn.workitem.id.x()110 store i32 %id, ptr addrspace(1) %out111 ret void112}113 114; ALL-LABEL: {{^}}test_workitem_id_y_func:115; HSA: v_bfe_u32 v2, v31, 10, 10116; MESA: v_bfe_u32 v2, v31, 10, 10117define void @test_workitem_id_y_func(ptr addrspace(1) %out) #1 {118 %id = call i32 @llvm.amdgcn.workitem.id.y()119 store i32 %id, ptr addrspace(1) %out120 ret void121}122 123; ALL-LABEL: {{^}}test_workitem_id_z_func:124; HSA: v_bfe_u32 v2, v31, 20, 10125; MESA: v_bfe_u32 v2, v31, 20, 10126define void @test_workitem_id_z_func(ptr addrspace(1) %out) #1 {127 %id = call i32 @llvm.amdgcn.workitem.id.z()128 store i32 %id, ptr addrspace(1) %out129 ret void130}131 132; FIXME: Should be able to avoid enabling in kernel inputs133; FIXME: Packed tid should avoid the and134; ALL-LABEL: {{^}}test_reqd_workgroup_size_x_only:135; MESA3D: enable_vgpr_workitem_id = 0136 137; ALL-DAG: v_mov_b32_e32 [[ZERO:v[0-9]+]], 0{{$}}138; UNPACKED-DAG: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, v0139 140; PACKED: v_and_b32_e32 [[MASKED:v[0-9]+]], 0x3ff, v0141; PACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[MASKED]]142 143; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]144; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]145define amdgpu_kernel void @test_reqd_workgroup_size_x_only(ptr %out) !reqd_work_group_size !0 {146 %id.x = call i32 @llvm.amdgcn.workitem.id.x()147 %id.y = call i32 @llvm.amdgcn.workitem.id.y()148 %id.z = call i32 @llvm.amdgcn.workitem.id.z()149 store volatile i32 %id.x, ptr %out150 store volatile i32 %id.y, ptr %out151 store volatile i32 %id.z, ptr %out152 ret void153}154 155; ALL-LABEL: {{^}}test_reqd_workgroup_size_y_only:156; MESA3D: enable_vgpr_workitem_id = 1157 158; ALL: v_mov_b32_e32 [[ZERO:v[0-9]+]], 0{{$}}159; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]160 161; UNPACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, v1162 163; PACKED: v_bfe_u32 [[MASKED:v[0-9]+]], v0, 10, 10164; PACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[MASKED]]165 166; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]167define amdgpu_kernel void @test_reqd_workgroup_size_y_only(ptr %out) !reqd_work_group_size !1 {168 %id.x = call i32 @llvm.amdgcn.workitem.id.x()169 %id.y = call i32 @llvm.amdgcn.workitem.id.y()170 %id.z = call i32 @llvm.amdgcn.workitem.id.z()171 store volatile i32 %id.x, ptr %out172 store volatile i32 %id.y, ptr %out173 store volatile i32 %id.z, ptr %out174 ret void175}176 177; ALL-LABEL: {{^}}test_reqd_workgroup_size_z_only:178; MESA3D: enable_vgpr_workitem_id = 2179 180; ALL: v_mov_b32_e32 [[ZERO:v[0-9]+]], 0{{$}}181; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]182; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]183 184; UNPACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, v2185 186; PACKED: v_bfe_u32 [[MASKED:v[0-9]+]], v0, 10, 20187; PACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[MASKED]]188define amdgpu_kernel void @test_reqd_workgroup_size_z_only(ptr %out) !reqd_work_group_size !2 {189 %id.x = call i32 @llvm.amdgcn.workitem.id.x()190 %id.y = call i32 @llvm.amdgcn.workitem.id.y()191 %id.z = call i32 @llvm.amdgcn.workitem.id.z()192 store volatile i32 %id.x, ptr %out193 store volatile i32 %id.y, ptr %out194 store volatile i32 %id.z, ptr %out195 ret void196}197 198attributes #0 = { nounwind readnone }199attributes #1 = { nounwind }200 201!0 = !{i32 64, i32 1, i32 1}202!1 = !{i32 1, i32 64, i32 1}203!2 = !{i32 1, i32 1, i32 64}204 205!llvm.module.flags = !{!99}206!99 = !{i32 1, !"amdhsa_code_object_version", i32 CODE_OBJECT_VERSION}207