Vizuara AI Labs

4-Bit Training, Tested on a Real B200

Thinking Machines trained a large model with a 4-bit forward pass. We explain every idea in their article from the ground up, then rebuild the recipe with real 4-bit multiplies on a rented Blackwell GPU and train a GPT with it.

By Dr. Raj Dandekar and the team at Vizuara AI Labs · 15 November 2026 · about 40 minutes

On 13 November 2026, Jack Cook and his colleagues at Thinking Machines published End-to-End 4-Bit Training. It describes how their model, Inkling 2 Small, was trained with most of its heavy multiplications done on numbers that are only four bits wide. A normal training run today uses sixteen bits for the same numbers. They report that the recipe made their largest pretraining jobs 8% faster and their reinforcement learning jobs 20% faster than an 8-bit run, at a cost of about 2% in how efficiently the model learns.

The article is short and dense. It assumes that the reader already knows what a block scale is, why a transposed matrix is a problem, and what stochastic rounding does to a gradient. This write-up assumes none of that, and it does not take the result on trust. The first half builds every idea from a single rounded number up to the full layer. The second half is an experiment on one NVIDIA B200. We wrote the formats, the quantize kernels and the layer ourselves, put real 4-bit and 8-bit tensor-core multiplies inside PyTorch, and trained GPTs of up to 3.3 billion parameters on TinyStories against the clock.

1. Why use fewer bits at all
2. The machine: one B200
3. What a floating-point number is
4. Precision against range
5. Why four bits need a scale
6. Block scaling, NVFP4 and MXFP8
7. A linear layer in 4 bits
8. The backward pass and the transpose
9. Bias, noise and stochastic rounding
10. The whole layer in one picture
11. Experts on other GPUs
12. A real 4-bit multiply inside PyTorch
13. Kernels: when the multiply stops being the slow part
14. One layer, then one training step
15. TinyStories on the B200, against the clock
16. How this compares with the article
17. What this does not show
18. A 90-minute workshop plan

1Why use fewer bits at all

Training a language model is mostly two activities: storing numbers and multiplying them. Every weight, every activation and every gradient is a number that has to live somewhere in memory, travel to the processor, be multiplied, and travel back. The width of that number, counted in bits, affects all three steps. A 4-bit number takes a quarter of the memory of a 16-bit number, and a quarter of the bytes move when it is loaded. The multiplier circuit for two 4-bit numbers is also far smaller than the circuit for two 16-bit numbers, so a chip designer can fit many more of them into the same area. The article gives the figures for NVIDIA's GB300 hardware: 2.5 PFLOPS in BF16, 5 in MXFP8 and 15 in NVFP4.

The price is that fewer bits can name fewer numbers. Four bits give sixteen patterns, so a 4-bit format can represent sixteen different values and nothing else. Anything in between has to be rounded to one of them. The figure below shows the article's own example.

Storing a value in a small format is called quantization. Reading it back as an ordinary number is dequantization. The round trip does not return what we put in: 3.2 goes in and 3 comes out. One such error is harmless. A training run makes this kind of error billions of times in every step, in every layer, and the errors feed into each other. The whole article is about arranging the rounding so that the errors stay small where they matter and cancel where they can.

2The machine: one B200

The article's numbers come from GB300 racks, in which 72 Blackwell GPUs are wired together. We rent the smallest piece of that world: a single NVIDIA B200, by the second, inside a container on Modal. Every measurement in this write-up was taken on it. It helps to know what is inside before asking it to do anything.

Three parts of the GPU matter for this story.

• The tensor cores are multipliers built for one job, the matrix multiply. Each of the 148 multiprocessors has them. Their rate depends on the width of the numbers: by the chip's own counters, a multiprocessor completes 8,192 multiply-add operations per clock cycle in BF16, 16,384 in FP8 and 32,768 in FP4. Halving the bits doubles the rate. This is the physical reason the article exists.
• The ordinary arithmetic units, which the article calls ALUs, do everything that is not a matrix multiply: comparing, rounding, computing scales. They are far slower per number than the tensor cores, and section 13 is about what happens when they have too much to do.
• The memory holds every tensor, and narrower numbers take less of it. A single activation tensor of 32,768 tokens by 32,768 features is 4.3 GB in FP32, 2.1 GB in BF16 and 0.6 GB in NVFP4. Moving fewer bytes is the second source of speed.

One GPU also sets a limit on what we can test. There is no second GPU to send tokens to, so the part of the article about communication between GPUs (section 11) is explained and checked for its effect on the numbers, but not timed.

3What a floating-point number is

Before four bits, it helps to recall how any float works. In school we write large and small numbers in scientific notation, such as 6.02 × 10²³. That notation has three parts: a sign, a few significant digits, and a power of ten that says how big the number is. A binary float stores the same three parts, with a power of two in place of the power of ten.

• The sign takes one bit.
• The exponent picks the power of two. It decides the neighbourhood the number lives in: between 1 and 2, between 2 and 4, between 4 and 8, and so on.
• The mantissa picks a position inside that neighbourhood. With one mantissa bit there are two positions, the start and the middle. With three bits there are eight.

