AMDGPU: Fix test relying on incompatible attributes

[arsenm]

Summary

This combination of amdgpu-waves-per-eu and
amdgpu-flat-work-group-size cannot be satisfied at the same time, so
this was using the default.

[arsenm]

Add test, although it's hard to directly observe this. Also add some documentation

[rampitec]

llvm/docs/AMDGPUUsage.rst

872

This needs reformat to keep column width.

[arsenm]

Correct patch

[rampitec]

A test only NFC change? LGTM anyway.

[arsenm]

4c2ee571482e066c0a524bd393a4c61aeaecff8a

llvm/docs/AMDGPUUsage.rst

.. table:: AMDGPU LLVM IR Attributes
      :name: amdgpu-llvm-ir-attributes-table
 
      ======================================= ==========================================================
      LLVM Attribute                          Description
      ======================================= ==========================================================
      "amdgpu-flat-work-group-size"="min,max" Specify the minimum and maximum flat work group sizes that
                                              will be specified when the kernel is dispatched. Generated
                                              by the ``amdgpu_flat_work_group_size`` CLANG attribute [CLANG-ATTR]_.
+                                             The implied default value is 1,1024.
+
      "amdgpu-implicitarg-num-bytes"="n"      Number of kernel argument bytes to add to the kernel
                                              argument block size for the implicit arguments. This
                                              varies by OS and language (for OpenCL see
                                              :ref:`opencl-kernel-implicit-arguments-appended-for-amdhsa-os-table`).
      "amdgpu-num-sgpr"="n"                   Specifies the number of SGPRs to use. Generated by
                                              the ``amdgpu_num_sgpr`` CLANG attribute [CLANG-ATTR]_.
      "amdgpu-num-vgpr"="n"                   Specifies the number of VGPRs to use. Generated by the
                                              ``amdgpu_num_vgpr`` CLANG attribute [CLANG-ATTR]_.
      "amdgpu-waves-per-eu"="m,n"             Specify the minimum and maximum number of waves per
                                              execution unit. Generated by the ``amdgpu_waves_per_eu``
-                                             CLANG attribute [CLANG-ATTR]_.
+                                             CLANG attribute [CLANG-ATTR]_. This is an optimization hint, and
+                                             the backend may not be able to satisfy the request. If the specified range
+                                             is incompatible with the function's "amdgpu-flat-work-group-size" value,
+                                             the implied occupancy bounds by the workgroup size takes precedence.
+
      "amdgpu-ieee" true/false.               Specify whether the function expects the IEEE field of the
                                              mode register to be set on entry. Overrides the default for
                                              the calling convention.
      "amdgpu-dx10-clamp" true/false.         Specify whether the function expects the DX10_CLAMP field of
                                              the mode register to be set on entry. Overrides the default
                                              for the calling convention.
 
      "amdgpu-no-workitem-id-x"               Indicates the function does not depend on the value of the

llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp

