Skip to content

fix(CUDA): weights_owned before the H2D copy — rebase of #1108 by @monotophic - #1115

Merged
JustVugg merged 3 commits into
devfrom
fix/cuda-weights-owned-rebased
Aug 19, 2026
Merged

JustVugg merged 3 commits into
devfrom
fix/cuda-weights-owned-rebased

Conversation

@JustVugg

Copy link
Copy Markdown
Owner

@monotophic's #1108, rebased onto dev by the maintainers — commits and authorship theirs. It conflicted once #1098 and #1101 landed (all add a test to the same cuda-test dependency list).

Conflict resolution: pure union on c/Makefile — the cuda-test target now lists test_absorb_determinism.cu (#1098), test_fp8_cuda.cu (#1114) and test_weights_owned_cuda.cu (this PR), each recipe body preserved in its intended order. Fix and regression test unchanged.

Targeting v1.7.0. Supersedes #1108 on merge.

@JustVugg

Copy link
Copy Markdown
Owner Author

@monotophic — disposition of all eight of your PRs, and a thank-you that should not wait for the release notes.

In v1.7.0:

Deferred to v1.8.0, on size not merit: #1102 (fmt=8 on the kv_b absorb path, +548), #1105 (FP8 container mint, +3007), #1107 (logprob channel + token-ID intake, +818). All three are features that deserve a real review rather than a release-eve skim; the release was already cut when they arrived. They are at the front of the queue after the tag.

The rebases were ours to do: your three CUDA fixes each add a test to the same cuda-test dependency list, so merging the first one conflicted the other two. Pure union in every case, fixes and tests untouched — but review the resolutions if you want, they are stated in each PR body.

And the part that matters: eight PRs in thirteen hours, every one of them carrying its own regression test — test_absorb_determinism.cu, test_fp8_cuda.cu, test_weights_owned_cuda.cu, test_798_guards.c, test_dup_name_refusal.c — and every one green on first submission. The __syncthreads() report that started this came with a three-site table and compute-sanitizer racecheck output, which turned an investigation into a five-minute fix. That is not the normal standard for drive-by contributions to a C codebase, and this repo is measurably better for it. Thank you.

@JustVugg

Copy link
Copy Markdown
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 CUDA syntax check, which downloads the CUDA network installer from NVIDIA's CDN, and that download has been hanging on GitHub's runners all afternoon (three separate jobs stuck on third-party downloads today — the Vulkan job's apt install has the same problem). Retrying twice changed nothing, because the stall is not ours.

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 timeout-minutes and toolkit caching on these jobs (#953), so a slow mirror fails fast instead of blocking a release.

…ee their buffer

coli_cuda_tensor_upload marked ownership only after the weight memcpy
succeeded, so a malloc-success/memcpy-failure upload freed the tensor
while weights_owned was still 0 and coli_cuda_tensor_free's ownership
gate skipped the cudaFree: every failed upload leaked weight_bytes of
device memory, invisible to the ledger (tracked=0). Ownership is a fact
of the allocation, not the copy — set it between the two.

tests/test_weights_owned_cuda.cu pins the contract by including the
backend directly and renaming its cudaMemcpy call sites to a hook that
fails H2D copies of the target byte count: 8 injected create-fail
cycles must leave cudaMemGetInfo free memory unchanged (previously a
monotonic ~64 MiB/cycle loss), the success-path control must round-trip
to zero delta, and the ledger must never see a failed upload. Wired
into cuda-test and gpu-compile alongside the other direct-include
tests.
…h's mapped surface

The hook named cudaMemcpyKind and cudaErrorInvalidValue, which the compat
header does not alias (the backend never uses them), so the cuda-test
recipe under HIP=1 failed to compile. Compat's precedent is to map only
what the backend needs; alias the kind type and the injected error value
locally per vendor instead of widening the product header for a
test-only seam. CUDA behavior unchanged.
… the HIP build

backend_gpu_compat.h maps only the CUDA runtime names backend_cuda.cu
itself uses; the injection hook's kind type and injected error value are
not among them, so hipcc -- which compiles this file in the HIP
syntax-check lane via gpu-compile -- finds them undefined. Alias the two
locally under the HIP arm and keep the shared hook body in the cuda
spelling, the arrangement the repo's other direct-include tests use for
their off-surface names. The hook's forward declaration also no longer
leans on hipcc pre-including the runtime header: the HIP arm now includes
<hip/hip_runtime.h> explicitly, the same header backend_gpu_compat.h
pulls, unreachable under nvcc. CUDA behavior unchanged.
@JustVugg
JustVugg force-pushed the fix/cuda-weights-owned-rebased branch from 14a7934 to 97bf97f Compare August 19, 2026 17:30
@JustVugg
JustVugg merged commit e0e3cc4 into dev Aug 19, 2026
22 checks passed
@JustVugg
JustVugg deleted the fix/cuda-weights-owned-rebased branch August 20, 2026 00:57
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants