Skip to content

[AI][BUG][Iluvatar] Non-deterministic results across repeated identical forward passes (vision tower reproduces it standalone; NaN on some samples) #617

Open

Description

@tonyh168

Summary

On Iluvatar BI-V150 (iluvatar backend, torch_fl 2.10.0+corex4.5.0), a repeated forward pass over the same input tensors in the same process does not return the same result. Running the same sample 4 times in one process gives e.g.

0.99418   0.87493   0.99427   0.99427        # prob_1, same batch, same weights, same input tensors

The same model, same weights, same input on the same machine with the vendor DTK torch (2.10.0+corex.4.5.0) is bit-identical across repeated runs (0.99539042 × 3). End to end this costs us accuracy: 4 of 12 reference samples fall outside tolerance, including two NaNs.

I isolated it further: the vision tower alone reproduces it (visual(...) called 4× with the same pixel_values), while a synthetic transformer built from the same primitives is bit-identical, and each individual op I tested (mm / bmm / softmax / layer_norm / rms_norm / conv1d / cumsum / sdpa / exp / sum / mean) is bit-identical across repeats.

The model weights are internal, so I cannot hand you a one-command repro. I have the hardware and can run any diagnostic probe you ask for (see "What I can run for you").


Environment

Item Value
Host iluvatar-186, 16 × Iluvatar BI-V150 (32 GiB each)
Kernel module iluvatar.ko 4.5.0
Host CoreX /usr/local/corex → corex-4.4.0; /usr/local/corex-4.5.0.20260509 also present
Host ixsmi IX-ML 4.4.0 / Driver 4.5.0
Image harbor.baai.ac.cn/flagos-dev/pytorch-plugin-fl@sha256:bc1b3c0beb22b301ec501896ea1c77903879526fb7303358669321d609458023 (the CI image named in docs/vendors/iluvatar/installation.md)
Python 3.12.11 (venv /opt/torch-fl-iluvatar-venv)
torch 2.10.0+cpu
torch_fl 2.10.0+corex4.5.0 (built from source, see below)
flagtree 0.6.2a3+iluvatar3.6 (Triton 3.6)
flag_gems 5.4.0
flagcx not installed (single card)
transformers 5.17.0
device torch.flagos.is_available() == True, device_count() == 16, flagos:0 = Iluvatar BI-V150, 32.0 GiB

Container (three bind-mounts are needed on this host — see notes at the end):

H=/usr/local/corex-4.5.0.20260509
IMG=harbor.baai.ac.cn/flagos-dev/pytorch-plugin-fl@sha256:bc1b3c0beb22b301ec501896ea1c77903879526fb7303358669321d609458023

docker run -d --name tfl --privileged --ipc=host --net=host --shm-size 64g \
  -v /dev:/dev -v /lib/modules:/lib/modules -v /sys:/sys \
  -v $H:/opt/host-corex45:ro \
  -v $H/lib64/libcuda.so.1:/usr/local/corex/lib64/libcuda.so.1:ro \
  -v $H/lib64/libixml.so:/usr/local/corex/lib64/libixml.so:ro \
  -v $H/bin/ixsmi:/usr/local/corex/bin/ixsmi:ro \
  -e LD_LIBRARY_PATH=/opt/host-corex45/lib64:/usr/local/corex/lib64:/usr/local/openmpi/lib:/usr/local/lib \
  --entrypoint /bin/bash $IMG -c 'sleep infinity'

Build torch_fl (the image does not ship it):

docker exec tfl bash -lc '
  source /opt/torch-fl-iluvatar-venv/bin/activate
  cd /workspace/Torch-FL        # GitHub clone of flagos-ai/Torch-FL
  FLAGOS_ACCELERATOR=iluvatar COREX_PATH=/usr/local/corex \
    pip install --no-build-isolation --no-deps -e .'
# → Successfully installed torch_fl-2.10.0+corex4.5.0

Model: a 2B Qwen3.5-VL checkpoint (Qwen3_5ForSequenceClassification) in bf16, single forward pass, attn_implementation="sdpa", two images + ~240 text tokens per sample.


Repro shape

import torch_fl, torch
from transformers import AutoProcessor, Qwen3_5ForSequenceClassification

proc = AutoProcessor.from_pretrained(CKPT)
model = Qwen3_5ForSequenceClassification.from_pretrained(
    CKPT, dtype=torch.bfloat16, attn_implementation="sdpa", device_map=None
).to("flagos:0").eval()

inputs = proc(text=[prompt], images=[img1, img2], return_tensors="pt")
inputs = {k: v.to("flagos:0") for k, v in inputs.items() if hasattr(v, "to")}

with torch.inference_mode():
    for _ in range(4):                       # same `inputs` object, no resampling
        logits = model(**inputs).logits.float()
        print(torch.softmax(logits, -1)[0, 1].item())

Full model — repeated runs, bf16, one process

run1 run2 run3 run4 vendor torch (3 runs)
0.99418 0.87493 0.99427 0.99427 0.99539042 / 0.99539042 / 0.99539042

Other repeats on another sample (s12, reference 0.9124):

0.70335  0.71143  0.71143
0.75092  0.70335  0.70335

Note it is not always the first call that differs — in the first table it is run2, in the vision-tower test below it is run0. There is no stable ordering; that is what makes it look like stale/uninitialised memory rather than a warm-up path.

12 reference samples (bf16)

sample reference Torch-FL Δ
s01 0.026759 0.034359 +7.60e-03
s02 0.995095 0.994089 −1.01e-03
s03 0.000029 0.000030 +9.30e-07
s04 0.996777 0.995532 −1.25e-03
s05 0.001016 0.001455 +4.39e-04
s06 0.987179 0.984811 −2.37e-03
s07 0.000030 0.000030 −1.82e-12
s08 0.874077 0.815233 −5.88e-02
s09 0.000028 0.000057 +2.90e-05
s10 0.988580 NaN —
s11 0.000001 NaN —
s12 0.912436 0.735262 −1.77e-01

The four bad samples are s08 (formula, 246 tokens), s10 (table, 696), s11 (table, 428), s12 (table, 551) — so prompt length alone does not predict failure; all 8 samples with 222–434 tokens except s08/s11 are fine. The two NaNs appear only when samples are run in sequence; running s10 alone gives 0.983 — i.e. there is also state leaking across calls.

Vendor torch on the same machine, same batch: 12/12 within tolerance, max |Δ| = 8.65e-3.


Isolation so far

Vision tower alone is non-deterministic

mm = proc.image_processor(images=[img1, img2], return_tensors="pt")
mm = {k: v.to("flagos:0") for k, v in mm.items() if hasattr(v, "to")}

with torch.inference_mode():
    for i in range(6):
        out = model.model.visual(mm["pixel_values"], grid_thw=mm["image_grid_thw"])
        e = out.last_hidden_state
        torch.flagos.synchronize()
        f = e.float().cpu()
        print(i, f.mean(), f.std())
run 0   mean 1.11808   std 79.563
run 1   mean 1.10983   std 79.137
run 2   mean 1.10983   std 79.137     ← runs 1..5 are bit-identical to each other
...

grid_thw = [[1,10,28],[1,6,46]], pixel_values shape (556, 1536).

This goes through the vision tower's variable-length attention path: transformers splits q/k/v per image with torch.split(tensor, lengths.tolist(), dim=2) and runs one attention call per image. With attn_implementation="eager" the same non-determinism appears with one image (run0 1.09946 vs runs 1–3 1.09230), and with sdpa + one image all four runs differ (1.09019 / 1.05524 / 1.06326 / 1.05995) — so the pattern is not tied to sdpa and not strictly "first call only".

Vision tower also deviates from vendor torch

Same pixel_values, one forward after a warm-up call, element-wise vs vendor torch:

shape (556, 1024)
vendor: mean 1.08775  std 77.352  absmax 5952
tfl   : mean 1.11767  std 79.532  absmax 6016
|Δ|   : max 5207   mean 2.079   (2.7 % of std)
per-token mean error: median 0.495, max 29.5   (std 77.4)
tokens with per-token mean error > 0.5·std: 0 / 556

So the vision tower output differs by ~2.7 % of its std, spread over all 556 tokens and all 1024 channels — not a localized channel/block. For bf16 I would expect this to be near the rounding floor.

What is not affected

