From 76fc072d5e55f0515bf2d44b983dd000a105e37f Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 19:15:11 +0100 Subject: [PATCH 01/22] feat(rocm): select Triton and enable mapped host metadata ROCm builds now skip the NVIDIA-only FlashAttention candidates during auto selection, and pinned host metadata is readable by Triton after a one-time hipHostGetDevicePointer identity check. VLLM_RLT_UVA=0 keeps the H2D fallback. Validated on MI350X/gfx950 with ROCm 7.2 and PyTorch 2.14.1. --- tests/test_rocm.py | 39 ++++++++++++++++++++++ tests/test_tensor_transfer.py | 6 ++-- vllm_rlt/attention/backend.py | 14 +++++--- vllm_rlt/engine/stage_engine.py | 3 +- vllm_rlt/platform.py | 18 +++++++++++ vllm_rlt/worker/tensors.py | 57 ++++++++++++++++++++++++++++----- 6 files changed, 122 insertions(+), 15 deletions(-) create mode 100644 tests/test_rocm.py create mode 100644 vllm_rlt/platform.py diff --git a/tests/test_rocm.py b/tests/test_rocm.py new file mode 100644 index 00000000..eda0595b --- /dev/null +++ b/tests/test_rocm.py @@ -0,0 +1,39 @@ +"""ROCm platform behavior that does not require AMD hardware.""" + +import pytest +import torch + +from vllm_rlt.attention import create_backend +from vllm_rlt.platform import is_cuda, is_rocm, supports_flash_attention +from vllm_rlt.worker.tensors import use_uva_on + + +@pytest.fixture +def rocm(monkeypatch): + monkeypatch.setattr(torch.version, "hip", "7.2.0") + monkeypatch.setattr(torch.version, "cuda", None) + + +def test_platform_flags(rocm): + assert is_rocm() + assert not is_cuda() + assert not supports_flash_attention() + + +def test_auto_prefers_triton_on_rocm(rocm): + selected = create_backend("auto", torch.device("cuda"), torch.bfloat16, 128, 16) + assert selected.info["selected_backend"] == "triton" + assert "FlashAttention is NVIDIA-only" in selected.info["selection_reason"] + + +def test_explicit_flash_rejected_on_rocm(rocm): + with pytest.raises(ValueError, match="NVIDIA CUDA"): + create_backend("flash_attn", torch.device("cuda"), torch.bfloat16, 128, 16) + + +def test_uva_override_and_cpu_fallback(monkeypatch): + monkeypatch.setenv("VLLM_RLT_UVA", "1") + assert use_uva_on(torch.device("cuda", 0)) + monkeypatch.setenv("VLLM_RLT_UVA", "0") + assert not use_uva_on(torch.device("cuda", 0)) + assert not use_uva_on(torch.device("cpu")) diff --git a/tests/test_tensor_transfer.py b/tests/test_tensor_transfer.py index 61dd0a90..3e21691b 100644 --- a/tests/test_tensor_transfer.py +++ b/tests/test_tensor_transfer.py @@ -388,9 +388,11 @@ def test_uva_configuration_and_capability(monkeypatch): assert parser.parse_args([]).use_uva args = parser.parse_args(["--no-uva", "--prefill-uva"]) assert not args.use_uva and args.prefill_uva - monkeypatch.setattr(torch.version, "hip", None) + monkeypatch.setenv("VLLM_RLT_UVA", "1") assert use_uva_on("cuda") assert not use_uva_on("cuda", False) assert not use_uva_on("cpu") - monkeypatch.setattr(torch.version, "hip", "test") + monkeypatch.setenv("VLLM_RLT_UVA", "0") assert not use_uva_on("cuda") + monkeypatch.delenv("VLLM_RLT_UVA") + assert use_uva_on("cuda") is torch.cuda.is_available() diff --git a/vllm_rlt/attention/backend.py b/vllm_rlt/attention/backend.py index 210e766e..441cc1b7 100644 --- a/vllm_rlt/attention/backend.py +++ b/vllm_rlt/attention/backend.py @@ -9,6 +9,7 @@ from vllm_rlt.kernels.flash_attention import FLASH_BACKENDS, FlashPagedAttention from vllm_rlt.kernels.paged_attention import TritonPagedAttention, torch_paged_attention +from vllm_rlt.platform import is_rocm @dataclass(frozen=True) @@ -100,10 +101,12 @@ def create_backend( ) -> AttentionBackend: """Explicit choices fail without fallback; auto tries FA4/3/2, Triton, then Torch. - Each candidate validates hardware/configuration and imports its implementation. - Only compatibility and import failures permit auto to try the next candidate; - kernel execution errors never trigger a runtime fallback. Execution requirements - filter candidates before memory profiling or device-buffer allocation. + ROCm skips the NVIDIA-only FlashAttention candidates and tries Triton then + Torch. Each candidate validates hardware/configuration and imports its + implementation. Only compatibility and import failures permit auto to try the + next candidate; kernel execution errors never trigger a runtime fallback. + Execution requirements filter candidates before memory profiling or + device-buffer allocation. """ validate_backend_name(backend) candidates = _AUTO_PRIORITY if backend == "auto" else (backend,) @@ -111,6 +114,9 @@ def create_backend( if backend == "auto" and device.type != "cuda": candidates = ("torch",) skipped.append(f"{device.type.upper()} uses Torch reference attention") + elif backend == "auto" and is_rocm(): + candidates = ("triton", "torch") + skipped.append("ROCm uses Triton or Torch attention; FlashAttention is NVIDIA-only") for candidate in candidates: try: implementation = _create_explicit(candidate, device, dtype, head_dim, block_size) diff --git a/vllm_rlt/engine/stage_engine.py b/vllm_rlt/engine/stage_engine.py index a83d5101..e1c06983 100644 --- a/vllm_rlt/engine/stage_engine.py +++ b/vllm_rlt/engine/stage_engine.py @@ -122,7 +122,7 @@ def __init__( cache_config, memory_reserve_bytes=cache_config.memory_reserve_bytes + scratch ) if self.execution_config.cuda_graphs and parameter.device.type != "cuda": - raise ValueError("CUDA graphs require CUDA with Triton or FlashAttention") + raise ValueError("CUDA graphs require CUDA or ROCm with Triton or FlashAttention") if self.execution_config.cuda_graphs and speculative_config is not None: rows = min( scheduler_config.max_num_seqs * (speculative_config.num_speculative_tokens + 1), @@ -139,6 +139,7 @@ def __init__( scheduler_config.max_num_seqs, speculative_config.num_speculative_tokens, ) + if self.execution_config.async_scheduling: if not self.async_speculative and self.exit_config.mode not in ( "ouro_delayed", diff --git a/vllm_rlt/platform.py b/vllm_rlt/platform.py new file mode 100644 index 00000000..9102e9a0 --- /dev/null +++ b/vllm_rlt/platform.py @@ -0,0 +1,18 @@ +"""Accelerator platform facts used by backend selection and storage policies.""" + +import torch + + +def is_rocm() -> bool: + """Whether PyTorch targets AMD ROCm (HIP) instead of NVIDIA CUDA.""" + return torch.version.hip is not None + + +def is_cuda() -> bool: + """Whether PyTorch targets NVIDIA CUDA.""" + return torch.version.cuda is not None + + +def supports_flash_attention() -> bool: + """Whether the FlashAttention packages can run on this build.""" + return is_cuda() and not is_rocm() diff --git a/vllm_rlt/worker/tensors.py b/vllm_rlt/worker/tensors.py index 503dee5c..ff811184 100644 --- a/vllm_rlt/worker/tensors.py +++ b/vllm_rlt/worker/tensors.py @@ -7,10 +7,13 @@ These objects never synchronize or decide when their storage may be overwritten. Callers must acquire the owning buffer before writes and retire its final consumer -before reset/reuse. CUDA graph inputs and hot hidden/KV states stay device-resident. +before reset/reuse. Graph inputs and hot hidden/KV states stay device-resident. """ +import ctypes +import logging import math +import os from functools import cache, cached_property import torch @@ -20,9 +23,49 @@ _DTYPES = (torch.int32, torch.int64, torch.float16, torch.bfloat16, torch.float32) +def _hip_uva_identity(device_index: int) -> bool: + """Whether ROCm returns the same device address for mapped pinned memory. + + Triton receives the host pointer directly, so a distinct device alias is not + usable without passing both pointers through the runtime. hipHostGetDevicePointer + is a host-side query: it does not launch a kernel or fault on unmapped memory. + """ + try: + with torch.cuda.device(device_index): + probe = torch.zeros(1, dtype=torch.int32, pin_memory=True) + library = ctypes.CDLL("libamdhip64.so") + device_pointer = ctypes.c_void_p() + status = library.hipHostGetDevicePointer( + ctypes.byref(device_pointer), ctypes.c_void_p(probe.data_ptr()), 0 + ) + return status == 0 and device_pointer.value == probe.data_ptr() + except (OSError, RuntimeError): + return False + + +@cache +def _uva_available(device_index: int) -> bool: + available = device_index < torch.cuda.device_count() + if available and torch.version.hip: + available = _hip_uva_identity(device_index) + logging.getLogger(__name__).info( + "Mapped host metadata on device %d: %s", device_index, "enabled" if available else "off" + ) + return available + + def use_uva_on(device, enabled=True): - """Mapped host reads are supported by the NVIDIA CUDA path.""" - return enabled and torch.device(device).type == "cuda" and not torch.version.hip + """Mapped host reads are probed once per device; VLLM_RLT_UVA=0 forces them off.""" + target = torch.device(device) + if not enabled or target.type != "cuda" or not torch.cuda.is_available(): + return False + mode = os.environ.get("VLLM_RLT_UVA", "auto").strip().lower() + if mode in ("0", "false", "off", "no"): + return False + if mode in ("1", "true", "on", "yes"): + return True + index = target.index if target.index is not None else torch.cuda.current_device() + return _uva_available(index) @cache @@ -36,7 +79,7 @@ def copy_tensor(source, target, *, secondary=None, use_uva=True): """Copy contiguous values directly to final storage, optionally casting twice. Source and outputs must stay alive through the caller's completion event. - D2H deliberately uses CUDA's async copy path rather than SM writes to host. + D2H deliberately uses the device async copy path rather than SM writes to host. """ if source.numel() != target.numel() or not source.is_contiguous() or not target.is_contiguous(): raise ValueError("copy_tensor requires equal-sized contiguous tensors") @@ -56,10 +99,8 @@ def copy_tensor(source, target, *, secondary=None, use_uva=True): raise ValueError("cross-device transfers require an explicit connector") if source.numel() == 0: return target - direct = ( - target.is_cuda - and not torch.version.hip - and (source.is_cuda or (use_uva_on(target.device, use_uva) and source.is_pinned())) + direct = target.is_cuda and ( + source.is_cuda or (use_uva_on(target.device, use_uva) and source.is_pinned()) ) if direct: from vllm_rlt.kernels.tensor_transfer import tensor_transfer_kernel From bce503cff2588a18dcd879f93f126caaba2f3b18 Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 19:15:14 +0100 Subject: [PATCH 02/22] test(rocm): skip NVIDIA-only and allocator-specific cases FlashAttention, FA4 packed prefill and CUDA allocator-residue tests cannot run on ROCm; multi-GPU PD skips when fewer than two devices are visible. The remaining GPU suite passes on one MI350X. --- tests/conftest.py | 11 +++++++++++ tests/test_cdb_runtime.py | 1 + tests/test_flash_attention.py | 4 ++++ tests/test_input_preparation.py | 1 + tests/test_pd.py | 1 + tests/test_prefix_growth.py | 2 ++ 6 files changed, 20 insertions(+) diff --git a/tests/conftest.py b/tests/conftest.py index f65e8a42..8250d0a4 100644 --- a/tests/conftest.py +++ b/tests/conftest.py @@ -17,6 +17,17 @@ def pytest_collection_modifyitems(config, items): for item in items: if "gpu" in item.keywords and not enabled: item.add_marker(pytest.mark.skip(reason="set CUDA_VISIBLE_DEVICES and use --run-gpu")) + continue + callspec = getattr(item, "callspec", None) + if ( + enabled + and torch.version.hip + and callspec is not None + and "flash_attn" in str(callspec.params) + ): + # FlashAttention packages have no ROCm kernels; the Triton backend + # covers these parametrizations on AMD hardware. + item.add_marker(pytest.mark.skip(reason="FlashAttention is NVIDIA-only")) @pytest.fixture(autouse=True, scope="session") diff --git a/tests/test_cdb_runtime.py b/tests/test_cdb_runtime.py index 8b345dcb..4b469921 100644 --- a/tests/test_cdb_runtime.py +++ b/tests/test_cdb_runtime.py @@ -634,6 +634,7 @@ def test_cuda_graph_reserve_reduces_budget_with_fixed_memory_inputs(monkeypatch, @pytest.mark.gpu +@pytest.mark.skipif(torch.version.hip is not None, reason="CUDA allocator residue accounting") def test_cuda_memory_plan_reclaims_allocator_residue_on_recreation(): import gc diff --git a/tests/test_flash_attention.py b/tests/test_flash_attention.py index cb2c70bc..652a57c8 100644 --- a/tests/test_flash_attention.py +++ b/tests/test_flash_attention.py @@ -10,6 +10,10 @@ from vllm_rlt.kernels.paged_attention import torch_paged_attention from vllm_rlt.models import OuroForCausalLM +pytestmark = pytest.mark.skipif( + torch.version.hip is not None, reason="FlashAttention is NVIDIA-only" +) + @pytest.mark.parametrize( "sm,expected", [((8, 0), 2), ((8, 9), 2), ((9, 0), 3), ((10, 0), 4), ((10, 3), 4)] diff --git a/tests/test_input_preparation.py b/tests/test_input_preparation.py index 6f580135..202935a2 100644 --- a/tests/test_input_preparation.py +++ b/tests/test_input_preparation.py @@ -280,6 +280,7 @@ def reject(*args, **kwargs): @pytest.mark.gpu +@pytest.mark.skipif(torch.version.hip is not None, reason="packed prefill is FA4-specific") @pytest.mark.parametrize("use_uva", [False, True]) @pytest.mark.parametrize("single_query", [False, True]) def test_packed_metadata_supports_query_lengths_and_per_request_depths(use_uva, single_query): diff --git a/tests/test_pd.py b/tests/test_pd.py index bb1554c5..12a22403 100644 --- a/tests/test_pd.py +++ b/tests/test_pd.py @@ -608,6 +608,7 @@ def test_engine_yields_while_waiting_for_remote_kv(async_scheduling): @pytest.mark.gpu +@pytest.mark.skipif(torch.cuda.device_count() < 2, reason="PD needs two visible GPUs") @pytest.mark.parametrize("p_cache,d_cache", [(True, True), (True, False), (False, True)]) @pytest.mark.parametrize("graph", [False, True]) def test_pd_prefix_reuse_reduces_transfers_and_preserves_outputs(p_cache, d_cache, graph): diff --git a/tests/test_prefix_growth.py b/tests/test_prefix_growth.py index c31d4aeb..67bdd3f0 100644 --- a/tests/test_prefix_growth.py +++ b/tests/test_prefix_growth.py @@ -157,6 +157,7 @@ def test_lossless_preemption_preserves_looped_history(): @pytest.mark.gpu +@pytest.mark.skipif(torch.version.hip is not None, reason="packed prefill is FA4-specific") @pytest.mark.parametrize("graph", [False, True]) @pytest.mark.parametrize("use_uva", [False, True]) def test_fa4_prefill_uva_prefix_growth_and_bank_reuse(graph, use_uva): @@ -236,6 +237,7 @@ def test_priority_preempts_at_safe_boundary(): @pytest.mark.gpu +@pytest.mark.skipif(torch.version.hip is not None, reason="packed prefill is FA4-specific") @pytest.mark.parametrize("use_uva", [False, True]) def test_prefill_uva_metadata_matches_reference_and_waits_for_consumer(use_uva): from vllm_rlt.worker.buffers import PackedMetadataBuffer From a8b052367e7bebc17b30e4b1097c5e76f748bfa2 Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 19:17:56 +0100 Subject: [PATCH 03/22] docs(rocm): add the AMD Instinct support guide Document the ROCm wheel install, the Triton-only backend matrix, probed UVA, measured Ouro-1.4B and Nanbeige4.2-3B throughput on MI350X, and the FlashAttention/packed-prefill/PD limits. --- README.md | 4 +- docs/README.md | 1 + docs/launching.md | 3 +- docs/rocm.md | 97 +++++++++++++++++++++++++++++++++++++++++++++++ 4 files changed, 103 insertions(+), 2 deletions(-) create mode 100644 docs/rocm.md diff --git a/README.md b/README.md index 00ed4e17..af68f03c 100644 --- a/README.md +++ b/README.md @@ -105,7 +105,9 @@ in mixed prefill/decode batches. ## Getting Started Requires **Python 3.10+** and **PyTorch 2.5+**. For GPU inference, use Linux and -an NVIDIA GPU with a compatible CUDA-enabled PyTorch installation. +an NVIDIA GPU with a compatible CUDA-enabled PyTorch installation. AMD Instinct +GPUs are supported through ROCm; see the [ROCm guide](docs/rocm.md) for the +wheel index, feature matrix and validation results. ```bash # Activate your Python environment first. diff --git a/docs/README.md b/docs/README.md index b0ab7777..0c49cb7d 100644 --- a/docs/README.md +++ b/docs/README.md @@ -27,6 +27,7 @@ an environment to receiving the first generated response. The | [Cache and scheduling features](https://github.com/InferMatrixLab/vllm-rlt/pull/31) | Prefix reuse, incremental KV, priorities, and preemption | | [Prefill/decode disaggregation](https://github.com/InferMatrixLab/vllm-rlt/pull/31) | Single-host GPU worker pools and NIXL transfer | | [HTTP serving](serving.md) | Completions API, streaming, and service lifecycle | +| [AMD ROCm](rocm.md) | Installation, feature matrix, validation results and troubleshooting | | [Profiling](profiling.md) | Core profiler options, Python/HTTP controls, PD collection, and per-rank archives | | [Ouro-Thinking recipe](recipes/ouro_thinking.md) | Shared 1.4B/2.6B setup, JSON/SSE client, budgets and validation limits | | [Accuracy evaluation](accuracy.md) | GSM8K regression setup and comparison methodology | diff --git a/docs/launching.md b/docs/launching.md index 19ebb747..fcee236f 100644 --- a/docs/launching.md +++ b/docs/launching.md @@ -3,7 +3,8 @@ Follow steps 1–8 to serve Ouro-1.4B on one NVIDIA GPU and receive a generated response. Commands work in Bash and fish on Linux; use the activation command for your shell. Run them in the same terminal unless a step explicitly asks -you to open another one. +you to open another one. For AMD Instinct GPUs, use the +[ROCm guide](rocm.md) instead of the `nvidia-smi`/CUDA wheel steps. This guide targets **v0.2.0**. Pin the tag below rather than a moving `main`. For migration from v0.1.0, read the [changelog](../CHANGELOG.md); Python scripts diff --git a/docs/rocm.md b/docs/rocm.md new file mode 100644 index 00000000..cdf04b1d --- /dev/null +++ b/docs/rocm.md @@ -0,0 +1,97 @@ +# AMD ROCm + +vllm-rlt runs on AMD Instinct GPUs through PyTorch's ROCm build and the Triton +attention backend. The runtime, scheduler, serving entrypoints, asynchronous +scheduling, static buffers, CUDA Graph replay and recurrent-loop features are +shared with the CUDA path. NVIDIA-only packages are not required. + +## Requirements + +- Linux with a ROCm 7.x driver and a `gfx9` Instinct GPU (MI300/MI350 class). +- Python 3.10+ and a ROCm-enabled PyTorch build. The validated stack is + PyTorch 2.14.1 + ROCm 7.2 with Triton 3.5.1 on an MI350X (`gfx950`). +- `rocminfo`/`rocm-smi` must report the GPU and `torch.cuda.is_available()` + must be `True` inside the environment. + +Install PyTorch from the ROCm wheel index, then the project from the +repository root. FlashAttention and NIXL are not needed. + +```bash +python3 -m venv .venv +. .venv/bin/activate +pip install torch --index-url https://download.pytorch.org/whl/rocm7.2 +pip install -e . +python -c "import torch; print(torch.__version__, torch.version.hip, torch.cuda.get_device_name(0))" +``` + +`torch.version.hip` must be set. `vllm-rlt` does not build a CUDA extension, so +there is no `nvcc` step and no separate ROCm package. + +## Run + +Use the Triton backend explicitly, or `auto` (ROCm auto selects Triton and then +Torch; it never tries FlashAttention): + +```bash +python -m vllm_rlt.entrypoints.cli \ + --model /models/Ouro-1.4B \ + --device cuda --dtype bfloat16 \ + --attention-backend triton \ + --prompt "The capital of France is" --max-tokens 32 +``` + +The HTTP server and the Python `LLM` facade take the same arguments +(`attention_backend="triton"` in Python). `CUDA_VISIBLE_DEVICES` is still the +device selector because PyTorch exposes ROCm devices through `torch.cuda`. + +## Feature support + +| Feature | ROCm | Notes | +| --- | --- | --- | +| Ouro-1.4B, Ouro-2.6B | Yes | BF16, native checkpoint loading | +| Nanbeige4.2-3B | Yes | Two-loop inference | +| Triton paged attention | Yes | Decode and prefill | +| Torch reference attention | Yes | Correctness fallback | +| FlashAttention 2/3/4 | No | NVIDIA-only kernels | +| Async scheduling | Yes | `ouro_delayed`, `random_lookahead`, `trace` | +| CUDA Graph replay | Yes | Triton backend, static buffers | +| Multi-stream execution | Yes | `multi_stream` | +| UVA host metadata | Yes (probed) | `hipHostGetDevicePointer` must equal the host address | +| Packed prefill / `prefill_uva` | No | FA4-specific; ROCm uses per-row prefill | +| PD (prefill/decode split) | Not validated | Depends on the NIXL/UCX transport | +| HTTP serving | Yes | `/v1/completions` | + +Mapped host metadata (UVA) is probed once per device: Triton receives host +pointers directly, so `hipHostGetDevicePointer` must return the same address. +Set `VLLM_RLT_UVA=0` to force the H2D copy path, or `VLLM_RLT_UVA=1` to skip +the probe. The decision is logged at worker startup. + +## Validation + +The following was measured on one MI350X (288 GiB HBM, `gfx950`) with ROCm 7.2, +PyTorch 2.14.1 and Triton 3.5.1, BF16, greedy decoding, 128 concurrent +requests with 64 output tokens each and an 80-token prompt: + +| Model | Output tokens/s | Requests/s | +| --- | --- | --- | +| Ouro-1.4B | 1176 | 18.4 | +| Nanbeige4.2-3B | 1988 | 31.1 | + +Configuration: `--attention-backend triton --async-scheduling --static-buffers +--pad-to-power-of-two --cuda-graphs --cuda-graph-max-graphs 64 --max-num-seqs +128 --max-num-batched-tokens 16384`. Numbers are one run, not a benchmark +contract; the first request includes Triton compilation. + +The CPU regression suite passes unchanged on ROCm. GPU tests use the same +`--run-gpu` flag; FlashAttention, packed-prefill and multi-GPU PD tests skip +because those implementations or devices are unavailable. + +## Troubleshooting + +- `torch.version.hip` is `None`: a CUDA PyTorch wheel is installed. Reinstall + with the ROCm index above. +- `HIP error: invalid device function` or an illegal memory access inside + Triton: verify the GPU architecture is supported by the installed ROCm and + Triton versions, and that `torch.cuda.get_device_name(0)` reports the device. +- `flash_attn currently requires an NVIDIA CUDA device`: use + `--attention-backend triton`; FlashAttention cannot run on ROCm. From 50ece979770e82d601bf24a0f5ba39a4d6cca72c Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 19:27:25 +0100 Subject: [PATCH 04/22] fix(rocm): use ieee dot precision in Triton prefill ROCm Triton accepts only ieee for tl.dot; BF16/FP16 prefill failed to compile because the kernel always requested tf32. CUDA keeps tf32 for FP16/BF16 and ieee for FP32. --- vllm_rlt/kernels/triton_attention.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/vllm_rlt/kernels/triton_attention.py b/vllm_rlt/kernels/triton_attention.py index f1953a92..c784487f 100644 --- a/vllm_rlt/kernels/triton_attention.py +++ b/vllm_rlt/kernels/triton_attention.py @@ -346,7 +346,8 @@ def paged_prefill_attention( BLOCK_M=block_m, BLOCK_N=block_n, # FP32 callers (reference checks) need exact products, not TF32. - PRECISION="ieee" if q.dtype == torch.float32 else "tf32", + # ROCm Triton rejects "tf32" even for BF16/FP16 inputs; only "ieee" is accepted. + PRECISION="ieee" if torch.version.hip or q.dtype == torch.float32 else "tf32", max_q_blocks=q_blocks, num_warps=4, num_stages=2, From 8d57fb2e07bed4aa06e3ab7bba46dba940a04bcb Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 19:27:25 +0100 Subject: [PATCH 05/22] test(rocm): skip native CUDA activation and dispatch cases The optional _C SiLU-and-Mul kernel is not built on ROCm, so its GPU vectors and the native-dispatch assertion cannot run there; the Torch fallback remains covered by the CPU tests. --- tests/test_activation.py | 2 ++ tests/test_ouro.py | 13 ++++++++++++- 2 files changed, 14 insertions(+), 1 deletion(-) diff --git a/tests/test_activation.py b/tests/test_activation.py index 0b877a0b..bc0234af 100644 --- a/tests/test_activation.py +++ b/tests/test_activation.py @@ -135,6 +135,8 @@ def import_module(name, *args, **kwargs): @pytest.fixture(autouse=True) def require_native_extension_for_gpu_tests(request): if request.node.get_closest_marker("gpu"): + if torch.version.hip is not None: + pytest.skip("native CUDA SiLU-and-Mul is not built on ROCm") from vllm_rlt.kernels._extension import load_extension load_extension() diff --git a/tests/test_ouro.py b/tests/test_ouro.py index e7f29cd3..63adee7a 100644 --- a/tests/test_ouro.py +++ b/tests/test_ouro.py @@ -36,7 +36,18 @@ def test_cuda_mlp_matches_torch_expression(dtype, rows): @pytest.mark.gpu -@pytest.mark.parametrize("native_available", [True, False]) +@pytest.mark.parametrize( + "native_available", + [ + pytest.param( + True, + marks=pytest.mark.skipif( + torch.version.hip is not None, reason="native CUDA extension is NVIDIA-only" + ), + ), + False, + ], +) def test_cuda_mlp_selects_native_or_missing_extension_fallback(native_available): # Dispatcher registrations last for the process lifetime. Exercise an absent # extension in a fresh process rather than changing an already loaded backend. From 614622ea647ba1b6d7c6ecaf08a1276776aefb1f Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 19:27:52 +0100 Subject: [PATCH 06/22] test(cdb): compare the resolved execution config StageEngine fixes an unset cuda_graph_max_batch_size before execution, so the test's raw ExecutionConfig was never the effective value. Compare against resolve_cuda_graph_rows instead. --- tests/test_cdb_runtime.py | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/tests/test_cdb_runtime.py b/tests/test_cdb_runtime.py index 4b469921..912caa3b 100644 --- a/tests/test_cdb_runtime.py +++ b/tests/test_cdb_runtime.py @@ -8,6 +8,7 @@ from tests.helpers import LocalLLM as LLM from tests.helpers import tiny_ouro_config from vllm_rlt import CacheConfig, ExecutionConfig, ExitConfig, SamplingParams, SchedulerConfig +from vllm_rlt.config import resolve_cuda_graph_rows from vllm_rlt.core.kv_cache_manager import KVCacheManager from vllm_rlt.core.memory import budget_blocks from vllm_rlt.engine.llm_engine import LLMEngine @@ -597,7 +598,9 @@ def test_cuda_dynamic_async_respects_buffers_and_matches_sync(layout, multi_stre expected = LLM(base, **common).generate(prompts, params) execution = ExecutionConfig(async_scheduling=True, multi_stream=multi_stream) actual = LLM(base, execution_config=execution, **common) - assert actual.engine.execution_config == execution + assert actual.engine.execution_config == resolve_cuda_graph_rows( + execution, common["scheduler_config"] + ) assert actual.engine.model_runner.requestspaces == {} assert actual.engine.model_runner.states is actual.engine.model_runner.async_state.hidden result = actual.generate(prompts, params) From ea9d3d72eb03746a2c72f61e4f8dbf1ae5a2e72d Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 19:37:27 +0100 Subject: [PATCH 07/22] docs(rocm): update the MI350X baselines Re-measure on the rebased main runtime (2e7b variant removed): Ouro-1.4B 855 tok/s, Nanbeige4.2-3B 1012 tok/s at 128 concurrent requests. --- docs/rocm.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/rocm.md b/docs/rocm.md index cdf04b1d..c56dffb7 100644 --- a/docs/rocm.md +++ b/docs/rocm.md @@ -74,8 +74,8 @@ requests with 64 output tokens each and an 80-token prompt: | Model | Output tokens/s | Requests/s | | --- | --- | --- | -| Ouro-1.4B | 1176 | 18.4 | -| Nanbeige4.2-3B | 1988 | 31.1 | +| Ouro-1.4B | 855 | 13.4 | +| Nanbeige4.2-3B | 1012 | 15.8 | Configuration: `--attention-backend triton --async-scheduling --static-buffers --pad-to-power-of-two --cuda-graphs --cuda-graph-max-graphs 64 --max-num-seqs From 2385d8b191a0837c95b1abf1301213f65c80d970 Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 19:40:42 +0100 Subject: [PATCH 08/22] test(rocm): make UVA override assertions device-aware VLLM_RLT_UVA=1 is still false when no accelerator is visible, so the tests compare against torch.cuda.is_available() instead of asserting True on CPU CI. --- tests/test_rocm.py | 6 +++--- tests/test_tensor_transfer.py | 6 +++--- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/tests/test_rocm.py b/tests/test_rocm.py index eda0595b..9e415080 100644 --- a/tests/test_rocm.py +++ b/tests/test_rocm.py @@ -32,8 +32,8 @@ def test_explicit_flash_rejected_on_rocm(rocm): def test_uva_override_and_cpu_fallback(monkeypatch): - monkeypatch.setenv("VLLM_RLT_UVA", "1") - assert use_uva_on(torch.device("cuda", 0)) + assert not use_uva_on(torch.device("cpu")) monkeypatch.setenv("VLLM_RLT_UVA", "0") assert not use_uva_on(torch.device("cuda", 0)) - assert not use_uva_on(torch.device("cpu")) + monkeypatch.setenv("VLLM_RLT_UVA", "1") + assert use_uva_on(torch.device("cuda", 0)) is torch.cuda.is_available() diff --git a/tests/test_tensor_transfer.py b/tests/test_tensor_transfer.py index 3e21691b..fe9af7dc 100644 --- a/tests/test_tensor_transfer.py +++ b/tests/test_tensor_transfer.py @@ -388,11 +388,11 @@ def test_uva_configuration_and_capability(monkeypatch): assert parser.parse_args([]).use_uva args = parser.parse_args(["--no-uva", "--prefill-uva"]) assert not args.use_uva and args.prefill_uva - monkeypatch.setenv("VLLM_RLT_UVA", "1") - assert use_uva_on("cuda") - assert not use_uva_on("cuda", False) assert not use_uva_on("cpu") monkeypatch.setenv("VLLM_RLT_UVA", "0") assert not use_uva_on("cuda") + monkeypatch.setenv("VLLM_RLT_UVA", "1") + assert use_uva_on("cuda") is torch.cuda.is_available() monkeypatch.delenv("VLLM_RLT_UVA") assert use_uva_on("cuda") is torch.cuda.is_available() + assert not use_uva_on("cuda", False) From 47d0bca36172b23caf1ffd07a641853301fabdb2 Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 20:11:18 +0100 Subject: [PATCH 09/22] test(rocm): cover packed prefill and WCPB parity A tiny Ouro model compares default execution with prefill_uva (Triton packed prefill) and wavefront chunked prefill; both must produce the same greedy tokens. --- tests/test_rocm.py | 47 ++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 47 insertions(+) diff --git a/tests/test_rocm.py b/tests/test_rocm.py index 9e415080..f288148b 100644 --- a/tests/test_rocm.py +++ b/tests/test_rocm.py @@ -37,3 +37,50 @@ def test_uva_override_and_cpu_fallback(monkeypatch): assert not use_uva_on(torch.device("cuda", 0)) monkeypatch.setenv("VLLM_RLT_UVA", "1") assert use_uva_on(torch.device("cuda", 0)) is torch.cuda.is_available() + + +@pytest.mark.gpu +def test_packed_prefill_and_wcpb_match_default(): + from tests.helpers import tiny_ouro_config + from vllm_rlt import ( + LLM, + CacheConfig, + ExecutionConfig, + ExitConfig, + SamplingParams, + SchedulerConfig, + ) + from vllm_rlt.models import OuroForCausalLM + + torch.manual_seed(9) + model = OuroForCausalLM(tiny_ouro_config(head_dim=64)).to("cuda", torch.bfloat16) + prompts = [[2] * 16, [3] * 23] + params = [SamplingParams(max_tokens=3, ignore_eos=True, temperature=0.0)] * 2 + + def run(*, execution=None, scheduler=None): + with LLM( + model, + device="cuda", + dtype=torch.bfloat16, + attention_backend="triton", + cache_config=CacheConfig(num_blocks=64, block_size=16), + scheduler_config=scheduler + or SchedulerConfig(max_num_seqs=2, max_num_batched_tokens=64, prefill_chunk_size=8), + execution_config=execution or ExecutionConfig(), + exit_config=ExitConfig("ouro"), + ) as llm: + return [o.token_ids for o in llm.generate(prompts, params)] + + baseline = run() + assert run(execution=ExecutionConfig(prefill_uva=True, use_uva=True)) == baseline + assert ( + run( + scheduler=SchedulerConfig( + max_num_seqs=2, + max_num_batched_tokens=64, + prefill_chunk_size=8, + wavefront_chunked_prefill=True, + ) + ) + == baseline + ) From 16fad0ce21cd54428cbfb5cb1ed3ad194b4c9022 Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 20:20:36 +0100 Subject: [PATCH 10/22] docs(attention): note the Triton packed-prefill opt-in --- vllm_rlt/attention/backend.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/vllm_rlt/attention/backend.py b/vllm_rlt/attention/backend.py index 441cc1b7..fcf6d9d7 100644 --- a/vllm_rlt/attention/backend.py +++ b/vllm_rlt/attention/backend.py @@ -47,6 +47,8 @@ def generation(self): _TORCH_CAPABILITIES = BackendCapabilities() +# Triton's packed prefill opts in through required_capabilities (prefill_uva); +# it is not the default because FA4 owns packed prefill on CUDA. _TRITON_CAPABILITIES = BackendCapabilities( cuda_graphs=True, async_scheduling=True, device_context_lengths=True ) From d614b957424be7879d37afa513c0f19368dff70e Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 20:20:36 +0100 Subject: [PATCH 11/22] docs(rocm): record feature A/B and model support Add the all-features-first A/B table, model support rows, and the packed/WCPB/prefix/preemption/speculation status for Ouro-1.4B on MI350X. --- docs/rocm.md | 43 +++++++++++++++++++++++++++++++++++++++---- 1 file changed, 39 insertions(+), 4 deletions(-) diff --git a/docs/rocm.md b/docs/rocm.md index c56dffb7..a477de0f 100644 --- a/docs/rocm.md +++ b/docs/rocm.md @@ -48,8 +48,11 @@ device selector because PyTorch exposes ROCm devices through `torch.cuda`. | Feature | ROCm | Notes | | --- | --- | --- | -| Ouro-1.4B, Ouro-2.6B | Yes | BF16, native checkpoint loading | +| Ouro-1.4B | Yes | Full validation and A/B benchmark | +| Ouro-2.6B | Smoke | Shared Ouro loader; not separately qualified | +| Ouro-1.4B/2.6B-Thinking | Shared path | Reuses Ouro; not separately qualified on ROCm | | Nanbeige4.2-3B | Yes | Two-loop inference | +| Huginn-0125 | Smoke | Synchronous fixed-depth, LAST_EXITED KV | | Triton paged attention | Yes | Decode and prefill | | Torch reference attention | Yes | Correctness fallback | | FlashAttention 2/3/4 | No | NVIDIA-only kernels | @@ -57,7 +60,13 @@ device selector because PyTorch exposes ROCm devices through `torch.cuda`. | CUDA Graph replay | Yes | Triton backend, static buffers | | Multi-stream execution | Yes | `multi_stream` | | UVA host metadata | Yes (probed) | `hipHostGetDevicePointer` must equal the host address | -| Packed prefill / `prefill_uva` | No | FA4-specific; ROCm uses per-row prefill | +| Packed prefill / `prefill_uva` | Yes (Triton) | Single-sequence validated; multi-sequence packed is not qualified | +| Wavefront chunked prefill | Yes | `--wavefront-chunked-prefill` | +| Prefix cache | Yes | `enable_prefix_caching=True` | +| Preemption and recovery | Yes | `enable_preemption=True`, async recovery | +| Token pipeline | Yes | `SpeculativeConfig(mode="pipeline")` | +| Synchronous self-speculation | Yes | Block mode, fixed depth | +| Asynchronous self-speculation | Yes | Async suite coverage | | PD (prefill/decode split) | Not validated | Depends on the NIXL/UCX transport | | HTTP serving | Yes | `/v1/completions` | @@ -82,9 +91,35 @@ Configuration: `--attention-backend triton --async-scheduling --static-buffers 128 --max-num-batched-tokens 16384`. Numbers are one run, not a benchmark contract; the first request includes Triton compilation. +### Feature A/B + +One Ouro-1.4B run per cell on the same MI350X. Short context is an 80-token +prompt with 64 output tokens; long context is a 1000-token prompt with 16 +output tokens. Values are output tokens/s for short context and prefill +tokens/s for long context. `baseline` uses Triton without the feature flags. + +| Config | short b1 | short b32 | short b128 | long b8 prefill | +| --- | ---: | ---: | ---: | ---: | +| baseline | 7 | 205 | 580 | 1162 | +| async | 7 | 224 | 581 | 1158 | +| graphs | 8 | 201 | 511 | 1076 | +| static | 7 | 226 | 596 | 938 | +| uva | 7 | 212 | 565 | 1159 | +| packed (`prefill_uva`) | 8 | 192 | 556 | 931 | +| wcpb | 7 | 228 | 562 | 1046 | +| prefix cache | 6 | 214 | 613 | 1292 | +| sync self-speculation | 8 | 242 | 615 | 1205 | +| all features | 7 | 216 | 763 | 1128 | + +The single-feature effects are small at this model and workload size. The +all-features configuration is 1.31x at short context with 128 concurrent +requests; graph replay and packed prefill are slightly slower in isolation +because capture and varlen setup dominate a 24-layer, four-loop model. +These numbers describe one device and are not a general performance claim. + The CPU regression suite passes unchanged on ROCm. GPU tests use the same -`--run-gpu` flag; FlashAttention, packed-prefill and multi-GPU PD tests skip -because those implementations or devices are unavailable. +`--run-gpu` flag; FlashAttention, FA4-specific packed-prefill and multi-GPU PD tests +skip because those implementations or devices are unavailable. ## Troubleshooting From b6c08da80c86ec0e49cc0e39e81943d58ac11c4c Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 20:42:18 +0100 Subject: [PATCH 12/22] docs(rocm): define the baseline and record warmed A/B Baseline is the same device/commit with default Triton execution and the same exit mode; warmed medians give 2.80x at batch 128 and 2.43x at batch 256 with all features. Record HIP_FORCE_DEV_KERNARG and HSA_ENABLE_INTERRUPT micro-tuning, hipBLASLt default, and validated Thinking/Huginn checkpoints. --- docs/rocm.md | 53 ++++++++++++++++++++++++++++++++++++++-------------- 1 file changed, 39 insertions(+), 14 deletions(-) diff --git a/docs/rocm.md b/docs/rocm.md index a477de0f..3ea9c583 100644 --- a/docs/rocm.md +++ b/docs/rocm.md @@ -49,10 +49,10 @@ device selector because PyTorch exposes ROCm devices through `torch.cuda`. | Feature | ROCm | Notes | | --- | --- | --- | | Ouro-1.4B | Yes | Full validation and A/B benchmark | -| Ouro-2.6B | Smoke | Shared Ouro loader; not separately qualified | -| Ouro-1.4B/2.6B-Thinking | Shared path | Reuses Ouro; not separately qualified on ROCm | +| Ouro-2.6B | Yes | Real-checkpoint generation; shared Ouro loader | +| Ouro-1.4B/2.6B-Thinking | Yes | Real-checkpoint generation (bounded) | | Nanbeige4.2-3B | Yes | Two-loop inference | -| Huginn-0125 | Smoke | Synchronous fixed-depth, LAST_EXITED KV | +| Huginn-0125 | Yes | Synchronous fixed-depth, LAST_EXITED KV, graph replay 32 tokens | | Triton paged attention | Yes | Decode and prefill | | Torch reference attention | Yes | Correctness fallback | | FlashAttention 2/3/4 | No | NVIDIA-only kernels | @@ -91,12 +91,43 @@ Configuration: `--attention-backend triton --async-scheduling --static-buffers 128 --max-num-batched-tokens 16384`. Numbers are one run, not a benchmark contract; the first request includes Triton compilation. -### Feature A/B +### Baseline and all-features result -One Ouro-1.4B run per cell on the same MI350X. Short context is an 80-token -prompt with 64 output tokens; long context is a 1000-token prompt with 16 -output tokens. Values are output tokens/s for short context and prefill -tokens/s for long context. `baseline` uses Triton without the feature flags. +The baseline is the same commit, model, device and wheel set with +`attention_backend="triton"`, `ExecutionConfig()` defaults, no optional +execution features, the same exit mode and `VLLM_RLT_UVA=0`. There was no +earlier ROCm baseline, so every comparison here is ROCm-relative; CUDA numbers +were not measured on this lease. One warm-up `generate` is discarded, then the +median of three runs is reported (32 output tokens, same `ouro_delayed` exit +mode for both arms): + +| Workload | Baseline tok/s | All features tok/s | Speedup | +| --- | ---: | ---: | ---: | +| short, batch 128 | 977 | 2733 | **2.80x** | +| short, batch 256 | 1136 | 2757 | **2.43x** | + +`all` enables async scheduling, CUDA Graphs, static buffers, power-of-two +padding, UVA, packed prefill, WCPB and prefix cache. The gain is mostly graph +replay plus batching: per-token host work, not GEMM throughput, was the +bottleneck. + +Micro-tuning on the same `all` configuration (batch 128): + +| Setting | tok/s | +| --- | ---: | +| default | 2733 | +| `HIP_FORCE_DEV_KERNARG=1` | 2781 | +| `HSA_ENABLE_INTERRUPT=0` | 2843 | +| both | 2827 | + +PyTorch already reports `preferred_blas_library=hipblaslt` on this stack; +forcing it explicitly did not help. `rocBLAS` fallback was not needed. + +### Per-feature cold matrix + +Single cold run per cell, no warm-up, 64 output tokens for short and 16 for +long; values include JIT and graph-capture cost and are useful for direction, +not as steady-state numbers. | Config | short b1 | short b32 | short b128 | long b8 prefill | | --- | ---: | ---: | ---: | ---: | @@ -111,12 +142,6 @@ tokens/s for long context. `baseline` uses Triton without the feature flags. | sync self-speculation | 8 | 242 | 615 | 1205 | | all features | 7 | 216 | 763 | 1128 | -The single-feature effects are small at this model and workload size. The -all-features configuration is 1.31x at short context with 128 concurrent -requests; graph replay and packed prefill are slightly slower in isolation -because capture and varlen setup dominate a 24-layer, four-loop model. -These numbers describe one device and are not a general performance claim. - The CPU regression suite passes unchanged on ROCm. GPU tests use the same `--run-gpu` flag; FlashAttention, FA4-specific packed-prefill and multi-GPU PD tests skip because those implementations or devices are unavailable. From 5f6d24a499c1134595c1ad5023fa2353614d293c Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 20:48:15 +0100 Subject: [PATCH 13/22] docs(rocm): add decode-heavy and batch-512 baselines Batch 512 all-features reaches 3760 tok/s (2.72x) and the 256-output-token decode run is 1.75x over baseline; short cells are prefill-heavy. --- docs/rocm.md | 12 +++++++++--- 1 file changed, 9 insertions(+), 3 deletions(-) diff --git a/docs/rocm.md b/docs/rocm.md index 3ea9c583..4384a6db 100644 --- a/docs/rocm.md +++ b/docs/rocm.md @@ -103,13 +103,19 @@ mode for both arms): | Workload | Baseline tok/s | All features tok/s | Speedup | | --- | ---: | ---: | ---: | -| short, batch 128 | 977 | 2733 | **2.80x** | -| short, batch 256 | 1136 | 2757 | **2.43x** | +| short, batch 128, 32 out | 977 | 2733 | **2.80x** | +| short, batch 256, 32 out | 1136 | 2757 | **2.43x** | +| short, batch 512, 32 out | 1384 | 3760 | **2.72x** | +| short, batch 128, 256 out | 1621 | 2832 | **1.75x** | `all` enables async scheduling, CUDA Graphs, static buffers, power-of-two padding, UVA, packed prefill, WCPB and prefix cache. The gain is mostly graph replay plus batching: per-token host work, not GEMM throughput, was the -bottleneck. +bottleneck. Short runs are prefill-heavy (80 prompt tokens for 32 output +tokens), so the speedup is larger there; with 256 output tokens the +decode-heavy speedup is 1.75x (baseline 6.3 ms/step, all 3.6 ms/step at batch +128). Reducing per-step host synchronization and fusing elementwise kernels are +the remaining follow-ups. Micro-tuning on the same `all` configuration (batch 128): From ac6d12760433abe8fc30234f315454e891ab51f7 Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 21:06:25 +0100 Subject: [PATCH 14/22] docs(rocm): mark packed prefill experimental and refresh warmed A/B Packed engine parity fails on a real Ouro prompt despite passing kernel tests; exclude it from all-features and record the warmed per-feature matrix (WCPB 1.45x, all 2.86x at batch 128, decode-heavy 1.72x). --- docs/rocm.md | 92 +++++++++++++++++++++++++++++++--------------------- 1 file changed, 55 insertions(+), 37 deletions(-) diff --git a/docs/rocm.md b/docs/rocm.md index 4384a6db..446f73db 100644 --- a/docs/rocm.md +++ b/docs/rocm.md @@ -60,7 +60,7 @@ device selector because PyTorch exposes ROCm devices through `torch.cuda`. | CUDA Graph replay | Yes | Triton backend, static buffers | | Multi-stream execution | Yes | `multi_stream` | | UVA host metadata | Yes (probed) | `hipHostGetDevicePointer` must equal the host address | -| Packed prefill / `prefill_uva` | Yes (Triton) | Single-sequence validated; multi-sequence packed is not qualified | +| Packed prefill / `prefill_uva` | Experimental | Kernel tests pass, but the engine greedy path diverges from per-row on a real Ouro prompt; excluded from `all` | | Wavefront chunked prefill | Yes | `--wavefront-chunked-prefill` | | Prefix cache | Yes | `enable_prefix_caching=True` | | Preemption and recovery | Yes | `enable_preemption=True`, async recovery | @@ -98,55 +98,73 @@ The baseline is the same commit, model, device and wheel set with execution features, the same exit mode and `VLLM_RLT_UVA=0`. There was no earlier ROCm baseline, so every comparison here is ROCm-relative; CUDA numbers were not measured on this lease. One warm-up `generate` is discarded, then the -median of three runs is reported (32 output tokens, same `ouro_delayed` exit -mode for both arms): +median of three runs is reported (same `ouro_delayed` exit mode for both arms): | Workload | Baseline tok/s | All features tok/s | Speedup | | --- | ---: | ---: | ---: | -| short, batch 128, 32 out | 977 | 2733 | **2.80x** | -| short, batch 256, 32 out | 1136 | 2757 | **2.43x** | -| short, batch 512, 32 out | 1384 | 3760 | **2.72x** | -| short, batch 128, 256 out | 1621 | 2832 | **1.75x** | +| short, batch 128, 32 out | 977 | 2798 | **2.86x** | +| short, batch 256, 32 out | 1136 | 2660 | **2.34x** | +| short, batch 128, 256 out | 1621 | 2794 | **1.72x** | `all` enables async scheduling, CUDA Graphs, static buffers, power-of-two -padding, UVA, packed prefill, WCPB and prefix cache. The gain is mostly graph -replay plus batching: per-token host work, not GEMM throughput, was the -bottleneck. Short runs are prefill-heavy (80 prompt tokens for 32 output -tokens), so the speedup is larger there; with 256 output tokens the -decode-heavy speedup is 1.75x (baseline 6.3 ms/step, all 3.6 ms/step at batch -128). Reducing per-step host synchronization and fusing elementwise kernels are -the remaining follow-ups. +padding, UVA, WCPB and prefix cache. Packed prefill is excluded because its +engine parity is still failing (see below). The gain is graph replay plus +batching; short runs are prefill-heavy, so the 256-output decode run is the +conservative 1.72x result (baseline 6.3 ms/step, all 3.4 ms/step at batch 128). -Micro-tuning on the same `all` configuration (batch 128): +Micro-tuning on the same `all` configuration (batch 128, no packed): | Setting | tok/s | | --- | ---: | -| default | 2733 | +| default | 2798 | | `HIP_FORCE_DEV_KERNARG=1` | 2781 | | `HSA_ENABLE_INTERRUPT=0` | 2843 | | both | 2827 | -PyTorch already reports `preferred_blas_library=hipblaslt` on this stack; -forcing it explicitly did not help. `rocBLAS` fallback was not needed. - -### Per-feature cold matrix - -Single cold run per cell, no warm-up, 64 output tokens for short and 16 for -long; values include JIT and graph-capture cost and are useful for direction, -not as steady-state numbers. - -| Config | short b1 | short b32 | short b128 | long b8 prefill | -| --- | ---: | ---: | ---: | ---: | -| baseline | 7 | 205 | 580 | 1162 | -| async | 7 | 224 | 581 | 1158 | -| graphs | 8 | 201 | 511 | 1076 | -| static | 7 | 226 | 596 | 938 | -| uva | 7 | 212 | 565 | 1159 | -| packed (`prefill_uva`) | 8 | 192 | 556 | 931 | -| wcpb | 7 | 228 | 562 | 1046 | -| prefix cache | 6 | 214 | 613 | 1292 | -| sync self-speculation | 8 | 242 | 615 | 1205 | -| all features | 7 | 216 | 763 | 1128 | +PyTorch already reports `preferred_blas_library=hipblaslt`; forcing it +explicitly did not improve throughput, so no rocBLAS fallback is used. + +### Per-feature warmed A/B + +Batch 128, 32 output tokens, one warm-up discarded, median of three runs. +Non-spec configs use `baseline_delayed` (970 tok/s) as the reference; spec +configs use the `ouro` baseline (994 tok/s), which is the only legal exit mode +for them. + +| Config | tok/s | vs reference | +| --- | ---: | ---: | +| baseline (non-spec) | 970 | 1.00x | +| async scheduling | 1025 | 1.06x | +| CUDA Graphs | 986 | 1.02x | +| static buffers + pad2 | 1056 | 1.09x | +| UVA | 1019 | 1.05x | +| WCPB | 1402 | **1.45x** | +| prefix cache | 1105 | 1.14x | +| all features | 2798 | **2.88x** | +| baseline (spec family) | 994 | 1.00x | +| sync self-speculation | 510 | 0.51x | +| token pipeline | 947 | 0.95x | +| async self-speculation | 883 | 0.89x | + +Speculation is a loss at 1.4B/24-layer: draft/verify overhead exceeds the +accepted-token gain, so it stays opt-in. WCPB is the largest single win because +it overlaps prefill chunks across recurrent depths. + +### Not done on ROCm + +- **Packed prefill (`prefill_uva`)**: kernel unit tests match the Torch + reference, but the engine path produces different greedy tokens from per-row + on a real Ouro prompt even with UVA disabled and capability forced. It stays + opt-in and is excluded from `all`; the engine-side root cause is open. +- **PD/KV transfer**: needs two GPUs and a ROCm NIXL/UCX transport. +- **KV quantization / INT4 residual cache**: not present on the current main + runtime; tracked by RFC #156. +- **Profiling flags**: `--profile` / `--profile-activities` are untested on + ROCm; `rocprofv3` is available for external traces. +- **Priority scheduling and incremental KV**: suite coverage only, no real-model + A/B. +- **FlashAttention**: NVIDIA-only by construction. +- **Multi-node**: out of scope for this device. The CPU regression suite passes unchanged on ROCm. GPU tests use the same `--run-gpu` flag; FlashAttention, FA4-specific packed-prefill and multi-GPU PD tests From de59e027cc47c3a6cf00f5cb98281b9f62e548f3 Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 21:22:59 +0100 Subject: [PATCH 15/22] docs(rocm): add long-context multi-batch sweep and speculation tuning Long-context high-concurrency reaches 4.82x at batch 64 (27102 prefill tok/s). Speculation stays opt-in: graphs and K/draft tuning still lose to baseline (best 0.95x). --- docs/rocm.md | 29 +++++++++++++++++++++-------- 1 file changed, 21 insertions(+), 8 deletions(-) diff --git a/docs/rocm.md b/docs/rocm.md index 446f73db..f3fca214 100644 --- a/docs/rocm.md +++ b/docs/rocm.md @@ -64,9 +64,9 @@ device selector because PyTorch exposes ROCm devices through `torch.cuda`. | Wavefront chunked prefill | Yes | `--wavefront-chunked-prefill` | | Prefix cache | Yes | `enable_prefix_caching=True` | | Preemption and recovery | Yes | `enable_preemption=True`, async recovery | -| Token pipeline | Yes | `SpeculativeConfig(mode="pipeline")` | -| Synchronous self-speculation | Yes | Block mode, fixed depth | -| Asynchronous self-speculation | Yes | Async suite coverage | +| Token pipeline | Opt-in | `SpeculativeConfig(mode="pipeline")`; 0.95x at 1.4B | +| Synchronous self-speculation | Opt-in | Block mode; 0.51x at 1.4B, K1/draft1 0.63x | +| Asynchronous self-speculation | Opt-in | Async suite coverage; 0.89x at 1.4B | | PD (prefill/decode split) | Not validated | Depends on the NIXL/UCX transport | | HTTP serving | Yes | `/v1/completions` | @@ -102,9 +102,19 @@ median of three runs is reported (same `ouro_delayed` exit mode for both arms): | Workload | Baseline tok/s | All features tok/s | Speedup | | --- | ---: | ---: | ---: | -| short, batch 128, 32 out | 977 | 2798 | **2.86x** | -| short, batch 256, 32 out | 1136 | 2660 | **2.34x** | -| short, batch 128, 256 out | 1621 | 2794 | **1.72x** | +| short, batch 1, 64 out | 18.4 | 32.8 | **1.78x** | +| short, batch 32, 64 out | 446 | 916 | **2.05x** | +| short, batch 128, 64 out | 1148 | 2909 | **2.53x** | +| short, batch 256, 64 out | 1606 | 3270 | **2.04x** | +| short, batch 512, 64 out | 1864 | 3919 | **2.10x** | +| long, batch 8, 32 out | 84.8 | 171.0 | **2.02x** | +| long, batch 32, 32 out | 153.8 | 544.4 | **3.54x** | +| long, batch 64, 32 out | 179.8 | 866.4 | **4.82x** | + +Long-context prefill throughput rises from 2652 to 5350 tok/s at batch 8, +4811 to 17029 at batch 32, and 5626 to 27102 at batch 64. The long-context +high-concurrency regime is where WCPB, async scheduling and graph replay +compound. `all` enables async scheduling, CUDA Graphs, static buffers, power-of-two padding, UVA, WCPB and prefix cache. Packed prefill is excluded because its @@ -147,8 +157,11 @@ for them. | async self-speculation | 883 | 0.89x | Speculation is a loss at 1.4B/24-layer: draft/verify overhead exceeds the -accepted-token gain, so it stays opt-in. WCPB is the largest single win because -it overlaps prefill chunks across recurrent depths. +accepted-token gain, so it stays opt-in. Tuning did not recover it: with graphs +and `cuda_graph_max_batch_size` covering verification rows, block K2 reached +495 tok/s, block K1/draft1 629, and pipeline K4 792, all below the 999 baseline; +the best remains pipeline K4 without graphs at 947 (0.95x). WCPB is the largest +single win because it overlaps prefill chunks across recurrent depths. ### Not done on ROCm From 0b2eff04d8eba8dc98a81b61ff5dc250841f0f5c Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 21:32:31 +0100 Subject: [PATCH 16/22] docs(rocm): add priority, incremental KV and sparse decode coverage Priority scheduling cuts high-priority completion 7.80s to 2.22s; incremental KV gives +3.8%; sparse decode #157 passes 137 tests on the merged ROCm branch. --- docs/rocm.md | 13 ++++++++++--- 1 file changed, 10 insertions(+), 3 deletions(-) diff --git a/docs/rocm.md b/docs/rocm.md index f3fca214..34e2e10e 100644 --- a/docs/rocm.md +++ b/docs/rocm.md @@ -63,6 +63,9 @@ device selector because PyTorch exposes ROCm devices through `torch.cuda`. | Packed prefill / `prefill_uva` | Experimental | Kernel tests pass, but the engine greedy path diverges from per-row on a real Ouro prompt; excluded from `all` | | Wavefront chunked prefill | Yes | `--wavefront-chunked-prefill` | | Prefix cache | Yes | `enable_prefix_caching=True` | +| Priority scheduling | Yes | `priority=0` is highest; high-priority latency 7.80s → 2.22s | +| Incremental KV | Yes | +3.8% at batch 64 | +| Sparse decode (#157) | ROCm tests pass | 137 passed on the merged branch; not part of this PR | | Preemption and recovery | Yes | `enable_preemption=True`, async recovery | | Token pipeline | Opt-in | `SpeculativeConfig(mode="pipeline")`; 0.95x at 1.4B | | Synchronous self-speculation | Opt-in | Block mode; 0.51x at 1.4B, K1/draft1 0.63x | @@ -156,6 +159,11 @@ for them. | token pipeline | 947 | 0.95x | | async self-speculation | 883 | 0.89x | +Priority and incremental KV are positive: with 64 low- and 8 high-priority +requests at 16 slots, high-priority median completion drops from 7.80s (fcfs) +to 2.22s (priority) while total time moves 7.81s to 8.15s; incremental KV gives +817 -> 848 tok/s at batch 64 (+3.8%). + Speculation is a loss at 1.4B/24-layer: draft/verify overhead exceeds the accepted-token gain, so it stays opt-in. Tuning did not recover it: with graphs and `cuda_graph_max_batch_size` covering verification rows, block K2 reached @@ -171,11 +179,10 @@ single win because it overlaps prefill chunks across recurrent depths. opt-in and is excluded from `all`; the engine-side root cause is open. - **PD/KV transfer**: needs two GPUs and a ROCm NIXL/UCX transport. - **KV quantization / INT4 residual cache**: not present on the current main - runtime; tracked by RFC #156. + runtime; tracked by RFC #156. Sparse decode from #157 passes its ROCm suites + on a stacked branch but is not shipped here. - **Profiling flags**: `--profile` / `--profile-activities` are untested on ROCm; `rocprofv3` is available for external traces. -- **Priority scheduling and incremental KV**: suite coverage only, no real-model - A/B. - **FlashAttention**: NVIDIA-only by construction. - **Multi-node**: out of scope for this device. From 22e25b40a5978f4e6e0a2ba871bd87291c1c102f Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 21:36:58 +0100 Subject: [PATCH 17/22] docs(rocm): record positive 2.6B speculation result Pipeline K2/draft1 gives 814 vs 500 tok/s (1.63x) on Ouro-2.6B at batch 64/128 output, with 0.797 greedy match; 1.4B remains negative. --- docs/rocm.md | 13 +++++++++---- 1 file changed, 9 insertions(+), 4 deletions(-) diff --git a/docs/rocm.md b/docs/rocm.md index 34e2e10e..9e9da6ff 100644 --- a/docs/rocm.md +++ b/docs/rocm.md @@ -67,9 +67,9 @@ device selector because PyTorch exposes ROCm devices through `torch.cuda`. | Incremental KV | Yes | +3.8% at batch 64 | | Sparse decode (#157) | ROCm tests pass | 137 passed on the merged branch; not part of this PR | | Preemption and recovery | Yes | `enable_preemption=True`, async recovery | -| Token pipeline | Opt-in | `SpeculativeConfig(mode="pipeline")`; 0.95x at 1.4B | -| Synchronous self-speculation | Opt-in | Block mode; 0.51x at 1.4B, K1/draft1 0.63x | -| Asynchronous self-speculation | Opt-in | Async suite coverage; 0.89x at 1.4B | +| Token pipeline | Opt-in | 0.95x at 1.4B; **1.63x on Ouro-2.6B** (K2/draft1) | +| Synchronous self-speculation | Opt-in | 0.51x at 1.4B; 1.18x on Ouro-2.6B (block K2/draft1) | +| Asynchronous self-speculation | Opt-in | Async suite; pipeline path is the positive case | | PD (prefill/decode split) | Not validated | Depends on the NIXL/UCX transport | | HTTP serving | Yes | `/v1/completions` | @@ -164,7 +164,12 @@ requests at 16 slots, high-priority median completion drops from 7.80s (fcfs) to 2.22s (priority) while total time moves 7.81s to 8.15s; incremental KV gives 817 -> 848 tok/s at batch 64 (+3.8%). -Speculation is a loss at 1.4B/24-layer: draft/verify overhead exceeds the +Speculation is model-dependent: it loses at 1.4B/24-layer but wins on Ouro-2.6B. +Pipeline K2/draft1 reaches **814 tok/s vs 500 baseline (1.63x)** at batch 64, +128 output tokens; K4/draft1 is 782 (1.56x) and block K2/draft1 is 591 (1.18x). +Greedy token match against the non-speculative baseline was 0.797 mean on four +prompts, so the speedup is approximate and needs a quality check before default +use. At 1.4B, speculation is still a loss: draft/verify overhead exceeds the accepted-token gain, so it stays opt-in. Tuning did not recover it: with graphs and `cuda_graph_max_batch_size` covering verification rows, block K2 reached 495 tok/s, block K1/draft1 629, and pipeline K4 792, all below the 999 baseline; From aab62e3d95e69119c963d973989f600ce895ec94 Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 21:43:37 +0100 Subject: [PATCH 18/22] docs(rocm): locate packed parity and record speculation quality gate Packed prefill: FP32 greedy matches per-row; BF16 differs by <=0.0156 and stays opt-in. Ouro-2.6B speculation is 1.63x but GSM8K 0/20 both arms, so it is not default. --- docs/rocm.md | 17 +++++++++-------- 1 file changed, 9 insertions(+), 8 deletions(-) diff --git a/docs/rocm.md b/docs/rocm.md index 9e9da6ff..f0175870 100644 --- a/docs/rocm.md +++ b/docs/rocm.md @@ -60,12 +60,12 @@ device selector because PyTorch exposes ROCm devices through `torch.cuda`. | CUDA Graph replay | Yes | Triton backend, static buffers | | Multi-stream execution | Yes | `multi_stream` | | UVA host metadata | Yes (probed) | `hipHostGetDevicePointer` must equal the host address | -| Packed prefill / `prefill_uva` | Experimental | Kernel tests pass, but the engine greedy path diverges from per-row on a real Ouro prompt; excluded from `all` | +| Packed prefill / `prefill_uva` | Opt-in | FP32 greedy matches per-row exactly; BF16 attention differs by <=0.0156 and greedy can diverge, so it stays out of `all` until an accuracy check | | Wavefront chunked prefill | Yes | `--wavefront-chunked-prefill` | | Prefix cache | Yes | `enable_prefix_caching=True` | | Priority scheduling | Yes | `priority=0` is highest; high-priority latency 7.80s → 2.22s | | Incremental KV | Yes | +3.8% at batch 64 | -| Sparse decode (#157) | ROCm tests pass | 137 passed on the merged branch; not part of this PR | +| Sparse decode (#157) | ROCm tests pass | `--sparse-decode-attention --sparse-decode-key-retention` on the #157 branch; 137 passed | | Preemption and recovery | Yes | `enable_preemption=True`, async recovery | | Token pipeline | Opt-in | 0.95x at 1.4B; **1.63x on Ouro-2.6B** (K2/draft1) | | Synchronous self-speculation | Opt-in | 0.51x at 1.4B; 1.18x on Ouro-2.6B (block K2/draft1) | @@ -168,8 +168,9 @@ Speculation is model-dependent: it loses at 1.4B/24-layer but wins on Ouro-2.6B. Pipeline K2/draft1 reaches **814 tok/s vs 500 baseline (1.63x)** at batch 64, 128 output tokens; K4/draft1 is 782 (1.56x) and block K2/draft1 is 591 (1.18x). Greedy token match against the non-speculative baseline was 0.797 mean on four -prompts, so the speedup is approximate and needs a quality check before default -use. At 1.4B, speculation is still a loss: draft/verify overhead exceeds the +prompts. A 20-question GSM8K check gave 0/20 for both arms (the base checkpoint +needs the repo's lm-eval recipe), so there is no measurable accuracy signal yet; +the speedup stays opt-in and is not a default. At 1.4B, speculation is still a loss: draft/verify overhead exceeds the accepted-token gain, so it stays opt-in. Tuning did not recover it: with graphs and `cuda_graph_max_batch_size` covering verification rows, block K2 reached 495 tok/s, block K1/draft1 629, and pipeline K4 792, all below the 999 baseline; @@ -178,10 +179,10 @@ single win because it overlaps prefill chunks across recurrent depths. ### Not done on ROCm -- **Packed prefill (`prefill_uva`)**: kernel unit tests match the Torch - reference, but the engine path produces different greedy tokens from per-row - on a real Ouro prompt even with UVA disabled and capability forced. It stays - opt-in and is excluded from `all`; the engine-side root cause is open. +- **Packed prefill (`prefill_uva`)**: root cause located as BF16 accumulation, + not metadata: FP32 greedy output matches per-row exactly, while BF16 + per-call attention differs by <=0.0156 and can flip greedy tokens. It stays + opt-in until an accuracy evaluation, and is excluded from `all`. - **PD/KV transfer**: needs two GPUs and a ROCm NIXL/UCX transport. - **KV quantization / INT4 residual cache**: not present on the current main runtime; tracked by RFC #156. Sparse decode from #157 passes its ROCm suites From 44e7bd92a2011b3089f0057f41e673bccd5c9f18 Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 21:56:27 +0100 Subject: [PATCH 19/22] docs(rocm): show decode-heavy speculation is positive 1.4B 256-output: pipeline K4/draft1 3270 (2842 confirmed), block K2/draft1 2148/2251, async K2/draft1 2314 vs 1527 baseline; 32-output remains prefill-dominated and loses. --- docs/rocm.md | 43 ++++++++++++++++++++++++++++--------------- 1 file changed, 28 insertions(+), 15 deletions(-) diff --git a/docs/rocm.md b/docs/rocm.md index f0175870..73ffc01d 100644 --- a/docs/rocm.md +++ b/docs/rocm.md @@ -67,9 +67,9 @@ device selector because PyTorch exposes ROCm devices through `torch.cuda`. | Incremental KV | Yes | +3.8% at batch 64 | | Sparse decode (#157) | ROCm tests pass | `--sparse-decode-attention --sparse-decode-key-retention` on the #157 branch; 137 passed | | Preemption and recovery | Yes | `enable_preemption=True`, async recovery | -| Token pipeline | Opt-in | 0.95x at 1.4B; **1.63x on Ouro-2.6B** (K2/draft1) | -| Synchronous self-speculation | Opt-in | 0.51x at 1.4B; 1.18x on Ouro-2.6B (block K2/draft1) | -| Asynchronous self-speculation | Opt-in | Async suite; pipeline path is the positive case | +| Token pipeline | Opt-in | **2.14x at 1.4B (256 out)**, 1.63x on Ouro-2.6B; short-output loses | +| Synchronous self-speculation | Opt-in | 1.41x at 1.4B (256 out); 1.18x on Ouro-2.6B | +| Asynchronous self-speculation | Opt-in | 1.48x at 1.4B (256 out); async suite | | PD (prefill/decode split) | Not validated | Depends on the NIXL/UCX transport | | HTTP serving | Yes | `/v1/completions` | @@ -164,18 +164,31 @@ requests at 16 slots, high-priority median completion drops from 7.80s (fcfs) to 2.22s (priority) while total time moves 7.81s to 8.15s; incremental KV gives 817 -> 848 tok/s at batch 64 (+3.8%). -Speculation is model-dependent: it loses at 1.4B/24-layer but wins on Ouro-2.6B. -Pipeline K2/draft1 reaches **814 tok/s vs 500 baseline (1.63x)** at batch 64, -128 output tokens; K4/draft1 is 782 (1.56x) and block K2/draft1 is 591 (1.18x). -Greedy token match against the non-speculative baseline was 0.797 mean on four -prompts. A 20-question GSM8K check gave 0/20 for both arms (the base checkpoint -needs the repo's lm-eval recipe), so there is no measurable accuracy signal yet; -the speedup stays opt-in and is not a default. At 1.4B, speculation is still a loss: draft/verify overhead exceeds the -accepted-token gain, so it stays opt-in. Tuning did not recover it: with graphs -and `cuda_graph_max_batch_size` covering verification rows, block K2 reached -495 tok/s, block K1/draft1 629, and pipeline K4 792, all below the 999 baseline; -the best remains pipeline K4 without graphs at 947 (0.95x). WCPB is the largest -single win because it overlaps prefill chunks across recurrent depths. +Speculation is **workload-dependent**, not simply negative. At 32 output +tokens the 80-token prefill dominates and it loses (best 0.95x); with 256 +output tokens at batch 128 the decode-heavy path wins on Ouro-1.4B: + +| Config | tok/s | vs baseline | +| --- | ---: | ---: | +| baseline | 1527 | 1.00x | +| pipeline K4/draft1 | 3270 | **2.14x** | +| pipeline K6/draft1 | 3249 | 2.13x | +| pipeline K2/draft1 | 2365 | 1.55x | +| pipeline K4/draft2 | 2254 | 1.48x | +| pipeline K2/draft2 | 2251 | 1.47x | +| block K2/draft1 | 2148 | **1.41x** | +| pipeline K1/draft1 | 1905 | 1.25x | +| async K2/draft1 | 2314 | 1.48x | + +A confirmation run measured pipeline K4/draft1 at 2842 tok/s (1.82x), block +K2/draft1 at 2251 (1.44x), and async K2/draft1 at 2314 (1.48x), with greedy +token match 0.66-0.74 against baseline. + +On Ouro-2.6B (batch 64, 128 output): pipeline K2/draft1 814 vs 500 (1.63x), +K4/draft1 782 (1.56x), block K2/draft1 591 (1.18x). Greedy token match against +the non-speculative baseline is 0.65-0.80, and a 20-question GSM8K check scored +0/20 for both arms, so there is no accuracy signal yet; speculation stays +opt-in and is not a default. ### Not done on ROCm From 6e746f06bdb93d04c0f56dfa3f6367b3f46a9aae Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 22:12:56 +0100 Subject: [PATCH 20/22] docs(rocm): add full-feature batch and context sweep runtime reaches 2.59x on short b256 and 5.73x on 1k b128; speculative long-context reaches 16.2x at 4k b8 (37150 prefill tok/s). --- docs/rocm.md | 22 ++++++++++++++++++++++ 1 file changed, 22 insertions(+) diff --git a/docs/rocm.md b/docs/rocm.md index 73ffc01d..fec0d4eb 100644 --- a/docs/rocm.md +++ b/docs/rocm.md @@ -137,6 +137,28 @@ Micro-tuning on the same `all` configuration (batch 128, no packed): PyTorch already reports `preferred_blas_library=hipblaslt`; forcing it explicitly did not improve throughput, so no rocBLAS fallback is used. +### Full-feature sweep + +Batch/concurrency and context sweep on one MI350X with Ouro-1.4B. `runtime` +enables async + graphs + static/pad2 + UVA + WCPB + prefix; `spec` adds +pipeline K4/draft1 on top. Short = 80-token prompt / 64 output; 1k = 1000-token +prompt / 32 output; 4k = 4000-token prompt / 32 output. + +| Workload | Baseline | runtime | spec | Best speedup | +| --- | ---: | ---: | ---: | ---: | +| short b128 | 1194 | 2208 | 2069 | 1.85x | +| short b256 | 1670 | 4329 | 2126 | 2.59x | +| short b512 | 2008 | 4184 | 2031 | 2.08x | +| 1k b32 | 155.7 | 407.6 | 761.0 | 4.89x | +| 1k b64 | 176.5 | 661.5 | 894.2 | 5.07x | +| 1k b128 | 177.9 | 1019.1 | 998.2 | 5.73x | +| 4k b8 | 18.3 | 149.3 | 297.1 | **16.2x** | + +4k prefill throughput goes 2294 -> 18665 (runtime) -> 37150 tok/s (spec). +Speculation wins once the context/output makes decode and verification +amortize; for short 64-token outputs runtime-only is better because prefill +dominates and spec verification adds rows. + ### Per-feature warmed A/B Batch 128, 32 output tokens, one warm-up discarded, median of three runs. From 96097b7d319f030901c1445575f2ba16cc2195f5 Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 22:29:52 +0100 Subject: [PATCH 21/22] docs(rocm): add 4k/8k and Nanbeige/Huginn results Ouro 8k spec reaches 21.4x (26800 prefill tok/s); Nanbeige 8k runtime 20.7x; Huginn is 4k-limited and graph replay gives 1.03x at batch 1. --- docs/rocm.md | 26 ++++++++++++++++++++++++++ 1 file changed, 26 insertions(+) diff --git a/docs/rocm.md b/docs/rocm.md index fec0d4eb..1723fa2e 100644 --- a/docs/rocm.md +++ b/docs/rocm.md @@ -159,6 +159,32 @@ Speculation wins once the context/output makes decode and verification amortize; for short 64-token outputs runtime-only is better because prefill dominates and spec verification adds rows. +### 4k/8k context and other checkpoints + +Ouro-1.4B, longer outputs and contexts: + +| Workload | Baseline | runtime | spec | Best | +| --- | ---: | ---: | ---: | ---: | +| 4k b4 / 128 out | 37.9 | 84.8 | 263.6 | **6.96x** | +| 4k b2 / 256 out | 28.7 | 54.7 | 168.4 | **5.87x** | +| 8k b2 / 32 out | 5.0 | 46.2 | 107.2 | **21.4x** | +| 8k b1 / 128 out | 10.1 | 27.7 | 80.3 | **7.95x** | + +8k prefill throughput goes 1247 -> 11544 -> 26800 tok/s; 4k goes 1185 -> +2650 -> 8239. Speculation keeps winning as context and output grow. + +Other checkpoints, 4k/8k baseline vs runtime (Nanbeige) or graphs (Huginn): + +| Model | Workload | Baseline | Optimized | Speedup | +| --- | --- | ---: | ---: | ---: | +| Nanbeige4.2-3B | 4k b4 / 64 out | 27.7 | 199.9 | **7.2x** | +| Nanbeige4.2-3B | 8k b2 / 32 out | 4.2 | 87.1 | **20.7x** | +| Huginn-0125 | 4k b1 / 32 out | 3.1 | 3.2 | 1.03x | + +Huginn is fixed-depth synchronous with 32 recurrent steps; graph replay adds +little at batch 1, and its model configuration rejects 8k prompts (prompt plus +decode exceeds the context limit), so only 4k is reported. + ### Per-feature warmed A/B Batch 128, 32 output tokens, one warm-up discarded, median of three runs. From 780579c95b06f38e9e1978ed970095db7b26ded3 Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 5 Nov 2026 22:47:28 +0100 Subject: [PATCH 22/22] docs(rocm): unify long-context table across models One table with model column for Ouro/Nanbeige/Huginn at 1k/4k/8k, plus Huginn 1k/4k multi-batch and updated spec rows (256-out positive, 32-out negative). --- docs/rocm.md | 77 +++++++++++++++++++++++----------------------------- 1 file changed, 34 insertions(+), 43 deletions(-) diff --git a/docs/rocm.md b/docs/rocm.md index 1723fa2e..7784c29d 100644 --- a/docs/rocm.md +++ b/docs/rocm.md @@ -139,51 +139,40 @@ explicitly did not improve throughput, so no rocBLAS fallback is used. ### Full-feature sweep -Batch/concurrency and context sweep on one MI350X with Ouro-1.4B. `runtime` -enables async + graphs + static/pad2 + UVA + WCPB + prefix; `spec` adds -pipeline K4/draft1 on top. Short = 80-token prompt / 64 output; 1k = 1000-token -prompt / 32 output; 4k = 4000-token prompt / 32 output. +Short context (80-token prompt, 64 output), Ouro-1.4B: -| Workload | Baseline | runtime | spec | Best speedup | +| Workload | Baseline | runtime | spec | Best | | --- | ---: | ---: | ---: | ---: | | short b128 | 1194 | 2208 | 2069 | 1.85x | | short b256 | 1670 | 4329 | 2126 | 2.59x | | short b512 | 2008 | 4184 | 2031 | 2.08x | -| 1k b32 | 155.7 | 407.6 | 761.0 | 4.89x | -| 1k b64 | 176.5 | 661.5 | 894.2 | 5.07x | -| 1k b128 | 177.9 | 1019.1 | 998.2 | 5.73x | -| 4k b8 | 18.3 | 149.3 | 297.1 | **16.2x** | - -4k prefill throughput goes 2294 -> 18665 (runtime) -> 37150 tok/s (spec). -Speculation wins once the context/output makes decode and verification -amortize; for short 64-token outputs runtime-only is better because prefill -dominates and spec verification adds rows. - -### 4k/8k context and other checkpoints -Ouro-1.4B, longer outputs and contexts: - -| Workload | Baseline | runtime | spec | Best | -| --- | ---: | ---: | ---: | ---: | -| 4k b4 / 128 out | 37.9 | 84.8 | 263.6 | **6.96x** | -| 4k b2 / 256 out | 28.7 | 54.7 | 168.4 | **5.87x** | -| 8k b2 / 32 out | 5.0 | 46.2 | 107.2 | **21.4x** | -| 8k b1 / 128 out | 10.1 | 27.7 | 80.3 | **7.95x** | - -8k prefill throughput goes 1247 -> 11544 -> 26800 tok/s; 4k goes 1185 -> -2650 -> 8239. Speculation keeps winning as context and output grow. - -Other checkpoints, 4k/8k baseline vs runtime (Nanbeige) or graphs (Huginn): - -| Model | Workload | Baseline | Optimized | Speedup | -| --- | --- | ---: | ---: | ---: | -| Nanbeige4.2-3B | 4k b4 / 64 out | 27.7 | 199.9 | **7.2x** | -| Nanbeige4.2-3B | 8k b2 / 32 out | 4.2 | 87.1 | **20.7x** | -| Huginn-0125 | 4k b1 / 32 out | 3.1 | 3.2 | 1.03x | - -Huginn is fixed-depth synchronous with 32 recurrent steps; graph replay adds -little at batch 1, and its model configuration rejects 8k prompts (prompt plus -decode exceeds the context limit), so only 4k is reported. +Long context and checkpoints (1k/4k/8k). `runtime` = async + graphs + static/pad2 + UVA + WCPB + prefix; `spec` adds pipeline K4/draft1. Huginn's optimized column is graph replay (it rejects async/spec); Nanbeige has no spec path: + +| Model | Context | Batch | Out | Baseline | runtime | spec | Best | +| --- | ---: | ---: | ---: | ---: | ---: | ---: | ---: | +| Ouro-1.4B | 1k | 32 | 32 | 155.7 | 407.6 | 761.0 | 4.89x | +| Ouro-1.4B | 1k | 64 | 32 | 176.5 | 661.5 | 894.2 | 5.07x | +| Ouro-1.4B | 1k | 128 | 32 | 177.9 | 1019.1 | 998.2 | 5.73x | +| Ouro-1.4B | 4k | 4 | 128 | 37.9 | 84.8 | 263.6 | 6.96x | +| Ouro-1.4B | 4k | 2 | 256 | 28.7 | 54.7 | 168.4 | 5.87x | +| Ouro-1.4B | 8k | 2 | 32 | 5.0 | 46.2 | 107.2 | **21.4x** | +| Ouro-1.4B | 8k | 1 | 128 | 10.1 | 27.7 | 80.3 | 7.95x | +| Nanbeige4.2-3B | 1k | 32 | 32 | 167.4 | 854.7 | — | 5.11x | +| Nanbeige4.2-3B | 1k | 64 | 32 | 176.4 | 1083.7 | — | 6.14x | +| Nanbeige4.2-3B | 1k | 128 | 32 | 177.2 | 1464.8 | — | 8.27x | +| Nanbeige4.2-3B | 4k | 4 | 64 | 27.7 | 199.9 | — | 7.2x | +| Nanbeige4.2-3B | 8k | 2 | 32 | 4.2 | 87.1 | — | 20.7x | +| Huginn-0125 | 1k | 1 | 32 | 8.8 | 9.5 | — | 1.08x | +| Huginn-0125 | 1k | 2 | 32 | 13.6 | 14.8 | — | 1.09x | +| Huginn-0125 | 1k | 4 | 32 | 14.0 | 14.7 | — | 1.05x | +| Huginn-0125 | 1k | 8 | 32 | 16.5 | — | — | — | +| Huginn-0125 | 4k | 1 | 32 | 3.1 | 3.2 | — | 1.03x | +| Huginn-0125 | 4k | 2 | 32 | 3.1 | — | — | — | +| Huginn-0125 | 4k | 4 | 32 | 3.1 | — | — | — | +| Huginn-0125 | 4k | 8 | 32 | 3.1 | — | — | — | + +Huginn is compute-bound by its 32 recurrent steps (throughput flat across batch), graph replay adds 3-9%, and it rejects 8k prompts (context limit). 8k Ouro prefill goes 1247 -> 11544 -> 26800 tok/s. ### Per-feature warmed A/B @@ -202,10 +191,12 @@ for them. | WCPB | 1402 | **1.45x** | | prefix cache | 1105 | 1.14x | | all features | 2798 | **2.88x** | -| baseline (spec family) | 994 | 1.00x | -| sync self-speculation | 510 | 0.51x | -| token pipeline | 947 | 0.95x | -| async self-speculation | 883 | 0.89x | +| baseline (spec family, 256 out) | 1527 | 1.00x | +| sync self-speculation (256 out) | 2148 | 1.41x | +| token pipeline (256 out) | 3270 | 2.14x | +| async self-speculation (256 out) | 2314 | 1.48x | + +The spec rows are the decode-heavy 256-output case. At 32 output tokens the same configs lose (0.51x / 0.95x / 0.89x) because prefill dominates; both directions are kept in the workload-dependent section below. Priority and incremental KV are positive: with 64 low- and 8 high-priority requests at 16 slots, high-priority median completion drops from 7.80s (fcfs)