fix(conv): bounds-guard all grid_1d convolution kernels + run tests on macOS CI - #77
Conversation
The forward convolution kernels index a flat output element by
`program_id::<0>()` and stored to `out[idx]` with no range check. Test
and production dispatch goes through `grid_1d(n_out, 256)`, which
`div_ceil`s the launch up to a whole threadgroup — so whenever `n_out`
is not a multiple of 256 the tail threads run with `idx` past the last
output, decode a channel/batch index beyond range, and read `input`/
`weight` and write `out` out of bounds. On Apple GPUs that reads/writes
whatever memory neighbours the buffers: nondeterministic inf/NaN.
Surfaced as a flaky `test_depthwise_conv1d_strided [f16]` (max|Δ|=inf,
channels·out_len = 8·15 = 120, 136 tail threads); the other kernels'
existing test shapes happened to be exact multiples of 256 and masked
it. CI never caught the class — the clippy/test job runs on
ubuntu-latest with no Metal GPU, so the correctness harness is skipped
there.
Guard, matching `winograd_conv`'s existing `if idx < total`: clamp tail
threads onto element 0 (always in-bounds) for the reads and skip their
store. Applied to all 11 kernel fns across the convolution family:
depthwise_conv1d, conv2d{,_generic,_grouped}, conv3d{_generic,_grouped},
depthwise_conv2d{,_nhwc}, and the three *_block_scaled variants.
conv2d_grouped and conv3d_{generic,grouped} lacked a `batch` constexpr
needed to compute the bound, so it was added (they are codegen-only with
no production caller; all test/bench call sites updated to pass it).
Adds `test_depthwise_conv1d_odd_tail` (3·17 = 51) as a second, small
non-threadgroup-multiple shape so the guard path stays covered.
Workspace suite: 938 passed / 0 failed (incl. Metal GPU correctness);
clippy + fmt clean.
The `clippy-test` job ran `cargo test --workspace` (via nextest) on
`ubuntu-latest`. But the GPU correctness sweep —
`crates/wh-iron-std/tests/kernel_tests_harness.rs`, the
`all_registered_kernel_tests_pass` test that dispatches every registered
`#[test_kernel]` on Metal and checks it against a CPU oracle — is gated
`#![cfg(target_os = "macos")]`. On a Linux runner it compiles to nothing,
so the workspace test job passed on every PR WITHOUT running a single GPU
kernel. That is how the conv-kernel out-of-bounds bug (this PR's other
commit) reached dev: no PR gate ever executed the harness.
Split the job:
- `clippy` (ubuntu-latest): `cargo clippy --all-targets --all-features`.
`--all-features` pulls in cuda/hip/vulkan, which only compile-check on
Linux, so this must stay off macOS.
- `test` (macos-26): `cargo nextest run --workspace` with default
(Metal) features, so the GPU harness and the codegen-every-benchspec
test actually run on real hardware, on every PR.
macOS runner selection (`runs-on: macos-26`) and the audit-mode
harden-runner match the existing coverage.yml / iron.yml jobs.
TheTom
left a comment
TheTom
left a comment
There was a problem hiding this comment.
LGTM. Verified the load-bearing details rather than trusting the prose:
Guard correctness
select(cond, a, b)argument order: cond-first matches the existing winograd_conv usage (select(row_ok_0, pr_0 - pad_h, 0u32)), so the clamp is right. The new odd-tail tests would hard-fail all dtypes if the order were flipped, and the harness is green, which independently confirms it.- Per-kernel bounds check out:
depthwise_conv1dcorrectly omitsbatch(no batch dim in its flat index), the NHWC variant's bound matches its channel-fastest layout, and the block_scaled variants' scale reads are also safe at the clamped index 0. - Clamp-reads-plus-guarded-store is the right idiom here vs early return; tail threads do dead arithmetic but that is negligible against the conv inner loop.
Signature changes (batch constexpr on conv2d_grouped / conv3d_generic / conv3d_grouped)
- Checked butter: no production callers, only OpsCoverageNotes mentions. The emitted wrapper signatures will change on butter's next
make regenerate-kernels; mechanical, but worth a line in whichever butter PR does the regenerate. That regenerate also needs #78's BK32 kernel in the tree (see the cross-repo note on butter #76), so one combined regenerate after both land is the clean path.
CI split
- This is the important half of the PR. Real-world confirmation from the other direction: #78's current "Clippy + Tests" green never dispatched a kernel, exactly the bypass described here. Suggest landing #77 first, then I'll rebase #78 so its GPU harness actually runs on the new macOS job before merge.
runs-on: macos-26precedent confirmed (coverage.yml, iron.yml).- One watch item: the old comment noted the codegen test alone took ~14 min and clippy+tests no longer fit in 15. macOS runner + build + nextest + GPU harness under a 30 min timeout could be tight as the kernel registry grows (each new #[test_kernel] adds to the sweep). If the first few runs land near the limit, bump to 45 rather than letting it flake.
Also lines up with the F-85 lesson from the Laguna campaign: the dual-slab BK64 body was unit-green but empty-gen in the full model. Getting the harness onto every PR closes the cheapest half of that gap.
…e MPP kernel
Moving `cargo nextest run --workspace` to the macOS CI runner (previous
commit) surfaced a real failure the ubuntu job hid: the `macos-26` GitHub
runner's Metal compiler rejects `iron_gdn_wy_plan` at PSO creation —
unsupported deferred-static-alloca-size function body in
metal::cooperative_tensor<...> mpp::tensor_ops::__mutmul2d ...
The kernel lowers a *dynamic-extent* `mpp::tensor_ops::matmul2d`
cooperative tensor (`extents<i, SIZE_MAX, SIZE_MAX>`), which that
toolchain image can't compile. It is a valid kernel — it builds and runs
on every dev machine, verified locally on an Apple7 M1 Max (family < 10)
through the M5 Max — so this is a toolchain gap, not a kernel or
numerical bug.
Gate on the ACTUAL blocker, not GPU family: `run_pipeline` now returns
`Option`, mapping the plan kernel's `IronError::PipelineCreation`
(matched on the `cooperative_tensor` / `deferred-static-alloca` reason)
to a skip; every other dispatch error still panics. Each test skips via
`let Some(..) = run_pipeline(..) else { return }`.
Deliberately NOT `skip_unless_apple10` (as the sibling moe_mpp tests use):
those kernels need Apple10 hardware features, but `iron_gdn_wy_plan` runs
on Apple7 — a family gate would wrongly disable this test on M1/M2/M3 dev
hardware where it passes. The toolchain probe skips only where the kernel
genuinely cannot be built, so it still runs for real everywhere it can.
…t build
The macOS test job surfaced more of the same class as the gdn-wy fix: the
GitHub `macos-26` runner's Metal toolchain rejects the whole MPP / bgemm
cooperative-tensor kernel family at runtime PSO creation ("unsupported
deferred-static-alloca-size ... cooperative_tensor"). `iron build`'s
offline `xcrun metal` compiles them fine; only the runtime
`makeLibrary(source:)` JIT path fails. The kernels are valid and run on
every dev machine (verified from an Apple7 M1 Max through the M5 Max).
Shared, precise gate — `common::is_unsupported_coop_tensor` (matches the
unique `deferred-static-alloca` marker, so it never masks a real
numerical failure):
- `kernel_tests_harness.rs` (`all_registered_kernel_tests_pass`): skip and
count such kernels instead of failing the whole harness, so the non-MPP
correctness sweep — conv (this PR's fix), norm, sdpa, sampling, … — still
gates every PR on real GPU. Reports the skip count.
- `moe_bm64_ragged_correctness.rs`: `run_bm64` returns `Option`, skipping
when its bm64 MPP kernel can't build.
Also `--no-fail-fast` on the macOS nextest run so a single ~15-min job
reports every failure rather than cancelling on the first — cheaper than
iterating one straggler per round.
Verified locally on Apple7 (M1 Max): both still RUN and pass (the skip
path is inert wherever the toolchain can build the kernel).
|
Reviewed the two follow-up commits (59357b1, 1800479), approval stands. The gate is precise: matches only the unique deferred-static-alloca marker, every other dispatch error still panics/fails, and the registered-kernel-count assertion still catches link regressions. Skipping on the actual toolchain error instead of a family gate is the right call given gdn_wy runs fine on Apple7. One consequence worth tracking rather than forgetting: on the macos-26 image the entire MPP cooperative-tensor family is now skipped, so a green Tests (macOS) job no longer implies any MPP kernel ran. The skip count in the log is the only tell. Suggest a tracking issue for re-enabling when the runner's Metal toolchain catches up (or a self-hosted M-series canary), so the skip does not quietly become permanent. |
The `macos-26` image ships Xcode without the Metal Toolchain component,
so anything invoking the `metal` compiler fails:
error: cannot execute tool 'metal' due to missing Metal Toolchain;
use: xcodebuild -downloadComponent MetalToolchain
This broke the Iron workflow's `Build` job (`iron build`'s MSL codegen
shells out to `xcrun metal`) and is the likely root cause of the runtime
`makeLibrary(source:)` failures on the MPP / bgemm cooperative-tensor
kernels in the new macОS test job — the JIT path needs the same toolchain.
Add `xcodebuild -downloadComponent MetalToolchain` before every Metal step:
check.yml `Tests (macOS)`, and iron.yml `Build` + both `Bench` shards.
With the toolchain present the GPU tests should compile and RUN (not skip),
so a breaking kernel change surfaces on the PR instead of hiding.
…ileplan, view_u16) Complete the toolchain-probe fallback across the last two ungated MPP integration files that the --no-fail-fast run surfaced: - moe_gather_qmm_tileplan_correctness.rs (run_tileplan -> Option) - moe_view_u16_correctness.rs (dispatch closure -> Option) Both use the shared `common::is_unsupported_coop_tensor` gate. This is a safety net: with the Metal Toolchain now installed (previous commit) these should compile and run for real; if a given runner still can't build the cooperative-tensor body, they skip rather than red the suite. Verified locally on Apple7 (M1 Max): both RUN and pass — the skip path is inert wherever the kernel builds.
…l is the real fix Drops the toolchain-probe skip fallbacks across kernel_tests_harness, gated_delta_wy_pipeline_integration, moe_bm64_ragged_correctness, moe_gather_qmm_tileplan_correctness, moe_view_u16_correctness, and the shared common::is_unsupported_coop_tensor helper. Those skips were added while the macos-26 failures looked like a permanent toolchain limitation. The real cause was the missing Metal Toolchain component; installing it (`xcodebuild -downloadComponent MetalToolchain`, prior commit) lets every MPP/bgemm cooperative-tensor kernel compile and run at runtime — the last green CI run executed all 937 test binaries with ZERO skips. With the toolchain present the skips are dead code and would only ever mask that one compiler error, so remove them: the GPU suite is now strict again, and any real kernel compile/correctness regression reds the PR immediately.
…-26 runners
Restores the skip fallbacks removed in the prior commit. Removing them was
premature: it was based on ONE green run that happened to land on a
`macos-26` runner whose Metal runtime could JIT-compile the MPP / bgemm
cooperative-tensor kernels. The next run (skips gone) hit a runner that
can't — 596 `deferred-static-alloca` errors across exactly the 5 MPP test
files and NOTHING else (zero real regressions).
The `Build` job (offline `xcrun metal`) passes on both, so the downloaded
Metal Toolchain fixes AOT compilation — but the runtime `makeLibrary(source:)`
path depends on the runner's Metal framework, which is not uniform across
GitHub's macos-26 pool. So:
- without the skips: CI is FLAKY (green or red by luck of the runner);
- with the skips: CI is STABLE — MPP kernels run wherever the runner can
build them and skip (counted, logged) where it can't, while the whole
non-MPP sweep (conv, norm, sdpa, sampling, …) always gates the PR.
The skips match only the unique `deferred-static-alloca` marker, so a
genuine numeric or non-toolchain compile regression still reds the PR.
|
Deep-dive on the cooperative-tensor CI failures. TL;DR: it is not your hardware and not the kernels. Two separate layers, and only one of them is fixable by us. Root cause1. The compile failure you are seeing is an OS-version gap, not a GPU issue. The Demangling the failing symbol makes the trigger visible: Known prior art with the exact same error: metaltile PR #257 hit it on macos-26 runners (26.4 fails, 26.5 passes, on M1 Max no less) and established that pinning DEVELOPER_DIR/Xcode does not help because 2. Separately: even after the image updates, hosted runners will never really execute MPP. The paravirtualized GPU reports a low family, and per the comment in our own For calibration: 922/937 tests passed on the paravirt GPU with the harness doing 450s of real work, so the runner is genuinely useful for everything outside the MPP family. Recommendations
Sources: runner image OS from the job setup log of the failed run; metaltile PRs #257/#258; butter PR #25; runner-images issues #7085 (paravirtual GPU limits) and #14172/#14344 (image toolchain cadence); Apple MPP Programming Guide. |
|
Addendum from the second research pass (hosted-runner GPU capabilities and how peer projects handle this). One correction to my comment above and some concrete options. Correction/nuance: I said PSO compilation should self-heal once the image rolls to macOS 26.5. That is only proven for bare metal (metaltile's data point is a local M1 Max on 26.5). In the VM, the guest talks to How the big Metal projects run CI (verified from their live workflow files):
The common shape is exactly where this PR landed: hosted VM for build + everything-but-the-special-paths, bare metal for the rest. MLX's artifact handoff (build hosted, test self-hosted) is the cleanest template if we add a box. Options for real MPP coverage, ranked: (1) self-hosted runner on an owned Mac mini, restricted to same-repo branches, nightly or label-gated (M4-class covers coop-tensor correctness; M5-class also covers the neural-accelerator paths); (2) MacStadium bare-metal tier (their Orka product is Virtualization.framework VMs, same paravirt GPU, avoid for this); (3) EC2 Mac is bare metal but has a 24-hour minimum host allocation, poor fit for CI; (4) Tart-based providers (Cirrus, Namespace, WarpBuild) are all VZ VMs with the same One cheap hardening: alongside the error-string gate, MLX's exact predicate |
2d7b5c9 to
837c0f2
Compare
Agreed. I think this is the right approach and yes we'll need hosted bare metal runners eventually. I'm updating CI stuff in a follow up PR #80. Will potentially include similar MLX I think this PR is good to go. Thanks for looking into that @TheTom. |
2 participants
What
Two related changes, discovered together:
1. Bounds-guard every
grid_1dconvolution kernel (535403a1)The forward convolution kernels index a flat output element by
program_id::<0>()and store toout[idx]with no range check. Dispatch goes throughgrid_1d(n_out, 256), whichdiv_ceils the launch up to a whole threadgroup — so whenevern_outisn't a multiple of 256, the tail threads run withidxpast the last output, decode a channel/batch index beyond range, and readinput/weight+ writeoutout of bounds. On Apple GPUs that touches whatever memory neighbours the buffers → nondeterministicinf/NaN.Surfaced as a flaky
test_depthwise_conv1d_strided [f16](max|Δ|=inf;channels·out_len = 8·15 = 120, 136 tail threads). The other kernels' existing test shapes happened to be exact multiples of 256, masking it.Fix (matches
winograd_conv's existingif idx < totalidiom): clamp tail threads onto element 0 (always in-bounds) for the reads and skip their store. Applied to all 11 kernel fns:depthwise_conv1d,conv2d{,_generic,_grouped},conv3d{_generic,_grouped},depthwise_conv2d{,_nhwc}, and the three*_block_scaledvariants.conv2d_groupedandconv3d_{generic,grouped}lacked abatchconstexpr to compute the bound, so it was added (they're codegen-only, no production caller; all test/bench call sites updated). Addedtest_depthwise_conv1d_odd_tail(3·17 = 51) as a second non-threadgroup-multiple regression shape.2. Run the test suite on macOS CI (
e9b7293c)This bug reached
devbecause no PR gate ever ran the GPU tests. Theclippy-testjob rancargo test --workspaceonubuntu-latest, but the GPU correctness sweep (kernel_tests_harness.rs,all_registered_kernel_tests_pass) is#![cfg(target_os = "macos")]— on Linux it compiles to nothing, so the job passed without dispatching a single kernel.Split
check.yml's job:clippy(ubuntu-latest):cargo clippy --all-targets --all-features—--all-featurespulls in cuda/hip/vulkan, which only compile-check on Linux, so this must stay off macOS.test(macos-26):cargo nextest run --workspacewith default (Metal) features, so the GPU harness +every_registered_benchspec_codegensactually run on real hardware, every PR.Runner selection (
runs-on: macos-26) and audit-mode harden-runner match the existingcoverage.yml/iron.ymljobs.Verification
cargo test -p wh-iron-std --test kernel_tests_harness— GPU harness green incl. the new tail shapes (rebased on currentdev).cargo fmt --checkclean;actionlintclean on the workflow change.