AMDGPU Reduce reported maximum group size to 1024

[arsenm]

Summary

While some targets allow encoding 2048, this was never tested or
supported.

[mareko]

It's supported by Mesa on gfx6-8. I'm OK with this though.

[arsenm]

In D66812#1647850, @mareko wrote:

It's supported by Mesa on gfx6-8. I'm OK with this though.

Hardware people told me they never tested it and they don’t expect it to be used

[arsenm]

25c5da5a426168b38fb3e9baa918faa75e4a92b4

lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp

return alignTo(getMaxWavesPerCU(STI, FlatWorkGroupSize),
                  getEUsPerCU(STI)) / getEUsPerCU(STI);
 }
 
 unsigned getMinFlatWorkGroupSize(const MCSubtargetInfo *STI) {
   return 1;
 }
 
 unsigned getMaxFlatWorkGroupSize(const MCSubtargetInfo *STI) {
-  return 2048;
+  // Some subtargets allow encoding 2048, but this isn't tested or supported.
+  return 1024;
 }
 
 unsigned getWavesPerWorkGroup(const MCSubtargetInfo *STI,
                               unsigned FlatWorkGroupSize) {
   return alignTo(FlatWorkGroupSize, getWavefrontSize(STI)) /
                  getWavefrontSize(STI);
 }
 

test/CodeGen/AMDGPU/attr-amdgpu-flat-work-group-size-v3.ll

 ; CHECK: NumSGPRsForWavesPerEU: 1
 ; CHECK: NumVGPRsForWavesPerEU: 1
 define amdgpu_kernel void @min_128_max_128() #2 {
 entry:
   ret void
 }
 attributes #2 = {"amdgpu-flat-work-group-size"="128,128"}
 
-; CHECK-LABEL: {{^}}min_1024_max_2048
-; CHECK: SGPRBlocks: 1
-; CHECK: VGPRBlocks: 7
-; CHECK: NumSGPRsForWavesPerEU: 12
-; CHECK: NumVGPRsForWavesPerEU: 32
+; CHECK-LABEL: {{^}}min_1024_max_1024
+; CHECK: SGPRBlocks: 0
+; CHECK: VGPRBlocks: 10
+; CHECK: NumSGPRsForWavesPerEU: 2{{$}}
+; CHECK: NumVGPRsForWavesPerEU: 43
 @var = addrspace(1) global float 0.0
-define amdgpu_kernel void @min_1024_max_2048() #3 {
+define amdgpu_kernel void @min_1024_max_1024() #3 {
   %val0 = load volatile float, float addrspace(1)* @var
   %val1 = load volatile float, float addrspace(1)* @var
   %val2 = load volatile float, float addrspace(1)* @var
   %val3 = load volatile float, float addrspace(1)* @var
   %val4 = load volatile float, float addrspace(1)* @var
   %val5 = load volatile float, float addrspace(1)* @var
   %val6 = load volatile float, float addrspace(1)* @var
   %val7 = load volatile float, float addrspace(1)* @var
define amdgpu_kernel void @min_1024_max_1024() #3 {
   store volatile float %val36, float addrspace(1)* @var
   store volatile float %val37, float addrspace(1)* @var
   store volatile float %val38, float addrspace(1)* @var
   store volatile float %val39, float addrspace(1)* @var
   store volatile float %val40, float addrspace(1)* @var
 
   ret void
 }
-attributes #3 = {"amdgpu-flat-work-group-size"="1024,2048"}
+attributes #3 = {"amdgpu-flat-work-group-size"="1024,1024"}
 
 ; CHECK: amdhsa.kernels:
 ; CHECK:   .max_flat_workgroup_size: 64
 ; CHECK:   .name:                 min_64_max_64
 ; CHECK:   .max_flat_workgroup_size: 128
 ; CHECK:   .name:                 min_64_max_128
 ; CHECK:   .max_flat_workgroup_size: 128
 ; CHECK:   .name:                 min_128_max_128
-; CHECK:   .max_flat_workgroup_size: 2048
-; CHECK:   .name:                 min_1024_max_2048
+; CHECK:   .max_flat_workgroup_size: 1024
+; CHECK:   .name:                 min_1024_max_1024
 ; CHECK: amdhsa.version:
 ; CHECK:   - 1
 ; CHECK:   - 0
 
 ; PARSER: AMDGPU HSA Metadata Parser Test: PASS

test/CodeGen/AMDGPU/attr-amdgpu-flat-work-group-size.ll

 ; CHECK: NumSGPRsForWavesPerEU: 1
 ; CHECK: NumVGPRsForWavesPerEU: 1
 define amdgpu_kernel void @min_128_max_128() #2 {
 entry:
   ret void
 }
 attributes #2 = {"amdgpu-flat-work-group-size"="128,128"}
 
-; CHECK-LABEL: {{^}}min_1024_max_2048
-; CHECK: SGPRBlocks: 1
-; CHECK: VGPRBlocks: 7
-; CHECK: NumSGPRsForWavesPerEU: 12
-; CHECK: NumVGPRsForWavesPerEU: 32
+; CHECK-LABEL: {{^}}min_1024_max_1024
+; CHECK: SGPRBlocks: 0
+; CHECK: VGPRBlocks: 10
+; CHECK: NumSGPRsForWavesPerEU: 2{{$}}
+; CHECK: NumVGPRsForWavesPerEU: 43
 @var = addrspace(1) global float 0.0