The 4-bit format is called E2M1, and it is small enough to list completely. Take the pattern 0111, which the article's interactive tool also shows. The sign bit is 0, so the number is positive. The exponent bits are 11, which is 3; the format subtracts 1 from it, giving 2² = 4. The mantissa bit is 1, which means "one and a half times". The value is 4 × 1.5 = 6. The eight sizes are:

With the sign bit, that makes the sixteen values ±{0, 0.5, 1, 1.5, 2, 3, 4, 6} quoted in the article. The first two entries are the exception to the rule above: when the exponent bits are all zero there is no "one and" in front of the mantissa, which is how the format reaches 0 and 0.5.

4Precision against range

Why split four bits this way and not another? Leave the sign aside and consider the three remaining bits. They can name eight sizes. The question is where on the number line to put them, and there are three natural answers.

If all three bits go to the mantissa, the values are evenly spaced, half a unit apart, and stop at 3.5. This is a fixed-point number, and we will meet one again later under the name S2F6. If all three bits go to the exponent, every value is double the previous one. The format reaches 32, but it cannot tell 5 from 4. E2M1 is the compromise: steps of 0.5 up to 2, steps of 1 up to 4, one step of 2 up to 6.

The useful property of this compromise is that the gap grows with the number. Near 1 the grid is fine, near 6 it is coarse, so the error stays roughly proportional to the size of the value. Neural network tensors suit this well, because they contain many small values and a few large ones.

The 8-bit format E4M3 follows the same design with more of everything: 4 exponent bits and 3 mantissa bits. Its largest value is 448, its smallest non-zero value is 2⁻⁹, and there are eight steps inside every power of two. Our code counts 127 distinct sizes including zero. We will use E4M3 in two roles: as a number format in its own right, and as the scale that makes four bits usable.

5Why four bits need a scale

Here is the first practical problem. The weights of a trained network are small numbers. In the model we train on the B200 later in this write-up, the weights of the MLP layers are spread around zero with a standard deviation of about 0.02. The smallest non-zero size E2M1 can hold is 0.5. If we round these weights straight onto the grid, every one of them lands on zero and the layer is gone.

The fix is to divide by a scale before rounding and multiply by it again after reading. If the largest weight in a group has size 0.046, we choose the scale so that it lands on 6, the top of the grid. The other weights then spread over the grid below it.

A single scale for a whole tensor is not enough, though. A tensor has millions of entries, and its largest entry can be many times bigger than a typical one. If that entry sets the scale, the grid is stretched to reach it, and ordinary values fall below half a step and round to zero. How much this hurts depends on the tensor. The table in the next section shows that for our weights the loss is moderate, while for activations and gradients it wipes out most of the tensor.

6Block scaling, NVFP4 and MXFP8

The answer is to give every small group of neighbours its own scale. NVFP4 cuts a tensor into blocks of 16 consecutive values and stores one scale per block. A large value now spoils only the 15 values that share its block. The article's example is a block whose largest value is −54. Since the top of the grid is 6, the scale is 54 ÷ 6 = 9, and −54 is stored as −6. Every other value in the block is divided by 9 and rounded.

The scale has to be stored too, and it must be cheap, otherwise the saving is lost. NVFP4 stores it as one E4M3 number, 8 bits. That choice matters. An E4M3 scale has three mantissa bits, so it can be 8, 9, 10, 11, 12 and so on. The competing family of formats, the MX formats, store the scale as a bare power of two, called E8M0: eight exponent bits and no mantissa. For our block the nearest usable power of two would be 16, and then the largest value would only reach 54 ÷ 16 ≈ 3.4 on a grid that goes up to 6. Part of the grid would be wasted on every block. The article credits NVFP4's finer scale for its lower error on language model data.

The cost of the scale is spread over the block: sixteen 4-bit values and one 8-bit scale come to 72 bits, or 4.5 bits per value. That is the "3.6 times fewer bits than BF16" in the article, since 16 ÷ 4.5 ≈ 3.6.

One limit remains. An E4M3 scale tops out at 448, so a block scale combined with a code can represent sizes from 2⁻¹⁰ up to 6 × 448 = 2,688. Tensors do not always fit in that window, so NVFP4 is used with a third number: one ordinary FP32 scale shared by a very large group, such as a whole tensor. A value is rebuilt as

value  =  4-bit code  ×  8-bit block scale  ×  FP32 outer scale

The outer scale only slides the window up or down. Setting it to 2 turns the range [2⁻¹⁰, 2688] into [2⁻⁹, 5376]. In the recipe, weights get one outer scale per expert and activations get one per token.

MXFP8 is the 8-bit relative. Each value is a full E4M3 number, so the grid inside a block is much finer, and a block of 32 shares one power-of-two scale. It costs 8.25 bits per value. In the recipe it plays the role of the careful format: slower than NVFP4, but accurate enough for gradients.

The table below is the first measurement from our own model. We took real tensors from the trained network and stored them in each format. The number is the size of the error relative to the size of the tensor.

