# Tensor cores and low precision · Kernel optimization Every kernel in the previous section topped out well short of the GPU’s advertised peak because none of them touched a **tensor core** (A hardware unit that performs a small matrix multiply-accumulate as a single instruction, instead of a sequence of individual fused multiply-adds.). On the same A6000 used above, switching cuBLAS from plain fp32 to TF32 or BF16 precision, so it can dispatch to tensor cores, raises its measured throughput by 2.5× and 3.5× respectively, entirely from using different hardware for the same multiply-accumulate. Going narrower than bf16 means an 8-bit floating-point format. The interchange format published by NVIDIA, Arm, and Intel defines two 8-bit encodings rather than one: **E4M3** (An 8-bit floating-point encoding with a 4-bit exponent and 3-bit mantissa, recommended for weight and activation tensors.), with a 4-bit exponent and 3-bit mantissa, and E5M2, with a 5-bit exponent and 2-bit mantissa. The two trade off differently: E4M3 gives up representing infinity and most NaN bit-patterns to extend its range, topping out at a maximum normal value of 448, while E5M2 keeps full IEEE 754 behavior for those special values and reaches a maximum normal value of 57,344. The paper’s own recommendation is E4M3 for weight and activation tensors and E5M2 for gradients, since gradients tend to need the wider dynamic range. > Figure. Two ways to spend eight bits. Both FP8 formats spend one bit on the sign and split the remaining seven differently. E4M3 gives up infinity and most NaN bit-patterns to push its maximum normal value to 448; E5M2 keeps full IEEE 754 special values and reaches 57,344, the wider dynamic range the paper recommends for gradients. Exponent bits are shaded, mantissa bits highlighted.Constants from the source FP8’s narrow range means a tensor has to be rescaled before it is cast down: a **scaling factor** (A per-tensor multiplier applied before casting a higher-precision value down to FP8, chosen so the tensor's largest magnitude lands close to the format's representable maximum.) is applied first so a tensor’s largest values sit near 448 or 57,344 instead of overflowing or clustering near zero, then removed again once the FP8 matrix multiply has produced a higher-precision result. NVIDIA’s TransformerEngine library implements this bookkeeping so it does not have to be hand-rolled per model: its DelayedScaling recipe wraps a forward pass in FP8 at a chosen format, E4M3 by default, and it targets Hopper, Ada, and Blackwell GPUs, the same generations whose tensor cores can execute FP8 matrix multiplies directly. Because the boundary between compute-bound and memory-bound work depends on flops per byte, halving a tensor’s size in memory does two things at once: it doubles the arithmetic a tensor core can push through in the same span of time, and it halves the bytes that have to move to feed it. That is why FP8, and the narrower MXFP8 and NVFP4 formats TransformerEngine has since added for Blackwell, keep showing up wherever a kernel is waiting on bytes rather than flops, the same trade-off the KV cache made in the previous chapter, applied to the weights and activations themselves. Feeding tensor cores at these rates also changed how data moves. The Hopper generation added the **Tensor Memory Accelerator** (A hardware unit introduced in NVIDIA's Hopper architecture that copies tiles of multi-dimensional arrays between global and shared memory asynchronously, driven by a descriptor rather than by per-thread address arithmetic.) (TMA), a dedicated unit that copies tiles of multi-dimensional arrays between global and shared memory asynchronously. A single thread issues the copy against a descriptor that carries the tensor’s shape and strides, so address computation and out-of-bounds predication happen in one place instead of consuming registers in every thread, and the copy’s asynchrony is what makes warp-specialized kernel schedules practical. TMA also emits the swizzled shared-memory layouts Hopper’s tensor-core instructions expect, layouts too intricate to be worth loading by hand. The tensor-core instruction itself went asynchronous in the same generation. Hopper’s `wgmma.mma_async` family is issued not by a thread or a warp but by a **warp group** (A group of four warps, 128 threads, that cooperatively executes Hopper's asynchronous wgmma tensor-core instructions and jointly holds the accumulator in its registers.) of four warps, because the accumulator no longer fits anywhere smaller: an m64n16k16 multiply accumulates a 64×16 tile of fp32 values, 1,024 registers, where a single thread may hold at most 256. Spread across the warp group’s 128 threads that is 8 registers each, and the instruction variants scale the N dimension from 8 all the way to 256. A worklog in the same spirit as the A6000 one above shows what those two units are worth on an H100. The A6000-style kernel, ported to bf16, reaches 32 TFLOP/s where cuBLAS reaches 716, because on Hopper the tensor cores are not optional. Rewriting the inner loop around TMA loads and wgmma lifts it to 317 TFLOP/s, larger 128×128 tiles to 423, and splitting each block into a producer warp group that issues TMA loads into a circular buffer of shared-memory tiles and a consumer warp group that drains it through the tensor cores, the pattern called warp specialization, to 498. The finished kernel, several optimizations later, outruns cuBLAS by 7% at N=4096. That machinery is no longer exotic: DeepSeek’s open-source DeepGEMM is a tensor-core kernel library for exactly these SM90 and SM100 GPUs, covering FP8, FP4, and BF16 GEMMs plus fused mixture-of-experts kernels, compiled at runtime by a lightweight JIT module so installation needs no CUDA compilation. Its README describes borrowing concepts from CUTLASS and CuTe while avoiding heavy reliance on their templates, keeping a small set of core kernel functions readable as a learning resource, and reaching up to 1550 TFLOP/s on an H800 while matching or exceeding expert-tuned libraries across matrix shapes. Source E4M3/E5M2 bit layout, exponent bias, and recommended usage from [FP8 Formats for Deep Learning](https://arxiv.org/abs/2209.05433). TransformerEngine’s API and supported GPU generations from the [TransformerEngine](https://github.com/NVIDIA/TransformerEngine) README. Tensor core speedup figures from the same A6000 benchmark as the matrix multiplication section. TMA’s design and its single-threaded descriptor-driven copies from Colfax’s [CUTLASS tutorial on TMA](https://research.colfax-intl.com/tutorial-hopper-tma/). The wgmma register arithmetic and the H100 kernel progression from [Outperforming cuBLAS on H100: a Worklog](https://cudaforfun.substack.com/p/outperforming-cublas-on-h100-a-worklog). DeepGEMM’s scope and performance from the [DeepGEMM](https://github.com/deepseek-ai/DeepGEMM) README. More formats and hardware detail in [the reading list](https://learn-kernels.com/chapters/reading#tensor-cores-low-precision).