-define amdgpu_kernel void @min_1024_max_2048() #3 {
+define amdgpu_kernel void @min_1024_max_1024() #3 {
   %val0 = load volatile float, float addrspace(1)* @var
   %val1 = load volatile float, float addrspace(1)* @var
   %val2 = load volatile float, float addrspace(1)* @var
   %val3 = load volatile float, float addrspace(1)* @var
   %val4 = load volatile float, float addrspace(1)* @var
   %val5 = load volatile float, float addrspace(1)* @var
   %val6 = load volatile float, float addrspace(1)* @var
   %val7 = load volatile float, float addrspace(1)* @var
define amdgpu_kernel void @min_1024_max_1024() #3 {
   store volatile float %val36, float addrspace(1)* @var
   store volatile float %val37, float addrspace(1)* @var
   store volatile float %val38, float addrspace(1)* @var
   store volatile float %val39, float addrspace(1)* @var
   store volatile float %val40, float addrspace(1)* @var
 
   ret void
 }
-attributes #3 = {"amdgpu-flat-work-group-size"="1024,2048"}
+attributes #3 = {"amdgpu-flat-work-group-size"="1024,1024"}
 
 ; HSAMD: NT_AMD_AMDGPU_HSA_METADATA (HSA Metadata)
 ; HSAMD: Version: [ 1, 0 ]
 ; HSAMD: Kernels:
 ; HSAMD: - Name:                 min_64_max_64
 ; HSAMD:   MaxFlatWorkGroupSize: 64
 ; HSAMD: - Name:                 min_64_max_128
 ; HSAMD:   MaxFlatWorkGroupSize: 128
 ; HSAMD: - Name:                 min_128_max_128
 ; HSAMD:   MaxFlatWorkGroupSize: 128
-; HSAMD: - Name:                 min_1024_max_2048
-; HSAMD:   MaxFlatWorkGroupSize: 2048
+; HSAMD: - Name:                 min_1024_max_1024
+; HSAMD:   MaxFlatWorkGroupSize: 1024

test/CodeGen/AMDGPU/large-work-group-promote-alloca.ll

entry:
   store i32 %2, i32 addrspace(1)* %out, align 4
   %arrayidx12 = getelementptr inbounds [5 x i32], [5 x i32]* %stack, i32 0, i32 1
   %3 = load i32, i32* %arrayidx12
   %arrayidx13 = getelementptr inbounds i32, i32 addrspace(1)* %out, i32 1
   store i32 %3, i32 addrspace(1)* %arrayidx13
   ret void
 }
 
-; SICI: @promote_alloca_size_1600.stack = internal unnamed_addr addrspace(3) global [1600 x [5 x i32]] undef, align 4
-; GFX10: alloca [5 x i32]
+; SI-NOT: @promote_alloca_size_1600.stack
+; CI: @promote_alloca_size_1600.stack = internal unnamed_addr addrspace(3) global [1024 x [5 x i32]] undef, align 4
+; GFX10: @promote_alloca_size_1600.stack = internal unnamed_addr addrspace(3) global [1024 x [5 x i32]] undef, align 4
 
 define amdgpu_kernel void @promote_alloca_size_1600(i32 addrspace(1)* nocapture %out, i32 addrspace(1)* nocapture %in) #2 {
 entry:
   %stack = alloca [5 x i32], align 4
   %0 = load i32, i32 addrspace(1)* %in, align 4
   %arrayidx1 = getelementptr inbounds [5 x i32], [5 x i32]* %stack, i32 0, i32 %0
   store i32 4, i32* %arrayidx1, align 4
   %arrayidx2 = getelementptr inbounds i32, i32 addrspace(1)* %in, i32 1
entry:
   %tmp3 = load i8, i8* %arrayidx12, align 1
   %arrayidx13 = getelementptr inbounds i8, i8 addrspace(1)* %out, i64 1
   store i8 %tmp3, i8 addrspace(1)* %arrayidx13, align 1
   ret void
 }
 
 attributes #0 = { nounwind "amdgpu-flat-work-group-size"="63,63" }
 attributes #1 = { nounwind "amdgpu-waves-per-eu"="1,3" "amdgpu-flat-work-group-size"="256,256" }
-attributes #2 = { nounwind "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="1600,1600" }
+attributes #2 = { nounwind "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="1024,1024" }
 attributes #3 = { nounwind "amdgpu-waves-per-eu"="1,10" }
 attributes #4 = { nounwind "amdgpu-waves-per-eu"="1,10" }
 attributes #5 = { nounwind "amdgpu-waves-per-eu"="1,6" "amdgpu-flat-work-group-size"="64,64" }
 attributes #6 = { nounwind "amdgpu-waves-per-eu"="1,8" "amdgpu-flat-work-group-size"="64,64" }
 attributes #7 = { nounwind "amdgpu-waves-per-eu"="1,9" "amdgpu-flat-work-group-size"="64,64" }