Skip to content

NVLS bf16/fp16 sum reductions on B300 are not round-to-nearest, and each result depends on the other elements of its 8-byte group #132

Open

Description

@0z5a

Environment: one node, 8× B300 (SM103), NVSwitch, CUDA multicast supported; NVSHMEM devel @ 2dca770 built for sm_103 with CUDA 13.3; 2 PEs (one GPU each), MPI bootstrap.

What I ran: each PE fills 2^20 elements from a hash of (PE, index): random sign and mantissa, exponent in [-4, 4], so the exact sum of the two inputs is representable in fp32. PE 0 compares the result of a sum reduction over NVSHMEM_TEAM_WORLD with the exact sum rounded to nearest even. About 293k of the 2^20 sums are exactly representable in the 16-bit type; the table counts only the inexact ones.

Path NVSHMEM_DISABLE_NVLS inexact sums round-to-nearest toward zero away from zero 5 repetitions
nvshmemx_bfloat16_sum_reduce_on_stream 0 755,732 69.6% 17.8% 12.7% bitwise identical
nvshmemx_bfloat16_sum_reduce_on_stream 1 755,732 100% 0 0 bitwise identical
nvshmemx_half_sum_reduce_on_stream 0 754,993 69.6% 17.6% 12.7% bitwise identical
nvshmemx_half_sum_reduce_on_stream 1 754,993 100% 0 0 bitwise identical

Observations:

  • The result is always one of the two representable neighbours of the exact sum, and for a given input the choice is deterministic.

  • It is not specific to how the reduction is issued. nvshmemx_*_sum_reduce_block, and kernels that run multimem.ld_reduce.relaxed.sys.global on nvshmemx_mc_ptr() with add.bf16x2, add.acc::f32.bf16x2, add.v4.bf16x2 and add.acc::f32.v4.bf16x2 (and the f16x2 forms), all give bitwise the same output as the host API.

  • The choice is predicted by one rule for 99.95% of the inexact sums (bf16 and fp16 alike). Number the elements of each aligned 8-byte group 0 to 3, let rem_q be the fraction of the way the exact sum of element q lies from the lower-magnitude neighbour to the higher one, and let b22(r), b21(r) be the top two mantissa bits of the fp32 exact sum of element r. Element q rounds away from zero iff

    rem_q + b22(q+1)/2 + b22(q+2)/4 + b22(q+3)/8 + b21(q+1)/16 + b21(q+2)/32 + b21(q+3)/64 >= 1 (indices mod 4)

    So it behaves like stochastic rounding whose random bits are the high mantissa bits of the other three elements of the group: the error is at most 1 ulp and roughly unbiased, but the result is not round-to-nearest and depends on the neighbours.

  • Consistent with that, shifting the whole input by k elements leaves every result unchanged for the same data when k = 4, 8 or 16, but changes 15–19% of them for k = 1 or 2.

Questions:

  1. Is this the expected rounding of NVLS reductions of 16-bit floats on this platform? If so, could the documentation of the NVLS collectives say that bf16/fp16 results are not round-to-nearest and do not match the NVSHMEM_DISABLE_NVLS=1 path bitwise?
  2. Because the rounding uses the other elements of the 8-byte group, the last word of a range whose byte length is not a multiple of 8 is rounded with whatever lies past the end of the range. NCCL hit the same effect and now zero-pads its scatter buffers ([Issue]: BF16/FP16 NVLS AllReduce and ReduceScatter are not repeatable when count % 4 != 0 nccl#2479). In one odd-length run I did see the second-to-last element change by one ulp when only the data past the end changed. That run also hit the unwritten-slice problem fixed in coll: keep NVLS reductions on 4-byte aligned 16-bit slices #131, though, so I would not lean on it. Should the NVLS paths avoid groups that extend past nelems, or is this acceptable given that the result is not round-to-nearest anyway?

While testing this I found that NVLS reductions of 16-bit types do not write slices that start on an odd element. That is a separate problem, with a fix in #131.

The probe is a single .cu file (about 350 lines) built against the installed NVSHMEM; I can attach it.

Activity

  1. 0z5a commented on Oct 25, 2026

    @0z5a
    ContributorAuthor

    A cleaner answer to question 2, now with #131 applied so that nothing is left unwritten. Same setup (2 PEs, B300), nreduce = 1026 and 2046. Both are below the one-shot threshold, and their byte length is not a multiple of 8, so the last 8-byte group is half outside the range. I ran each case twice and changed only the 16 source elements past the end (different random values in each run), then compared the in-range results:

    nreduce type in-range outputs that changed, NVLS same, NVSHMEM_DISABLE_NVLS=1
    1026 bf16 1 (index 1025) 0
    1026 fp16 1 (index 1024) 0
    2046 bf16 1 (index 2045) 0
    2046 fp16 1 (index 2045) 0

    *_sum_reduce_on_stream and *_sum_reduce_block give the same result. In both runs every output is one of the two neighbours of the exact sum, so the change is a one-ulp flip inside the last group. In short, the last elements of a 16-bit NVLS sum whose byte length is not a multiple of 8 depend on memory past nelems. In nvshmemi_add_reduce_mcast_threadroup() the last two elements of these ranges go through a 32-bit multimem.ld_reduce (after the 16-byte, and for 2046 the 8-byte, words), so the instruction itself does not cover the bytes past the end. If that is right, keeping operations inside nelems would not change this, and only zero padding (as NCCL now does) would make the result independent of the neighbouring memory. I only tried two tail patterns, so the count of one changed element per case means only that it happens.

    The longer odd lengths I tried on the patched build, 1048573 to 1048575, showed no tail dependence. In those runs the CTA that holds the end of the range takes the non-NVLS path, so they do not test this.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Labels

    No labels
    No labels

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions