140 lines · plain
1; RUN: opt -mtriple=amdgcn-- -passes=amdgpu-attributor -o %t.bc %s2; RUN: llc -mtriple=amdgcn -mcpu=hawaii < %t.bc | FileCheck --check-prefixes=ALL,MESA,UNPACKED %s3; RUN: llc -mtriple=amdgcn -mcpu=tonga -mattr=-flat-for-global < %t.bc | FileCheck --check-prefixes=ALL,MESA,UNPACKED %s4; RUN: llc -mtriple=amdgcn-unknown-mesa3d -mcpu=hawaii < %t.bc | FileCheck -check-prefixes=ALL,MESA3D,UNPACKED %s5; RUN: llc -mtriple=amdgcn-unknown-mesa3d -mcpu=tonga -mattr=-flat-for-global < %t.bc | FileCheck -check-prefixes=ALL,MESA3D,UNPACKED %s6; RUN: llc -mtriple=amdgcn-unknown-amdhsa -mcpu=gfx90a < %t.bc | FileCheck -check-prefixes=ALL,PACKED-TID %s7; RUN: llc -mtriple=amdgcn-unknown-amdhsa -mcpu=gfx1100 -amdgpu-enable-vopd=0 < %t.bc | FileCheck -check-prefixes=ALL,PACKED-TID %s8 9declare i32 @llvm.amdgcn.workitem.id.x() #010declare i32 @llvm.amdgcn.workitem.id.y() #011declare i32 @llvm.amdgcn.workitem.id.z() #012 13; MESA: .section .AMDGPU.config14; MESA: .long 4718015; MESA-NEXT: .long 132{{$}}16 17; ALL-LABEL: {{^}}test_workitem_id_x:18; MESA3D: enable_vgpr_workitem_id = 019 20; ALL-NOT: v021; ALL: {{buffer|flat|global}}_store_{{dword|b32}} {{.*}}v022 23; PACKED-TID: .amdhsa_system_vgpr_workitem_id 024define amdgpu_kernel void @test_workitem_id_x(ptr addrspace(1) %out) #1 {25 %id = call i32 @llvm.amdgcn.workitem.id.x()26 store i32 %id, ptr addrspace(1) %out27 ret void28}29 30; MESA: .section .AMDGPU.config31; MESA: .long 4718032; MESA-NEXT: .long 2180{{$}}33 34; ALL-LABEL: {{^}}test_workitem_id_y:35; MESA3D: enable_vgpr_workitem_id = 136; MESA3D-NOT: v137; MESA3D: {{buffer|flat}}_store_dword {{.*}}v138 39; PACKED-TID: v_bfe_u32 [[ID:v[0-9]+]], v0, 10, 1040; PACKED-TID: {{buffer|flat|global}}_store_{{dword|b32}} {{.*}}[[ID]]41; PACKED-TID: .amdhsa_system_vgpr_workitem_id 142define amdgpu_kernel void @test_workitem_id_y(ptr addrspace(1) %out) #1 {43 %id = call i32 @llvm.amdgcn.workitem.id.y()44 store i32 %id, ptr addrspace(1) %out45 ret void46}47 48; MESA: .section .AMDGPU.config49; MESA: .long 4718050; MESA-NEXT: .long 4228{{$}}51 52; ALL-LABEL: {{^}}test_workitem_id_z:53; MESA3D: enable_vgpr_workitem_id = 254; MESA3D-NOT: v255; MESA3D: {{buffer|flat}}_store_dword {{.*}}v256 57; PACKED-TID: v_bfe_u32 [[ID:v[0-9]+]], v0, 20, 1058; PACKED-TID: {{buffer|flat|global}}_store_{{dword|b32}} {{.*}}[[ID]]59; PACKED-TID: .amdhsa_system_vgpr_workitem_id 260define amdgpu_kernel void @test_workitem_id_z(ptr addrspace(1) %out) #1 {61 %id = call i32 @llvm.amdgcn.workitem.id.z()62 store i32 %id, ptr addrspace(1) %out63 ret void64}65 66; FIXME: Packed tid should avoid the and67; ALL-LABEL: {{^}}test_reqd_workgroup_size_x_only:68; MESA3D: enable_vgpr_workitem_id = 069 70; ALL-DAG: v_mov_b32_e32 [[ZERO:v[0-9]+]], 0{{$}}71; UNPACKED-DAG: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, v072 73; PACKED: v_and_b32_e32 [[MASKED:v[0-9]+]], 0x3ff, v074; PACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[MASKED]]75 76; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]77; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]78define amdgpu_kernel void @test_reqd_workgroup_size_x_only(ptr %out) !reqd_work_group_size !0 {79 %id.x = call i32 @llvm.amdgcn.workitem.id.x()80 %id.y = call i32 @llvm.amdgcn.workitem.id.y()81 %id.z = call i32 @llvm.amdgcn.workitem.id.z()82 store volatile i32 %id.x, ptr %out83 store volatile i32 %id.y, ptr %out84 store volatile i32 %id.z, ptr %out85 ret void86}87 88; ALL-LABEL: {{^}}test_reqd_workgroup_size_y_only:89; MESA3D: enable_vgpr_workitem_id = 190 91; ALL: v_mov_b32_e32 [[ZERO:v[0-9]+]], 0{{$}}92; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]93 94; UNPACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, v195 96; PACKED: v_bfe_u32 [[MASKED:v[0-9]+]], v0, 10, 1097; PACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[MASKED]]98 99; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]100define amdgpu_kernel void @test_reqd_workgroup_size_y_only(ptr %out) !reqd_work_group_size !1 {101 %id.x = call i32 @llvm.amdgcn.workitem.id.x()102 %id.y = call i32 @llvm.amdgcn.workitem.id.y()103 %id.z = call i32 @llvm.amdgcn.workitem.id.z()104 store volatile i32 %id.x, ptr %out105 store volatile i32 %id.y, ptr %out106 store volatile i32 %id.z, ptr %out107 ret void108}109 110; ALL-LABEL: {{^}}test_reqd_workgroup_size_z_only:111; MESA3D: enable_vgpr_workitem_id = 2112 113; ALL: v_mov_b32_e32 [[ZERO:v[0-9]+]], 0{{$}}114; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]115; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]116 117; UNPACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, v2118 119; PACKED: v_bfe_u32 [[MASKED:v[0-9]+]], v0, 10, 20120; PACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[MASKED]]121define amdgpu_kernel void @test_reqd_workgroup_size_z_only(ptr %out) !reqd_work_group_size !2 {122 %id.x = call i32 @llvm.amdgcn.workitem.id.x()123 %id.y = call i32 @llvm.amdgcn.workitem.id.y()124 %id.z = call i32 @llvm.amdgcn.workitem.id.z()125 store volatile i32 %id.x, ptr %out126 store volatile i32 %id.y, ptr %out127 store volatile i32 %id.z, ptr %out128 ret void129}130 131attributes #0 = { nounwind readnone }132attributes #1 = { nounwind }133 134!llvm.module.flags = !{!3}135 136!0 = !{i32 64, i32 1, i32 1}137!1 = !{i32 1, i32 64, i32 1}138!2 = !{i32 1, i32 1, i32 64}139!3 = !{i32 1, !"amdhsa_code_object_version", i32 400}140