[llvm] [AMDGPU] Add documentation for scheduler intrinsics (PR #69854)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Oct 31 13:10:38 PDT 2023
https://github.com/bcahoon updated https://github.com/llvm/llvm-project/pull/69854
>From 9dc033d8caa5d9b3dd8ac5fde18c5bc4bd77ac19 Mon Sep 17 00:00:00 2001
From: Brendon Cahoon <brendon.cahoon at amd.com>
Date: Sat, 21 Oct 2023 13:01:24 -0500
Subject: [PATCH 1/2] [AMDGPU] Add documentation for scheduler intrinsics
Adding sched_barrier, sched_group_barrier, and iglp_opt.
---
llvm/docs/AMDGPUUsage.rst | 47 +++++++++++++++++++++++++++++++++++++++
1 file changed, 47 insertions(+)
diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index 9427df94e128e28..2e8991066d840f6 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -1098,6 +1098,53 @@ The AMDGPU backend implements the following LLVM IR intrinsics.
with the fifth i32 operand. The i1 sixth operand is used to clamp
the output. The i1s preceding the vector operands decide the signedness.
+ llvm.amdgcn.sched_barrier Controls the types of instructions that may be allowed to cross the intrinsic
+ during instruction scheduling. The parameter is a mask for the instruction types
+ that can cross the intrinsic.
+
+ - 0x0000: No instructions may be scheduled across sched_barrier.
+ - 0x0001: All, non-memory, non-side-effect producing instructions may be
+ scheduled across sched_barrier, *i.e.* allow ALU instructions to pass.
+ - 0x0002: VALU instructions may be scheduled across sched_barrier.
+ - 0x0004: SALU instructions may be scheduled across sched_barrier.
+ - 0x0008: MFMA/WMMA instructions may be scheduled across sched_barrier.
+ - 0x0010: All VMEM instructions may be scheduled across sched_barrier.
+ - 0x0020: VMEM read instructions may be scheduled across sched_barrier.
+ - 0x0040: VMEM write instructions may be scheduled across sched_barrier.
+ - 0x0080: All DS instructions may be scheduled across sched_barrier.
+ - 0x0100: All DS read instructions may be scheduled accoss sched_barrier.
+ - 0x0200: All DS write instructions may be scheduled across sched_barrier.
+
+ llvm.amdgcn.sched_group_barrier Creates schedule groups with specific properties to create custom scheduling
+ pipelines. The ordering between groups is enforced by the instruction scheduler.
+ The intrinsic applies to the code that preceeds the intrinsic. The intrinsic
+ takes three values that control the behavior of the schedule groups.
+
+ - Mask : Classify instruction groups using the llvm.amdgcn.sched_barrier mask values.
+ - Size : The number of instructions that are in the group.
+ - SyncID : Order is enforced between groups with matching values.
+
+ Combining multiple sched_group_barrier intrinsics enables an ordering of specific
+ instruction types during instruction scheduling. For example, the following enforces
+ a sequence of 1 VMEM read, 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)``
+
+ llvm.amdgcn.iglp_opt An **experimental** intrinsic for instruction group level parallelism. The intrinsic
+ implements predefined intruction scheduling orderings. The intrinsic applies to the
+ code that appears after the intrinsic. The intrinsic takes a value that specifies the
+ strategy. The compiler implements two strategies.
+
+ 0. Interleave DS and MFMA instructions for small GEMM kernels.
+ 1. Interleave DS and MFMA instructions for single wave small GEMM kernels.
+
+ The iglp_opt strategy implementations are subject to change.
============================================== ==========================================================
>From f3e7e24e8221864ee2a8936d40e06c046ae2d6f6 Mon Sep 17 00:00:00 2001
From: Brendon Cahoon <brendon.cahoon at amd.com>
Date: Tue, 31 Oct 2023 15:05:07 -0500
Subject: [PATCH 2/2] fixup! [AMDGPU] Add documentation for scheduler
intrinsics
---
llvm/docs/AMDGPUUsage.rst | 8 +++++++-
1 file changed, 7 insertions(+), 1 deletion(-)
diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index 2e8991066d840f6..2ab45d455279734 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -1124,6 +1124,9 @@ The AMDGPU backend implements the following LLVM IR intrinsics.
- Size : The number of instructions that are in the group.
- SyncID : Order is enforced between groups with matching values.
+ The mask can include multiple instruction types. It is undefined behavior to set
+ values beyond the range of valid masks.
+
Combining multiple sched_group_barrier intrinsics enables an ordering of specific
instruction types during instruction scheduling. For example, the following enforces
a sequence of 1 VMEM read, followed by 1 VALU instruction, followed by 5 MFMA
@@ -1138,12 +1141,15 @@ The AMDGPU backend implements the following LLVM IR intrinsics.
llvm.amdgcn.iglp_opt An **experimental** intrinsic for instruction group level parallelism. The intrinsic
implements predefined intruction scheduling orderings. The intrinsic applies to the
- code that appears after the intrinsic. The intrinsic takes a value that specifies the
+ surrounding scheduling region. The intrinsic takes a value that specifies the
strategy. The compiler implements two strategies.
0. Interleave DS and MFMA instructions for small GEMM kernels.
1. Interleave DS and MFMA instructions for single wave small GEMM kernels.
+ Only one iglp_opt intrinsic may be used in a scheduling region. The iglp_opt intrinsic
+ cannot be combined with sched_barrier or sched_group_barrier.
+
The iglp_opt strategy implementations are subject to change.
============================================== ==========================================================
More information about the llvm-commits
mailing list