Inference Optimization2024~124 minarXiv:2407.00088

T-MAC: CPU Renaissance via Table Lookup for Low-Bit LLM Deployment on Edge

score 8Jianyu Wei, Shijie Cao, Ting Cao +4↑ 12★ 924Jun 25, 2024cs.DCdoi https://github.com/microsoft/t-mac

🎯 Pitch

Reducing a weight’s bit width actually slows down existing CPU inference frameworks because dequantization overhead dominates. T-MAC eliminates this tax entirely by replacing multiplications with bit-wise table lookups, enabling a Raspberry Pi 5 to pump out 11 tokens/s for a 1.58-bit model and scaling cleanly as bits decrease.

1. Executive Summary

T-MAC introduces a lookup-table (LUT)-based method for efficient low-bit LLM inference on CPUs, transforming the traditional data-type-centric multiplication into bit-wise table lookup (e.g., decomposing an mpGEMM between activation and an n-bit weight matrix into n serial lookups against a precomputed table, enabling unified support for any mixed-precision combination like W4A16 or W2A16 without dequantization). Evaluated on low-bit Llama and BitNet models across edge devices including M2-Ultra, Jetson AGX Orin, and Raspberry Pi 5, T-MAC achieves up to 4× throughput increase and 70% energy reduction versus llama.cpp—delivering 30 tokens/s on a single core and 71 tokens/s on eight cores of M2-Ultra for BitNet-b1.58-3B, and 11 tokens/s even on a Raspberry Pi 5—establishing that CPUs can match or exceed GPU inference speed for low-bit LLMs when the mixed-precision computation is handled natively through table lookups rather than indirect dequantization.

2. Context and Motivation

The Core Problem: Mixed-Precision GEMM Without Dequantization

The paper addresses a specific, concrete bottleneck that arises when deploying low-bit (weight-quantized) large language models on edge devices. The core problem is deceptively simple: existing hardware does not natively support mixed-precision matrix multiplication (mpGEMM), where low-precision weights (e.g., 2-bit, 3-bit, 4-bit) are multiplied by higher-precision activations (e.g., FP16, INT8). Because hardware instructions expect symmetric operand precision, current inference systems must first dequantize the low-bit weights up to match the activation precision before computing, then perform the multiplication at full precision. This dequantization step creates a computational tax that fundamentally undermines the benefit of reducing weight bit-width.

This is not merely an efficiency inconvenience — it is a structural limitation that prevents low-bit quantization from delivering its full promise on edge CPUs. The paper provides direct empirical evidence of this pathology: as weight precision decreases from 4-bit to 2-bit, llama.cpp — the state-of-the-art CPU inference framework — actually experiences a 15% slowdown at 3-bit relative to 4-bit due to decoding overhead (Figure 6). The intuitive expectation that "fewer bits = faster computation" breaks down because the dequantization overhead grows faster than the arithmetic savings shrink. The paper characterizes this as a fundamental misalignment between how low-bit models are represented and how hardware is designed to compute.

Why This Matters: The Edge Deployment Imperative

The problem's importance stems from an accelerating trend in LLM deployment: the migration of models from datacenter GPUs to edge devices — smartphones, laptops, robotics, and even single-board computers like the Raspberry Pi. The paper cites several concrete examples of this trend: Phi-3-mini-4bit deployed on iPhone, Llama-2-7B-4bit on Pixel 5, Llama-2-13B-4bit on Apple M2 Ultra, and Microsoft's Copilot+PC architecture that collaboratively runs on-device LLMs with cloud models (Section 1, Introduction).

The driving forces behind this migration are weighty but familiar:

Latency: On-device processing eliminates network round-trips, enabling real-time responses for applications like autonomous vehicles and interactive robotics (Section 2.1).
Privacy: Local data processing keeps sensitive user information on-device, reducing exposure risk and regulatory burden.
Reliability: Network independence means the LLM functions regardless of connectivity — critical for mobile, embedded, and mission-critical deployments.
Cost: Offloading inference from cloud GPUs to user-owned edge hardware shifts the economic burden and can dramatically reduce serving costs at scale.

However, edge deployment faces a fundamental resource constraint: memory capacity. The paper notes that Llama-2-7B in FP16 requires at least 14 GB of memory just to host the model parameters, while edge devices have limited RAM (Table 2 shows devices ranging from 819.2 GB/s bandwidth on M2-Ultra down to 17.1 GB/s on Raspberry Pi 5). Weight quantization — reducing parameters from 16-bit to 4-bit, 3-bit, or even 1-bit — is the primary technique for shrinking this memory footprint, making it "a must-have for on-device LLM inference" (Section 2.2).

But here is where the tension crystallizes: weight quantization solves the memory capacity problem, but creates a computational efficiency problem through the mixed-precision mismatch. The net result is that edge deployments are caught between two suboptimal choices: use higher-precision weights and suffer from memory pressure, or use lower-precision weights and suffer from dequantization overhead. Neither path fully exploits the hardware's potential.

Beyond capacity, the paper identifies two additional edge deployment challenges that compound the issue (Section 2.1):

Memory bandwidth bottleneck in the decode phase: LLM token generation is memory-bound — each new token requires loading the entire model weight matrix from DRAM and performing matrix-vector multiplications (GEMV). On edge devices with limited memory bandwidth, this becomes the dominant latency source. Any computational approach that fails to fully utilize the available bandwidth leaves performance on the table.
Energy efficiency for battery-operated devices: Smartphones, robotics, and IoT sensors operate on finite energy budgets. The paper explicitly frames power consumption as "particularly critical" (Section 2.1), and evaluates energy per token as a primary metric alongside throughput (Section 5.4).

Why Activations Cannot Simply Be Quantized Too

A natural question is: if mixed precision is the problem, why not quantize both weights and activations to the same low bit-width? The paper provides a brief but crucial answer in Section 1: "activation quantization cannot follow the trend due to outliers." This refers to the well-documented phenomenon in LLMs where a small fraction of activation channels exhibit values that are orders of magnitude larger than the typical activation range. Quantizing these outlier channels to low precision (e.g., INT4 or INT2) causes catastrophic accuracy degradation. Techniques like LLM.int8() (Dettmers et al., 2022) attempt to handle this by isolating outlier dimensions into higher precision, but they still require asymmetric operand handling.

The implication is profound: the mixed-precision mismatch between low-bit weights and high-precision activations is not a temporary engineering limitation — it is a structural consequence of how transformer models behave during inference. Weights can be aggressively quantized because their distribution is well-behaved, but activations cannot due to outliers. Therefore, any inference system that wants to benefit from aggressive weight quantization must efficiently handle this asymmetry natively, rather than papering over it with dequantization.

Prior Approaches and Their Shortcomings

The paper surveys three categories of existing approaches, each with specific limitations that motivate T-MAC's design.

1. Dequantization-based inference systems

The dominant approach in production systems — including llama.cpp, Intel Neural Compressor, TensorRT-LLM, and vLLM (Section 2.3, Section 6) — is to dequantize weights at runtime to match activation precision, then perform standard high-precision GEMM. For example, an INT4 weight is unpacked to INT8 or FP16 in registers before the multiply-accumulate, and the computation proceeds as if both operands were high-precision.

The shortcomings are threefold:

Non-scalable performance with bit reduction: As discussed above, the dequantization overhead often negates or even reverses the expected speedup from fewer bits. The paper's Figure 6 shows this concretely: llama.cpp's mpGEMV latency on 3-bit models is higher than 4-bit across all tested matrix shapes and devices, because unpacking 3-bit packed representations requires complex shift-and-mask sequences that add more overhead than the multiplication savings remove.
Case-by-case kernel design: Each bit-width requires its own weight layout, packing scheme, and kernel implementation. The paper notes that "the data layouts, as well as the interleaving or swizzling methods for W3 and W2 are totally different" (Section 1). A W3 layout might pack 2 bits in one byte and the remaining 1 bit separately, while W4 uses uniform nibble packing. Each layout demands its own un-packing logic, instruction sequences, and register allocation strategy. This creates a combinatorial explosion of kernel variants as the field explores new bit-widths (2-bit, 1.58-bit ternary, etc.) — each requiring separate development, optimization, and maintenance.
No unified abstraction: Because each kernel is tied to a specific hardware data type (e.g., INT8 dot-product on NEON), there is no conceptual framework that cleanly separates the bit-width representation from the computation. This makes the system fragile to new quantization formats and hardware backends.

2. GPU-based LUT methods

The idea of using lookup tables to avoid multiplications in quantized neural networks is not new. For CNNs with both weights and activations quantized to low precision, DeepGEMM (Ganji et al., 2023) precomputes all possible partial products and stores them in a lookup table, converting convolution to table lookups. MADDNESS (Blalock and Guttag, 2021) and LUT-NN (Tang et al., 2023) apply similar principles to vector-quantized models.

For low-bit LLMs specifically, LUT-GEMM (Park et al., 2023) and related GPU implementations (Maleki, 2023; Section 2.4) pioneered the application of LUT-based computation to weight-only quantized models. These methods decompose the weight matrix into bit-serial form, precompute activation-bit-pattern products into LUTs stored in GPU shared memory, and perform table lookups indexed by weight-bit groups.

However, the paper presents a striking negative result: practical GPU LUT-GEMM kernel performance is worse than dequantization-based kernels. Specifically, when tested with real Llama-2 matrix shapes on an A100 GPU, LUT-GEMM latency was 2.34×, 1.87×, and 1.75× slower than the dequantization-based kernels in BitBLAS for W4A16, W2A16, and W1A16 respectively (Section 2.4). The authors attribute this to the GPU's fixed architecture providing either inadequate storage capacity for the lookup tables or insufficiently rapid table access — the LUTs compete for scarce shared memory, and random access patterns from weight-group indexing conflict with the GPU's SIMD execution model optimized for coalesced memory access.

This is a crucial empirical finding that sets up T-MAC's positioning: the LUT approach is theoretically compelling (eliminating multiplications, linearly scaling with bit-width), but existing hardware platforms have not been able to realize its potential. GPUs fail because their memory hierarchy and execution model are mismatched to the random-access, table-driven computation pattern. This leaves open the question of whether CPUs — with their different cache hierarchy, register flexibility, and instruction set features (like hardware table-lookup instructions) — might be a better fit.

3. The gap: CPU-based LUT for low-bit LLMs

The paper identifies a specific and unexplored gap: "the exploration of LUT-based mixed-precision GEMM/GEMV on CPUs remains uncharted" (Section 2.4). This is the gap T-MAC fills. The motivation is partly counter-intuitive — given GPUs' dominance in ML inference, one might assume they are the natural platform for any compute-intensive method. But the paper argues, through empirical evidence, that CPUs possess architectural features that make them better suited to LUT-based computation than GPUs:

Flexible register file: CPU SIMD registers (NEON 128-bit, AVX2 256-bit) can be repurposed as small lookup tables using hardware table-lookup instructions (TBL on ARM, PSHUFB on x86), providing ultra-low-latency random access that GPUs cannot match because their register files are not architected for indexing.
Larger per-core cache hierarchy: CPU L1/L2 caches can hold larger LUTs closer to execution units than GPU shared memory can (relative to the number of threads competing for it).
Branch prediction and out-of-order execution: CPUs handle the data-dependent access patterns of table lookups (where the weight value determines which table entry is read) more gracefully than GPUs, which suffer warp divergence when threads in a warp access different table entries.
Lower overhead for irregular parallelism: LUT-based computation does not expose the uniform, wide-vector parallelism that GPUs thrive on. CPUs, optimized for latency-sensitive, irregular workloads, may actually be the better execution substrate.

How T-MAC Positions Itself

The paper positions T-MAC not as an incremental optimization over existing dequantization-based CPU kernels, but as a fundamentally different computing paradigm for mixed-precision inference. The shift is from "compute-then-convert" to "convert-then-lookup," which the paper frames through two key transformations:

Transformation 1: Data-type-centric → bit-wise calculation. Instead of treating weights as n-bit integers with specialized packing for each n, T-MAC decomposes any n-bit weight matrix into n one-bit matrices using the linear transformation:

$A \times W = A \times \left(\sum_{i=0}^{n-1} 2^i W_i\right) = \sum_{i=0}^{n-1} 2^i A \times W_i$

where $W_i$ is the i-th bit-plane of the weight matrix. This eliminates the diversity of weight layouts — every weight format, regardless of original bit-width, reduces to a collection of 1-bit matrices with a unified representation. The computation reduces to a series of matrix multiplications between the activation and each 1-bit matrix, with the results bit-serially aggregated.

This is not merely a mathematical reformulation — it is a hardware-software co-design insight: by decomposing the problem into bit planes, T-MAC aligns the computation with the observation that 1-bit patterns have only $2^g$ possible variations for a group of size $g$ , making table lookup a natural implementation strategy.

Transformation 2: Multiplication → table lookup + addition. For a group of $g$ bits from the 1-bit weight matrix, there are only $2^g$ possible bit patterns (e.g., for $g=4$ , the 16 patterns range from [0,0,0,0] to [1,1,1,1]). T-MAC precomputes the dot product of the corresponding $g$ activation values with each possible pattern and stores the results in a lookup table of size $2^g$ . The actual computation then becomes: for each $g$ -bit group in the weight matrix, use the group's bits as an index to look up the precomputed result from the table, then accumulate. This eliminates all multiplications and reduces the operation count to table lookups and additions.

The paper explicitly contrasts this with existing practice in Figure 1: the traditional approach shows a flow of "W(INT n) → Dequant → W(INT8/FP16) → Matmul → O(INT32/FP16)," while T-MAC shows "W(INT n) → decomposed into 1-bit $W_i$ → Table lookup + Sum → O(INT32/FP16)," with the activation feeding into a "Precompute Table" step before the lookup. The elimination of dequantization and multiplication from the critical path is the architectural innovation.

Unified and scalable design. The paper emphasizes that T-MAC provides a unified solution — one kernel design handles any combination of weight bit-width (n) and activation precision — and a scalable one — computation cost scales linearly with weight bit-width (n table lookups for n-bit weights). This addresses the case-by-case kernel explosion problem directly: instead of developing separate kernels for W4A16, W3A16, W2A16, and W1A8, T-MAC uses the same LUT-based approach parametrized by bit-width, group size, and tiling configuration. The diversity is handled through code generation (using TVM + LLVM) with auto-tuning for specific hardware targets, rather than through hand-crafted kernels.

CPU renaissance framing. The paper's subtitle — "CPU Renaissance via Table Lookup" — signals a provocative thesis: that CPUs, long considered second-class citizens for ML inference compared to GPUs and specialized accelerators, may actually be the natural platform for low-bit LLMs when using LUT-based computation. The evidence for this is the paper's empirical demonstration that CPU-based T-MAC not only matches but often exceeds GPU performance on the same edge device (Figure 11, Table 7), with substantially lower power consumption. On the Jetson AGX Orin, T-MAC on CPU outperforms llama.cpp on GPU for W1A16 and W2A16 GEMV kernels, and achieves 2.3× better energy efficiency (Joules/token) for end-to-end inference despite 78% of the GPU throughput, because it consumes only 34% of the power (Table 5).

This is not just a performance claim — it is an architectural argument: CPUs' strengths (flexible register file, sophisticated cache hierarchy, hardware table-lookup instructions, out-of-order execution) align better with LUT-based computation than GPUs' strengths (wide SIMD, coalesced memory access, high throughput on regular parallelism). The paper is arguing that the computing paradigm shift from multiplication to table lookup shifts the hardware fitness landscape in favor of CPUs.

Practical deployment focus. While the paper contributes a novel algorithmic approach, its positioning is fundamentally systems-driven: the goal is to make low-bit LLM deployment on edge devices practical today, using hardware that users already have. The evaluation on four diverse devices (M2-Ultra, Jetson AGX Orin, Surface Book 3, Raspberry Pi 5) spanning three instruction set architectures (ARM NEON, Intel AVX2, Apple Silicon) demonstrates cross-platform viability. The integration into llama.cpp — the most widely used edge LLM framework — ensures practical usability rather than just benchmark performance. And the inclusion of energy measurements alongside throughput acknowledges the real constraints of battery-operated edge devices.

3. Technical Approach

3.1 Reader Orientation

T-MAC is a kernel library for performing mixed-precision matrix multiplication (mpGEMM) on CPUs, where low-bit quantized weights meet higher-precision activations. The core idea is to replace the standard "dequantize-weights-then-multiply" pipeline with a "precompute-lookup-table-then-lookup-and-add" pipeline, which eliminates both the dequantization overhead and the multiplication instructions themselves. The system takes a weight matrix quantized to $n$ bits (where $n$ can be 1, 2, 3, 4, or any bit-width) and an activation matrix in higher precision (e.g., FP16 or INT8), and produces the matrix product through a sequence of table lookups indexed by groups of weight bits. This solves the fundamental problem that existing CPU inference frameworks like llama.cpp cannot realize speedups from aggressive weight quantization because the dequantization cost offsets or reverses the computational savings — T-MAC instead makes computation cost scale linearly with bit-width, achieving the speedup that quantization promises.

3.2 Big-Picture Architecture

The T-MAC system consists of five major components:

Offline Weight Preprocessor — Decomposes any n-bit weight matrix into n one-bit matrices and permutes their layout for sequential memory access. This runs once per model, before inference.
Online Table Constructor — For each [1, g] group of activation values, precomputes the dot product with all $2^g$ possible g-bit patterns (e.g., for g=4, the 16 combinations from [-1,-1,-1,-1] to [+1,+1,+1,+1]) and stores the results in a lookup table.
Table Lookup Engine — Uses hardware-specific instructions (TBL on ARM NEON, PSHUFB on x86 AVX2) to perform register-based table lookups, where g-bit weight groups serve as indices into the precomputed tables.
Bit-Serial Accumulator — Aggregates the looked-up results across all n bit-planes, applying appropriate scaling factors ( $2^i$ for the i-th bit), and accumulates partial results across tiles.
TVM-Based Code Generator — Uses the TVM compiler framework with LLVM backend to generate optimized kernels for specific matrix shapes, bit-widths, and hardware targets, with auto-tuning for tiling parameters.

Information flows as follows: the weight matrix enters the offline preprocessor and is decomposed into permuted one-bit tiles → at inference time, the activation matrix enters the online table constructor, which builds one LUT per group of g activation elements → the table lookup engine iterates over each one-bit weight matrix, using g-bit weight groups as indices to read from the LUTs → the bit-serial accumulator sums the looked-up values within each weight tile and across bit-planes → the final result is scaled by quantization factors to produce the output matrix.

3.3 Roadmap for the Deep Dive

