This AMD GPUOpen blog post introduces the utilization of the 3rd-generation Matrix Cores in AMD RDNA 4 architecture GPUs to significantly accelerate Generalized Matrix Multiplication (GEMM) operations. It highlights substantial performance improvements over previous generations and details the updated Wave Matrix Multiply Accumulate (WMMA) intrinsics, specifically tailored for RDNA 4 with a simpler VGPR layout that is not backward compatible with RDNA 3. The article provides practical guidance on integrating these new intrinsics within HIP kernels, including a simplified data mapping scheme. It further demonstrates the chaining of WMMA operations by implementing a Multi-Layer Perceptron (MLP) inference, emphasizing the RDNA 4's streamlined approach to sequential matrix computations without the need for inter-lane data shuffling.
Using the Matrix Cores of AMD RDNA 4 architecture GPUs: A Deep Dive into WMMA Operations
Enhanced Matrix Core Performance: AMD RDNA 4 architecture GPUs feature 3rd-generation Matrix Cores, which significantly boost Generalized Matrix Multiplication (GEMM) operations. These show considerable performance improvements, with FP16 FLOPS/clock/CU doubling from RDNA 3 and quadrupling from RDNA 2.
Wave Matrix Multiply Accumulate (WMMA) Operations: WMMA operations are designed for 16x16 matrix multiplications (D=AB+C). For larger matrices, decomposition is required, while smaller ones need padding. Unlike traditional HIP kernel operations, WMMA intrinsics are designed to be executed across an entire wavefront simultaneously.
New VGPR Layout for RDNA 4: RDNA 4 GPUs introduce a new VGPR (Vector General Purpose Register) layout for WMMA intrinsics, identifiable by a _gfx12 postfix. This layout is not backward compatible with RDNA 3 but is much simpler, eliminating the need for duplicating elements for A and B matrices.
Reduced VGPR Pressure: In the RDNA 4 WMMA scheme, each lane in a wavefront allocates smaller VGPRs to store only a part of a matrix (8 elements). This design allows a 32-lane wavefront to load an entire 16x16 matrix precisely once, reducing overall VGPR pressure.
Implementing WMMA in HIP: The article provides example HIP kernel code demonstrating how to use the new WMMA intrinsics (e.g., __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12). It highlights the specific matrix mapping where B, C, and D matrices are row major, while A is transposed (column major).
Efficient Data Packing: A helper intrinsic, packFp32s, is introduced to efficiently convert and pack two 32-bit floats into a single 32-bit VGPR, which is particularly useful for optimizing the conversion of 32-bit float matrices to 16-bit float matrices in contexts like matrix multiplication chains.
Simplified WMMA Chaining: A significant advantage of the RDNA 4 WMMA design is that it simplifies chaining matrix operations. Unlike RDNA 3, which required data shuffling among lanes to convert output (D matrix) for input (B matrix) in subsequent operations, RDNA 4's consistent layout eliminates this need.
MLP Inference Application: The practical utility of WMMA intrinsics is demonstrated through the implementation of a simple Multi-Layer Perceptron (MLP) inference kernel, showcasing how sequential matrix multiplications can be efficiently mapped and executed.
Further Learning: The article encourages readers to explore how to apply these WMMA intrinsics to GEMM operations involving larger matrices and refers to the AMD ISA guide for more in-depth technical details.