diff --git a/.gitignore b/.gitignore index bf804e07..0fb695f6 100644 --- a/.gitignore +++ b/.gitignore @@ -6,6 +6,13 @@ __pycache__/ # C extensions *.so +# torch.utils.cpp_extension's ROCm auto-hipify writes translated copies next to the +# CUDA sources it translates (kernel/csrc/gguf/*.cu -> *.hip, *.cuh -> *_hip.cuh); +# regenerated on every build, never hand-edited. +*.hip +*_hip.cuh +*_hip.h + # Distribution / packaging .Python build/ diff --git a/benchmarks/README.md b/benchmarks/README.md index 6218903f..e0d622e4 100644 --- a/benchmarks/README.md +++ b/benchmarks/README.md @@ -11,6 +11,11 @@ include the full serving path. AIME-25 prompt, checkpoint-recommended sampling. python benchmarks/bench_decode_moe.py --model /path/to/model --backend offload,cpu,hybrid ``` +Use `--cache N` to pin the expert-cache slot count and +`--num-token-override N` to pin the server KV-token pool. Supply both when +comparing cache policies so automatic spare-VRAM allocation does not change the +tested context capacity. + **`bench_load_weight_generic.py`** — expert-bank load time: serial vs parallel O_DIRECT vs pre-repacked FTW, each mode in its own subprocess. Linux-only; stages the FTW under `/var/tmp` (`--ftw-dir` overrides; roughly checkpoint-sized). diff --git a/benchmarks/bench_decode_moe.py b/benchmarks/bench_decode_moe.py index 56621792..313af333 100644 --- a/benchmarks/bench_decode_moe.py +++ b/benchmarks/bench_decode_moe.py @@ -100,6 +100,15 @@ def parse_args(argv: list[str] | None = None) -> argparse.Namespace: help="hybrid: max PCIe fetches/layer; -1 = auto (benched pcie/cpu bandwidth fraction)", ) p.add_argument("--mem-ratio", type=float, default=0.9, help="target VRAM utilization") + p.add_argument( + "--num-token-override", + type=int, + default=None, + help=( + "pin the server KV-token pool capacity instead of accepting its automatic " + "allocation; use this to compare cache policies at the same context capacity" + ), + ) p.add_argument("--gpu", default=None, help="GPU for the serve: a UUID or nvidia-smi index (as ft serve --gpu)") p.add_argument("--no-graph", action="store_true", help="eager decode instead of CUDA graph") @@ -187,6 +196,13 @@ def serve_cmd(args: argparse.Namespace, backend: str, port: int) -> list[str]: ] if args.gpu: cmd += ["--gpu", args.gpu] + # An explicit token-pool size makes cache-policy comparisons fair: auto cache + # sizing otherwise consumes the remaining VRAM for KV pages, while a fixed + # expert cache leaves the server's conservative default KV allocation intact. + if args.num_token_override is not None: + # The benchmark names the value after the Engine field, while the public + # CLI intentionally exposes it as the concise ``--num-tokens`` flag. + cmd += ["--num-tokens", str(args.num_token_override)] if args.cache > 0: cmd += ["--moe-cache-size", str(args.cache)] elif args.cache_rate is not None: diff --git a/benchmarks/bench_gguf_q4_dense_kernel.py b/benchmarks/bench_gguf_q4_dense_kernel.py new file mode 100644 index 00000000..c8327740 --- /dev/null +++ b/benchmarks/bench_gguf_q4_dense_kernel.py @@ -0,0 +1,137 @@ +"""Measure the dense native-GGUF Q4_0 vector kernels used by Gemma 4 on LAN-223. + +The Gemma 4 26B A4B Q4_0 checkpoint has four recurring dense projection +geometries. They are supplied as defaults here so a HIP optimization can be +measured before it is allowed into the full OpenAI-compatible server benchmark: + +* 2816 by 4096 attention output projection; +* 8192 by 2816 full-attention QKV projection; +* 4224 by 2816 fused shared-MLP gate/up projection; and +* 10240 by 2816 sliding-window QKV projection. + +Like ``bench_gguf_q4_moe_kernel.py``, this tool creates valid packed Q4_0 +weights on the accelerator and measures only post-warm-up GPU event time. It +does not claim an end-to-end serving rate and must be paired with the API +benchmark before a kernel candidate is accepted. +""" + +from __future__ import annotations + +import argparse +import json +from pathlib import Path + +import torch + +from freetoken.kernel.gguf import ggml_mul_mat_vec_a8 +from freetoken.models.gguf.dequant import GGML_Q4_0, row_bytes + + +# Output rows and input columns, recovered from the exact LAN-223 Gemma GGUF. +DEFAULT_SHAPES = ((2816, 4096), (8192, 2816), (4224, 2816), (10240, 2816)) + + +def _parse_shape(value: str) -> tuple[int, int]: + """Parse a ``ROWSxCOLS`` override and validate its Q4_0 block alignment.""" + try: + rows_text, cols_text = value.lower().split("x", 1) + rows, cols = int(rows_text), int(cols_text) + except ValueError as error: + raise argparse.ArgumentTypeError("shape must be ROWSxCOLS, for example 2816x4096") from error + if rows <= 0 or cols <= 0 or cols % 32: + raise argparse.ArgumentTypeError("rows must be positive and cols must be a positive multiple of 32") + return rows, cols + + +def _parse_args() -> argparse.Namespace: + """Parse reproducible dense-kernel benchmark controls.""" + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument( + "--shape", + action="append", + type=_parse_shape, + help="repeatable ROWSxCOLS override; defaults to all production shapes", + ) + parser.add_argument("--vectors", type=int, default=1, help="input rows per kernel call") + parser.add_argument("--warmup", type=int, default=20) + parser.add_argument("--repetitions", type=int, default=200) + parser.add_argument("--seed", type=int, default=20260828) + parser.add_argument("--json", type=Path, help="write one JSON artifact") + return parser.parse_args() + + +def _make_q4_weight(rows: int, cols: int, device: torch.device) -> torch.Tensor: + """Create finite, contiguous ``[rows, row_bytes(cols)]`` Q4_0 packed weights.""" + blocks = cols // 32 + weight = torch.randint(0, 256, (rows, blocks, 18), dtype=torch.uint8, device=device) + # Q4_0 starts each 18-byte block with an FP16 scale. Use 1/32 rather than + # arbitrary random bytes so the measured real kernel cannot create NaNs. + scale_bytes = torch.tensor([1.0 / 32.0], dtype=torch.float16, device=device).view(torch.uint8) + weight[..., :2] = scale_bytes.reshape(1, 1, 2) + return weight.reshape(rows, row_bytes(cols, GGML_Q4_0)).contiguous() + + +def _average_event_us(kernel, repetitions: int, device: torch.device) -> float: + """Measure an already-warmed kernel with GPU events and return microseconds/call.""" + start, end = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True) + torch.cuda.synchronize(device) + start.record() + for _ in range(repetitions): + kernel() + end.record() + end.synchronize() + return start.elapsed_time(end) * 1000.0 / repetitions + + +def main() -> int: + """Run the selected dense projection shapes and write a durable JSON result.""" + args = _parse_args() + if args.vectors <= 0 or args.warmup <= 0 or args.repetitions <= 0: + raise ValueError("--vectors, --warmup, and --repetitions must be positive") + if not torch.cuda.is_available(): + raise RuntimeError("this benchmark requires a CUDA or HIP PyTorch device") + torch.manual_seed(args.seed) + device = torch.device("cuda") + shapes = args.shape or DEFAULT_SHAPES + measurements = [] + + for rows, cols in shapes: + weight = _make_q4_weight(rows, cols, device) + x = torch.randn(args.vectors, cols, dtype=torch.bfloat16, device=device) + + def call() -> torch.Tensor: + return ggml_mul_mat_vec_a8(weight, x, int(GGML_Q4_0), rows) + + for _ in range(args.warmup): + result = call() + torch.cuda.synchronize(device) + if not torch.isfinite(result).all(): + raise RuntimeError(f"non-finite result for dense Q4_0 shape {rows}x{cols}") + measurements.append( + { + "rows": rows, + "cols": cols, + "vectors": args.vectors, + "output_shape": list(result.shape), + "average_us": _average_event_us(call, args.repetitions, device), + } + ) + + output = { + "device": torch.cuda.get_device_name(device), + "hip": torch.version.hip, + "torch": torch.__version__, + "quant_type": "Q4_0", + "warmup": args.warmup, + "repetitions": args.repetitions, + "measurements": measurements, + } + print(json.dumps(output, indent=2, sort_keys=True)) + if args.json is not None: + args.json.parent.mkdir(parents=True, exist_ok=True) + args.json.write_text(json.dumps(output, indent=2, sort_keys=True) + "\n", encoding="utf-8") + return 0 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/benchmarks/bench_gguf_q4_moe_kernel.py b/benchmarks/bench_gguf_q4_moe_kernel.py new file mode 100644 index 00000000..1a5a8a76 --- /dev/null +++ b/benchmarks/bench_gguf_q4_moe_kernel.py @@ -0,0 +1,184 @@ +"""Measure FreeToken's native GGUF Q4_0 MoE vector kernels in isolation. + +This benchmark deliberately uses the Gemma 4 26B A4B Q4_0 expert geometry +observed on LAN-223: 128 routed experts, top-k 8, hidden width 2816, and MoE +intermediate width 704. It is not a replacement for the end-to-end OpenAI API +benchmark. Instead, it supplies the kernel-level evidence needed before a HIP +port changes Q4_0 launch geometry, indexing, or register use. + +The benchmark creates valid packed Q4_0 rows directly on the GPU. Every block +has a finite FP16 scale and random packed nibbles, so the real production +``ggml_moe_a8_vec`` path, including activation quantization, runs without model +loading, host-cache copying, scheduler work, or HTTP overhead. CUDA events are +used only after warm-up and synchronization; compilation and allocation are not +included in the reported microseconds. +""" + +from __future__ import annotations + +import argparse +import json +from pathlib import Path +from time import perf_counter + +import torch + +from freetoken.kernel.gguf import ggml_moe_a8_vec +from freetoken.models.gguf.dequant import GGML_Q4_0, row_bytes + + +# These defaults are the verified LAN-223 Gemma 4 26B A4B Q4_0 dimensions. +DEFAULT_EXPERTS = 128 +DEFAULT_TOP_K = 8 +DEFAULT_HIDDEN = 2816 +DEFAULT_INTERMEDIATE = 704 + + +def _parse_args() -> argparse.Namespace: + """Parse only parameters that preserve a reproducible kernel experiment.""" + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--experts", type=int, default=DEFAULT_EXPERTS) + parser.add_argument("--top-k", type=int, default=DEFAULT_TOP_K) + parser.add_argument("--hidden", type=int, default=DEFAULT_HIDDEN) + parser.add_argument("--intermediate", type=int, default=DEFAULT_INTERMEDIATE) + parser.add_argument("--tokens", type=int, default=1, help="decoded token rows per call") + parser.add_argument("--warmup", type=int, default=20, help="unmeasured calls per kernel") + parser.add_argument("--repetitions", type=int, default=200, help="timed calls per kernel") + parser.add_argument("--seed", type=int, default=20260828) + parser.add_argument("--json", type=Path, help="write one reproducible JSON result") + return parser.parse_args() + + +def _require_valid_geometry(args: argparse.Namespace) -> None: + """Reject shapes that cannot be represented by the Q4_0 block format.""" + for name in ("hidden", "intermediate"): + value = getattr(args, name) + if value <= 0 or value % 32: + raise ValueError(f"--{name} must be a positive multiple of 32, got {value}") + for name in ("experts", "top_k", "tokens", "warmup", "repetitions"): + if getattr(args, name) <= 0: + raise ValueError(f"--{name.replace('_', '-')} must be positive") + if args.top_k > args.experts: + raise ValueError("--top-k cannot exceed --experts") + if not torch.cuda.is_available(): + raise RuntimeError("this benchmark requires a CUDA or HIP PyTorch device") + + +def _q4_scale_bytes(device: torch.device) -> torch.Tensor: + """Return little-endian bytes for a finite FP16 Q4_0 scale of 1/32. + + Q4_0 stores two FP16 scale bytes before every 16-byte packed-nibble payload. + A constant finite scale is sufficient for performance work and avoids random + bit patterns that could otherwise create NaNs during the warm-up kernel. + """ + scale = torch.tensor([1.0 / 32.0], dtype=torch.float16, device=device) + return scale.view(torch.uint8).reshape(2) + + +def _make_q4_bank(experts: int, rows: int, columns: int, device: torch.device) -> torch.Tensor: + """Create a contiguous GPU Q4_0 bank shaped exactly like an expert cache. + + The byte layout is ``[expert, output_row, columns//32, 18]`` before the + final view. Byte positions zero and one receive the valid scale, while the + remaining sixteen bytes contain arbitrary Q4_0 nibbles. The final shape + mirrors the packed tensors passed by the Gemma GGUF offload cache. + """ + packed_row_bytes = row_bytes(columns, GGML_Q4_0) + blocks = columns // 32 + bank = torch.randint( + 0, + 256, + (experts, rows, blocks, 18), + dtype=torch.uint8, + device=device, + ) + bank[..., :2] = _q4_scale_bytes(device) + return bank.reshape(experts, rows, packed_row_bytes).contiguous() + + +def _make_topk_ids(tokens: int, top_k: int, experts: int, device: torch.device) -> torch.Tensor: + """Create deterministic valid expert selections without invoking router code.""" + ids = torch.arange(tokens * top_k, dtype=torch.int32, device=device) + return (ids.remainder(experts)).reshape(tokens, top_k).contiguous() + + +def _event_time_us(callable_kernel, repetitions: int, device: torch.device) -> float: + """Return average GPU elapsed time per invocation after explicit synchronization.""" + start = torch.cuda.Event(enable_timing=True) + end = torch.cuda.Event(enable_timing=True) + torch.cuda.synchronize(device) + start.record() + for _ in range(repetitions): + callable_kernel() + end.record() + end.synchronize() + return start.elapsed_time(end) * 1000.0 / repetitions + + +def main() -> int: + """Build the two production-shaped calls, warm them, measure them, and emit JSON.""" + args = _parse_args() + _require_valid_geometry(args) + torch.manual_seed(args.seed) + device = torch.device("cuda") + + # Gate/up maps H to 2I and consumes one routing row for every selected expert. + hidden = torch.randn(args.tokens, args.hidden, device=device, dtype=torch.bfloat16) + gate_up = _make_q4_bank(args.experts, 2 * args.intermediate, args.hidden, device) + route_ids = _make_topk_ids(args.tokens, args.top_k, args.experts, device) + + def gate_up_call() -> torch.Tensor: + return ggml_moe_a8_vec( + hidden, gate_up, route_ids, args.top_k, int(GGML_Q4_0), 2 * args.intermediate, args.tokens + ) + + # Down maps I to H. Its input and routing layout match fused_q4_0.py exactly. + inter = torch.randn(args.tokens * args.top_k, args.intermediate, device=device, dtype=torch.bfloat16) + down = _make_q4_bank(args.experts, args.hidden, args.intermediate, device) + + def down_call() -> torch.Tensor: + return ggml_moe_a8_vec( + inter, down, route_ids, 1, int(GGML_Q4_0), args.hidden, args.tokens * args.top_k + ) + + # Materialize the extension and check that valid Q4_0 data produces finite outputs. + for _ in range(args.warmup): + gate_result = gate_up_call() + down_result = down_call() + torch.cuda.synchronize(device) + if not torch.isfinite(gate_result).all() or not torch.isfinite(down_result).all(): + raise RuntimeError("synthetic Q4_0 data produced a non-finite kernel result") + + wall_start = perf_counter() + gate_up_us = _event_time_us(gate_up_call, args.repetitions, device) + down_us = _event_time_us(down_call, args.repetitions, device) + torch.cuda.synchronize(device) + + result = { + "device": torch.cuda.get_device_name(device), + "hip": torch.version.hip, + "torch": torch.__version__, + "quant_type": "Q4_0", + "experts": args.experts, + "top_k": args.top_k, + "hidden": args.hidden, + "intermediate": args.intermediate, + "tokens": args.tokens, + "warmup": args.warmup, + "repetitions": args.repetitions, + "gate_up_us": gate_up_us, + "down_us": down_us, + "pair_us": gate_up_us + down_us, + "wall_seconds": perf_counter() - wall_start, + "gate_up_shape": list(gate_result.shape), + "down_shape": list(down_result.shape), + } + print(json.dumps(result, indent=2, sort_keys=True)) + if args.json is not None: + args.json.parent.mkdir(parents=True, exist_ok=True) + args.json.write_text(json.dumps(result, indent=2, sort_keys=True) + "\n", encoding="utf-8") + return 0 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/benchmarks/bench_rocm_gqa_attention.py b/benchmarks/bench_rocm_gqa_attention.py new file mode 100644 index 00000000..503b64d3 --- /dev/null +++ b/benchmarks/bench_rocm_gqa_attention.py @@ -0,0 +1,114 @@ +"""Benchmark the Gemma 4 ROCm GQA decode tile without changing serving defaults. + +Gemma 4's sliding attention is 16 query heads by 8 KV heads at head dimension +256. ROCm serving pads this group-of-two GQA tile to 16 query-head lanes so +Triton can lower ``tl.dot`` to RDNA WMMA. This tool calls the same attention +function twice on identical tensors: once with the default tile and once with +an explicitly requested HIP probe tile. It checks numerical agreement before +reporting GPU-event latency, so a compilation success alone is never treated +as an optimization result. +""" + +from __future__ import annotations + +import argparse +import json +from pathlib import Path + +import torch + +from freetoken.kernel.triton.attention import decode_paged_attention + + +def _parse_args() -> argparse.Namespace: + """Parse reproducible ROCm GQA tile benchmark controls.""" + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--probe-block-h", type=int) + parser.add_argument("--probe-block-n", type=int) + parser.add_argument("--probe-num-warps", type=int) + parser.add_argument("--sequence-length", type=int, default=1024) + parser.add_argument("--kv-heads", type=int, default=8) + parser.add_argument("--head-dim", type=int, default=256) + parser.add_argument("--sliding-window", type=int, default=1024, help="zero means full attention") + parser.add_argument("--warmup", type=int, default=20) + parser.add_argument("--repetitions", type=int, default=200) + parser.add_argument("--seed", type=int, default=20260829) + parser.add_argument("--json", type=Path) + return parser.parse_args() + + +def _event_us(call, repetitions: int) -> float: + """Return post-warm-up accelerator event time for an already-built kernel.""" + start, end = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True) + torch.cuda.synchronize() + start.record() + for _ in range(repetitions): + call() + end.record() + end.synchronize() + return start.elapsed_time(end) * 1000.0 / repetitions + + +def main() -> int: + """Execute the exact Gemma sliding-GQA decode comparison on the HIP device.""" + args = _parse_args() + if not torch.cuda.is_available() or torch.version.hip is None: + raise RuntimeError("this benchmark requires a HIP PyTorch device") + if args.sequence_length <= 0 or args.warmup <= 0 or args.repetitions <= 0: + raise ValueError("sequence length, warmup, and repetitions must be positive") + torch.manual_seed(args.seed) + device = torch.device("cuda") + batch, query_heads, kv_heads, head_dim, splits = 1, 16, args.kv_heads, args.head_dim, 8 + if query_heads % kv_heads: + raise ValueError("--kv-heads must divide Gemma's 16 query heads") + q = torch.randn(batch, query_heads, head_dim, dtype=torch.bfloat16, device=device) + k = torch.randn(args.sequence_length, kv_heads, head_dim, dtype=torch.bfloat16, device=device) + v = torch.randn_like(k) + indptr = torch.tensor([0, args.sequence_length], dtype=torch.int32, device=device) + indices = torch.arange(args.sequence_length, dtype=torch.int32, device=device) + positions = torch.tensor([args.sequence_length - 1], dtype=torch.int64, device=device) + mid_o = torch.empty(batch, query_heads, splits, head_dim, dtype=torch.float32, device=device) + mid_lse = torch.empty(batch, query_heads, splits, dtype=torch.float32, device=device) + num_splits = torch.full((batch,), splits, dtype=torch.int32, device=device) + + def call(probe_h: int | None, probe_n: int | None, probe_warps: int | None) -> torch.Tensor: + return decode_paged_attention( + q, k, v, indptr, indices, positions, mid_o, mid_lse, num_splits, + splits, head_dim**-0.5, + sliding_window=args.sliding_window or None, + rocm_block_h_probe=probe_h, + rocm_block_n_probe=probe_n, rocm_num_warps_probe=probe_warps, + ) + + for _ in range(args.warmup): + default = call(None, None, None) + for _ in range(args.warmup): + candidate = call(args.probe_block_h, args.probe_block_n, args.probe_num_warps) + torch.cuda.synchronize() + torch.testing.assert_close(candidate.float(), default.float(), atol=2e-2, rtol=2e-2) + result = { + "device": torch.cuda.get_device_name(device), + "hip": torch.version.hip, + "geometry": {"q_heads": query_heads, "kv_heads": kv_heads, "head_dim": head_dim}, + "sequence_length": args.sequence_length, + "sliding_window": args.sliding_window or None, + "probe_block_h": args.probe_block_h, + "probe_block_n": args.probe_block_n, + "probe_num_warps": args.probe_num_warps, + "default_us": _event_us(lambda: call(None, None, None), args.repetitions), + "probe_us": _event_us( + lambda: call(args.probe_block_h, args.probe_block_n, args.probe_num_warps), + args.repetitions, + ), + "warmup": args.warmup, + "repetitions": args.repetitions, + } + print(json.dumps(result, indent=2, sort_keys=True)) + if args.json: + args.json.parent.mkdir(parents=True, exist_ok=True) + args.json.write_text(json.dumps(result, indent=2, sort_keys=True) + "\n", encoding="utf-8") + return 0 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/docs/amd-rocm-gfx1151.md b/docs/amd-rocm-gfx1151.md new file mode 100644 index 00000000..00c14aa0 --- /dev/null +++ b/docs/amd-rocm-gfx1151.md @@ -0,0 +1,132 @@ +# FreeToken AMD ROCm on Radeon 8060S `gfx1151` + +## Purpose + +This branch ports the FreeToken serving runtime to native AMD ROCm and HIP on +the AMD Ryzen AI Max+ 395 with Radeon 8060S (`gfx1151`). The port preserves +the NVIDIA implementation as a separate runtime path. It does not use Vulkan +or a CPU-only runner as a substitute for native GPU execution. + +The intended first deployment host is LAN-223. It serves the same local API +surface as upstream FreeToken, including OpenAI-compatible endpoints, while +using HIP-compiled extensions and AMD Triton kernels. + +## Scope and parity contract + +The port is complete only when the target model can load and serve through +`ft serve`, return a coherent streamed and non-streamed OpenAI-compatible +response, and exercise the applicable FreeToken cache and MoE paths. The +initial full-model validation set is: + +1. `Qwen/Qwen3.6-35B-A3B`, FreeToken's primary consumer-hardware MoE + benchmark model. +2. The current Gemma 4 MoE GGUF accepted by FreeToken's native Gemma loader. + +The project records correctness, stability, API behavior, GPU memory, host +memory, prefill throughput, decode throughput, TTFT, temperature, clocks, and +throttling. NVIDIA GPU tokens per second are context, not an AMD acceptance +threshold: LAN-223 uses a shared-memory APU rather than discrete VRAM and +PCIe. + +## What this branch changes + +The code is deliberately gated at the narrowest possible boundary so CUDA +behavior stays unchanged. + +- `setup.py` detects a ROCm PyTorch build and links the two native extensions + to `libamdhip64` instead of `libcudart`. +- `kernel/csrc/hip_compat.h` maps the small CUDA Runtime API subset used by + FreeToken's pinned-memory and CPU MoE extensions to HIP equivalents. +- CUDA JIT compilation removes NVCC-only flags on HIP and replaces CUDA-only + launch behavior with compatible HIP launch behavior. +- Triton paths avoid NVIDIA PTX inline assembly, Hopper Programmatic Dependent + Launch controls, and CUDA tile assumptions when PyTorch reports HIP. +- CUDA-only optional package probes are suppressed on HIP. The pure Triton + implementations remain the portable GPU fast path. +- NVIDIA SM feature gates reject ROCm before numerical capability comparison. + This matters because PyTorch presents HIP devices under `torch.cuda` for + compatibility, and `gfx1151` must never be interpreted as a new NVIDIA SM. + +## Clean LAN-223 installation + +Do not install into system Python, an existing llama.cpp environment, or the +existing vLLM environment. The reference layout is intentionally isolated: + +```text +/home/david/freetoken-amd/ + source/ this Git checkout + .venv/ Python 3.12, ROCm PyTorch, AMD Triton, FreeToken + artifacts/ commands, environment manifests, tests, logs, telemetry + models/ optional links to read-only local model storage +``` + +The exact PyTorch ROCm wheel must be selected after validating its compatible +Triton build on LAN-223. FreeToken's upstream CUDA package set must not be +installed on AMD: `flashinfer`, `sglang-kernel`, CUDA-indexed Torch wheels, and +the CUDA kernel-cache wheel are NVIDIA binaries. + +The initial build command is run from `source` only after the isolated Python +environment has a working HIP PyTorch import: + +```bash +python -m pip install -e . --no-build-isolation --no-deps +``` + +Use `hipcc --version`, `rocminfo`, and a small PyTorch HIP allocation before +the FreeToken build. Record outputs in `artifacts/environment/`, with secrets +and access tokens removed. + +## Persistent GGUF HIP JIT cache + +The native Gemma GGUF extension is compiled once per combination of FreeToken +source, PyTorch and HIP version, compiler flags, Python ABI, and GPU target. +`torch.utils.cpp_extension` reuses the resulting shared object on later +process starts. Normal serving must not delete that cache. + +The default cache is `$HOME/.cache/torch_extensions/`. For a deliberate, +portable installation-specific location, set this before every `ft serve` +launch and keep the directory across reboots and service restarts: + +```bash +export TORCH_EXTENSIONS_DIR=/home/david/freetoken-amd/cache/torch_extensions +mkdir -p "$TORCH_EXTENSIONS_DIR" +``` + +After an intentional FreeToken source or ROCm toolchain update, one rebuild is +expected. Deleting this directory is a recovery action only. It was cleared +during the original port investigation to force revised HIP sources to build; +that development step is not part of normal operation. + +## Required validation sequence + +1. Verify the host's `gfx1151` device, HIP runtime, PyTorch HIP build, and + AMD Triton version. +2. Build and import `_pinned_tensor` and `_cpu_moe` from the isolated + environment. +3. Run the ROCm gate unit tests plus the relevant CPU and Triton tests. +4. Run Qwen3.6-35B-A3B through `ft serve` on a non-conflicting local port. +5. Test `/v1/models`, non-streaming `/v1/chat/completions`, and streamed + `/v1/chat/completions` with fixed requests. +6. Run `ft bench bw` on LAN-223. Treat its recommendation as a measured + candidate, then verify it with full serving workloads. +7. Repeat the same API and stability checks for the supported Gemma 4 MoE + GGUF. +8. Save raw command output, service logs, request responses, profiler output, + and hardware telemetry under `artifacts/`. + +No llama-swap service, model configuration, or existing port is modified by +these commands. Service packaging happens only after the full validation set +passes. + +## Provenance + +This branch incorporates the focused current-main ROCm work from FreeToken +pull request #241, preserving its commits and authorship. It adds explicit +`gfx1151` safety coverage and project-specific validation documentation. +Upstream review should receive a focused pull request containing code plus +tests. LAN-223 environment reports and benchmark artifacts belong in this +fork unless the upstream maintainers request them. + +The completed 2026-08-28 native HIP validation, exact LAN-223 environment, +API evidence, command shapes, and known limitations are documented in +[`lan223-rocm-validation-2026-08-28.md`](lan223-rocm-validation-2026-08-28.md). diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md new file mode 100644 index 00000000..d5824a9d --- /dev/null +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -0,0 +1,1245 @@ +# LAN-223 native ROCm validation, 2026-08-28 + +## Result + +This validation passed the first release gate for the AMD port. FreeToken +served both required MoE models through the OpenAI-compatible API on LAN-223's +Radeon 8060S (`gfx1151`) using a native HIP and ROCm execution path. + +This is not a CPU fallback or a Vulkan result. The serving process uses the +ROCm PyTorch wheel, HIP-compiled native extensions, and Triton GPU kernels. +CUDA graphs were deliberately disabled for this validation because the MVP +needs correctness before graph capture tuning. + +## Reproducibility record + +| Item | Value | +| --- | --- | +| Host | LAN-223, `david-Gmktec-x2-2` | +| GPU | AMD Radeon 8060S Graphics, `gfx1151`, 40 CUs | +| System ROCm installation | ROCm 10.0 at `/opt/rocm-10.0` | +| PyTorch wheel | `2.13.0+rocm10.0.0` | +| HIP reported by PyTorch | `7.15.26333` | +| FreeToken branch | `amd-rocm-gfx1151` | +| Validation commit | `065d806` | +| API exposure | loopback-only ports, not llama-swap | + +The isolated validation layout was `/home/david/freetoken-amd/`; no existing +llama-swap service, model configuration, or production endpoint was changed. + +## Models and API evidence + +| Model | Source revision | Backend selection | Non-streaming result | Streaming result | +| --- | --- | --- | --- | --- | +| `nvidia/Qwen3.6-35B-A3B-NVFP4` | vendor model snapshot used for this run | Triton attention, MoE offload, native Triton NVFP4, serial expert load | HTTP 200, `AMD ROCm FreeToken ready.` in 1.54 s | HTTP 200, SSE chunks and `[DONE]` | +| `google/gemma-4-26B-A4B-it-qat-q4_0-gguf` | `d1c082be9cf3c8a514acf63b8761f4b41935842e` | Triton attention, MoE offload, serial expert load, HIP GGUF JIT | HTTP 200, `native hip api works` in 341.304 ms | HTTP 200, SSE chunks and `[DONE]` | + +Raw evidence remains on LAN-223 in these isolated artifact directories: + +```text +/home/david/freetoken-amd/artifacts/qwen36-nvfp4-serial-hip-prefill/ +/home/david/freetoken-amd/artifacts/gemma4-q4-rocm-thrust-system/ +``` + +The Gemma telemetry captured immediately after the API tests identified the +same `gfx1151` device, 33 percent GPU utilization, 46 percent allocated VRAM, +and a 40 C edge temperature. The model uses the APU's shared-memory design; +the tool's VRAM label is therefore only its standard telemetry label. + +## Warm single-request throughput + +The following measurements use one fixed 733-token prompt, greedy sampling, +and a one-sentence answer that produced 26 completion tokens. `TTFT` is the +client-observed time to the first non-empty SSE text chunk. Prompt throughput +is the end-to-end prompt-token count divided by TTFT, so it includes normal +API and scheduler overhead. Output throughput is completion tokens divided +by the interval from that first chunk through `[DONE]`. + +| Model | Prompt tokens | Completion tokens | TTFT | Prompt TPS | Generation interval | Output TPS | +| --- | ---: | ---: | ---: | ---: | ---: | ---: | +| Qwen3.6-35B-A3B NVFP4 | 733 | 26 | 4.976 s | 147.3 | 0.899 s | 28.9 | +| Gemma 4 26B A4B Q4_0 GGUF | 733 | 26 | 3.244 s | 226.0 | 0.581 s | 44.8 | + +These are warm, single-request measurements, not concurrency or maximum +throughput claims. The Qwen configuration uses the native Triton serial +NVFP4 prefill route selected for ROCm correctness. Its approximately +seven-minute cold initialization is expert-bank preparation and cache +allocation, not inference time. + +## Same-model llama.cpp Vulkan comparison + +To compare the usable Strix Halo serving baseline rather than an unrelated +model, the exact Gemma GGUF was served by llama.cpp Vulkan build `b10141` +(`0d47ea742`) on a separate loopback port. Both servers used one slot, +8,192-token context, greedy sampling, and the same repeated scheduler prompt. +The model SHA-256 was +`3eca3b8f6d7baf218a7dd6bba5fb59a56ee25fe2d567b6f5f589b4f697eca51d`. + +| Runtime | GPU backend | Prompt tokens | Completion tokens | TTFT | Client prompt TPS | Client output TPS | Runtime prompt TPS | Runtime output TPS | +| --- | --- | ---: | ---: | ---: | ---: | ---: | ---: | ---: | +| FreeToken | ROCm/HIP | 733 | 26 | 3.244 s | 226.0 | 44.8 | not exposed | not exposed | +| llama.cpp `b10141` | Vulkan | 758, 7 template tokens cached | 128 | 0.855 s | 886.7 | 63.2 | 1,078.4 | 61.7 | + +For this isolated, single-request Gemma workload, llama.cpp Vulkan reached +first output about 3.8 times sooner, delivered about 3.9 times the +client-observed prompt rate, and delivered about 1.4 times the client-observed +generation rate. llama.cpp's internal timing excludes ordinary API and +scheduler overhead, so its 1,078.4 prompt TPS and 61.7 output TPS must not be +compared directly with FreeToken's client-observed rates. + +The completion lengths differ because llama.cpp exposed Gemma's reasoning +stream and consumed the 128-token cap, whereas FreeToken's parser emitted the +final concise answer and stopped at 26 tokens. That makes the output-rate +comparison useful as a warm streaming rate, but not a quality or exact +end-to-end task comparison. The raw llama.cpp evidence is retained under +`/home/david/freetoken-amd/artifacts/llamacpp-vulkan-gemma4-q4-tps/` on +LAN-223. + +## Same-model ROCm 10 and HIP comparison + +The Vulkan baseline above answers a practical deployment question, but it is +not a backend-for-backend comparison. This follow-up rebuilt the same +llama.cpp source revision, `b10141` (`0d47ea742`), with HIP for `gfx1151` and +ran it under the same ROCm 10 installation used by FreeToken. The compiler +was ROCm 10 HIP `7.15.26333` with AMD Clang 23.0.0. At runtime, llama.cpp's +`libamdhip64`, `libhipblas`, `librocblas`, `libamd_comgr`, and HSA runtime +libraries all resolved from `/opt/rocm-10.0`, not the older ROCm installation. + +Both runners used the identical 14 GB Gemma 4 26B A4B Q4_0 GGUF, SHA-256 +`3eca3b8f6d7baf218a7dd6bba5fb59a56ee25fe2d567b6f5f589b4f697eca51d`, one +request at a time, an 8,192-token context, greedy sampling, `max_tokens: 128`, +and a 48-times repeated scheduler prompt. Each measurement used a distinct +nonce, preventing prompt-cache reuse. The token totals differ by one because +the two runners tokenize and render Gemma's chat template differently. + +| Runtime | HIP and ROCm stack | Prompt tokens | Completion tokens | TTFT | Client prompt TPS | Client output TPS | Runtime prompt TPS | Runtime output TPS | +| --- | --- | ---: | ---: | ---: | ---: | ---: | ---: | ---: | +| FreeToken, steady state | PyTorch `2.13.0+rocm10.0.0`, HIP `7.15.26333`, native HIP GGUF extension | 772 | 20 | 2.863 s | 269.6 | 46.1 | not exposed | not exposed | +| llama.cpp `b10141` | ROCm 10 HIP, `gfx1151` | 771 | 128 | 0.850 s | 906.6 | 58.3 | 1,011.6 | 56.2 | + +On this uncached, single-request workload, llama.cpp ROCm 10 reached first +text about 3.4 times sooner, supplied about 3.4 times the client-observed +prompt rate, and supplied about 1.3 times the client-observed output rate. +llama.cpp's internal numbers exclude HTTP, SSE, and scheduling overhead and +therefore are only comparable to another internal timing source, not directly +to FreeToken's client values. + +The FreeToken request that triggered a fresh GGUF HIP extension build is kept +as a separate cold-start measurement: 768 prompt tokens, 21 completion tokens, +109.938 s TTFT, 6.99 client prompt TPS, and 27.47 client output TPS. It +contains HIP compilation and must not be presented as inference throughput. +The subsequent steady-state run above was made after the extension completed, +using a fresh nonce and no prompt cache hit. FreeToken's extension compiler +was `/opt/rocm-10.0/bin/hipcc` targeting `gfx1151`, and its runtime libraries +came from the ROCm 10 PyTorch SDK packages. Its existing JIT command also +passed `/opt/rocm-7.2.4/include` as a supplemental include path. That does not +change the ROCm 10 compiler or loaded runtime libraries, but it prevents this +FreeToken build from being described as a strictly ROCm 10-only header build. + +The llama.cpp response used all 128 allowed tokens because it exposed Gemma +reasoning text. FreeToken stopped after a concise 20-token answer. This +makes the output-rate comparison a useful streaming measurement, but it is +not an exact answer-quality or equal-completion-length evaluation. + +Raw artifacts are retained only on LAN-223: + +```text +/home/david/freetoken-amd/artifacts/llamacpp-rocm10-gemma4-q4-tps/ +/home/david/freetoken-amd/artifacts/freetoken-rocm10-gemma4-q4-tps/ +``` + +## AMD TPS optimization campaign + +The first configuration optimization pass used the same warm AIME-25 problem +and a 128-token greedy completion for both FreeToken and the ROCm 10 HIP build +of llama.cpp `b10141`. Each runner received the identical user message, used +a warm identical request before the measured request, and ran one stream at a +time. Both rendered 63 prompt tokens; FreeToken's measured request reused 62 +prompt tokens and llama.cpp's reused 58. + +| Runtime and candidate | Decode TPS | TTFT | Result | +| --- | ---: | ---: | --- | +| FreeToken, offload, eager | 54.89 | 267.9 ms | Baseline | +| FreeToken, offload, HIP graph capture at batch size 1 | 55.73 | 259.3 ms | Best observed safe configuration | +| FreeToken, HIP graph plus experimental `-ffast-math` GGUF extension | 55.65 | 261.6 ms | Rejected: no gain, despite matching output hash | +| FreeToken, final target-specific `gfx1151` GGUF extension plus graph capture | 55.44 | 263.6 ms | Validated shipping configuration; normal run-to-run variation | +| FreeToken, experimental two-row Q4_0 MoE block | 55.30 | 291.6 ms | Rejected: slower with identical output hash | +| FreeToken, experimental Q4_0 MoE two-block residency hint | 55.08 | 294.9 ms | Rejected: slower with identical output hash | +| FreeToken, HIP Q4_0 MoE one-wave/two-row specialization | 55.89 median, 55.91 mean | 262.2 ms mean | Accepted: five independent API runs, identical output hash | +| FreeToken, full 4,096-slot expert cache and pinned 8,320-token KV pool | 60.11 median, 58.61 mean | 260.3 ms mean | Accepted configuration; four of five runs at 60.06 to 60.20 TPS, one host-contention outlier at 52.50 TPS | +| llama.cpp `b10141`, ROCm 10 HIP, earlier matched reference | 60.42 client, 58.88 internal | 128.6 ms | Historical reference | +| llama.cpp `b10141`, ROCm 10 HIP, current-host five-run control | 62.44 median, 62.13 mean client TPS | 111.5 ms mean | Same prompt, greedy decode, five fresh servers, requested 8,320-token context | + +The graph configuration removes approximately 1.5 percent of the eager decode +cost. The capacity-aware resident-expert configuration below then removes the +dominant configuration gap without changing the model, server API, or HIP +kernel arithmetic. Its uncontended median is within 0.51 percent of the +60.42 client-TPS llama.cpp reference, but its five-run arithmetic mean remains +below that reference because one run experienced external host stalls. The +criterion of meeting or exceeding llama.cpp is therefore not yet claimed as a +fully repeatable mean result. + +### Accepted full-expert-cache and fixed-KV configuration + +The original automatic offload configuration sized 3,840 GPU expert slots and +then assigned the remaining memory budget to a very large KV pool. That pool +is not required by the fixed 8,320-token operating target and lowered the +observed decode rate. A fixed expert-cache configuration leaves the same +native Q4_0 GGUF, HIP extension, graph-captured decode, OpenAI-compatible API, +and `offload` backend intact while making the capacity choices explicit: + +```bash +python benchmarks/bench_decode_moe.py \ + --model /home/david/freetoken-amd/models/Gemma-4-26B-A4B-it-qat-q4_0-gguf/gemma-4-26B_q4_0-it.gguf \ + --backend offload --cache 4096 --num-token-override 8320 \ + --mem-ratio 0.50 --decode 128 --greedy +``` + +`4096` is the complete 32-layer by 128-expert cache domain. A 3,840-slot +control preserved the fixed KV allocation but produced two severe decode-tail +events, confirming that leaving any of the 4,096 slots uncached can still +exercise the miss path. The explicit 4,096-slot configuration was therefore +retained. The new benchmark option maps `--num-token-override` to the public +server flag `--num-tokens`, so experiments can pin KV capacity without a +private wrapper. + +Five independent API runs used the fixed 63-token AIME request, 126 measured +decode steps, greedy sampling, `0.50` memory ratio, and the deterministic +output SHA-1 `abeee5e73e89`: + +| Run | Decode TPS | ms/token | TTFT | Event p50 / p99 | +| --- | ---: | ---: | ---: | --- | +| 1 | 60.063 | 16.649 | 259.7 ms | 16.929 / 17.841 ms | +| 2 | 52.499 | 19.048 | 261.4 ms | 16.928 / 120.541 ms | +| 3 | 60.203 | 16.610 | 259.6 ms | 16.855 / 17.657 ms | +| 4 | 60.114 | 16.635 | 260.3 ms | 16.937 / 17.619 ms | +| 5 | 60.183 | 16.616 | 260.6 ms | 16.876 / 17.522 ms | +| Aggregate | **58.613 mean, 60.114 median** | 17.112 mean | 260.3 ms mean | 16.928 / 17.657 ms median | + +The four normal runs are within 60.063 to 60.203 TPS and have p99 latency at +or below 17.841 ms. The one low-throughput run kept the same output, VRAM, +TTFT, and p50 latency, but had isolated 120.541 ms decode events. Kernel logs +recorded `kfd_process_wq_release` holding CPU for more than 10 ms and the host +showed full I/O pressure. Read-only inspection also found two long-running, +blocked user-owned filesystem scans. They were not stopped by this campaign. +This is host contention evidence, not a FreeToken numerical or API failure. + +Capacity was tested through the public OpenAI-compatible API, not merely at +startup. A request with 7,619 prompt tokens plus one completion token ran +inside the pinned 8,320-token pool, returned exactly `OK`, and completed in +22.352 seconds. The server then exited cleanly with no KFD processes. + +The retained raw evidence is: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/full-expert-cache-4096-20260829T010740Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/fixed-expert-cache-3840-control-20260829T011537Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/full-cache-4096-context8320-20260829T012342Z/ +``` + +### Current-host ROCm llama.cpp control + +The historical llama.cpp reference was useful for identifying the original +gap, but it was not collected alongside the accepted 4,096-slot FreeToken +configuration. A new five-run control was therefore run immediately after +that configuration investigation, without changing LAN-223, stopping any +user process, or enabling a production service. Each trial launched a fresh +`llama-server` from the ROCm 10 `b10141` build with all layers on `gfx1151`, +Flash Attention enabled, one parallel slot, and `-c 8320`. The server reports +an 8,448-token slot after its own request reserve is added. This is a llama.cpp +internal allocation detail; the requested application context target was +8,320 tokens in both runners. + +Both runners used the same cached AIME-25 problem 0, a warmed streamed +OpenAI-compatible `/v1/chat/completions` request, greedy sampling, and a +128-token completion. The metric in this table is client-observed decode +throughput: `(completion_tokens - 1)` divided by elapsed time from the first +to last SSE token event. It includes HTTP and SSE delivery for both runners. + +| Runtime | Five client decode TPS | Mean | Median | Mean TTFT | p99 event gap median | +| --- | --- | ---: | ---: | ---: | ---: | +| FreeToken, 4,096 experts, 8,320-token KV pool | 60.06, 52.50, 60.20, 60.11, 60.18 | 58.61 | 60.11 | 260.3 ms | 17.66 ms, excluding the host-stalled run 120.54 ms | +| llama.cpp `b10141`, ROCm 10 HIP | 61.04, 62.03, 62.44, 62.57, 62.56 | 62.13 | 62.44 | 111.5 ms | 16.63 ms | + +llama.cpp leads FreeToken by 3.7 percent on median client decode TPS +(`62.44 / 60.11 - 1`) and 5.7 percent on the unfiltered five-run mean +(`62.13 / 58.61 - 1`). It also has lower warm TTFT. FreeToken produced the +same deterministic output hash in every measured run; llama.cpp produced the +same deterministic output hash in every one of its own runs. The hashes are +not compared across runtimes because their tokenizers and chat-template +implementations differ. + +This is a close result for decode rate, but it does **not** meet the stated +criterion of meeting or exceeding llama.cpp. The remaining performance work +is therefore directed at the HIP decode path and the source of the FreeToken +tail stall, rather than a claim of parity. The raw llama.cpp evidence is +retained on LAN-223 at: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/llamacpp-current-host-context8320-20260829T013730Z/ +``` + +### Current-upstream rebase and full API revalidation + +After the comparison, upstream `main` advanced from `9ef3651` to `a05c265` +with Qwen 3.8 support and engine or cache changes. The AMD branch was rebased +onto that current upstream revision without a conflict, rather than leaving a +performance result attached to an obsolete upstream base. The rebased branch +was then installed into the isolated LAN-223 virtual environment so its native +HIP pinned-memory extension was built from the rebased source. The source +checkout used for that validation was deliberately separate from the earlier +test checkout, preventing an uncommitted working-tree change from becoming +test evidence. + +The first complete Gemma launch from the rebased checkout rebuilt the target- +specific GGUF HIP extension and matching graph helper because the source path +is part of their cache identity. The build used ROCm 10 `hipcc`, `-O3`, and +`--offload-arch=gfx1151`; a later process restart can reuse that cache. The +server then completed its normal graph capture, exposed `/v1/models`, and +served two streamed OpenAI-compatible chat completions before clean shutdown. + +| Check | Observed value | +| --- | --- | +| Upstream revision in branch history | `a05c265` | +| HIP build and ROCm-runtime tests | 4 passed | +| MoE configuration | `offload`, 4,096 expert slots, 8,320 KV tokens, graph batch size 1 | +| Warm streamed API decode | 60.07 client TPS, 16.648 ms/token | +| Warm TTFT | 258.2 ms | +| Prompt and completion tokens | 63 and 127, respectively | +| Output SHA-1 | `abeee5e73e89` | +| Server VRAM | 15.66 GiB | +| Post-run process state | Server shut down; no serving process remained | + +The response ended at 127 tokens despite the requested 128-token limit, so +the benchmark harness emitted its explicit token-count warning. The request +was otherwise successful, deterministic, and had the expected response hash. +This one-run revalidation is evidence that rebasing did not break native HIP +serving. It is intentionally not folded into the five-run performance score. +Its raw logs and result are retained at: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/rebased-current-main-api-retry-20260829T014741Z/ +``` + +### Rebased Qwen3.6 NVFP4 API revalidation + +The other MoE model retained in the isolated FreeToken inventory is +`Qwen3.6-35B-A3B-NVFP4`. It was revalidated after the upstream rebase through +the same loopback OpenAI-compatible streaming API, using the native Triton +NVFP4 expert path, graph batch size 1, automatic expert-cache sizing, a 0.35 +memory ratio, and greedy 128-token AIME decoding. This exercised its complete +21.8 GiB parallel expert-bank load, cache allocation, graph capture, warm +request, measured request, and cleanup. + +| Check | Observed value | +| --- | --- | +| Resolved expert cache | 9,499 slots and 8,255 KV tokens | +| Warm streamed API decode | 28.93 client TPS, 34.560 ms/token | +| Warm TTFT | 404.7 ms | +| Prompt and completion tokens | 54 and 127, respectively | +| Output SHA-1 | `0acef4eab6f4` | +| Server VRAM | 19.12 GiB | +| Post-run process state | Server shut down; no serving process remained | + +As with the rebased Gemma validation, the response ended at 127 tokens and the +harness recorded its explicit limit-warning rather than silently treating it as +a 128-token result. The API transaction and deterministic response succeeded. +The server also reported that `triton_kernels` was absent and selected the +numerically equivalent pure-PyTorch router fallback. That is a documented +performance limitation, not a functional failure. A native ROCm-compatible +fused-router installation must be independently verified before it can be +considered an optimization. + +The raw evidence is retained at: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/rebased-current-main-qwen36-api-20260829T015240Z/ +``` + +### Rejected ROCm vendored-Triton router candidate + +The Qwen revalidation exposed a pure-PyTorch router fallback because OpenAI's +`triton_kernels` package contains CUDA-only binaries. Current upstream also +contains an in-tree Triton router, so it was evaluated as a ROCm-only candidate +before any production use. On the actual Radeon 8060S it selected the same +expert set as PyTorch; BF16 equal-logit ties can have a different internal +ordering, while FP32 indices matched exactly. Its selected routing weights +matched PyTorch within `2.98e-8`, and the isolated one-token, 128-expert, +top-8 router time improved from 21.14 us to 14.96 us. + +That microbenchmark improvement was insufficient. The complete Qwen API run +with the candidate reached 30.26 client TPS, but its deterministic greedy +response SHA-1 was `cd580f4978fb`, not the reference `0acef4eab6f4`. Small +router differences therefore accumulated into a different generated response. +The candidate was reverted and ROCm continues to use the reference PyTorch +router. This keeps quality behavior stable even though the fused alternative +is faster in isolation. The fallback warning now explicitly distinguishes +intentional ROCm behavior from a missing CUDA Linux package. + +The rejected candidate evidence is retained at: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/qwen36-vendored-router-api-20260829T015937Z/ +``` + +The restored branch was then revalidated through the full Qwen API path. It +returned to the reference SHA-1 `0acef4eab6f4` at 28.96 client TPS, with the +same 19.12 GiB VRAM use and clean shutdown. Its 1,272.2 ms warm TTFT is not a +performance regression claim: the model's 21.8 GiB expert-bank load was +concurrently slowed by the documented host I/O pressure, taking 3 minutes and +37 seconds instead of about 2 minutes. The final exact-path artifact is: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/qwen36-router-revert-api-20260829T020626Z/ +``` + +### Accepted HIP Q4_0 one-wave/two-row MoE specialization + +The first two-row experiment did not reproduce llama.cpp's execution shape: it +used two independent 32-thread waves. Commit `d1de602` instead adds a ROCm-only +Q4_0 kernel in which one 32-thread wave accumulates two adjacent output rows. +It preserves FreeToken's flattened token/top-k route IDs, packed expert-bank +layout, Q8_1 activation layout, and BF16 public output contract. CUDA retains +the established generic path. + +The dedicated LAN-223 microbenchmark uses the verified Gemma 4 26B A4B Q4_0 +geometry: 128 experts, top-k 8, hidden width 2816, intermediate width 704, and +one decode token. Five runs with 2,000 timed calls each measured a 73.509 us +baseline median for the gate/up plus down pair and a 64.340 us candidate median, +a 12.5 percent kernel-pair reduction. ROCprof recorded a 32-thread wave, zero +LDS and scratch allocation, and half the former row-block grid. The compiler +still allocated 48 VGPRs, so future work must target register pressure +separately rather than claiming it was resolved by this change. + +The end-to-end gate was five independent loopback OpenAI-compatible API server +runs, each using the exact Gemma GGUF SHA-256, offload backend, 0.50 memory +ratio, HIP graph capture, greedy AIME-25 problem 0, and 128-token decode +procedure. All five emitted the original deterministic output SHA-1 +`abeee5e73e89` and retained 27.52 GiB server-reported VRAM use. + +| Metric | Five-run result | +| --- | --- | +| Decode TPS | 55.713 to 56.071 | +| Decode TPS median / mean | **55.894 / 55.905** | +| Decode ms/token median / mean | **17.891 / 17.887** | +| TTFT mean | 262.2 ms | +| Output SHA-1 | `abeee5e73e89` in every run | +| Compared shipping configuration | 55.44 TPS single verified run | +| Matched llama.cpp ROCm 10 reference | 60.42 client TPS | + +The candidate is accepted because it produces a repeatable FreeToken gain of +approximately 0.8 percent over the prior shipping result while preserving the +observable API result. It remains approximately 7.5 percent below the +matched llama.cpp client-TPS reference, so it is an incremental port +improvement rather than completion of the performance objective. + +Artifacts are retained on LAN-223: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/q4-moe-microbench-20260828T231332Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/q4-moe-two-row-wave-20260828T231950Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/q4-moe-two-row-wave-20260828T231950Z/api-repeats-20260828T232646Z/ +``` + +### Rejected two-row MoE Q8 activation-reuse candidate + +The accepted HIP Q4_0 one-wave/two-row MoE kernel computes two adjacent output +rows from the same Q8_1 activation block. The generic dot helper loads the +four packed Q8 activation words independently for each row. Commit `684148d` +tested a ROCm-only helper that loads those four words once and supplies them to +both row dot products, while retaining the same Q4 nibble order, DP4A order, +scales, BF16 public-output contract, and all CUDA code. + +The shape-accurate Gemma routed-expert microbenchmark improved from about +64.34 us to **61.54 us per gate/up plus down pair**. That local result did not +translate to a material full-server result. Five independent OpenAI-compatible +API runs, each using the fixed 63-token prompt and 126 measured decode steps, +all returned greedy output SHA-1 `abeee5e73e89`: + +| Run | Decode TPS | ms/token | TTFT | Event p50 / p99 | +| --- | ---: | ---: | ---: | --- | +| 1 | 55.993 | 17.859 | 264.2 ms | 18.133 / 18.753 ms | +| 2 | 55.970 | 17.867 | 262.1 ms | 18.171 / 18.853 ms | +| 3 | 55.958 | 17.871 | 259.9 ms | 18.183 / 18.823 ms | +| 4 | 56.012 | 17.853 | 260.8 ms | 18.086 / 18.990 ms | +| 5 | 56.155 | 17.808 | 262.9 ms | 18.018 / 18.853 ms | +| Aggregate | **56.018 mean, 55.993 median, 0.080 stddev** | 17.851 mean | 262.0 ms mean | 18.133 / 18.853 ms median | + +This is only 0.20 percent above the accepted 55.905 TPS mean, materially below +the campaign's repeatable-improvement threshold and far below the 60.42 client +TPS matched llama.cpp ROCm 10 reference. The candidate was therefore reverted +in `a237b12`; the accepted one-wave/two-row implementation remains active. +The benchmark sequence also ended with no KFD GPU processes, confirming that +the service was torn down cleanly. + +The retained raw evidence is: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/moe-q8-reuse-20260829T005332Z/microbench.json +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/moe-q8-reuse-20260829T005332Z/api-first.jsonl +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/moe-q8-reuse-20260829T005332Z/api-repeats.jsonl +``` + +### Rejected dense Q4_0 one-wave/two-row specialization + +The dense Q4_0 vector path uses the same older one-row scheduling structure as +the routed-expert path. Commit `b4a53d1` applied the accepted one-wave/two-row +pattern to that dense kernel, while leaving CUDA unchanged. A new shape-aware +microbenchmark covered the exact Gemma projection dimensions recovered from the +GGUF: 2816x4096, 8192x2816, 4224x2816, and 10240x2816. In isolation it reduced +the measured GPU event time for every shape, including 10240x2816 from 41.34 us +to 28.12 us. + +That synthetic gain did not survive the real graph-captured serving path. The +fixed loopback API workload compiled the candidate from a fresh HIP extension +cache, returned the exact deterministic output SHA-1 `abeee5e73e89`, and used +the same 27.52 GiB of server-reported VRAM, but measured only **54.71 TPS** or +18.277 ms/token. This is below the 55.89 TPS accepted MoE-specialization +median and below the prior 55.04 TPS repaired baseline. The dense candidate +was therefore reverted. It proves that isolated event timing alone is not an +acceptance metric for graph-captured end-to-end decode. + +The retained raw evidence is: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-q4-microbench-20260828T233506Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-q4-two-row-wave-20260828T233930Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-two-row-wave-api-20260828T234147Z/ +``` + +### Rejected dense Q4_0 FP32-output hypothesis + +llama.cpp's corresponding vector kernel stores FP32 values, whereas the +FreeToken GGUF adapter normally returns the input dtype, BF16 for this Gemma +run. That difference was a plausible explanation for the profiler contrast: +FreeToken's generic Q4_0 dense kernel reported 48 architectural VGPRs and the +llama.cpp reference reported 24. Commit `e77a44b` added a deliberately +benchmark-only Q4_0 flag that changed only the destination tensor to FP32. +It was never wired to the GGUF model layers, and a source guard ensured the +normal serving call retained its BF16 contract. + +The result rejects that explanation. In the first independent event run, the +four exact Gemma projection geometries measured 18.28 us, 33.59 us, 18.16 us, +and 35.98 us respectively. The profiler trace showed the FP32 specialization +still at **48 VGPRs**, 128 SGPRs, no LDS, and no scratch. It therefore did not +match llama.cpp's 24-VGPR code shape. Its profile-run event values also showed +no consistent gain. The experiment was reverted in `203062f`; no public or +model-serving API changed. + +The retained raw evidence is: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-q4-fp32-output-20260828T234953Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-q4-fp32-output-rocprof-20260828T235258Z/ +``` + +### Rejected dense HIP launch-bound candidate + +Commit `c7009a9` tested the other conspicuous structural difference from the +matched llama.cpp Q4_0 vector kernel: a HIP-only `__launch_bounds__(32, 1)` +constraint for FreeToken's one-wave dense GEMV. CUDA was unchanged. The +candidate compiled cleanly for `gfx1151` and kept the same 48 VGPRs, 128 +SGPRs, zero LDS, and zero scratch as the generic FreeToken kernel. It did +improve three of the four isolated Gemma projection measurements, but did not +reduce the compiler resource gap against llama.cpp. + +Five independent graph-captured loopback API runs produced **55.90 TPS mean** +and **55.88 TPS median**, with deterministic output SHA-1 `abeee5e73e89` in +every run. The accepted MoE-only specialization measured 55.91 TPS mean and +55.89 TPS median under the same workload. The candidate therefore has no +meaningful decode gain and its 274.7 ms mean TTFT was worse than the accepted +candidate's 262.2 ms mean. It was reverted in `1d555e3`; the upstream-ready +path remains unchanged by this experiment. + +The retained raw evidence is: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-q4-launch-bounds-20260828T235459Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-q4-launch-bounds-rocprof-20260828T235726Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-launch-bounds-api-20260828T235756Z/ +``` + +### Rejected indexed Q4_0 dense HIP kernel + +The dominant llama.cpp Q4_0 trace was rechecked before this experiment. Its +main kernel uses the same 32-thread by 1-row workgroup as FreeToken, but +reports 24 VGPRs versus FreeToken's 48. Commit `a5e04d1` isolated the remaining +source-level difference: a Q4_0-only HIP kernel that keeps the base weight +pointer and block index separate until the vector-dot helper. It preserved +the generic FreeToken launch geometry and left CUDA and every non-Q4_0 type +unchanged. + +The isolated evidence was favorable but insufficient. The four exact Gemma +dense projections measured 16.49 us, 28.69 us, 17.83 us, and 35.52 us, and the +profile trace measured 18.16 us, 22.88 us, 13.17 us, and 29.22 us. The compiler +still used 48 VGPRs, 128 SGPRs, no LDS, and no scratch. The first full +graph-captured API run preserved the deterministic output SHA-1 +`abeee5e73e89`, but collapsed to **15.21 TPS**, 65.73 ms/token, 3014.5 ms TTFT, +and 1674.5 ms p99 event latency. This is a functional result but a clear +performance failure. It was reverted in `b281e0e` and must not be retried +without an explanation for the end-to-end stalls. + +The retained raw evidence is: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-q4-indexed-pointer-20260829T000901Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-q4-indexed-pointer-rocprof-20260829T001126Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-indexed-pointer-api-20260829T001156Z/ +``` + +### Rejected scalarized dense Q4_0 dot-product candidate + +The source and trace audit found that the historical FreeToken dense Q4_0 +helper materializes two packed Q4 words and four Q8 words in short local +arrays before issuing four DP4A operations. llama.cpp's newer HIP path does +not share FreeToken's old wrapper structure, so a HIP-only candidate replaced +only that dense helper with named scalar values. It retained the original +packed GGUF layout, four DP4A operations in the same order, scale formula, and +BF16 output contract. CUDA and the separately accepted routed-MoE kernel were +unchanged. + +The candidate compiled for `gfx1151`, passed the targeted HIP build and +attention tests, and produced finite results for all four exact Gemma dense +projection shapes. Its isolated event times were 17.38 us for 2816x4096, +28.38 us for 8192x2816, 19.20 us for 4224x2816, and 35.84 us for 10240x2816. +That showed useful synthetic movement, especially for the second shape, but +was not enough to accept it. + +Five independent graph-captured API runs all returned the deterministic +SHA-1 `abeee5e73e89`, retained 27.52 GiB server-reported VRAM, and left no KFD +process after shutdown. Their TPS range was 55.812 to 56.155, with **55.946 +TPS mean** and **55.919 TPS median**. Those figures differ from the accepted +Q4_0 MoE baseline by only 0.041 TPS mean and 0.025 TPS median, while mean TTFT +increased from 262.2 ms to 269.7 ms. This is normal run-to-run noise, not a +repeatable end-to-end improvement, so it was reverted in `d9ce2c5`. + +The retained raw evidence is: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-q4-scalarized-20260829T003720Z/microbench.json +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-q4-scalarized-20260829T003720Z/microbench-rocprof.json +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-q4-scalarized-20260829T003720Z/api-first.jsonl +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-q4-scalarized-20260829T003720Z/api-repeats.jsonl +``` + +### Rejected Q4_0 MoE route-grouping candidate + +The source comparison showed that llama.cpp places the eight routed experts in +separate waves of one multi-wave MoE workgroup, while FreeToken's accepted HIP +specialization uses one workgroup per route. Commit `dc73e8e` tested that +topology directly: a HIP-only Q4_0 kernel with eight independent 32-lane route +waves in a 256-thread workgroup, retaining the accepted two-output-row +arithmetic inside each wave. CUDA and all non-Q4_0 formats remained unchanged. + +The candidate compiled for `gfx1151` and passed the targeted HIP build tests, +but failed the shape-accurate microbenchmark gate. For Gemma's eight-route +decode geometry it measured 34.93 us gate/up plus 30.69 us down, or **65.62 us +per pair**, versus the accepted two-row kernel's 64.25 us mean pair time. Since +the grouped workgroup was slower before the API workload, no server benchmark +was run. It was reverted in `96c51f9` and the accepted one-wave/two-row MoE +kernel remains active. + +The retained raw evidence is: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/q4-moe-route-group8-20260829T001802Z/ +``` + +### Rejected Triton GQA attention eight-warp candidate + +Gemma's sliding decode attention has 16 query heads, 8 KV heads, a 256-wide +head dimension, and a 1,024-token sliding window. The HIP production path +uses a 16-head padded tile, 32-token KV blocks, and four Triton warps. A +shape-accurate benchmark tested head tiles of 2, 4, and 8, a 64-token KV +block, and two or eight warps. All tile and 64-token-block alternatives were +slower. Eight warps was faster in isolation: 39.51 us versus 41.87 us for +sliding attention, and 49.84 us versus 93.17 us for Gemma's 2-KV-head, +512-wide full-attention geometry. + +That microbenchmark win did not survive the full serving workload. An +otherwise identical graph-captured loopback OpenAI-compatible API run with +eight warps returned the expected deterministic SHA-1 `abeee5e73e89`, but +measured **53.76 TPS**, 18.600 ms/token, 281.2 ms TTFT, and 20.227 ms p99 +event latency. This is below the accepted five-run 55.91 TPS mean. The +production override was removed, so normal HIP serving remains at four warps; +the benchmark-only probe parameters remain available for future controlled +research. This result is a second independent example of why isolated GPU +event timings cannot be used as a serving-performance acceptance criterion. + +The retained raw evidence is: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/gqa-attention-blockh2-20260829T002315Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/gqa-attention-blockh4-8-20260829T002334Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/gqa-attention-blockn64-20260829T002442Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/gqa-attention-warps2-8-20260829T002459Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/gqa-attention-global-warps8-20260829T002558Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/attention-warps8-api-20260829T002618Z/ +``` + +The best verified FreeToken command shape is: + +```bash +export ROCM_PATH=/opt/rocm-10.0 +export HIP_PATH=/opt/rocm-10.0 +export TORCH_EXTENSIONS_DIR=/home/david/freetoken-amd/cache/torch_extensions + +ft serve --model-path /home/david/freetoken-amd/models/Gemma-4-26B-A4B-it-qat-q4_0-gguf/gemma-4-26B_q4_0-it.gguf \ + --attention-backend triton --moe-backend offload --moe-cache-size 4096 \ + --num-tokens 8320 --memory-ratio 0.50 --max-running-requests 1 \ + --max-seq-len-override 8320 --cuda-graph-max-bs 1 +``` + +The port now derives and exports `PYTORCH_ROCM_ARCH=gfx1151` before the GGUF +extension is compiled when the operator did not set an explicit architecture. +This avoids compiling for unnecessary visible targets and makes the extension +cache target-specific. It does not itself increase steady-state TPS because +the original HIP build already selected `gfx1151` on this single-GPU host. + +The remaining gap is not an untested cache or residency setting: Gemma's GGUF +adapter only supports the native Q4_0 offload implementation, and the accepted +configuration keeps all 4,096 routed-expert slots resident while retaining a +verified 8,320-token KV pool. Closing the gap requires a profile-guided +improvement to the HIP GGUF decode kernels or another proven ROCm attention or +quantized-linear implementation. + +The initial direct `rocprofv3` attempts did fail because the host profiler +injected a second LLVM and rocprofiler SDK beside the SDK bundled with the +PyTorch ROCm wheel. That historical failure is retained in +`rocprof-gfx1151*/` and `rocprof-launch-gfx1151-v2/` under the raw artifact +directory. It was subsequently repaired by +[`scripts/lan223-rocprof-wheel-sdk.sh`](../scripts/lan223-rocprof-wheel-sdk.sh), +which directs the host profiler front end to the wheel's matching SDK. The +repaired launch produced FreeToken kernel traces, including the active +`moe_vec_q4_0_hip_two_rows` kernel. Traces are diagnostic evidence only and +are never used as TPS scoring because profiling changes execution timing. + +### Current-source trace and rejected RDNA4 dense Q4_0 eight-wave candidate + +After an unprofiled final-source warm run rebuilt the path-specific native HIP +extension, the complete loopback API workload returned the established greedy +Gemma SHA-1 `abeee5e73e89` at **60.16 TPS**, 16.623 ms per token, 259.0 ms +TTFT, 16.877 ms p50 event latency, and 17.931 ms p99 event latency. It kept +the full 4,096-slot expert cache and 8,320-token KV budget. This is the +current unprofiled checkpoint for the accepted source path. + +The repaired profiler wrapper then traced that already-built final source. +The trace also preserved the output SHA-1, but measured 41.41 TPS and a 216.2 +ms p99 because tracing changes dispatch timing. It is not a performance +result. Its kernel statistics do identify the next work order: routed Q4_0 +MoE vector work consumed 31.05 percent of GPU kernel time, dense Q4_0 vector +work 28.18 percent, and dense Q6_K vector work 15.64 percent. The active +MoE kernel name was `moe_vec_q4_0_hip_two_rows`, proving that the trace covers +the accepted HIP specialization rather than the earlier generic path. + +Current llama.cpp source uses an RDNA4-specific eight-wave policy for simple +one-vector Q4_0 matvecs. Candidate commit `1bf9489` applied that scheduling +policy only to FreeToken's dense HIP Q4_0 launcher. It deliberately retained +the generic dot product, Q4_0 and Q8_1 packing, BF16 result contract, CUDA +path, all non-Q4_0 types, and the separately accepted routed-MoE kernel. +This made the candidate distinct from the already rejected dense two-row and +launch-bound experiments. + +The four shape-accurate dense microbenchmarks were mixed when rerun with 10 +warmups and 100 repetitions: the candidate improved 8,192 by 2,816 from +28.29 to 27.49 microseconds and 4,224 by 2,816 from 26.89 to 17.26 +microseconds, but regressed 2,816 by 4,096 from 21.30 to 23.91 microseconds +and 10,240 by 2,816 from 33.04 to 34.18 microseconds. Because Gemma uses all +four projections, this was insufficient to accept the launch policy. + +The full graph-captured API result confirmed rejection. It returned the +exact established SHA-1, but reached only **59.33 TPS**, 16.856 ms per token, +290.7 ms TTFT, and 18.011 ms p99 event latency. This is below the current +60.16 TPS final-source checkpoint and below the established accepted five-run +60.11 TPS median. The candidate remains on its separate branch and is not +part of the upstream-review branch. Its isolated worktree initially lacked +the unchanged native pinned-memory extension; the test setup copied the +validated extension only after SHA-256 and byte-for-byte equality checks. +That repair affected no source logic and the resulting API run is the only +performance outcome used for this decision. + +The retained raw evidence is: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/gemma-final-path-warm-20260829T021243Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/gemma-final-current-kernel-trace-20260829T021738Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-q4-current-baseline-micro-20260829T022723Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-q4-rdna4-eightwaves-micro-20260829T022528Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-q4-rdna4-eightwaves-api-repaired-20260829T023038Z/ +``` + +### Rejected RDNA4 dense Q6_K eight-wave candidate + +The current final-source trace showed dense Q6_K vector work at 15.64 percent +of total traced GPU kernel time. Gemma uses Q6_K for its tied token embedding +and LM head, and the matching current llama.cpp RDNA4 policy selects eight +waves for one-vector Q6_K matvec. Candidate commit `ebf3f06` therefore +changed only FreeToken's HIP dense Q6_K wrapper to launch the established +generic Q6_K dot-product kernel with eight independent row waves per block. +It retained Q6_K and Q8_1 packing, reduction arithmetic, the BF16 result +contract, CUDA behavior, Q4_0 dense behavior, and all routed-MoE behavior. + +The full graph-captured loopback API workload returned the exact established +Gemma SHA-1 `abeee5e73e89`, used 15.66 GiB VRAM, and reported 259.6 ms TTFT +with a 17.663 ms p99 event latency. Its decode result was **59.98 TPS** or +16.672 ms per token. That is close to, but below, the 60.16 TPS accepted +final-source checkpoint. A one-run result without a TPS improvement does not +justify a second specialized scheduling path, so the candidate remains on its +separate experiment branch and is not part of the upstream-review branch. + +The test used the unchanged native pinned-memory extension after SHA-256 and +byte-for-byte equality checks against the validated final source. The raw +evidence is retained at: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-q6-rdna4-eightwaves-api-20260829T023720Z/ +``` + +### Current host-interference qualifier + +A read-only LAN-223 health capture at 2026-08-29T02:41:15Z found no GPU reset, +thermal problem, or active FreeToken server. The Radeon 8060S was idle at +30 C after the test. It did, however, identify two pre-existing user-owned +filesystem scans in uninterruptible `D` state: one scanning `/home/david`, +`/mnt`, and `/data` for large GGUF or SafeTensors files, and one scanning +`/home/david` and `/media/david` for Gemma GGUF files. At capture time they +had been alive for approximately 8.8 and 6.1 hours respectively. + +The same capture reported I/O full-pressure at 0.61 percent over ten seconds +and retained kernel warnings that `kfd_process_wq_release` and +`svm_range_deferred_list_work` had exceeded their CPU workqueue budget. These +facts do not prove that a particular FreeToken result is invalid, but they +provide a concrete explanation for occasional multi-millisecond dispatch +outliers and the isolated 52.50 TPS baseline run. They can affect both +FreeToken and llama.cpp under a matched test. + +No process priority, service state, kernel option, ROCm installation, or +hardware component was changed by this investigation. Any decision to stop +or otherwise alter the two user-owned scans requires explicit operator +authorization. Until then, accepted performance claims remain based on +multiple clean launches and retain raw tail-latency data rather than hiding +the interference. + +### Current review-branch static validation + +The current upstream-review commit `6c6198b10d9fb6a9c93e0aa94a05ac4144ec061d` +was validated directly on LAN-223 after the I/O evidence capture tooling was +added. The check completed without starting an inference server or changing +host state: + +```text +python -m compileall -q python benchmarks passed +pytest -q tests/kernels/test_gguf_hip_build_flags.py \ + tests/utils/test_rocm_runtime.py 4 passed +``` + +The raw output and commit metadata are retained at: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/current-review-static-validation-20260829T024456Z/ +``` + +A temporary high-performance DPM governor test could not be run because the +non-root LAN-223 account cannot write `power_dpm_force_performance_level`; +automatic mode was unchanged. + +Raw campaign artifacts are retained on LAN-223: + +```text +/home/david/freetoken-amd/artifacts/amd-optimization-2026-08-28/ +``` + +## Deep-investigation baseline and profiler repair + +The reproducible read-only baseline is captured by +[`../scripts/lan223-capture-baseline.sh`](../scripts/lan223-capture-baseline.sh). +The first baseline was written to: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/baseline-20260828T220753Z/ +``` + +### Test-checkout repair and revalidated shipping baseline + +During the follow-on investigation, the isolated LAN-223 source checkout was +found at `61a1505`. That commit contained the subsequently rejected +two-block-residency Q4_0 MoE experiment. The authoritative branch had already +reverted that experiment at `b77825d` and documented the rejection at +`222cbd3`. Using the stale checkout for another benchmark would have made the +result impossible to attribute to the branch under review. + +The checkout was clean, so it was repaired with a fast-forward only update to +`origin/amd-rocm-gfx1151`, reaching `222cbd3`. No production process, +llama-swap configuration, or other LAN host was touched. The next run used a +new, dated `TORCH_EXTENSIONS_DIR`, forcing a fresh native HIP binary rather +than reusing the binary compiled from the stale source. + +| Item | Revalidated value | +| --- | --- | +| Artifact directory | `/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/repaired-baseline-20260828T224642Z/` | +| Source commit | `222cbd3` | +| Model SHA-256 | `3eca3b8f6d7baf218a7dd6bba5fb59a56ee25fe2d567b6f5f589b4f697eca51d` | +| Extension build | Fresh ROCm 10 `hipcc`, `--offload-arch=gfx1151`, `-O3` | +| API and workload | Loopback FreeToken API, greedy AIME-25 problem 0, one warm and one measured request | +| Measured completion | 127 tokens, 126 decode intervals | +| Client decode throughput | **55.04 TPS** or **18.169 ms/token** | +| TTFT | 295.2 ms | +| Event p50 / p99 | 18.454 ms / 19.509 ms | +| Output SHA-1 | `abeee5e73e89`, identical to the earlier shipping-configuration run | +| Post-run ROCm process check | No KFD PIDs | + +The single revalidation is consistent with the existing 55.44 TPS shipping +baseline and remains below the 60.42 TPS matched llama.cpp reference. It is a +provenance repair, not a new performance claim and not a substitute for the +planned repeated candidate measurements. + +### Rejected Q4_0 aligned-load candidate + +The matched traces showed FreeToken's Q4_0 vector kernels using 48 VGPRs per +thread, while llama.cpp's corresponding generic Q4 vector kernel reported 24 +VGPRs. Both used a 32-thread workgroup with zero LDS and scratch allocation. +As a narrow, low-risk test, commit `b8de163` replaced only the Q4_0 packed-load +helper expressions with the aligned `get_int_b2` and `get_int_b4` expressions +used by the current llama.cpp HIP source. The dot-product arithmetic, output +type, data layout, model, workload, and launch geometry were otherwise +unchanged. + +The target-host HIP build-configuration tests passed, the extension rebuilt +for `gfx1151`, and the output SHA-1 remained `abeee5e73e89`. However, the +candidate measured 54.99 TPS or 18.185 ms/token, versus 55.04 TPS or 18.169 +ms/token for the immediately preceding repaired baseline. That difference is +well inside normal run variation and does not improve the runner. The +candidate was therefore reverted by `c1899a0`; it is not part of the shipping +configuration. + +The raw candidate evidence is retained at: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/q4-load-alignment-20260828T225237Z/ +``` + +This eliminates aligned helper spelling as the explanation for the measured +register and throughput gap. The next candidate must change a more material +component: the Q4_0 vector-kernel execution structure, MoE expert dispatch, +or intermediate BF16 output path. + +### Rejected Q4_0 FP32-intermediate candidate + +llama.cpp's HIP Q4 vector paths use an FP32 destination, while FreeToken's +normal Q4_0 MoE path returns an activation-typed BF16 tensor after each vector +product. Commit `5bbe10f` added a deliberately opt-in experiment that used +FP32 only for the two Q4_0 MoE vector-product temporaries, then converted the +final MoE result back to the original BF16 public contract. It was activated +only with `FREETOKEN_GGUF_MOE_FP32_INTERMEDIATE=1`; normal launches stayed on +the existing dtype-preserving path. The target-host build-flag and opt-in +contract tests passed before the full-model run. + +The first launch under this candidate used the literal placeholder `model` +instead of the local GGUF path and exited during model resolution. It did not +reach HIP compilation, graph capture, or an API request. The failed artifact +is retained as a labelled harness error and is excluded from every comparison. +The corrected launch used the exact Gemma GGUF checksum, offload backend, +0.50 memory ratio, graph capture, greedy AIME-25 problem 0, and 128-token +decode procedure used by the repaired baseline. + +| Candidate | Decode TPS | ms/token | TTFT | Output SHA-1 | Decision | +| --- | ---: | ---: | ---: | --- | --- | +| Repaired BF16 baseline | 55.04 | 18.169 | 295.2 ms | `abeee5e73e89` | Reference | +| FP32 intermediates | 55.11 | 18.144 | 292.6 ms | `ce247609d76c` | Rejected | + +The 0.14 percent TPS change is smaller than the observed run-to-run variation, +does not close the gap to the 60.42 client TPS ROCm 10 llama.cpp reference, +and changes the deterministic greedy response hash. The candidate was +therefore reverted and is not a shipping option. Raw evidence remains on +LAN-223 at: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/fp32-intermediate-20260828T230126Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/fp32-intermediate-retry-20260828T230209Z/ +``` + +It confirms the active device is `gfx1151`, PyTorch is +`2.13.0+rocm10.0.0` with HIP `7.15.26333`, and `/opt/rocm` resolves to +`/opt/rocm-10.0`. It also records that the system package database retains +ROCm 7.2 development packages. This alone does not prove an application +runtime conflict, so library maps were collected before changing any host +component. + +The maps show that the PyTorch wheel loads its own ROCm SDK, including LLVM 23 +and rocprofiler-sdk 1.3.5, from `_rocm_sdk_core` in the virtual environment. +The host `rocprofv3` launch initially injected a second LLVM 23 and profiler +SDK from `/opt/rocm-10.0`, causing `import torch` to abort with duplicate LLVM +registration for `spirv-expand-step`. The failure was reproduced with a +minimal PyTorch import, so it is not caused by FreeToken. + +`scripts/lan223-rocprof-wheel-sdk.sh` repairs the launch path without editing +the host installation. It keeps the host `rocprofv3` front end but passes +`--rocm-root` for the wheel's `_rocm_sdk_core`, making the profiler use the +same library identities as PyTorch. The repair was validated by profiling a +small HIP allocation and reduction. ROCm emitted `kernel_trace.csv` and +`kernel_stats.csv` with the expected GPU dispatches. Use this wrapper only +for profiling, never for TPS scoring because tracing alters execution time. + +The first full FreeToken trace launch passed PyTorch import and model loading, +then reached the GGUF JIT compiler. The profiler environment is inherited by +that compiler subprocess, so the run was stopped before a request was sent. +The next trace must warm the GGUF extension unprofiled, then profile the +already-built decode path, or explicitly prevent profiler injection into JIT +child processes. This avoids treating compile activity as token-generation +performance. + +## GGUF extension reuse validation + +The first Gemma request after the original source change built the native HIP +GGUF extension. A subsequent complete server restart retained the existing +Torch extension cache. Its first API request returned HTTP 200 and Ninja +reported `no work to do`, proving the compiled shared module was reused. +Torch still runs a lightweight hipify and dependency check before loading the +cached module; it did not run `hipcc` compilation or shared-library linking. +See the persistent-cache operating procedure in +[`amd-rocm-gfx1151.md`](amd-rocm-gfx1151.md#persistent-gguf-hip-jit-cache). + +## Commands used + +Qwen was started in the isolated environment with this functional shape: + +```bash +ft serve --model-path /home/david/freetoken-amd/models/Qwen3.6-35B-A3B-NVFP4 \ + --served-model-name qwen3.6-35b-a3b-nvfp4-amd --host 127.0.0.1 --port 18501 \ + --attention-backend triton --moe-backend offload --nvfp4-backend triton \ + --expert-load serial --moe-cache-auto --memory-ratio 0.35 \ + --max-seq-len-override 8192 --kv-reserve-tokens 2048 \ + --cuda-graph-max-bs 0 --disable-pynccl --disable-moe-prefill-overlap +``` + +Gemma used the native GGUF model file and its own loopback port: + +```bash +ft serve --model-path /home/david/freetoken-amd/models/Gemma-4-26B-A4B-it-qat-q4_0-gguf/gemma-4-26B_q4_0-it.gguf \ + --served-model-name gemma-4-26b-a4b-q4-amd --host 127.0.0.1 --port 18502 \ + --attention-backend triton --moe-backend offload --expert-load serial \ + --moe-cache-auto --memory-ratio 0.50 --max-seq-len-override 8192 \ + --kv-reserve-tokens 2048 --cuda-graph-max-bs 0 --disable-pynccl +``` + +The API checks used `/v1/models` and `/v1/chat/completions`, both with normal +JSON responses and with `stream: true`. The front-end port can answer before +the worker finishes loading, so the successful tests waited for the server log +line `API server is ready to serve` before submitting requests. + +## AMD-specific corrections verified here + +1. ROCm detection is explicit, preventing `gfx1151` from being treated as an + NVIDIA SM 11.5 capability. +2. CUDA-only optional backends are not selected on HIP. +3. DLPack and fast indexed-copy tensor handling accepts HIP tensors. +4. HIP avoids the unsafe grouped NVFP4 prefill kernel and uses the native + Triton serial expert implementation instead. This trades prompt prefill + speed for correctness on the current Strix Halo stack. +5. The Gemma GGUF JIT discovers a system Thrust include directory when the + PyTorch wheel omits Thrust. It passes that path as a compiler system + include, avoiding an attempted hipify write into the ROCm installation. +6. The same JIT adds a system ROCm library directory only when the wheel SDK + lacks the unversioned `libamdhip64.so` linker name. On LAN-223 this allowed + the native `gfx1151` object and shared module to compile and link. + +## Known limitations and follow-up work + +- This is a functional API validation, not a performance benchmark. The + recorded request timings include the chosen small fixed requests and are not + tokens-per-second claims. +- CUDA graph capture remains disabled for the HIP MVP. +- Qwen's HIP prefill deliberately uses the safe serial Triton route instead of + the grouped NVFP4 prefill route that produced an HSA aperture violation on + this machine. +- The first Gemma request compiles its GGUF HIP extension and has a substantial + cold-start cost. Later requests use the cached module. +- llama-swap integration is intentionally outside this release gate. + +## Local checks completed + +```bash +python -m compileall -q python +git diff --check +``` + +The port's HIP gate tests are retained under `tests/utils/test_rocm_runtime.py`. +The live end-to-end checks above are the required full-model validation for +this change. + +## Clean-host ROCm 10 comparison after I/O remediation + +The earlier five-run comparison was repeated after the two identified +user-space filesystem scans had been stopped with the operator's explicit +authorization. This is the decision-quality comparison: it uses the same +LAN-223 `gfx1151` device, ROCm 10 runtime, 14 GB Gemma 4 26B A4B Q4_0 GGUF, +cached AIME-25 problem 0, greedy OpenAI-compatible streamed request, and +128-token generation limit on each runner. Every scored sample starts a +fresh server, makes one excluded warm request, then makes one scored request. +Decode TPS is `(completion_tokens - 1)` divided by the client-observed time +between the first and last text SSE events. + +FreeToken uses its accepted native HIP configuration: offload backend, 4,096 +expert-cache slots, 8,320-token KV pool, 0.50 memory ratio, and graph batch +size 1. llama.cpp uses its fixed ROCm 10 `b10141` release with all layers on +the GPU (`-ngl 999`), `-c 8320`, one parallel slot, and Flash Attention on. +Both use loopback only and no production service was enabled. + +| Runtime | Five decode TPS samples | Mean TPS | Median TPS | Mean TTFT | Median p99 event gap | +| --- | --- | ---: | ---: | ---: | ---: | +| FreeToken native HIP | 60.25, 60.20, 60.17, 60.13, 60.33 | 60.21 | 60.20 | 250.0 ms | 17.74 ms | +| llama.cpp `b10141` ROCm 10 HIP | 60.24, 60.50, 62.25, 61.71, 61.83 | 61.31 | 61.71 | 113.3 ms | 16.88 ms | + +All five FreeToken completions had hash `abeee5e73e89`; all five llama.cpp +completions had hash `63a18854de72`. The two hashes are intentionally not +compared to each other because the independent implementations render their +chat templates and tokenize internally. They establish deterministic output +within each runner. llama.cpp leads by 1.8 percent on the five-run mean and +2.5 percent on the median decode rate. FreeToken's mean warm TTFT is 120.6 +percent higher. Therefore the AMD port is proven functional and stable but +does not yet meet the requested requirement to match or exceed the optimized +llama.cpp control. + +The raw, per-run result and server-log bundles remain on LAN-223: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/clean-host-freetoken-matrix-20260829T030633Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/clean-host-llamacpp-matrix-20260829T031840Z/ +``` + +The llama.cpp bundle records zero blocked (`D`) processes before and after all +five samples. The FreeToken bundle was generated after the same scan removal +and has a 0.08 TPS standard deviation, so it is the more stable side of this +matrix. The remaining performance investigation should prioritize the +observed HIP decode hot spots already captured by rocprof: Q4 MoE vector +decode first, then dense Q4 and Q6_K matrix-vector kernels. New candidates +must retain deterministic API output and be accepted only when a five-fresh- +server clean-host matrix matches or exceeds the llama.cpp median, rather than +on an isolated best run. + +### Post-matrix Q4 MoE launch experiments + +The following HIP-only experiments were run after the clean-host matrix. They +are not shipping changes. Each used the accepted Gemma workload and produced +the expected deterministic greedy response hash `abeee5e73e89`; the throughput +result, not merely successful compilation, determines rejection. + +| Candidate | Change from accepted two-row kernel | Result | Decision | +| --- | --- | ---: | --- | +| `2f019fa` | Raise the wave32 launch minimum from one to eight resident workgroups per CU | 59.95 TPS | Rejected: 0.44 percent below the accepted 60.21 TPS mean. | +| `ddaf194` | Raise the same minimum from one to two resident workgroups per CU | 60.20 TPS | Rejected: no improvement and no progress toward the 61.71 TPS gate. | +| `3f57285` | Have one wave32 calculate four rows per route instead of two | 59.22 TPS | Rejected: 1.64 percent below the accepted mean. | + +The first two experiments initially used the shared persistent extension +directory. The eight-workgroup result compiled its own source successfully; +the two-workgroup source reused the existing shared module, so its numerical +result is recorded only as a directional screen rather than a source-binary +proof. A four-row run also detected this cache reuse before it was interpreted +and is explicitly excluded. The valid four-row result then set +`TORCH_EXTENSIONS_DIR` to an artifact-local directory, rebuilt the native +`gfx1151` shared module there, and recorded that module alongside the raw logs. + +This establishes a stricter rule for all remaining performance work: every +source-changing HIP candidate must compile in a unique extension-cache path, +and the artifact must contain the resulting shared module before API timing is +accepted. The immutable raw bundles are on LAN-223: + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/hip-moe-q4-occupancy-retry-20260829T032503Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/hip-moe-q4-occupancy-two-20260829T032852Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/hip-moe-q4-four-rows-20260829T033123Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/hip-moe-q4-four-rows-isolated-cache-20260829T033230Z/ +``` + +### Isolated dense Q4 four-wave experiment + +Commit `9e36b2d` keeps the CUDA generic path intact and adds a HIP-only dense +Q4_0 dispatch wrapper that launches four wave32 rows in one workgroup. The +change addresses the second-largest measured decode hot spot, rather than the +already-tested MoE-vector kernel. The candidate passed the static ROCm gate +(`4 passed`) before the live run. + +Its live test used a fresh artifact-local `TORCH_EXTENSIONS_DIR`. The log +contains both the `hipcc --offload-arch=gfx1151` compile invocation and the +successful shared-module link. The resulting module is +`freetoken_gguf_kernels.so`, SHA-256 +`8c363e3c9345b9ab03bda75a7660d4a284642c908a03f3faeb7b21f6f078e61d`. +This proves the timing used the candidate source rather than a shared cached +extension. + +The candidate produced the expected deterministic FreeToken output hash +`abeee5e73e89` and measured **60.67 decode TPS** with a 241.9 ms warm TTFT. +That is a 0.75 percent single-run improvement over FreeToken's clean-host +60.21 TPS five-run mean, but it remains 1.68 percent below the llama.cpp +61.71 TPS five-run median acceptance gate. It is therefore retained only as +an evidence-backed non-shipping experiment, not promoted to the AMD branch or +given a five-run matrix. + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/hip-dense-q4-four-waves-isolated-cache-20260829T033846Z/ +``` + +### Isolated dense Q4 two-wave experiment + +Commit `56caf3b` tested the only remaining small workgroup-size point: two +wave32 rows per HIP Q4_0 dense-matrix-vector workgroup. As with the four-wave +experiment, the CUDA route remains unchanged and the candidate retains the +generic arithmetic, row mapping, and partial-row bounds check. The static +ROCm gate passed (`4 passed`) before the live run. + +The live run used a new artifact-local extension directory and logged a native +`hipcc --offload-arch=gfx1151` build plus link of +`freetoken_gguf_kernels.so`. It returned the exact expected output hash +`abeee5e73e89`, but measured **60.64 decode TPS** with 242.3 ms warm TTFT. +This is statistically indistinguishable from the four-wave single-run screen +(60.67 TPS), below the llama.cpp 61.71 TPS median gate, and insufficient to +justify a clean-host five-run matrix. The candidate is rejected and remains +outside the shipping AMD branch. + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/hip-dense-q4-two-waves-isolated-cache-20260829T034349Z/ +``` + +### Accepted gfx1151 RDNA3 dot-product intrinsic selection + +The current llama.cpp source was examined at immutable revision +`d7bd3bfcad3e29c7e49fd26f38c79ee3e9a3fd6b`. Its HIP helper chooses +`__builtin_amdgcn_sudot4(true, a, true, b, c, false)` on RDNA3 and RDNA4, +whereas FreeToken's copied GGUF helper always selected `sdot4` whenever that +builtin was available. Both forms implement the same signed four-byte dot +product for this Q4_0 plus Q8_1 route, but the source-level intrinsic choice +changes the gfx1151 compiler's generated code. + +Commit `5f040ba` adds a strictly scoped branch before FreeToken's existing +`sdot4` fallback. It is enabled only when the compiler exposes `sudot4` and +the device macro is `__gfx1100__`, `__gfx1150__`, or `__gfx1151__`. The +CUDA implementation, all non-RDNA3-family AMD targets, packing, scale math, +and BF16 output contract remain unchanged. The candidate passed the static +ROCm gate (`4 passed`) and built a separate native gfx1151 module with +`hipcc --offload-arch=gfx1151`. Its matrix module SHA-256 is +`5683cf07a9a081dbaf51c857757ce2307daa6822c6da4102908eb00ecd08ee3c`. + +The first five fresh-server executions all produced the expected FreeToken +output hash `abeee5e73e89`. Two were conservatively excluded because a +transient blocked process was present at their preflight snapshot, even though +none remained afterward. Two replacements explicitly waited for zero blocked +processes. The final acceptance set therefore uses runs 1, 2, 3, 6, and 7, +all with zero blocked processes before and after execution: + +| Runtime | Five clean decode TPS samples | Mean TPS | Median TPS | TPS stdev | Mean warm TTFT | +| --- | --- | ---: | ---: | ---: | ---: | +| FreeToken gfx1151 `sudot4` HIP | 61.68, 61.80, 61.84, 62.05, 61.94 | **61.86** | **61.84** | 0.14 | 241.7 ms | +| Previous FreeToken native HIP | 60.25, 60.20, 60.17, 60.13, 60.33 | 60.21 | 60.20 | 0.08 | 250.0 ms | +| llama.cpp `b10141` ROCm 10 HIP control | 60.24, 60.50, 62.25, 61.71, 61.83 | 61.31 | 61.71 | 0.88 | 113.3 ms | + +This is a 2.74 percent FreeToken mean-decode improvement over the prior clean +matrix. It exceeds the matched llama.cpp control by 0.56 TPS or 0.91 percent +on mean decode throughput, and by 0.14 TPS or 0.22 percent on median decode +throughput. The FreeToken warm TTFT remains higher, so this acceptance is +specifically for the requested sustained decode-TPS requirement. The API +remains OpenAI-compatible and deterministic for the workload. + +```text +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/hip-rdna35-sudot4-isolated-cache-20260829T035002Z/ +/home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/hip-rdna35-sudot4-clean-host-matrix-20260829T035209Z/ +``` diff --git a/pyproject.toml b/pyproject.toml index 8bd653f8..a7276fd8 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -3,7 +3,7 @@ # own, and a mismatch links the C++ extensions against the wrong libtorch. # setuptools floor: 77 is the first release that understands the PEP 639 `license` # SPDX string and `license-files` below. -requires = ["setuptools>=77", "torch>=2.11,<2.12", "wheel"] +requires = ["setuptools>=77", "torch>=2.11,<2.14", "wheel"] build-backend = "setuptools.build_meta" [project] @@ -22,6 +22,7 @@ classifiers = [ "Intended Audience :: Developers", "Intended Audience :: Science/Research", "Operating System :: POSIX :: Linux", + "Environment :: GPU :: AMD ROCm", "Environment :: GPU :: NVIDIA CUDA", "Programming Language :: Python :: 3", "Programming Language :: Python :: 3.10", @@ -54,10 +55,13 @@ dependencies = [ # floor+ceiling: sglang-kernel 0.4.5 links libtorch symbols only 2.11 has. # PyPI's torch 2.11.0 wheel is itself the cu130 build, so plain pip resolves # correctly from PyPI alone; uv additionally pins the index below. - "torch>=2.11,<2.12", + # ROCm note: on AMD (rocm.nightlies.amd.com builds) this range is intentionally + # loosened -- those wheels report their own local version segment + # (2.13.0a0+rocm...) which the sglang-kernel/cu130 constraint above doesn't apply to. + "torch>=2.11,<2.14", "tqdm>=4.66,<5", "transformers>=5.5,<6", - "triton==3.6.0; platform_system == 'Linux'", + "triton>=3.6,<3.8; platform_system == 'Linux'", "uvicorn>=0.30,<1", ] diff --git a/python/freetoken/engine/engine.py b/python/freetoken/engine/engine.py index 73dc7688..b157abb1 100644 --- a/python/freetoken/engine/engine.py +++ b/python/freetoken/engine/engine.py @@ -423,7 +423,22 @@ def __init__(self, config: EngineConfig): ) if config.attention_backend.split(",")[0] == "triton": # Prefill runs on the first comma part; warm its autotune cache. - self._warmup_prefill() + # ROCm's HIP graph and large-prompt warmup path is exercised by the first + # real request just like CUDA. Do not force that optional precompile on + # HIP at server construction: current AMD Triton releases can reject the + # synthetic 80/128-token NVFP4 MoE launch before the API becomes ready. + # Inference itself remains native HIP and eager prefill still compiles on + # demand. Operators may set this explicit opt-in for targeted testing. + should_warmup_prefill = torch.version.hip is None or os.environ.get( + "FREETOKEN_ROCM_PREFILL_WARMUP", "" + ).lower() in ("1", "true", "yes", "on") + if should_warmup_prefill: + self._warmup_prefill() + else: + logger.info_rank0( + "Skipping optional Triton prefill warmup on ROCm; " + "set FREETOKEN_ROCM_PREFILL_WARMUP=1 to enable it." + ) def _init_communication(self, config: EngineConfig) -> torch.distributed.ProcessGroup: if config.tp_info.size == 1 or config.use_pynccl: diff --git a/python/freetoken/kernel/backend.py b/python/freetoken/kernel/backend.py index 3037ad8d..137177df 100644 --- a/python/freetoken/kernel/backend.py +++ b/python/freetoken/kernel/backend.py @@ -11,6 +11,20 @@ import importlib.util +@functools.cache +def is_rocm_runtime() -> bool: + """Return whether PyTorch is backed by HIP rather than NVIDIA CUDA. + + Optional packages in this module publish CUDA binaries. Import discovery + alone is insufficient on ROCm because a stale CUDA package may be present + in an otherwise healthy environment. Returning ``False`` from each + CUDA-only capability probe preserves the existing pure-Triton fallback. + """ + import torch + + return bool(getattr(torch.version, "hip", None)) + + def _importable(name: str) -> bool: # find_spec normally returns None when a package is absent, but it can raise # (broken parent package, or a meta_path finder that blocks the name); treat @@ -23,12 +37,12 @@ def _importable(name: str) -> bool: @functools.cache def is_flashinfer_installed() -> bool: - return _importable("flashinfer") + return not is_rocm_runtime() and _importable("flashinfer") @functools.cache def is_sgl_kernel_installed() -> bool: - return _importable("sgl_kernel") + return not is_rocm_runtime() and _importable("sgl_kernel") @functools.cache @@ -39,7 +53,7 @@ def is_triton_kernels_installed() -> bool: source tree and has no Windows wheel. It is also not one of the six ops ``freetoken.kernel.triton`` reimplements, so its call-site carries its own fallback. """ - return _importable("triton_kernels") + return not is_rocm_runtime() and _importable("triton_kernels") @functools.cache @@ -50,6 +64,8 @@ def driver_cuda_version() -> int | None: toolkit version. Resolved through the ``_pinned_tensor`` extension's link-time cudart, so it works wherever the extension builds (including Windows) -- no dlopen by soname.""" + if is_rocm_runtime(): + return None try: from freetoken.kernel.pinned import _load_pinned_extension diff --git a/python/freetoken/kernel/csrc/cpu_moe/cpu_moe_ext.cpp b/python/freetoken/kernel/csrc/cpu_moe/cpu_moe_ext.cpp index 880e8637..d90e7ab2 100644 --- a/python/freetoken/kernel/csrc/cpu_moe/cpu_moe_ext.cpp +++ b/python/freetoken/kernel/csrc/cpu_moe/cpu_moe_ext.cpp @@ -29,7 +29,7 @@ #include #include -#include +#include "../hip_compat.h" #include #if defined(__linux__) diff --git a/python/freetoken/kernel/csrc/gguf/dispatch.h b/python/freetoken/kernel/csrc/gguf/dispatch.h index f42a2163..17d2a0db 100644 --- a/python/freetoken/kernel/csrc/gguf/dispatch.h +++ b/python/freetoken/kernel/csrc/gguf/dispatch.h @@ -11,6 +11,20 @@ #endif // Warp-shuffle wrappers the donor pulls from sgl-kernel's utils.h (CUDA variants). +// HIP's __shfl_xor_sync requires a 64-bit mask unconditionally (amd_warp_sync_functions.h +// static_asserts sizeof(mask) == 8) regardless of actual wavefront width; the donor's +// CUDA-style callers pass a 32-bit `unsigned int` mask (e.g. 0xffffffff), so widen it here +// rather than touching every call site. +#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) +#ifndef SGLANG_SHFL_XOR_SYNC +#define SGLANG_SHFL_XOR_SYNC(mask, var, lane_mask) \ + __shfl_xor_sync((unsigned long long)(mask), (var), (lane_mask)) +#endif +#ifndef SGLANG_SHFL_XOR_SYNC_WIDTH +#define SGLANG_SHFL_XOR_SYNC_WIDTH(mask, var, lane_mask, width) \ + __shfl_xor_sync((unsigned long long)(mask), (var), (lane_mask), (width)) +#endif +#else #ifndef SGLANG_SHFL_XOR_SYNC #define SGLANG_SHFL_XOR_SYNC(mask, var, lane_mask) __shfl_xor_sync((mask), (var), (lane_mask)) #endif @@ -18,6 +32,7 @@ #define SGLANG_SHFL_XOR_SYNC_WIDTH(mask, var, lane_mask, width) \ __shfl_xor_sync((mask), (var), (lane_mask), (width)) #endif +#endif #define DISPATCH_CASE_FLOAT_TYPES(...) \ AT_DISPATCH_CASE(at::ScalarType::Float, __VA_ARGS__) \ diff --git a/python/freetoken/kernel/csrc/gguf/ggml-common.h b/python/freetoken/kernel/csrc/gguf/ggml-common.h index 88c21a4a..2cd81e54 100644 --- a/python/freetoken/kernel/csrc/gguf/ggml-common.h +++ b/python/freetoken/kernel/csrc/gguf/ggml-common.h @@ -1004,7 +1004,15 @@ static __device__ __forceinline__ int __vsubss4(const int a, const int b) { } static __device__ __forceinline__ int __dp4a(const int a, const int b, int c) { -#if __has_builtin(__builtin_amdgcn_sdot4) +#if __has_builtin(__builtin_amdgcn_sudot4) && (defined(__gfx1100__) || defined(__gfx1150__) || defined(__gfx1151__)) + // RDNA3-family HIP compilers can lower the signed-dot form through sudot4. + // The two `true` operand flags preserve the signed four-byte dot-product + // semantics of sdot4, while matching the intrinsic selection in the current + // llama.cpp HIP implementation. This branch is intentionally limited to + // gfx1100/gfx1150/gfx1151 so older AMD targets and every CUDA build retain + // their proven implementation below. + c = __builtin_amdgcn_sudot4(true, a, true, b, c, false); +#elif __has_builtin(__builtin_amdgcn_sdot4) c = __builtin_amdgcn_sdot4(a, b, c, false); #else const int8x4_t va = reinterpret_cast(a); diff --git a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh index 8cef9e08..3d12f1e9 100644 --- a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh +++ b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh @@ -51,6 +51,23 @@ static __global__ void moe_vec_q( } } +#if defined(USE_ROCM) +// The HIP launcher is defined after the CUDA-compatible wrapper so the +// generic wrappers remain grouped by quantization format below. +template +static void moe_vec_q4_0_q8_1_hip_two_rows_cuda( + const void* vx, + const void* vy, + scalar_t* dst, + const int* topk_ids, + const int top_k, + const int tokens, + const int ncols, + const int nrows, + const int token_stride, + cudaStream_t stream); +#endif + template static void moe_vec_q4_0_q8_1_cuda( const void* vx, @@ -63,12 +80,103 @@ static void moe_vec_q4_0_q8_1_cuda( const int nrows, const int token_stride, cudaStream_t stream) { +#if defined(USE_ROCM) + // Route AMD builds through the one-wave/two-row specialization above. CUDA + // retains the established generic implementation until it has independent + // NVIDIA evidence, so this HIP experiment cannot alter CUDA behavior. + moe_vec_q4_0_q8_1_hip_two_rows_cuda( + vx, vy, dst, topk_ids, top_k, tokens, ncols, nrows, token_stride, stream); +#else const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y; const dim3 block_nums(block_num_y, 1, tokens * top_k); const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1); moe_vec_q <<>>(vx, vy, dst, topk_ids, top_k, ncols, nrows, token_stride); +#endif +} + +#if defined(USE_ROCM) +// HIP Q4_0 MoE specialization derived from the current llama.cpp MMVQ row +// structure. Unlike the older GGML_CUDA_MMV_Y=2 experiment, this launch uses +// one 32-lane wave for two rows, rather than two independent waves. The two +// float accumulators share the same packed Q4_0 activation block and expert +// selection, reducing grid work while preserving FreeToken's existing packed +// bank layout, route indexing, and BF16 output contract. +template +__launch_bounds__(WARP_SIZE, 1) +static __global__ void moe_vec_q4_0_hip_two_rows( + const void* __restrict__ vx, + const void* __restrict__ vy, + scalar_t* __restrict__ dst, + const int* __restrict__ topk_ids, + const int topk, + const int ncols, + const int nrows, + const int token_stride) { + // X indexes adjacent pairs of output rows. Y is the flattened + // token/top-k route index, matching the former Z dimension exactly. + const int row0 = 2 * blockIdx.x; + const int route = blockIdx.y; + if (row0 >= nrows) { + return; + } + + const int token = route / topk; + const int expert = topk_ids[route]; + const int blocks_per_row = ncols / QK4_0; + const int blocks_per_wave = VDR_Q4_0_Q8_1_MMVQ * WARP_SIZE / QI4_0; + const block_q4_0* x = ((const block_q4_0*)vx) + expert * nrows * blocks_per_row; + const block_q8_1* y = (const block_q8_1*)(((const int*)vy) + token * token_stride); + + // Each lane owns the same packed-Q4 range for both rows. Keeping the + // reductions separate preserves the original arithmetic for each result. + float tmp0 = 0.0f; + float tmp1 = 0.0f; + for (int i = threadIdx.x / (QI4_0 / VDR_Q4_0_Q8_1_MMVQ); i < blocks_per_row; + i += blocks_per_wave) { + const int iby = i * (QK4_0 / QK8_1); + const int iqs = VDR_Q4_0_Q8_1_MMVQ * (threadIdx.x % (QI4_0 / VDR_Q4_0_Q8_1_MMVQ)); + tmp0 += vec_dot_q4_0_q8_1(&x[row0 * blocks_per_row + i], &y[iby], iqs); + if (row0 + 1 < nrows) { + tmp1 += vec_dot_q4_0_q8_1(&x[(row0 + 1) * blocks_per_row + i], &y[iby], iqs); + } + } + + // A wave-level XOR reduction leaves the same sum in every lane. Lane zero + // writes row zero and lane one writes row one, avoiding shared memory. +#pragma unroll + for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) { + tmp0 += SGLANG_SHFL_XOR_SYNC(uint32_t(-1), tmp0, mask); + tmp1 += SGLANG_SHFL_XOR_SYNC(uint32_t(-1), tmp1, mask); + } + if (threadIdx.x == 0) { + dst[route * nrows + row0] = tmp0; + } + if (threadIdx.x == 1 && row0 + 1 < nrows) { + dst[route * nrows + row0 + 1] = tmp1; + } +} + +template +static void moe_vec_q4_0_q8_1_hip_two_rows_cuda( + const void* vx, + const void* vy, + scalar_t* dst, + const int* topk_ids, + const int top_k, + const int tokens, + const int ncols, + const int nrows, + const int token_stride, + cudaStream_t stream) { + // One block now covers two rows and one route. ``tokens * top_k`` remains + // the complete flattened routing domain used by the original launcher. + const dim3 block_nums((nrows + 1) / 2, tokens * top_k, 1); + const dim3 block_dims(WARP_SIZE, 1, 1); + moe_vec_q4_0_hip_two_rows + <<>>(vx, vy, dst, topk_ids, top_k, ncols, nrows, token_stride); } +#endif template static void moe_vec_q4_1_q8_1_cuda( diff --git a/python/freetoken/kernel/csrc/hip_compat.h b/python/freetoken/kernel/csrc/hip_compat.h new file mode 100644 index 00000000..1acd5b29 --- /dev/null +++ b/python/freetoken/kernel/csrc/hip_compat.h @@ -0,0 +1,55 @@ +#pragma once + +// Lets pinned_tensor.cpp and cpu_moe_ext.cpp call the CUDA Runtime API names they +// were written against while actually linking HIP on ROCm builds. Only the calls +// those two files use are covered -- this is not a general CUDA/HIP compat layer. +#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) +#include + +// CUDA's host-callback calling-convention annotation; empty on POSIX (matches +// cuda_runtime_api.h's own definition there). hipHostFn_t has no such annotation. +#define CUDART_CB + +using cudaError_t = hipError_t; +using cudaStream_t = hipStream_t; +constexpr hipError_t cudaSuccess = hipSuccess; +constexpr unsigned int cudaHostAllocPortable = hipHostMallocPortable; +constexpr unsigned int cudaHostAllocMapped = hipHostMallocMapped; +constexpr unsigned int cudaHostRegisterPortable = hipHostRegisterPortable; +constexpr unsigned int cudaHostRegisterMapped = hipHostRegisterMapped; +constexpr hipDeviceAttribute_t cudaDevAttrUnifiedAddressing = + hipDeviceAttributeUnifiedAddressing; +constexpr hipDeviceAttribute_t cudaDevAttrCanUseHostPointerForRegisteredMem = + hipDeviceAttributeCanUseHostPointerForRegisteredMem; + +inline hipError_t cudaMallocHost(void **ptr, size_t size) { + return hipHostMalloc(ptr, size, hipHostMallocDefault); +} +inline hipError_t cudaFreeHost(void *ptr) { return hipHostFree(ptr); } +inline hipError_t cudaHostAlloc(void **ptr, size_t size, unsigned int flags) { + return hipHostMalloc(ptr, size, flags); +} +inline hipError_t cudaGetDevice(int *device) { return hipGetDevice(device); } +inline hipError_t cudaDeviceGetAttribute(int *value, hipDeviceAttribute_t attr, + int device) { + return hipDeviceGetAttribute(value, attr, device); +} +inline hipError_t cudaHostGetDevicePointer(void **devPtr, void *hostPtr, + unsigned int flags) { + return hipHostGetDevicePointer(devPtr, hostPtr, flags); +} +inline hipError_t cudaHostRegister(void *ptr, size_t size, unsigned int flags) { + return hipHostRegister(ptr, size, flags); +} +inline hipError_t cudaDriverGetVersion(int *v) { return hipDriverGetVersion(v); } +inline const char *cudaGetErrorString(hipError_t e) { return hipGetErrorString(e); } +inline hipError_t cudaStreamSynchronize(hipStream_t s) { + return hipStreamSynchronize(s); +} +inline hipError_t cudaLaunchHostFunc(hipStream_t s, hipHostFn_t fn, void *data) { + return hipLaunchHostFunc(s, fn, data); +} + +#else +#include +#endif diff --git a/python/freetoken/kernel/csrc/include/freetoken/utils.cuh b/python/freetoken/kernel/csrc/include/freetoken/utils.cuh index 8e917832..ea472b63 100644 --- a/python/freetoken/kernel/csrc/include/freetoken/utils.cuh +++ b/python/freetoken/kernel/csrc/include/freetoken/utils.cuh @@ -10,6 +10,50 @@ #include #include +// nvcc implicitly pulls in the CUDA runtime for .cu translation units; hipcc does +// not do the equivalent for HIP, so it must be included explicitly here. On the +// HIP path there is no cudaLaunchKernelEx/cudaLaunchConfig_t equivalent (that API +// is Hopper PDL-specific), so LaunchKernel gets its own HIP-side definition below +// instead of a name-aliasing shim -- see PDL below for why that also means +// with_attr(true) is a no-op on this path. +#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) +#include + +using cudaError_t = hipError_t; +constexpr hipError_t cudaSuccess = hipSuccess; +using cudaStream_t = hipStream_t; + +inline const char *cudaGetErrorString(hipError_t e) { return hipGetErrorString(e); } +inline hipError_t cudaGetLastError() { return hipGetLastError(); } +inline hipError_t cudaFuncSetAttribute(const void *func, hipFuncAttribute attr, + int value) { + return hipFuncSetAttribute(func, attr, value); +} +constexpr hipFuncAttribute cudaFuncAttributeMaxDynamicSharedMemorySize = + hipFuncAttributeMaxDynamicSharedMemorySize; + +inline hipError_t cudaGetDevice(int *device) { return hipGetDevice(device); } +inline hipError_t cudaDeviceGetAttribute(int *value, hipDeviceAttribute_t attr, + int device) { + return hipDeviceGetAttribute(value, attr, device); +} +inline hipError_t cudaHostGetDevicePointer(void **devPtr, void *hostPtr, + unsigned int flags) { + return hipHostGetDevicePointer(devPtr, hostPtr, flags); +} +constexpr hipDeviceAttribute_t cudaDevAttrUnifiedAddressing = + hipDeviceAttributeUnifiedAddressing; +constexpr hipDeviceAttribute_t cudaDevAttrCanUseHostPointerForRegisteredMem = + hipDeviceAttributeCanUseHostPointerForRegisteredMem; + +// CUDA-only kernel-parameter annotation (passes large by-value params via constant +// memory instead of copying them into local/generic memory first); HIP has no +// equivalent attribute, so this just falls back to an ordinary by-value parameter. +#define __grid_constant__ +#else +#include +#endif + namespace device { inline constexpr auto kWarpThreads = 32u; @@ -42,16 +86,24 @@ __always_inline __device__ auto offset(const T *ptr, U... offset) -> const namespace PDL { +// Programmatic Dependent Launch is a Hopper-only CUDA hardware feature; the PTX +// below has no HIP/ROCm equivalent. Callers gate kUsePDL off for non-Hopper CUDA +// targets already, and LaunchKernel::with_attr is a no-op on HIP (see below), so +// this stays unconditionally a no-op there rather than a compile failure. template __always_inline __device__ void wait() { +#if !(defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)) if constexpr (kUsePDL) { asm volatile("griddepcontrol.wait;" ::: "memory"); } +#endif } template __always_inline __device__ void launch() { +#if !(defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)) if constexpr (kUsePDL) { asm volatile("griddepcontrol.launch_dependents;" :::); } +#endif } } // namespace PDL @@ -88,6 +140,50 @@ template inline void set_smem_once(std::size_t smem_size) { last_smem_size, " bytes"); } +#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) + +// HIP has no cudaLaunchKernelEx/cudaLaunchConfig_t analog (that API only exists to +// carry Hopper PDL attributes, which ROCm hardware has no equivalent for), so this +// launches via the plain triple-chevron form instead. with_attr(true) is therefore +// a no-op here -- there is no attribute to carry. +struct LaunchKernel { +public: + explicit LaunchKernel(dim3 grid_dim, dim3 block_dim, DLDevice device, + std::size_t dynamic_shared_mem_bytes = 0) noexcept + : m_grid_dim(grid_dim), m_block_dim(block_dim), + m_smem(dynamic_shared_mem_bytes), m_stream(resolve_device(device)) {} + + explicit LaunchKernel(dim3 grid_dim, dim3 block_dim, cudaStream_t stream, + std::size_t dynamic_shared_mem_bytes = 0) noexcept + : m_grid_dim(grid_dim), m_block_dim(block_dim), + m_smem(dynamic_shared_mem_bytes), m_stream(stream) {} + + static auto resolve_device(DLDevice device) -> cudaStream_t { + return static_cast( + ::TVMFFIEnvGetStream(device.device_type, device.device_id)); + } + + LaunchKernel(const LaunchKernel &) = delete; + LaunchKernel &operator=(const LaunchKernel &) = delete; + + template + auto operator()(T &&kernel, Args &&...args) const -> void { + kernel<<>>( + std::forward(args)...); + CUDA_CHECK(::cudaGetLastError()); + } + + auto with_attr(bool /*use_pdl*/) -> LaunchKernel & { return *this; } + +private: + dim3 m_grid_dim; + dim3 m_block_dim; + std::size_t m_smem; + cudaStream_t m_stream; +}; + +#else + struct LaunchKernel { public: explicit LaunchKernel(dim3 grid_dim, dim3 block_dim, DLDevice device, @@ -141,4 +237,6 @@ private: cudaLaunchAttribute m_attr_cache; }; +#endif + } // namespace host diff --git a/python/freetoken/kernel/csrc/jit/fast_index_copy.cuh b/python/freetoken/kernel/csrc/jit/fast_index_copy.cuh index bb83c23e..2d1dbc05 100644 --- a/python/freetoken/kernel/csrc/jit/fast_index_copy.cuh +++ b/python/freetoken/kernel/csrc/jit/fast_index_copy.cuh @@ -33,6 +33,38 @@ inline constexpr auto get_mem_package() { } } +// The ld.global.L1::no_allocate / st.global.wt PTX below are cache-policy hints +// (skip L1 allocate on read, write-through on store) with no HIP equivalent -- AMD +// ROCm builds fall back to plain loads/stores. Correctness is unchanged; only the +// cache-policy hint is lost. +#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) + +__always_inline __device__ auto load_nc(const uint1* __restrict__ src) -> uint1 { + return *src; +} + +__always_inline __device__ auto load_nc(const uint2* __restrict__ src) -> uint2 { + return *src; +} + +__always_inline __device__ auto load_nc(const uint4* __restrict__ src) -> uint4 { + return *src; +} + +__always_inline __device__ void store_nc(uint1* __restrict__ dst, const uint1& value) { + *dst = value; +} + +__always_inline __device__ void store_nc(uint2* __restrict__ dst, const uint2& value) { + *dst = value; +} + +__always_inline __device__ void store_nc(uint4* __restrict__ dst, const uint4& value) { + *dst = value; +} + +#else + __always_inline __device__ auto load_nc(const uint1* __restrict__ src) -> uint1 { uint32_t tmp; asm volatile("ld.global.L1::no_allocate.b32 %0,[%1];" : "=r"(tmp) : "l"(src)); @@ -70,6 +102,8 @@ __always_inline __device__ void store_nc(uint4* __restrict__ dst, const uint4& v asm volatile("st.global.wt.v4.b32 [%0],{%1,%2,%3,%4};" ::"l"(dst), "r"(tmp0), "r"(tmp1), "r"(tmp2), "r"(tmp3)); } +#endif + __always_inline __device__ void wait_flag_clear(const int32_t* __restrict__ flag_ptr) { // Exponential backoff to avoid hammering a global atomic in a tight loop. auto* flag = reinterpret_cast(const_cast(flag_ptr)); @@ -344,17 +378,17 @@ struct FastIndexCopyKernel { TensorMatcher({-1, D}) .with_dtype(data_dtype) - .with_device() + .with_device() .verify(src); TensorMatcher({-1, D}) .with_dtype(data_dtype) - .with_device() + .with_device() .verify(dst); TensorMatcher({L}) .with_dtype(indices_dtype) - .with_device(device) + .with_device(device) .verify(src_indices) .verify(dst_indices); @@ -363,7 +397,7 @@ struct FastIndexCopyKernel { const auto num_indices_tensor = num_indices.value(); TensorMatcher({1}) .with_dtype(num_indices_dtype) - .with_device(device) + .with_device(device) .verify(num_indices_tensor); num_indices_data_ptr = static_cast(num_indices_tensor.data_ptr()); diff --git a/python/freetoken/kernel/csrc/jit/index.cu b/python/freetoken/kernel/csrc/jit/index.cu index ca0e1db2..aca58383 100644 --- a/python/freetoken/kernel/csrc/jit/index.cu +++ b/python/freetoken/kernel/csrc/jit/index.cu @@ -114,15 +114,15 @@ struct IndexKernel { TensorMatcher({-1, D}) // .with_dtype(weights_dtype_) - .with_device(device_) + .with_device(device_) .verify(weights); TensorMatcher({L, D}) // .with_dtype(weights_dtype_) - .with_device(device_) + .with_device(device_) .verify(output); TensorMatcher({L}) // .with_dtype(indices_dtype_) - .with_device(device_) + .with_device(device_) .verify(indices); const auto device = device_.unwrap(); diff --git a/python/freetoken/kernel/csrc/jit/store.cu b/python/freetoken/kernel/csrc/jit/store.cu index 8d84d76e..162dfdfe 100644 --- a/python/freetoken/kernel/csrc/jit/store.cu +++ b/python/freetoken/kernel/csrc/jit/store.cu @@ -72,18 +72,18 @@ struct StoreKernel { TensorMatcher({-1, D}) // .with_strides({X, 1}) - .with_device(device_) + .with_device(device_) .with_dtype(dtype_) .verify(k_cache) .verify(v_cache); TensorMatcher({L, D}) // .with_strides({Y, 1}) - .with_device(device_) + .with_device(device_) .with_dtype(dtype_) .verify(k) .verify(v); TensorMatcher({L}) // - .with_device(device_) + .with_device(device_) .with_dtype(indices_dtype_) .verify(indices); diff --git a/python/freetoken/kernel/csrc/pinned_tensor.cpp b/python/freetoken/kernel/csrc/pinned_tensor.cpp index c3947adf..4cb983f2 100644 --- a/python/freetoken/kernel/csrc/pinned_tensor.cpp +++ b/python/freetoken/kernel/csrc/pinned_tensor.cpp @@ -1,5 +1,5 @@ #include -#include +#include "hip_compat.h" #include namespace { diff --git a/python/freetoken/kernel/gguf.py b/python/freetoken/kernel/gguf.py index 04a16560..60dbbb19 100644 --- a/python/freetoken/kernel/gguf.py +++ b/python/freetoken/kernel/gguf.py @@ -22,6 +22,86 @@ _CSRC = pathlib.Path(__file__).parent / "csrc" / "gguf" +def _hip_target_arch() -> str | None: + """Return the active AMD GPU target in ``gfxNNNN`` form when HIP exposes it. + + PyTorch's extension builder otherwise emits code for every visible AMD target. + A one-GPU serving process only needs the active target, so preserving an explicit + user selection or deriving the target from the active device avoids unnecessary + JIT work and records the architecture in the extension build key. + """ + explicit = os.environ.get("PYTORCH_ROCM_ARCH", "").strip() + if explicit: + return explicit.split(";", 1)[0].strip() + if not torch.cuda.is_available(): + return None + arch = getattr(torch.cuda.get_device_properties(0), "gcnArchName", "") + return str(arch).split(":", 1)[0] or None + + +def _hip_gguf_cflags() -> list[str]: + """Build conservative HIP GGUF flags for the active AMD GPU target. + + The architecture environment variable is set before PyTorch asks hipcc to + compile, which makes the cache target-specific without overriding a deployment's + explicit multi-target configuration. Keep floating-point flags conservative: + the native GGUF kernels must preserve model output, and unsupported aggressive + math flags belong only in isolated benchmark experiments. + """ + target = _hip_target_arch() + if target and not os.environ.get("PYTORCH_ROCM_ARCH"): + os.environ["PYTORCH_ROCM_ARCH"] = target + return ["-O3"] + + +def _hip_thrust_include() -> str | None: + """Return a ROCm developer include directory that exposes ``thrust/complex.h``. + + The PyTorch ROCm wheel bundles hipcc but may omit the header-only Thrust + dependency required by libtorch's HIP complex header. Prefer explicitly + configured ROCm homes, then inspect the standard versioned installation + layout. Returning ``None`` leaves hosts with a complete wheel toolchain + unchanged. + """ + candidates = [ + os.environ.get("ROCM_HOME"), + os.environ.get("ROCM_PATH"), + "/opt/rocm", + ] + candidates.extend(str(path) for path in sorted(pathlib.Path("/opt").glob("rocm-*"), reverse=True)) + for root in candidates: + if not root: + continue + include = pathlib.Path(root) / "include" + if (include / "thrust" / "complex.h").is_file(): + return str(include) + return None + + +def _hip_runtime_library_dir() -> str | None: + """Return a ROCm library directory that can satisfy ``-lamdhip64``. + + Some PyTorch ROCm wheels ship ``libamdhip64.so.7`` but not the unversioned + linker name that ``torch.utils.cpp_extension`` emits. A regular ROCm + installation supplies that linker name under its ``lib`` directory. Keep + this discovery separate from the Thrust fallback so a host can provide one + dependency through the wheel and the other through its ROCm installation. + """ + candidates = [ + os.environ.get("ROCM_HOME"), + os.environ.get("ROCM_PATH"), + "/opt/rocm", + ] + candidates.extend(str(path) for path in sorted(pathlib.Path("/opt").glob("rocm-*"), reverse=True)) + for root in candidates: + if not root: + continue + for lib_dir in (pathlib.Path(root) / "lib", pathlib.Path(root) / "lib64"): + if (lib_dir / "libamdhip64.so").is_file(): + return str(lib_dir) + return None + + def _host_compiler() -> str | None: """A host compiler nvcc + libtorch headers accept. @@ -51,24 +131,50 @@ def _c_compiler_for(cxx: str) -> str: def _module(): from torch.utils.cpp_extension import load - extra_cuda_cflags = ["-O3", "--expt-relaxed-constexpr"] - host_cxx = _host_compiler() - if host_cxx is not None: - # Point both nvcc's host pass (-ccbin) and torch's C++ compile (CXX) at a - # libtorch/nvcc-compatible compiler. Force (not setdefault): the system - # default (CXX unset -> g++) can be a gcc too new for the torch headers. - cxx_path = shutil.which(host_cxx) or host_cxx - extra_cuda_cflags += ["-ccbin", cxx_path] - os.environ["CXX"] = cxx_path - os.environ["CC"] = _c_compiler_for(cxx_path) + if torch.version.hip is not None: + # Neither issue -ccbin works around applies under hipcc: it has no separate + # nvcc-style host pass (its own bundled clang IS the host compiler), and + # --expt-relaxed-constexpr is an nvcc-only flag hipcc/clang rejects outright. + extra_cuda_cflags = _hip_gguf_cflags() + # The minimal PyTorch ROCm SDK can omit Thrust while libtorch's HIP + # headers include it. Add a real system ROCm developer include only + # when present, retaining the wheel-only build on complete installs. + # This must be a compiler flag, not ``extra_include_paths``: PyTorch's + # hipify pass recursively rewrites every extension include path and + # cannot write beneath the read-only system ROCm installation. + hip_thrust_include = _hip_thrust_include() + hip_runtime_library_dir = _hip_runtime_library_dir() + extra_include_paths = [str(_CSRC)] + extra_ldflags: list[str] = [] + if hip_thrust_include is not None: + extra_cuda_cflags += ["-isystem", hip_thrust_include] + if hip_runtime_library_dir is not None: + # The extension linker uses ``-lamdhip64``. Add a real ROCm + # library directory only when the wheel SDK lacks its unversioned + # linker symlink, preserving self-contained wheel installations. + extra_ldflags += [f"-L{hip_runtime_library_dir}"] + else: + extra_cuda_cflags = ["-O3", "--expt-relaxed-constexpr"] + host_cxx = _host_compiler() + if host_cxx is not None: + # Point both nvcc's host pass (-ccbin) and torch's C++ compile (CXX) at a + # libtorch/nvcc-compatible compiler. Force (not setdefault): the system + # default (CXX unset -> g++) can be a gcc too new for the torch headers. + cxx_path = shutil.which(host_cxx) or host_cxx + extra_cuda_cflags += ["-ccbin", cxx_path] + os.environ["CXX"] = cxx_path + os.environ["CC"] = _c_compiler_for(cxx_path) + extra_include_paths = [str(_CSRC)] + extra_ldflags = [] # gguf_kernel.cu carries its own PYBIND11_MODULE (appended at the end), so a # plain `load` of the single source compiles + binds the ggml_* ops. return load( name="freetoken_gguf_kernels", sources=[str(_CSRC / "gguf_kernel.cu")], - extra_include_paths=[str(_CSRC)], + extra_include_paths=extra_include_paths, extra_cuda_cflags=extra_cuda_cflags, + extra_ldflags=extra_ldflags, verbose=True, ) diff --git a/python/freetoken/kernel/triton/activation.py b/python/freetoken/kernel/triton/activation.py index 2c38b533..c26354ec 100644 --- a/python/freetoken/kernel/triton/activation.py +++ b/python/freetoken/kernel/triton/activation.py @@ -23,6 +23,13 @@ from freetoken.utils.arch import is_sm90_supported +# _fast_tanh/_fast_ex2 below inline raw PTX text (tanh.approx.f32, ex2.approx.f32) via +# tl.inline_asm_elementwise. HIP's inline-asm path doesn't reject PTX outright -- it +# fails much later, deep in register allocation ("couldn't allocate output register +# for constraint 'f'"), since the constraint syntax is generic LLVM inline-asm but the +# instruction text is NVIDIA-ISA-only. Route ROCm through portable tl/libdevice ops. +_IS_HIP = tl.constexpr(torch.version.hip is not None) + SILU = 0 GELU = 1 GELU_TANH = 2 @@ -48,6 +55,8 @@ def _pdl_supported() -> bool: @triton.jit def _fast_tanh(x): + if _IS_HIP: + return libdevice.tanh(x) # PTX tanh.approx.f32 — single HW op, matches flashinfer math::tanh. return tl.inline_asm_elementwise( "tanh.approx.f32 $0, $1;", "=f,f", [x], @@ -57,6 +66,8 @@ def _fast_tanh(x): @triton.jit def _fast_ex2(x): + if _IS_HIP: + return tl.exp2(x) # PTX ex2.approx.f32 — matches __expf fast path used by flashinfer silu. return tl.inline_asm_elementwise( "ex2.approx.f32 $0, $1;", "=f,f", [x], @@ -129,12 +140,16 @@ def _act_and_mul( M = x2.shape[0] grid = lambda meta: (M, triton.cdiv(d, meta["BLOCK_D"])) pdl = _pdl_supported() + # launch_pdl is a CUDA-Hopper-only Triton launch kwarg; the AMD backend's + # arg-packer rejects it outright (KeyError) even when passed as False, so it + # is only included on the one backend/arch combination that ever sets pdl=True. + pdl_kwargs = {"launch_pdl": pdl} if pdl else {} # Fixed via H100 sweep (72-config grid; 512/w4/s3 within 11% everywhere, # 1024/w4/s2 best at rows>=4096). block_d = min(triton.next_power_of_2(d), 1024 if M >= 4096 else 512) num_stages = 2 if block_d == 1024 else 3 _act_and_mul_kernel[grid]( - o2, x2, d, alpha, limit, ACT=kind, ENABLE_PDL=pdl, launch_pdl=pdl, + o2, x2, d, alpha, limit, ACT=kind, ENABLE_PDL=pdl, **pdl_kwargs, BLOCK_D=block_d, num_warps=4, num_stages=num_stages, ) return out diff --git a/python/freetoken/kernel/triton/attention.py b/python/freetoken/kernel/triton/attention.py index c2358d84..cbf36637 100644 --- a/python/freetoken/kernel/triton/attention.py +++ b/python/freetoken/kernel/triton/attention.py @@ -364,8 +364,17 @@ def decode_paged_attention( sliding_window: int | None = None, sinks: torch.Tensor | None = None, out: torch.Tensor | None = None, + rocm_block_h_probe: int | None = None, + rocm_block_n_probe: int | None = None, + rocm_num_warps_probe: int | None = None, ) -> torch.Tensor: - """SGLang-style split-k grouped decode attention for one query per request.""" + """SGLang-style split-k grouped decode attention for one query per request. + + The ``rocm_*_probe`` arguments are benchmark-only HIP controls. They let + LAN-223 measure a query-head tile, KV block length, or launch warp count + without changing the serving defaults. Normal callers leave every probe + argument ``None`` and preserve the established ROCm configuration. + """ assert q.is_cuda and k_cache.is_cuda and v_cache.is_cuda assert q.dim() == 3 and k_cache.dim() == 3 and v_cache.dim() == 3 @@ -396,8 +405,35 @@ def decode_paged_attention( # (e.g. 6), where block_h rounds up and the kernel masks the extra lanes. valid_block_h = min(16, group) block_h = triton.next_power_of_2(valid_block_h) + if rocm_block_h_probe is not None: + if torch.version.hip is None: + raise ValueError("rocm_block_h_probe is only valid for HIP builds") + if rocm_block_h_probe < valid_block_h or rocm_block_h_probe & (rocm_block_h_probe - 1): + raise ValueError("rocm_block_h_probe must be a power of two at least valid_block_h") + block_h = rocm_block_h_probe + elif torch.version.hip is not None: + # RDNA WMMA has no matrix-core instruction below a 16x16 tile, so a decode + # GQA group smaller than 16 (e.g. 4 here) leaves tl.dot's M dim too small to + # lower on this backend. The kernel already masks lanes >= VALID_BLOCK_H + # (it does this for non-power-of-two groups too), so padding BLOCK_H up to + # 16 is safe -- it only adds masked-out, discarded head lanes. + block_h = max(block_h, 16) block_d = triton.next_power_of_2(head_dim) block_dv = triton.next_power_of_2(head_dim) + block_n = 32 + num_warps = 4 + if rocm_block_n_probe is not None: + if torch.version.hip is None: + raise ValueError("rocm_block_n_probe is only valid for HIP builds") + if rocm_block_n_probe < 16 or rocm_block_n_probe & (rocm_block_n_probe - 1): + raise ValueError("rocm_block_n_probe must be a power of two at least 16") + block_n = rocm_block_n_probe + if rocm_num_warps_probe is not None: + if torch.version.hip is None: + raise ValueError("rocm_num_warps_probe is only valid for HIP builds") + if rocm_num_warps_probe not in (1, 2, 4, 8): + raise ValueError("rocm_num_warps_probe must be one of 1, 2, 4, or 8") + num_warps = rocm_num_warps_probe _decode_grouped_stage1_kernel[ (batch, triton.cdiv(num_q_heads, valid_block_h), max_kv_splits) @@ -428,14 +464,14 @@ def decode_paged_attention( NUM_Q_HEADS=num_q_heads, BLOCK_D=block_d, BLOCK_DV=block_dv, - BLOCK_N=32, + BLOCK_N=block_n, BLOCK_H=block_h, VALID_BLOCK_H=valid_block_h, MIN_BLOCK_KV=_MIN_BLOCK_KV, D=head_dim, DV=head_dim, SLIDING_WINDOW=sliding_window or 0, - num_warps=4, + num_warps=num_warps, num_stages=2, ) _decode_stage2_kernel[(batch, num_q_heads)]( diff --git a/python/freetoken/kernel/triton/e4m3_compat.py b/python/freetoken/kernel/triton/e4m3_compat.py index 1d9f744c..d5c923ad 100644 --- a/python/freetoken/kernel/triton/e4m3_compat.py +++ b/python/freetoken/kernel/triton/e4m3_compat.py @@ -59,6 +59,14 @@ def e4m3_native() -> bool: if _native is None: if FORCE_EMU: _native = False + elif torch.version.hip is not None: + # torch.cuda.get_device_capability() on a HIP build returns the gfx/RDNA + # generation number (e.g. (11, 5) for gfx1150), not a CUDA compute + # capability -- comparing it against (8, 9) below is a tuple comparison + # over two unrelated numbering schemes and can false-positive (11 > 8). + # No AMD GPU has this fp8e4nv unit; e4m3_native_cx() (Triton's own, + # backend-aware check) already agrees this must be False. + _native = False else: from freetoken.gpu_select import assigned_visible_gpu diff --git a/python/freetoken/kernel/triton/norm.py b/python/freetoken/kernel/triton/norm.py index 3f95c29f..62bd6382 100644 --- a/python/freetoken/kernel/triton/norm.py +++ b/python/freetoken/kernel/triton/norm.py @@ -142,9 +142,13 @@ def _rmsnorm(input, weight, eps, out, gemma: bool): # PDL only on the contiguous (decode-replay) path: on the strided qk-norm's # 32k-CTA prefill grids the per-CTA gdc_wait poll costs more than it hides. pdl = contig and is_sm90_supported() + # launch_pdl is a CUDA-Hopper-only Triton launch kwarg; the AMD backend's + # arg-packer rejects it outright (KeyError) even when passed as False, so it + # is only included on the one backend/arch combination that ever sets pdl=True. + pdl_kwargs = {"launch_pdl": pdl} if pdl else {} _rmsnorm_kernel[(A, B)]( out, input, weight, eps, H, sxa, sxb, soa, sob, - CONTIG=contig, ENABLE_PDL=pdl, launch_pdl=pdl, GEMMA=gemma, + CONTIG=contig, ENABLE_PDL=pdl, GEMMA=gemma, **pdl_kwargs, num_warps=_num_warps(A * B), num_stages=1, ) return out @@ -170,9 +174,10 @@ def _fused_add_rmsnorm(input, residual, weight, eps, gemma: bool): _, _, sra, srb = _leading(residual) contig = input.ndim == 2 and input.is_contiguous() and residual.is_contiguous() pdl = contig and is_sm90_supported() + pdl_kwargs = {"launch_pdl": pdl} if pdl else {} _fused_add_rmsnorm_kernel[(A, B)]( input, residual, weight, eps, H, sxa, sxb, sra, srb, - CONTIG=contig, ENABLE_PDL=pdl, launch_pdl=pdl, GEMMA=gemma, + CONTIG=contig, ENABLE_PDL=pdl, GEMMA=gemma, **pdl_kwargs, num_warps=_num_warps(A * B), num_stages=1, ) diff --git a/python/freetoken/kernel/utils.py b/python/freetoken/kernel/utils.py index 7a0164b5..42f15a5b 100644 --- a/python/freetoken/kernel/utils.py +++ b/python/freetoken/kernel/utils.py @@ -30,7 +30,13 @@ def _cuda_cflags(extra: List[str]) -> List[str]: PTX→SASS JIT (driver-only, no CUDA toolkit). One top PTX suffices: the loader always JIT-forwards from the highest compatible PTX. When the env is unset (runtime JIT), this is a no-op and tvm-ffi targets only the local GPU.""" - flags = DEFAULT_CUDA_CFLAGS + extra + import torch + + flags = list(DEFAULT_CUDA_CFLAGS) + if torch.version.hip is not None: + # nvcc-only: hipcc/clang rejects it outright. + flags = [f for f in flags if f != "--expt-relaxed-constexpr"] + flags = flags + extra arch_list = os.getenv("TVM_FFI_CUDA_ARCH_LIST", "").split() if arch_list: def _rank(a: str) -> int: diff --git a/python/freetoken/moe/fused.py b/python/freetoken/moe/fused.py index 4b9a4875..d33655b4 100644 --- a/python/freetoken/moe/fused.py +++ b/python/freetoken/moe/fused.py @@ -44,21 +44,28 @@ def fused_topk( ) -> Tuple[torch.Tensor, torch.Tensor]: assert hidden_states.shape[0] == gating_output.shape[0], "Number of tokens mismatch" - from freetoken.kernel.backend import is_triton_kernels_installed + from freetoken.kernel.backend import is_rocm_runtime, is_triton_kernels_installed - # triton_kernels ships no Windows wheel, and unlike flashinfer/sgl_kernel it is not one - # of the six ops the in-repo triton kernels cover -- so this router needs its own fallback. + # OpenAI's triton_kernels package distributes CUDA-only binaries. The + # in-tree Triton router is useful for research on HIP, but it has not yet + # met this runner's exact greedy end-to-end output contract on ROCm, so + # production HIP retains the reference PyTorch router below. if not is_triton_kernels_installed(): global _warned_torch_topk if not _warned_torch_topk: _warned_torch_topk = True - # Once, not per call: this runs every MoE forward. On Linux a missing - # triton_kernels used to fail fast with ImportError; keep the misconfiguration - # visible without giving up the fallback that Windows needs. + # Once, not per call: this runs every MoE forward. ROCm has no + # supported triton_kernels package, while CUDA Linux may restore + # the optimized package by installing it. Keep the distinction + # explicit so an AMD operator is not told to install CUDA binaries. + reason = ( + "ROCm keeps the reference pure-torch router" + if is_rocm_runtime() + else "triton_kernels is not installed" + ) logger.warning_rank0( - "fused_topk: triton_kernels is not installed -> pure-torch router fallback " - "(numerically equivalent, slower). Expected on Windows (no wheel); on Linux " - "install triton_kernels to restore the fused router." + f"fused_topk: {reason} -> pure-torch router fallback " + "(numerically equivalent, slower)." ) return _torch_fused_topk(gating_output, topk, renormalize, num_token_non_padded) diff --git a/python/freetoken/moe/fused_nvfp4.py b/python/freetoken/moe/fused_nvfp4.py index 29a561e6..4b54ae6f 100644 --- a/python/freetoken/moe/fused_nvfp4.py +++ b/python/freetoken/moe/fused_nvfp4.py @@ -323,6 +323,29 @@ def fused_experts_nvfp4( """Prefill inline-NVFP4 MoE. ``topk_ids`` index rows of the bank tensors in ``[0, num_experts)``: full-layer banks with position == expert id (the materialized ``[:E]`` slot view or the overlap double buffer), raw ids.""" + if torch.version.hip is not None: + # The grouped prefill kernel below currently trips an HSA memory-aperture + # violation on gfx1151. The serial Triton kernel is already FreeToken's + # native inline-dequant implementation and accepts an arbitrary M, so it + # preserves HIP GPU inference and model results without materializing BF16 + # experts. It is intentionally slower for prompt prefill than the CUDA + # grouped kernel, but is safe until the grouped launch is ROCm-qualified. + return fused_experts_decode_nvfp4_serial( + hidden_states, + gate_up_packed, + gate_up_scale, + gate_up_global, + down_packed, + down_scale, + down_global, + topk_weights, + topk_ids, + activation, + apply_router_weight_on_input, + act_alpha, + act_limit, + ) + M, H = hidden_states.shape top_k = topk_ids.shape[1] two_i = gate_up_packed.shape[1] diff --git a/python/freetoken/utils/__init__.py b/python/freetoken/utils/__init__.py index 2e4ad15f..bcd2d544 100644 --- a/python/freetoken/utils/__init__.py +++ b/python/freetoken/utils/__init__.py @@ -1,5 +1,6 @@ from .arch import ( is_arch_supported, + is_rocm_runtime, is_sm90_family, is_sm90_supported, is_sm100_family, @@ -35,6 +36,7 @@ "load_toolcall_anchor_id", "init_logger", "is_arch_supported", + "is_rocm_runtime", "is_sm90_family", "is_sm90_supported", "is_sm100_family", diff --git a/python/freetoken/utils/arch.py b/python/freetoken/utils/arch.py index 8c1c6c3d..422cdce0 100644 --- a/python/freetoken/utils/arch.py +++ b/python/freetoken/utils/arch.py @@ -4,12 +4,30 @@ from typing import Tuple +@functools.cache +def is_rocm_runtime() -> bool: + """Return whether the active PyTorch build uses AMD's HIP runtime. + + PyTorch intentionally preserves the ``torch.cuda`` namespace on ROCm for + source compatibility. Consequently, a Radeon architecture such as + ``gfx1151`` can be reported as a numeric capability that superficially + resembles a newer NVIDIA SM version. Architecture gates in this module + control NVIDIA-only features such as Programmatic Dependent Launch, so + they must reject HIP before comparing those numeric values. + """ + import torch + + return bool(getattr(torch.version, "hip", None)) + + @functools.cache def _get_torch_cuda_version() -> Tuple[int, int] | None: import torch import torch.version - if not torch.cuda.is_available() or not torch.version.cuda: + # ROCm retains torch.cuda APIs, but neither CUDA SM feature checks nor the + # numeric capability ordering below are meaningful for an AMD GPU. + if is_rocm_runtime() or not torch.cuda.is_available() or not torch.version.cuda: return None return torch.cuda.get_device_capability() diff --git a/scripts/lan223-capture-baseline.sh b/scripts/lan223-capture-baseline.sh new file mode 100644 index 00000000..42ffe9a3 --- /dev/null +++ b/scripts/lan223-capture-baseline.sh @@ -0,0 +1,138 @@ +#!/usr/bin/env bash +# Capture a secret-free, read-only LAN-223 ROCm baseline for a FreeToken run. +# +# The script intentionally does not start a server, alter GPU clocks, install +# packages, delete cache entries, or edit system configuration. It records +# the environment that makes a later throughput claim reproducible. + +# Fail on an unset variable, an unsuccessful command in a pipeline, or a +# command error. Individual optional probes use `|| true` so that a missing +# diagnostic utility is recorded without invalidating the whole manifest. +set -euo pipefail + +# Keep the output path explicit. A caller may pass a unique campaign folder; +# the default is suitable only for a one-off local capture. +output_dir="${1:-./artifacts/lan223-baseline-$(date -u +%Y%m%dT%H%M%SZ)}" + +# Accept the GGUF path as an optional second argument. Hashing the exact +# payload prevents a same-name but different model file from contaminating a +# benchmark comparison. +model_path="${2:-}" + +# Accept the llama.cpp executable as an optional third argument. Its checksum +# establishes the comparison binary without assuming a particular install path. +llama_binary="${3:-}" + +# Resolve this script's repository root. This makes the Git metadata capture +# independent of the shell's starting directory. +repo_root="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)" + +# Prefer an explicit virtual environment, then support FreeToken's LAN-223 +# layout where the environment is a sibling of the source checkout, and finally +# support a conventional in-repository `.venv`. Resolving this once prevents +# later runtime probes from silently using the system Python. +venv_python="${FREETOKEN_VENV:-}" +if [[ -z "$venv_python" && -x "$(dirname "$repo_root")/.venv/bin/python" ]]; then + venv_python="$(dirname "$repo_root")/.venv/bin/python" +elif [[ -z "$venv_python" && -x "$repo_root/.venv/bin/python" ]]; then + venv_python="$repo_root/.venv/bin/python" +fi + +# Create the requested artifact directory without overwriting prior captures. +mkdir -p "$output_dir" + +# Write one command's standard output and standard error to a named text file. +# The function returns success even for unavailable optional commands so the +# artifact shows the diagnostic failure instead of silently omitting it. +capture_command() { + local name="$1" + shift + { + printf '$' + printf ' %q' "$@" + printf '\n\n' + "$@" + } >"$output_dir/$name" 2>&1 || true +} + +# Record Git identity and local changes before inspecting the host. Later +# benchmark reports use these files to prove which source was executed. +capture_command git-status.txt git -C "$repo_root" status --short +capture_command git-head.txt git -C "$repo_root" rev-parse HEAD +capture_command git-branch.txt git -C "$repo_root" branch --show-current +capture_command git-remotes.txt git -C "$repo_root" remote -v + +# Record kernel, distribution, CPU, memory, and mount information. These are +# read-only inputs that can affect JIT compilation and UMA decode performance. +capture_command uname.txt uname -a +capture_command os-release.txt cat /etc/os-release +capture_command cpu.txt lscpu +capture_command memory.txt free -h +capture_command mounts.txt findmnt -D + +# Record Linux pressure-stall information before a benchmark starts. UMA +# inference shares system memory and storage paths, so I/O pressure can create +# latency outliers even when the GPU, model, and launch command are unchanged. +capture_command io-pressure.txt cat /proc/pressure/io +capture_command memory-pressure.txt cat /proc/pressure/memory + +# Record only blocked filesystem scans, not every blocked process. This keeps +# the artifact focused on a known source of benchmark interference and avoids +# collecting unrelated command-line arguments from other user applications. +capture_command blocked-find-scans.txt bash -lc "ps -eo pid,ppid,state,etimes,ni,pcpu,pmem,comm,args --sort=pid | awk 'NR == 1 || (\$3 == \"D\" && \$8 == \"find\")'" + +# Preserve recent AMDGPU and KFD warnings as read-only context. The command +# deliberately tolerates missing journal permissions and records an empty file +# when no relevant warnings occurred in the preceding two hours. +capture_command recent-amdgpu-kfd-warnings.txt bash -lc "journalctl -k --since '2 hours ago' --no-pager 2>/dev/null | grep -Ei 'amdgpu|kfd|xgmi|gpu reset|ring timeout|ras|fault' || true" + +# Record the ROCm installation selected by the shell and the compiler version. +# Resolving symlinks exposes mixed ROCm installations before profiling begins. +capture_command rocm-links.txt readlink -f /opt/rocm +capture_command rocm-tree.txt find -L /opt/rocm -maxdepth 2 -type f -name 'hipcc' -o -type l -name 'libamdhip64.so*' +capture_command hipcc-version.txt /opt/rocm/bin/hipcc --version +capture_command rocprof-version.txt /opt/rocm/bin/rocprofv3 --version +capture_command rocm-packages.txt bash -lc "dpkg-query -W -f='\${Package}\t\${Version}\n' 'rocm*' 'hip*' 'rocprofiler*' 'llvm*' 2>/dev/null | sort" + +# Record the active AMD device, dynamic power policy, and thermal state without +# attempting to change privileged DPM controls. +capture_command rocm-smi.txt rocm-smi --showproductname --showuniqueid --showmeminfo vram --showuse --showtemp --showclocks --showpower +capture_command dpm-policy.txt bash -lc "for f in /sys/class/drm/card*/device/power_dpm_force_performance_level /sys/class/drm/card*/device/pp_dpm_sclk; do printf '%s\n' \"### \$f\"; cat \"\$f\" 2>&1; done" + +# Record only performance-relevant environment names. Filtering avoids +# accidentally writing credentials or unrelated user environment variables. +capture_command performance-environment.txt bash -lc "env | LC_ALL=C sort | grep -E '^(ROCM|HIP|HSA|PYTORCH|TORCH|TRITON|LD_LIBRARY_PATH|PATH|FREETOKEN)=' || true" + +# Ask the exact FreeToken virtual environment which HIP runtime and device it +# sees. This detects a wheel whose embedded runtime differs from host ROCm. +if [[ -n "$venv_python" && -x "$venv_python" ]]; then + capture_command pytorch-runtime.txt "$venv_python" -c "import json, torch; p=torch.cuda.get_device_properties(0); print(json.dumps({'torch':torch.__version__,'hip':torch.version.hip,'cuda_available':torch.cuda.is_available(),'device':p.name,'gcnArchName':getattr(p,'gcnArchName',None),'total_memory':p.total_memory}, indent=2, sort_keys=True))" + capture_command python-ldd.txt ldd "$venv_python" +else + printf 'FreeToken virtual-environment Python not found. FREETOKEN_VENV=%s\n' "${FREETOKEN_VENV:-}" >"$output_dir/pytorch-runtime.txt" + cp "$output_dir/pytorch-runtime.txt" "$output_dir/python-ldd.txt" +fi + +# Record loaded-library resolution for the Python interpreter and rocprofv3. +# This is the primary evidence for a mixed LLVM or ROCm profiler environment. +capture_command rocprof-ldd.txt ldd /opt/rocm/bin/rocprofv3 + +# Hash optional comparison artifacts only when the caller supplied a readable +# path. The explicit messages make missing input obvious in the manifest. +if [[ -n "$model_path" && -r "$model_path" ]]; then + capture_command model-sha256.txt sha256sum "$model_path" +else + printf 'Model path not supplied or unreadable: %s\n' "$model_path" >"$output_dir/model-sha256.txt" +fi + +if [[ -n "$llama_binary" && -x "$llama_binary" ]]; then + capture_command llama-binary-sha256.txt sha256sum "$llama_binary" + capture_command llama-version.txt "$llama_binary" --version +else + printf 'llama.cpp binary not supplied or not executable: %s\n' "$llama_binary" >"$output_dir/llama-binary-sha256.txt" +fi + +# Create a deterministic inventory of every captured file and its SHA256. The +# final line is a simple completion marker for automation and human review. +(cd "$output_dir" && find . -maxdepth 1 -type f ! -name SHA256SUMS -printf '%P\0' | LC_ALL=C sort -z | xargs -0 sha256sum) >"$output_dir/SHA256SUMS" +printf 'Baseline capture complete: %s\n' "$output_dir" diff --git a/scripts/lan223-rocprof-wheel-sdk.sh b/scripts/lan223-rocprof-wheel-sdk.sh new file mode 100644 index 00000000..eab4fb0c --- /dev/null +++ b/scripts/lan223-rocprof-wheel-sdk.sh @@ -0,0 +1,46 @@ +#!/usr/bin/env bash +# Launch rocprofv3 against the ROCm SDK bundled with the active PyTorch wheel. +# +# On LAN-223, FreeToken's PyTorch ROCm wheel loads its own LLVM and +# rocprofiler-sdk libraries. Launching rocprofv3 against /opt/rocm injects a +# second copy of LLVM, which aborts during `import torch` because LLVM command +# line options are registered twice. This wrapper selects the wheel's matching +# SDK so the profiler and application load one library identity. + +# Stop on programming errors. The wrapped application exit status is preserved +# so callers can distinguish profiler setup failures from application failures. +set -euo pipefail + +# Require the application separator used by rocprofv3. Keeping profiler flags +# before `--` makes arbitrary HIP applications usable without hard-coding a +# FreeToken server command in this helper. +if [[ "$#" -lt 1 ]]; then + printf 'Usage: %s [rocprofv3 options] -- application [arguments...]\n' "$0" >&2 + exit 64 +fi + +# Prefer an explicit virtual environment and otherwise use the LAN-223 layout +# where `.venv` is adjacent to the source checkout that contains this script. +repo_root="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)" +venv_root="${FREETOKEN_VENV_ROOT:-$(dirname "$repo_root")/.venv}" + +# Locate the wheel-owned ROCm SDK rather than assuming a Python minor version. +# The glob is validated to prevent a shell literal from being passed to rocprof. +sdk_candidates=("$venv_root"/lib/python*/site-packages/_rocm_sdk_core) +if [[ ! -d "${sdk_candidates[0]}" ]]; then + printf 'Cannot find PyTorch wheel ROCm SDK under %s. Set FREETOKEN_VENV_ROOT.\n' "$venv_root" >&2 + exit 66 +fi +sdk_root="${sdk_candidates[0]}" + +# Verify the two libraries needed by rocprofv3 exist in the selected SDK. This +# catches an incomplete or non-ROCm PyTorch wheel before it starts an app. +if [[ ! -r "$sdk_root/lib/librocprofiler-sdk.so.1" || ! -r "$sdk_root/lib/rocprofiler-sdk/librocprofiler-sdk-tool.so.1" ]]; then + printf 'The selected SDK lacks rocprofiler-sdk 1.3 components: %s\n' "$sdk_root" >&2 + exit 66 +fi + +# Use the host's rocprofv3 front end but direct every profiler library lookup to +# the exact SDK already used by PyTorch. Do not set LD_PRELOAD here: rocprofv3 +# owns its preload order and forwards the selected tool to the child process. +exec /opt/rocm/bin/rocprofv3 --rocm-root "$sdk_root" "$@" diff --git a/setup.py b/setup.py index cfe41b7d..faf9fc3f 100644 --- a/setup.py +++ b/setup.py @@ -4,13 +4,18 @@ from pathlib import Path from setuptools import setup -from torch.utils.cpp_extension import BuildExtension, CUDA_HOME, CppExtension +from torch.utils.cpp_extension import BuildExtension, CUDA_HOME, ROCM_HOME, CppExtension ROOT = Path(__file__).parent +IS_ROCM = CUDA_HOME is None and ROCM_HOME is not None def _check_toolchain() -> None: + if IS_ROCM: + # nvcc/CUDA-major checks below are meaningless on a ROCm torch build + # (torch.version.cuda is None there), so _toolchain.py's check is a no-op. + return path = ROOT / "python" / "freetoken" / "kernel" / "_toolchain.py" spec = importlib.util.spec_from_file_location("_freetoken_toolchain", path) module = importlib.util.module_from_spec(spec) @@ -18,20 +23,39 @@ def _check_toolchain() -> None: module.check_nvcc_matches_torch() -def _cuda_runtime_paths() -> tuple[list[str], list[str]]: +def _gpu_runtime_paths() -> tuple[list[str], list[str], list[str], list[str]]: + """Returns (include_dirs, library_dirs, libraries, extra_link_args).""" + if IS_ROCM: + rocm_home = Path(ROCM_HOME) + library_dirs = [d for d in (rocm_home / "lib64", rocm_home / "lib") if d.exists()] + # The pip-vendored rocm-sdk-core ships versioned sonames (libamdhip64.so.7) + # without the bare .so dev symlink `-lamdhip64` needs, so link the exact + # file. At runtime the dynamic linker dedupes on SONAME, so this resolves + # to whichever libamdhip64 torch itself already loaded into the process. + hip_lib = next( + (f for d in library_dirs for f in sorted(d.glob("libamdhip64.so*"))), None + ) + if hip_lib is None: + raise RuntimeError(f"libamdhip64.so* not found under {library_dirs}") + return ( + [str(rocm_home / "include")], + [str(d) for d in library_dirs], + [], + [f"-l:{hip_lib.name}"], + ) if CUDA_HOME is None: raise RuntimeError( - "CUDA_HOME is required to build freetoken.kernel._pinned_tensor " - "because it links against the CUDA runtime API." + "CUDA_HOME (or ROCM_HOME) is required to build freetoken.kernel._pinned_tensor " + "because it links against the CUDA/HIP runtime API." ) cuda_home = Path(CUDA_HOME) library_dirs = [str(cuda_home / "lib64")] if (cuda_home / "lib").exists(): library_dirs.append(str(cuda_home / "lib")) - return [str(cuda_home / "include")], library_dirs + return [str(cuda_home / "include")], library_dirs, ["cudart"], [] -cuda_include_dirs, cuda_library_dirs = _cuda_runtime_paths() +cuda_include_dirs, cuda_library_dirs, cuda_libraries, cuda_extra_link_args = _gpu_runtime_paths() _check_toolchain() @@ -44,7 +68,8 @@ def _cuda_runtime_paths() -> tuple[list[str], list[str]]: ], include_dirs=cuda_include_dirs, library_dirs=cuda_library_dirs, - libraries=["cudart"], + libraries=cuda_libraries, + extra_link_args=cuda_extra_link_args, extra_compile_args=["-O3", "-std=c++17"], ), # CPU-compute MoE executor for --moe-backend cpu. Links cudart for the @@ -59,7 +84,8 @@ def _cuda_runtime_paths() -> tuple[list[str], list[str]]: ], include_dirs=cuda_include_dirs, library_dirs=cuda_library_dirs, - libraries=["cudart"], + libraries=cuda_libraries, + extra_link_args=cuda_extra_link_args, extra_compile_args=["-O3", "-std=c++17", "-pthread"], ), ], diff --git a/tests/kernels/test_gguf_hip_build_flags.py b/tests/kernels/test_gguf_hip_build_flags.py new file mode 100644 index 00000000..942ec9e6 --- /dev/null +++ b/tests/kernels/test_gguf_hip_build_flags.py @@ -0,0 +1,35 @@ +"""Unit coverage for the HIP-only GGUF extension compiler configuration. + +These tests do not invoke hipcc. They verify the environment that is prepared +before PyTorch's extension builder computes its target-specific cache key. +""" + +import os +from types import SimpleNamespace + +from freetoken.kernel import gguf + + +def test_hip_gguf_flags_pin_the_active_gfx_target(monkeypatch): + """A single-GPU HIP process derives gfx1151 when no target was configured.""" + monkeypatch.delenv("PYTORCH_ROCM_ARCH", raising=False) + monkeypatch.delenv("FREETOKEN_HIP_GGUF_FAST_MATH", raising=False) + monkeypatch.setattr(gguf.torch.cuda, "is_available", lambda: True) + monkeypatch.setattr( + gguf.torch.cuda, + "get_device_properties", + lambda _index: SimpleNamespace(gcnArchName="gfx1151:sramecc-:xnack-"), + ) + + assert gguf._hip_gguf_cflags() == ["-O3"] + assert gguf._hip_target_arch() == "gfx1151" + assert os.environ["PYTORCH_ROCM_ARCH"] == "gfx1151" + + +def test_hip_gguf_flags_preserve_an_explicit_multi_target_choice(monkeypatch): + """An explicit multi-target deployment choice is never replaced by auto-detection.""" + monkeypatch.setenv("PYTORCH_ROCM_ARCH", "gfx1100;gfx1151") + + assert gguf._hip_gguf_cflags() == ["-O3"] + assert gguf._hip_target_arch() == "gfx1100" + assert os.environ["PYTORCH_ROCM_ARCH"] == "gfx1100;gfx1151" diff --git a/tests/kernels/test_pinned_tensor.py b/tests/kernels/test_pinned_tensor.py index e61108fd..eaf3b338 100644 --- a/tests/kernels/test_pinned_tensor.py +++ b/tests/kernels/test_pinned_tensor.py @@ -122,7 +122,11 @@ def test_host_device_ptr_is_identity_under_uva(): pytest.skip("non-UVA platform: host_device_ptr rejects unregistered memory instead") # Under UVA cudaHostGetDevicePointer degenerates to identity for any host pointer # (no registration validation); rejection of pageable memory only exists on - # non-identity platforms (Windows/WDDM), where the translation is real. + # non-identity CUDA platforms (Windows/WDDM), where the translation is real. + # HIP validates registration even when registered memory has an identity + # address on Linux. The pinned identity check above is the relevant test. + if torch.version.hip is not None: + return pageable = torch.empty(64, dtype=torch.uint8) ext = _load_pinned_extension() assert ext.host_device_ptr(pageable.data_ptr()) == pageable.data_ptr() diff --git a/tests/moe/test_fused_moe.py b/tests/moe/test_fused_moe.py index 41a75c20..cd34fc21 100644 --- a/tests/moe/test_fused_moe.py +++ b/tests/moe/test_fused_moe.py @@ -2,6 +2,34 @@ import torch +def test_fused_topk_keeps_reference_router_on_rocm(monkeypatch): + """HIP keeps the exact PyTorch router until a Triton path passes API parity.""" + from freetoken.kernel import backend + from freetoken.moe import fused + + weights = torch.tensor([[0.7, 0.3]], dtype=torch.float32) + ids = torch.tensor([[4, 9]], dtype=torch.int32) + calls = [] + + monkeypatch.setattr(backend, "is_rocm_runtime", lambda: True) + monkeypatch.setattr( + fused, + "_torch_fused_topk", + lambda logits, topk, renormalize, limit: ( + calls.append((logits, topk, renormalize, limit)) or (weights, ids) + ), + ) + + got_weights, got_ids = fused.fused_topk( + torch.empty((1, 3)), torch.empty((1, 16)), topk=2, renormalize=True + ) + + assert len(calls) == 1 + assert calls[0][1:] == (2, True, None) + assert got_weights is weights + assert got_ids is ids + + def _activation_and_mul(gate_up: torch.Tensor, activation: str) -> torch.Tensor: gate, up = gate_up.chunk(2, dim=-1) if activation == "silu": diff --git a/tests/utils/test_rocm_runtime.py b/tests/utils/test_rocm_runtime.py new file mode 100644 index 00000000..13021dfa --- /dev/null +++ b/tests/utils/test_rocm_runtime.py @@ -0,0 +1,54 @@ +"""Regression coverage for the CUDA-namespace compatibility boundary on ROCm. + +PyTorch exposes AMD devices through ``torch.cuda`` so CUDA-oriented Python +programs can run on HIP. FreeToken must not mistake a ``gfx11xx`` capability +for a newer NVIDIA SM capability, nor select optional CUDA binaries merely +because a stale package happens to be installed in the environment. +""" + +import torch + +from freetoken.kernel import backend +from freetoken.utils import arch + + +def test_rocm_never_satisfies_nvidia_sm_gates(monkeypatch): + """HIP hardware is excluded before numerical NVIDIA capability comparison.""" + monkeypatch.setattr(torch.version, "hip", "7.15") + monkeypatch.setattr(torch.version, "cuda", None) + monkeypatch.setattr(torch.cuda, "is_available", lambda: True) + monkeypatch.setattr(torch.cuda, "get_device_capability", lambda: (11, 5)) + arch.is_rocm_runtime.cache_clear() + arch._get_torch_cuda_version.cache_clear() + + try: + assert arch.is_rocm_runtime() is True + assert arch._get_torch_cuda_version() is None + assert arch.is_sm90_supported() is False + assert arch.is_sm100_supported() is False + finally: + # Cached runtime detection must not leak the synthetic HIP state into + # unrelated test modules that run later in the same interpreter. + arch.is_rocm_runtime.cache_clear() + arch._get_torch_cuda_version.cache_clear() + + +def test_rocm_disables_cuda_only_optional_backends(monkeypatch): + """Triton remains available, while CUDA binary packages are bypassed on HIP.""" + monkeypatch.setattr(backend, "is_rocm_runtime", lambda: True) + monkeypatch.setattr(backend, "_importable", lambda _name: True) + backend.is_flashinfer_installed.cache_clear() + backend.is_sgl_kernel_installed.cache_clear() + backend.is_triton_kernels_installed.cache_clear() + backend.driver_cuda_version.cache_clear() + + try: + assert backend.is_flashinfer_installed() is False + assert backend.is_sgl_kernel_installed() is False + assert backend.is_triton_kernels_installed() is False + assert backend.driver_cuda_version() is None + finally: + backend.is_flashinfer_installed.cache_clear() + backend.is_sgl_kernel_installed.cache_clear() + backend.is_triton_kernels_installed.cache_clear() + backend.driver_cuda_version.cache_clear()