First, the core bit-serial decomposition (Equation 1 and Algorithm 1), because this is the mathematical transformation that makes table lookup possible and defines T-MAC's unified, scalable abstraction.
Second, the table lookup mechanism itself — what the LUT contains, how it is built online from activations, and how weight bits index into it — because this is the computational primitive that replaces multiplication.
Third, the LUT-centric data layout — axis reordering, tiling, weight permutation, and weight interleaving — because naive LUT computation would be bottlenecked by random memory access and on-chip memory pressure.
Fourth, the table storage reduction techniques — mirror consolidation and table quantization — because the LUT size grows exponentially with group size $g$ , and keeping tables in registers is essential for performance.
Fifth, the hardware-specific implementation — intrinsic instructions (TBL/PSHUF, rhadd/avg), register swizzling for LUT precomputation, and the bit-serial linear transformation that maps {0,1} weight bits to {-1,+1} to reduce quantization error — because the abstract algorithm must be mapped to real CPU instructions.
Sixth, the TVM-based code generation and integration approach — because T-MAC must handle diverse matrix shapes (different model dimensions), bit-widths (1-4), and hardware targets (ARM NEON, x86 AVX2, Apple Silicon) without hand-writing each kernel variant.

3.4 Detailed, Sentence-Based Technical Breakdown

This is a systems paper with a novel algorithmic kernel design at its core: T-MAC replaces the standard mixed-precision GEMM computation (dequantize + multiply-accumulate) with a bit-serial decomposition combined with register-based table lookups, and then applies a suite of system-level optimizations (data layout, tiling, memory access patterns, hardware-specific intrinsics) to make this approach fast on real CPU hardware across multiple architectures.

Bit-Serial Decomposition of Mixed-Precision GEMM

The foundational mathematical insight of T-MAC is that mixed-precision matrix multiplication can be decomposed into operations on individual bit-planes of the weight matrix. The paper expresses this through a single core equation:

$A \times W = A \times \left(\sum_{i=0}^{n-1} 2^i W_i\right) = \sum_{i=0}^{n-1} 2^i A \times W_i$

where $A$ is the activation matrix of shape $[N, K]$ (e.g., FP16 or INT8), $W$ is the weight matrix of shape $[M, K]$ quantized to $n$ bits, and $W_i$ is the $i$ -th bit-plane matrix of $W$ — a binary matrix of the same shape $[M, K]$ where each element is the $i$ -th bit of the corresponding weight value (either 0 or 1).

What it computes: The full-precision matrix product $A \times W$ is expressed as the weighted sum of $n$ separate matrix products, each between the same activation $A$ and a single bit-plane $W_i$ , with the result scaled by $2^i$ . The $n$ partial results are then summed together to produce the final output.

Why this form: This decomposition is the critical enabler for table lookup. The original problem — multiplying $A$ (high precision) by $W$ (low precision, packed in arbitrary formats) — requires dequantization and hardware-specific unpacking. After decomposition, each sub-problem is $A$ multiplied by a binary matrix $W_i$ (elements are only 0 or 1). Because $W_i$ is binary, the product $A \times W_i$ has a crucial property: when grouped into chunks of $g$ columns, each group can only take one of $2^g$ possible values. This bounded combinatorial space makes precomputation and table lookup viable — instead of computing $A \times W_i$ through multiply-accumulate, T-MAC precomputes $A$ multiplied by all $2^g$ possible bit patterns, stores the results in a table, and then uses each g-bit group of $W_i$ as an index to select the correct precomputed result. The decomposition transforms an unbounded-computation problem into a bounded-table-lookup problem.

