A few months back a client came to me with a server farm full of RTX 3080s and a real headache. LLM inference was crawling and there was no budget for H100s. We sat down with Software FP8 performance benchmarks, and the gap between his Ampere cards and the newer Ada Lovelace hardware was ugly. Textbook memory bandwidth bottleneck. The cores were plenty fast, they were just idle, waiting for data to crawl across the bus.
I will admit my first guess was wrong. I figured a simple FP16 cast would do it and spent a long weekend optimizing the weights with standard PyTorch quantization. It saved some VRAM and the speedup was negligible, because I was still loading the same number of elements per cycle. The problem was never precision, it was data movement. Fixing slow training or inference on older cards means changing how the data is packed.
How bit-packing boosts software FP8 performance
The point is not lower precision on its own. It is getting the GPU to load four values in the time it normally spends loading one. The RTX 30-series has no FP8 support in hardware, so we use bitwise operators to pack four FP8 values into a single FP32 container. FlashAttention does something comparable with tiling to stay in SRAM. Here the goal is compression, cutting the traffic between HBM and the compute cores.
You pay a little math on the unpacking side and get a lot of bandwidth back. For memory-bound work like GEMV, take that trade every time. Here is a simplified version of the packing utility I wrote for it:
/**
* bbioon_pack_fp16_to_fp32
* Packs two FP16 values into one FP32 container to optimize bandwidth.
*/
def bbioon_pack_fp16_to_fp32(val_a, val_b):
# Cast to 16-bit unsigned integers
u16_a = val_a.astype(np.uint16).astype(np.uint32)
u16_b = val_b.astype(np.uint16).astype(np.uint32)
# Shift and combine using bitwise OR
packed = (u16_a << 16) | u16_b
return packedTriton kernels instead of raw CUDA
Writing raw CUDA for this is miserable. One bad pointer and you spend days chasing memory corruption. So I moved the project to Triton, which lets you write GPU kernels in Python at close to native speed. The kernel loads the packed FP32, unpacks it into FP8 right away, then upcasts to FP32 for the accumulation. That last step is what keeps low-precision math from overflowing.
On the final benchmarks, software FP8 performance on the client’s 3080s came in at nearly 3.3x. The memory savings were beside the point. The hardware was finally being used. It is the same idea as optimizing frontend animations: get the friction out from between the data and the processor.
@triton.jit
def bbioon_fp8_gemv_kernel(matrix_ptr, vector_ptr, out_ptr):
row_id = tl.program_id(0)
accumulator = 0.0
# Load 4 values at once in a single FP32 slot
packed_data = tl.load(matrix_ptr + row_id)
# Unpack logic (simplified)
# v1, v2, v3, v4 = bbioon_unpack_fp8(packed_data)
# Accumulate with upcasting to maintain precision
# accumulator += (v1 * vec_val) ...
tl.store(out_ptr + row_id, accumulator)Where software FP8 performance stops helping
This is not a magic bullet. If your workload is compute-bound, meaning the math is the bottleneck rather than the loading, you will not see any of this. For inference and large matrix-vector products it saves the day. The NVIDIA FP8 Whitepaper is clear that native support wins, but if you are stuck a generation behind, software packing is what you have.
This gets complicated fast. GPU memory hierarchies are finicky and one badly shaped kernel launch will tank your throughput. If you are tired of debugging somebody else’s mess and want your application performing like it is on native hardware, drop me a line. I have probably met this exact bottleneck already.
Are you getting software FP8 performance gains in production, or still stuck against the VRAM wall?
