[all-commits] [llvm/llvm-project] f5b216: [AMDGPU] Add amdgcn_sched_group_barrier builtin

[Austin Kerbow via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: f5b21680d1221d7acaa1b174d0b86fa907c71eb8
      https://github.com/llvm/llvm-project/commit/f5b21680d1221d7acaa1b174d0b86fa907c71eb8
  Author: Austin Kerbow <Austin.Kerbow at amd.com>
  Date:   2022-07-28 (Thu, 28 Jul 2022)

  Changed paths:
    M clang/include/clang/Basic/BuiltinsAMDGPU.def
    M clang/test/CodeGenOpenCL/builtins-amdgcn.cl
    M clang/test/SemaOpenCL/builtins-amdgcn-error.cl
    M llvm/include/llvm/IR/IntrinsicsAMDGPU.td
    M llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
    M llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp
    M llvm/lib/Target/AMDGPU/SIInstructions.td
    M llvm/lib/Target/AMDGPU/Utils/AMDGPUMemoryUtils.cpp
    A llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sched.group.barrier.ll
    A llvm/test/CodeGen/AMDGPU/sched-group-barrier-pre-RA.mir

  Log Message:
  -----------
  [AMDGPU] Add amdgcn_sched_group_barrier builtin

This builtin allows the creation of custom scheduling pipelines on a per-region
basis. Like the sched_barrier builtin this is intended to be used either for
testing, in situations where the default scheduler heuristics cannot be
improved, or in critical kernels where users are trying to get performance that
is close to handwritten assembly. Obviously using these builtins will require
extra work from the kernel writer to maintain the desired behavior.

The builtin can be used to create groups of instructions called "scheduling
groups" where ordering between the groups is enforced by the scheduler.
__builtin_amdgcn_sched_group_barrier takes three parameters. The first parameter
is a mask that determines the types of instructions that you would like to
synchronize around and add to a scheduling group. These instructions will be
selected from the bottom up starting from the sched_group_barrier's location
during instruction scheduling. The second parameter is the number of matching
instructions that will be associated with this sched_group_barrier. The third
parameter is an identifier which is used to describe what other
sched_group_barriers should be synchronized with. Note that multiple
sched_group_barriers must be added in order for them to be useful since they
only synchronize with other sched_group_barriers. Only "scheduling groups" with
a matching third parameter will have any enforced ordering between them.

As an example, the code below tries to create a pipeline of 1 VMEM_READ
instruction followed by 1 VALU instruction followed by 5 MFMA instructions...
// 1 VMEM_READ
__builtin_amdgcn_sched_group_barrier(32, 1, 0)
// 1 VALU
__builtin_amdgcn_sched_group_barrier(2, 1, 0)
// 5 MFMA
__builtin_amdgcn_sched_group_barrier(8, 5, 0)
// 1 VMEM_READ
__builtin_amdgcn_sched_group_barrier(32, 1, 0)
// 3 VALU
__builtin_amdgcn_sched_group_barrier(2, 3, 0)
// 2 VMEM_WRITE
__builtin_amdgcn_sched_group_barrier(64, 2, 0)

Reviewed By: jrbyrnes

Differential Revision: https://reviews.llvm.org/D128158