Probe Result
torch.mm bf16, M∈{256…4096}, K∈{512…6144}, 6 repeats each bit-identical
torch.mm bf16 accuracy vs fp32 reference, flagos vs CPU, M∈{1,256,556,696,782,1024}, K∈{1024,2048,6144} identical max error (ratio 1.0×)
torch.bmm, F.linear, addmm bit-identical
torch.softmax on (6,556,556), (2,780,780), (6,1024,1024), (6,556,1024) bit-identical, ≤6e-8 vs CPU
F.layer_norm / F.rms_norm bf16 on (139/556/780/1668, 1024) bit-identical, matches CPU
F.conv1d, torch.cumsum, F.scaled_dot_product_attention, exp, sum(-1), mean(-1), pad, gather, sort, roll bit-identical
A synthetic 4–8 block transformer (LayerNorm + q/k/v + softmax(qkᵀ/√d) + o + silu MLP), bf16, 1024 hidden, seq 256–1024, 5 repeats bit-identical
Vendor torch, same real model, 3 repeats bit-identical
Adding torch.flagos.synchronize() after every block still non-deterministic (0.99409 / 0.86154 / 0.99427)
fp32 instead of bf16 still off (s12 = 0.6879 vs 0.9124)
torch.tensor([<0-d flagos tensor>]) works fine on this backend (it does not on DCU, for reference)

What I suspect / what would help

  1. Stale or partially-initialised output buffer. Simple ops are stable under repetition, but a long alternating sequence of different kernels is not — consistent with an allocator/stream reusing device memory before the producing kernel has retired (despite synchronize() between blocks), or a kernel that does not write every output element. Can you add a debug mode that forces per-op synchronisation and disables memory pooling (e.g. FLAGOS_DETERMINISTIC=1) so we can bisect without guessing?
  2. Numerical divergence in the vision tower. A ~2.7 %-of-std error over 24 blocks is much larger than bf16 should give. If the Iluvatar backend routes some of those ops to lowered/approximated implementations (e.g. matmul with reduced accumulator precision, or a softmax/layernorm variant using fast math), that would explain both the magnitude and the amplification.
  3. The failure is input-length dependent: the 4 bad samples (551–782 tokens, grid_thw up to [1,6,46]) are the longest. Short samples (222–434 tokens) are all fine.

What I can run for you

The machine is idle and I have the weights. Please tell me which of these is most useful:

  • FLAGOS_LOG=dispatch dump of the op routing for a single failing forward, so you can see the backend chosen per op;
  • per-layer dumps (I already have a harness that hooks every block and dumps mean/std/absmax per layer, vendored vs Torch-FL, JSON out);
  • a bisect over the vision tower only (24 blocks, feed the same token tensor through block k repeatedly);
  • binary-search the op set by pinning individual ops to another backend via FLAGOS_OP_<name>=..., if the Iluvatar config supports it;
  • a reduced standalone repro if you can suggest a public Qwen3.5-VL-shaped checkpoint that triggers it (the two internals I cannot share are the weights and the exact prompt data).

Minor environment notes (may be useful for the Iluvatar guide)

While bringing this up on iluvatar-186 I hit four things that cost time; none of them are the bug above:

  1. docs/vendors/iluvatar/installation.md mounts /usr/local/corex-4.5.0/lib64/libcuda.so.1; on this host /usr/local/corex actually symlinks to corex-4.4.0, so the container's bundled libcuda.so.1 must be overridden with the host's 4.5.0 one, otherwise torch.flagos.device_count() returns 0.
  2. In the CI image /usr/local/corex/bin/ixsmi is a 0-byte placeholder; ixsmi prints nothing until the host binary is bind-mounted over it.
  3. libixthunk.so only exists under the host CoreX 4.5.0 tree, so LD_LIBRARY_PATH must include it or import torch_fl dies with ImportError: libixthunk.so.
  4. The image ships transformers 4.57.6 + huggingface_hub 2.x; to use a newer transformers you also need huggingface_hub>=1.5,<2, httpx, typer, psutil, accelerate and a matching torchvision (transformers >= 5.10.4 is required for Qwen3_5ForSequenceClassification).

