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

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