fix(CUDA): fmt=8/fmt=6 scale-byte accounting — rebase of #1100 by @monotophic - #1114
Merged
Conversation
coli_cuda_tensor_free recomputed freed scale bytes as O*ng*sizeof(float), but coli_cuda_tensor_upload charges tensor->scale_count*sizeof(float) -- and for fmt=8, scale_count is ((O+127)/128)*ng block-scales, not O*ng. Free therefore overcounted, the tensor_bytes >= bytes guard silently declined the subtraction, and coli_cuda_stats' diagnostic counter stuck non-zero forever after freeing any tracked fmt=8 tensor. Diagnostic-only: cudaFree still runs unconditionally, so there was no real leak. Fix: free() now charges tensor->scale_count*sizeof(float) directly -- the exact expression upload uses -- so the two can't drift independently again. Verified this reproduces the compressed-path charge (archive_bytes + O*sizeof(float)) and the fmt=4-grouped charge unchanged. Added a byte-accounting regression to test_fp8_cuda.cu: create+free a tracked fmt=8 dense tensor and a fmt=4-grouped tensor, assert coli_cuda_stats returns to its pre-upload value exactly. Confirmed FAILS on the pre-fix code (count 0->0 bytes 40432->38152) and PASSES with the fix, plus full `make cuda-test` on spark1 (NVIDIA GB10, sm_121). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…scale bytes coli_cuda_tensor_bytes() used the same buggy O*ng scale-byte shape that 44e1d48 already fixed in tensor_free(): for fmt=8 the real footprint is scale_count = ((O+127)/128)*ng 128x128-block scales, not O*ng, and the bare `tensor->fmt ?` guard (vs. upload/free's `fmt && fmt != 6`) charged fmt=6 a phantom scale buffer it never allocates. Replaced with the same scale_count*sizeof(float) expression upload and free already use, so all three can no longer drift independently. Not diagnostic-only: c/colibri.c's two callers (repin refresh accounting and initial VRAM-tier placement) feed the value into m->gpu_expert_bytes and the per-device `remaining[]` capacity gate that decides how many experts get placed on GPU and which device -- a real, if conservative (under-placement, not OOM), budget decision. Extended test_fp8_cuda.cu with a tensor_bytes() regression: fmt=8 dense against hand-computed weight_bytes+scale_count*4 (exact equality), fmt=6 (must equal weight_bytes alone, scale_count==0), and fmt=4-grouped as a regression (scale_count=O*ng, value unchanged by this fix). Verified on spark1 (NVIDIA GB10, sm_121, CUDA 13.0): reverting only this hunk reproduces the failure (got 63744 want 61464), restoring it passes; full make cuda-test (COLI_GPU=0) passes all four stages, make_rc=0. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…tprint oracle Wire tests/test_fp8_cuda.cu into the cuda-test recipe (compiled alone like test_fp8_warp_cuda.cu since it includes backend_cuda.cu, run before the MXFP4 test) so the fmt=8 regression cases actually execute. Add an independent assert to the tensor_bytes fmt=8 check: expected footprint computed from quant.h's documented block layout (O*I e4m3 bytes + ceil(O/128)*ceil(I/128) f32 scales) rather than the struct fields the accessor itself reads.
Owner
Author
|
Deferred to the next release (v1.7.1/v1.8.0) — not on merit, on infrastructure. This PR is 21/22 green; the one job left is Rather than hold v1.7.0 — which ships a sixth engine with its GPU tier — on NVIDIA's CDN, this goes in the moment the runners recover. The branch is already rebased and ready; it needs a green, nothing else. Follow-up on our side: your PRs are the argument for putting |
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
2 participants
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.
@monotophic's #1100, rebased onto
devby the maintainers — commits and authorship theirs. It went conflicting the moment #1098 (their own__syncthreads()fix) merged, since both add a test to thecuda-testtarget.Conflict resolution: pure union, both mechanical — the
cuda-testdependency line now lists bothtest_absorb_determinism.cu(from #1098) andtest_fp8_cuda.cu(this PR), and both recipe bodies are kept in their intended order (fp8 before MXFP4, absorb-determinism last, per each test's own comment about schedule sensitivity).Fix and its regression test unchanged. Targeting v1.7.0. Supersedes #1100 on merge.