gpucheck v1.1 — per-(kernel,dtype) MPS overlay + Apple-tile fuzzing + xfail expansion + memory-loop closure - #3
Draft
Akasxh wants to merge 7 commits into
Draft
gpucheck v1.1 — per-(kernel,dtype) MPS overlay + Apple-tile fuzzing + xfail expansion + memory-loop closure#3Akasxh wants to merge 7 commits into
Akasxh wants to merge 7 commits into
Conversation
Per IMPLEMENTATION_PLAN_v1.1.md and AMENDMENTS — this branch is the target for the 27-task v1.1 work. Substantive commits land in subsequent sessions. Audit corpus at .claude/teams/audit/v1.1/.
…pher-v3)
Wires Apple-canonical tile constants (8x8 simdgroup_matrix MMA fragment,
BK=16 GEMM block, FA head_dim=80) into the shape fuzzer behind a new
device_type parameter. CUDA path is bit-for-bit preserved by the default
arg.
- TILE_SIZES_MPS = (8, 16, 32, 64, 128)
- POWER_OF_2_BOUNDARIES_MPS adds {7,8,9}, {15,16,17}, {31,32,33},
{63,64,65}, {79,80,81} on top of CUDA's 127..513 boundaries
- _tile_set / _pow2_boundary_set helpers select per device_type
- fuzz_shapes() and ShapeStrategy take device_type="cuda"|"mps"
- 3 property tests in tests/test_fuzz_shapes_mps_tiles.py:
* 8x8-MMA-fragment boundary surfaces in MPS pool
* head_dim=80 boundary surfaces in MPS pool
* CUDA deterministic pool unchanged (no MPS-only dim leak)
Sources: cartographer-v2.md §S4 (mlx mma.h), §S6 (steel_gemm_fused.metal),
§S8 (steel_attention.metal). Diff drafted in cartographer-v3-fuzzer-patch.md.
…T-25, github-miner-v3)
Adds the Round-2/Round-3 long-tail PyTorch issue audit findings to
[tool.gpucheck.mps.xfail] in pyproject.toml. The block now ships
12 v1 entries plus the v1.1 long-tail batch, organized into 8 tiers
(silent-correctness, OOB-zero, validation-missing, NaN/drift,
copy_/argmax/clamp edges, backward/functorch, rank-5 crashes,
hard crashes & specialty). Each line cites the PyTorch issue
number and last-update stamp so future re-mines have traceability.
New tests/test_mps_xfail_v11_expansion.py covers:
- post-expansion entry total
- op-name uniqueness inside the TOML block
- live registration of the high-value entries
(copy_.strided_view_offset_2pow32_wrap,
binary_ops.uint16_uint32_uint64,
avg_pool1d.prefix_sum_drift_long_seq,
scatter_add_.nonzero_offset_slice,
model.voxtral_asr_full_pipeline)
- is_mps_xfailed("copy_.strided_view_offset_2pow32_wrap") returns True
- apply_mps_xfail_config round-trip from a fresh tomllib parse
Source: .claude/teams/research/v1.0/EVIDENCE/github-miner-v3-xfail-config.md §A.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
…ration (T-24, calibration-final)
Replace the flat per-dtype `_MPS_TOLERANCE_MULTIPLIERS` with a layered
overlay: the v1.0 flat table is preserved verbatim as the
DEFAULT-kernel-class fallback, and a new
`_MPS_KERNEL_DTYPE_MULTIPLIERS: dict[tuple[KernelClass, str], float]`
holds the v1.1 5K-iter Apple-M5 calibrated multipliers.
Multipliers (from `.claude/teams/audit/v1.1/EVIDENCE/calibration-final.md`
and `drift_histogram_5k.json`, sized to cover measured P99 with safety):
- MATMUL fp32/16/bf16 = 16/20/32 (5K P99 = 13.4×/17.7×/27.6×)
- CONV2D fp32/16/bf16 = 4/ 8/12 (5K P99 = 1.98×/6.71×/10.44×;
fp32 P99.9=2.52× crosses FA-2×)
- NORM/REDUCTION/POINTWISE = 2.0× via dtype wildcard rows.
`compute_tolerance` gains an opt-in keyword-only
`kernel_class: KernelClass | None = None`. When omitted, resolution
short-circuits to the v1.0 flat path — backward-compatible byte-for-byte.
When supplied, `_resolve_mps_multiplier()` walks five steps:
exact → wildcard → DEFAULT-class-exact → flat table → hard 2.0.
`KernelClass` is a `str`-Enum exported from `gpucheck.assertions`.
Tests: `tests/test_per_kernel_tolerance_overlay.py` pins each measured
cell, the wildcard rows, the v1.0 backward-compat parity, the fallback
ordering, and the dict-cardinality.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
… pytorch#162872 still open) Adds gpucheck.diagnostics.mps_event_deadlock with three public helpers: - probe_mps_event_deadlock(timeout_ms=2000) -> "healthy" | "deadlocked" | "skipped" Runs the unsafe Event.synchronize + elapsed_time trigger pattern in a daemon=True worker thread; returns the result of the timeout race. - assert_no_event_deadlock(timeout_ms=2000) Thin assertion wrapper that raises RuntimeError on "deadlocked". - mps_event_deadlock_status_fixture(timeout_ms=2000) Generator body for users to wrap in their own session-scoped fixture. The probe is an opt-in companion to the in-tree MPSBackend.event_timer guarantee (which already routes around pytorch#162872 by using device-level torch.mps.synchronize). It exists so users can detect the bug if a future contributor or custom fixture drifts back to the unsafe pattern, and so M-silicon CI can xfail tests that depend on working Event.elapsed_time until upstream lands a fix. Tracer-v3 verified the MPSEvent.mm:230 deadlock geometry against pytorch/main HEAD as of 2026-05-01; PR #162874's one-line fix was closed without merge. Tests cover the three return states (mocked torch via sys.modules shim), the assert wrapper across all states (parametrized), the session-fixture contract, the daemon=True invariant on the leaked thread, and a real-MPS opt-in integration test gated on GPUCHECK_RUN_REAL_MPS_PROBE=1 (the real probe leaks a daemon thread inside libtorch's pthread_cond_wait, which on broken MPS may not release the GIL — tracer-v3 §"Uncovered ground"). Files: - src/gpucheck/diagnostics/__init__.py (new) - src/gpucheck/diagnostics/mps_event_deadlock.py (new, ~234 LOC) - src/gpucheck/__init__.py (lazy-export wiring) - tests/test_mps_event_deadlock_probe.py (new, 9 tests) Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
…1/12/13, mutator-survivors) Add tests/test_mutation_killers.py — a dedicated home for the three high-leverage clusters identified by the v1.1 mutmut audit (EVIDENCE/mutator-survivors.md). 29 new tests, no source changes. T-11. Pin numeric fields in format_mismatch_report (~30 reporting.py mutants). Canonical zeros((4,4)) / eye(4) fixture pins exact rendered strings: max-abs "1.000000e+00", mean-abs "2.500000e-01", mismatch count "4 / 16 (25.00%)", location "(0, 0)", tolerances line "atol=0.00e+00, rtol=0.00e+00", "Error Histogram" panel header, bucket "[1e+0, 1e+1)" with count 4. A non-symmetric (2, 4) fixture also pins (1, 3) vs (3, 1) so any axis-swap in np.unravel_index fails. Substring "NaN" in report (which matched "XXNaNXX") is replaced by ANSI-stripped equality. T-12. Round-trip the tolerance config loader (~17 tolerances.py mutants). apply_config_tolerances → compute_tolerance → reset → compute_tolerance proves the overlay actually applies and the reset clears it. Edge cases cover: missing/empty section returns None, malformed entries (missing atol, missing rtol, non-dict) silently skipped, literal "tool"/"gpucheck"/"tolerances" key path required. T-13. Hard-coded parametrize over _DEFAULT_TOLERANCES (~12 dict-value mutants). The existing test_known_dtypes iterates the dict against itself and is tautological — every mutated value still passes. New parametrize uses a hard-coded ground-truth list of 7 (dtype, (atol, rtol)) pairs and runs both as a dict-literal pin and as a compute_tolerance() return-value pin. ruff + mypy clean on the new file. Full pytest suite still green (330 passed, 3 skipped). Targets ~59 surviving mutants without any source code changes. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
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
This PR opens the gpucheck v1.1 development branch off post-v1.0 main. Work is planned per the v1.1 audit cycle that ran 2026-05-07.
Plan
Two binding documents at the repo root:
IMPLEMENTATION_PLAN_v1.1.md— engineering-lead's 27-atomic-task plan across 5 phases (~6h with 4-way parallelism)IMPLEMENTATION_PLAN_v1.1_AMENDMENTS.md— skeptic's C1-C5 conditions formalised + adversary's 3 re-source itemsAudit corpus
.claude/teams/audit/v1.1/contains 18 specialist evidence files + 17 summaries:Key v1.1 deliverables
.claude/teams/audit/v1.1/drift_histogram_5k.json(105K measurements on M5)Test baseline
v1.0.0 ships at 253 tests passing, 2 skipped. v1.1 should land ≥30 net new tests.
Status
🟡 Draft / scaffolding. Substantive commits land in subsequent sessions.
🤖 Generated with Claude Code