AMDGPU: Avoid enabling kernel workitem IDs with reqd_work_group_size

[arsenm]

de1600a1d946d1bd9ca5f7b0711012a8f31a66e2

llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp

if (!F.hasFnAttribute("amdgpu-no-workgroup-id-y"))
       WorkGroupIDY = true;
 
     if (!F.hasFnAttribute("amdgpu-no-workgroup-id-z"))
       WorkGroupIDZ = true;
 
     if (IsKernel || !F.hasFnAttribute("amdgpu-no-workitem-id-x"))
       WorkItemIDX = true;
 
-    if (!F.hasFnAttribute("amdgpu-no-workitem-id-y"))
+    if (!F.hasFnAttribute("amdgpu-no-workitem-id-y") &&
+        ST.getMaxWorkitemID(F, 1) != 0)
       WorkItemIDY = true;
 
-    if (!F.hasFnAttribute("amdgpu-no-workitem-id-z"))
+    if (!F.hasFnAttribute("amdgpu-no-workitem-id-z") &&
+        ST.getMaxWorkitemID(F, 2) != 0)
       WorkItemIDZ = true;
 
     if (!F.hasFnAttribute("amdgpu-no-dispatch-ptr"))
       DispatchPtr = true;
 
     if (!F.hasFnAttribute("amdgpu-no-queue-ptr"))
       QueuePtr = true;
 

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

define void @test_workitem_id_z_func(i32 addrspace(1)* %out) #1 {
   %id = call i32 @llvm.amdgcn.workitem.id.z()
   store i32 %id, i32 addrspace(1)* %out
   ret void
 }
 
 ; FIXME: Should be able to avoid enabling in kernel inputs
 ; FIXME: Packed tid should avoid the and
 ; ALL-LABEL: {{^}}test_reqd_workgroup_size_x_only:
-; CO-V2: enable_vgpr_workitem_id = 2
+; CO-V2: enable_vgpr_workitem_id = 0
 
 ; ALL-DAG: v_mov_b32_e32 [[ZERO:v[0-9]+]], 0{{$}}
 ; UNPACKED-DAG: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, v0
 
 ; PACKED: v_and_b32_e32 [[MASKED:v[0-9]+]], 0x3ff, v0
 ; PACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[MASKED]]
 
 ; ALL: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
 ; ALL: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
 define amdgpu_kernel void @test_reqd_workgroup_size_x_only(i32* %out) !reqd_work_group_size !0 {
   %id.x = call i32 @llvm.amdgcn.workitem.id.x()
   %id.y = call i32 @llvm.amdgcn.workitem.id.y()
   %id.z = call i32 @llvm.amdgcn.workitem.id.z()
   store volatile i32 %id.x, i32* %out
   store volatile i32 %id.y, i32* %out
   store volatile i32 %id.z, i32* %out
   ret void
 }
 
 ; ALL-LABEL: {{^}}test_reqd_workgroup_size_y_only:
-; CO-V2: enable_vgpr_workitem_id = 2
+; CO-V2: enable_vgpr_workitem_id = 1
 
 ; ALL: v_mov_b32_e32 [[ZERO:v[0-9]+]], 0{{$}}
 ; ALL: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
 
 ; UNPACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, v1
 
 ; PACKED: v_bfe_u32 [[MASKED:v[0-9]+]], v0, 10, 10
 ; PACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[MASKED]]

llvm/test/CodeGen/AMDGPU/call-reqd-group-size.ll

 ; CHECK-NEXT:    s_getpc_b64 s[4:5]
 ; CHECK-NEXT:    s_add_u32 s4, s4, callee@rel32@lo+4
 ; CHECK-NEXT:    s_addc_u32 s5, s5, callee@rel32@hi+12
 ; CHECK-NEXT:    s_swappc_b64 s[30:31], s[4:5]
 ; CHECK-NEXT:    s_endpgm
   call void @callee()
   ret void
 }
-; CHECK: .amdhsa_system_vgpr_workitem_id 2
+; CHECK: .amdhsa_system_vgpr_workitem_id 1
 
 define amdgpu_kernel void @known_yz_0(i32 addrspace(1)* %out) !reqd_work_group_size !3 {
 ; CHECK-LABEL: known_yz_0:
 ; CHECK:       ; %bb.0:
 ; CHECK-NEXT:    s_add_u32 flat_scratch_lo, s6, s9
 ; CHECK-NEXT:    s_addc_u32 flat_scratch_hi, s7, 0
 ; CHECK-NEXT:    s_add_u32 s0, s0, s9
 ; CHECK-NEXT:    s_addc_u32 s1, s1, 0
 ; CHECK-NEXT:    v_mov_b32_e32 v31, v0
 ; CHECK-NEXT:    s_mov_b32 s32, 0
 ; CHECK-NEXT:    s_getpc_b64 s[4:5]
 ; CHECK-NEXT:    s_add_u32 s4, s4, callee@rel32@lo+4
 ; CHECK-NEXT:    s_addc_u32 s5, s5, callee@rel32@hi+12
 ; CHECK-NEXT:    s_swappc_b64 s[30:31], s[4:5]
 ; CHECK-NEXT:    s_endpgm
   call void @callee()
   ret void
 }
