Floating point and reproducibility

Day 26 ended with a table it did not fix:

kernel             distinct                min                max      max - min
----------------   --------   ----------------   ----------------   ------------
global atomic            20          6289136.5            6289304          167.5
shared atomic             6          6289251.5            6289254            2.5
tree + atomic             3          6289252.5          6289253.5              1

Twenty runs of one program on one Tesla T4, twenty different answers from the global atomic, and nothing crashed, raced or misread memory (code/day26-atomics/evidence/run-2026-08-30.txt, 2026-08-30). Every one of those twenty numbers is a legal output of that kernel.

This page explains why, then builds a reduction that returns the same bit pattern on every run. It checks the result by hashing the bits across 100 runs and states what that promise does not cover.

Twenty answers, none of them wrong

Float addition is not associative. NVIDIA's floating point guide states it as arithmetic, not as a GPU quirk: "the results corresponding to the sum rn(rn(A + B) + C) and the sum rn(A + rn(B + C)) are different from each other" (section 2.2, https://docs.nvidia.com/cuda/floating-point/index.html , checked 2026-09-01).

Each addition rounds to the nearest representable float, and where the rounding lands depends on what has already been accumulated. A sum of n floats is really a choice of parenthesisation, and different parenthesisations are different numbers in the last bits.

A sequential loop makes that choice once, in program order, so it returns the same bits every time. atomicAdd on a float hands the choice to the scheduler. Day 27 established what a relaxed atomic promises: atomicity at a scope, and nothing about order.

When 1,024 blocks each add their partial sum to one float, the additions happen in arrival order, arrival order depends on which blocks landed on which of the card's SMs first, and the result is a different parenthesisation per run. The result loses floating point determinism because atomic contention leaves the order to the scheduler. No memory-model rule is broken.

The numeric difference is small, but it blocks some checks. A test cannot use ==, and CI cannot compare output bits with an earlier run.

A changing baseline also makes a debugging comparison less useful. In these cases, a tolerance is not enough; the bit pattern must stay fixed.

Diagram: three orders for the same 1,024 partial sums. Left column, 1,024 partial-sum boxes; right, the single float they end in; arrows are additions, numbered by when they happened. Band 1, run 1 of the atomic version: arrows arrive in one shuffled order. Caption "run 1: one arrival order out of 1024! possible ones." Band 2, run 2 of the atomic version: the same boxes, a different shuffle, the result box highlighted as differing in the last bits. Caption "run 2: a different order, a different last bit, both legal." Band 3, the tree: arrows drawn as a fixed bracket over pairs, ten rounds from 1,024 to 1. Caption "the tree: ten fixed rounds, the same parenthesisation every run, the same bits." Alt text: "One thousand twenty-four partial sums reach one float three ways. Two atomic runs add them in different arrival orders and disagree in the last bits. A fixed tree adds them in ten pinned rounds and cannot."

Contracts from tolerance to fixed bits

People often pick a float tolerance without checking the source of the error. A tolerance is right when the addition order can vary, which is why day 5's harness uses one.

Each fixed decision gives a stronger contract:

1. Fixed-shape tree: same bits every run, one binary, one card. Day 24's tree already had this property and never claimed credit for it: which pairs meet in which round is decided by tid and the halving loop, so the parenthesisation is frozen in the code. Day 26's tree variant pinned its intra-block pairing the same way, and still logged three distinct totals in twenty runs, because its atomic tail handed the final order back to the scheduler.
2. Fixed launch shape. The tree's parenthesisation is a function of the grid. Derive the block count from the element count, an occupancy call or multiProcessorCount, and the same binary returns different bits at a different size or on a different card. Reproducibility means the shape is a constant, not a computation.
3. Kahan improves accuracy, not reproducibility. Compensated summation tightens the error of a sequential accumulation you control. It cannot touch an atomic tail, because the compensation term needs the running sum and atomicAdd never shows it to you. A Kahan-corrected atomic reduction can still change between runs.
4. Cross-GPU reproducibility is a different contract. cuBLAS defines the line precisely: its routines "generate the same bit-wise results at every run when executed on GPUs with the same architecture and the same number of SMs" (section 2.1.4, https://docs.nvidia.com/cuda/cublas/index.html , checked 2026-09-01), and it withdraws even that for the atomics-based variants of cublas<t>symv(). Same bits on a T4 and an H100 means pinning shape, order, math-function implementations and code generation across two instruction sets. This course does not promise it and neither should you, without pricing it first.

The compiler adds a fifth decision. Day 47 showed that nvcc may contract x * x + y into one FMA with one rounding, and that -fmad=false splits it back into two roundings and changed a residual from exactly right to exactly wrong. Contraction is a per-build decision, so a reduction that is bitwise stable across runs can still change bits after a rebuild.

Fast math can also change the generated instructions and result bits. A reproducibility claim must name its input, shape, order, and build.

Fix the order of every addition

Full program in code/day68-reproducibility/reproducibility.cu. It sums f(x) = x * x + 0.25f over 4,194,915 floats filled with day 26's inexact pattern, 100 runs per kernel, and records the 32-bit pattern of every result. Each design choice below pins one more decision.

Every kernel accumulates a multiply-add, so compiler settings affect it. A bare sum has no multiply and nothing for -fmad to change; squaring each element gives every accumulation step one contraction decision, which the second build takes away.

__device__ float squarePlusQuarter(float x) {
    return x * x + kQuarter;
}

The nondeterministic kernels are reported, never gated. Ordering nondeterminism is legal, so a gate demanding that the bits move would fail on scheduler luck, the same reasoning day 27 applied to its racy kernel. sumSquaresAtomic walks a grid-stride loop, reduces the block in shared memory with day 24's tree, and then does the one thing this page is about:

    if (threadIdx.x == 0) {
        // Deliberate (day 68): arrival order decides the addition order.
        atomicAdd(total, blockSum);
    }

sumSquaresKahanAtomic is the same kernel with Kahan compensation on the per-thread accumulator. This makes the accuracy result measurable while the order can still vary. The fixed version replaces the atomic with a second launch. Pass 1 writes one partial per block; pass 2 folds the 1,024 partials in a single block:

__global__ void sumPartials(const float* __restrict__ partials,
                            float* __restrict__ total) {
    __shared__ float tile[kThreadsPerBlock];

    float acc = 0.0f;
    for (int p = static_cast<int>(threadIdx.x); p < kBlocks;
         p += kThreadsPerBlock) {
        acc += partials[p];
    }

    const float treeSum = reduceBlock(tile, acc);
    if (threadIdx.x == 0) {
        *total = treeSum;
    }
}

The claim is a hash, so two builds can be compared with grep. Each kernel's 100 bit patterns are folded through FNV-1a; equal hashes mean all 100 runs matched bit for bit.

// FNV-1a over the whole run sequence, in order. Two binaries printed the
// same hash for a kernel if and only if all 100 runs matched bit for bit,
// which is the comparison the -fmad=false build exists for.
static uint64_t hashRuns(const uint32_t* bits, int runs) {
    uint64_t h = kFnvOffset;
    for (int i = 0; i < runs; ++i) {
        for (int b = 0; b < 4; ++b) {
            h ^= (bits[i] >> (8 * b)) & 0xffu;
            h *= kFnvPrime;
        }
    }
    return h;
}

The tree returns a fixed answer, which may still differ from the reference. All three kernels are checked against a double Kahan reference with a loose relative tolerance.

The tree has one extra check: its bits must not change. The output reports accuracy and determinism in separate columns.

Results

Measured on the project's Tesla T4 with driver 595.84 and CUDA 12.6 (V12.6.85) on 2026-09-01. Both binaries passed every gate.

Both CUDA 13.0 builds also passed. The two-pass tree retained exactly one pattern, 0x4b25426d, the same 1.739e-09 maximum relative error and the same 0x8189389fe1bc0a65 hash.

Atomic distinct counts and hashes moved between runs, as the experiment predicts: the default build reported 5 and 6, while -fmad=false reported 5 and 5. Those reported values are nondeterministic and remain outside the gates.

Default build:

kernel	runs	distinct patterns	first bits	max rel err	run hash
sumSquaresAtomic	100	6	0x4b25426e	2.753e-07	0x27e53ec1e5219b2d
sumSquaresKahanAtomic	100	5	0x4b25426e	1.864e-07	0xdc79f9deb705e63d
two-pass tree	100	1	0x4b25426d	1.739e-09	0x8189389fe1bc0a65

The -fmad=false build kept the same distinct counts and errors. Its atomic hashes moved to 0xcfde1f533972dd30 and 0xfdc309f74baa4250, but the tree stayed at the same first bits and the same 0x8189389fe1bc0a65 hash.

The central prediction held: scheduler-ordered atomic tails produced 6 and 5 patterns, while the fixed tree produced exactly 1. The build-specific tree prediction did not hold.

Disabling FMA did not change this input and reduction's tree result. A compiler flag is part of the reproducibility conditions, but changing it need not change the answer.

Constant bits show that these conditions produced the same result on each run.

Run it yourself

Use a supported CUDA GPU. The exercise needs no counters, root access, or tool beyond nvcc because it produces two binaries and their printed tables.

Both build lines are in code/day68-reproducibility/README.md; the second adds -fmad=false. The program needs about 17 MiB of device memory and takes no arguments, so the page and the binary cannot disagree about sizes.

Compiler Explorer's runner can hold the file, but 300 kernel launches may exceed its 20 second limit. Use a notebook or a local GPU instead.

Exercise

Make three edits, one at a time, rebuild, and record the two-pass tree's bit pattern and run hash after each: raise kBlocks from 1,024 to 2,048; reverse the direction of pass 2's per-thread loop; rebuild unchanged with -fmad=false. All gates must stay green throughout.

Time: 25 to 40 minutes. Submit: the three bit patterns next to the baseline's, and one sentence naming which part of the pinned tuple each edit moved.

Check: the program's own gates, as real branches returning EXIT_FAILURE: the tree must report exactly 1 distinct pattern over 100 runs, and every kernel must stay within 1e-5 relative of the double Kahan reference. A failure names the kernel, the distinct count and the error. The atomic kernels' distinct counts are printed and never gated, so a change there is a result to report, not a pass or a fail.

Hint 1

None of the three edits should break the gate. So what did each one change, if not whether the runs agree? Look at what the runs agree on.

Hint 2

The bits are a function of four things: the input, the launch shape, the addition order, and the instructions the compiler emitted. Assign each edit to one of the four.

Solution

All three edits leave the tree bitwise reproducible. Doubling kBlocks changes the shape, so different elements meet in different rounds. Reversing the pass 2 loop keeps the shape and changes the order within each thread's slice.

-fmad=false keeps shape and order but changes the generated instructions, using two roundings where there was one. Only returning the order to the scheduler through the atomicAdd tail breaks run-to-run agreement.

A reproducible result depends on the input, shape, order, and build. Each edit changes one of those conditions.

Pitfalls

You made it deterministic and started calling it correct. A fixed answer can be wrong on every run. The tree still differs from the double reference. Keep the accuracy and determinism checks separate.

You tested determinism on friendly data. Day 24 and day 26's timing sweep used inputs whose sums are exact integers below 2^24, where every addition order returns identical bits and even the global atomic looks reproducible. That is the same shape as day 27's lesson: a hundred green runs on inputs that cannot expose the property prove nothing about it. Test on values that round.

Your grid comes from the device, so your bits do. A block count derived from multiProcessorCount or an occupancy query gives a different tree shape on a different card, and the same binary stops agreeing with itself across machines. That is the exact condition cuBLAS attaches to its own guarantee: same architecture, same number of SMs.

You reached for -fmad=false to make the GPU match the CPU. Day 47's verdict stands: it usually makes both sides equally wrong instead of one of them right, and it re-prices every kernel in the file. It is legitimate here as a pin, a way to freeze the contraction decision across toolkit updates, if you pay for it knowingly and write down why.

You pinned your kernels and a library call unpinned the total. Library routines may use atomics internally; cuBLAS documents that the faster cublas<t>symv() path "is not guaranteed to be bit-wise reproducible because atomics are used for the computation" (section 2.1.4, checked 2026-09-01). A reproducibility claim covers the whole pipeline or it covers nothing.

Go deeper

• Floating Point and IEEE 754 Compliance for NVIDIA GPUs, sections 2.2 and 4.3: https://docs.nvidia.com/cuda/floating-point/index.html (checked 2026-09-01)
• cuBLAS Library User Guide, 2.1.4 "Results Reproducibility": https://docs.nvidia.com/cuda/cublas/index.html (checked 2026-09-01)
• cuda-samples, cpp/2_Concepts_and_Techniques/reduction: https://github.com/NVIDIA/cuda-samples/tree/master/cpp/2_Concepts_and_Techniques/reduction (checked 2026-09-01)
• Programming Massively Parallel Processors, 4th edition, chapter 10, on reduction and its ordering: https://shop.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-323-91231-0

Next

Day 69 takes the tuple's last entry further: one binary that carries code for several GPU generations and picks a path at run time, which is what shipping to a T4 and an H100 at once actually requires. Day 70 closes the module with the bug catalog, and this page's twenty-answers table earns its place in it as the bug that passes every test.