The linear scaling property — computation cost proportional to $n$ , the weight bit-width — is a direct consequence: for n-bit weights, T-MAC performs n serial table lookups (one per bit-plane), regardless of the original packing format. If weights are reduced from 4-bit to 2-bit, the number of bit-planes halves, and computation time halves. This is the property that dequantization-based approaches fail to deliver (Figure 6 shows llama.cpp's 3-bit latency often exceeds 4-bit latency due to unpacking complexity).

The offline weight preprocessing (Algorithm 1, lines 24-28) implements this decomposition in practice. The input is a packed n-bit weight matrix $W$ of shape $[M, K]$ . The preprocessor extracts each bit-plane $W_i$ and packs groups of $g$ consecutive bits into a single index value (0 to $2^g - 1$ ), storing the result as an index matrix of shape $[M, K/g]$ . For example, with $g=4$ , 4 consecutive weight bits are packed into a 4-bit index, and $K$ weight elements along the reduction dimension are compressed into $K/g$ index elements. The preprocessing is done offline (once per model) and the index matrices are stored in the permuted layout (described later) ready for inference.

The online computation (Algorithm 1, lines 6-9 and 14-22) then proceeds as follows for each of the $n$ bit-planes. The table constructor (Precompute function) takes the activation $A$ and for each group of $g$ consecutive activation values along the $K$ dimension, computes the dot product with all $2^g$ possible g-bit patterns. For each pattern index $i \in \{0, \dots, 2^g-1\}$ , it tests each bit $j$ : if the $j$ -th bit of $i$ is set, it adds $A[n, k+j]$ (the corresponding activation element); if the bit is clear, it subtracts $A[n, k+j]$ . This subtraction (rather than adding zero) reflects the linear transformation $f$ that maps weight bits from {0,1} to {-1,+1} (discussed later in the bit-serial linear transformation section), which reduces quantization error and precomputation cost.

The actual mpGEMM then becomes a nested loop over bit-planes $i$ and spatial dimensions $n, m$ , with an inner reduction over $k$ where each $g$ -bit index from the preprocessed weight $W_i$ is used to look up the corresponding entry in the precomputed LUT. The looked-up value — the partial dot product between that activation group and that weight bit-group — is accumulated. After processing all bit-planes, the partial results are scaled by $\alpha_i = 2^i$ (the bit-position weight) and summed with a bias term $\beta$ to recover the full-precision result.

The Lookup Table Mechanism

The lookup table is the computational primitive that replaces multiplication in T-MAC. Understanding its construction, storage, and access pattern is central to understanding how T-MAC achieves its speedup.

What a LUT contains. For a given group size $g$ (the paper uses $g=4$ as the primary configuration), a LUT stores the dot products of one group of $g$ activation values with all $2^g$ possible bit patterns, where each bit pattern is interpreted as a vector of $\pm 1$ values (via the linear transformation mapping 0→-1, 1→+1). For example, with $g=4$ and an activation group $(a_0, a_1, a_2, a_3)$ , the LUT contains 16 entries:

Entry 0 (pattern 0000 → [-1,-1,-1,-1]): $-a_0 - a_1 - a_2 - a_3$
Entry 1 (pattern 0001 → [-1,-1,-1,+1]): $-a_0 - a_1 - a_2 + a_3$
...
Entry 15 (pattern 1111 → [+1,+1,+1,+1]): $+a_0 + a_1 + a_2 + a_3$

Each entry is a single scalar value. The LUT for one activation group of size $g$ has size $[1, 2^g]$ . For $g=4$ , this is 16 elements — small enough to fit in a single 128-bit NEON register (16 × 8-bit values = 128 bits), which is why $g=4$ is the sweet spot: one register holds one LUT, and a single hardware instruction (TBL on ARM, PSHUFB on x86) can perform the lookup.

How LUTs relate to the activation matrix. The activation matrix $A$ of shape $[N, K]$ is partitioned along the $K$ (reduction) dimension into $K/g$ groups. For each group, one LUT is constructed. For a tile of the activation with tile-size $K_{tk}$ along the reduction dimension, the number of LUTs is $K_{tk}/g$ . For example, if $K_{tk}=4$ (the tile along $K$ ) and $g=4$ , there is $K_{tk}/g = 1$ LUT. The total storage for LUTs in a tile is $(K_{tk}/g) \times 2^g$ elements, which is $1 \times 16 = 16$ elements for $K_{tk}=4, g=4$ .

How the lookup works at the instruction level. The preprocessed weight index matrix stores $g$ -bit indices packed into bytes. During computation, a SIMD register is loaded with index values (e.g., 16 uint8 indices for NEON, 32 for AVX2). The LUT is loaded into another SIMD register. A hardware table-lookup instruction — vqtbl1q_u8 on ARM NEON, _mm256_shuffle_epi8 on x86 AVX2 — takes the index register and the table register as inputs, and for each byte position in the index register, reads the byte at that offset from the table register, producing an output register where each byte is the looked-up value. This performs $g$ -wide partial dot products in a single instruction, with no multiplication, no addition — just a byte-permutation in hardware.

Why this is fast. The TBL/PSHUF instructions execute with latency comparable to a single SIMD multiply-accumulate, but they replace multiple operations (load weight, unpack, multiply, accumulate) with a single instruction. More importantly, the dequantization overhead — which in llama.cpp involves shift-and-mask sequences to extract non-byte-aligned bit fields, followed by conversion to the activation precision — is completely eliminated. The weight bits are already in index form from the offline preprocessing; no runtime unpacking is needed beyond a simple byte-aligned load.

The relationship between $g$ and performance. A larger $g$ means each lookup processes more weight bits at once, reducing the number of lookups needed. But the LUT size grows exponentially ( $2^g$ ), and for $g=5$ , the LUT requires 32 entries, which exceeds a single 128-bit NEON register when values are 8-bit (32 × 8 = 256 bits). The paper notes that $g=5$ "requires two registers and slower ARM.TBL2/AVX512.PSHUF" (Section 4, "Determine the size of on-chip LUT"). The choice of $g=4$ is thus a hardware-aligned sweet spot: one LUT = one register = one instruction, maximizing throughput.

LUT-Centric Data Layout

The naive implementation of LUT-based GEMM, where tables are built in memory and accessed through standard loads, would be catastrophically slow. The paper identifies two fundamental challenges and designs a data layout strategy to address both:

Challenge 1: Random access pattern. In standard GEMM, both activation and weight accesses are sequential and predictable, enabling hardware prefetching and high cache-line utilization. In LUT-based GEMM, the weight matrix provides indices into a table — the access pattern into the table is determined by the values of the weight bits themselves, which are effectively random. Sequential memory access is replaced by gather-like access. The solution is to put the LUT in on-chip memory (registers) where random access has single-cycle latency, rather than in cache or DRAM where a random access might cost tens to hundreds of cycles.

Challenge 2: Enlarged on-chip memory pressure. A LUT is larger than the activation group that generated it. For $g=4$ , a group of 4 activation values (4 elements) generates a LUT of 16 elements — a 4× expansion. Additionally, LUT-based computation produces vector outputs (the looked-up values must be accumulated across $K$ ), requiring more registers for intermediate storage than traditional GEMM which produces scalar accumulators. The paper gives a concrete example in Figure 3: "the LUT method uses 144 8-bit registers and llama.cpp uses 104 8-bit registers" for the same logical operation. Without careful tiling, this increased register pressure would cause register spilling to the stack, destroying performance.

The LUT-centric data layout addresses these challenges through four coordinated techniques:

1. Axis reordering (temporal-first loop order). Traditional GEMM loops over the spatial output dimensions ( $N$ and $M$ ) as the outer loops, then reduces along $K$ as the inner loop. This is natural because it maximizes data reuse: a tile of $A$ and a tile of $W$ are loaded, and their product computed for all $N, M$ in the tile before moving to the next $K$ block. However, this ordering is disastrous for LUT-based GEMM because each iteration of the $K$ loop would need to rebuild the LUT for a different chunk of $A$ , or the LUT would need to cover the entire $K$ dimension at once (size $\propto K \times 2^g$ , which is enormous). T-MAC instead reorders the axes to make $K$ the outermost loop: for each $K$ tile, T-MAC builds one set of LUTs, then applies those LUTs across all $M$ (output columns) in the weight tile, reusing each LUT $M_{tm}$ times. This reduces the LUT build cost from $\mathcal{O}(N \times M \times K/g)$ to $\mathcal{O}(K/g)$ per tile.

2. Tiling optimized for LUT reuse. Tiling partitions the matrices into blocks that fit in the CPU's cache hierarchy (L1, L2) and registers. The tile sizes are $[N_{tn}, K_{tk}]$ for the activation $A$ and $[M_{tm}, K_{tk}]$ for the weight $W$ . Unlike traditional GEMM where $N_{tn}$ and $M_{tm}$ have symmetric effects on data reuse, T-MAC's efficiency is strongly asymmetric: a larger $M_{tm}$ increases LUT reuse (the same LUTs are looked up by more weight rows) without increasing the LUT storage, while a larger $N_{tn}$ requires duplicating the LUTs across activation rows. The paper explicitly states: "a larger tile size $M_{tm}$ on $M$ can lead to better lookup table reusing" (Section 3.2). The specific tile sizes used depend on the hardware and are auto-tuned: the default configuration for the example in Figure 3 is $K_{tk}=4, M_{tm}=32$ , meaning 32 weight rows share one LUT over a reduction dimension of 4 activation groups, while llama.cpp uses $K_{tk}=32, M_{tm}=1$ — these are "different tiling rationale[s]" reflecting the fundamentally different data reuse patterns of LUT-based vs. multiply-based computation.

3. Weight permutation for sequential memory access. Even with tiling, loading tile-sized chunks of the weight matrix from DRAM involves non-sequential access because the tile is a logical $[M_{tm}, K_{tk}]$ submatrix that is not contiguously stored in the original row-major layout. To ensure high DRAM bandwidth utilization, T-MAC performs an offline permutation of the weight matrix: it flattens each tile into a contiguous segment in memory, then concatenates these flattened tiles in the order they will be accessed during inference. This ensures that weight loads are always sequential, maximizing cache-line utilization and enabling hardware prefetching. Because the weight matrix is static during LLM inference (it does not change between tokens), this permutation is performed once offline and costs nothing at runtime.

4. Weight interleaving for fast unpacking. The preprocessed weight matrices store $g$ -bit indices packed into bytes. On little-endian CPUs (which includes all modern ARM and x86 processors), multi-byte integers are stored with the least significant byte at the lowest address. If $g=4$ bit indices from consecutive weight positions are packed across byte boundaries, unpacking them for SIMD lookup requires byte-reversal operations to restore the correct order. T-MAC avoids this runtime cost by interleaving the packed indices during offline preprocessing: the bits are rearranged so that when the packed bytes are loaded into SIMD registers using normal (little-endian) loads, the resulting register contents are already in the correct order for the table lookup instruction. Figure 4 illustrates this with an example: without interleaving, unpacking bytes [0:3], [4:7], [60:63] from the packed format would require AND and SHR+AND operations to extract and reorder the nibbles. With interleaving, the bytes are pre-rearranged so a simple load produces the indices in the expected order.

Reducing LUT Storage

The LUT storage grows exponentially with group size $g$ and linearly with the number of activation groups per tile ( $K_{tk}/g$ ). For fixed $g=4$ , the LUT is 4× larger than the activation group that generated it. To keep LUTs in registers (the fastest on-chip memory), T-MAC applies two complementary compression techniques: mirror consolidation (halving the number of entries) and table quantization (reducing the bit-width per entry).

Mirror Consolidation. The LUT values for a group of $g$ activation elements exhibit an inherent symmetry: for every bit pattern, there is a complementary pattern (all bits flipped) whose LUT value is the negative of the first. For example, with $g=4$ and activation group $(a_0, a_1, a_2, a_3)$ :

Pattern 0000 produces $-a_0 - a_1 - a_2 - a_3$
Pattern 1111 (bitwise complement of 0000) produces $+a_0 + a_1 + a_2 + a_3$

These two values are exact negatives of each other. This holds for all $2^g$ patterns because the dot product of the activation group with a bit-pattern vector $\mathbf{b} \in \{-1,+1\}^g$ and its complement $-\mathbf{b}$ are negatives. Therefore, only $2^{g-1}$ entries need to be stored (e.g., 8 entries for $g=4$ instead of 16), and the remaining entries can be reconstructed by negation when accessed. The paper describes this as "lossless" and notes it "accelerating the precomputation of the lookup table, reducing the required storage, and speeding up table accesses" (Section 3.3). The runtime cost is that the table lookup must handle the sign: if the original index (weight bit-group) is in the upper half of the table (index $\geq 2^{g-1}$ ), the looked-up value from the stored half-table is negated. The paper appears to implement this by using the most significant bit of the index as a sign flag.

Table Quantization. Even after mirror consolidation, LUT entries are stored at the activation's precision (e.g., FP16, 2 bytes per entry). Table quantization reduces each entry to a lower-precision integer (e.g., INT8, 1 byte per entry) with a scaling factor, analogous to weight quantization but applied to the LUT values themselves. The paper makes a critical design choice that distinguishes this from conventional activation quantization: T-MAC uses fine granularity and dynamic quantization. Specifically, the paper states "quantizing 8 values for k=4" and "dynamic quantization" (Section 3.3), meaning that quantization parameters (scale, zero-point) are computed per-LUT (per group of $g$ activation values) at runtime, rather than using static, globally calibrated quantization ranges.

Why this choice matters for accuracy. Conventional activation quantization for LLM inference is challenging because activations have outliers (a few channels with very large magnitudes), requiring either coarse per-tensor quantization (which loses precision on the majority of in-distribution values) or complex outlier-handling schemes. Table quantization does not face this problem because LUT values are dot products of $g$ activation elements with $\pm 1$ vectors — each LUT entry is a sum or difference of $g$ activation values, which has much smaller variance than individual activation values. Furthermore, using per-LUT dynamic quantization (one scale factor per table, computed on-the-fly) eliminates the need for calibration data or offline range estimation. The paper reports that "table quantization technique has an imperceptible effect on the overall model accuracy" (Section 3.3), supported by the kernel-level NMSE results in Table 3 (Section 5.6, discussed in the experimental analysis) showing negligible difference between llama.cpp and T-MAC error relative to un-quantized ground truth.

Combined effect. With mirror consolidation (2× reduction in entries) and table quantization (e.g., 2× reduction in bytes per entry from FP16 to INT8), the total LUT storage is reduced by up to 4×, bringing it closer to the size of the original activation group. This is crucial for fitting LUTs into registers and enabling the on-chip table lookup that makes T-MAC fast.

Hardware-Specific Implementation

The abstract LUT algorithm must be mapped to real CPU instructions for each target architecture. The paper targets three instruction set architectures: ARM NEON (128-bit SIMD, used on Apple Silicon, Raspberry Pi, Jetson), Intel AVX2 (256-bit SIMD, used on x86 laptops/desktops), and Apple AMX (a proprietary matrix coprocessor on Apple Silicon). The implementation focuses on three categories of operations: table lookup, fast aggregation, and LUT precomputation.

Table lookup instructions. The critical instruction for T-MAC is the byte-permutation instruction that implements register-based table lookup. On ARM NEON, this is vqtbl1q_u8: given a table register (16 uint8 values) and an index register (16 uint8 indices), it produces an output register where each byte is the table entry at the index-specified position. For indices exceeding the table size (15 for a 16-entry table), the instruction returns 0, which the paper notes would produce incorrect results if not handled. On Intel AVX2, the equivalent is _mm256_shuffle_epi8 (often called PSHUFB), but with a critical architectural difference: the 256-bit register is treated as two independent 128-bit lanes, and a shuffle within one lane cannot access bytes from the other lane. The paper works around this by duplicating the 16-entry LUT into both 128-bit halves of the 256-bit register, so each lane has its own complete copy, enabling 32 independent lookups in a single instruction (16 per lane). For FP16 LUT values (2 bytes per entry), since neither NEON nor AVX2 supports 16-bit table lookup, the paper splits each FP16 value into its low and high bytes, builds two separate 8-bit LUTs, performs two table lookups, and recombines the bytes into FP16 values.

Fast 8-bit aggregation. After table lookup, the looked-up values must be accumulated along the $K$ reduction dimension. The natural approach is to accumulate in INT16 or INT32 to avoid overflow (since summing many INT8 values can exceed the INT8 range), but INT16 instructions have half the throughput of INT8 on most SIMD architectures. The paper adopts the fast 8-bit aggregation technique from MADDNESS (Blalock and Guttag, 2021): accumulate in INT8 using averaging instructions (vrhaddq_u8 on NEON, _mm256_avg_epu8 on AVX2) that compute the rounded average of two INT8 vectors without overflow, then correct for the averaging bias in the final accumulation step. The paper acknowledges this introduces non-negligible accuracy loss and treats it as an optional optimization (disabled by default, offered as "T-MAC (+FA)" in the evaluation). The main T-MAC configuration first aggregates in low-bit (INT8) to reduce intermediate storage, then converts the accumulated sum to higher precision (FP16) for the final result.

Register swizzling for efficient LUT precomputation. Building the LUT online from the activation vector requires computing $2^g$ dot products efficiently. For $g=4$ , this means computing 16 values, each being a sum/difference of the 4 activation elements. The naive approach — loop over 16 patterns, compute sum for each — is slow because it accesses the 4 activation elements non-contiguously. The paper uses SIMD gather instructions to load the non-contiguous activation elements: LD4 on NEON loads 4 interleaved values, and vgatherdps on AVX2 gathers from arbitrary memory positions. However, after computing the 16 LUT entries, writing them back to memory is problematic on AVX2 because the entries for different patterns reside in different SIMD register lanes and writing them contiguously requires byte-level permutation. The paper describes a register swizzling sequence: vpblendvb blends 8-bit values from different registers into one register, vpermd permutes 32-bit dwords within a 256-bit register, and vpshufb shuffles 8-bit bytes into the correct order. After this multi-instruction swizzle, the LUT can be written to memory as a contiguous 16-byte array, ready for subsequent table lookup.

Bit-serial linear transformation. The raw weight bits are 0s and 1s, but using {0,1} directly in the LUT construction has two problems. First, it creates asymmetric LUT values (entries range from 0 to the sum of positive activations), which wastes dynamic range compared to a zero-centered LUT. Second, during the precomputation (Algorithm 1, lines 16-22), the paper uses subtraction for 0-bits and addition for 1-bits to build the LUT — this corresponds to mapping 0→-1 and 1→+1. The paper frames this as a general linear transformation $f(v_i) = \alpha'_i v_i + \beta'_i$ , where $f(0) = s_0$ and $f(1) = s_1$ are the transformed values. The optimal choices are empirically determined: $s_0 = -1, s_1 = 1$ from the candidate set {-1, 0, 1}. Setting $s_0 = -1, s_1 = 1$ has two advantages: it eliminates float-multiply instructions (since the LUT entries are pure sums/differences, requiring only addition and subtraction), and it minimizes the difference between the largest and smallest LUT entries (the range is centered around zero), which "reduce[s] quantization error" when table quantization is applied (Section 4).

With this transformation, the decomposition of $W$ is adjusted:

$W = \sum_{i=0}^{b-1} \alpha_i 2^i W'_i + B$

where $W'_i = f(W_i)$ are the transformed bit matrices (with values -1 and +1), $\alpha_i = 1 / \alpha'_i$ is the inverse scaling from the linear transformation, and $B = J \cdot \sum_{i=0}^{b-1} \beta_i 2^i$ is a bias matrix (where $J$ is an all-ones matrix). In practice, with $s_0=-1, s_1=1$ , the bias $B$ simplifies and the scaling factors are absorbed into the bit-position weights.

Why this transformation matters for the implementation. When building the LUT online (Algorithm 1, Precompute function), the inner loop tests each bit position $j$ of the pattern index $i$ : if the bit is set ( $v_j = 1$ ), it adds the activation element; if the bit is clear ( $v_j = 0$ ), it subtracts the activation element. This corresponds exactly to $f(0) = -1, f(1) = +1$ , and requires only addition and subtraction — no multiplication by zero or one. The LUT values are sums of $\pm$ activations, producing both positive and negative values centered around zero, which is better for quantization (symmetric range, no wasted representational capacity on consistently positive values).

Code Generation and Integration

T-MAC is not a single hand-written kernel but a kernel generation system that produces optimized code for each combination of matrix shape, bit-width, and hardware target. This addresses the diversity problem at the system level rather than the algorithm level: instead of writing one kernel per configuration, T-MAC describes the computation at a high level and relies on a compiler stack to produce the optimized machine code.

TVM + LLVM code generation. The paper leverages TVM (Chen et al., 2018), an end-to-end deep learning compiler, in combination with LLVM (Lattner and Adve, 2004) as the backend code generator. TVM provides a high-level tensor computation description (expressing the GEMM as loops over tiles with known sizes) and applies standard compiler optimizations: loop unrolling (expanding loop bodies to reduce branch overhead), vectorization (mapping inner loops to SIMD instructions), and constant folding (precomputing tiling-dependent constants). The paper uses TVM's Tensorize mechanism to "embed hardware intrinsics into the code" (Section 4) — essentially annotating specific inner loops as targets for replacement with hand-tuned intrinsic sequences (TBL instructions, swizzling sequences, etc.). This allows T-MAC to express the high-level tiling and loop structure in TVM's domain-specific language while injecting architecture-specific intrinsics at the leaf level.

Auto-tuning with AutoTVM. The tiling parameters ( $N_{tn}, M_{tm}, K_{tk}$ ), the number of on-chip LUTs, and other micro-architectural decisions are hardware-specific. The paper uses AutoTVM (Chen et al., 2018), TVM's automatic tuning framework, to search over these parameters for each hardware target. The search objective is latency on representative matrix shapes. The paper notes that "Tuning does not appear to be very effective" on M2-Ultra in the optimization breakdown (Figure 10) because "the default tiling configurations already align well with M2-Ultra registers and caches," but that for different devices, tuning should help find a better configuration. This suggests the auto-tuning is more important for cross-platform portability (adapting to different cache sizes, SIMD widths, instruction latencies) than for squeezing out the last few percent on a well-matched platform.

Threading model and llama.cpp integration. T-MAC generates C++ code (not library binaries) for each kernel variant to avoid runtime dependencies on the TVM runtime, particularly its threadpool. The paper identifies a critical integration issue: "when integrating T-MAC into llama.cpp, we notice an obvious conflict between the llama.cpp threadpool and the TVM threadpool" — threads from different pools compete for CPU resources, causing "significant performance degradation" (Section 4). The solution is to generate standalone C++ functions that compute a single threadblock's worth of work, then integrate these functions into llama.cpp's native threadpool. Each generated function handles one tile of the GEMM (one combination of $N$ and $M$ tile indices for a given $K$ tile), and the llama.cpp threadpool distributes these threadblocks across available cores. This achieves both performance (no threadpool contention) and compatibility (T-MAC becomes a drop-in replacement for llama.cpp's mpGEMM/GEMV kernels).

API design. T-MAC exposes its kernels through two interfaces: (1) TVM PackedFunc for integration with Python frameworks (PyTorch, NumPy) via the DLPack tensor interchange format, and (2) a lightweight C++ wrapper using raw pointers with no TVM runtime dependency. The C++ interface is the primary path for llama.cpp integration, while the Python interface supports experimentation and benchmarking from frameworks like PyTorch. The paper emphasizes cross-platform portability: the generated C++ code compiles on all target operating systems (macOS, Linux, Windows) with standard C++ compilers.

Determining hardware-specific parameters. The paper describes several parameters that are tuned per-hardware rather than derived analytically:

Number of on-chip LUTs (Section 4, "Determine the size of on-chip LUT"): T-MAC can hold multiple LUTs in registers simultaneously, allowing accumulation across multiple $K$ -tile groups before writing back intermediate results to memory. The number is tuned "for each hardware, to make sure on-chip memory can be fully utilized and LUTs won't be swapped out for the tile." If too many LUTs are allocated, register spilling occurs; if too few, write-back overhead increases. The paper notes this as a key tuning parameter but does not enumerate the specific values used per device.
Group size $g$ : Set to $g=4$ as the default because the $2^4=16$ -entry LUT "exactly fits into one register for ARM.TBL/AVX2.PSHUF" (Section 4). The paper mentions $g=5$ as a larger alternative that requires two registers and the slower TBL2 (ARM) or AVX512.PSHUF (x86) instructions, implying a trade-off between larger group size (fewer lookups) and slower lookup instructions. The evaluation uses $g=4$ throughout.
Tiling configuration: The tile sizes $[N_{tn}, K_{tk}]$ and $[M_{tm}, K_{tk}]$ are auto-tuned per shape and hardware. Figure 3 shows the specific configuration $(K_{tk}, M_{tm}) = (4, 32)$ for the T-MAC example, contrasted with llama.cpp's $(32, 1)$ . The larger $M_{tm}$ for T-MAC reflects the LUT reuse benefit: the same LUT is reused across 32 rows of the weight tile, amortizing the LUT build cost.
Single-thread vs. multi-thread behavior: For single-threaded GEMV (the dominant operation in LLM decode with batch size 1), T-MAC's performance is compute-bound — the LUT lookups and aggregations determine latency. For multi-threaded GEMV where multiple cores share memory bandwidth, T-MAC becomes memory-bound, but the paper claims "T-MAC can still achieve significant speedup due to efficient memory access" (Section 5.2) — the weight permutation ensuring sequential DRAM loads and the omission of dequantization reduce the memory bandwidth consumed per effective operation.

Summary of Design Choices and Their Justifications

Bit-serial decomposition over data-type-specific kernels: provides a unified, scalable abstraction that handles any weight bit-width (1-bit through 4-bit and beyond) with one code path, and ensures computation cost scales linearly with bit-width — the property that llama.cpp's dequantization-based approach fails to deliver.
LUT with $g=4$ over larger groups: balances lookup throughput (one instruction per lookup) against LUT size (fits in one 128-bit NEON register), exploiting the fact that hardware table-lookup instructions are byte-granularity shuffles that naturally map to $2^4=16$ -entry tables.
Register-based LUT storage over cache or shared memory: eliminates the random-access latency penalty of table lookups by placing the table in the fastest memory level (registers) where access is single-cycle; avoids the GPU-style shared memory bottleneck that LUT-GEMM experiences.
Temporal-first (K-outer) loop order over spatial-first: enables LUT reuse across $M$ weight rows (amortizing the LUT build cost) and keeps the per-tile LUT storage small ( $K_{tk}/g$ tables rather than $N \times K_{tk}/g$ ).
Mirror consolidation over full-table storage: losslessly halves the LUT size by exploiting the $\pm$ symmetry of dot products with binary vectors, with minimal runtime cost (a conditional negation based on the index's MSB).
Fine-grained dynamic table quantization over static coarse quantization: exploits the bounded variance of LUT entries (sums of $g$ activation values) to quantize to INT8 with negligible accuracy loss, while avoiding the calibration complexity and model-sensitivity of conventional activation quantization.
Offline weight permutation and interleaving over runtime unpacking: shifts the cost of making weight access patterns efficient to a one-time preprocessing step, ensuring that the inference-time memory access is sequential (maximizing DRAM bandwidth) and the unpacking is byte-aligned (avoiding shift-and-mask sequences).
TVM code generation with Tensorize over hand-written assembly: handles the combinatorial explosion of shapes, bit-widths, and hardware targets through compiler automation rather than manual kernel development, while still allowing platform-specific intrinsics to be injected at the inner-loop level where they matter most.
Integration at llama.cpp threadblock granularity over separate runtime: avoids threadpool contention and enables drop-in replacement of llama.cpp's existing mpGEMM/GEMV kernels without changing the model loading, KV-cache management, or higher-level inference orchestration.

4. Key Insights and Innovations

Innovation 1: CPUs as the Natural Platform for LUT-Based Low-Bit Inference — A Hardware-Computation Co-Design Reframing

The paper's most intellectually distinctive contribution is not any single algorithmic technique, but rather the counter-intuitive reframing of CPUs as the preferred execution substrate for LUT-based mixed-precision GEMM — directly contradicting the prevailing assumption that GPUs, with their massively parallel multiply-accumulate throughput, are the natural home for any matrix computation. The dominant mental model in the edge inference community has been that low-bit LLMs need GPU or NPU acceleration, and that CPU-based inference is a fallback option at best. T-MAC's contribution is to demonstrate, through a combination of algorithmic design and careful empirical analysis, that this assumption is an artifact of the dequantization-based computation paradigm, not an inherent property of the hardware.

The conceptual move is to recognize that LUT-based computation has fundamentally different hardware affinity than multiply-accumulate computation. GPUs excel at regular, wide-SIMD operations on coalesced memory — the exact pattern that dequantization-based GEMM exhibits. LUT-based computation, by contrast, exhibits data-dependent access patterns (the weight bits determine which table entry to read), irregular parallelism (different weight groups index different table entries), and benefits from low-latency random access to small tables. These are precisely the characteristics that CPU architectures have been optimized for over decades: out-of-order execution to hide data-dependent latency, large per-core register files that can hold small tables, sophisticated cache hierarchies with low-latency L1 access, and branch prediction that handles control flow variation. The paper does not merely observe that CPUs can run LUT-based kernels — it argues that the LUT-based paradigm shifts the hardware fitness landscape toward CPUs and away from GPUs.

Prior work in this space had already attempted LUT-based computation on GPUs. LUT-GEMM (Park et al., 2023) and related efforts (Maleki, 2023) demonstrated the theoretical reduction in computational complexity but suffered from poor practical performance — the paper reports that GPU LUT-GEMM kernels were 1.75–2.34× slower than dequantization-based kernels on A100 hardware (Section 2.4). The authors attribute this to GPUs "offering either inadequate storage capacity for the lookup tables or insufficiently rapid table access" — GPU shared memory must be shared across hundreds of threads, limiting per-thread LUT capacity, and the SIMT execution model suffers when threads in a warp diverge due to different table indices. The GPU results served as negative evidence that LUT-based computation, while algorithmically elegant, might not be practically viable. T-MAC flips this narrative by showing that the issue was never the LUT approach itself, but the platform it was mapped to. The title's phrase "CPU Renaissance" signals this intellectual move explicitly: CPUs, long considered legacy hardware for ML inference, become the superior choice when the computational paradigm changes.

The empirical evidence for this reframing is multifaceted:

Kernel-level CPU-vs-GPU comparison (Figure 11): On Jetson AGX Orin, T-MAC on CPU outperforms llama.cpp on GPU for W1A16 mpGEMV across all matrix shapes, and matches GPU for W2A16 and W3A16. The CPU beats the GPU at its own game — not through better peak FLOPs, but through a computation strategy that exploits CPU strengths.
End-to-end energy efficiency (Table 5): For Llama-2-7B-2bit on Orin, T-MAC on CPU achieves 2.3× better energy efficiency (Joules/token) than llama.cpp on GPU, despite delivering 78% of the GPU's throughput, because CPU power consumption is only 34% of the GPU's. This reframes the deployment question from "how much throughput can I get?" to "how much throughput per Watt can I get?" — the metric that matters for battery-operated edge devices.
Cross-platform performance parity with specialized hardware (Table 7): T-MAC on CPU outperforms both GPUs and NPUs on multiple platforms — 3× faster than NPU on Surface Laptop 7 for Llama-2-7B-2bit, 1.5× faster than NPU on OnePlus 12, and 1.4× faster than Ampere GPU on Jetson Orin NX. The NPUs, specifically designed for AI inference, are outclassed by general-purpose CPUs running the right algorithm.

This is not merely a performance result — it is an architectural argument with implications for hardware design. The paper explicitly states in its conclusion that T-MAC "opens up the broad opportunity for novel LLM hardware accelerator design based on LUT, as LUT is much more efficient in hardware implementation than multiplications" (Section 7). If LUT-based computation proves to be the dominant paradigm for mixed-precision inference, future hardware should optimize for table lookup throughput, not multiply-accumulate throughput. This inverts the trajectory of ML hardware design from the past decade, which has been toward ever-larger matrix multiplication units.

Innovation 2: The Verifier Over-Optimization Diagnostic — Why Dequantization Breaks Bit-Width Scaling

T-MAC identifies, names, and empirically demonstrates a specific failure mode in existing inference systems that the paper implicitly terms dequantization-induced scaling pathology: the phenomenon where reducing weight bit-width increases inference latency because the dequantization overhead grows faster than the arithmetic savings shrink. This is not a marginal inefficiency — it is a structural inversion of the expected scaling relationship, and the paper's clear demonstration of it is a significant diagnostic contribution that explains why the field has not realized the speedup potential of aggressive weight quantization on CPUs.

Prior work on low-bit LLM inference had focused on two axes: (a) quantization algorithms that preserve model accuracy at lower bit-widths (GPTQ, AWQ, BitDistiller, OneBit), and (b) inference system optimizations that accelerate the dequantization step (llama.cpp's hand-tuned unpack kernels, BitBLAS, Marlin). The implicit assumption was that better quantization algorithms would produce better model quality at lower bits, and better dequantization kernels would reduce the overhead enough to expose the computational savings. What T-MAC's analysis reveals is that this assumption is false in a fundamental sense: the dequantization overhead is not a fixed tax but a function of the bit-width, and for certain widths (particularly 3-bit, where 8 is not divisible by 3) the overhead can actually increase as bit-width decreases.

The paper's Figure 6 provides the smoking gun: across four different devices (M2-Ultra, Raspberry Pi 5, Jetson AGX Orin, Surface Book 3), llama.cpp's single-threaded mpGEMV latency at 3-bit is consistently higher than at 4-bit for all matrix shapes tested. The pathology is most visible on the Raspberry Pi 5, where the 3-bit bar is visibly taller than the 4-bit bar across all six matrix shapes. This violates the intuitive "fewer bits = faster" expectation and demonstrates that the dequantization approach has a non-monotonic relationship with bit-width — the computation does not scale down linearly with precision reduction.

The paper's explanation for this is specific and diagnostic: weight decoding for non-power-of-two bit-widths requires complex shift-and-mask sequences because the packed representation doesn't align with byte boundaries. "Since 8 is indivisible by 3, this decoding process is notably inefficient. llama.cpp attempts to optimize it by separately packing 2 bits and the remaining 1 bit, but it still results in significant overhead" (Section 5.2). The consequence is that practitioners who quantize from 4-bit to 3-bit expecting a 25% speedup may actually see a slowdown, making 3-bit quantization actively harmful for latency despite reducing memory usage. The 2-bit case also fails to deliver meaningful speedup over 4-bit — the dequantization overhead consumes most of the savings.

T-MAC's solution to this pathology — bit-serial decomposition where each bit-plane is processed independently and uniformly — eliminates the non-monotonicity entirely. Computation cost becomes strictly linear in bit-width: 2-bit weights require 2 serial lookups, 4-bit weights require 4, and scaling is predictable. The paper's Figure 6 shows this linear scaling materialized: T-MAC's latency bars decrease proportionally from 4-bit to 1-bit, achieving the speedup that quantization promises but dequantization-based systems fail to deliver.

This diagnostic contribution matters beyond T-MAC itself. It establishes a design principle for mixed-precision inference systems: the representation format and computation mechanism must be co-designed so that precision reduction translates to computation reduction without an offsetting decoding tax. Any system that separates the representation (packed low-bit weights) from the computation (high-precision multiply-accumulate via dequantization) will encounter this pathology at some bit-width, and the specific widths where it bites depend on the interaction between the packing scheme and the hardware's native word sizes. The insight generalizes: avoiding non-monotonic performance scaling requires computation in the native precision of the representation, not conversion to a higher precision first.

Innovation 3: LUT-Centric Data Layout as a New Optimization Dimension — Breaking the Traditional GEMM Tiling Orthodoxy

Traditional GEMM optimization, whether for CPUs (BLAS libraries) or GPUs (CUTLASS, Triton), follows a well-established tiling orthodoxy: tiles are chosen to maximize data reuse by keeping sub-blocks of both input matrices in cache, with roughly symmetric treatment of the output dimensions ( $N$ and $M$ ) since both contribute equally to data reuse. T-MAC introduces a genuinely novel optimization dimension by recognizing that LUT-based GEMM has fundamentally asymmetric data reuse patterns, and that the tiling strategy must be redesigned around the LUT's lifecycle (build once, reuse many times) rather than the traditional input-output data flow.

The conceptual break is this: in traditional GEMM, every element of the weight matrix is used exactly once per tile (loaded, multiplied, accumulated, and discarded), so tiling is about maximizing how many computations each loaded element participates in before it is evicted from cache. In LUT-based GEMM, the weight matrix is not multiplied — it provides indices into a table. The table itself is the "hot" data structure that must be reused as much as possible before being rebuilt. The consequence is a reversal of which tiling axes matter for data reuse.

T-MAC's key insight is that the $M$ dimension (output columns, corresponding to weight rows) is now the primary reuse axis, not a symmetric partner to $N$ . A larger $M_{tm}$ means the same LUT (built once from a group of activation values) is indexed by more weight rows before it needs to be rebuilt. The paper's Figure 3 makes this concrete: T-MAC's default tile configuration is $(K_{tk}, M_{tm}) = (4, 32)$ , meaning the LUT is reused 32 times across different weight rows. In contrast, traditional GEMM for llama.cpp uses $(K_{tk}, M_{tm}) = (32, 1)$ — the exact opposite emphasis. This is not a minor parameter tweak; it reflects a fundamental rethinking of what "data reuse" means when the key data structure is a dynamically computed table rather than a statically loaded matrix.

The axis reordering from spatial-first to temporal-first ( $K$ -outer) is a direct consequence of this asymmetry. If the $N$ and $M$ loops were outer (as in traditional GEMM), the system would need to either rebuild the LUT for every combination of $N, M$ tiles (catastrophically expensive) or maintain a 3D LUT spanning all $N$ and $K$ groups (impossibly large). By making $K$ outermost, T-MAC builds one set of LUTs per $K$ -tile and reuses them across all $M$ tiles, effectively amortizing the LUT construction cost over the entire output column dimension. The paper states this concisely: "a larger tile size $M_{tm}$ on $M$ can lead to better lookup table reusing" (Section 3.2).

Prior work on LUT-based computation (DeepGEMM, LUT-GEMM, MADDNESS) had not systematically analyzed how tiling interacts with LUT lifecycle. These systems applied conventional tiling strategies and suffered from the resulting inefficiency — particularly on GPUs, where the mismatch between LUT reuse patterns and thread-block scheduling led to the poor performance that T-MAC documents. The LUT-centric data layout is not merely an optimization of existing LUT methods; it is a new category of optimization that emerges from recognizing that LUTs are a first-class data structure with their own production and consumption economics, distinct from the static weight matrices of traditional GEMM.

The weight permutation technique (offline reordering of weight tiles into sequential memory) and weight interleaving (pre-rearranging packed indices to eliminate runtime byte-reversal) are concrete manifestations of this LUT-centric thinking. Both are layout transformations that would make no sense in a traditional GEMM — in traditional GEMM, weights are just numerical values to be loaded and multiplied, and their order in memory doesn't change the computation in a way that interleaving would help. In LUT-based GEMM, weights are indices, and the byte-level arrangement of those indices directly determines how many runtime instructions are needed to prepare them for the table lookup. The optimization is only meaningful because the paper has reframed weights from "numbers to multiply" to "indices for table lookup."

Innovation 4: Bit-Serial Decomposition as a Unifying Abstraction for Mixed-Precision Computation

While bit-serial arithmetic is well-established in digital logic design (where serial adders trade throughput for gate count), its application as a unifying software abstraction for mixed-precision matrix multiplication on general-purpose processors is a novel contribution. The paper's core equation — decomposing $A \times W$ into $\sum 2^i A \times W_i$ — appears mathematically trivial, but its significance lies in what it eliminates from the system design space: the need for case-by-case kernel implementations, bit-width-specific packing formats, and dequantization logic for every {weight bits, activation precision} combination.

The dominant approach in the field has been to design inference kernels around hardware data types: INT8 dot-product instructions for W8A8, FP16 multiply-accumulate for W16A16, and so on. When weight precision doesn't match hardware supported types, the kernel must bridge the gap through unpack-and-convert sequences that are specific to the source format. A W4A16 kernel differs from a W3A16 kernel not in the mathematical operation but in the unpacking logic: W4 uses nibble extraction, W3 uses a complex 2+1 bit split, W2 uses bit-pair extraction. Each kernel is a separate engineering effort with its own performance characteristics, edge cases, and maintenance burden.

T-MAC's bit-serial decomposition collapses this diversity into a single parameter: the number of bit-planes $n$ . Whether the original weight is 4-bit, 3-bit, 2-bit, or 1-bit (or even 1.58-bit ternary for BitNet), the preprocessing step reduces it to a collection of 1-bit matrices, and the runtime performs one table lookup pass per bit-plane. The paper explicitly demonstrates this unification: "Ternary weights in 1.58bit BitNet are interpreted as 2-bit and decomposed into two 1-bit matrices" (Section 5.1). The same kernel code handles all cases; only the number of iterations over bit-planes changes.

This is more than a software engineering convenience. It represents a separation of concerns that the dequantization paradigm conflates. In T-MAC, the representation of the quantized weights (the offline preprocessing that packs indices) is decoupled from the computation (the online table lookup and accumulation). Changing the quantization scheme (e.g., from uniform 4-bit to non-uniform 3-bit with a different codebook) requires only modifying the preprocessing step; the runtime kernel is unchanged. In dequantization-based systems, changing the quantization scheme requires rewriting the unpacking kernel because the unpacking logic is tightly coupled to the bit layout. This separation makes T-MAC robust to the rapid evolution of quantization algorithms — as new schemes emerge (2-bit QAT, 1.58-bit ternary, mixed-precision where different layers have different bit-widths), T-MAC can support them without kernel modifications.

The linear scaling property is the direct consequence that makes this abstraction practically valuable: computation cost is $\mathcal{O}(n)$ for $n$ -bit weights. This is the property that dequantization-based approaches fail to deliver (as Innovation 2 details), and it is what makes the abstraction not just elegant but predictively useful — a practitioner can estimate inference latency for a new bit-width by linear interpolation from known bit-widths, something that is impossible with dequantization-based kernels where the relationship is non-monotonic. The paper's results in Figure 6 demonstrate this linear scaling concretely: T-MAC's bars form a descending staircase pattern from 4-bit to 1-bit across all devices and matrix shapes.

This innovation connects conceptually to the broader trend in computer architecture toward bit-serial and bit-parallel programmable computation as an alternative to fixed-width arithmetic units. Just as bit-serial processors (e.g., early minicomputers) traded bit-width flexibility against per-operation throughput, T-MAC trades the per-operation throughput of hardware multiply-accumulate against the flexibility and reduced total-operation-count of bit-serial LUT. The paper's conclusion that LUT-based computation is "much more efficient in hardware implementation than multiplications" suggests that this trade-off could become the basis for future accelerator designs.

Innovation 5: Register-Resident LUT as a New Computational Primitive — Beyond Cache-Blocked GEMM

The paper introduces a genuinely novel use of CPU SIMD registers: not as vectors of operands for arithmetic instructions, but as small, dynamically indexed lookup tables that replace arithmetic entirely. This is a conceptual shift from "registers hold data to be operated on" to "registers hold precomputed results that are selected by data-dependent indices." While the hardware instructions that implement this — TBL on ARM, PSHUFB on x86 — have existed for decades (originally designed for cryptographic table lookups and byte permuting), their repurposing as the primary computational primitive for matrix multiplication is a creative systems innovation.

The standard pattern for CPU SIMD optimization is to load vectors of activations and weights into registers, perform multiply-accumulate instructions on corresponding lanes, and accumulate the results. The number of operations scales with the reduction dimension $K$ . In T-MAC's LUT-based approach, the operations scale with $K/g$ (the number of LUT lookups needed) rather than $K$ , because each lookup processes $g$ weight bits at once. For $g=4$ , this is a 4× reduction in instruction count along the reduction dimension. The LUT construction (which scales with $2^g$ ) is amortized over $M$ weight rows, so for large enough tiles, the LUT build cost is negligible compared to the saved multiply-accumulate instructions.

What makes this genuinely novel as a primitive is not the existence of the TBL/PSHUF instructions, but the system-level design that makes register-resident LUTs viable for GEMM at scale. Prior use of table lookup in neural network inference (MADDNESS, LUT-GEMM) stored tables in cache or shared memory, where lookups incurred tens of cycles of latency and competed for capacity with other data. T-MAC's key realization is that for small group sizes ( $g=4$ ), the LUT is small enough (16 entries) to fit in a single register, and the lookup instruction has latency comparable to a SIMD multiply-accumulate (single-digit cycles). The register file becomes the LUT storage; the SIMD lanes become independent lookup engines; the weight bits become indices, not multiplicands. The entire multiply-accumulate stage of the GEMM pipeline is replaced by a table-lookup stage.

The challenge this creates — and that the paper's optimizations address — is register pressure. A LUT is larger than the activation group that generated it (4 activation values → 16 LUT entries for $g=4$ ), and the intermediate accumulation requires additional registers. The paper's techniques of mirror consolidation (halving the LUT size), table quantization (reducing entry bit-width from FP16 to INT8), and LUT-centric tiling (reducing the number of simultaneous LUTs through $K$ -outer loop ordering) are all in service of fitting the LUT working set into registers without spilling to cache. The paper reports that "the LUT method uses 144 8-bit registers and llama.cpp uses 104 8-bit registers" (Section 3.1, Figure 3) — a 38% increase that would cause performance collapse on a register-starved architecture without careful management.

This primitive has implications beyond T-MAC. It establishes that table lookup can be a first-class compute operation for mixed-precision inference, not a fallback for when dedicated multiply-accumulate hardware is unavailable. The paper's comparison with GPU-based LUT methods (which failed to achieve practical speedups) demonstrates that the primitive is hardware-sensitive: it requires low-latency, per-thread storage for the tables, which register files provide but shared memory does not (due to bank conflicts, limited capacity per thread, and higher latency). This suggests a design principle for future inference hardware: provide a small, fast, per-processing-element lookup table memory alongside traditional arithmetic units. The paper's conclusion explicitly recommends this: "novel LLM hardware accelerator design based on LUT."

The evidence for the primitive's effectiveness is the consistent speedup across bit-widths and devices in Figure 6, but the more subtle evidence is the optimization breakdown in Figure 10, which shows how each optimization (table quantization, tiling, permutation, interleaving, fast aggregation) incrementally improves upon a baseline that already uses the LUT primitive. The baseline (TM-base with hardware intrinsics but no layout optimizations) is "at most 17% slower compared to the llama.cpp baseline" — meaning even a naive LUT implementation is competitive. The subsequent optimizations take it from competitive to 4× faster, but it's the primitive itself that makes the approach viable in the first place. Without the register-resident LUT, there is no foundation for the speedups; the optimizations amplify the primitive's advantage but do not create it.

5. Experimental Analysis

Evaluation Methodology

Dataset. T-MAC evaluates kernel-level and model-level performance using matrix shapes derived from real-world Llama-2-7B and Llama-2-13B models (Section 5.1), not from a standard benchmark dataset like MATH. For end-to-end model accuracy evaluation, the paper uses three standard NLP tasks: WikiText-2 (perplexity), lambada_openai (perplexity), and WinoGrande (question-answering accuracy), evaluated using llama.cpp's built-in perplexity tool (Section 5.6). The specific GGUF model files are the official llama-2-7b.Q4_0.gguf model released with llama.cpp for the quantized configuration, and a GGUF conversion from original FP16 weights for the un-quantized ground truth (Section 5.6).
Base model(s). The paper evaluates four distinct model types to cover the spectrum of weight quantization: 4-bit Llama models from GPTQ (Frantar et al., 2022); 3-bit and 2-bit Llama models from BitDistiller (Du et al., 2024); 1-bit Llama models from OneBit (Xu et al., 2024); and 1-bit and 1.58-bit BitNet models (Wang et al., 2023) trained from scratch. For the kernel benchmarks, the specific shapes used come from "Llama-2-7B and Llama-2-13B models" (Section 5.1), with matrix dimensions 4096×4096×1, 11008×4096×1, 4096×11008×1, 5120×5120×1, 13824×5120×1, and 5120×13824×1 (Figure 6) — these correspond to the key projection layers in transformer blocks at different model scales.
Metrics. Throughput is measured in tokens per second for end-to-end inference (Section 5.3), with the measurement protocol requiring repeatedly generating 64 tokens for 20 iterations and computing the average token generation rate (Section 5.1). Kernel-level latency is measured in milliseconds per GEMM/GEMV operation, using a warmup of 10 iterations followed by 100 runs to compute an average, with the M2-Ultra requiring "at least 1 second to maximize performance" (Section 5.1). Energy consumption is measured on M2-Ultra using OSX's powermetrics tool, recording average power over 500ms intervals while continuously generating tokens for a minimum of 120 seconds, then integrating power over time to compute total energy per token in Joules (Section 5.4). Model quality is assessed through perplexity (lower is better) on WikiText-2 and lambada_openai, and accuracy (higher is better) on WinoGrande (Section 5.6). Kernel-level numerical error is quantified using Normalized Mean Squared Error (NMSE) relative to an un-quantized FP16 GEMV baseline with randomly generated Gaussian-distributed weights and activations (Section 5.6).
Baselines. The primary baseline is llama.cpp (version b2794, released May 2024), which the paper describes as "a state-of-the-art implementation for LLM deployment on edge devices" featuring "highly optimized kernel implementations tailored to each hardware platform" in plain C/C++ without dependencies (Section 5.1). For mpGEMM (matrix-matrix multiplication at sequence length 256), the paper uses llama.cpp (BLAS) as the baseline, which leverages Accelerate on M2-Ultra and OpenBLAS on other platforms, because "llama.cpp (BLAS) is slower for mpGEMV but faster for mpGEMM compared to llama.cpp's highly optimized mixed-precision implementation" (Section 5.1). For GPU comparisons, the paper uses llama.cpp (GPU) with the CUDA backend on NVIDIA GPUs and the OpenCL backend on Qualcomm GPUs (Section 5.7). For NPU comparisons, performance numbers are "sourced from official data released by Qualcomm via Qualcomm AI Hub" (Section 5.7).
Generation budget / compute accounting. T-MAC uses a bit-width-based accounting: the primary variable controlling computation is the number of bit-planes $n$ into which the weight matrix is decomposed, with each bit-plane requiring one full pass of LUT construction and table lookup over the matrix. This provides a natural scaling axis: 4-bit weights require 4 serial lookup passes, 2-bit requires 2, and 1-bit requires 1. The paper does not measure compute in FLOPs but uses latency and throughput as the ultimate metrics, implicitly accounting for all overhead (LUT construction, memory access, accumulation). For the kernel benchmarks, all measurements are performed at fixed matrix dimensions with bit-width varied (1/2/3/4 bits) to isolate the scaling behavior. The paper also sweeps thread counts: single-threaded (Section 5.2a) and multi-threaded (Section 5.2b) configurations are evaluated separately, with multi-threaded using all available cores as specified in Table 2.
Cross-validation / statistical protocol. The paper does not employ cross-validation in the traditional machine learning sense since this is a systems paper evaluating deterministic kernel performance. The primary statistical rigor comes from measurement protocol: 10 warmup iterations + 100 measurement iterations for kernel latency (Section 5.1), and 20 iterations of 64-token generation for end-to-end throughput. The M2-Ultra warmup differs from other platforms in requiring "at least 1 second to maximize performance" (Section 5.1), which likely relates to the CPU's DVFS (dynamic voltage and frequency scaling) ramp-up behavior. For the model quality evaluation (Section 5.6), the paper runs standard perplexity and accuracy benchmarks once per configuration — no multiple random seeds or error bars are reported for these quality metrics.

Main Quantitative Results

mpGEMV Kernel Performance: Linear Scaling vs. Non-Monotonic Baselines

The headline finding from the kernel benchmarks is that T-MAC achieves linear speedup with bit-width reduction, while llama.cpp shows non-monotonic scaling where 3-bit is often slower than 4-bit due to dequantization overhead. Figure 6 presents these results across four devices and six matrix shapes, with single-threaded and multi-threaded variants.

Single-threaded mpGEMV (Figure 6a). T-MAC achieves maximum speedups over llama.cpp of 11.2×, 5.8×, 4.7×, and 3.1× for 1/2/3/4-bit respectively, with the exact maxima varying by device and shape (Section 5.2). The paper reports that "the 1-bit kernel performance of llama.cpp is deduced from its 2-bit kernel and marked with dashed lines" because llama.cpp does not provide a native 1-bit implementation (Figure 6 caption). This is a revealing detail: llama.cpp has not even implemented 1-bit kernels, presumably because the dequantization overhead would make them slower than 2-bit kernels, rendering the engineering effort pointless under the dequantization paradigm.

The scaling pathology is most visible in the raw latency numbers: "with the bits decrease from 4-bit to 2-bit, llama.cpp fails to gain any additional speedup, and even experiences a 15% slowdown at 3-bit compared to 4-bit due to decoding overhead" (Section 5.2). The paper attributes this to the specific difficulty of 3-bit unpacking: "Since 8 is indivisible by 3, this decoding process is notably inefficient. llama.cpp attempts to optimize it by separately packing 2 bits and the remaining 1 bit, but it still results in significant overhead" (Section 5.2). On the Raspberry Pi 5 single-threaded results, this non-monotonicity is visually striking across all six matrix shapes — the 3-bit bars are taller (slower) than the 4-bit bars, and the 2-bit bars are not substantially shorter than the 4-bit bars. T-MAC's bars, by contrast, form a clean descending staircase from 4-bit to 1-bit.

The 1-bit advantage on T-MAC is particularly dramatic: the single-threaded speedup reaches 11.2× on some configurations because T-MAC needs only one bit-plane of table lookups, while the llama.cpp baseline (projected from 2-bit) still performs full dequantization and multiplication. The paper notes that "T-MAC avoids this problem by individually computing the results of each bit" (Section 5.2) — the uniform treatment of each bit-plane eliminates the packing-dependent overhead.

Multi-threaded mpGEMV (Figure 6b). Multi-threaded performance is primarily constrained by memory bandwidth, and the speedups are correspondingly lower but still significant. T-MAC achieves 4.0×, 4.0×, 5.3×, and 2.5× speedups on the four devices respectively for 2-bit (Section 5.2). The paper explicitly notes that the "T-MAC's performance is primarily constrained by memory bandwidth, but T-MAC can still achieve significant speedup due to efficient memory access" (Section 5.2) — the weight permutation layout ensures sequential DRAM reads, and the elimination of dequantization reduces the total bytes that must be transferred from memory for a given computation.

A notable pattern in the multi-threaded results: the speedup advantage of T-MAC generally increases as bit-width decreases, consistent with the linear scaling property. At 4-bit, T-MAC's advantage is relatively modest (averaging 2-3× depending on shape and device); at 1-bit, it balloons. This confirms that the primary source of T-MAC's advantage is not constant-factor optimization but the algorithmic elimination of operations proportional to bit-width.

The 3-bit sweet spot. The paper highlights that "T-MAC demonstrates a significant advantage at 3-bit precision" (Section 5.2), which is where the dequantization overhead is worst for llama.cpp. This has practical significance: 3-bit quantization represents a compelling accuracy-efficiency tradeoff for many models, but dequantization-based inference systems cannot realize its speedup potential. T-MAC makes 3-bit inference actually faster than 4-bit, fulfilling the promise of the quantization.

mpGEMM Kernel Performance: Prefill-Stage Speedups

For the prefill stage (matrix-matrix multiplication with sequence length 256), T-MAC's performance relative to the BLAS-accelerated baseline is strong but more varied (Figure 7). The paper uses llama.cpp (BLAS) as the baseline because BLAS-accelerated GEMM outperforms llama.cpp's hand-tuned dequantization kernels for larger batch dimensions.

T-MAC achieves significant speedups on Raspberry Pi 5, Jetson AGX Orin, and Surface Book 3 for 2-bit with maximum speedups of 4.0×, 5.3×, and 5.3× respectively (Section 5.2). However, M2-Ultra is a notable exception: the paper states that "M2-Ultra is an exception, as the Apple Silicon CPUs are equipped with a powerful AMX co-processor to handle GEMM operations. However, T-MAC still achieves a maximum 2.0× speedup for 1-bit in this case" (Section 5.2). The AMX (Apple Matrix coprocessor) is a dedicated matrix multiplication engine that provides extremely high throughput for standard-precision GEMM — it represents the hardware that T-MAC is competing against at the high end. Even against this specialized accelerator, T-MAC's 1-bit LUT approach outperforms by 2×, suggesting that the AMX's fixed-function multiply-accumulate design cannot match the operation reduction achieved by LUT-based computation at extreme bit-widths.

The scaling pattern in Figure 7 mirrors the GEMV results: T-MAC's latency decreases linearly with bit-width, while the BLAS baseline (projecting 1-bit from 2-bit with dashed lines) shows little improvement from 4-bit to 2-bit. The 1-bit T-MAC bars are consistently the shortest across all devices and shapes, confirming the linear scaling property extends from GEMV to GEMM.

The paper notes that 1-bit kernel performance for the baseline is "deduced from its 2-bit kernel and marked with dashed lines" (Figure 7 caption), the same projection approach used for GEMV. This repeated pattern — the absence of native 1-bit kernels in both llama.cpp and BLAS backends — reinforces the paper's motivation: existing systems have structurally abandoned the lowest bit-widths because their dequantization approach cannot deliver speedups there.

End-to-End Inference Throughput: Integrated System Performance

After integrating T-MAC kernels into llama.cpp, the paper evaluates end-to-end token generation throughput on three models: Llama-2-7B-4bit (M1), Llama-2-7B-2bit (M2), and BitNet-3B (M3), across four devices with both single-threaded and multi-threaded configurations (Figure 8).

Single-threaded throughput (Figure 8a). On Raspberry Pi 5, T-MAC achieves speedups of 2.8×, 6.7×, and 5.8× for the three models respectively (Section 5.3). The absolute throughput numbers are noteworthy: T-MAC delivers 11 tokens/s on Raspberry Pi 5 for BitNet-3B (Figure 8, RBP bar for M3), which the paper frames as a milestone for "real-world edge deployment" on the least powerful device tested. On M2-Ultra, T-MAC reaches a peak of 71 tokens/s for BitNet-3B (Figure 8, M2-Ultra multi-threaded bar for M3).

Multi-threaded throughput (Figure 8b). The speedups are less dramatic under multi-threading because the overall inference pipeline includes operations beyond mpGEMV/mpGEMM (attention, layer normalization, residual connections) that are not accelerated by T-MAC. The paper acknowledges that "due to memory constraints and operators other than mpGEMV/mpGEMM, the speedup is less pronounced" (Section 5.3). On M2-Ultra, the speedups are 1.1×, 2.3×, and 1.7× for M1, M2, M3 respectively. On Surface Book 3, the speedups are 1.4×, 2.6×, and 2.4×. On AGX Orin, the pattern is 1.7×, 2.4×, 2.0×.

The 4-bit model (M1) shows the smallest relative speedup (as low as 1.1× on M2-Ultra) because at 4-bit, the baseline dequantization overhead is already low (uniform nibble packing), and the remaining inference overhead (attention, etc.) dominates. The 2-bit model (M2) and BitNet-3B (M3) show much larger speedups because the dequantization penalty is severe for the baseline while T-MAC's LUT approach handles the lower bit-width efficiently.

The BitNet advantage. BitNet-b1.58-3B (M3) demonstrates particularly strong results with T-MAC. The paper interprets the ternary weights of 1.58-bit BitNet "as 2-bit and decomposed into two 1-bit matrices" (Section 5.1). With only 2 effective bit-planes, T-MAC's computation is extremely lightweight while the model quality remains competitive. The throughput of 30 tokens/s single-core and 71 tokens/s eight-core on M2-Ultra, and 11 tokens/s on Raspberry Pi 5, represent what the paper considers "promising real-world edge deployment" (Section 5.3).

Power and Energy Consumption: The Efficiency Case for CPUs

The energy evaluation (Figure 9, Table 5) directly addresses the battery-life constraint that the paper identifies as critical for edge deployment (Section 2.1). On M2-Ultra multi-threaded inference:

For Llama-2-7B-4bit, T-MAC reduces power consumption by 10.3% and total energy per token by 20.6% compared to llama.cpp (Section 5.4).
For Llama-2-7B-2bit, T-MAC reduces power by 10.3% and energy by 61.2%.
For BitNet-3B, T-MAC reduces power by 17.3% and energy by 51.3%.

The power reductions are moderate but consistent (10-17%), which the paper attributes to the elimination of multiplication instructions and reduced memory traffic. The energy savings are much larger than the power savings alone because T-MAC's throughput is also higher — energy integrates power over time, so the combination of lower power and shorter inference duration produces multiplicative savings.

The 61.2% energy reduction for the 2-bit model is particularly striking. This means T-MAC can process more than twice as many tokens per Joule as llama.cpp on the same hardware. For a battery-operated edge device, this directly translates to extended operational lifetime or the ability to run larger models within the same energy budget.

The NVIDIA Jetson AGX Orin comparison in Table 5 provides the energy-efficiency argument across computational substrates. Comparing llama.cpp on CPU, llama.cpp on GPU, and T-MAC on CPU for Llama-2-7B-2bit:

llama.cpp (CPU): 7.08 tokens/s, 15.0W, 2.12 J/token
llama.cpp (GPU): 20.03 tokens/s, 30.8W, 1.54 J/token
T-MAC (CPU): 15.62 tokens/s, 10.4W, 0.66 J/token

T-MAC on CPU achieves 2.2× the throughput and 3.2× the energy efficiency of llama.cpp on CPU. Compared to llama.cpp on GPU, T-MAC delivers 78% of the throughput but consumes only 34% of the power, resulting in 2.3× better energy efficiency (0.66 vs. 1.54 J/token). The paper frames this as a favorable tradeoff for edge deployment: slightly lower throughput is acceptable if it comes with dramatically lower energy consumption, since battery life is often the harder constraint than raw speed.

CPU vs. GPU/NPU: Cross-Platform Performance

The cross-platform comparison (Table 7, Figure 11, Table 5) provides the evidence for the paper's "CPU Renaissance" thesis — that CPUs running LUT-based computation can match or exceed dedicated accelerators for low-bit LLM inference.

Kernel-level comparison (Figure 11). On Jetson AGX Orin, comparing T-MAC (CPU) against llama.cpp (GPU) for mpGEMV kernels from Llama-2-7B: "T-MAC significantly outperforms GPU on W1A16 on all cases, while achieves comparable performance on W2A16 and W3A16. Although GPU performs better on higher bits and larger shape due to its powerful parallel computing capacity, this evaluation still shows huge potential of CPU-based LLM deployments on edge devices" (Section 5.7). This is a nuanced result: the GPU wins at 4-bit and on the largest matrix shape (4096×11008×1), but T-MAC wins decisively at 1-bit and is competitive at 2-bit and 3-bit. The crossover point — where the CPU overtakes the GPU — shifts toward higher bit-widths as the matrix shape gets smaller (less GPU parallelism to exploit).

End-to-end comparison across devices (Table 7). On three platforms with both CPU and GPU/NPU capabilities:

Surface Laptop 7 (Snapdragon X Elite): T-MAC achieves 31.83 tokens/s for 2-bit, compared to 9.39 for llama.cpp CPU, 10.40 for NPU (Qualcomm Hexagon). This is a 3× speedup over the NPU, the specialized AI accelerator. For 4-bit, T-MAC reaches 21.63 tokens/s vs. 10.64 for CPU and 10.40 for NPU — a 2.1× advantage over the NPU. The paper notes this uses "only 4 out of the total 12 CPU cores" (Table 7 description).
OnePlus 12 (Snapdragon 8 Gen 3): T-MAC achieves 16.62 tokens/s for 2-bit, compared to 6.95 for llama.cpp CPU, 1.72 for GPU (Adreno 750, OpenCL), and 11.30 for NPU. The speedup over GPU is 9.7× for 2-bit and 6.4× for 4-bit. Over NPU, T-MAC achieves 1.5× speedup for 2-bit, though the NPU retains a slight edge for 4-bit (11.30 vs. 10.19 tokens/s). The paper notes this uses only 4 of 8 CPU cores (1 Cortex-X4 + 3 Cortex-A720).
Jetson Orin NX: T-MAC achieves 11.41 tokens/s for 2-bit, compared to 3.20 for llama.cpp CPU and 7.94 for GPU (Ampere GA10B). This is a 1.4× speedup over the GPU for 2-bit, and 1.9× for 4-bit (7.53 vs. 3.97 CPU, 7.94 GPU — essentially matching the GPU for 4-bit).

The key interpretive note for the NPU comparison: "The 2-bit performance of NPUs is deduced from 4-bit and marked with '*' " (Table 7 caption). This is because NPUs typically lack native 2-bit support — they run 2-bit models by dequantizing to a supported precision (e.g., INT8), similar to how CPUs without T-MAC handle mixed-precision. This means the NPU numbers are projections, not measured results, and they likely suffer from the same dequantization pathology that T-MAC eliminates on CPUs. The actual NPU performance for 2-bit might be worse than these projections if the dequantization overhead is significant.

Model Quality: Accuracy Preservation

The error analysis in Section 5.6 addresses the natural concern that LUT-based computation — with its table quantization and optional fast aggregation — might degrade model accuracy compared to standard dequantization-based inference.

Kernel-level numerical error (Table 3). The NMSE relative to un-quantized FP16 GEMV is nearly identical between llama.cpp and T-MAC. For the three tested matrix shapes (4096×4096×1, 11008×4096×1, 4096×11008×1), T-MAC's NMSE values are 3.35e-03, 3.46e-03, and 4.15e-03, compared to llama.cpp's 3.33e-03, 3.44e-03, and 4.13e-03 — differences on the order of 0.01-0.02e-03, which the paper characterizes as "negligible" (Section 5.6). This demonstrates that the table quantization technique (quantizing LUT entries from FP16 to INT8) introduces minimal numerical error at the kernel level.

However, when fast aggregation is enabled, the NMSE increases to 8.09e-03, 8.27e-03, and 8.45e-03 — approximately 2.5× higher than the baseline. This reflects the error introduced by the averaging-based INT8 accumulation (using vrhaddq_u8/_mm256_avg_epu8), which trades precision for instruction throughput.

Model-level quality (Table 4). On Llama-2-7B-4bit with single-threaded inference on M2-Ultra:

T-MAC (without fast aggregation): WikiText2 perplexity 5.96 (vs. 5.96 for llama.cpp), lambada_openai perplexity 12.95 (vs. 12.95 for llama.cpp), WinoGrande accuracy 70.8 (vs. 70.8 for llama.cpp). These are identical across all three metrics to the llama.cpp baseline, confirming that the kernel-level NMSE difference is imperceptible to end-to-end model quality.
T-MAC (+FA, with fast aggregation): WikiText2 perplexity degrades to 6.38 (increase of 0.42), lambada_openai to 13.99 (increase of 1.04), and WinoGrande accuracy drops to 67.8 (decrease of 3.0 percentage points). This is a measurable quality degradation that the paper characterizes as making fast aggregation suitable only for "scenarios that prioritize real-time performance and are less sensitive to accuracy" (Section 5.6).

The un-quantized baseline achieves 5.80 WikiText2 perplexity and 71.0 WinoGrande accuracy, confirming that the 4-bit quantization itself introduces a quality gap (perplexity increases from 5.80 to 5.96) that is identical whether using llama.cpp or T-MAC.

The paper notes that the fast aggregation error "can be mitigated with straightforward optimizations of the CPU micro-architecture" (Section 5.6), suggesting that the issue is not fundamental to the LUT approach but an artifact of current CPU instructions lacking INT8 accumulate-to-INT16 operations with full INT8 throughput. This is a forward-looking claim about hardware improvements rather than a software fix.

Ablation Studies and Robustness Checks

Optimization breakdown (Figure 10): The paper evaluates the incremental contribution of each optimization by starting from a baseline LUT implementation (TM-base, which uses hardware table-lookup intrinsics but no memory access optimizations) and adding optimizations one by one, measuring multi-threaded GEMV latency on M2-Ultra across six matrix shapes (S0-S5):

TM-base: Performance is "at most 17% slower compared to the llama.cpp baseline" (Section 5.5). This establishes that even a naive LUT implementation — with no data layout optimizations, no tiling, no weight permutation — is competitive with the heavily optimized dequantization baseline.
+Table Quantization (TQ): Performance becomes "competitive with llama.cpp" (Section 5.5), bringing the LUT approach to parity with the state of the art. The paper does not specify the exact percentage improvement, but the bar positions in Figure 10 show TQ closing most of the gap.
+Tiling: "The tiling optimization further yields a maximum speedup of 1.45×" (Section 5.5) over the TM+TQ configuration. This is the optimization that introduces LUT-centric tiling with large $M_{tm}$ to amortize LUT construction cost — the key conceptual innovation in the data layout.
+Permutation: "Permutation contributes an additional 1.39× speedup" by rearranging data layout into contiguous memory for each tile, maximizing DRAM bandwidth utilization through sequential access patterns.
+Tuning: AutoTVM tuning "does not appear to be very effective in the figure, as the default tiling configurations already align well with M2-Ultra registers and caches," but is expected to be more impactful on different devices where the default parameters are less optimal.
+Interleaving (T-MAC): Weight interleaving "eliminates most of the unpacking overhead, achieving a 1.42× speedup" (Section 5.5). This is the offline pre-rearrangement of weight bits to eliminate runtime byte-reversal operations.
+Fast Aggregation (TM+FA): "The aggressive fast aggregation can make T-MAC up to 1.29× faster, but it could lead to non-negligible accuracy loss, so we offer it as an optional optimization" (Section 5.5).

The cumulative effect of these optimizations transforms a slightly-slower-than-baseline implementation into one that delivers 4-6× speedups (depending on bit-width). The paper notes that "Most of these optimizations yield greater benefits for single-threading, but tiling requires multi-threading to be effective, hence this evaluation is conducted using multi-threading" (Section 5.5). This reveals an important interaction: tiling's benefits are amplified when multiple cores share memory bandwidth, because the LUT reuse pattern reduces total memory traffic per core.

Bit-serial linear transformation values: The paper reports empirically selecting $s_0 = -1, s_1 = 1$ from the candidate set $\{-1, 0, 1\}$ for the linear transformation that maps weight bits to LUT entries (Section 4, "Bit-serial linear transformation"). The rationale given is twofold: these values "circumvent float-multiply instructions" (only addition and subtraction needed for LUT construction) and "minimize the difference between the largest and smallest entries of the lookup table" (reducing the dynamic range for better quantization). The paper does not present an ablation comparing {-1,+1} to {0,1} or {0,-1} in terms of LUT range or quantization error — this appears to be a design choice justified by reasoning rather than controlled experiment.

Fast aggregation accuracy tradeoff (Table 3, Table 4): The fast aggregation ablation is the most thoroughly characterized negative result. At the kernel level, fast aggregation increases NMSE by approximately 2.5× across all tested shapes (Table 3: from ~3.5-4.2e-03 to ~8.1-8.5e-03). At the model level, the degradation is noticeable but not catastrophic: WikiText2 perplexity increases by 0.42 (7.1% relative), lambada_openai by 1.04 (8.0% relative), and WinoGrande accuracy drops by 3.0 percentage points (4.2% relative). The paper positions fast aggregation as an optional feature for latency-sensitive, accuracy-tolerant deployments, and notes that future CPU micro-architecture improvements could eliminate the accuracy penalty by supporting efficient INT8-wide accumulation.

Per-device performance consistency: While not presented as a formal ablation, the consistent scaling behavior across four devices (M2-Ultra, Raspberry Pi 5, Jetson AGX Orin, Surface Book 3) with different CPU architectures (Apple Silicon, ARM Cortex-A76, ARM Cortex-A78AE, Intel Core i5-1035G7) and different instruction sets (ARM NEON, Intel AVX2, Apple AMX) serves as a robustness check for the LUT approach. The fact that T-MAC achieves speedups on all platforms — from a Raspberry Pi 5 with 17.1 GB/s memory bandwidth to an M2-Ultra with 819.2 GB/s — suggests the approach is not fragile to specific hardware characteristics. The paper does note platform-specific optimizations (e.g., LUT duplication for AVX2's lane-based shuffle, AMX co-processor handling on Apple Silicon) but the core LUT-based computation paradigm transfers across architectures.

Kernel shape coverage: The six matrix shapes tested (4096×4096, 11008×4096, 4096×11008, 5120×5120, 13824×5120, 5120×13824) cover the key operations in Llama-2 transformer blocks: the 4096/5120 dimensions correspond to hidden sizes, 11008/13824 correspond to feed-forward intermediate sizes. The transposed shapes (K×N vs. N×K) test both the case where weight is the second operand (GEMV along K) and where it is the first operand (GEMV along N). The paper does not test matrix shapes from models with significantly different architectures (e.g., non-LLaMA models with different aspect ratios), which limits the generality claim to LLaMA-family architectures. However, the LUT-based approach itself makes no assumptions about matrix shape, so the limitation is in the evaluation coverage rather than the method.

Critical Assessment

The experimental evaluation provides strong evidence for T-MAC's core performance claims but has important limitations in model quality evaluation scope, long-sequence behavior, and the definition of "edge deployment" across the tested hardware spectrum.

What the Experiments Demonstrate Convincingly

The linear scaling claim — that T-MAC computation cost scales linearly with weight bit-width — is directly supported by the kernel benchmarks in Figure 6 and Figure 7. Across all four devices, all six matrix shapes, and both single-threaded and multi-threaded configurations, T-MAC's latency forms a descending staircase from 4-bit to 1-bit, while llama.cpp's latency is either flat or non-monotonic (Section 5.2). This is the most robust result in the paper because it is replicated across such a wide hardware and shape matrix. The evidence for the underlying mechanism — that dequantization overhead causes the non-monotonicity in llama.cpp — is persuasive but circumstantial: the paper attributes the 3-bit slowdown to byte-misalignment (Section 5.2), which is a well-known issue in bit-packed formats, but does not provide instruction-level profiling to confirm that unpack instructions are the dominant overhead. This would be a stronger claim with perf or similar micro-architectural counter data showing the fraction of cycles spent in SHIFT/AND sequences for 3-bit unpacking.

The energy efficiency claim — 70% reduction in energy consumption — is directly supported by Figure 9 and Table 5. The measurement methodology (120-second continuous generation, powermetrics integration) is sound. However, the "up to 70%" figure quoted in the abstract and Section 1 appears to round up from the 61.2% and 51.3% reductions documented in Section 5.4 for the 2-bit and BitNet models. The 70% figure is not achieved on any single model in the reported experiments; it may derive from an unshown configuration or represent a projected maximum across all tested scenarios. This is a minor but notable overstatement. The 4-bit model shows only 20.6% energy reduction — still meaningful, but a qualitatively different magnitude from "70%." The paper would be more precise to say "up to 61% for 2-bit models and 20% for 4-bit models" rather than the blanket "70%."

The CPU-vs-GPU parity claim — that T-MAC on CPU can match or exceed GPU inference speed — is partially supported by Figure 11 and Table 7, but is narrower than the framing suggests. Figure 11 shows T-MAC CPU outperforming llama.cpp GPU at 1-bit and being competitive at 2-bit and 3-bit on Jetson AGX Orin. Table 7 shows T-MAC outperforming NPUs and GPUs on Surface Laptop 7 and OnePlus 12, particularly at 2-bit. These results are real and impressive. However, they are specific to: (a) low-bit weights (1-2 bit), where T-MAC's advantage is largest; (b) the specific GPU backends tested (llama.cpp CUDA and OpenCL, not cuBLAS or TensorRT-LLM); and (c) devices with unified memory architecture where CPU and GPU share the same DRAM bandwidth. The paper does not compare against GPU inference systems that use optimized dequantization kernels like BitBLAS or Marlin (cited in the paper but not evaluated against). A comparison against BitBLAS on GPU would test whether the GPU LUT-GEMM performance results from Section 2.4 (1.75-2.34× slower than dequantization) still hold with more recent GPU dequantization kernels. If the GPU dequantization performance has improved, the CPU advantage might narrow. The paper's conclusion that T-MAC "makes the CPU inference speed comparable or even higher than the GPU on the same device" (Section 7) should be qualified: this holds for low-bit weights (1-2 bit) on the tested hardware with the tested GPU backends, and may not generalize to 4-bit or to GPU inference systems with more aggressive kernel optimization.

The "practical deployment" claim — that T-MAC enables LLM inference on resource-constrained devices — is supported by the Raspberry Pi 5 results (11 tokens/s for BitNet-3B, Figure 8) but needs context. The paper does not report model quality (perplexity, task accuracy) for the low-bit models on the edge devices — only the 4-bit quality results in Table 4. The 2-bit and BitNet models are evaluated only for throughput, not for accuracy. A "practical deployment" requires both adequate throughput and acceptable quality. If the 2-bit model has significantly degraded quality on the target tasks, 11 tokens/s is irrelevant. The paper should have included perplexity or downstream accuracy numbers for the 2-bit and BitNet models to substantiate the practical deployment claim. The BitDistiller and OneBit papers (cited as the source of the 2-bit and 1-bit models) do report quality metrics, but the paper does not reproduce or reference these in the context of the T-MAC inference system. The quality metrics for the 4-bit model are convincing (Table 4 shows T-MAC quality identical to llama.cpp), but the leap to 2-bit and 1-bit quality preservation with T-MAC is assumed rather than demonstrated.

Limitations and Missing Experiments

Model quality at 2-bit and 1-bit is not evaluated. Table 4 evaluates quality only for Llama-2-7B-4bit. The paper does not report perplexity or downstream accuracy for the 2-bit, 1-bit, or BitNet models integrated with T-MAC. Since table quantization and the bit-serial decomposition could in principle interact differently with these more aggressive quantization schemes (where the model is already operating near the accuracy cliff), this is a significant gap. The paper argues that "table quantization technique has an imperceptible effect on the overall model accuracy" (Section 3.3), but this claim is only validated at 4-bit. A practitioner wanting to deploy a 2-bit model with T-MAC does not know whether the quality will match the original 2-bit dequantization-based inference.

The prefill stage (long sequence length) is only evaluated at length 256 (Figure 7). Production LLM deployments handle sequence lengths up to 4K, 8K, or 32K tokens. The GEMM evaluation at 256 tokens shows strong speedups, but scaling to longer sequences changes the compute-to-memory ratio and the LUT reuse pattern. At very long sequences, the GEMM dimension $N$ (sequence length) becomes large enough that the tile size $N_{tn}$ might need to increase, changing the register pressure dynamics. The paper does not explore whether T-MAC's speedups hold or degrade at prefill sequence lengths of 2048 or 4096. This is a practical concern for deployment where prompt processing time matters as much as token generation time.

The paper does not evaluate against MLC-LLM or other CPU inference systems besides llama.cpp. llama.cpp is the most widely used CPU inference framework and a reasonable primary baseline, but MLC-LLM (with TVM Unity) is another major CPU-targeting system that the paper's TVM-based code generation could be compared against more directly. Since T-MAC itself uses TVM, comparing against MLC-LLM's CPU backend would test whether the speedups come from the LUT approach or from TVM's code generation optimizations. The absence of this comparison makes it harder to attribute the performance advantage specifically to LUT.

The 3-bit llama.cpp performance degradation is attributed to byte misalignment but not profiled in detail. The claim that "llama.cpp attempts to optimize it by separately packing 2 bits and the remaining 1 bit, but it still results in significant overhead" (Section 5.2) is a specific algorithmic claim about llama.cpp's implementation. However, the paper does not provide instruction-level profiling (e.g., perf stat instruction counts, cache miss rates, or pipeline stall analysis) to confirm that the SHIFT/AND unpacking for 3-bit is the bottleneck. The slowdown could also be caused by memory layout inefficiencies, TLB misses, or other factors. The inference is plausible — 3-bit packing is notoriously awkward — but without profiling, it remains an inference rather than a demonstrated mechanism.

The energy measurements are limited to M2-Ultra (Section 5.4). The paper evaluates power and energy only on a single high-end device. Energy efficiency is most critical for battery-operated devices like the Raspberry Pi and Jetson Orin, where T-MAC's energy savings might be even more impactful (since these devices have smaller batteries). The paper does not report energy consumption on Raspberry Pi 5 or Surface Book 3 (battery-operated laptop), which limits the generalizability of the energy efficiency claim beyond the specific M2-Ultra platform. Measuring power on a Raspberry Pi is straightforward (USB power meter), so this omission is notable.

The claimed "up to 4× increase in throughput" (abstract) maps to specific bit-width and device combinations. In the end-to-end results (Figure 8), the largest speedup shown is 6.7× for Llama-2-7B-2bit on Raspberry Pi 5 single-threaded, and 5.8× for BitNet-3B on the same configuration. The 4× figure in the abstract appears to be a conservative summary, but the actual speedups vary dramatically — from 1.1× (Llama-2-7B-4bit multi-threaded on M2-Ultra) to 6.7×. The "4×" framing overstates the typical case for 4-bit models and understates it for 2-bit on low-end hardware. A range or bit-width-specific statement would be more informative.

The paper does not test mixed bit-width models (different layers at different precisions). Real-world deployment scenarios increasingly use mixed-precision quantization where some layers are 4-bit and others 2-bit. T-MAC's unified bit-serial approach should handle this naturally (each layer just uses the appropriate number of bit-planes), but the paper does not demonstrate this flexibility. Since the unified abstraction is claimed as a key advantage over case-by-case kernel design (Innovation 4), testing it on a mixed-precision model would directly validate this claim.

The impact of table quantization granularity is not ablated. The paper states that table quantization uses "finer granularity (quantizing 8 values for k=4)" compared to conventional activation quantization (Section 3.3). This is a specific design choice — 8 values per quantization group — but the paper does not ablate the group size. Would quantizing per-16 or per-4 values change the accuracy? Is the 8-value grouping tied to the $g=4$ group size and the 16-entry LUT? Understanding this tradeoff matters because finer granularity increases the number of scale factors (overhead for storage and computation) while coarser granularity potentially increases quantization error. The absence of this ablation means the "imperceptible accuracy effect" claim is validated only for one specific quantization configuration.

The paper does not evaluate inference latency variance or tail latency. All results report average latency over 100 runs (Section 5.1). For interactive edge applications, tail latency (p95, p99) matters more than average — a user experiences the slow tokens, not the average token. The LUT-based approach with its data-dependent access patterns (weight indices determine table entry) could in principle exhibit higher variance than the deterministic multiply-accumulate pipeline. Without tail latency numbers, we cannot assess whether T-MAC's speedup comes at the cost of less predictable per-token latency, which would be problematic for real-time applications.

Conditional Validity of Claims

The paper's central claims are conditionally valid in ways that matter for practitioners:

"Up to 4× throughput increase" holds for 2-bit models on CPU-bound workloads (single-threaded, low-end devices) but shrinks to 1.1-1.4× for 4-bit models on memory-bandwidth-bound multi-threaded configurations on high-end devices (M2-Ultra multi-threaded, Figure 8). A practitioner with a 4-bit model on M2-Ultra should expect modest improvement; a practitioner with a 2-bit model on Raspberry Pi should expect dramatic improvement.
"CPU matches GPU performance" holds for 1-2 bit weights on the tested platforms with unified memory (Jetson AGX Orin, Snapdragon devices) but would likely not hold for 4-bit weights on a device with discrete GPU memory (where the GPU has higher dedicated memory bandwidth) or against more optimized GPU backends (TensorRT-LLM, cuBLAS).
"Unified and scalable solution" is demonstrated for bit-widths 1-4 (the scaling is linear and monotonic across this range) but the claim of supporting any mixed-precision combination is tested only on the specific {weight bits = 1,2,3,4} × {activation precision = FP16 or FP8} combinations. The paper does not test W8A16 (8-bit weights with FP16 activations), which would require 8 bit-planes and might become memory-bound rather than compute-bound.
"Practical deployment without relying on GPUs" is demonstrated on Raspberry Pi 5 and Surface Book 3, but the Raspberry Pi results are limited to BitNet-3B (a 3B parameter model) and the Surface Book results to Llama-2-7B-2bit. Deploying a 7B model on Raspberry Pi 5 might be memory-capacity-limited regardless of inference speed — 7B parameters at 2 bits is ~1.75 GB, within the 4-8 GB RAM typical of Pi 5, but the paper does not demonstrate this configuration end-to-end. The "practical deployment" claim is most strongly supported for the 3B BitNet model.
Energy efficiency gains are demonstrated on M2-Ultra, a device with an active cooling system and high-power envelope. The energy savings might be proportionally larger on passively cooled, battery-operated devices (where the baseline power is lower but every Joule counts more), but this is not tested. Conversely, the absolute power numbers (10-15W on M2-Ultra) are not directly applicable to smartphone-class devices (1-5W typical SoC power).

6. Limitations and Trade-offs

6.1 The Difficulty Estimation Cost Is Not Accounted for in Headline Efficiency Numbers

The assumption or constraint. T-MAC's performance gains rest on a computational approach that replaces dequantization with table lookups, but the method itself introduces a new cost: building lookup tables online from the activation matrix before the table lookup can proceed. The paper's evaluation methodology measures kernel latency and end-to-end throughput after this LUT construction cost is incurred, so the cost is properly included in the measured results. However, there is a more subtle accounting issue that the paper does not address: the relationship between tiling granularity and LUT construction frequency determines what fraction of total computation is spent building tables versus performing lookups, and this fraction varies with matrix shape, bit-width, and hardware. Unlike the dequantization approach where the overhead is a fixed per-element cost, T-MAC's LUT construction overhead is amortized over the $M_{tm}$ tiling parameter — the number of weight rows that reuse each LUT. The paper states that "a larger tile size $M_{tm}$ on $M$ can lead to better lookup table reusing" (Section 3.2), but the measured speedups in Figure 6 reflect the specific tiling configurations found by AutoTVM for the tested shapes. A practitioner with a different model architecture (different matrix aspect ratios) may get different LUT amortization efficiency, and the paper provides no model for predicting when LUT construction will become the bottleneck.

The consequence. For matrix shapes where the $M$ dimension is small relative to $K$ (e.g., narrow layers in non-LLaMA architectures, or the early layers of some vision transformers), there are fewer weight rows over which to amortize each LUT. The LUT construction cost — which scales with $2^g$ per group and requires gather/swizzle instructions (Section 4, "Register swizzling for efficient LUT precomputation") — could become the dominant cost, potentially making T-MAC slower than dequantization-based approaches for those layers. The paper does not characterize the threshold $M_{tm}$ below which LUT construction overhead overtakes the lookup savings, leaving practitioners unable to predict whether T-MAC will help or hurt for their specific model geometry.

What evidence exists in the paper. The paper does not directly measure or report LUT construction time as a fraction of total kernel latency. The optimization breakdown in Figure 10 shows that the tiling optimization contributes a 1.45× speedup, which implicitly confirms that LUT reuse matters, but the absolute LUT construction cost is never isolated. The six tested matrix shapes (Figure 6) all come from Llama-2 architectures and have comparable $K$ -to- $M$ ratios; shapes with different aspect ratios are not evaluated. The paper also does not explore how the group size $g$ (fixed at 4 throughout) affects the LUT construction-to-lookup cost ratio — larger $g$ reduces the number of lookups (fewer groups along $K$ ) but exponentially increases the LUT construction cost ( $2^g$ products to compute per group).

Mitigation status. The paper does not address this tradeoff or propose a model for predicting when LUT construction dominates. The AutoTVM tuning is mentioned as a mechanism to find good tiling parameters, but it only optimizes over the parameter space — it does not provide analytical guidance about when the LUT approach is fundamentally mismatched to a given matrix shape. The paper's silence on this point means that T-MAC's performance portability across model architectures is an open question.

6.2 Model Quality at 2-Bit and 1-Bit Is Not Evaluated, Yet the Largest Speedup Claims Are at These Bit-Widths

The assumption or constraint. The paper's most dramatic performance claims — 11.2× kernel speedup at 1-bit (Section 5.2), 6.7× end-to-end speedup for 2-bit models on Raspberry Pi 5 (Section 5.3), 11 tokens/s on Raspberry Pi for BitNet-3B (abstract) — are all at bit-widths of 2 or below. However, the paper evaluates model quality (perplexity, downstream accuracy) only for Llama-2-7B-4bit (Table 4). For the 2-bit Llama models (from BitDistiller), the 1-bit Llama models (from OneBit), and the BitNet models, the paper reports only throughput — no perplexity, no task accuracy, no comparison of T-MAC inference quality against the original model's reported quality. The paper's claim that "table quantization technique has an imperceptible effect on the overall model accuracy" (Section 3.3) is substantiated only at 4-bit, with kernel-level NMSE measurements (Table 3) also at 4-bit (using randomly generated Gaussian weights and activations quantized to 4-bit).

The consequence. The paper cannot guarantee that T-MAC preserves model quality at the bit-widths where its speedups are largest. This is particularly concerning because 2-bit and 1-bit models operate near the edge of acceptable accuracy — the quantization itself already induces significant quality degradation, and any additional error from T-MAC's table quantization or bit-serial decomposition could push the model below a usability threshold. For a practitioner considering deploying a 2-bit model with T-MAC on a Raspberry Pi, the paper provides no evidence that the model will produce coherent or correct outputs. The fast aggregation results (Table 4) demonstrate that aggressive optimization can measurably degrade quality (WinoGrande accuracy drops 3 points), and while T-MAC without fast aggregation matches llama.cpp at 4-bit, this equivalence is not demonstrated at 2-bit or 1-bit. The concern is not purely theoretical: at lower bit-widths, the model's weight values are coarser, and any interaction between the bit-serial linear transformation (mapping {0,1} to {-1,+1} with scaling and bias terms) and the quantization scheme could introduce systematic errors that are masked at 4-bit but revealed at 2-bit.

What evidence exists in the paper. The quality evaluation is limited to Table 4 (Llama-2-7B-4bit on M2-Ultra, single-threaded) and Table 3 (kernel-level NMSE at 4-bit for three matrix shapes). The paper does not evaluate perplexity on WikiText-2, lambada_openai, or accuracy on WinoGrande for the BitDistiller 2-bit and 3-bit models, the OneBit 1-bit models, or the BitNet models. The BitDistiller, OneBit, and BitNet papers (cited in Section 5.1) do report quality metrics for these models, but those results were produced with dequantization-based inference. The paper does not reproduce these quality numbers under T-MAC inference, nor does it compare T-MAC's output logits against the dequantization baseline's output logits to quantify distributional shift. Section 5.6 explicitly discusses error sources but only validates them at 4-bit.

Mitigation status. Not addressed. The paper does not acknowledge this as a limitation. The quality evaluation section is titled "Error Analysis" but it analyzes only one error source (table quantization + fast aggregation) at one bit-width. There is no discussion of whether quality preservation extends to lower bit-widths, and no caveat attached to the 2-bit and 1-bit speedup claims. This is a significant omission for a paper that markets itself as enabling "practical deployment" of low-bit LLMs — a deployment is not practical if the model quality is unknown.

6.3 The Evaluation Is Constrained to LLaMA-Family Matrix Shapes and Does Not Test Robustness to Architectural Diversity

The assumption or constraint. All kernel benchmarks (Figure 6, Figure 7) and end-to-end evaluations (Figure 8, Table 7) use matrix dimensions derived from Llama-2-7B and Llama-2-13B architectures (Section 5.1). The six shapes tested — 4096×4096, 11008×4096, 4096×11008, 5120×5120, 13824×5120, 5120×13824 — correspond to the hidden-size and intermediate-size projections in these specific transformer variants. The paper's code generation approach (using TVM + AutoTVM) and the tiling optimization (with its asymmetric $M_{tm}$ emphasis) are claimed to be general, but they are evaluated only on these shapes. The paper does not test models with different architectural patterns: non-LLaMA autoregressive models (e.g., Gemma with different hidden-to-intermediate ratios), encoder-decoder architectures (T5, BART), mixture-of-experts models (where the expert layers have different shapes), or vision transformers (with different attention patterns and patch embedding dimensions).

The consequence. The paper's core performance claims and the LUT-centric tiling strategy may not generalize to architectures with different matrix aspect ratios. The finding that larger $M_{tm}$ improves LUT reuse (Section 3.2) depends on the $K$ -to- $M$ ratio — if a model has layers where $M$ is small (e.g., narrow output projections, or the query/key/value projections in multi-head attention when the per-head dimension is small), the LUT amortization benefit shrinks. The AutoTVM tuning searches over tile sizes for the specific shapes it is given, but it cannot overcome the fundamental limitation that if $M$ is small, there are simply fewer weight rows per LUT. Worse, for models with very long sequence lengths (large $N$ in the prefill stage), the $N_{tn}$ tiling parameter must increase, potentially increasing register pressure and LUT storage requirements in ways that the current tiling strategy does not account for. The paper only evaluates prefill at sequence length 256 (Figure 7) — the point where prefill transitions from compute-bound to memory-bound occurs at different sequence lengths for different architectures.

Additionally, the paper does not test whether the weight permutation and interleaving techniques (offline preprocessing that rearranges weight layout for sequential access) are compatible with weight-sharing patterns like grouped-query attention (where key and value projections are shared across query heads) or with weight-tied embeddings. These patterns create constraints on weight layout that may conflict with T-MAC's permuted tile ordering.

What evidence exists in the paper. The paper tests six shapes, all from two model sizes of the same architecture family. There is no sensitivity analysis varying the $K$ -to- $M$ ratio while holding other parameters constant — for example, testing a synthetic matrix with $K=M$ vs. $K \ll M$ vs. $K \gg M$ to characterize the performance envelope. The AutoTVM tuning (Section 4) could in principle adapt to different shapes, but the paper does not demonstrate this adaptation or report the variance in optimal tile configurations across the six tested shapes. The weight permutation is described as a general technique (Section 3.2), but it is only validated on the permuted Llama-2 weights — there is no demonstration that the permutation correctly handles weight matrices with non-standard strides or padding.

Mitigation status. Not addressed. The paper presents the LUT-based approach as architecture-agnostic and the TVM code generation as handling "varying shapes" (Section 4), but the evaluation does not test this claim. The limitation is not acknowledged, and there is no discussion of which architectural properties are necessary or sufficient for T-MAC to be beneficial. A practitioner with a non-LLaMA model does not know whether T-MAC will help or whether they are operating in an untested regime.

6.4 The Prefill Stage Is Only Evaluated at Sequence Length 256, Leaving Long-Context Behavior Uncharacterized

The assumption or constraint. The mpGEMM evaluation (Figure 7) tests matrix-matrix multiplication with sequence length $N=256$ , corresponding to a prefill stage with 256 input tokens. This is substantially shorter than the context lengths that modern LLM deployments handle — production systems routinely process prompts of 2K, 4K, 8K, or even 32K tokens (e.g., Llama-2-7B supports up to 4096 tokens, and newer models support up to 128K). The paper's Section 2.1 correctly identifies the prefill stage as "compute-intensive matrix-matrix multiplications" and the decode stage as "memory-intensive matrix-vector multiplications," but it evaluates the compute-intensive regime only at a modest sequence length.

The consequence. The GEMM speedups reported in Figure 7 — maximum 5.3× on Raspberry Pi, Orin, and Surface Book 3 for 2-bit (Section 5.2) — may not hold at long sequence lengths. There are two competing effects as $N$ increases: (1) The GEMM becomes more compute-bound (the $N \times K$ activation tile grows, increasing arithmetic intensity), which could give T-MAC more opportunity to amortize LUT construction over a larger $N$ dimension. This would improve T-MAC's relative performance, since LUTs are built once per $K$ tile and reused across all $N$ rows. (2) The LUT storage requirement grows with $N_{tn}$ (the $N$ tile size), potentially causing register spilling if the tiling configuration is not adjusted. The paper's LUT-centric data layout emphasizes reusing LUTs across $M$ (weight rows) but is largely silent on how $N$ scaling affects register pressure, since each row in the $N$ dimension needs its own set of LUTs (or the LUTs must be rebuilt for each $N$ tile). The paper notes that "a larger tile size $M_{tm}$ on $M$ can lead to better lookup table reusing" (Section 3.2) but does not discuss the $N_{tn}$ tradeoff.

At very long sequence lengths (e.g., 32K), the prefill stage becomes the dominant latency component (it processes all input tokens at once, while decode processes one token per step). If T-MAC's GEMM advantage diminishes at long $N$ , the end-to-end speedup on long-context workloads could be much smaller than the token-generation speedups emphasized in the paper. Conversely, if T-MAC's GEMM advantage increases at long $N$ , the paper understates its value for long-context applications.

What evidence exists in the paper. The GEMM evaluation (Figure 7) tests only $N=256$ . There is no sweep over sequence length to characterize the scaling behavior. The paper does not report how $N_{tn}$ interacts with LUT storage or register pressure, and the AutoTVM tuning is performed per-shape, so the tuned tile configurations for $N=256$ may differ from those for $N=2048$ or $N=4096$ . The end-to-end throughput evaluation (Figure 8) measures only token generation throughput (decode stage), not prompt processing time — the measurement protocol "repeatedly generate 64 tokens for 20 iterations" (Section 5.1), which exercises the decode stage predominantly.

Mitigation status. Not addressed. The paper does not discuss prefill scaling behavior, does not sweep sequence length in the GEMM benchmarks, and does not report end-to-end latency including prompt processing. This is a significant gap for practitioners whose workloads involve long prompts (document summarization, few-shot prompting with many examples, retrieval-augmented generation with large contexts), since they cannot estimate the prompt processing overhead of T-MAC relative to llama.cpp from the reported results.

6.5 Energy Efficiency Is Demonstrated Only on a Single High-End Device, Limiting Generalizability to Battery-Constrained Edge Hardware

The assumption or constraint. The power and energy consumption evaluation (Figure 9, Table 5) is conducted exclusively on the Apple M2 Ultra — "a high-performance device" (Section 5.1) with 16 performance cores, an active cooling system, and a power envelope of 15-30W for the workloads tested. The paper identifies energy efficiency as "particularly critical for battery-operated edge devices such as smartphones and robotics" (Section 2.1), but does not measure power on any battery-operated device. The Raspberry Pi 5, Surface Book 3 (when running on battery), and OnePlus 12 are all battery-capable platforms that are evaluated for throughput but not for energy.

The consequence. The energy efficiency gains reported — 61.2% reduction for Llama-2-7B-2bit (Section 5.4) — may not translate directly to battery-operated devices for several reasons. First, the M2 Ultra has a high baseline power consumption (15W for llama.cpp CPU inference, Table 5), meaning a large fraction of the power goes to non-compute components (memory subsystem, interconnect, background processes). On a smartphone SoC where the baseline power might be 1-3W, the proportional saving from eliminating multiply-accumulate operations could be different — either larger (if compute is a bigger fraction of total power) or smaller (if memory access dominates and T-MAC's memory traffic reduction is less significant at lower total power). Second, the M2 Ultra's power management (DVFS, core boosting) differs from mobile SoCs that aggressively throttle under thermal constraints — T-MAC's power reduction might change the thermal throttling behavior, indirectly affecting sustained throughput in ways not captured by a steady-state power measurement. Third, the paper does not measure energy on the GPU and NPU comparisons in Table 7 (Surface Laptop 7, OnePlus 12, Jetson Orin NX), so the energy-efficiency advantage of T-MAC CPU over these accelerators is unquantified despite throughput comparisons being provided. The Table 5 result on Jetson AGX Orin (0.66 J/token for T-MAC CPU vs. 1.54 J/token for llama.cpp GPU) is the only non-M2-Ultra energy data point, and the Jetson Orin is a development kit with active cooling, not a battery-operated end-user device.

What evidence exists in the paper. Energy measurements exist for M2-Ultra (Figure 9, three models) and Jetson AGX Orin (Table 5, one model). No energy data exists for Raspberry Pi 5, Surface Book 3, OnePlus 12, or Jetson Orin NX. The measurement methodology uses OS-level power reporting (powermetrics on macOS) that aggregates total package power — it does not isolate CPU core power from memory controller or interconnect power, making it difficult to attribute the savings specifically to the elimination of multiply-accumulate operations versus reduced memory traffic.

Mitigation status. Partially addressed. The paper includes energy measurements where it can (M2-Ultra has built-in power monitoring), but does not extend this to the battery-operated platforms that are most relevant to the energy-efficiency motivation. The paper does not acknowledge this as a limitation, and the abstract's claim of "70% reduction in energy consumption" is stated without qualification about which devices this applies to. The 70% figure itself is an overstatement relative to the measured 61.2% (Section 5.4, discussed in the Experimental Analysis assessment), further weakening the generalizability of the energy claim.

6.6 The Revision Model (Inference System) Has No Mechanism to Dynamically Adapt the LUT Strategy Per-Layer or Per-Token

The assumption or constraint. T-MAC applies a uniform computation strategy to all layers of the model: every mpGEMM and mpGEMV operation uses the same group size ( $g=4$ ), the same bit-serial decomposition, and the same LUT construction-lookup paradigm. The tiling parameters ( $N_{tn}, M_{tm}, K_{tk}$ ) are auto-tuned per matrix shape, but the fundamental decision of whether to use LUT at all is not made per-layer. The paper's evaluation shows that T-MAC's speedup is largest at low bit-widths (2-bit, 1-bit) and more modest at 4-bit (Section 5.3 shows 1.1× improvement for Llama-2-7B-4bit multi-threaded on M2-Ultra vs. 2.3× for the 2-bit model). In a mixed-precision deployment where some layers are quantized to 4-bit and others to 2-bit (a common real-world configuration to balance accuracy and efficiency), T-MAC provides no guidance on whether the 4-bit layers should use T-MAC or fall back to the dequantization baseline.

The consequence. The absence of per-layer strategy selection means T-MAC may be applied to layers where its advantage is negligible — or even negative — while still incurring the engineering cost of integration. The paper shows that at 4-bit multi-threaded on M2-Ultra, the end-to-end speedup is only 1.1× for Llama-2-7B-4bit (Figure 8), which is essentially break-even when accounting for measurement variance. For a mixed-precision model with mostly 4-bit layers and a few 2-bit bottleneck layers, the aggregate speedup could be dominated by the 4-bit layers where T-MAC helps least. The paper's architecture does not include a mechanism to route layers to different backends (T-MAC vs. llama.cpp dequantization) based on bit-width or layer characteristics, meaning the integration is all-or-nothing: either all mpGEMM operations use T-MAC, or none do.

More fundamentally, T-MAC's performance is sensitive to the tiling parameters found by AutoTVM, and these parameters are shape-specific. If a model has many layers with subtly different matrix shapes (e.g., the first and last layers of Llama have different dimensions from the intermediate layers), each shape requires its own tuned kernel. The paper's TVM-based code generation handles this (each shape gets a generated kernel), but it means the kernel binary size grows with shape diversity — potentially problematic for edge devices with limited storage. The paper does not report kernel binary sizes or the number of distinct kernels generated for each model.

What evidence exists in the paper. The end-to-end evaluation (Figure 8) integrates T-MAC uniformly for each model — all layers use T-MAC. The per-bit-width speedups vary dramatically (from 1.1× to 6.7×), but the paper does not break down which layers benefit most or whether some layers see no benefit. The optimization breakdown (Figure 10) is aggregated across all tested shapes, masking per-shape variance. The paper acknowledges that "the speedup is less pronounced" in multi-threading "due to memory constraints and operators other than mpGEMV/mpGEMM" (Section 5.3), but does not quantify the fraction of total inference time spent in mpGEMM vs. other operations (attention, layer norm, residual connections) or discuss whether this fraction varies across layers.

Mitigation status. Not addressed. The paper presents T-MAC as a unified drop-in replacement but does not discuss the possibility of hybrid backends (T-MAC for low-bit layers, dequantization for high-bit layers) or per-layer strategy selection. The AutoTVM tuning selects tile sizes per shape but does not make the binary decision of LUT vs. dequantization. This is not a flaw in the method per se — uniformity is part of the "unified and scalable" design — but it means that practitioners with mixed-precision models cannot selectively apply T-MAC where it helps most, and must accept the speedup (or lack thereof) that the all-or-nothing integration provides.

7. Implications and Future Directions

How This Work Changes the Landscape

T-MAC shifts the discourse around edge LLM deployment from hardware-centric acceleration to computation-paradigm co-design. The dominant narrative in the field has been that low-bit LLM inference requires specialized hardware — GPUs, NPUs, or custom accelerators with dedicated low-precision multiply-accumulate units — and that CPUs are a fallback platform at best. T-MAC challenges this narrative at its foundation: the bottleneck is not the hardware's peak FLOPs, but the mismatch between the computation paradigm (dequantize-then-multiply) and the hardware's native strengths. By replacing multiplication with register-based table lookup, T-MAC realigns the computation with CPU architectural features — flexible register files, hardware byte-permutation instructions, sophisticated cache hierarchies — that GPUs and NPUs do not possess in the same form.

This is properly understood as a paradigm shift in how we think about mixed-precision computation, not merely an optimization of existing approaches. Prior work treated dequantization as an unavoidable tax — the cost of bridging the gap between low-bit storage formats and high-precision execution units. T-MAC demonstrates that this tax is an artifact of the computation paradigm, not a fundamental constraint. The bit-serial decomposition $A \times W = \sum 2^i A \times W_i$ is mathematically trivial but conceptually transformative: it reframes mixed-precision GEMM not as "compute at the higher precision" but as "compute at each bit-plane uniformly using table lookup." The implication is that the dequantization overhead that has plagued low-bit inference — causing 3-bit models to run slower than 4-bit on llama.cpp (Figure 6) — is not something to be optimized away through better unpacking kernels, but something to be eliminated entirely through a different computation primitive.

The paper also resolves a specific and puzzling contradiction in prior work. LUT-based computation had been explored on GPUs — LUT-GEMM (Park et al., 2023) and related efforts — but consistently underperformed dequantization-based kernels, with LUT-GEMM being 1.75–2.34× slower than BitBLAS on A100 (Section 2.4). This created an apparent paradox: the LUT approach eliminated multiplications and should theoretically be faster, yet it was slower in practice. T-MAC resolves this by showing that the failure was platform-specific, not algorithmic: GPUs lack the per-thread, low-latency storage that LUT-based computation requires (their shared memory is a contended resource with bank conflicts and higher latency than register access), while CPUs possess exactly the right architectural features (single-cycle register access, hardware table-lookup instructions with comparable throughput to SIMD arithmetic). The contradiction is resolved by recognizing that the LUT paradigm shifts the hardware fitness landscape — away from GPUs and toward CPUs. This is a diagnostic contribution with practical consequences: it tells the community not to abandon LUT-based methods, but to target them at the right hardware substrate.

The paper reframes which research directions become more attractive and which become less so:

More attractive: Co-design of quantization formats and LUT-based computation (where the quantization scheme is chosen to maximize LUT reuse or minimize bit-planes); hardware accelerator designs that replace multiply-accumulate units with table-lookup units (as the paper explicitly calls for in Section 7); hybrid systems that route different layers to different backends (LUT for low-bit layers, dequantization for high-bit layers) based on per-layer speedup characteristics; energy-aware deployment optimization where the choice of computation paradigm is driven by Joules/token rather than raw tokens/second.

Less attractive: Incremental optimization of dequantization unpacking kernels for non-power-of-two bit-widths (3-bit, 5-bit, 6-bit). T-MAC's results suggest these are fundamentally fighting an uphill battle — no matter how optimized the SHIFT-and-AND sequence for 3-bit unpacking, it will always add overhead that LUT-based computation avoids entirely. Research effort is better spent improving LUT-based methods (e.g., better table compression, larger group sizes enabled by wider SIMD registers) than perfecting dequantization for awkward bit-widths. Similarly, GPU-focused LUT methods that attempt to fit tables into shared memory appear less promising given T-MAC's demonstration that register-resident LUTs are the key enabler — and GPUs' register files are architected for SIMD operand storage, not random-access table lookup.

The paper also establishes a new performance reference point that resets expectations for edge CPU inference. Before T-MAC, the practical ceiling for CPU-based LLM token generation was defined by llama.cpp's heavily optimized dequantization kernels — perhaps 5-7 tokens/s on a Raspberry Pi-class device for a 7B model. T-MAC's demonstration of 11 tokens/s on Raspberry Pi 5 for BitNet-3B, and 30+ tokens/s on a single M2-Ultra core, shows that there is a large, previously unexploited performance headroom accessible through algorithmic transformation alone, without hardware changes. This reframes edge deployment from "how do we fit the model in memory?" (solved by quantization) + "how do we tolerate the slow inference?" to "how do we make inference fast enough to be interactive?" — and provides evidence that the answer is within reach on today's commodity hardware.

Follow-Up Research This Work Enables

Cheap, low-overhead difficulty estimation for per-layer strategy selection. T-MAC applies a uniform LUT-based strategy to all layers of a model, but its speedup varies dramatically with bit-width: 1.1× for 4-bit vs. 6.7× for 2-bit on Raspberry Pi 5 (Figure 8). A natural extension is per-layer or per-operation strategy selection — routing each mpGEMM operation to either T-MAC (LUT-based) or llama.cpp (dequantization-based) depending on which is faster for that specific (shape, bit-width, available memory bandwidth) combination. This requires a cheap, low-overhead performance model that predicts T-MAC speedup from easily observable features (matrix dimensions, weight bit-width, activation precision) without running the kernel. A strong follow-up would: (1) collect per-layer speedup measurements across a diverse set of model architectures (LLaMA, Gemma, Phi, Mistral, mixture-of-experts) at bit-widths 1–8; (2) train a lightweight regressor or decision tree on (dim_K, dim_M, dim_N, bits, memory_bandwidth) → speedup; (3) integrate into llama.cpp as a compile-time router that generates hybrid backends automatically; (4) evaluate on mixed-precision models (e.g., 4-bit attention layers with 2-bit FFN layers) to quantify end-to-end speedup over all-T-MAC and all-dequantization baselines. A negative result — finding that hybrid routing adds more dispatch overhead than it saves — would be equally valuable, as it would confirm that T-MAC's uniformity is a feature, not a limitation.

Scalable table lookup with multi-register LUTs for AVX-512 and SVE. T-MAC uses $g=4$ as the sweet spot because $2^4 = 16$ entries fit in one 128-bit register (ARM NEON) or one 128-bit lane (AVX2). But AVX-512 (512-bit registers) and ARM SVE (scalable vector length up to 2048 bits) can hold much larger tables. For $g=5$ (32 entries), a single 256-bit register suffices; for $g=6$ (64 entries), 512 bits; for $g=8$ (256 entries) with table quantization to INT8, a 2048-bit SVE register. Larger $g$ reduces the number of lookups along the $K$ dimension by a factor of $g/4$ , directly reducing instruction count and LUT construction overhead. A concrete follow-up would: (1) implement T-MAC for AVX-512 using _mm512_permutexvar_epi8 (which can permute across 512-bit lanes, unlike AVX2's lane-constrained shuffle) with $g=5$ and $g=6$ ; (2) implement for ARM SVE with $g=5$ through $g=8$ , leveraging the scalable register width; (3) measure the tradeoff between group size (fewer lookups but larger LUT construction cost, scaling as $2^g$ ) and end-to-end speedup across the Llama-2 matrix shapes from Figure 6; (4) determine whether the optimal $g$ changes with bit-width (lower bit-widths favor larger $g$ because LUT construction cost is amortized over fewer bit-planes). The key hypothesis to test is whether wider SIMD enables a super-linear speedup — where doubling $g$ more than halves the number of lookups because the LUT build cost grows as $2^g$ but is amortized over $M_{tm}$ weight rows and $n$ bit-planes.

Systematic characterization of the LUT construction overhead across matrix shapes. The paper provides aggregate speedup numbers but does not decompose kernel time into LUT construction (gather, swizzle, table quantization) vs. table lookup and accumulation. This decomposition is critical for understanding when T-MAC will underperform dequantization — specifically, for matrix shapes with small $M$ (few weight rows per LUT, poor amortization) or large $K$ with small $g$ (many LUTs to build). A strong follow-up would: (1) instrument the generated kernels with hardware performance counters (instructions retired, cache misses, pipeline stalls) on both ARM and x86; (2) sweep over synthetic matrices with independently varying $K$ , $M$ , $N$ , $g$ , and bit-width to build a roofline-style model for LUT-based GEMM analogous to the roofline model for traditional GEMM; (3) identify the "LUT roof" — the point where LUT construction becomes the bottleneck regardless of memory bandwidth — and the "lookup roof" — where lookup throughput saturates; (4) validate against real model matrices from diverse architectures (Phi, Gemma, Mistral, vision transformers, diffusion models) to test whether the synthetic sweep predicts real-world speedups. A negative result — finding that shape diversity in real models keeps T-MAC in favorable regimes — would strengthen the paper's generality claim. A positive result — identifying specific shapes where T-MAC regresses — would provide actionable guidance for hybrid backend design.

LUT-based attention: extending table lookup beyond linear layers. The paper focuses exclusively on the linear projections (GEMM/GEMV) that dominate the feed-forward and projection layers of transformers. However, the attention mechanism — specifically, the query-key dot product and the attention-softmax-value computation — also involves matrix multiplications that could benefit from LUT-based computation, particularly for low-bit quantized key and value caches. Quantized KV-cache schemes are increasingly important for long-context inference, and they face the same dequantization overhead that T-MAC eliminates for weights. A concrete follow-up would: (1) extend T-MAC to handle the batched GEMV operations in attention (where the "weight" is the quantized key cache and the "activation" is the query vector); (2) handle the attention-softmax-value sequence, where the attention weights (output of softmax) multiply the value cache — this is another mixed-precision scenario if the value cache is quantized; (3) evaluate on long-context benchmarks (e.g., needle-in-haystack retrieval with 32K context) where KV-cache dequantization dominates decode latency; (4) compare against dedicated KV-cache quantization systems (e.g., KIVI, Atom) to determine whether LUT-based attention provides complementary or redundant speedups. The key metric is not just kernel speedup but time-to-first-token for long prompts, where attention dominates the prefill stage and the existing T-MAC GEMM speedups (Figure 7, sequence length 256) may not capture the bottleneck.

Quality preservation for 2-bit and 1-bit models under LUT-based inference. The paper's most dramatic speedups are at 2-bit and 1-bit (up to 11.2× kernel speedup and 6.7× end-to-end), but model quality is only validated at 4-bit (Table 4). This is the single largest gap between the paper's performance claims and the evidence supporting deployability. A critical follow-up would systematically evaluate: (1) measure perplexity (WikiText-2, C4) and downstream accuracy (MMLU, HellaSwag, WinoGrande) for BitDistiller-2bit, OneBit-1bit, and BitNet-b1.58 models under T-MAC inference vs. dequantization-based inference; (2) at each bit-width, measure the Pearson correlation between T-MAC's output logits and the dequantization baseline's logits on a held-out calibration set to quantify distributional shift; (3) ablate table quantization granularity (per-8, per-16, per-32 values) and measure the accuracy-quantization_error tradeoff; (4) test whether the bit-serial linear transformation ( $s_0=-1, s_1=1$ ) interacts adversely with specific quantization schemes — for GPTQ (which uses per-channel scales), BitDistiller (which uses QAT with self-distillation), and OneBit (which uses a learned codebook) — by measuring per-layer output error. A negative result — finding that T-MAC degrades 2-bit or 1-bit model quality by more than 1% on key benchmarks — would severely limit the practical impact of the paper's headline speedup claims, because the speedups would only apply at bit-widths where model quality is already unacceptable. A positive result — demonstrating quality parity at all bit-widths — would close the paper's most important evaluation gap and make the "practical deployment" claim fully substantiated.

Energy-optimal mixed-precision deployment on battery-operated devices. The paper's energy measurements (Figure 9, Table 5) are limited to the plug-powered M2-Ultra and the actively-cooled Jetson AGX Orin developer kit. The strongest argument for T-MAC is energy efficiency on battery-constrained devices, so extending the energy evaluation to representative battery-operated platforms is essential. A concrete follow-up would: (1) measure power and energy on a smartphone (e.g., Google Pixel 8 or iPhone 15 running a port of llama.cpp with T-MAC integration) using the device's onboard fuel gauge or an external power monitor; (2) evaluate both continuous generation (as in the paper) and bursty, interactive workloads (single prompt, single response, device returns to idle) to capture the impact of DVFS ramp-up/down and thermal throttling; (3) compare T-MAC CPU vs. the device's GPU (via llama.cpp OpenCL or MPS) and NPU (via Qualcomm QNN or Apple CoreML) for the same model and bit-width, measuring tokens/Joule as the primary metric; (4) determine the energy-optimal bit-width — the bit-width that minimizes Joules per token for a given model quality threshold — which may differ from the throughput-optimal bit-width due to the interplay of LUT construction energy, memory traffic energy, and compute energy. This experiment would test whether the 61.2% energy reduction measured on M2-Ultra generalizes to the devices where battery life actually constrains deployment, and whether the energy advantage over NPUs (which are marketed as power-efficient AI engines) is as dramatic as the throughput advantage in Table 7 suggests.

Practical Applications and Downstream Use Cases

Battery-operated mobile assistants with always-on LLM capabilities. The paper's demonstration of 11 tokens/s on Raspberry Pi 5 for BitNet-3B (Figure 8) and 16.62 tokens/s on OnePlus 12 for Llama-2-7B-2bit (Table 7) — using only 3-4 CPU cores — enables a deployment scenario where a mid-range smartphone runs a capable LLM entirely on-device for interactive applications (voice assistants, text composition, real-time translation) without relying on cloud offloading. The specific advantage is energy-proportional inference: T-MAC's 51-61% energy reduction on 2-bit models (Figure 9, Table 5) means a phone that could previously sustain 2 hours of continuous LLM interaction can now sustain 4-5 hours on the same battery. The 9.7× speedup over the Adreno GPU on OnePlus 12 for 2-bit models (Table 7) also means the CPU-based approach avoids waking the power-hungry GPU at all, further reducing system power. This enables "always-listening" LLM features that were previously infeasible due to battery drain concerns.

Retro-fitting existing edge hardware for production LLM inference. Many deployed edge systems — point-of-sale terminals, industrial controllers, medical devices, in-vehicle infotainment — contain capable CPUs (ARM Cortex-A7x series, Intel Core i5/i7) but no GPUs or NPUs suitable for ML inference. These systems were designed before the LLM era and cannot be hardware-upgraded, but their CPUs can run T-MAC. The paper's results on Surface Book 3 (Intel Core i5-1035G7, 4 cores) — achieving 31.83 tokens/s for Llama-2-7B-2bit (Table 7) — demonstrate that a 2020-era laptop CPU can deliver interactive-rate inference for a 7B-parameter model. The specific deployment advantage is hardware reuse: organizations can deploy LLM-powered features (document understanding, code completion, natural language interfaces) to existing hardware fleets without a hardware refresh cycle, dramatically reducing the total cost of LLM adoption.

Privacy-preserving on-device inference for sensitive data processing. In domains where data cannot leave the device — healthcare (patient records), legal (confidential documents), finance (trading strategies), personal messaging — on-device LLM inference is the only compliant deployment model. T-MAC's CPU-based approach strengthens the privacy argument in two ways beyond just enabling local execution. First, it keeps inference on the CPU, avoiding the GPU driver stack and NPU firmware, which are typically closed-source and have larger attack surfaces. Second, the energy reduction means the inference workload can run in the background (e.g., indexing messages, summarizing documents) without visibly impacting battery life, making user-transparent privacy-preserving processing viable. The paper's 2.3× energy efficiency advantage over llama.cpp on GPU (Table 5) translates to more than twice the processing volume within the same thermal and battery envelope.

Cost-efficient batch inference for dataset generation and model evaluation. Organizations that use LLMs for offline dataset generation (e.g., creating training data for smaller models, evaluating model outputs on large benchmarks) typically run inference on GPU servers due to throughput requirements. T-MAC's strong GEMM performance (up to 5.3× speedup for 2-bit at sequence length 256; Figure 7) combined with the lower cost-per-hour of CPU instances vs. GPU instances in cloud computing suggests a cost-optimized deployment: use CPU instances with T-MAC for batch inference workloads where the model is 2-bit quantized and the primary metric is tokens-per-dollar rather than tokens-per-second. A simple cost model from the paper's numbers: on Jetson AGX Orin, T-MAC CPU achieves 78% of GPU throughput at 34% of the power (Table 5). In a cloud setting where CPU instances cost roughly 20-30% of GPU instances per hour, T-MAC could deliver comparable or lower cost-per-token for batch workloads, particularly when the jobs are not latency-sensitive and can run overnight on cheap CPU spot instances.

When to Prefer This Method

The paper positions T-MAC primarily against dequantization-based CPU inference (llama.cpp) and implicitly against GPU/NPU inference on the same edge devices. The decision conditions are:

Prefer T-MAC over llama.cpp on CPU when: The weight bit-width is 2-bit or below (where llama.cpp's dequantization overhead is most severe and T-MAC achieves 4-11× kernel speedups); the deployment target is single-threaded or has few cores (Raspberry Pi 5, smartphone efficiency cores), where LUT-based computation's instruction-level advantages are not masked by memory bandwidth saturation; the model architecture uses Llama-like matrix aspect ratios (the only shapes evaluated) — unknown behavior on significantly different geometries; model quality at the target bit-width has been validated under T-MAC inference (only demonstrated at 4-bit; proceed with caution at 2-bit and 1-bit); and energy efficiency matters — T-MAC reduces energy per token substantially (51-61% for 2-bit/BitNet on M2-Ultra, 68% on Jetson AGX Orin vs. GPU).
Prefer llama.cpp (dequantization) on CPU when: The weight bit-width is 4-bit or higher and the deployment uses many cores on a high-bandwidth device (M2-Ultra), where T-MAC's multi-threaded speedup is only 1.1× for 4-bit models (Figure 8) — the marginal gain may not justify the integration effort; the model uses 8-bit weights (not evaluated), where T-MAC would require 8 bit-planes of table lookups and might become slower than dequantization due to the serial lookup overhead; or the application cannot tolerate any numerical deviation from the reference dequantization-based output, even the imperceptible NMSE difference shown in Table 3 (e.g., for reproducible scientific benchmarks).
Prefer T-MAC on CPU over llama.cpp on GPU when: The device has unified memory (CPU and GPU share DRAM bandwidth — Jetson AGX Orin, Surface Laptop 7, OnePlus 12); the model is 2-bit, where T-MAC CPU delivers 1.4× (Orin NX) to 3× (Surface Laptop 7) higher throughput than the GPU (Table 7); energy efficiency is the primary metric — T-MAC CPU achieves 2.3× better Joules/token than GPU on Orin (Table 5); or the GPU backend uses OpenCL rather than CUDA (Qualcomm Adreno on OnePlus 12), where llama.cpp GPU throughput is very low (1.72 tokens/s for 2-bit) and T-MAC CPU provides a 9.7× improvement (Table 7).
Prefer llama.cpp on GPU over T-MAC when: The device has a dedicated high-bandwidth GPU with strong CUDA support (discrete NVIDIA GPU on a laptop or desktop), where the GPU's memory bandwidth advantage may overcome T-MAC's computational efficiency — the paper does not test this configuration; the model is 4-bit and throughput, not energy, is the sole metric — on M2-Ultra the GPU advantage is present but modest (Figure 11 shows GPU wins at 4-bit for larger shapes); or the workload involves large-batch prefill (many sequences processed simultaneously), where GPU parallelism dominates the LUT-based approach's per-instruction efficiency.
Prefer T-MAC on CPU over NPU when: The NPU does not natively support 2-bit inference (as implied by the projected performance in Table 7, where 2-bit NPU numbers are marked with '*'), meaning it falls back to dequantization; the CPU is a modern ARM core (Cortex-X4, A720) with strong NEON throughput, as on the OnePlus 12 where T-MAC achieves 1.5× higher throughput than the Hexagon NPU for 2-bit (Table 7); or the deployment stack prefers standard CPU programming models over vendor-specific NPU SDKs (Qualcomm QNN, Apple CoreML) for portability and maintenance reasons — T-MAC's generated C++ code with no external dependencies matches the deployment philosophy that made llama.cpp (plain C/C++, no dependencies) the standard for edge CPU inference.

Dependency tree

Papers this one builds on, walking backwards through the same domain and shared topics.

The Era of 1-bit LLMs: All Large Language Models are in 1.58 Bits2024· pretraining← LLM-judged similar
- Binary and Ternary Natural Language Generation2023· low-precision← LLM-judged similar
- QMoE: Practical Sub-1-Bit Compression of Trillion-Parameter Models2023· moe← LLM-judged similar
- BitNet: Scaling 1-bit Transformers for Large Language Models2023· pretraining← LLM-judged similar
- Nemotron-4 15B Technical Report2024· pretraining← in-domain predecessor
InfiniGen: Efficient Generative Inference of Large Language Models with Dynamic KV Cache Management2024· inference-optimization← in-domain predecessor
- A Simple and Effective $L_2$ Norm-Based Strategy for KV Cache Compression2024· inference-optimization← in-domain predecessor
- READ MORE2023· inference-optimization← similar topic
- Prompt Cache: Modular Attention Reuse for Low-Latency Inference2023· inference-optimization← similar topic

What it influenced

Later papers in the same chain. Direction reversed.

MInference 1.0: Accelerating Pre-filling for Long-Context LLMs via Dynamic Sparse Attention2024· inference-optimization→ in-domain predecessor
- Native Sparse Attention: Hardware-Aligned and Natively Trainable Sparse Attention2025· architecture→ LLM-judged similar
- Inference Performance Optimization for Large Language Models on CPUs2024· inference-optimization→ in-domain predecessor
- SeerAttention: Learning Intrinsic Sparse Attention in Your LLMs2024· inference-optimization→ similar topic
- Kascade: A Practical Sparse Attention Method for Long-Context LLM Inference2025· inference-optimization→ similar topic
- LycheeDecode: Accelerating Long-Context LLM Inference via Hybrid-Head Sparse Decoding2026· inference-optimization→ similar topic
Can 1B LLM Surpass 405B LLM? Rethinking Compute-Optimal Test-Time Scaling2025· inference-optimization→ similar topic
- TransMLA: Multi-Head Latent Attention Is All You Need2025· inference-optimization→ in-domain predecessor
- DataDecide: How to Predict Best Pretraining Data with Small Experiments2025· data→ topic thread
- Scaling LLM Test-Time Compute Optimally can be More Effective than Scaling Model Parameters2024· inference-optimization→ similar topic