-; CHECK: .amdhsa_system_vgpr_workitem_id 2
+; CHECK: .amdhsa_system_vgpr_workitem_id 0
 
 define amdgpu_kernel void @known_xz_0(i32 addrspace(1)* %out) !reqd_work_group_size !4 {
 ; CHECK-LABEL: known_xz_0:
 ; CHECK:       ; %bb.0:
 ; CHECK-NEXT:    s_add_u32 flat_scratch_lo, s6, s9
 ; CHECK-NEXT:    s_addc_u32 flat_scratch_hi, s7, 0
 ; CHECK-NEXT:    s_add_u32 s0, s0, s9
 ; CHECK-NEXT:    s_addc_u32 s1, s1, 0
 ; CHECK-NEXT:    v_lshlrev_b32_e32 v31, 10, v1
 ; CHECK-NEXT:    s_mov_b32 s32, 0
 ; CHECK-NEXT:    s_getpc_b64 s[4:5]
 ; CHECK-NEXT:    s_add_u32 s4, s4, callee@rel32@lo+4
 ; CHECK-NEXT:    s_addc_u32 s5, s5, callee@rel32@hi+12
 ; CHECK-NEXT:    s_swappc_b64 s[30:31], s[4:5]
 ; CHECK-NEXT:    s_endpgm
   call void @callee()
   ret void
 }
-; CHECK: .amdhsa_system_vgpr_workitem_id 2
+; CHECK: .amdhsa_system_vgpr_workitem_id 1
 
 
 define amdgpu_kernel void @known_xyz_0(i32 addrspace(1)* %out) !reqd_work_group_size !5 {
 ; CHECK-LABEL: known_xyz_0:
 ; CHECK:       ; %bb.0:
 ; CHECK-NEXT:    s_add_u32 flat_scratch_lo, s6, s9
 ; CHECK-NEXT:    s_addc_u32 flat_scratch_hi, s7, 0
 ; CHECK-NEXT:    s_add_u32 s0, s0, s9
 ; CHECK-NEXT:    s_addc_u32 s1, s1, 0
 ; CHECK-NEXT:    v_mov_b32_e32 v31, 0
 ; CHECK-NEXT:    s_mov_b32 s32, 0
 ; CHECK-NEXT:    s_getpc_b64 s[4:5]
 ; CHECK-NEXT:    s_add_u32 s4, s4, callee@rel32@lo+4
 ; CHECK-NEXT:    s_addc_u32 s5, s5, callee@rel32@hi+12
 ; CHECK-NEXT:    s_swappc_b64 s[30:31], s[4:5]
 ; CHECK-NEXT:    s_endpgm
   call void @callee()
   ret void
 }
-; CHECK: .amdhsa_system_vgpr_workitem_id 2
+; CHECK: .amdhsa_system_vgpr_workitem_id 0
 
 attributes #0 = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" }
 
 !0 = !{i32 1, i32 64, i32 64}
 !1 = !{i32 64, i32 1, i32 64}
 !2 = !{i32 64, i32 64, i32 1}
 !3 = !{i32 64, i32 1, i32 1}
 !4 = !{i32 1, i32 64, i32 1}
 !5 = !{i32 1, i32 1, i32 1}

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

 ; PACKED-TID: {{buffer|flat|global}}_store_dword {{.*}}[[ID]]
 ; PACKED-TID: .amdhsa_system_vgpr_workitem_id 2
 define amdgpu_kernel void @test_workitem_id_z(i32 addrspace(1)* %out) #1 {
   %id = call i32 @llvm.amdgcn.workitem.id.z()
   store i32 %id, i32 addrspace(1)* %out
   ret void
 }
 
-; FIXME: Should be able to avoid enabling in kernel inputs
 ; FIXME: Packed tid should avoid the and
 ; ALL-LABEL: {{^}}test_reqd_workgroup_size_x_only:
-; CO-V2: enable_vgpr_workitem_id = 2
+; CO-V2: enable_vgpr_workitem_id = 0
 
 ; ALL-DAG: v_mov_b32_e32 [[ZERO:v[0-9]+]], 0{{$}}
 ; UNPACKED-DAG: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, v0
 
 ; PACKED: v_and_b32_e32 [[MASKED:v[0-9]+]], 0x3ff, v0
 ; PACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[MASKED]]
 
 ; ALL: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
 ; ALL: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
 define amdgpu_kernel void @test_reqd_workgroup_size_x_only(i32* %out) !reqd_work_group_size !0 {
   %id.x = call i32 @llvm.amdgcn.workitem.id.x()
   %id.y = call i32 @llvm.amdgcn.workitem.id.y()
   %id.z = call i32 @llvm.amdgcn.workitem.id.z()
   store volatile i32 %id.x, i32* %out
   store volatile i32 %id.y, i32* %out
   store volatile i32 %id.z, i32* %out
   ret void
 }
 
 ; ALL-LABEL: {{^}}test_reqd_workgroup_size_y_only:
-; CO-V2: enable_vgpr_workitem_id = 2
+; CO-V2: enable_vgpr_workitem_id = 1
 
 ; ALL: v_mov_b32_e32 [[ZERO:v[0-9]+]], 0{{$}}
 ; ALL: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[ZERO]]
 
 ; UNPACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, v1
 
 ; PACKED: v_bfe_u32 [[MASKED:v[0-9]+]], v0, 10, 10
 ; PACKED: flat_store_dword v{{\[[0-9]+:[0-9]+\]}}, [[MASKED]]