Add concurrent ganged-wave graph replay - #37
Merged
Conversation
…ions, and structure benchmarking/diagnostic infrastructure.
…frastructure while hardening kernel launch scheduling
…t-kernel-cache (gfx1030, relocated) Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…ssions Concurrent multi-session graph capture on gfx1030/RDNA2 could corrupt output (NaN) because ROCm HIP stream-capture bookkeeping is process-global even with hipStreamCaptureModeThreadLocal. Add a process-wide reader/writer gate in Rust (rs_gfxgraph_core::capture_gate, exposed via PyO3 as CaptureLock/ReplayLock): capture takes the exclusive write lock; replay takes the shared read lock. Both LLMs still capture (serialized) and replay concurrently — high-level capture is preserved. Lock acquisition releases the GIL; degrades to a no-op on pure-Python installs. Wired into every capture site (capture_begin/end, standard capture context, shape-bucket lazy capture, conditional branches) and replay sites. Bump version 1.0.1 -> 1.1.0. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…r-hipGraph garble Wires the previously-vestigial hgb_shape_pool_t::static_bufs idea (allocated-but-never-written) the right way. The "aiter garbles under gfxGRAPH capture" report is decode-attn-under-graph replay-data-unsafety, NOT JIT-during-capture: a captured decode graph that bakes the per-step seq-len / block-table by value attends the capture-time length on every replay -> garble as the sequence grows. - hgb_decode_pool_create/replay/destroy + hgb_decode_capture_fn (src/decode_pool.hip, header decls): owns PERSISTENT device metadata buffers (seq_lens, block_tables) at fixed addresses; one graph per bucket reads them; replay = hipMemcpyAsync(live -> persistent) + hipGraphLaunch. No hipGraphExec*SetParams — only the buffer CONTENTS change, so the captured node reads fresh data. - tests/test_decode.hip: capture once, replay the SAME bucket graph at different seq-lens via in-place refresh, assert correct at both + cross-bucket + a negative control that the lengths genuinely differ. Validated on gfx1030 (RX 6700 XT): ctest `decode` PASSED. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…iter works on gfx1030 The aiter_on_rdna diagnostic asserted "AITER attention routes to (slower) Triton on RDNA; sglang auto-routes aiter->triton on gfx10xx" — FALSE for our patched build (the active sglang fork has no arch gate and prefers aiter). Rewrote summary/explanation/detail/fix: only the CK/ASM-only ops (flydsl moe_common, CK rmsnorm) are CDNA-only; native AITER flash-attention + JIT is built and working on gfx1030 (enable via SGLANG_USE_AITER=1). Added the real capture-safety guidance: the decode-under-graph garble was stale baked-by-value seq metadata, fixed by the device-resident-metadata path (hgb_decode_pool_*), not an aiter defect. Repointed the 2 bridge docs to the renamed gfx1030 frozen-cache dirs. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…riton from the safe-launch example Native AITER attention works on the patched gfx1030 (SGLANG_USE_AITER=1). The safe-launch example + SKILL no longer hard-code --attention-backend triton or the false comment; backend left unforced so sglang prefers aiter, with a note to keep cuda-graph on triton until the aiter+hgb_decode_pool capture lane lands. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…eSafeDecode) Opaque-handle C-API hgb_decode_pool_handle_create/replay/destroy (runtime_handles.cpp + header, mirroring hgb_composed_handle_*) + Rust bindings in rs_gfxgraph_native: the HgbDecodeCaptureFn callback type, NativeApi fields + dlsym loads, and NativeBridge::capture_safe_decode -> NativeCaptureSafeDecode<'a> with replay() + RAII Drop. The .so re-exports the 3 symbols; test_decode still passes; the native crate compiles clean. Completes A2 — the capture-safe decode pool is now usable from Rust consumers. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…ve import-time init init_environment_and_affinity already programmatically sets HSA_OVERRIDE_GFX_VERSION=10.3.0 when it detects the RDNA2 PCI device; now it also sets HSA_ENABLE_SDMA=1 (CPU-free SDMA copy engines) and HIP_HOST_COHERENT=0 (non-coherent host memory so the Infinity Cache can cache system RAM) — both only if the user hasn't set them. Propagates to every gfxGRAPH consumer at import. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
|
You have reached your Codex usage limits for code reviews. You can see your limits in the Codex usage dashboard. |
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.
Runtime changes
Validation