feat: ZTENSOR_DETERMINISTIC mode -- deterministic GPU training runs (T4.1)#179
Open
dndungu wants to merge 4 commits into
Open
feat: ZTENSOR_DETERMINISTIC mode -- deterministic GPU training runs (T4.1)#179dndungu wants to merge 4 commits into
dndungu wants to merge 4 commits into
Conversation
Add internal/cuda.DeterministicEnabled(), read once from ZTENSOR_DETERMINISTIC=1 at process init, mirroring the existing ZTENSOR_ARENA_POISON pattern. Off by default. Also add SetDeterministicEnabledForTesting so other packages (cublas, compute) can flip the flag in unit tests without a real process environment. This is the scope-gate for the deterministic-reductions debug mode (plan-gpu-training-hardening.md T4.1): the nondeterminism inventory and per-op disposition are documented in docs/design.md.
…TENSOR_DETERMINISTIC (T4.1) CreateHandle now sets CUBLAS_WORKSPACE_CONFIG (if the process has not already set one) and calls cublasSetMathMode(CUBLAS_PEDANTIC_MATH) on new handles when ZTENSOR_DETERMINISTIC=1 -- the two documented NVIDIA/PyTorch levers for bit-reproducible cuBLAS GEMMs: they disable TF32 tensor-core downcasting and any workspace-dependent split-K/atomics reduction algorithm cuBLAS's heuristic might otherwise pick. cublasSetMathMode is resolved best-effort at load time so a cuBLAS build stripped of it cannot break ordinary (non-deterministic-mode) BLAS loading; both levers degrade to a stderr warning rather than a hard failure if unavailable, since the handle is still needed outside the debug mode.
…RMINISTIC (T4.1) fused_encoder_bwd.cu's dScale/dBias LayerNorm-style bias-gradient accumulation uses atomicAdd across row blocks -- order-dependent, with no deterministic variant. Rather than silently return order-dependent gradients when ZTENSOR_DETERMINISTIC=1 is set, FusedEncoderBackward now checks the flag first and returns a descriptive error naming the exclusion. This is not reachable from zerfoo's current PatchTST training path (the Go-side wiring is a stub that always falls back to the unfused per-op backward, per the T135.4 audit / issue #522), so the exclusion does not block the GB10 bitwise-identity proof on that model; it will block any future consumer that wires the fused path in, on purpose, until a deterministic dScale/dBias reduction exists.
Document the deterministic-reductions debug mode's nondeterminism inventory and per-op disposition: CPU reductions and GPU custom kernels were already deterministic (T135.2's audit found no cross-block atomics in the reduction path); cuBLAS GEMMs are now routed to a deterministic configuration; fused_encoder_bwd.cu's dScale/dBias atomicAdd path is a documented, honest exclusion that errors under the flag instead of silently returning order-dependent gradients. Also records what the mode does NOT guarantee (cross-GPU/driver/cuBLAS- version identity, CPU/GPU parity) and its cost (CUBLAS_PEDANTIC_MATH forgoes TF32 tensor cores).
Merged
4 tasks
Contributor
Author
|
GB10 bitwise proof landed: zerfoo/zerfoo#936 (2 pods x 2 processes, 3/3 epoch losses bitwise-identical under ZTENSOR_DETERMINISTIC=1; bits 0x3fe8024d7e12aa6c / 0x3fe066b05c164401 / 0x3fd18a52c2f44ad5; pods zerfoo-validate-wave3taskT13-1783115032 and -1783115121). Evidence in zerfoo docs/devlog.md 2026-07-03. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Implements the env-gated deterministic-reductions debug mode from zerfoo
docs/plan-gpu-training-hardening.mdT4.1 (zerfoo plan.md T135.5).ZTENSOR_DETERMINISTIC=1(read once at process init, mirroringZTENSOR_ARENA_POISON; off by default):internal/cublas.CreateHandlesetsCUBLAS_WORKSPACE_CONFIG(if theprocess has not already set one) and calls
cublasSetMathMode(CUBLAS_PEDANTIC_MATH)on new handles -- the twodocumented NVIDIA/PyTorch levers that disable TF32 downcasting and any
workspace-dependent split-K/atomics reduction algorithm. Best-effort:
degraded determinism warns on stderr, never fails handle creation.
fused_encoder_bwd.cu'sdScale/dBias accumulation uses
atomicAddacross row blocks with nodeterministic variant;
compute.GPUEngine.FusedEncoderBackwardrefusesto run under the flag (descriptive error) instead of silently returning
order-dependent gradients.
docs/design.md("ZTENSOR_DETERMINISTIC scope"):a full nondeterminism inventory with per-source disposition (deterministic
already / routed to deterministic config / NOT covered), what the mode
does NOT guarantee (cross-GPU/driver/cuBLAS-version identity, CPU-vs-GPU
bit parity), and the cost (PEDANTIC math forgoes TF32 tensor cores).
Notably, the inventory found most of the training path was already
deterministic run-to-run: CPU fp32 reductions are unconditionally fixed-order
since T135.2, and the custom GPU kernels (softmax, RMSNorm, GEMV family,
flash attention/decode, fused AdamW) reduce via warp-shuffle/tree patterns
that are pure functions of input shape, with no cross-block atomics.
GB10 proof
The bitwise double-run proof runs through zerfoo's harness (ztensor CI has
no GPU): zerfoo PR (branch
wave-3-task-T135.5) addsTestPatchTSTTrainGPUDeterministicDoubleRun, which re-execs itself as twochild processes with
ZTENSOR_DETERMINISTIC=1set at process init, trainsidentically-seeded PatchTST models, and asserts per-epoch losses are
bitwise-identical (exact float64 bit patterns). Evidence recorded in
zerfoo's devlog with pod names and bit patterns.
Test plan
go build ./...,go vet(no new findings), gofmt clean on touched filesgo test ./internal/cuda/... ./internal/cublas/... ./compute/...greenSetMathModenil-handle validation;workspace-config set-when-unset and never-overwrites-operator-value;
FusedEncoderBackwardrefuses under flag / falls through without itTestCreateHandle_DeterministicMode(GPU) exercises the livecublasSetMathMode(CUBLAS_PEDANTIC_MATH)path -- runs on the GB10 viazerfoo's validation harness (skips without CUDA)
wave-3-task-T135.5Refs zerfoo docs/plan-gpu-training-hardening.md T4.1 / zerfoo plan.md T135.5.