Conversation
4 tasks
1 task
thynics
pushed a commit
to thynics/cudnn-frontend
that referenced
this pull request
Jul 29, 2026
…robe) Instrumentation-only fork of v12 (NOT a perf candidate). Adjudicates the decisive fork: is the W17 supply chain (~6.6us/tile) intrinsic TMA/L2 latency or ROUND_STAGES credit back-pressure. - W17 wait split: RK_ACQ(i) around ROUTE_K's two kdq credit acquires; MAT_ACQ(m,q) around each quadrant gen's producer_acquire (payload t*8+q); MAT_WAIT(m,q) around each raw round_tma_mbars completion wait (payload t*8+q-1; drain q=7). MAT_QDO envelope kept for continuity. - W19 (IDLE_WARP) observe-only hardware-completion probe: bounded-spin parity reads (in-file _mbarrier_try_wait PTX wrapper, never arrives) on round EMPTY mbars (PROBE_GEN(i,g): 10 per-gen completion edges per grads block) and s_done/dp_done/dkv_done FULL + pds EMPTY mbars (PROBE_DONE(i,k), k=0..4; k=4 pds release == whole-block hardware completion, the ring edge with zero wake latency). Parities are pure index functions: s/dp (t//2)%2 at stage t%2; round (b+g//2)%2 at stage g%2; dkv/pds b%2. Fail-soft: first timeout (PROBE_SPIN_CAP 2^22 polls) permanently disarms the probe; kernel never hangs. - IKET name budget: retired 6 settled math-side spans (WAIT_S, T2R_S, WAIT_dP, T2R_dP, MATH_PDS_ACQ, MATH_BAR1; wrapped ops byte-identical), added the 5 above: 30/31 names. dsa_v2_native span-table comparators will fail BY DESIGN -- raw extraction only, see VM5PROBE_RUNNER_NOTES.md (it also OBSOLETES the earlier mbarrier_try_wait import precheck). - Adversarially reviewed (protocol/parity refuter + DSL/scope refuter + merge): flip order verified against _issue_dq_rounds_v2 and grads head/tail helpers line-by-line; review fix NVIDIA#1 applied (probe mbar pointers hoisted to the kernel prologue -- storage struct must not cross a dynamic-if region). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
thynics
pushed a commit
to thynics/cudnn-frontend
that referenced
this pull request
Jul 30, 2026
Four env-bisectable levers on the v12 base (all-off reproduces v12); composition targets the two hardware-adjudicated ring legs from the vm5probe run2 evidence (MAT_ACQ:MAT_WAIT 70:30 back-pressure; leader per-pass software ~3us/tile is the largest single item; exchange seam ~1.5-1.7us; L2 read hit 95.4% killed the re-warm direction). - DSA_V15_L2X (default 1): P/dS exchange re-staged through an HBM ring (workspace_pds, [slot 2][cta 2][kind 2] x 8KB per token; hot set ~4.9MB at 74 resident clusters). Math publishes whole P/dS images in two waves through a new 8,192B pds_stage funded byte-exactly by deleting p_xchg + the write-dead ds_xchg; W18 becomes the pds producer (acquire moved from math -- the 2.2us MATH_PDS_ACQ park dies) and converts the ring into SMEM operands with five 1D bulk G2S fills, committing the leader's dQ gate only after ds_image bytes LAND (strictly stronger than v12). DSM sends, the count-128 pds_ready handshake and math's ring residence all retire. Leader code unchanged. workspace_pds is a trailing default-None param: pre-v15 harness call sites keep working; L2X=1 without the ring raises at build time (see V15_RUNNER_NOTES.md for the 2-line harness patch). - DSA_V15_REGSWAP (default 2): warpgroup-uniform variants only (setmaxnreg is warpgroup-granular): B = W16-19 uniform 64 / reduce 120; C = W16-19 uniform 56 / gather 40. Pool exactly 61,440 in all settings. - DSA_V15_DQ_MERGE (default 1): both kdq credits waited back-to-back (ROUTE_K acquires AND commits them upfront -- no producer cycle, unlike the condemned dVdK merge), releases staggered per round. - DSA_V15_ALLTMA (default 1): own-half DSM bulk retired via the pre-existing pure-TMA fallback (run2: DSM legs 1.7x slower mean, 2.25x worse p90; also retires the spec-UB own-CTA DSM bulk). Adversarially reviewed (protocol + DSL refuters + merge, zero fatals): byte identity of all three operand buffers re-derived from layouts; wait_group(0) write-completion ordering for the cross-proxy S2G->G2S edge confirmed; review fix NVIDIA#1 applied (gen_ready mbars slot ping-pong -- restores the structural <=1-ahead publish bound v12 had); review fix NVIDIA#2 applied (batch>1 ring aliasing guard). IKET at L2X=1: 28 names (retired MATH_PDS_ACQ/MATH_BAR1/ROUTE_P/ROUTE_dS, added W18_PDS); L2X=0 bisect builds are release/correctness-only (31 names > 29 cap). Expected: realistic 8.5-8.9ms, optimistic 7.7-8.1 (v12 = 11.654, baseline = 8.14). Runner protocol: V15_RUNNER_NOTES.md (v15_run1). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
thynics
pushed a commit
to thynics/cudnn-frontend
that referenced
this pull request
Aug 4, 2026
…17 zone unfrozen) Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
YangXu1990uiuc
added a commit
that referenced
this pull request
Aug 20, 2026
…n both sides of the boundary (#626) * fix(device): ensure the RIGHT context, not merely a context ensure_current_context returned as soon as ANY context was current, so a thread already bound to another GPU's context kept it. The legacy default stream (handle 0) carries no context of its own -- it resolves against whatever is current -- so under a foreign context the work runs on THAT context's GPU, where the pointers are invalid: an async fault at some later sync rather than an error at the launch. A real stream does carry its context and a cross-context launch is rejected outright, so only the stream-0 path is silent, and stream 0 is exactly what torch's default stream is. Resolve the target rather than accept the incumbent: the stream's context when the stream names one, else the caller's device. execute() passes the handle's ordinal, so the FE path no longer asks the runtime which GPU it is on -- Handle.device owns that since #612 -- and cudaGetDevice() stays only as the fallback for a caller that cannot name a device. Cost measured on parley: 59 ns per execute in situ, 0.1% of a 79 us GDN op. test_ensure_current_context.py covers the cold thread, the foreign-device replacement, stream-wins-over-device, and the steady-state no-op. * fix(execute): bind a context on the calling thread in C++ too The same gap on the other side of the boundary. cuDNN's runtime-compiled engines launch through the driver, which reads the CALLING thread's context stack, and a thread that has done no CUDA work has nothing on it. Measured on a matmul+relu+relu graph (which routes to those engines), on a thread where cuDNN is the first CUDA call: before: cold #0/#1/#2 ctx 0x0 -> 0x0 cuCtxGetLimit returned error invalid device context (201) after: cold #0/#1/#2 ctx 0x0 -> 0x3f67f670 OK Deterministic both ways, and one torch op on the thread beforehand hides it entirely -- the CUDA runtime binds the primary context as a side effect, and something normally does, which is why no framework has run into this. The precompiled engines launch with <<<>>> and are unaffected for the same reason. Placed at execute_plan_at_index, which the file already documents as the point all execute overloads funnel through, so backend and OSS paths are both covered once. Driver entry points are resolved through the runtime (cudaGetDriverEntryPointByVersion), so the front end still never links libcuda -- the approach cu_tensor_map_encode_tiled already uses and documents. Same rung order as the Python side: a bound context is left alone, a real stream names its own context, and the default-stream handles name none, so the runtime's device decides there. Review: the dynamic-loading lookup is non-throwing (get_cuda_symbol throws when the library or symbol is missing, and this runs in a static initializer), guarded the way the rest of the headers guard exceptions; the backend test skips instead of failing where no engine serves the fused graph. Comments trimmed throughout -- the rationale and the measurements live in the PR. * build: declare the cuTile runtime as a [cutile] extra The cuTile linear-attention engines import cuda.tile, which was not declared anywhere -- not an extra, not requirements.txt -- so whether they run at all depended on the environment happening to have it. That is also why test_execute_from_a_thread_with_no_cuda_context silently covers only the FROST half on most machines: the cuTile engines decline in check_support when the import fails, and the test skips. Base cuda-tile only. Its [tileiras] extra pins cuda-toolkit>=13.2,<13.4, and that upper bound would cap the whole environment's toolkit and shut out CUDA 12 entirely -- the same reason nvidia-cutlass-dsl is not pinned to the FROST floor here. Without it cuda.tile falls back to a system tileiras, consistent with this package already leaving GPU wheels to the user. Resolution checked: `.[cutedsl,cutile]` resolves in one pass and adds exactly one package (cuda-tile 1.5.0) with nothing downgraded -- base cuda-tile requires only typing-extensions. The python_version marker keeps the extra resolvable on the declared 3.9 floor, which cuda-tile itself does not support. * perf(execute): probe the context before querying the stream The guard ran cudnnGetStream on every execute, though only the cold path needs a stream. Probe with one cuCtxGetCurrent instead and fetch the stream only when a context actually has to be established. Backend graph.execute() host time on parley, rebuilding the module for each: develop, no guard 10.805 us this PR, probe first 10.498 us / 10.683 us (two builds) The PR measured faster than develop both times, so the difference between builds is noise -- run-to-run spread alone is ~0.35 us across the 15 samples, and the guard's one cuCtxGetCurrent is ~106 ns. The unconditional version measured 10.788 us, i.e. also within noise: cudnnGetStream from C++ is nowhere near the ~1.5 us the Python path costs through pybind. Probing first is still the right shape, but it was not buying back a visible regression.
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.
No description provided.