Mma Class Reference#
Mma Class Reference
Driver for the wave-tile Mma operation. Given a backend block-wise MmaOp implementation (e.g., mfma or wmma), this class performs block-wise decomposition to matrix-multiply input fragments of (A: FragM x FragK) x (B: FragK x FragN) and accumulates results into output fragment (C: FragM x FragN). More...
#include <mma.hpp>
Detailed Description
Driver for the wave-tile Mma operation. Given a backend block-wise MmaOp implementation (e.g., mfma or wmma), this class performs block-wise decomposition to matrix-multiply input fragments of (A: FragM x FragK) x (B: FragK x FragN) and accumulates results into output fragment (C: FragM x FragN).
- Template Parameters
-
ADataType Data type of input fragment A BDataType Data type of input fragment B CDataType Data type of input/output fragment C (accumulator) FragM Mma fragment M dimension FragN Mma fragment K dimension FragK Mma fragment M dimension AccumPolicy The block order of the accumulation registers (row major or col major block order) CompilerTarget The compiler target MmaOp The backend wrapper class that will perform block-wise mma op (e.g., mfma or wmma) MmaTransforms The set of transforms to be applied to input/output fragments
- This is an example of an Mma decomposition driver class that can be used in a wave-tile
- context. Given a fragment size, we can decompose the fragment into smaller block-wise mma ops that are natively supported by the hardware (e.g., mfma or wmma). The class also supports applying transforms to the input/output fragments as needed (e.g., layout conversions, data type conversions, etc.). We may also specify the accumulation order (row-major or col-major) for the output fragment. This is a powerful example of how to build a flexible and reusable mma driver that can adapt to different hardware capabilities and requirements.
The documentation for this class was generated from the following file:
- /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/arch/mma/mma.hpp