1; RUN: llc -march=amdgcn -mtriple=amdgcn-unknown-amdhsa --amdhsa-code-object-version=2 -mcpu=kaveri -verify-machineinstrs < %s | FileCheck --check-prefixes=ALL,CO-V2,UNPACKED %s 2; RUN: llc -march=amdgcn -mtriple=amdgcn-unknown-amdhsa --amdhsa-code-object-version=2 -mcpu=carrizo -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck --check-prefixes=ALL,CO-V2,UNPACKED %s 3; RUN: llc -march=amdgcn -mcpu=hawaii -verify-machineinstrs < %s | FileCheck --check-prefixes=ALL,MESA,UNPACKED %s 4; RUN: llc -march=amdgcn -mcpu=tonga -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck --check-prefixes=ALL,MESA,UNPACKED %s 5; RUN: llc -mtriple=amdgcn-unknown-mesa3d -mcpu=hawaii -verify-machineinstrs < %s | FileCheck -check-prefixes=ALL,CO-V2,UNPACKED %s 6; RUN: llc -mtriple=amdgcn-unknown-mesa3d -mcpu=tonga -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck -check-prefixes=ALL,CO-V2,UNPACKED %s 7; RUN: llc -march=amdgcn -mtriple=amdgcn-unknown-amdhsa -mcpu=gfx90a -verify-machineinstrs < %s | FileCheck -check-prefixes=ALL,PACKED-TID %s 8; RUN: llc -march=amdgcn -mtriple=amdgcn-unknown-amdhsa -mcpu=gfx1100 -verify-machineinstrs -amdgpu-enable-vopd=0 < %s | FileCheck -check-prefixes=ALL,PACKED-TID %s 9 10declare i32 @llvm.amdgcn.workitem.id.x() #0 11declare i32 @llvm.amdgcn.workitem.id.y() #0 12declare i32 @llvm.amdgcn.workitem.id.z() #0 13 14; MESA: .section .AMDGPU.config 15; MESA: .long 47180 16; MESA-NEXT: .long 132{{$}} 17 18; ALL-LABEL: {{^}}test_workitem_id_x: 19; CO-V2: enable_vgpr_workitem_id = 0 20 21; ALL-NOT: v0 22; ALL: {{buffer|flat|global}}_store_{{dword|b32}} {{.*}}v0 23 24; PACKED-TID: .amdhsa_system_vgpr_workitem_id 0 25define amdgpu_kernel void @test_workitem_id_x(i32 addrspace(1)* %out) #1 { 26 %id = call i32 @llvm.amdgcn.workitem.id.x() 27 store i32 %id, i32 addrspace(1)* %out 28 ret void 29} 30 31; MESA: .section .AMDGPU.config 32; MESA: .long 47180 33; MESA-NEXT: .long 2180{{$}} 34 35; ALL-LABEL: {{^}}test_workitem_id_y: 36; CO-V2: enable_vgpr_workitem_id = 1 37; CO-V2-NOT: v1 38; CO-V2: {{buffer|flat}}_store_dword {{.*}}v1 39 40; PACKED-TID: v_bfe_u32 [[ID:v[0-9]+]], v0, 10, 10 41; PACKED-TID: {{buffer|flat|global}}_store_{{dword|b32}} {{.*}}[[ID]] 42; PACKED-TID: .amdhsa_system_vgpr_workitem_id 1 43define amdgpu_kernel void @test_workitem_id_y(i32 addrspace(1)* %out) #1 { 44 %id = call i32 @llvm.amdgcn.workitem.id.y() 45 store i32 %id, i32 addrspace(1)* %out 46 ret void 47} 48 49; MESA: .section .AMDGPU.config 50; MESA: .long 47180 51; MESA-NEXT: .long 4228{{$}} 52 53; ALL-LABEL: {{^}}test_workitem_id_z: 54; CO-V2: enable_vgpr_workitem_id = 2 55; CO-V2-NOT: v2 56; CO-V2: {{buffer|flat}}_store_dword {{.*}}v2 57 58; PACKED-TID: v_bfe_u32 [[ID:v[0-9]+]], v0, 20, 10 59; PACKED-TID: {{buffer|flat|global}}_store_{{dword|b32}} {{.*}}[[ID]] 60; PACKED-TID: .amdhsa_system_vgpr_workitem_id 2 61define amdgpu_kernel void @test_workitem_id_z(i32 addrspace(1)* %out) #1 { 62 %id = call i32 @llvm.amdgcn.workitem.id.z() 63 store i32 %id, i32 addrspace(1)* %out 64 ret void 65} 66 67; FIXME: Packed tid should avoid the and 68; ALL-LABEL: {{^}}test_reqd_workgroup_size_x_only: 69; CO-V2: enable_vgpr_workitem_id = 0 70 71; ALL-DAG: v_mov_b32_e32 [[ZERO:v[0-9]+]], 0{{$}} 72; UNPACKED-DAG: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, v0 73 74; PACKED: v_and_b32_e32 [[MASKED:v[0-9]+]], 0x3ff, v0 75; PACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[MASKED]] 76 77; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]] 78; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]] 79define amdgpu_kernel void @test_reqd_workgroup_size_x_only(i32* %out) !reqd_work_group_size !0 { 80 %id.x = call i32 @llvm.amdgcn.workitem.id.x() 81 %id.y = call i32 @llvm.amdgcn.workitem.id.y() 82 %id.z = call i32 @llvm.amdgcn.workitem.id.z() 83 store volatile i32 %id.x, i32* %out 84 store volatile i32 %id.y, i32* %out 85 store volatile i32 %id.z, i32* %out 86 ret void 87} 88 89; ALL-LABEL: {{^}}test_reqd_workgroup_size_y_only: 90; CO-V2: enable_vgpr_workitem_id = 1 91 92; ALL: v_mov_b32_e32 [[ZERO:v[0-9]+]], 0{{$}} 93; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]] 94 95; UNPACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, v1 96 97; PACKED: v_bfe_u32 [[MASKED:v[0-9]+]], v0, 10, 10 98; PACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[MASKED]] 99 100; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]] 101define amdgpu_kernel void @test_reqd_workgroup_size_y_only(i32* %out) !reqd_work_group_size !1 { 102 %id.x = call i32 @llvm.amdgcn.workitem.id.x() 103 %id.y = call i32 @llvm.amdgcn.workitem.id.y() 104 %id.z = call i32 @llvm.amdgcn.workitem.id.z() 105 store volatile i32 %id.x, i32* %out 106 store volatile i32 %id.y, i32* %out 107 store volatile i32 %id.z, i32* %out 108 ret void 109} 110 111; ALL-LABEL: {{^}}test_reqd_workgroup_size_z_only: 112; CO-V2: enable_vgpr_workitem_id = 2 113 114; ALL: v_mov_b32_e32 [[ZERO:v[0-9]+]], 0{{$}} 115; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]] 116; ALL: flat_store_{{dword|b32}} v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]] 117 118; UNPACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, v2 119 120; PACKED: v_bfe_u32 [[MASKED:v[0-9]+]], v0, 10, 20 121; PACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[MASKED]] 122define amdgpu_kernel void @test_reqd_workgroup_size_z_only(i32* %out) !reqd_work_group_size !2 { 123 %id.x = call i32 @llvm.amdgcn.workitem.id.x() 124 %id.y = call i32 @llvm.amdgcn.workitem.id.y() 125 %id.z = call i32 @llvm.amdgcn.workitem.id.z() 126 store volatile i32 %id.x, i32* %out 127 store volatile i32 %id.y, i32* %out 128 store volatile i32 %id.z, i32* %out 129 ret void 130} 131 132attributes #0 = { nounwind readnone } 133attributes #1 = { nounwind } 134 135!0 = !{i32 64, i32 1, i32 1} 136!1 = !{i32 1, i32 64, i32 1} 137!2 = !{i32 1, i32 1, i32 64} 138