std::pair<unsigned, unsigned> AMDGPUSubtarget::getWavesPerEU(
 
   // If minimum/maximum flat work group sizes were explicitly requested using
   // "amdgpu-flat-work-group-size" attribute, then set default minimum/maximum
   // number of waves per execution unit to values implied by requested
   // minimum/maximum flat work group sizes.
   unsigned MinImpliedByFlatWorkGroupSize =
     getWavesPerEUForWorkGroup(FlatWorkGroupSizes.second);
   Default.first = MinImpliedByFlatWorkGroupSize;
-  bool RequestedFlatWorkGroupSize =
-      F.hasFnAttribute("amdgpu-flat-work-group-size");
 
   // Requested minimum/maximum number of waves per execution unit.
   std::pair<unsigned, unsigned> Requested = AMDGPU::getIntegerPairAttribute(
     F, "amdgpu-waves-per-eu", Default, true);
 
   // Make sure requested minimum is less than requested maximum.
   if (Requested.second && Requested.first > Requested.second)
     return Default;
 
   // Make sure requested values do not violate subtarget's specifications.
   if (Requested.first < getMinWavesPerEU() ||
       Requested.second > getMaxWavesPerEU())
     return Default;
 
   // Make sure requested values are compatible with values implied by requested
   // minimum/maximum flat work group sizes.
-  if (RequestedFlatWorkGroupSize &&
-      Requested.first < MinImpliedByFlatWorkGroupSize)
+  if (Requested.first < MinImpliedByFlatWorkGroupSize)
     return Default;
 
   return Requested;
 }
 
 static unsigned getReqdWorkGroupSize(const Function &Kernel, unsigned Dim) {
   auto Node = Kernel.getMetadata("reqd_work_group_size");
   if (Node && Node->getNumOperands() == 3)

llvm/test/CodeGen/AMDGPU/default-flat-work-group-size-overrides-waves-per-eu.ll

+; RUN: opt -S -mtriple=amdgcn-unknown-unknown -mcpu=tahiti -amdgpu-promote-alloca -disable-promote-alloca-to-vector < %s | FileCheck %s
+
+; Both of these kernels have the same value for
+; amdgpu-flat-work-group-size, except one explicitly sets it. This is
+; a program visible property which should always take precedence over
+; the amdgpu-waves-per-eu optimization hint.
+;
+; The range is incompatible with the amdgpu-waves-per-eu value, so the
+; flat work group size should take precedence implying a requirement
+; to support 1024 size workgroups (which exceeds the available LDS
+; amount).
+
+; CHECK-NOT: @no_flat_workgroup_size.stack
+; CHECK-NOT: @explicit_default_workgroup_size.stack
+
+; CHECK-LABEL: @no_flat_workgroup_size(
+; CHECK: alloca [5 x i32]
+; CHECK: store i32 4, i32 addrspace(5)* %arrayidx1, align 4
+define amdgpu_kernel void @no_flat_workgroup_size(i32 addrspace(1)* nocapture %out, i32 addrspace(1)* nocapture %in) #0 {
+entry:
+  %stack = alloca [5 x i32], align 4, addrspace(5)
+  %0 = load i32, i32 addrspace(1)* %in, align 4
+  %arrayidx1 = getelementptr inbounds [5 x i32], [5 x i32] addrspace(5)* %stack, i32 0, i32 %0
+  store i32 4, i32 addrspace(5)* %arrayidx1, align 4
+  %arrayidx2 = getelementptr inbounds i32, i32 addrspace(1)* %in, i32 1
+  %1 = load i32, i32 addrspace(1)* %arrayidx2, align 4
+  %arrayidx3 = getelementptr inbounds [5 x i32], [5 x i32] addrspace(5)* %stack, i32 0, i32 %1
+  store i32 5, i32 addrspace(5)* %arrayidx3, align 4
+  %arrayidx10 = getelementptr inbounds [5 x i32], [5 x i32] addrspace(5)* %stack, i32 0, i32 0
+  %2 = load i32, i32 addrspace(5)* %arrayidx10, align 4
+  store i32 %2, i32 addrspace(1)* %out, align 4
+  %arrayidx12 = getelementptr inbounds [5 x i32], [5 x i32] addrspace(5)* %stack, i32 0, i32 1
+  %3 = load i32, i32 addrspace(5)* %arrayidx12
+  %arrayidx13 = getelementptr inbounds i32, i32 addrspace(1)* %out, i32 1
+  store i32 %3, i32 addrspace(1)* %arrayidx13
+  ret void
+}
+
+; CHECK-LABEL: @explicit_default_workgroup_size(
+; CHECK: alloca [5 x i32]
+; CHECK: store i32 4, i32 addrspace(5)* %arrayidx1, align 4
+define amdgpu_kernel void @explicit_default_workgroup_size(i32 addrspace(1)* nocapture %out, i32 addrspace(1)* nocapture %in) #1 {
+entry:
+  %stack = alloca [5 x i32], align 4, addrspace(5)
+  %0 = load i32, i32 addrspace(1)* %in, align 4
+  %arrayidx1 = getelementptr inbounds [5 x i32], [5 x i32] addrspace(5)* %stack, i32 0, i32 %0
+  store i32 4, i32 addrspace(5)* %arrayidx1, align 4
+  %arrayidx2 = getelementptr inbounds i32, i32 addrspace(1)* %in, i32 1
+  %1 = load i32, i32 addrspace(1)* %arrayidx2, align 4
+  %arrayidx3 = getelementptr inbounds [5 x i32], [5 x i32] addrspace(5)* %stack, i32 0, i32 %1
+  store i32 5, i32 addrspace(5)* %arrayidx3, align 4
+  %arrayidx10 = getelementptr inbounds [5 x i32], [5 x i32] addrspace(5)* %stack, i32 0, i32 0
+  %2 = load i32, i32 addrspace(5)* %arrayidx10, align 4
+  store i32 %2, i32 addrspace(1)* %out, align 4
+  %arrayidx12 = getelementptr inbounds [5 x i32], [5 x i32] addrspace(5)* %stack, i32 0, i32 1
+  %3 = load i32, i32 addrspace(5)* %arrayidx12
+  %arrayidx13 = getelementptr inbounds i32, i32 addrspace(1)* %out, i32 1
+  store i32 %3, i32 addrspace(1)* %arrayidx13
+  ret void
+}
+
+attributes #0 = { "amdgpu-waves-per-eu"="1,1" }
+attributes #1 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="1,1024" }