llvm/test/CodeGen/AMDGPU/llvm.amdgcn.workitem.id.ll

   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
   9 declare i32 @llvm.amdgcn.workitem.id.x() #0
  10 declare i32 @llvm.amdgcn.workitem.id.y() #0
  11 declare i32 @llvm.amdgcn.workitem.id.z() #0
  12
  13 ; MESA: .section .AMDGPU.config
  14 ; MESA: .long 47180
  15 ; MESA-NEXT: .long 132{{$}}
  16
  17 ; ALL-LABEL: {{^}}test_workitem_id_x:
  18 ; CO-V2: enable_vgpr_workitem_id = 0
  19
  20 ; ALL-NOT: v0
  21 ; ALL: {{buffer|flat|global}}_store_dword {{.*}}v0
  22
  23 ; PACKED-TID: .amdhsa_system_vgpr_workitem_id 0
  24 define amdgpu_kernel void @test_workitem_id_x(i32 addrspace(1)* %out) #1 {
  25   %id = call i32 @llvm.amdgcn.workitem.id.x()
  26   store i32 %id, i32 addrspace(1)* %out
  27   ret void
  28 }
  29
  30 ; MESA: .section .AMDGPU.config
  31 ; MESA: .long 47180
  32 ; MESA-NEXT: .long 2180{{$}}
  33
  34 ; ALL-LABEL: {{^}}test_workitem_id_y:
  35 ; CO-V2: enable_vgpr_workitem_id = 1
  36 ; CO-V2-NOT: v1
  37 ; CO-V2: {{buffer|flat}}_store_dword {{.*}}v1
  38
  39 ; PACKED-TID: v_bfe_u32 [[ID:v[0-9]+]], v0, 10, 10
  40 ; PACKED-TID: {{buffer|flat|global}}_store_dword {{.*}}[[ID]]
  41 ; PACKED-TID: .amdhsa_system_vgpr_workitem_id 1
  42 define amdgpu_kernel void @test_workitem_id_y(i32 addrspace(1)* %out) #1 {
  43   %id = call i32 @llvm.amdgcn.workitem.id.y()
  44   store i32 %id, i32 addrspace(1)* %out
  45   ret void
  46 }
  47
  48 ; MESA: .section .AMDGPU.config
  49 ; MESA: .long 47180
  50 ; MESA-NEXT: .long 4228{{$}}
  51
  52 ; ALL-LABEL: {{^}}test_workitem_id_z:
  53 ; CO-V2: enable_vgpr_workitem_id = 2
  54 ; CO-V2-NOT: v2
  55 ; CO-V2: {{buffer|flat}}_store_dword {{.*}}v2
  56
  57 ; PACKED-TID: v_bfe_u32 [[ID:v[0-9]+]], v0, 20, 10
  58 ; PACKED-TID: {{buffer|flat|global}}_store_dword {{.*}}[[ID]]
  59 ; PACKED-TID: .amdhsa_system_vgpr_workitem_id 2
  60 define amdgpu_kernel void @test_workitem_id_z(i32 addrspace(1)* %out) #1 {
  61   %id = call i32 @llvm.amdgcn.workitem.id.z()
  62   store i32 %id, i32 addrspace(1)* %out
  63   ret void
  64 }
  65
  66 ; FIXME: Packed tid should avoid the and
  67 ; ALL-LABEL: {{^}}test_reqd_workgroup_size_x_only:
  68 ; CO-V2: enable_vgpr_workitem_id = 0
  69
  70 ; ALL-DAG: v_mov_b32_e32 [[ZERO:v[0-9]+]], 0{{$}}
  71 ; UNPACKED-DAG: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, v0
  72
  73 ; PACKED: v_and_b32_e32 [[MASKED:v[0-9]+]], 0x3ff, v0
  74 ; PACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[MASKED]]
  75
  76 ; ALL: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
  77 ; ALL: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
  78 define amdgpu_kernel void @test_reqd_workgroup_size_x_only(i32* %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, i32* %out
  83   store volatile i32 %id.y, i32* %out
  84   store volatile i32 %id.z, i32* %out
  85   ret void
  86 }
  87
  88 ; ALL-LABEL: {{^}}test_reqd_workgroup_size_y_only:
  89 ; CO-V2: enable_vgpr_workitem_id = 1
  90
  91 ; ALL: v_mov_b32_e32 [[ZERO:v[0-9]+]], 0{{$}}
  92 ; ALL: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
  93
  94 ; UNPACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, v1
  95
  96 ; PACKED: v_bfe_u32 [[MASKED:v[0-9]+]], v0, 10, 10
  97 ; PACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[MASKED]]
  98
  99 ; ALL: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
 100 define amdgpu_kernel void @test_reqd_workgroup_size_y_only(i32* %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, i32* %out
 105   store volatile i32 %id.y, i32* %out
 106   store volatile i32 %id.z, i32* %out
 107   ret void
 108 }
 109
 110 ; ALL-LABEL: {{^}}test_reqd_workgroup_size_z_only:
 111 ; CO-V2: enable_vgpr_workitem_id = 2
 112
 113 ; ALL: v_mov_b32_e32 [[ZERO:v[0-9]+]], 0{{$}}
 114 ; ALL: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
 115 ; ALL: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
 116
 117 ; UNPACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, v2
 118
 119 ; PACKED: v_bfe_u32 [[MASKED:v[0-9]+]], v0, 10, 20
 120 ; PACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[MASKED]]
 121 define amdgpu_kernel void @test_reqd_workgroup_size_z_only(i32* %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, i32* %out
 126   store volatile i32 %id.y, i32* %out
 127   store volatile i32 %id.z, i32* %out
 128   ret void
 129 }
 130
 131 attributes #0 = { nounwind readnone }
 132 attributes #1 = { nounwind }
 133
 134 !0 = !{i32 64, i32 1, i32 1}
 135 !1 = !{i32 1, i32 64, i32 1}
 136 !2 = !{i32 1, i32 1, i32 64}