Breaking the Hardware Barrier: Software FP8 for Older GPUs

Hardware Level Support

Like the AVX-512 instructions, which are only supported on a limited number of hardware platforms, the FP8 and FP16 instructions and registers are also limited by hardware and are only available in the latter. If you're on the RTX-30 or RTX-20 series of GPUs from Nvidia, you won't be able to take advantage of this low-precision FP8 model. This is exactly the problem He is a feather settlement efforts.

Packaging method

By using bitwise operators, one can easily pack the FP16 type into FP32. The algorithm is described below.

Packaging FP16

• Throw the input of FP32 to FP16; this step can be easily done using numpy's astype work.
• Throw them in U16 and then U32; this sets the upper 16 bits to 0 and the lower 16 bits of the actual FP16.
• Shift one of them by 16 using bitwise LSHIFT operator, and combine them both bitwise OR the operator.

Unpacking FP16

• Extract the lower 16 bits using bitwise AND operator and mask 0xFFFF.
• Extract the upper 16 bits using RSHIFT working at 16 and doing it gradually AND working with mask 0xFFFF.
• Broadcast both U16 values ​​back to FP16 and FP32 if needed.

FP8 packaging

FP8 has two widely used formats – E5M2 & E4M3. One cannot use the same algorithm used to pack two FP16s into FP32 because the CPU does not natively support FP8 types, but supports FP16 (partial precision); this is the reason that np.float8 there is no.

Casting FP16 to FP8-E5M2 is straightforward, as seen in the figure, because both have the same number of exponent bits and differ only in their component.

FP8-E5M2 packaging

• Throw the input of FP32 to FP16; this step can be easily done using numpy's astype function, or get input yourself like FP16.
• Broadcast to U16, LSHIFT at 8, then RSHIFT by 8 to divide the upper 8 bits
• Do this for every FP32 or FP16.
• Now use the LSHIFT operator, subtract them by 0, 8, 16 and 24 units and add them bitwise OR the operator.

Again, the unpacking should be straightforward; it's the exact opposite of packaging.

Packing the FP8-E4M3 is not as simple and straightforward as packing the FP16 or FP8-E5M2, due to the difference in the exponent bits.

Instead of using it from scratch, the library uses the ml_d types library, which already performs broadcast calculations.

I ml_d types the library provides support for commonly used FP8 standards, such as E5M2 and E4M3 casting, for NumPy arrays. Using the same astype function, we can do the casting as we do for FP16 types. The algorithm is exactly the same as how we package FP16, so I skip it here.

Triton GPU Kernels

After packing, we need an algorithm (kernel) to use this type of packed data and perform calculations. Passing a packed data type to a kernel implemented by FP32 or FP64 will result in an undefined calculation because we have already corrupted the FP32 or FP64 being passed. Writing a kernel that takes a packed data type as input in CUDA is not a straightforward task and is prone to errors. That's right there Triton it shines; is a Domain-specific language library that implements a custom intermediate representation of GPU kernels. In layman's terms, it allows one to write GPU scripts in Python itself without the need to write CUDA scripts in C.

Triton characters do what was said before; the algorithm is as follows:

• Load the full list into memory
• Remove the memory and upload it to the FP32 for accumulator functions
• Do the math

The reader should note that when doing computing, upcasting is used to prevent overflow. So, from an accounting point of view, there is no benefit. However, from a bandwidth perspective, we load the memory twice or four times without compromising the bandwidth.

Triton Kernel Implementation (pseudocode)

@triton.jit
def gemv_fp8_kernel(packed_matrix_ptr, packed_vector_ptr, out_ptr): 
    # Get current row to process
    row_id = get_program_id()
    
    # Initialize accumulator for dot product
    accumulator = 0
    
    # Iterate over row in blocks
    for each block in row:
        # Load packed FP32 values (each contains 4 FP8s)
        packed_matrix = load(packed_matrix_ptr)
        packed_vector = load(packed_vector_ptr)
        
        # Unpack the FP32 into 4 FP8 values
        m_a, m_b, m_c, m_d = unpack_fp8(packed_matrix)
        v_a, v_b, v_c, v_d = unpack_fp8(packed_vector)
        
        # Upcast to FP32 and compute partial dot products
        accumulator += (m_a * v_a) + (m_b * v_b) + (m_c * v_c) + (m_d * v_d)
    
    # Store final result
    store(out_ptr, accumulator)

Results

Hardware: NVIDIA GeForce RTX 3050 6GB VRAM

CUDA version: 13.0

Python version: 3.13.9

GEMV Benchmark (M = 16384, N = 16384) (MxN matrix)

A theoretical performance improvement that can be reached by 4x; 3.3x is pretty good in comparison, with the remaining overhead mainly coming from packet/unpack operations and kernel initialization costs.

The E5M2 is faster than the E4M3 due to ease of release, but the E4M3 offers better accuracy. However, it is more complicated to decode (Feather uses a different GPU kernel to decode the E4M3 format).

Flash Attention Benchmark (Sequence length = 8192, Embed size = 512)

Accuracy and Precision

Testing with random matrices (numerical distributions in the range [-3, 3] and normal normal distribution) show that both E4M3 and E5M2 keep numerical results within realistic tolerances for deep learning performance. Accumulation errors remain manageable for typical workload sizes; however, users who require strict numerical accuracy should ensure their exact use.