Repository navigation
fix(cuda): honor explicit cuDNN algorithms and accumulate half convolutions in float - #14
Merged
Merged
Conversation
jamesbrink
added a commit
to utensils/mold
that referenced
this pull request
Sep 8, 2026
…in (#1658) Moves the candle pin from `c34de9b7` to `bf2cd29a` (utensils/candle#14), which carries two cuDNN convolution corrections mold's code and docs already assume. ## The problem PR #1620 shipped the consumer side of the Hunyuan3D paint campaign's cuDNN work — `rrdbnet.rs`'s explicit algorithm selection, the `.claude/rules/inference.md` invariant, the changelog note, and the campaign qualification record — but pinned candle at `c34de9b7`, which contains neither correction. The pin never pointed at the revision the work was qualified against, so both behaviors were inert on every CUDA build since. `docs/qualification/hunyuan3d-campaign.md:1082` states "all dependency declarations, both Candle-containing lockfiles, and H3's backend provenance point to that commit". They did not. ### 1. The explicit cuDNN algorithm dispatched nothing `rrdbnet.rs:145` asks for `ImplicitPrecompGemm` on Real-ESRGAN's first convolution to avoid an im2col rounding difference that, in its own comment, "amplif[ies] through all 23 RRDBs". At the pinned revision the CUDA dispatch arm never read the field: ```rust if crate::cudnn_policy::is_enabled() && cudnn_is_worth_it_2d(params, self.dtype()) { ``` That convolution is 3-channel with a 3x3 kernel, so its column buffer is `27 x H x W` — 7.1M elements at the default 512-pixel tile, against a 35M half-precision threshold. It cannot clear it at any supported tile size, so the request was discarded on every run. The upscaler runs F16 on GPU (`upscaler/engine.rs:636`), so the second defect compounded on the same path. ### 2. Half convolutions accumulated in half Torch promotes HALF convolution compute to FLOAT while keeping HALF storage on both of its cuDNN routes (`aten/src/ATen/cudnn/Descriptors.h:202-205`, `native/cudnn/Conv_v8.cpp:124-129`). Candle passed `f16` as the compute type; a tensor-op math mode does not promote the compute descriptor. ## Verification Candle's CI is CPU-only `cargo check` and never reaches these paths, so the suite was run on hardware — 4x NVIDIA L40S, driver 595.71.05. Against the **old pinned** `c34de9b7`, with only the new tests applied and no source changes: ``` assertion `left == right` failed: explicit algorithm was ignored below the threshold left: 0 right: 1 F16 cuDNN accumulation relative deviation: 0.010695187 test result: FAILED. 9 passed; 3 failed ``` `left: 0` is the substantive result: an explicitly requested algorithm dispatched cuDNN zero times. The `.010695187` deviation reproduces the figure recorded at `docs/qualification/hunyuan3d-campaign.md:630` when this was first qualified. Against the **new pin**: `test result: ok. 12 passed; 0 failed`. In this repo: `cargo check --workspace --all-targets` clean, `scripts/tests/candle-single-identity.sh` passes on all three lockfiles. ## Blast radius `ConvPolicy::Video` (Wan, LTX) and `ConvPolicy::Paint` select cuDNN; `ConvPolicy::Image` stays on im2col. So convolution numerics move toward the Torch reference for video VAE decode and Hunyuan3D paint on CUDA builds, and still-image generation is byte-for-byte unchanged. ## Not closed by this The paint qualification gates were recorded against these commits, so this restores the configuration they were measured under — it does not re-run them. The real512 VAE bound was still failing in those notes and remains open. ## Docs corrected here - `docs/architecture/candle-extension.md:107` named the old revision as "the current compatibility source" — the only stale SHA left in the repo. - `docs/qualification/hunyuan3d-campaign.md` asserted twice that "all dependency declarations, both Candle-containing lockfiles, and H3's backend provenance point to that commit", naming two Nix hashes that were never shipped. Those claims held on the campaign branch only. Rather than rewrite an evidence record, a dated correction is appended to each. ## No changelog fragment `changelog.d/hunyuan3d-campaign.md` is still unassembled and already carries both notes ("CUDA half-precision convolutions", "Real-ESRGAN cuDNN dispatch"). Neither has reached a released section, so no user ever saw the broken claim, and a "Fixed" bullet here would contradict two "Added" bullets in the same release. Labelled `skip-changelog`; this PR makes the existing notes true. ## Nix hash `sha256-yCwX+XdtJKY6oXp++9d9kw6opSJlRHV69Khc9LwaSQo=`, computed with `nix store prefetch-file --unpack` on the GitHub archive. The method was validated first by reproducing the outgoing `sha256-gInEf6…` for `c34de9b7`. https://claude.ai/code/session_01S1PVcju7evhKq3t9tEFXiW
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
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.
Two cuDNN convolution corrections from the Hunyuan3D paint campaign. Both were
qualified on branch and referenced by mold's checked-in campaign evidence, but
never merged here — so the revision mold pins (
c34de9b7) cannot perform thebehavior mold's code and docs already assume.
de478c1c— accumulate half convolutions in floatTorch promotes HALF convolution compute to FLOAT while keeping HALF tensor
storage, on both of its cuDNN routes:
aten/src/ATen/cudnn/Descriptors.h:202-205(v2.5.1)aten/src/ATen/native/cudnn/Conv_v8.cpp:124-129Candle passed
f16as the compute type for both Conv1D and Conv2D. Setting atensor-op math mode does not promote the compute descriptor — the accumulation
dtype is a separate choice — so half convolutions summed in half.
Tensor storage, the alpha/beta ABI, and every other dtype are unchanged.
5c74a518— honor an explicit algorithm below the size heuristicconv2d_with_algoplumbedcudnn_fwd_algoall the way to the CUDA backend,where the dispatch arm then ignored it:
An explicit algorithm is a numerical choice, not just a performance one — the
im2col shortcut sums differently — so a caller asking for one below the
automatic threshold silently got im2col instead. The heuristic still decides
when no algorithm was named, and the thread policy still wins over both: a
disabled policy takes im2col even with an explicit algorithm, and an
unsupported launch still falls back.
Verification
CI here is CPU-only
cargo check, so it never reaches these paths. Ran thesuite on real hardware instead — 4x NVIDIA L40S, driver 595.71.05, cuDNN 9.x:
Three of those are new and assert actual cuDNN dispatch (not merely that the
feature compiled):
cudnn_f16_convolution_preserves_float_accumulationcudnn_f16_conv1d_preserves_float_accumulationexplicit_algorithm_bypasses_size_heuristic_but_not_disabled_policyApplying those three tests alone to this branch's merge-base — the currently
pinned
c34de9b7, no source changes — reproduces the exact failures they werewritten against:
left: 0is the substantive point: below the threshold an explicitly requestedalgorithm dispatched cuDNN zero times. The
0.010695187deviation reproducesthe figure recorded when this was first qualified.
https://claude.ai/code/session_01S1PVcju7evhKq3t9tEFXiW