Integrate the Rust HIP graph runtime - #36
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>
|
Closing without merge: gfxGRAPH maintains separate release branches, and rust-hip-cpp is an independent release line rather than a feature branch targeting main. The branch remains committed and pushed. |
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 42ff08dc7e
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
| hipError_t err = hipMemcpyAsync(pool->d_seq_lens, h_seq_lens, (size_t)num_seqs * sizeof(int), | ||
| hipMemcpyHostToDevice, stream); | ||
| if (err != hipSuccess) return err; |
There was a problem hiding this comment.
When hgb_decode_pool_replay is called again on a different stream before the prior decode graph has finished, this async copy overwrites the single shared d_seq_lens/d_block_tables buffers while the in-flight graph may still be reading them. The function records per-bucket events after launch but never waits on any previous event before refreshing these shared buffers, so concurrent or cross-stream decode replays can read mixed metadata and produce corrupted results.
Useful? React with 👍 / 👎.
| return Err(PyValueError::new_err(format!( | ||
| "Input size {} exceeds largest bucket {}. Add a larger bucket.", | ||
| input_size, | ||
| self.buckets.last().unwrap_or(&0) | ||
| ))); |
There was a problem hiding this comment.
When the Rust router is installed, an input larger than the configured buckets now raises ValueError here instead of returning (-1, 2). ShapeBucketPool.__call__ still expects state 2 to trigger its documented eager fallback path, and the pure-Python router still returns that sentinel, so native installs regress oversized dynamic-shape inputs from fallback execution to an uncaught exception.
Useful? React with 👍 / 👎.
| hipError_t err = hgb_pipeline_update_kernel(pipe, node, params); | ||
| if (err != hipSuccess) return err; | ||
| return hgb_pipeline_launch(pipe); |
There was a problem hiding this comment.
For multi-threaded users of the public pipeline API, this releases the pipeline mutex after updating the active exec and reacquires it for launch, so another thread can launch in between and consume the just-updated buffer; this call then launches the other buffer with stale params. Since the helper is named/used as an atomic update+launch operation, it should hold the same lock across both steps or perform the combined logic directly.
Useful? React with 👍 / 👎.
| free(pool->graphs); | ||
| free(pool->execs); | ||
| free(pool->events); | ||
| memset(pool, 0, sizeof(*pool)); |
There was a problem hiding this comment.
Every decode-pool destroy path reaches this memset while pool->lock is still held, which zeros the pthread_mutex_t before the following pthread_mutex_unlock/pthread_mutex_destroy. Destroying a successfully created decode pool (or cleanup after a partial create failure) therefore unlocks a different, reinitialized mutex object, which is undefined behavior and can leave the original mutex locked or corrupt the pthread state.
Useful? React with 👍 / 👎.
1 participant
Integrates the Rust/HIP/C++ graph runtime, capture-safe decode pool, native handles and profiling, shape-aware routing, and the refreshed dependency lockfile.
Validation: all six workspace packages pass individual
cargo check.