Activity

  1. lvyufeng commented on Oct 22, 2026

    @lvyufeng
    Collaborator

    Investigated this from the source side. Summary: the reported facts already rule out everything this repository can inspect, and what is left needs a measurement on your box. #637 adds the instrument; this comment is what was checked and what to run.

    Ruled out in source

    • A missing barrier. torch.flagos.synchronize() is cudaDeviceSynchronize() on iluvatar (csrc/runtime/accelerator/iluvatar/device.cc). Your "synchronize after every block does not fix it" is therefore real evidence, not a weaker variant of the test. Only one stream exists: every vendor kernel takes IluvatarCurrentStream(device), and FlagGems/Triton read the same registry (torch_fl/accelerator/iluvatar/ilu_stream.py:266-271, csrc/runtime/accelerator/iluvatar/stream.cc:98).
    • An allocator double-handout. find_free_block erases from the pool (caching_device_allocator.cc:304) and both re-insert sites are !allocated-guarded (:262-283, :465).
    • Our FlagGems fast-launch path. _flaggems_fast_run falls back to the stock launcher for tuner constants, heuristic constants, launch hooks and callable grids, so it cannot launch with a wrong constexpr.
    • index_put_ staying on the CPU. That is real on your wheel but fix: install the index_put FlagGems route on every platform it can serve #636 fixes it, and it is deterministic, so it is not the run-to-run variation.

    One thing that is real and is not the nondeterminism

    bf16 mm/bmm/addmm take a CPU round-trip on iluvatar. IluvatarSupportsGemmDtype accepts only kHalf and kFloat (csrc/aten/backends/iluvatar/iluvatar_common.h:257-259), so a bf16 GEMM runs at::mm(self.cpu(), mat2.cpu()).to(self.device()). That matters for your report in a specific way: your vendor-torch comparison is not like-for-like. Vendor torch runs bf16 iXBLAS on device; Torch-FL runs an fp32 CPU GEMM and rounds. That alone can account for a stable part of the 2.7%-of-std gap, separate from the run-to-run variation. It is worth re-measuring that gap with an fp32 model on both sides before reading it as a numerical defect — you already have an fp32 datapoint (s12 0.6879 vs 0.9124), and if fp32 is also far off, the gap is not the bf16 CPU path.

    What is left, and what #637 does about it

    Eager differing only on run0 while sdpa differs on every run is the sharpest clue, and it has two shapes that produce identical symptoms and need different fixes:

    1. a kernel never wrote part of its output, or
    2. a kernel read memory it should not have.

    With a normal allocator these are indistinguishable, because an unwritten element reads back the previous block's contents — a plausible number. #637 adds FLAGOS_MEM_POISON, which fills every handed-out block with 0xFF (= NaN in every float dtype) before your kernels run. An element no kernel wrote then reaches the output as a NaN that no arithmetic on finite inputs could produce.

    What to run

    pip install git+https://github.com/lvyufeng/PyTorch-FL@feat/mem-poison-diagnostic   # or cherry-pick
    python tests/manual/probe_nondeterminism.py --repeats 6 --json-out probe.json

    Three arms, each in its own process (allocator, tuner cache and Triton cache are all process-global). The verdict maps to a next step:

    Verdict Meaning
    NO_REPEAT_VARIATION the path is not the bug
    UNWRITTEN_OUTPUT a kernel left elements unwritten — grid, mask or tail
    STALE_OR_ALIEN_READ a read of memory this op did not own
    REPEATS_DIFFER_NONDETERMINISTIC_ARITHMETIC atomics / split-k / tuner picking different configs

    The default workload is a synthetic RMSNorm/GEMM/softmax chain, so it answers whether the toolchain has this class of bug. For your model, the more decisive run is the same arms on the vision tower — please paste probe.json plus the two vision-tower mean/std pairs for the poison and no-poison arms, run back to back:

    FLAGOS_MEM_POISON=0 python vision.py   # your existing 6-repeat harness
    FLAGOS_MEM_POISON=1 python vision.py

    If the poison arm's output contains NaNs, the model's own kernels are leaving writes out and the per-op bisect can find which one. If it does not, but the two arms still differ, an op is reading memory it does not own. If the two arms agree exactly, the arithmetic is order-dependent and the next probe is config pinning, not memory.

    I could not run any of this — no BI-V150 here and no built extension on this host — so the probe's device arms are unexercised. If it fails to run for you, paste the traceback and I will fix it.

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