Relative error ‖stored − original‖ ÷ ‖original‖ on real tensors of the 840-million-parameter model after training on the B200. "FP4, one scale" uses a single scale for the entire tensor. NVFP4 + 4/6 is explained in section 10. Two of the tensors are held in BF16 by the model, so BF16 stores them exactly.

A single scale is poor for the weights, at 15% and 19%, and ruinous for the activations and the gradient, at 38% to 54%, where a few large entries stretch the grid and most values become zero. Block scaling brings every tensor to between 8% and 10%. MXFP8 sits at 2.7% everywhere, and BF16 at 0.17%.

Look at the NVFP4 column once more. Even with block scaling, every number is off by roughly a tenth of the typical size. NVFP4 is a rough format, and the rest of the article is about making training survive it.

E2M1 = NS(mbits=1, emin=0,  emax=2, vmax=6.0)
E4M3 = NS(mbits=3, emin=-6, emax=8, vmax=448.0)

def round_float(v, fmt, mode="rtn"):
    a = v.abs()
    _, ex = torch.frexp(a)                       # a = m * 2^ex, m in [0.5, 1)
    e = (ex - 1).clamp(fmt.emin, fmt.emax)       # which power of two
    step = torch.ldexp(torch.ones_like(a), e - fmt.mbits)   # gap between neighbours there
    t = a / step
    if mode == "rtn":  q = torch.round(t)                        # to nearest
    elif mode == "sr": q = torch.floor(t + torch.rand_like(t))   # stochastic (section 9)
    else:              q = torch.ceil(t)
    return torch.copysign((q * step).clamp(max=fmt.vmax), v)

