PacQ: A SIMT Microarchitecture for Efficient Dataflow in Hyper-asymmetric GEMMs

TL;DR

PacQ tackles the inefficiency of deploying weight-only quantized LLMs on SIMT hardware by keeping INT4/INT2 weights packed through the GEMM stack and avoiding unpacking/dequantization overhead. It introduces a hyper-asymmetric GEMM framework that packs INT weights along the output feature dimension and uses an output-stationary dataflow, complemented by a parallel FP-INT multiplier integrated into a new SIMT microarchitecture, PacQ. The key contributions include the packing/dataflow strategy that reduces register-file traffic, the parallel FP-INT unit that boosts throughput and energy efficiency, and extensive evaluation showing up to 1.99x throughput and up to 81.4% EDP reduction on representative LLM workloads, with insights into PTQ-grouping for iso-perplexity. Overall, PacQ provides a practical design-space exploration for efficient hyper-asymmetric GEMM acceleration on SIMT architectures, enabling more efficient deployment of weight-only quantized LLMs.

Abstract

Weight-only quantization has been widely explored in large language models (LLMs) to reduce memory storage and data loading overhead. During deployment on single-instruction-multiple-threads (SIMT) architectures, weights are stored in low-precision integer (INT) format, while activations remain in full-precision floating-point (FP) format to preserve inference accuracy. Although memory footprint and data loading requirements for weight matrices are reduced, computation performance gains remain limited due to the need to convert weights back to FP format through unpacking and dequantization before GEMM operations. In this work, we investigate methods to accelerate GEMM operations involving packed low-precision INT weights and high-precision FP activations, defining this as the hyper-asymmetric GEMM problem. Our approach co-optimizes tile-level packing and dataflow strategies for INT weight matrices. We further design a specialized FP-INT multiplier unit tailored to our packing and dataflow strategies, enabling parallel processing of multiple INT weights. Finally, we integrate the packing, dataflow, and multiplier unit into PacQ, a SIMT microarchitecture designed to efficiently accelerate hyper-asymmetric GEMMs. We show that PacQ can achieve up to 1.99x speedup and 81.4% reduction in EDP compared to weight-only quantized LLM workloads running on conventional SIMT baselines.

Figures (12)

• Figure 1: (a) Standard inference flow on SIMT architecture for weight-only quantized LLMs. Green denotes high-precision floating-point activations, and blue denotes low-precision integer weights. (b) Our proposed flow leverages the hyper-asymmetric GEMM to achieve the benefits of reduced loading, storage, and compute costs from low-precision integer weights throughout the entire GEMM computation stack.
• Figure 2: IEEE 754 standard FP16 format.
• Figure 3: (a) & (b): The tile mapping of the warp-level (a set of 32 threads) mma.sync.m16n16k16 instruction to different octets (a subset of 8 threads). Here, $A$ represents the input matrix with a shape of [$m,k$], $B$ represents the weight matrix with a shape of [$k,n$], and $C$ represents the output (partial-sum) matrix with a shape of [$m,n$]. (c) An example of a step-by-step breakdown of the octet-level data movement and computation flow. (d) The mapping of octet-level workloads onto the baseline tensor processing units. DP-4 stands for four-element dot-product units that performs $4\times4$ inner-product.
• Figure 4: (a) Illustration of multiple fetch instructions required by the hyper-asymmetric GEMM when packing weights along the $k$ dimension. (b) Illustration of the poor data reuse of $A$ in $k$ dimension packing. (c) Illustration of the improved data reuse and reduction in fetch instructions achieved by packing along the $n$ dimension. The tile size is identical to Figure. \ref{['fig:dataflow:baseline']}.
• Figure 5: (a) Standard FP multiplier design. $s$ stands for sign, $e$ represents exponent, and $m$ denotes the mantissa. (b) Overview of our proposed parallel FP-INT-16 multiplier. In the provided example, our design can generate four output mantissas in one cycle. (c) Detailed diagram of the modifications made to the original 11-bit integer multiplier. The elements in purple are part of the original multiplier, while the elements in white represent additional units. All bits of $B$'s mantissas perform a logical AND operation with shifted $A$'s mantissas before being fed into the adders. (d) Explanation of how the final output mantissa (shaded gray) is assembled from different value sources. $\{\}$ denotes the append operation.