def nvfp4(x, mode="rtn"):                        # blocks of 16 along the last dimension
    xb = x.reshape(-1, x.shape[-1] // 16, 16)
    amax_b = xb.abs().amax(-1, keepdim=True)                     # largest size in each block
    s_outer = xb.abs().max() / (6.0 * 448.0)                     # FP32 scale for the tensor
    sb = round_float(amax_b / 6.0 / s_outer, E4M3)               # 8-bit block scale
    scale = sb * s_outer
    return (round_float(xb / scale, E2M1, mode) * scale).reshape(x.shape)

7A linear layer in 4 bits

A linear layer multiplies a matrix of activations by a matrix of weights. Call the activations x, with one row per token and K numbers in each row. Call the weights w, with one row per output and K numbers in each row. The layer computes y = x wᵀ: every entry of the output is the dot product of one row of x with one row of w. This operation is called a GEMM, for general matrix multiply, and it is where nearly all the time of a large model goes.

In the forward pass of the recipe, both matrices are quantized to NVFP4 first, giving x_q and w_q, and the GEMM runs directly on the 4-bit codes. A dot product walks along the K direction, so both matrices are cut into blocks along K. The hardware takes a block of 16 codes from a row of x_q and the matching block from a row of w_q, multiplies the sixteen pairs, adds them up, and multiplies the sum once by the two block scales.

That last detail is what makes the format fast. The expensive part, sixteen multiplications and additions, happens on tiny 4-bit numbers. The scales are handled once per block. If the two operands were cut along different directions, the scales could not be pulled out of the sum, and the trick would not work. Keep this constraint in mind: a block-scaled GEMM needs its blocks to run along the dimension it sums over. The next section is about what happens when that dimension changes.

8The backward pass and the transpose

Training needs gradients, and for a linear layer the chain rule gives two of them. Let g be the gradient of the loss with respect to the output y, handed down from the layer above. It has one row per token and N numbers per row.

• DGRAD, the gradient for the layer below: dx = g w_q. This multiplies a (tokens × N) matrix by an (N × K) matrix, so the sum runs over N.
• WGRAD, the gradient for the weights: dw = gᵀ x_q. This multiplies an (N × tokens) matrix by a (tokens × K) matrix, so the sum runs over the tokens.

Both GEMMs reuse the tensors from the forward pass, but neither sums over K. DGRAD needs w_q cut into blocks along N, which means down its columns. WGRAD needs x_q cut into blocks across tokens. The article calls this "transposing the tensors". The numbers have not moved, but each one now shares a block with 15 different neighbours.

New neighbours mean a new scale, and a new scale means a new rounding. The value from Figure 6 was stored as 4.5 in the forward pass. Suppose its column block has a scale of 2. Then 4.5 ÷ 2 = 2.25, the nearest grid point is 2, and the backward pass reads 4.0. The forward pass computed the loss with one matrix, and the backward pass computes its gradient with a slightly different one. In the worst case the difference is not slight at all. The article illustrates it with a tile in which one whole row holds large values: every column block then contains one large value, every column scale is set by it, and most of the small values in the other rows fall to zero.

Our own trained weights are milder than that illustration, and it is worth seeing what really happens to them.

In this tile, the most uneven one in the matrix, a scale per row loses 22 of the 256 values and a scale per column loses 20. The two directions are equally good. They are not the same, though. Different cells are lost, and the cells that survive are rounded against different scales. Over the whole matrix, the forward copy and a copy re-blocked along the other dimension and rounded to nearest differ by 9.3% of the matrix's size. That is almost as large as the quantization error itself, which is 9.5%. The backward pass would be working with a matrix that is as far from the forward one as the forward one is from the truth.

NVIDIA's suggestion for weights is the last panel: compute one scale for each 16 × 16 tile. A square tile looks the same from both directions, so the forward and backward copies are identical bit for bit. The cost is visible in the figure. One scale now serves 256 values, and a single large value among them pushes the scale up for all the rest; this tile loses 45 values instead of 22. Thinking Machines report that they tried tiles and found them too coarse. Over whole tensors of our model the picture is the same:

Share of non-zero values that are rounded to exactly zero, on real tensors from layer 2.

For weights, tiles turn 10.1% of the values into zeros where 1 × 16 blocks lose 6.8%. For activations after the GELU it is 63% against 37%, and for the gradient 23% against 7%.

So the recipe needs to keep the fine 1 × 16 scales in the forward pass and still give the backward pass a matrix it can use. The way out comes from asking what the backward pass really requires.

9Bias, noise and stochastic rounding

There are two different kinds of error, and training reacts to them very differently.

Noise is error that is sometimes too high and sometimes too low, with no preference. A gradient is an average over thousands of tokens in a batch, and training takes hundreds of thousands of steps. Errors without a preferred direction cancel in those averages. Average four times as many and the leftover halves.

Bias is error that leans the same way every time. It does not cancel. Averaging more tokens leaves it exactly where it was, and it is passed from layer to layer and from step to step, so it can build up. The article's priority table says this in checkmarks: for backward passes, minimizing bias gets three, minimizing error gets one.

Ordinary rounding, round to nearest, is biased in this sense. It is deterministic: given 2.25 on a grid of 2 and 3, it returns 2 every time, and the error is −0.25 every time. In the transposed matrix above, every small value in a column with a large scale is pushed toward zero, on every step.

Stochastic rounding (SR) replaces the fixed choice with a weighted coin. The value is rounded up or down at random, and the odds are set by the distances, with the nearer neighbour more likely.

For 2.25, the lower neighbour is three times closer than the upper one, so SR returns 2 three quarters of the time and 3 one quarter of the time. The average is 0.75 × 2 + 0.25 × 3 = 2.25, which is exactly the value we started with. With the scale of 2 from the previous section, the backward pass reads 4.0 on most steps and 6.0 on the rest, and on average it reads 4.5: the number the forward pass used.

This is the key relaxation in the article. NVIDIA's tiles make the two copies identical, w_q = (w_qᵀ)ᵀ. The recipe only asks that they agree on average, E[w_q] = E[(w_qᵀ)ᵀ]. That weaker demand is enough to remove the bias, and it leaves the forward pass free to keep its 1 × 16 scales. The procedure for each saved tensor is:

1. Dequantize the forward tensor back to ordinary numbers. These are the values the forward pass really used, such as 4.5.
2. Transpose it, so that the blocks will run along the dimension the backward GEMM sums over.
3. Re-quantize it with stochastic rounding.

Stochastic rounding has a cost, which Figure 11 already shows: a single draw is further from the truth than round to nearest would be. In our tests, NVFP4 to nearest is off by 9.5% on Gaussian data, one stochastic draw is off by 14.9%, and the mean of 400 draws is off by 0.75%. SR trades a small bias for a larger noise. To keep that noise low, the recipe re-quantizes to MXFP8 instead of NVFP4. On the same data MXFP8 is off by 2.7% to nearest and 3.8% for one stochastic draw. Because the tensor was dequantized first, any target format can be chosen at this step, and the article reports that MXFP8 gave the best balance of accuracy and speed.

We measured both effects on our trained model. First, the weights alone: take the forward copy of a weight matrix, transpose it, re-quantize it, and compare.

Mismatch between the forward weights and their re-quantized transpose, an MLP weight matrix of the B200 model. "Average size kept" is the mean absolute value after re-quantizing to nearest, as a share of the original: below 100% means the weights were shrunk.

Rounded to nearest, the 4-bit transposed copy differs from the forward copy by 9.3%, and it is the same 9.3% every time, for as long as the weights stay where they are. One stochastic draw is further away, at 13.7%, but its error is noise: the mean of 16 draws is 3.4% away and the mean of 256 draws is 0.9% away, halving each time the number of draws is quadrupled. In MXFP8 everything is about three times smaller.

One thing we did not find is a direction. The article's illustration suggests that round to nearest shrinks small values toward zero. On our weights it keeps the average size within 1% of the original. The mismatch is fixed and repeatable, which is what makes it a bias in the statistical sense, but in this model it does not obviously push one way.

Second, the gradients, and this time on the real hardware path that the later sections build. We took one batch and ran the forward pass with real 4-bit multiplies on the tensor cores. We computed the exact gradient of that forward pass in high precision, and then computed it again with real 8-bit backward multiplies. With round to nearest there is only one answer. With SR every backward pass is different, so we averaged 1, 2, 4, up to 64 of them.

Rounded to nearest, the 8-bit backward pass gives a gradient that is 2.6% away from the exact one, and running it again gives the identical result, bit for bit. One stochastic pass is worse, at 3.7%. The mean of two passes already matches round to nearest, and the mean of 64 is at 0.47%. From 1 pass to 64 the error falls by a factor of 8.0, which is exactly what pure noise would give. Stochastic rounding leaves no fixed error behind.

For a 4-bit backward pass, the weight table above tells the rest of the story: the fixed mismatch is 9.3% and one stochastic draw is 13.7% away, three to four times the 8-bit figures. Stochastic rounding turns a fixed error into noise that averaging removes, and 8 bits make that noise several times smaller. Those two facts are the case for the recipe's backward pass.

Two notes on the measurement. The gradients are those of the last block's MLP, because PyTorch's own attention backward is slightly unrepeatable from run to run, which blurs the comparison for earlier blocks. And we averaged repeated passes over the same batch. In training the averaging happens over different batches and steps, and we did not measure how much of the round-to-nearest error survives there.

10The whole layer in one picture

We now have every piece of the article's central diagram. The figure below redraws it with the same colour for each precision.

Reading it from the top: activations arrive from the previous layer in BF16 and are quantized to NVFP4. The weights are kept by the optimizer in FP32 and are quantized to NVFP4 with 1 × 16 blocks. One 4-bit GEMM produces the output. The two quantized tensors are saved. On the way back, the gradient arrives in BF16 and is quantized to MXFP8 with SR twice, once as it is for DGRAD and once transposed for WGRAD. The saved weights and activations are each dequantized, transposed and re-quantized to MXFP8 with SR. DGRAD sends a gradient to the layer below, and WGRAD sends one to the optimizer, which updates the FP32 weights.

One consequence is easy to miss. The backward pass starts from the quantized activations, not the original BF16 ones. The high-precision activations are never needed again after the forward pass, so they do not have to be stored, or, as we will see, sent anywhere.

The article adds three smaller choices to this picture.

Four over six

The E2M1 grid has a weak spot at the top: nothing lies between 4 and 6. A value at 5/6 of the block's maximum has to move a long way. The technique called Four Over Six (4/6) tries a second scale for each block, one that places the largest value on 4 instead of 6. With that scale the top of the grid is unused, but the remaining points are spread differently across the block. The quantizer computes both versions and keeps whichever has the smaller error. Nothing extra is stored, because the choice is absorbed into the block scale. The recipe uses 4/6 for activations and not for weights. On our activations it lowers the error from 9.3% to 8.9% going into the MLP, and from 8.3% to 8.1% after the GELU.

One outer scale per token

For activations, the FP32 outer scale is computed separately for every token's row. This has two benefits beyond accuracy. The way a token is quantized does not depend on which other sequences happen to be in the batch, and it does not depend on tokens that come later in its own sequence. A scale shared across a sequence would let information about future tokens leak into the present one through the rounding.

The last layers stay in 8 bits

The last 16.6% of the layers are computed in MXFP8, not NVFP4. Our B200 model has four blocks, so a sixth of them is less than one layer, and we keep none in 8 bits.

class QLinearFn(torch.autograd.Function):
    # x is (M tokens, K), w is (N, K)
    @staticmethod
    def forward(ctx, x, w, c):
        xq = nvfp4(x, outer="row", four_six=True)     # per-token outer scale, blocks along K
        wq = nvfp4(w, outer="tensor")                 # 1 x 16 blocks along K
        ctx.save_for_backward(xq, wq)                 # only the quantized tensors are kept
        return xq @ wq.t()                            # FPROP

    @staticmethod
    def backward(ctx, g):                             # g is (M, N)
        xq, wq = ctx.saved_tensors
        w_t = mxfp8(wq.t().contiguous(), "sr")        # (K, N): blocks along N
        dx  = mxfp8(g, "sr") @ w_t.t()                # DGRAD sums over N
        g_t = mxfp8(g.t().contiguous(), "sr")         # (N, M): blocks along M
        x_t = mxfp8(xq.t().contiguous(), "sr")        # (K, M): blocks along M
        dw  = g_t @ x_t.t()                           # WGRAD sums over M
        return dx, dw, None

11Experts on other GPUs

The layers that Thinking Machines quantize are Mixture of Experts (MoE) layers. In an ordinary transformer block, every token passes through the same MLP. In an MoE block there are many MLPs, called experts, and a small router sends each token to a few of them. The experts together are very large, so they are spread over several GPUs. Running the layer then takes three steps: dispatch the tokens to the GPUs that hold their experts, compute, and combine the results back on the GPU the tokens came from.

Once the GEMM itself became about three times faster, the link between the GPUs became the slow part. The article gives the arithmetic. NVLink moves 900 GB per second in each direction, and the tensor cores perform up to 15 PFLOPS. Dividing one by the other, a GPU can do roughly 16,000 floating-point operations in the time it takes one byte to arrive. If the bytes do not come fast enough, the tensor cores wait.

Dispatch has a simple remedy that we have already met. The backward pass uses the quantized activations and never needs the BF16 originals, so the tokens can be quantized before they are sent. They travel at 4.5 bits per value.

Combine is harder. The outputs of the layer feed the rest of the network and need more precision; by default they are BF16, 16 bits per value. The obvious move is to send them as MXFP8. The article reports that this hurt the model, because MXFP8 represents outliers poorly. Their solution has two parts.

A rotation that removes outliers

A block-scaled format suffers when one value in a block is much larger than the rest, because the scale has to reach that value and the grid becomes coarse for everything else. A random Hadamard transform (RHT) is a way to get rid of such values without losing any information. It works on 32 numbers at a time in two steps: flip the signs of some of them according to a fixed random pattern, then replace the 32 numbers by 32 different sums and differences of all of them, each divided by √32.

Every output is a mix of all 32 inputs. A single input of size 100 no longer sits in one place; it contributes 100 ÷ √32 ≈ 17.7 to every output. The total energy of the block is unchanged, and the mixing matrix is its own inverse, so applying it again and undoing the sign flips returns the original numbers exactly. The receiver can always recover what was sent.

A format made for smooth data

After the rotation the values in a block are bell-shaped and of similar size. A float grid is not ideal for such data, since it spends many of its codes on values close to zero. The article uses S2F6, a fixed-point format: an 8-bit signed integer divided by 64. It covers −2 to 127/64 in equal steps of 1/64, which is the "all bits on mantissa" choice from Figure 4 at a larger size. With one power-of-two scale per 32 values, it costs 8.25 bits per value, about half of BF16.

We measured the wire formats on the outputs of the MLP layers of our B200-trained model. MXFP8 has an error of 2.66%. Rotated S2F6 has an error of 0.79%, more than three times lower at the same 8.25 bits. Most of that gain comes from the format: S2F6 without the rotation already reaches 0.92%, and the rotation takes it the rest of the way. These outputs do have outliers. Their kurtosis, a measure of how heavy the tails are, is 59 where a bell curve has 3, and the rotation brings it down to 7. Figure 15 shows the spikiest block of the batch, where the largest value falls from 5.6 times the typical size to 1.3 times.

The article gives outliers as the reason MXFP8 failed on their model. Fixed point with equal steps is the format an outlier hurts most, since one large value makes every step large, which is why the rotation and the format come as a pair. We have a single GPU, so we could measure what the wire format does to the numbers but not what it does to the clock.

Blackwell GPUs have instructions that convert to and from block-scaled S2F6 directly, which the article lists. The rotation and the conversion are done at the very end of the GEMM kernel, in what kernel writers call the epilogue, so the data never sits in memory in BF16 form. On the receiving side it is dequantized and added up as it arrives, and then the rotation is undone.

12A real 4-bit multiply inside PyTorch

Everything so far has been explanation. From here on we build the recipe on the B200 and time it. The goal is a training step in which the heavy multiplies of the MLP layers really are 4-bit going forward and 8-bit going back, on the tensor cores, with nothing emulated.

PyTorch cannot issue a 4-bit GEMM. NVIDIA's cuBLASLt library can, on Blackwell, if it is handed the operands already packed: 4-bit codes, sixteen to a 64-bit word, with the 8-bit block scales in a separate buffer arranged the way the tensor cores read them. We wrote a small C library, fb.cu, that sits between the two. It receives PyTorch's tensors by their memory address, runs our own quantize kernels to pack them, and calls the cuBLASLt multiply on the result. PyTorch loads it through ctypes, and the tensors never leave the GPU. The 8-bit multiplies of the backward pass go through the same library.

The first check is that the hardware computes what sections 6 to 10 say it should. We keep a plain PyTorch version of every format, written with ordinary tensor operations (q4.py, about 200 lines), as the reference. We ran one layer through the real 4-bit multiply and through that reference:

The differences are the rounding of the final FP32 sums. The tensor cores compute exactly code × block scale × outer scale, as the format says.

13Kernels: when the multiply stops being the slow part

With the plumbing in place we can time the multiply. For 32,768 tokens entering a layer of width 8,192 with 32,768 outputs, the BF16 GEMM takes 12.6 ms. The 8-bit GEMM takes 6.4 ms and the 4-bit GEMM 3.5 ms. The 4-bit multiply is 3.6 times faster than BF16, which is the hardware claim of the article, reproduced on our rented GPU.

The trouble starts next to the multiply. Before the second GEMM of an MLP, a billion activations have to pass through the GELU and be quantized. PyTorch's GELU alone takes 1.2 ms on that tensor. Our first kernel did the GELU and the 4-bit rounding with ordinary compares, shifts and masks. It took 2.8 ms, and 7.0 ms when it also wrote the 8-bit copy that the backward pass needs. A multiply that saves 9 ms is of little use if preparing its input costs 7.

This is the point the article makes about its own kernel stack, which the authors call BBBK. The tensor cores perform the GEMM. Everything else in the recipe, which includes finding the largest value of each block, computing scales, rounding, and drawing random numbers for stochastic rounding, is ordinary arithmetic done by the GPU's general-purpose units, the ALUs of Figure 2. When that work is slow, the tensor cores finish and wait. The article's remedy is to keep both busy at once, and to use the conversion instructions that Blackwell provides. Those instructions turn two floats into two 4-bit or 8-bit codes in a single step, rounding included. We rewrote our kernels around them.

The new kernels are two to five times faster. The fused GELU-and-quantize pass now takes 1.4 ms, close to what PyTorch needs for the GELU by itself, so the 4-bit conversion is nearly free when it rides along with work that had to be done anyway. That is a small-scale version of what the article calls fusing work into the GEMM's epilogue.

One detail surfaced on the way. The hardware instruction and our first kernel disagreed on about one block in a hundred. The cause is exact ties, such as a value that lands precisely on 0.75, halfway between the grid points 0.5 and 1. Our code breaks ties toward the even neighbour, and the hardware breaks them toward zero. Both are correct round-to-nearest. The 8-bit kernels agree with their software versions on every value.

__global__ void k2_quant_nvfp4(const void* xin, int is_bf16, long rows, int nb, float outer, uint64_t* codes,
                               uint8_t* scales, uint64_t* f8, float inv_s8, int do_gelu, uint32_t seed, int copy_sr) {
  long b = (long)blockIdx.x * blockDim.x + threadIdx.x;            // one thread, one block of 16
  if (b >= rows * nb) return;
  long r = b / nb; int kb = (int)(b - r * nb);
  float v[16], amax = 0.0f;
  load16(xin, is_bf16, b, v);                                      // 16 BF16 values arrive as four 64-bit loads
  if (do_gelu) for (int i = 0; i < 16; ++i) v[i] = gelu(v[i]);     // the activation rides along
  for (int i = 0; i < 16; ++i) amax = fmaxf(amax, fabsf(v[i]));
  uint8_t sc = max((int)enc_e4m3_mag(amax / 6.0f / outer, 0, 0.0f), 1);   // 8-bit block scale
  float scale = dec_ue4m3(sc) * outer, inv = 1.0f / scale;
  uint64_t w = 0;
  for (int j = 0; j < 8; ++j) {                                    // two values per conversion instruction
    uint32_t by = __nv_cvt_float2_to_fp4x2(make_float2(v[2*j] * inv, v[2*j+1] * inv), __NV_E2M1, cudaRoundNearest);
    w |= (uint64_t)by << (8 * j);
    /* ... optionally: decode, add the stochastic offset, convert to FP8 for the backward pass ... */
  }
  codes[b] = w;                                                    // 16 codes in one word
  scales[/* position in cuBLASLt's scale layout */] = sc;
}

14One layer, then one training step

With fast multiplies and fast kernels we can assemble a whole MLP layer as Figure 13 draws it. Going forward: quantize, 4-bit GEMM, GELU and quantize in one pass, 4-bit GEMM. Going back: four 8-bit GEMMs, fed by a gradient that is stochastically rounded to 8 bits and by the saved 4-bit tensors, stochastically re-quantized to 8 bits.

The forward pass falls from 28 ms to about 11, and the backward pass from 55 ms to about 31. The green segments are the price of low precision: work that BF16 does not have to do. They are small next to the multiplies they make possible, but only because of the rewrite in the previous section.

A layer is not a model. We built a GPT with four blocks, 16 heads and a context of 256 tokens, and timed complete training steps (forward pass, backward pass and the AdamW update) at three widths. The MLP layers use the real path above. Attention, normalization, the embedding and the optimizer are ordinary PyTorch in BF16 in all three versions. This mirrors the article, where only the expert layers change precision.

A full training step with the recipe is 1.18 times faster than BF16 at width 2,048, 1.28 times at 4,096 and 1.38 times at 8,192. The wider the layers, the larger the share of the step that is GEMM, and the more of the tensor cores' advantage reaches the clock.

The last column is the comparison the article itself makes: the 4-bit forward pass against a run that is already 8-bit everywhere. There the gain is 0.6%, 1.3% and 3.5%. To see why it is so much smaller than the 3.6 times of the bare multiply, we timed every operation inside a live step.

Where the time goes inside one training step at width 8,192, summed over the four blocks. Each timed operation was bracketed by a GPU synchronize, so the parts add up to slightly more than an untimed step.

The MLP layers take 312 ms of a BF16 step and 172 ms of a recipe step: the part of the model that the recipe touches runs 1.8 times faster. Its forward pass alone halves, from 100 ms to 50. Two things dilute that on the way to the bottom line. About half of the step is attention, normalization and the optimizer, which the recipe leaves alone. And of the 22 ms that the 4-bit multiplies save over 8-bit ones in the forward pass, 12 ms goes back into the extra quantization work that 4 bits require: block scales for every 16 values, and the stochastic 8-bit copies for the backward pass.

15TinyStories on the B200, against the clock

Speed per step is only worth having if the model still learns. We trained the wide model on TinyStories, a dataset of short stories written in the vocabulary of a small child, once in each precision, and recorded validation loss against the time spent training. If training works, the model writes a little story. The setup:

The recipe is faster than BF16 and ends at nearly the same loss. At width 4,096, averaged over two seeds, the BF16 run takes 336 seconds and the recipe 259: 1.30 times faster, or 23% less time. The final losses are 1.688 and 1.691. The gap of 0.003 is small next to anything else on this page; the two BF16 seeds differ from each other by 0.001. Read the other way, BF16 needs 326 seconds to reach the loss at which the recipe stops after 262.

The stories are the same kind of story. Each model continued the same prompt with the same sampling settings.

Against 8-bit, the 4-bit forward pass is a small step, not a leap. The all-8-bit run takes 274 seconds, which makes the recipe about 6% faster per step here. Each training run had its own B200, though, and machines differ by a few percent; timed back to back on one GPU, the gap at this width was 1.3% (Figure 18). The recipe also learns slightly less per step: the 8-bit run ends at 1.686. Measured as time to reach the same loss, the two are level. The 8-bit run passes the recipe's final loss at 264 seconds, against 262.

The wider model shows the same pattern, larger. At width 8,192, with 3.3 billion parameters, the recipe finishes its 1,000 steps in 438 seconds against 591 for BF16, which is 1.35 times faster, and 463 for 8-bit, which reaches the recipe's final loss at about the same moment.

In this short run the recipe is 0.03 behind BF16 at the same step count (2.40 against 2.37), a larger gap than at width 4,096. Against the clock it is still well ahead: BF16 needs 545 seconds to reach the loss the recipe has after 438.

Stochastic or nearest rounding in the backward pass made no difference to the loss. The run with every backward operand rounded to nearest ends at 1.690, against 1.691 for the recipe. Figure 12 shows that the fixed error is real at the level of one gradient. At 8 bits, over 2,000 steps in a four-block model, it is too small to move the result.

Trained in 4 bits, or quantized afterwards?

The article notes that training in the final numerics from the start removes the need to calibrate a quantized checkpoint at the end. We can test the idea directly: take the finished BF16 model and switch on the real 4-bit forward pass with no further training.

Validation loss on 40 held-out batches, the same for every row.

Quantizing the finished BF16 model raises its loss from 1.690 to 1.695. The model trained in 4 bits scores 1.692 on the same batches, so training in the final numerics recovers about half of what quantizing afterwards loses. The difference of 0.003 is two to three times the spread between seeds: small, and in the direction the article describes. The practical benefit is simplicity. No separate step is needed, and the checkpoint that comes out of training is already the 4-bit model.

16How this compares with the article

The direction of every claim we could test holds on our hardware. The size of the last step, 4-bit over 8-bit, is smaller for us than for them, and there are good reasons to expect that. Their kernels fuse the quantization into the GEMM itself, where ours run beside it. Their models spend most of their compute in the layers that change precision, where ours spend half. And part of their gain comes from sending fewer bits between GPUs, which a single GPU cannot show.

17What this does not show

• One GPU. The article's largest effects involve many GPUs: experts spread across them, and quantized data on the links between them. We have one B200, so nothing here is evidence about that half of the recipe beyond its arithmetic.
• A small task. The layers are as wide as a real language model's, but the model has four blocks and was trained on 33 to 66 million tokens of children's stories. The article is about a model trained on trillions of tokens. Effects that build up over depth and time, bias above all, have little room to appear here, which is probably why stochastic and nearest rounding ended at the same loss.
• The width-8,192 run is short. A 3.3-billion-parameter model was given 1,000 steps. It shows the clock clearly and the loss only roughly.
• Not the full recipe on hardware. The real runs use an 8-bit format with one scale per tensor in the backward pass, one outer scale per tensor, and no Four Over Six. The list is under section 16.
• Formats we did not run on the tensor cores. The numbers for a 4-bit backward pass, for MXFP8 with block scales, for 16 × 16 tiles and for the rotated wire format were computed on tensors from the B200 model with our reference arithmetic, q4.py. That arithmetic matches the tensor cores to 6 parts in 100 million where we could compare the two, on the NVFP4 multiply.
• Seeds. BF16 and the recipe were each run twice at width 4,096. Every other training result is a single run.
• The baseline. "BF16" means stock PyTorch with autocast, which is how most models are trained today, not a BF16 path tuned with the same care as our kernels.
• Someone else's GEMM. The multiplies are cuBLASLt's. Writing a tensor-core GEMM with the quantization fused into it, as the article's authors did, is the obvious next step and would move the 4-bit-over-8-bit number most.

18A 90-minute workshop plan

This write-up follows the order in which we would teach the article.

python3 mo.py modal deploy train.py && python3 mo.py call prep        # the data, once
python3 mo.py modal deploy real.py
export FOURBIT_APP=fourbit-real
python3 mo.py call probe                                            # hardware against arithmetic, and the first timings
python3 mo.py call bench                                            # full training steps at three widths
python3 mo.py spawn train_real arm=nvfp4 d=4096 steps=2000 lr=0.0002   # arm: bf16, fp8, nvfp4 or nvfp4_rtn