From 392748244aea3c23f4f59f6a6a9bdb56b4e6fa79 Mon Sep 17 00:00:00 2001 From: skywalk1411 <61213518+skywalk1411@users.noreply.github.com> Date: Thu, 27 Aug 2026 15:10:47 -0400 Subject: [PATCH 01/72] rocm: fix native extensions, kernel JIT builds, and Triton PTX fallbacks for gfx1150 FreeToken built and ran only against CUDA. On a native ROCm install (tested on a Ryzen AI 9 HX 470 / Radeon 890M, gfx1150) it failed at every stage: build, JIT compile, and finally a silent native crash mid-warmup with no Python traceback. Build system (setup.py, hip_compat.h): - _pinned_tensor and _cpu_moe link against HIP instead of cudart when ROCM_HOME is present (CUDA_HOME stays required on the CUDA path). - kernel/csrc/hip_compat.h aliases the CUDA Runtime API calls those two files use onto their HIP equivalents, including CUDART_CB (undefined under HIP, which otherwise corrupts the surrounding declaration's parse). - The pip-vendored ROCm SDK ships versioned sonames (libamdhip64.so.7) with no bare .so dev symlink, so link the exact file via -l:; the dynamic linker dedupes by SONAME at runtime against whatever libamdhip64 torch itself already loaded. Shared kernel header (kernel/csrc/include/freetoken/utils.cuh): - Explicit HIP runtime include + CUDA Runtime API aliases (nvcc pulls cuda_runtime.h in implicitly for .cu files; hipcc does not). - __grid_constant__ has no HIP equivalent; falls back to an ordinary by-value kernel parameter. - LaunchKernel has no HIP path for cudaLaunchKernelEx/cudaLaunchConfig_t (that API only exists to carry Hopper PDL attributes) -- added a HIP variant that launches via plain triple-chevron syntax instead, with with_attr() as a no-op since there is no attribute to carry. - The griddepcontrol PDL asm is now unconditionally a no-op under HIP, not just when kUsePDL is false, so a stray HIP-side call can't try to assemble Hopper-only PTX. Kernel JIT (kernel/utils.py): - Drop --expt-relaxed-constexpr on HIP; hipcc/clang rejects it outright. Triton kernels: - norm.py, activation.py: launch_pdl is a CUDA-Hopper-only kwarg; the AMD arg-packer raises KeyError on it even when passed as False, so it's only included when pdl is actually true (never on ROCm). - attention.py: the decode kernel's GQA head-tile floors at 16 under HIP (RDNA WMMA has no instruction below M=16); the kernel already masks padded head lanes for non-power-of-two groups, so this is a safe widening. Falls back to broadcast-multiply-reduce instead of tl.dot for that tile as a second-layer guard. The split extend/prefill kernel's tile shrinks from 128x64 to 64x32 under HIP -- running both the cached- and newly-computed-KV loops live at once is register-heavier than the plain extend kernel, and exhausts this GPU's VGPR file at the CUDA-tuned tile size. - activation.py (the actual root cause of the crash above): _fast_tanh and _fast_ex2 inline raw PTX text (tanh.approx.f32, ex2.approx.f32) via tl.inline_asm_elementwise. HIP's inline-asm path doesn't reject foreign PTX at parse time -- it fails much later in register allocation with a generic, misleading diagnostic ("couldn't allocate output register for constraint 'f'") that looks like a matrix-core or register-pressure issue and sent debugging down that path for a while. Routed through libdevice.tanh / tl.exp2 on HIP instead. pyproject.toml: loosen the torch/triton ceilings so ROCm builds (which carry a local version segment such as +rocm7.14.0...) can satisfy them. Every change is gated on HIP detection (torch.version.hip / ROCM_HOME / __HIP_PLATFORM_AMD__) at build or run time; the CUDA path is unchanged. Verified end to end on gfx1150: server boot, weight load, KV cache alloc, CUDA graph capture at bs=1/2/4, and real chat completions against Qwen/Qwen3-8B (bf16, triton attention backend). --- pyproject.toml | 9 +- .../kernel/csrc/cpu_moe/cpu_moe_ext.cpp | 2 +- python/freetoken/kernel/csrc/hip_compat.h | 55 ++++++++++++ .../kernel/csrc/include/freetoken/utils.cuh | 84 +++++++++++++++++++ .../freetoken/kernel/csrc/pinned_tensor.cpp | 2 +- python/freetoken/kernel/triton/activation.py | 17 +++- python/freetoken/kernel/triton/attention.py | 37 +++++++- python/freetoken/kernel/triton/norm.py | 9 +- python/freetoken/kernel/utils.py | 8 +- setup.py | 42 ++++++++-- 10 files changed, 246 insertions(+), 19 deletions(-) create mode 100644 python/freetoken/kernel/csrc/hip_compat.h diff --git a/pyproject.toml b/pyproject.toml index 8bd653f8..d7de9c15 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] @@ -54,10 +54,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/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/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..c24799fe 100644 --- a/python/freetoken/kernel/csrc/include/freetoken/utils.cuh +++ b/python/freetoken/kernel/csrc/include/freetoken/utils.cuh @@ -10,6 +10,36 @@ #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; + +// 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 +72,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 +126,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 +223,6 @@ private: cudaLaunchAttribute m_attr_cache; }; +#endif + } // namespace host 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/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..d6710811 100644 --- a/python/freetoken/kernel/triton/attention.py +++ b/python/freetoken/kernel/triton/attention.py @@ -34,6 +34,13 @@ def fits(block_m: int, block_n: int) -> bool: return (block_m + 2 * block_n) * block_d * 2 <= budget if head_dim <= 128: + # The split extend/prefill kernel (separate cached + newly-computed KV loops + # live at once) exhausts this GPU's VGPR file at the 128x64 tile -- register + # pressure, not shared memory, is the binding constraint here, and this + # function was only ever tuned against the latter. Shrink unconditionally on + # ROCm rather than trying to model AMD's register budget per-arch. + if torch.version.hip is not None: + return 64, 32 return 128, 64 if head_dim <= 256: return (128, 64) if fits(128, 64) else (64, 32) @@ -179,6 +186,7 @@ def _decode_grouped_stage1_kernel( D: tl.constexpr, DV: tl.constexpr, SLIDING_WINDOW: tl.constexpr, + USE_TL_DOT: tl.constexpr, ): batch_id = tl.program_id(0) head_block_id = tl.program_id(1) @@ -237,7 +245,18 @@ def _decode_grouped_stage1_kernel( mask=mask_n[None, :] & mask_d[:, None], other=0.0, ) - scores = tl.dot(q, k) * sm_scale + if USE_TL_DOT: + scores = tl.dot(q, k) * sm_scale + else: + # RDNA WMMA has no matrix-core instruction for M < 16, which this + # kernel's decode head-tile (BLOCK_H, often 4-8 real GQA heads + # padded to 16) hits reliably; the AMD Triton/LLVM backend fails + # instruction selection rather than falling back on its own. Sum + # of broadcast products is the plain-arithmetic equivalent of + # tl.dot -- slower, but sidesteps matrix-core lowering entirely. + scores = tl.sum( + q.to(tl.float32)[:, :, None] * k.to(tl.float32)[None, :, :], axis=1 + ) * sm_scale scores = tl.where(mask_h[:, None] & mask_n[None, :], scores, -float("inf")) v = tl.load( @@ -249,7 +268,13 @@ def _decode_grouped_stage1_kernel( m_new = tl.maximum(tl.max(scores, axis=1), m_i) alpha = tl.exp(m_i - m_new) p = tl.exp(scores - m_new[:, None]) - acc = acc * alpha[:, None] + tl.dot(p.to(v.dtype), v) + if USE_TL_DOT: + pv = tl.dot(p.to(v.dtype), v) + else: + pv = tl.sum( + p.to(tl.float32)[:, :, None] * v.to(tl.float32)[None, :, :], axis=1 + ) + acc = acc * alpha[:, None] + pv l_i = l_i * alpha + tl.sum(p, axis=1) m_i = m_new @@ -396,6 +421,13 @@ 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 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) @@ -435,6 +467,7 @@ def decode_paged_attention( D=head_dim, DV=head_dim, SLIDING_WINDOW=sliding_window or 0, + USE_TL_DOT=torch.version.hip is None, num_warps=4, num_stages=2, ) 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/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"], ), ], From 2614973fd5cf0f02bbf37075879bbc003a90eee4 Mon Sep 17 00:00:00 2001 From: skywalk1411 <61213518+skywalk1411@users.noreply.github.com> Date: Thu, 27 Aug 2026 16:02:38 -0400 Subject: [PATCH 02/72] rocm: fix fp8 native-capability detection and the offload cache's fast-copy kernel Two more real bugs found while running an actual MoE model (Qwen3.6-35B-A3B-FP8, --moe-backend offload) end to end on gfx1150, past what the first commit covered. e4m3_compat.py: e4m3_native() decides whether kernels get raw fp8 tensors or a uint8 view by checking torch.cuda.get_device_capability() >= (8, 9). On a HIP build that call returns the GPU's RDNA generation number, not a CUDA compute capability -- gfx1150 reports (11, 5), and (11, 5) >= (8, 9) is True by plain tuple comparison (11 > 8), so this incorrectly claimed native fp8 support on AMD. Triton's own compile-time twin, e4m3_native_cx() (target_info. cuda_capability_geq, which checks target.backend != "cuda" first), correctly said False, so the kernel compiled for the emulated uint8 path while the host side hands it an untouched fp8 tensor -- IncompatibleTypeErrorImpl inside e4m3_u8_to_f32's bitwise ops. Fixed by checking torch.version.hip first. fast_index_copy.cuh (the offload cache's fast host->device expert-copy kernel, only exercised once a real MoE model with --moe-backend offload actually streams experts): same two problems as the first commit's fixes elsewhere in this file family, just not caught until this path actually ran. - Missing HIP aliases for cudaGetDevice/cudaDeviceGetAttribute/ cudaHostGetDevicePointer/the two cudaDevAttr* constants it uses -- added to utils.cuh's existing HIP block alongside the ones from the first commit. - load_nc/store_nc inline raw PTX (ld.global.L1::no_allocate, st.global.wt -- cache-policy hints, no HIP equivalent). Falls back to plain loads/stores under HIP; correctness unchanged, only the cache hint is lost. Verified: Qwen3.6-35B-A3B-FP8 (256 experts/layer x 40 layers, 3B active) boots and serves real chat completions with --moe-backend offload --moe-cache-size 2560 (25% of the model's 10240 total experts resident, LRU-evicting the rest from host RAM on every miss) -- ft ctl cache confirms the pool is live at the requested size, not silently falling back to full residency. --- .../kernel/csrc/include/freetoken/utils.cuh | 14 ++++++++ .../kernel/csrc/jit/fast_index_copy.cuh | 34 +++++++++++++++++++ python/freetoken/kernel/triton/e4m3_compat.py | 8 +++++ 3 files changed, 56 insertions(+) diff --git a/python/freetoken/kernel/csrc/include/freetoken/utils.cuh b/python/freetoken/kernel/csrc/include/freetoken/utils.cuh index c24799fe..ea472b63 100644 --- a/python/freetoken/kernel/csrc/include/freetoken/utils.cuh +++ b/python/freetoken/kernel/csrc/include/freetoken/utils.cuh @@ -32,6 +32,20 @@ inline hipError_t cudaFuncSetAttribute(const void *func, hipFuncAttribute attr, 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. diff --git a/python/freetoken/kernel/csrc/jit/fast_index_copy.cuh b/python/freetoken/kernel/csrc/jit/fast_index_copy.cuh index bb83c23e..649c12dd 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)); 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 From fa75233d56baf78c8452531b0e0210c1793acdf3 Mon Sep 17 00:00:00 2001 From: skywalk1411 <61213518+skywalk1411@users.noreply.github.com> Date: Thu, 27 Aug 2026 17:32:47 -0400 Subject: [PATCH 03/72] rocm: fix the GGUF kernel JIT build (kernel/gguf.py, dispatch.h) Third loading path verified: google/gemma-4-26B-A4B-it-qat-q4_0-gguf (native GGUF, MoE offload) now boots and serves on gfx1150, alongside the dense bf16 and FP8 MoE paths from the earlier commits. kernel/gguf.py: same nvcc-only-flag problem as elsewhere in this port, in a third JIT mechanism (torch.utils.cpp_extension.load, distinct from both setup.py's CppExtension and the tvm-ffi JIT the rest of kernel/ uses). --expt-relaxed-constexpr is rejected outright, and the -ccbin/CXX-forcing block exists only to work around an nvcc+libtorch-headers compiler mismatch that doesn't apply under hipcc (its own bundled clang already is the host compiler). Both dropped on HIP. kernel/csrc/gguf/dispatch.h: the donor's SGLANG_SHFL_XOR_SYNC(_WIDTH) macros forward a CUDA-style 32-bit mask straight into __shfl_xor_sync. HIP's amd_warp_sync_functions.h static_asserts the mask must be 64 bits unconditionally (regardless of actual wavefront width) -- widened the cast on HIP only. .gitignore: torch's ROCm auto-hipify (a real, working translation pass built into torch.utils.cpp_extension -- unlike the other two JIT paths, this one needed no manual porting for the .cu/.cuh sources themselves) writes translated copies next to the CUDA sources it processes (gguf_kernel.cu -> .hip, *.cuh -> *_hip.cuh). Regenerated every build, never hand-edited; ignore rather than track. Not in this commit, environment-only: the pip ROCm nightly distribution used here (rocm.nightlies.amd.com) ships no thrust/rocprim headers, which torch's own extension headers pull in transitively. Ubuntu's librocthrust-dev is one fix, but it depends on libamdhip64-dev, which drops a second, conflicting HIP header set into /usr/include/hip that silently wins over the correct pip-bundled ones for any plain -I (though not -isystem) -- diagnosed by hand with `clang++ -v` and a minimal reproducer. Worked around locally by extracting just the thrust/rocprim headers (dpkg -x, no install) into the pip package's own include dir and removing the conflicting system packages; ROCM_PATH/HIP_PATH/HIP_DEVICE_LIB_PATH also had to point at the pip package for this JIT path's device-bitcode-library lookup. Left out of the diff since there's no source change to make -- noting it here for the next person on this distribution. --- .gitignore | 7 ++++++ python/freetoken/kernel/csrc/gguf/dispatch.h | 15 +++++++++++ python/freetoken/kernel/gguf.py | 26 ++++++++++++-------- 3 files changed, 38 insertions(+), 10 deletions(-) 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/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/gguf.py b/python/freetoken/kernel/gguf.py index 04a16560..dbe392c5 100644 --- a/python/freetoken/kernel/gguf.py +++ b/python/freetoken/kernel/gguf.py @@ -51,16 +51,22 @@ 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 = ["-O3"] + 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) # 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. From fd2b287d0d7efd539f74b04c60df01270612f3a6 Mon Sep 17 00:00:00 2001 From: skywalk1411 <61213518+skywalk1411@users.noreply.github.com> Date: Thu, 27 Aug 2026 17:47:52 -0400 Subject: [PATCH 04/72] rocm: revert two defensive fixes that turned out to be unnecessary Both were made mid-investigation, before the real cause of a since-fixed crash (the raw-PTX bug in activation.py, and separately the e4m3_native() tuple- comparison bug) was actually found. Re-tested each in isolation -- eager, batched, and inside real CUDA graph capture+replay -- now that those are fixed, and both work fine at the original, CUDA-tuned settings: - decode_paged_attention: the block_h>=16 floor (kept -- RDNA WMMA genuinely has no instruction below M=16, confirmed independently and matches upstream #137) was sufficient on its own. The USE_TL_DOT broadcast-sum fallback this PR had added on top was solving a problem that was actually in a different kernel; removed, restoring real matrix-core-accelerated decode attention. - _select_extend_tile: the 128x64 -> 64x32 shrink on HIP was diagnosed as a VGPR-exhaustion issue via a py-spy trace mid-investigation, before the session had isolated the actual crash to activation.py. Re-verified end-to-end against Qwen3.6-35B-A3B-FP8's GDN/split-extend path (the kernel this shrink targeted) at the original tile size: no crash, correct output. Reverted to the CUDA-tuned tile. Both re-verified against real chat completions (Qwen3-8B for the decode path, Qwen3.6-35B-A3B-FP8 for the extend/split path) after reverting, not just the isolated kernel tests. --- python/freetoken/kernel/triton/attention.py | 30 ++------------------- 1 file changed, 2 insertions(+), 28 deletions(-) diff --git a/python/freetoken/kernel/triton/attention.py b/python/freetoken/kernel/triton/attention.py index d6710811..bc1fd11a 100644 --- a/python/freetoken/kernel/triton/attention.py +++ b/python/freetoken/kernel/triton/attention.py @@ -34,13 +34,6 @@ def fits(block_m: int, block_n: int) -> bool: return (block_m + 2 * block_n) * block_d * 2 <= budget if head_dim <= 128: - # The split extend/prefill kernel (separate cached + newly-computed KV loops - # live at once) exhausts this GPU's VGPR file at the 128x64 tile -- register - # pressure, not shared memory, is the binding constraint here, and this - # function was only ever tuned against the latter. Shrink unconditionally on - # ROCm rather than trying to model AMD's register budget per-arch. - if torch.version.hip is not None: - return 64, 32 return 128, 64 if head_dim <= 256: return (128, 64) if fits(128, 64) else (64, 32) @@ -186,7 +179,6 @@ def _decode_grouped_stage1_kernel( D: tl.constexpr, DV: tl.constexpr, SLIDING_WINDOW: tl.constexpr, - USE_TL_DOT: tl.constexpr, ): batch_id = tl.program_id(0) head_block_id = tl.program_id(1) @@ -245,18 +237,7 @@ def _decode_grouped_stage1_kernel( mask=mask_n[None, :] & mask_d[:, None], other=0.0, ) - if USE_TL_DOT: - scores = tl.dot(q, k) * sm_scale - else: - # RDNA WMMA has no matrix-core instruction for M < 16, which this - # kernel's decode head-tile (BLOCK_H, often 4-8 real GQA heads - # padded to 16) hits reliably; the AMD Triton/LLVM backend fails - # instruction selection rather than falling back on its own. Sum - # of broadcast products is the plain-arithmetic equivalent of - # tl.dot -- slower, but sidesteps matrix-core lowering entirely. - scores = tl.sum( - q.to(tl.float32)[:, :, None] * k.to(tl.float32)[None, :, :], axis=1 - ) * sm_scale + scores = tl.dot(q, k) * sm_scale scores = tl.where(mask_h[:, None] & mask_n[None, :], scores, -float("inf")) v = tl.load( @@ -268,13 +249,7 @@ def _decode_grouped_stage1_kernel( m_new = tl.maximum(tl.max(scores, axis=1), m_i) alpha = tl.exp(m_i - m_new) p = tl.exp(scores - m_new[:, None]) - if USE_TL_DOT: - pv = tl.dot(p.to(v.dtype), v) - else: - pv = tl.sum( - p.to(tl.float32)[:, :, None] * v.to(tl.float32)[None, :, :], axis=1 - ) - acc = acc * alpha[:, None] + pv + acc = acc * alpha[:, None] + tl.dot(p.to(v.dtype), v) l_i = l_i * alpha + tl.sum(p, axis=1) m_i = m_new @@ -467,7 +442,6 @@ def decode_paged_attention( D=head_dim, DV=head_dim, SLIDING_WINDOW=sliding_window or 0, - USE_TL_DOT=torch.version.hip is None, num_warps=4, num_stages=2, ) From b6593ccd8695368bad2228b35db40bf62f287bd1 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 10:49:35 -0700 Subject: [PATCH 05/72] feat(rocm): harden HIP runtime gating for gfx1151 --- docs/amd-rocm-gfx1151.md | 107 +++++++++++++++++++++++++++++ pyproject.toml | 1 + python/freetoken/kernel/backend.py | 22 +++++- python/freetoken/utils/__init__.py | 2 + python/freetoken/utils/arch.py | 20 +++++- tests/utils/test_rocm_runtime.py | 54 +++++++++++++++ 6 files changed, 202 insertions(+), 4 deletions(-) create mode 100644 docs/amd-rocm-gfx1151.md create mode 100644 tests/utils/test_rocm_runtime.py diff --git a/docs/amd-rocm-gfx1151.md b/docs/amd-rocm-gfx1151.md new file mode 100644 index 00000000..f3e10c88 --- /dev/null +++ b/docs/amd-rocm-gfx1151.md @@ -0,0 +1,107 @@ +# 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. + +## 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. diff --git a/pyproject.toml b/pyproject.toml index d7de9c15..a7276fd8 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -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", 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/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/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() From b45e46144203390f2be97693beba59fe34a521b4 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 10:56:47 -0700 Subject: [PATCH 06/72] fix(rocm): accept HIP tensors in TVM JIT kernels --- python/freetoken/kernel/csrc/jit/index.cu | 6 +++--- python/freetoken/kernel/csrc/jit/store.cu | 6 +++--- tests/kernels/test_pinned_tensor.py | 6 +++++- 3 files changed, 11 insertions(+), 7 deletions(-) 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/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() From 25bf1c777e27421139836cecea21c19343721107 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 10:58:07 -0700 Subject: [PATCH 07/72] fix(rocm): support HIP fast-index copy tensors --- python/freetoken/kernel/csrc/jit/fast_index_copy.cuh | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/python/freetoken/kernel/csrc/jit/fast_index_copy.cuh b/python/freetoken/kernel/csrc/jit/fast_index_copy.cuh index 649c12dd..2d1dbc05 100644 --- a/python/freetoken/kernel/csrc/jit/fast_index_copy.cuh +++ b/python/freetoken/kernel/csrc/jit/fast_index_copy.cuh @@ -378,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); @@ -397,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()); From 34ab367e08666a30180dc249cf2a4289ab44854f Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 11:44:27 -0700 Subject: [PATCH 08/72] fix(rocm): skip unsafe optional NVFP4 prefill warmup --- python/freetoken/engine/engine.py | 17 ++++++++++++++++- 1 file changed, 16 insertions(+), 1 deletion(-) 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: From 5ab7e48d072b0ac64101ad31428497f21a9da40b Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 11:52:55 -0700 Subject: [PATCH 09/72] fix(rocm): use safe Triton NVFP4 prefill path --- python/freetoken/moe/fused_nvfp4.py | 23 +++++++++++++++++++++++ 1 file changed, 23 insertions(+) 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] From a482d395af807de8defec490ad01392c651a99a7 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 12:08:07 -0700 Subject: [PATCH 10/72] fix(rocm): locate Thrust headers for GGUF JIT --- python/freetoken/kernel/gguf.py | 34 ++++++++++++++++++++++++++++++++- 1 file changed, 33 insertions(+), 1 deletion(-) diff --git a/python/freetoken/kernel/gguf.py b/python/freetoken/kernel/gguf.py index dbe392c5..f6a5a12e 100644 --- a/python/freetoken/kernel/gguf.py +++ b/python/freetoken/kernel/gguf.py @@ -22,6 +22,30 @@ _CSRC = pathlib.Path(__file__).parent / "csrc" / "gguf" +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 _host_compiler() -> str | None: """A host compiler nvcc + libtorch headers accept. @@ -56,6 +80,13 @@ def _module(): # 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 = ["-O3"] + # 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. + hip_thrust_include = _hip_thrust_include() + extra_include_paths = [str(_CSRC)] + if hip_thrust_include is not None: + extra_include_paths.append(hip_thrust_include) else: extra_cuda_cflags = ["-O3", "--expt-relaxed-constexpr"] host_cxx = _host_compiler() @@ -67,13 +98,14 @@ def _module(): extra_cuda_cflags += ["-ccbin", cxx_path] os.environ["CXX"] = cxx_path os.environ["CC"] = _c_compiler_for(cxx_path) + extra_include_paths = [str(_CSRC)] # 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, verbose=True, ) From e9b1b67848f78a6ac4cdca36d980edb9d507490c Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 12:11:01 -0700 Subject: [PATCH 11/72] fix(rocm): pass GGUF Thrust headers as system include --- python/freetoken/kernel/gguf.py | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/python/freetoken/kernel/gguf.py b/python/freetoken/kernel/gguf.py index f6a5a12e..9a931fb4 100644 --- a/python/freetoken/kernel/gguf.py +++ b/python/freetoken/kernel/gguf.py @@ -83,10 +83,13 @@ def _module(): # 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() extra_include_paths = [str(_CSRC)] if hip_thrust_include is not None: - extra_include_paths.append(hip_thrust_include) + extra_cuda_cflags += ["-isystem", hip_thrust_include] else: extra_cuda_cflags = ["-O3", "--expt-relaxed-constexpr"] host_cxx = _host_compiler() From 6a4f7b4ce976c76c3825be89194bf243a3a9493c Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 12:17:05 -0700 Subject: [PATCH 12/72] fix(rocm): locate HIP runtime for GGUF JIT --- python/freetoken/kernel/gguf.py | 33 +++++++++++++++++++++++++++++++++ 1 file changed, 33 insertions(+) diff --git a/python/freetoken/kernel/gguf.py b/python/freetoken/kernel/gguf.py index 9a931fb4..d4a88c05 100644 --- a/python/freetoken/kernel/gguf.py +++ b/python/freetoken/kernel/gguf.py @@ -46,6 +46,30 @@ def _hip_thrust_include() -> str | None: 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. @@ -87,9 +111,16 @@ def _module(): # 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() @@ -102,6 +133,7 @@ def _module(): 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. @@ -110,6 +142,7 @@ def _module(): sources=[str(_CSRC / "gguf_kernel.cu")], extra_include_paths=extra_include_paths, extra_cuda_cflags=extra_cuda_cflags, + extra_ldflags=extra_ldflags, verbose=True, ) From 26501dbe33ad110a7f86bb8363f490f42b438cae Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 12:25:35 -0700 Subject: [PATCH 13/72] docs(rocm): record LAN-223 full model validation --- docs/amd-rocm-gfx1151.md | 4 + docs/lan223-rocm-validation-2026-08-28.md | 115 ++++++++++++++++++++++ 2 files changed, 119 insertions(+) create mode 100644 docs/lan223-rocm-validation-2026-08-28.md diff --git a/docs/amd-rocm-gfx1151.md b/docs/amd-rocm-gfx1151.md index f3e10c88..0f699c91 100644 --- a/docs/amd-rocm-gfx1151.md +++ b/docs/amd-rocm-gfx1151.md @@ -105,3 +105,7 @@ pull request #241, preserving its commits and authorship. It adds explicit 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..ab53aa20 --- /dev/null +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -0,0 +1,115 @@ +# 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. + +## 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. From 73f5f96168197528a13b97cd365d865fa9b9c80b Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 13:27:05 -0700 Subject: [PATCH 14/72] docs(rocm): retain the GGUF HIP extension cache --- docs/amd-rocm-gfx1151.md | 21 +++++++++++++++++++++ 1 file changed, 21 insertions(+) diff --git a/docs/amd-rocm-gfx1151.md b/docs/amd-rocm-gfx1151.md index 0f699c91..00c14aa0 100644 --- a/docs/amd-rocm-gfx1151.md +++ b/docs/amd-rocm-gfx1151.md @@ -76,6 +76,27 @@ 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 From 6dfa7074eddffcaad0639b0ef1c9a2e8f350b585 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 13:32:00 -0700 Subject: [PATCH 15/72] docs(rocm): add warm LAN-223 TPS and GGUF cache reuse --- docs/lan223-rocm-validation-2026-08-28.md | 31 +++++++++++++++++++++++ 1 file changed, 31 insertions(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index ab53aa20..bd719aeb 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -46,6 +46,37 @@ 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. + +## 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: From 03bec421e0d54c9465d75ce202bd9db84aed7843 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 13:38:37 -0700 Subject: [PATCH 16/72] docs(rocm): compare Gemma HIP throughput with llama.cpp Vulkan --- docs/lan223-rocm-validation-2026-08-28.md | 29 +++++++++++++++++++++++ 1 file changed, 29 insertions(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index bd719aeb..ed815df3 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -66,6 +66,35 @@ 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. + ## GGUF extension reuse validation The first Gemma request after the original source change built the native HIP From 4458a85ac73ff14db6946029a300087b711d1d31 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 14:01:04 -0700 Subject: [PATCH 17/72] docs(rocm): compare FreeToken with llama.cpp ROCm 10 --- docs/lan223-rocm-validation-2026-08-28.md | 53 +++++++++++++++++++++++ 1 file changed, 53 insertions(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index ed815df3..e4a16071 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -95,6 +95,59 @@ 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/ +``` + ## GGUF extension reuse validation The first Gemma request after the original source change built the native HIP From d3a2f57f3ca523b6c80935316efdfed8f3294e23 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 14:23:31 -0700 Subject: [PATCH 18/72] perf(rocm): target GGUF HIP extensions to active gfx --- python/freetoken/kernel/gguf.py | 39 +++++++++++++++++++++- tests/kernels/test_gguf_hip_build_flags.py | 36 ++++++++++++++++++++ 2 files changed, 74 insertions(+), 1 deletion(-) create mode 100644 tests/kernels/test_gguf_hip_build_flags.py diff --git a/python/freetoken/kernel/gguf.py b/python/freetoken/kernel/gguf.py index d4a88c05..458e463f 100644 --- a/python/freetoken/kernel/gguf.py +++ b/python/freetoken/kernel/gguf.py @@ -20,6 +20,43 @@ import torch _CSRC = pathlib.Path(__file__).parent / "csrc" / "gguf" +_TRUE_VALUES = {"1", "true", "yes", "on"} + + +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, with an explicit fast-math experiment. + + ``-O3`` is the normal portable optimization level. Fast math may improve an + AMD compile, but it can alter floating-point contraction and must therefore be + enabled only by ``FREETOKEN_HIP_GGUF_FAST_MATH=1`` while output equivalence is + benchmarked. 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. + """ + target = _hip_target_arch() + if target and not os.environ.get("PYTORCH_ROCM_ARCH"): + os.environ["PYTORCH_ROCM_ARCH"] = target + flags = ["-O3"] + if os.environ.get("FREETOKEN_HIP_GGUF_FAST_MATH", "").strip().lower() in _TRUE_VALUES: + flags.append("-ffast-math") + return flags def _hip_thrust_include() -> str | None: @@ -103,7 +140,7 @@ def _module(): # 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 = ["-O3"] + 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. 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..728b94ae --- /dev/null +++ b/tests/kernels/test_gguf_hip_build_flags.py @@ -0,0 +1,36 @@ +"""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_fast_math_is_explicit_and_preserves_user_target(monkeypatch): + """Fast math is opt-in and an explicit multi-target choice is never replaced.""" + monkeypatch.setenv("PYTORCH_ROCM_ARCH", "gfx1100;gfx1151") + monkeypatch.setenv("FREETOKEN_HIP_GGUF_FAST_MATH", "true") + + assert gguf._hip_gguf_cflags() == ["-O3", "-ffast-math"] + assert gguf._hip_target_arch() == "gfx1100" + assert os.environ["PYTORCH_ROCM_ARCH"] == "gfx1100;gfx1151" From a7e77b759a8c6e0df961ec60edac7be74e1d1566 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 14:27:24 -0700 Subject: [PATCH 19/72] perf(rocm): retain conservative GGUF math flags --- python/freetoken/kernel/gguf.py | 21 ++++++++------------- tests/kernels/test_gguf_hip_build_flags.py | 7 +++---- 2 files changed, 11 insertions(+), 17 deletions(-) diff --git a/python/freetoken/kernel/gguf.py b/python/freetoken/kernel/gguf.py index 458e463f..60dbbb19 100644 --- a/python/freetoken/kernel/gguf.py +++ b/python/freetoken/kernel/gguf.py @@ -20,7 +20,6 @@ import torch _CSRC = pathlib.Path(__file__).parent / "csrc" / "gguf" -_TRUE_VALUES = {"1", "true", "yes", "on"} def _hip_target_arch() -> str | None: @@ -41,22 +40,18 @@ def _hip_target_arch() -> str | None: def _hip_gguf_cflags() -> list[str]: - """Build conservative HIP GGUF flags, with an explicit fast-math experiment. - - ``-O3`` is the normal portable optimization level. Fast math may improve an - AMD compile, but it can alter floating-point contraction and must therefore be - enabled only by ``FREETOKEN_HIP_GGUF_FAST_MATH=1`` while output equivalence is - benchmarked. 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. + """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 - flags = ["-O3"] - if os.environ.get("FREETOKEN_HIP_GGUF_FAST_MATH", "").strip().lower() in _TRUE_VALUES: - flags.append("-ffast-math") - return flags + return ["-O3"] def _hip_thrust_include() -> str | None: diff --git a/tests/kernels/test_gguf_hip_build_flags.py b/tests/kernels/test_gguf_hip_build_flags.py index 728b94ae..942ec9e6 100644 --- a/tests/kernels/test_gguf_hip_build_flags.py +++ b/tests/kernels/test_gguf_hip_build_flags.py @@ -26,11 +26,10 @@ def test_hip_gguf_flags_pin_the_active_gfx_target(monkeypatch): assert os.environ["PYTORCH_ROCM_ARCH"] == "gfx1151" -def test_hip_gguf_fast_math_is_explicit_and_preserves_user_target(monkeypatch): - """Fast math is opt-in and an explicit multi-target choice is never replaced.""" +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") - monkeypatch.setenv("FREETOKEN_HIP_GGUF_FAST_MATH", "true") - assert gguf._hip_gguf_cflags() == ["-O3", "-ffast-math"] + assert gguf._hip_gguf_cflags() == ["-O3"] assert gguf._hip_target_arch() == "gfx1100" assert os.environ["PYTORCH_ROCM_ARCH"] == "gfx1100;gfx1151" From c867f6296e573a00af0b098ae4dd5433dfde8706 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 14:32:07 -0700 Subject: [PATCH 20/72] docs(rocm): record LAN-223 TPS optimization results --- docs/lan223-rocm-validation-2026-08-28.md | 57 +++++++++++++++++++++++ 1 file changed, 57 insertions(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index e4a16071..2e499249 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -148,6 +148,63 @@ Raw artifacts are retained only on LAN-223: /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 | +| llama.cpp `b10141`, ROCm 10 HIP | 60.42 client, 58.88 internal | 128.6 ms | Matched reference | + +The graph configuration removes approximately 1.5 percent of the eager decode +cost, but FreeToken still trails llama.cpp by 7.8 percent using client TPS and +by approximately 5.4 percent compared with llama.cpp's internal decode timing. +The requested criterion of meeting or exceeding llama.cpp is therefore **not +met** by the first configuration pass. + +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-auto \ + --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 automatic +cache selected all 3,840 routed-expert slots. Closing the gap requires a +profile-guided improvement to the HIP GGUF decode kernels or another proven +ROCm attention or quantized-linear implementation. `rocprofv3` ROCm 10 is +installed for that next phase. 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/ +``` + ## GGUF extension reuse validation The first Gemma request after the original source change built the native HIP From 2c12e6c1edbe0949664291b8ed46da8d29ee0dd7 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 14:40:54 -0700 Subject: [PATCH 21/72] docs(rocm): record profiler limitations --- docs/lan223-rocm-validation-2026-08-28.md | 13 ++++++++++--- 1 file changed, 10 insertions(+), 3 deletions(-) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index 2e499249..68028f58 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -194,9 +194,16 @@ 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 automatic cache selected all 3,840 routed-expert slots. Closing the gap requires a profile-guided improvement to the HIP GGUF decode kernels or another proven -ROCm attention or quantized-linear implementation. `rocprofv3` ROCm 10 is -installed for that next phase. A temporary high-performance DPM governor test -could not be run because the non-root LAN-223 account cannot write +ROCm attention or quantized-linear implementation. The available ROCm 10 +`rocprofv3` installation could not yet provide that kernel breakdown: attach +mode reports that the PyTorch process has no `rocp-bg-attach` registration +thread even when launched with `ROCP_TOOL_ATTACH=1`, while launch mode aborts +before FreeToken starts with LLVM's duplicate `spirv-expand-step` option. The +full error evidence is retained in `rocprof-gfx1151*/` and +`rocprof-launch-gfx1151-v2/` under the raw artifact directory. This is a +toolchain issue, not a FreeToken performance result, so no profiler-derived +optimization claim is made here. 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: From bfd0dde04f26a431d1aaedf5e6c1440767c09b6f Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 15:06:34 -0700 Subject: [PATCH 22/72] tools(rocm): capture reproducible LAN-223 baselines --- scripts/lan223-capture-baseline.sh | 106 +++++++++++++++++++++++++++++ 1 file changed, 106 insertions(+) create mode 100644 scripts/lan223-capture-baseline.sh diff --git a/scripts/lan223-capture-baseline.sh b/scripts/lan223-capture-baseline.sh new file mode 100644 index 00000000..e926c673 --- /dev/null +++ b/scripts/lan223-capture-baseline.sh @@ -0,0 +1,106 @@ +#!/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)" + +# 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 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. +capture_command pytorch-runtime.txt "$repo_root/.venv/bin/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))" + +# 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 python-ldd.txt ldd "$repo_root/.venv/bin/python" +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" From ff2a8ccd48219e8a242877a6eeac6fcb6f1441cd Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 15:07:34 -0700 Subject: [PATCH 23/72] fix(tools): resolve sibling FreeToken virtual environment --- scripts/lan223-capture-baseline.sh | 20 ++++++++++++++++++-- 1 file changed, 18 insertions(+), 2 deletions(-) diff --git a/scripts/lan223-capture-baseline.sh b/scripts/lan223-capture-baseline.sh index e926c673..a76aab82 100644 --- a/scripts/lan223-capture-baseline.sh +++ b/scripts/lan223-capture-baseline.sh @@ -27,6 +27,17 @@ llama_binary="${3:-}" # 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" @@ -78,11 +89,16 @@ capture_command performance-environment.txt bash -lc "env | LC_ALL=C sort | grep # Ask the exact FreeToken virtual environment which HIP runtime and device it # sees. This detects a wheel whose embedded runtime differs from host ROCm. -capture_command pytorch-runtime.txt "$repo_root/.venv/bin/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))" +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 python-ldd.txt ldd "$repo_root/.venv/bin/python" capture_command rocprof-ldd.txt ldd /opt/rocm/bin/rocprofv3 # Hash optional comparison artifacts only when the caller supplied a readable From 5c25f1b09e10be7d630889b2447b63e435bf1b8b Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 15:15:57 -0700 Subject: [PATCH 24/72] fix(rocm): profile PyTorch with its wheel SDK --- docs/lan223-rocm-validation-2026-08-28.md | 40 ++++++++++++++++++++ scripts/lan223-rocprof-wheel-sdk.sh | 46 +++++++++++++++++++++++ 2 files changed, 86 insertions(+) create mode 100644 scripts/lan223-rocprof-wheel-sdk.sh diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index 68028f58..73eaccdc 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -212,6 +212,46 @@ Raw campaign artifacts are retained on LAN-223: /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/ +``` + +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 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" "$@" From a6c41e7c3928a85ab828f1486d3494b8a16e7944 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 15:27:38 -0700 Subject: [PATCH 25/72] perf(rocm): process two MoE rows per quantized block --- python/freetoken/kernel/csrc/gguf/ggml-common.h | 7 +++++++ python/freetoken/kernel/csrc/gguf/moe_vec.cuh | 12 ++++++++++++ 2 files changed, 19 insertions(+) diff --git a/python/freetoken/kernel/csrc/gguf/ggml-common.h b/python/freetoken/kernel/csrc/gguf/ggml-common.h index 88c21a4a..d7ec066a 100644 --- a/python/freetoken/kernel/csrc/gguf/ggml-common.h +++ b/python/freetoken/kernel/csrc/gguf/ggml-common.h @@ -10,6 +10,13 @@ #define GGML_CUDA_DMMV_X 32 #define GGML_CUDA_MMV_Y 1 +// Keep the generic quantized matrix-vector launch at one row per block, but +// let the routed-expert path process two independent rows. The latter is the +// dominant LAN-223 decode kernel and matches the rows-per-block strategy used +// by the comparable llama.cpp MoE implementation. Each row occupies its own +// 32-lane thread-x group, so reductions and output addresses remain isolated. +#define GGML_CUDA_MOE_MMV_Y 2 + // Data Structures // QK = number of values after dequantization // QR = QK / number of values before dequantization diff --git a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh index 8cef9e08..8c3b834a 100644 --- a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh +++ b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh @@ -2,6 +2,14 @@ // https://github.com/vllm-project/vllm/blob/4492e3a55428e161ca8db381edc28263e5da4c8d/csrc/quantization/gguf/moe_vec.cuh // copied and adapted from // https://github.com/ggerganov/llama.cpp/blob/b2899/ggml-cuda/mmvq.cu +// +// This header is the routed-expert-only specialization. Temporarily select +// the two-row launch geometry defined in ggml-common.h, then restore the +// generic one-row setting after all MoE wrappers are declared below. Keeping +// the scope local avoids changing non-MoE quantized matrix-vector operations. +#undef GGML_CUDA_MMV_Y +#define GGML_CUDA_MMV_Y GGML_CUDA_MOE_MMV_Y + template static __global__ void moe_vec_q( const void* __restrict__ vx, @@ -411,3 +419,7 @@ static void moe_vec_iq3_s_q8_1_cuda( moe_vec_q <<>>(vx, vy, dst, topk_ids, top_k, ncols, nrows, token_stride); } + +// Restore the generic setting for headers included after this specialization. +#undef GGML_CUDA_MMV_Y +#define GGML_CUDA_MMV_Y 1 From 0fe9e28b7599e7d61eb824f0180817805e99a6e9 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 15:30:23 -0700 Subject: [PATCH 26/72] Revert "perf(rocm): process two MoE rows per quantized block" This reverts commit 4a2440c788120f8ad9cd1f024d762bece84d3aa2. --- python/freetoken/kernel/csrc/gguf/ggml-common.h | 7 ------- python/freetoken/kernel/csrc/gguf/moe_vec.cuh | 12 ------------ 2 files changed, 19 deletions(-) diff --git a/python/freetoken/kernel/csrc/gguf/ggml-common.h b/python/freetoken/kernel/csrc/gguf/ggml-common.h index d7ec066a..88c21a4a 100644 --- a/python/freetoken/kernel/csrc/gguf/ggml-common.h +++ b/python/freetoken/kernel/csrc/gguf/ggml-common.h @@ -10,13 +10,6 @@ #define GGML_CUDA_DMMV_X 32 #define GGML_CUDA_MMV_Y 1 -// Keep the generic quantized matrix-vector launch at one row per block, but -// let the routed-expert path process two independent rows. The latter is the -// dominant LAN-223 decode kernel and matches the rows-per-block strategy used -// by the comparable llama.cpp MoE implementation. Each row occupies its own -// 32-lane thread-x group, so reductions and output addresses remain isolated. -#define GGML_CUDA_MOE_MMV_Y 2 - // Data Structures // QK = number of values after dequantization // QR = QK / number of values before dequantization diff --git a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh index 8c3b834a..8cef9e08 100644 --- a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh +++ b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh @@ -2,14 +2,6 @@ // https://github.com/vllm-project/vllm/blob/4492e3a55428e161ca8db381edc28263e5da4c8d/csrc/quantization/gguf/moe_vec.cuh // copied and adapted from // https://github.com/ggerganov/llama.cpp/blob/b2899/ggml-cuda/mmvq.cu -// -// This header is the routed-expert-only specialization. Temporarily select -// the two-row launch geometry defined in ggml-common.h, then restore the -// generic one-row setting after all MoE wrappers are declared below. Keeping -// the scope local avoids changing non-MoE quantized matrix-vector operations. -#undef GGML_CUDA_MMV_Y -#define GGML_CUDA_MMV_Y GGML_CUDA_MOE_MMV_Y - template static __global__ void moe_vec_q( const void* __restrict__ vx, @@ -419,7 +411,3 @@ static void moe_vec_iq3_s_q8_1_cuda( moe_vec_q <<>>(vx, vy, dst, topk_ids, top_k, ncols, nrows, token_stride); } - -// Restore the generic setting for headers included after this specialization. -#undef GGML_CUDA_MMV_Y -#define GGML_CUDA_MMV_Y 1 From 652b53a91caf234290a639648e8787ef2f49166c Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 15:30:36 -0700 Subject: [PATCH 27/72] docs(rocm): record rejected MoE geometry experiment --- docs/lan223-rocm-validation-2026-08-28.md | 1 + 1 file changed, 1 insertion(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index 73eaccdc..da5e6916 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -163,6 +163,7 @@ prompt tokens and llama.cpp's reused 58. | 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 | | llama.cpp `b10141`, ROCm 10 HIP | 60.42 client, 58.88 internal | 128.6 ms | Matched reference | The graph configuration removes approximately 1.5 percent of the eager decode From b96257ec7a46a61a7d8f117bd53acbd83b722315 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 15:38:07 -0700 Subject: [PATCH 28/72] perf(rocm): test MoE route kernel residency --- python/freetoken/kernel/csrc/gguf/moe_vec.cuh | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh index 8cef9e08..463940e7 100644 --- a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh +++ b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh @@ -3,6 +3,11 @@ // copied and adapted from // https://github.com/ggerganov/llama.cpp/blob/b2899/ggml-cuda/mmvq.cu template +// The decode profile spends about forty percent of GPU kernel time here. Ask +// HIP to keep at least two 32-lane route blocks resident per compute unit so +// independent expert rows can hide memory latency. This does not alter the +// calculation, tensor layout, or one-warp reduction semantics. +__launch_bounds__(WARP_SIZE, 2) static __global__ void moe_vec_q( const void* __restrict__ vx, const void* __restrict__ vy, From 2a259539dfb0edc72c3e7a8c9bd17ffae69b0144 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 15:40:09 -0700 Subject: [PATCH 29/72] Revert "perf(rocm): test MoE route kernel residency" This reverts commit 61a1505b036a51ae8f5ebd92763bb29c7b89083d. --- python/freetoken/kernel/csrc/gguf/moe_vec.cuh | 5 ----- 1 file changed, 5 deletions(-) diff --git a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh index 463940e7..8cef9e08 100644 --- a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh +++ b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh @@ -3,11 +3,6 @@ // copied and adapted from // https://github.com/ggerganov/llama.cpp/blob/b2899/ggml-cuda/mmvq.cu template -// The decode profile spends about forty percent of GPU kernel time here. Ask -// HIP to keep at least two 32-lane route blocks resident per compute unit so -// independent expert rows can hide memory latency. This does not alter the -// calculation, tensor layout, or one-warp reduction semantics. -__launch_bounds__(WARP_SIZE, 2) static __global__ void moe_vec_q( const void* __restrict__ vx, const void* __restrict__ vy, From 43829ee42a57f030e8bb4a30fe5932ae7025e1a1 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 15:40:26 -0700 Subject: [PATCH 30/72] docs(rocm): record rejected occupancy experiment --- docs/lan223-rocm-validation-2026-08-28.md | 1 + 1 file changed, 1 insertion(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index da5e6916..4777cc7c 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -164,6 +164,7 @@ prompt tokens and llama.cpp's reused 58. | 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 | | llama.cpp `b10141`, ROCm 10 HIP | 60.42 client, 58.88 internal | 128.6 ms | Matched reference | The graph configuration removes approximately 1.5 percent of the eager decode From 422a1d6d0fd2746c0419cbeb3102b72dfe91136a Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 15:51:02 -0700 Subject: [PATCH 31/72] docs(rocm): record repaired LAN-223 baseline --- docs/lan223-rocm-validation-2026-08-28.md | 34 +++++++++++++++++++++++ 1 file changed, 34 insertions(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index 4777cc7c..f0bacb65 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -224,6 +224,40 @@ The first baseline was written to: /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. + 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 From b8f0d8aed302bcadbfec4f40ed43f4d1baf0074d Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 15:52:07 -0700 Subject: [PATCH 32/72] perf(rocm): align Q4 vector packed loads --- python/freetoken/kernel/csrc/gguf/vecdotq.cuh | 13 ++++++++++--- 1 file changed, 10 insertions(+), 3 deletions(-) diff --git a/python/freetoken/kernel/csrc/gguf/vecdotq.cuh b/python/freetoken/kernel/csrc/gguf/vecdotq.cuh index 08b4cd26..9902f91d 100644 --- a/python/freetoken/kernel/csrc/gguf/vecdotq.cuh +++ b/python/freetoken/kernel/csrc/gguf/vecdotq.cuh @@ -544,9 +544,16 @@ vec_dot_q4_0_q8_1(const void* __restrict__ vbq, const block_q8_1* __restrict__ b #pragma unroll for (int i = 0; i < VDR_Q4_0_Q8_1_MMVQ; ++i) { - v[i] = get_int_from_uint8(bq4_0->qs, iqs + i); - u[2 * i + 0] = get_int_from_int8_aligned(bq8_1->qs, iqs + i); - u[2 * i + 1] = get_int_from_int8_aligned(bq8_1->qs, iqs + i + QI4_0); + // Keep this Q4_0 load sequence in the same aligned form as the current + // llama.cpp HIP vector path. `qs` in block_q4_0 is halfword aligned, so + // get_int_b2 loads the two adjacent halfwords that form one packed int. + // The Q8_1 quant bytes are int aligned, so get_int_b4 loads each packed + // four-byte group directly. These expressions preserve the exact packed + // values used by the previous helpers while giving the AMD compiler the + // explicit alignment information needed to minimize register pressure. + v[i] = get_int_b2(bq4_0->qs, iqs + i); + u[2 * i + 0] = get_int_b4(bq8_1->qs, iqs + i); + u[2 * i + 1] = get_int_b4(bq8_1->qs, iqs + i + QI4_0); } return vec_dot_q4_0_q8_1_impl(v, u, __half2float(bq4_0->d), bq8_1->ds); From 20d8075d43ed045fcc5e26797fe55c929e2886d5 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 15:56:31 -0700 Subject: [PATCH 33/72] Revert "perf(rocm): align Q4 vector packed loads" This reverts commit b8de1636e639081aca7401ad8062dd89a6193ff5. --- python/freetoken/kernel/csrc/gguf/vecdotq.cuh | 13 +++---------- 1 file changed, 3 insertions(+), 10 deletions(-) diff --git a/python/freetoken/kernel/csrc/gguf/vecdotq.cuh b/python/freetoken/kernel/csrc/gguf/vecdotq.cuh index 9902f91d..08b4cd26 100644 --- a/python/freetoken/kernel/csrc/gguf/vecdotq.cuh +++ b/python/freetoken/kernel/csrc/gguf/vecdotq.cuh @@ -544,16 +544,9 @@ vec_dot_q4_0_q8_1(const void* __restrict__ vbq, const block_q8_1* __restrict__ b #pragma unroll for (int i = 0; i < VDR_Q4_0_Q8_1_MMVQ; ++i) { - // Keep this Q4_0 load sequence in the same aligned form as the current - // llama.cpp HIP vector path. `qs` in block_q4_0 is halfword aligned, so - // get_int_b2 loads the two adjacent halfwords that form one packed int. - // The Q8_1 quant bytes are int aligned, so get_int_b4 loads each packed - // four-byte group directly. These expressions preserve the exact packed - // values used by the previous helpers while giving the AMD compiler the - // explicit alignment information needed to minimize register pressure. - v[i] = get_int_b2(bq4_0->qs, iqs + i); - u[2 * i + 0] = get_int_b4(bq8_1->qs, iqs + i); - u[2 * i + 1] = get_int_b4(bq8_1->qs, iqs + i + QI4_0); + v[i] = get_int_from_uint8(bq4_0->qs, iqs + i); + u[2 * i + 0] = get_int_from_int8_aligned(bq8_1->qs, iqs + i); + u[2 * i + 1] = get_int_from_int8_aligned(bq8_1->qs, iqs + i + QI4_0); } return vec_dot_q4_0_q8_1_impl(v, u, __half2float(bq4_0->d), bq8_1->ds); From 5f68e1eda5b64199aa9b6a2953973eacbfc84c3c Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 15:57:04 -0700 Subject: [PATCH 34/72] docs(rocm): record rejected Q4 load experiment --- docs/lan223-rocm-validation-2026-08-28.md | 30 +++++++++++++++++++++++ 1 file changed, 30 insertions(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index f0bacb65..ec4a1d96 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -258,6 +258,36 @@ 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. + 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 From 0cc5e7765ca5e199c795f3d6f084b2ef93b9a868 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 16:00:30 -0700 Subject: [PATCH 35/72] perf(rocm): test Q4 MoE FP32 intermediates --- .../freetoken/kernel/csrc/gguf/gguf_kernel.cu | 35 ++++++++++++++- python/freetoken/kernel/gguf.py | 13 +++++- python/freetoken/moe/fused_q4_0.py | 44 +++++++++++++++++-- tests/moe/test_fused_q4_0_flags.py | 23 ++++++++++ 4 files changed, 108 insertions(+), 7 deletions(-) create mode 100644 tests/moe/test_fused_q4_0_flags.py diff --git a/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu b/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu index d88960d5..388872ad 100644 --- a/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu +++ b/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu @@ -545,15 +545,48 @@ torch::Tensor ggml_moe_a8_vec( int64_t top_k, int64_t type, int64_t row, - int64_t tokens) { + int64_t tokens, + bool output_fp32) { + // The normal public contract returns the same dtype as X. The opt-in + // Q4_0 experiment below instead uses a FP32 temporary to match the output + // representation used by llama.cpp's HIP MMVQ path while retaining BF16 at + // the Python MoE boundary. Keeping the decision here, next to allocation, + // prevents a mismatched pointer type from reaching a HIP kernel. int col = X.sizes()[1]; const int padded = (col + 512 - 1) / 512 * 512; const at::cuda::OptionalCUDAGuard device_guard(device_of(X)); auto options = torch::TensorOptions().dtype(X.dtype()).device(W.device()); + if (output_fp32) { + options = options.dtype(torch::kFloat); + } at::Tensor Y = torch::zeros({tokens * top_k, row}, options); cudaStream_t stream = at::cuda::getCurrentCUDAStream().stream(); options = torch::TensorOptions().dtype(torch::kInt32).device(W.device()); at::Tensor quant_X = torch::empty({tokens, padded / 32 * 9}, options); + + // The Q4_0 output scalar is independent of the scalar used to quantize X. + // Special-casing this branch avoids changing the other GGUF formats while + // allowing a HIP build to report whether an FP32 vector destination removes + // the register-pressure difference observed in the matched ROCm traces. + if (output_fp32 && type == 2) { + DISPATCH_FLOAT_TYPES(X.scalar_type(), "ggml_moe_vec_a8_fp32_q4_0", [&] { + quantize_row_q8_1_cuda( + (scalar_t*)X.data_ptr(), (void*)quant_X.data_ptr(), col, tokens, stream); + moe_vec_q4_0_q8_1_cuda( + (void*)W.data_ptr(), + (void*)quant_X.data_ptr(), + (float*)Y.data_ptr(), + (int*)topk_ids.data_ptr(), + top_k, + tokens, + col, + row, + quant_X.stride(0), + stream); + }); + return Y; + } + DISPATCH_FLOAT_TYPES(X.scalar_type(), "ggml_moe_vec_a8", [&] { quantize_row_q8_1_cuda((scalar_t*)X.data_ptr(), (void*)quant_X.data_ptr(), col, tokens, stream); switch (type) { diff --git a/python/freetoken/kernel/gguf.py b/python/freetoken/kernel/gguf.py index 60dbbb19..e319c367 100644 --- a/python/freetoken/kernel/gguf.py +++ b/python/freetoken/kernel/gguf.py @@ -229,9 +229,18 @@ def ggml_moe_a8_vec( quant_type: int, row: int, tokens: int, + output_fp32: bool = False, ) -> torch.Tensor: - """MMVQ grouped expert GEMV over stacked experts ``weight[E, row, *]``.""" - return _module().ggml_moe_a8_vec(x, weight, topk_ids, top_k, quant_type, row, tokens) + """MMVQ grouped expert GEMV over stacked experts ``weight[E, row, *]``. + + ``output_fp32`` is an opt-in Q4_0 investigation mode. It keeps activation + quantization in ``x.dtype`` but stores the HIP vector result in FP32 so the + caller can test llama.cpp-compatible intermediate precision. The normal + path remains dtype-preserving and is the only shipping behavior. + """ + return _module().ggml_moe_a8_vec( + x, weight, topk_ids, top_k, quant_type, row, tokens, output_fp32 + ) def ggml_moe_get_block_size(quant_type: int) -> int: diff --git a/python/freetoken/moe/fused_q4_0.py b/python/freetoken/moe/fused_q4_0.py index cdab82bf..cee1cab5 100644 --- a/python/freetoken/moe/fused_q4_0.py +++ b/python/freetoken/moe/fused_q4_0.py @@ -11,6 +11,8 @@ from __future__ import annotations +import os + import torch from freetoken.layers.activation import gelu_and_mul, gelu_tanh_and_mul, silu_and_mul @@ -19,6 +21,16 @@ _ACT = {"silu": silu_and_mul, "gelu": gelu_and_mul, "gelu_tanh": gelu_tanh_and_mul} +def _use_fp32_intermediate() -> bool: + """Return whether the explicitly experimental Q4_0 FP32 path was requested. + + The flag is deliberately strict and opt-in. A normal service launch never + changes precision or throughput behavior merely because the environment + contains an unrelated truthy-looking value. + """ + return os.environ.get("FREETOKEN_GGUF_MOE_FP32_INTERMEDIATE") == "1" + + def fused_experts_gguf_q4_0( hidden_states: torch.Tensor, gate_up_q: torch.Tensor, # [num_slots, 2I, H//32*18] uint8 @@ -38,16 +50,40 @@ def fused_experts_gguf_q4_0( h = down_q.shape[1] # hidden top_k = topk_ids.shape[1] qt = int(GGML_Q4_0) + use_fp32_intermediate = _use_fp32_intermediate() - # gate_up: [num_tokens*top_k, 2I] -> activation -> [num_tokens*top_k, I] - gate_up = ggml_moe_a8_vec(hidden_states, gate_up_q, topk_ids, top_k, qt, n2, num_tokens) + # gate_up: [num_tokens*top_k, 2I] -> activation -> [num_tokens*top_k, I]. + # The opt-in temporary is only for an AMD HIP experiment. It mirrors the + # FP32 vector destination used by llama.cpp without changing the public + # result dtype returned to the transformer layer. + gate_up = ggml_moe_a8_vec( + hidden_states, + gate_up_q, + topk_ids, + top_k, + qt, + n2, + num_tokens, + output_fp32=use_fp32_intermediate, + ) inter = act_fn(gate_up) # down: each of the num_tokens*top_k intermediate rows uses its own expert id. - out = ggml_moe_a8_vec(inter, down_q, topk_ids, 1, qt, h, num_tokens * top_k) + out = ggml_moe_a8_vec( + inter, + down_q, + topk_ids, + 1, + qt, + h, + num_tokens * top_k, + output_fp32=use_fp32_intermediate, + ) out = out.reshape(num_tokens, top_k, h) * topk_weights.reshape(num_tokens, top_k, 1).to( out.dtype ) - return out.sum(dim=1) + # Preserve the original BF16 caller contract even when the temporary + # candidate computed its two quantized vector products in FP32. + return out.sum(dim=1).to(hidden_states.dtype) __all__ = ["fused_experts_gguf_q4_0"] diff --git a/tests/moe/test_fused_q4_0_flags.py b/tests/moe/test_fused_q4_0_flags.py new file mode 100644 index 00000000..f13a0a44 --- /dev/null +++ b/tests/moe/test_fused_q4_0_flags.py @@ -0,0 +1,23 @@ +"""Contract tests for opt-in Q4_0 MoE HIP investigation flags. + +These unit tests intentionally do not allocate a GPU tensor or invoke hipcc. +They protect the important public guarantee that FP32 intermediates are an +explicit experiment and cannot be enabled by an arbitrary environment value. +""" + +from freetoken.moe import fused_q4_0 + + +def test_q4_fp32_intermediate_is_disabled_without_the_exact_opt_in(monkeypatch): + """Absent and loose truthy values preserve the dtype-stable shipping path.""" + monkeypatch.delenv("FREETOKEN_GGUF_MOE_FP32_INTERMEDIATE", raising=False) + assert fused_q4_0._use_fp32_intermediate() is False + + monkeypatch.setenv("FREETOKEN_GGUF_MOE_FP32_INTERMEDIATE", "true") + assert fused_q4_0._use_fp32_intermediate() is False + + +def test_q4_fp32_intermediate_requires_exact_one(monkeypatch): + """Only the documented value enables the temporary FP32 HIP experiment.""" + monkeypatch.setenv("FREETOKEN_GGUF_MOE_FP32_INTERMEDIATE", "1") + assert fused_q4_0._use_fp32_intermediate() is True From ec230c4481238ae6106774ca280ae85f9190bf0a Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 16:06:40 -0700 Subject: [PATCH 36/72] docs(rocm): record rejected FP32 MoE experiment --- docs/lan223-rocm-validation-2026-08-28.md | 35 +++++++++++++++++++++++ 1 file changed, 35 insertions(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index ec4a1d96..3b7debf8 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -288,6 +288,41 @@ 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 From 1be4c838483a11cdbaffa42333e4f8b43f66fd46 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 16:06:41 -0700 Subject: [PATCH 37/72] Revert "perf(rocm): test Q4 MoE FP32 intermediates" This reverts commit 5bbe10f508c5fddf15804bb72a28fce655902d19. --- .../freetoken/kernel/csrc/gguf/gguf_kernel.cu | 35 +-------------- python/freetoken/kernel/gguf.py | 13 +----- python/freetoken/moe/fused_q4_0.py | 44 ++----------------- tests/moe/test_fused_q4_0_flags.py | 23 ---------- 4 files changed, 7 insertions(+), 108 deletions(-) delete mode 100644 tests/moe/test_fused_q4_0_flags.py diff --git a/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu b/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu index 388872ad..d88960d5 100644 --- a/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu +++ b/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu @@ -545,48 +545,15 @@ torch::Tensor ggml_moe_a8_vec( int64_t top_k, int64_t type, int64_t row, - int64_t tokens, - bool output_fp32) { - // The normal public contract returns the same dtype as X. The opt-in - // Q4_0 experiment below instead uses a FP32 temporary to match the output - // representation used by llama.cpp's HIP MMVQ path while retaining BF16 at - // the Python MoE boundary. Keeping the decision here, next to allocation, - // prevents a mismatched pointer type from reaching a HIP kernel. + int64_t tokens) { int col = X.sizes()[1]; const int padded = (col + 512 - 1) / 512 * 512; const at::cuda::OptionalCUDAGuard device_guard(device_of(X)); auto options = torch::TensorOptions().dtype(X.dtype()).device(W.device()); - if (output_fp32) { - options = options.dtype(torch::kFloat); - } at::Tensor Y = torch::zeros({tokens * top_k, row}, options); cudaStream_t stream = at::cuda::getCurrentCUDAStream().stream(); options = torch::TensorOptions().dtype(torch::kInt32).device(W.device()); at::Tensor quant_X = torch::empty({tokens, padded / 32 * 9}, options); - - // The Q4_0 output scalar is independent of the scalar used to quantize X. - // Special-casing this branch avoids changing the other GGUF formats while - // allowing a HIP build to report whether an FP32 vector destination removes - // the register-pressure difference observed in the matched ROCm traces. - if (output_fp32 && type == 2) { - DISPATCH_FLOAT_TYPES(X.scalar_type(), "ggml_moe_vec_a8_fp32_q4_0", [&] { - quantize_row_q8_1_cuda( - (scalar_t*)X.data_ptr(), (void*)quant_X.data_ptr(), col, tokens, stream); - moe_vec_q4_0_q8_1_cuda( - (void*)W.data_ptr(), - (void*)quant_X.data_ptr(), - (float*)Y.data_ptr(), - (int*)topk_ids.data_ptr(), - top_k, - tokens, - col, - row, - quant_X.stride(0), - stream); - }); - return Y; - } - DISPATCH_FLOAT_TYPES(X.scalar_type(), "ggml_moe_vec_a8", [&] { quantize_row_q8_1_cuda((scalar_t*)X.data_ptr(), (void*)quant_X.data_ptr(), col, tokens, stream); switch (type) { diff --git a/python/freetoken/kernel/gguf.py b/python/freetoken/kernel/gguf.py index e319c367..60dbbb19 100644 --- a/python/freetoken/kernel/gguf.py +++ b/python/freetoken/kernel/gguf.py @@ -229,18 +229,9 @@ def ggml_moe_a8_vec( quant_type: int, row: int, tokens: int, - output_fp32: bool = False, ) -> torch.Tensor: - """MMVQ grouped expert GEMV over stacked experts ``weight[E, row, *]``. - - ``output_fp32`` is an opt-in Q4_0 investigation mode. It keeps activation - quantization in ``x.dtype`` but stores the HIP vector result in FP32 so the - caller can test llama.cpp-compatible intermediate precision. The normal - path remains dtype-preserving and is the only shipping behavior. - """ - return _module().ggml_moe_a8_vec( - x, weight, topk_ids, top_k, quant_type, row, tokens, output_fp32 - ) + """MMVQ grouped expert GEMV over stacked experts ``weight[E, row, *]``.""" + return _module().ggml_moe_a8_vec(x, weight, topk_ids, top_k, quant_type, row, tokens) def ggml_moe_get_block_size(quant_type: int) -> int: diff --git a/python/freetoken/moe/fused_q4_0.py b/python/freetoken/moe/fused_q4_0.py index cee1cab5..cdab82bf 100644 --- a/python/freetoken/moe/fused_q4_0.py +++ b/python/freetoken/moe/fused_q4_0.py @@ -11,8 +11,6 @@ from __future__ import annotations -import os - import torch from freetoken.layers.activation import gelu_and_mul, gelu_tanh_and_mul, silu_and_mul @@ -21,16 +19,6 @@ _ACT = {"silu": silu_and_mul, "gelu": gelu_and_mul, "gelu_tanh": gelu_tanh_and_mul} -def _use_fp32_intermediate() -> bool: - """Return whether the explicitly experimental Q4_0 FP32 path was requested. - - The flag is deliberately strict and opt-in. A normal service launch never - changes precision or throughput behavior merely because the environment - contains an unrelated truthy-looking value. - """ - return os.environ.get("FREETOKEN_GGUF_MOE_FP32_INTERMEDIATE") == "1" - - def fused_experts_gguf_q4_0( hidden_states: torch.Tensor, gate_up_q: torch.Tensor, # [num_slots, 2I, H//32*18] uint8 @@ -50,40 +38,16 @@ def fused_experts_gguf_q4_0( h = down_q.shape[1] # hidden top_k = topk_ids.shape[1] qt = int(GGML_Q4_0) - use_fp32_intermediate = _use_fp32_intermediate() - # gate_up: [num_tokens*top_k, 2I] -> activation -> [num_tokens*top_k, I]. - # The opt-in temporary is only for an AMD HIP experiment. It mirrors the - # FP32 vector destination used by llama.cpp without changing the public - # result dtype returned to the transformer layer. - gate_up = ggml_moe_a8_vec( - hidden_states, - gate_up_q, - topk_ids, - top_k, - qt, - n2, - num_tokens, - output_fp32=use_fp32_intermediate, - ) + # gate_up: [num_tokens*top_k, 2I] -> activation -> [num_tokens*top_k, I] + gate_up = ggml_moe_a8_vec(hidden_states, gate_up_q, topk_ids, top_k, qt, n2, num_tokens) inter = act_fn(gate_up) # down: each of the num_tokens*top_k intermediate rows uses its own expert id. - out = ggml_moe_a8_vec( - inter, - down_q, - topk_ids, - 1, - qt, - h, - num_tokens * top_k, - output_fp32=use_fp32_intermediate, - ) + out = ggml_moe_a8_vec(inter, down_q, topk_ids, 1, qt, h, num_tokens * top_k) out = out.reshape(num_tokens, top_k, h) * topk_weights.reshape(num_tokens, top_k, 1).to( out.dtype ) - # Preserve the original BF16 caller contract even when the temporary - # candidate computed its two quantized vector products in FP32. - return out.sum(dim=1).to(hidden_states.dtype) + return out.sum(dim=1) __all__ = ["fused_experts_gguf_q4_0"] diff --git a/tests/moe/test_fused_q4_0_flags.py b/tests/moe/test_fused_q4_0_flags.py deleted file mode 100644 index f13a0a44..00000000 --- a/tests/moe/test_fused_q4_0_flags.py +++ /dev/null @@ -1,23 +0,0 @@ -"""Contract tests for opt-in Q4_0 MoE HIP investigation flags. - -These unit tests intentionally do not allocate a GPU tensor or invoke hipcc. -They protect the important public guarantee that FP32 intermediates are an -explicit experiment and cannot be enabled by an arbitrary environment value. -""" - -from freetoken.moe import fused_q4_0 - - -def test_q4_fp32_intermediate_is_disabled_without_the_exact_opt_in(monkeypatch): - """Absent and loose truthy values preserve the dtype-stable shipping path.""" - monkeypatch.delenv("FREETOKEN_GGUF_MOE_FP32_INTERMEDIATE", raising=False) - assert fused_q4_0._use_fp32_intermediate() is False - - monkeypatch.setenv("FREETOKEN_GGUF_MOE_FP32_INTERMEDIATE", "true") - assert fused_q4_0._use_fp32_intermediate() is False - - -def test_q4_fp32_intermediate_requires_exact_one(monkeypatch): - """Only the documented value enables the temporary FP32 HIP experiment.""" - monkeypatch.setenv("FREETOKEN_GGUF_MOE_FP32_INTERMEDIATE", "1") - assert fused_q4_0._use_fp32_intermediate() is True From 45dfb43433f8e903cceb7e211b68192e8bd26bd4 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 16:13:04 -0700 Subject: [PATCH 38/72] bench(rocm): add Gemma Q4 MoE kernel probe --- benchmarks/bench_gguf_q4_moe_kernel.py | 184 +++++++++++++++++++++++++ 1 file changed, 184 insertions(+) create mode 100644 benchmarks/bench_gguf_q4_moe_kernel.py 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()) From b2e4271cf4e89c26d7824e8517c801c4c16ae1d9 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 16:19:23 -0700 Subject: [PATCH 39/72] perf(rocm): specialize Q4 MoE two-row wave --- python/freetoken/kernel/csrc/gguf/moe_vec.cuh | 108 ++++++++++++++++++ 1 file changed, 108 insertions(+) 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( From bac5ed1729aa3cae1f9537ffd047c0f8351ed5f6 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 16:34:07 -0700 Subject: [PATCH 40/72] docs(rocm): record accepted two-row MoE wave --- docs/lan223-rocm-validation-2026-08-28.md | 49 +++++++++++++++++++++++ 1 file changed, 49 insertions(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index 3b7debf8..2c260a3e 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -165,6 +165,7 @@ prompt tokens and llama.cpp's reused 58. | 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 | | llama.cpp `b10141`, ROCm 10 HIP | 60.42 client, 58.88 internal | 128.6 ms | Matched reference | The graph configuration removes approximately 1.5 percent of the eager decode @@ -173,6 +174,54 @@ by approximately 5.4 percent compared with llama.cpp's internal decode timing. The requested criterion of meeting or exceeding llama.cpp is therefore **not met** by the first configuration pass. +### 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/ +``` + The best verified FreeToken command shape is: ```bash From a0a4932518d79027ebf888b68bab0f6f2733aa20 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 16:37:04 -0700 Subject: [PATCH 41/72] bench(rocm): add Gemma dense Q4 kernel probe --- benchmarks/bench_gguf_q4_dense_kernel.py | 137 +++++++++++++++++++++++ 1 file changed, 137 insertions(+) create mode 100644 benchmarks/bench_gguf_q4_dense_kernel.py 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()) From 56cfadc08876e780f406508a7365c9c331617deb Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 16:39:04 -0700 Subject: [PATCH 42/72] perf(rocm): specialize dense Q4 two-row wave --- python/freetoken/kernel/csrc/gguf/mmvq.cuh | 94 ++++++++++++++++++++++ 1 file changed, 94 insertions(+) diff --git a/python/freetoken/kernel/csrc/gguf/mmvq.cuh b/python/freetoken/kernel/csrc/gguf/mmvq.cuh index 7331731a..8295bb47 100644 --- a/python/freetoken/kernel/csrc/gguf/mmvq.cuh +++ b/python/freetoken/kernel/csrc/gguf/mmvq.cuh @@ -47,6 +47,20 @@ static __global__ void mul_mat_vec_q( } } +#if defined(USE_ROCM) +// The HIP dense-Q4 launcher is defined below the generic wrapper so the other +// quantization wrappers retain their original ordering in this shared header. +template +static void mul_mat_vec_q4_0_q8_1_hip_two_rows_cuda( + const void* vx, + const void* vy, + scalar_t* dst, + const int ncols, + const int nrows, + const int nvecs, + cudaStream_t stream); +#endif + template static void mul_mat_vec_q4_0_q8_1_cuda( const void* vx, @@ -56,12 +70,92 @@ static void mul_mat_vec_q4_0_q8_1_cuda( const int nrows, const int nvecs, cudaStream_t stream) { +#if defined(USE_ROCM) + // Keep CUDA on the established generic path. The HIP route is an isolated + // one-wave/two-row specialization validated only for gfx1151 so far. + mul_mat_vec_q4_0_q8_1_hip_two_rows_cuda(vx, vy, dst, ncols, nrows, nvecs, stream); +#else const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y; const dim3 block_nums(block_num_y, nvecs, 1); const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1); mul_mat_vec_q <<>>(vx, vy, dst, ncols, nrows, nvecs); +#endif +} + +#if defined(USE_ROCM) +// Dense Q4_0 counterpart to the accepted routed-expert HIP specialization. +// Each 32-lane wave computes two adjacent matrix rows for one input vector. +// That differs from the earlier MMVQ geometry, which scheduled one independent +// wave per row. It does not change Q4_0 packing, Q8_1 activation packing, +// vector indexing, or the BF16 destination contract exposed to Python. +template +__launch_bounds__(WARP_SIZE, 1) +static __global__ void mul_mat_vec_q4_0_hip_two_rows( + const void* __restrict__ vx, + const void* __restrict__ vy, + scalar_t* __restrict__ dst, + const int ncols, + const int nrows, + const int nvecs) { + const int row0 = 2 * blockIdx.x; + const int vec = blockIdx.y; + if (row0 >= nrows || vec >= nvecs) { + return; + } + + const int blocks_per_row = ncols / QK4_0; + const int blocks_per_wave = VDR_Q4_0_Q8_1_MMVQ * WARP_SIZE / QI4_0; + const int padded_cols = (ncols + 512 - 1) / 512 * 512; + const block_q4_0* x = (const block_q4_0*)vx; + const block_q8_1* y = ((const block_q8_1*)vy) + vec * (padded_cols / QK8_1); + + // The same lane visits corresponding Q4_0 blocks in both rows, allowing two + // independent sums without altering the established per-row reduction order. + 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); + } + } + + // XOR reductions leave the complete value in every lane. Lanes zero and + // one write the two rows directly, so this specialization requires no LDS. +#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[vec * nrows + row0] = tmp0; + } + if (threadIdx.x == 1 && row0 + 1 < nrows) { + dst[vec * nrows + row0 + 1] = tmp1; + } +} + +template +static void mul_mat_vec_q4_0_q8_1_hip_two_rows_cuda( + const void* vx, + const void* vy, + scalar_t* dst, + const int ncols, + const int nrows, + const int nvecs, + cudaStream_t stream) { + // X is a pair of output rows, Y selects one input vector. This retains the + // public ``[nvecs, nrows]`` result ordering of ggml_mul_mat_vec_a8. + const dim3 block_nums((nrows + 1) / 2, nvecs, 1); + const dim3 block_dims(WARP_SIZE, 1, 1); + mul_mat_vec_q4_0_hip_two_rows + <<>>(vx, vy, dst, ncols, nrows, nvecs); } +#endif template static void mul_mat_vec_q4_1_q8_1_cuda( From c03fd57fffa3364a4ad43e662aaa9852cb8a4b58 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 16:44:12 -0700 Subject: [PATCH 43/72] docs(rocm): record rejected dense Q4 wave --- docs/lan223-rocm-validation-2026-08-28.md | 27 +++++++++++++++++++++++ 1 file changed, 27 insertions(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index 2c260a3e..496d41e4 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -222,6 +222,33 @@ Artifacts are retained on LAN-223: /home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/q4-moe-two-row-wave-20260828T231950Z/api-repeats-20260828T232646Z/ ``` +### 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/ +``` + The best verified FreeToken command shape is: ```bash From 5e9ed2756bb1301b9f8ba981da2fa9fe53a0a5b4 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 16:44:12 -0700 Subject: [PATCH 44/72] Revert "perf(rocm): specialize dense Q4 two-row wave" This reverts commit b4a53d11f48cb52d9baa1432e8302a5ba03c71dd. --- python/freetoken/kernel/csrc/gguf/mmvq.cuh | 94 ---------------------- 1 file changed, 94 deletions(-) diff --git a/python/freetoken/kernel/csrc/gguf/mmvq.cuh b/python/freetoken/kernel/csrc/gguf/mmvq.cuh index 8295bb47..7331731a 100644 --- a/python/freetoken/kernel/csrc/gguf/mmvq.cuh +++ b/python/freetoken/kernel/csrc/gguf/mmvq.cuh @@ -47,20 +47,6 @@ static __global__ void mul_mat_vec_q( } } -#if defined(USE_ROCM) -// The HIP dense-Q4 launcher is defined below the generic wrapper so the other -// quantization wrappers retain their original ordering in this shared header. -template -static void mul_mat_vec_q4_0_q8_1_hip_two_rows_cuda( - const void* vx, - const void* vy, - scalar_t* dst, - const int ncols, - const int nrows, - const int nvecs, - cudaStream_t stream); -#endif - template static void mul_mat_vec_q4_0_q8_1_cuda( const void* vx, @@ -70,92 +56,12 @@ static void mul_mat_vec_q4_0_q8_1_cuda( const int nrows, const int nvecs, cudaStream_t stream) { -#if defined(USE_ROCM) - // Keep CUDA on the established generic path. The HIP route is an isolated - // one-wave/two-row specialization validated only for gfx1151 so far. - mul_mat_vec_q4_0_q8_1_hip_two_rows_cuda(vx, vy, dst, ncols, nrows, nvecs, stream); -#else const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y; const dim3 block_nums(block_num_y, nvecs, 1); const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1); mul_mat_vec_q <<>>(vx, vy, dst, ncols, nrows, nvecs); -#endif -} - -#if defined(USE_ROCM) -// Dense Q4_0 counterpart to the accepted routed-expert HIP specialization. -// Each 32-lane wave computes two adjacent matrix rows for one input vector. -// That differs from the earlier MMVQ geometry, which scheduled one independent -// wave per row. It does not change Q4_0 packing, Q8_1 activation packing, -// vector indexing, or the BF16 destination contract exposed to Python. -template -__launch_bounds__(WARP_SIZE, 1) -static __global__ void mul_mat_vec_q4_0_hip_two_rows( - const void* __restrict__ vx, - const void* __restrict__ vy, - scalar_t* __restrict__ dst, - const int ncols, - const int nrows, - const int nvecs) { - const int row0 = 2 * blockIdx.x; - const int vec = blockIdx.y; - if (row0 >= nrows || vec >= nvecs) { - return; - } - - const int blocks_per_row = ncols / QK4_0; - const int blocks_per_wave = VDR_Q4_0_Q8_1_MMVQ * WARP_SIZE / QI4_0; - const int padded_cols = (ncols + 512 - 1) / 512 * 512; - const block_q4_0* x = (const block_q4_0*)vx; - const block_q8_1* y = ((const block_q8_1*)vy) + vec * (padded_cols / QK8_1); - - // The same lane visits corresponding Q4_0 blocks in both rows, allowing two - // independent sums without altering the established per-row reduction order. - 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); - } - } - - // XOR reductions leave the complete value in every lane. Lanes zero and - // one write the two rows directly, so this specialization requires no LDS. -#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[vec * nrows + row0] = tmp0; - } - if (threadIdx.x == 1 && row0 + 1 < nrows) { - dst[vec * nrows + row0 + 1] = tmp1; - } -} - -template -static void mul_mat_vec_q4_0_q8_1_hip_two_rows_cuda( - const void* vx, - const void* vy, - scalar_t* dst, - const int ncols, - const int nrows, - const int nvecs, - cudaStream_t stream) { - // X is a pair of output rows, Y selects one input vector. This retains the - // public ``[nvecs, nrows]`` result ordering of ggml_mul_mat_vec_a8. - const dim3 block_nums((nrows + 1) / 2, nvecs, 1); - const dim3 block_dims(WARP_SIZE, 1, 1); - mul_mat_vec_q4_0_hip_two_rows - <<>>(vx, vy, dst, ncols, nrows, nvecs); } -#endif template static void mul_mat_vec_q4_1_q8_1_cuda( From f9b41e02e88be409201f186d51784a014068248b Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 16:49:19 -0700 Subject: [PATCH 45/72] bench(rocm): probe dense Q4 FP32 output --- benchmarks/bench_gguf_q4_dense_kernel.py | 15 ++++++++++- .../freetoken/kernel/csrc/gguf/gguf_kernel.cu | 24 ++++++++++++++--- python/freetoken/kernel/gguf.py | 16 +++++++++--- tests/kernels/test_gguf_dense_fp32_probe.py | 26 +++++++++++++++++++ 4 files changed, 73 insertions(+), 8 deletions(-) create mode 100644 tests/kernels/test_gguf_dense_fp32_probe.py diff --git a/benchmarks/bench_gguf_q4_dense_kernel.py b/benchmarks/bench_gguf_q4_dense_kernel.py index c8327740..e76dc701 100644 --- a/benchmarks/bench_gguf_q4_dense_kernel.py +++ b/benchmarks/bench_gguf_q4_dense_kernel.py @@ -56,6 +56,11 @@ def _parse_args() -> argparse.Namespace: 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( + "--output-fp32", + action="store_true", + help="isolated Q4_0 destination-type probe; never enables it for model serving", + ) parser.add_argument("--json", type=Path, help="write one JSON artifact") return parser.parse_args() @@ -100,7 +105,13 @@ def main() -> int: 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) + return ggml_mul_mat_vec_a8( + weight, + x, + int(GGML_Q4_0), + rows, + output_fp32=args.output_fp32, + ) for _ in range(args.warmup): result = call() @@ -113,6 +124,7 @@ def call() -> torch.Tensor: "cols": cols, "vectors": args.vectors, "output_shape": list(result.shape), + "output_dtype": str(result.dtype), "average_us": _average_event_us(call, args.repetitions, device), } ) @@ -122,6 +134,7 @@ def call() -> torch.Tensor: "hip": torch.version.hip, "torch": torch.__version__, "quant_type": "Q4_0", + "output_fp32": args.output_fp32, "warmup": args.warmup, "repetitions": args.repetitions, "measurements": measurements, diff --git a/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu b/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu index d88960d5..53444e2d 100644 --- a/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu +++ b/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu @@ -95,12 +95,20 @@ torch::Tensor ggml_mul_mat_vec_a8( torch::Tensor W, // quant weight torch::Tensor X, // input int64_t type, - int64_t row) { + int64_t row, + bool output_fp32) { int col = X.sizes()[1]; int vecs = X.sizes()[0]; const int padded = (col + 512 - 1) / 512 * 512; const at::cuda::OptionalCUDAGuard device_guard(device_of(X)); - auto options = torch::TensorOptions().dtype(X.dtype()).device(W.device()); + // Production callers keep their input dtype output. ``output_fp32`` is an + // explicitly opt-in, Q4_0-only benchmark probe used to measure whether the + // destination type changes HIP register allocation. It must not silently + // change model-layer numerics or the public GGUF layer contract. + const bool q4_fp32_probe = output_fp32 && type == 2; + auto options = torch::TensorOptions() + .dtype(q4_fp32_probe ? torch::kFloat32 : X.scalar_type()) + .device(W.device()); at::Tensor Y = torch::empty({vecs, row}, options); cudaStream_t stream = at::cuda::getCurrentCUDAStream().stream(); options = torch::TensorOptions().dtype(torch::kInt32).device(W.device()); @@ -109,8 +117,16 @@ torch::Tensor ggml_mul_mat_vec_a8( quantize_row_q8_1_cuda((scalar_t*)X.data_ptr(), (void*)quant_X.data_ptr(), col, vecs, stream); switch (type) { case 2: - mul_mat_vec_q4_0_q8_1_cuda( - (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (scalar_t*)Y.data_ptr(), col, row, vecs, stream); + if (q4_fp32_probe) { + // The input still uses ``scalar_t`` during Q8_1 quantization. Only + // the final GEMV store changes type, isolating the code-generation + // question from all production inference behavior. + mul_mat_vec_q4_0_q8_1_cuda( + (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (float*)Y.data_ptr(), col, row, vecs, stream); + } else { + mul_mat_vec_q4_0_q8_1_cuda( + (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (scalar_t*)Y.data_ptr(), col, row, vecs, stream); + } break; case 3: mul_mat_vec_q4_1_q8_1_cuda( diff --git a/python/freetoken/kernel/gguf.py b/python/freetoken/kernel/gguf.py index 60dbbb19..b60f5909 100644 --- a/python/freetoken/kernel/gguf.py +++ b/python/freetoken/kernel/gguf.py @@ -190,10 +190,20 @@ def ggml_dequantize( def ggml_mul_mat_vec_a8( - weight: torch.Tensor, x: torch.Tensor, quant_type: int, row: int + weight: torch.Tensor, + x: torch.Tensor, + quant_type: int, + row: int, + *, + output_fp32: bool = False, ) -> torch.Tensor: - """MMVQ: small-batch GEMV with on-the-fly dequant. ``row`` = output features.""" - return _module().ggml_mul_mat_vec_a8(weight, x, quant_type, row) + """Run small-batch quantized GEMV; ``output_fp32`` is an isolated Q4 benchmark probe. + + Normal inference leaves ``output_fp32`` false, preserving the output dtype + expected by GGUF layers. The opt-in mode changes only Q4_0's destination + storage so LAN-223 profiling can compare register allocation with llama.cpp. + """ + return _module().ggml_mul_mat_vec_a8(weight, x, quant_type, row, output_fp32) def ggml_mul_mat_a8( diff --git a/tests/kernels/test_gguf_dense_fp32_probe.py b/tests/kernels/test_gguf_dense_fp32_probe.py new file mode 100644 index 00000000..f17bd4a2 --- /dev/null +++ b/tests/kernels/test_gguf_dense_fp32_probe.py @@ -0,0 +1,26 @@ +"""Guard the benchmark-only FP32 dense-output experiment's public boundary.""" + +from __future__ import annotations + +import ast +from pathlib import Path + + +REPOSITORY_ROOT = Path(__file__).resolve().parents[2] + + +def test_dense_probe_is_explicitly_opt_in_and_not_used_by_layers() -> None: + """Prevent an experiment flag from changing normal GGUF model execution.""" + wrapper_path = REPOSITORY_ROOT / "python" / "freetoken" / "kernel" / "gguf.py" + layer_path = REPOSITORY_ROOT / "python" / "freetoken" / "layers" / "gguf.py" + wrapper_module = ast.parse(wrapper_path.read_text(encoding="utf-8")) + function = next( + node + for node in wrapper_module.body + if isinstance(node, ast.FunctionDef) and node.name == "ggml_mul_mat_vec_a8" + ) + output_argument = next(arg for arg in function.args.kwonlyargs if arg.arg == "output_fp32") + default = function.args.kw_defaults[function.args.kwonlyargs.index(output_argument)] + + assert isinstance(default, ast.Constant) and default.value is False + assert "output_fp32" not in layer_path.read_text(encoding="utf-8") From 9f95e6c966747c0cab2c510baf2356ad1e38179b Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 16:53:44 -0700 Subject: [PATCH 46/72] Revert "bench(rocm): probe dense Q4 FP32 output" This reverts commit e77a44b27c9ea995dab8f2d9f361841be6c76ffa. --- benchmarks/bench_gguf_q4_dense_kernel.py | 15 +---------- .../freetoken/kernel/csrc/gguf/gguf_kernel.cu | 24 +++-------------- python/freetoken/kernel/gguf.py | 16 +++--------- tests/kernels/test_gguf_dense_fp32_probe.py | 26 ------------------- 4 files changed, 8 insertions(+), 73 deletions(-) delete mode 100644 tests/kernels/test_gguf_dense_fp32_probe.py diff --git a/benchmarks/bench_gguf_q4_dense_kernel.py b/benchmarks/bench_gguf_q4_dense_kernel.py index e76dc701..c8327740 100644 --- a/benchmarks/bench_gguf_q4_dense_kernel.py +++ b/benchmarks/bench_gguf_q4_dense_kernel.py @@ -56,11 +56,6 @@ def _parse_args() -> argparse.Namespace: 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( - "--output-fp32", - action="store_true", - help="isolated Q4_0 destination-type probe; never enables it for model serving", - ) parser.add_argument("--json", type=Path, help="write one JSON artifact") return parser.parse_args() @@ -105,13 +100,7 @@ def main() -> int: 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, - output_fp32=args.output_fp32, - ) + return ggml_mul_mat_vec_a8(weight, x, int(GGML_Q4_0), rows) for _ in range(args.warmup): result = call() @@ -124,7 +113,6 @@ def call() -> torch.Tensor: "cols": cols, "vectors": args.vectors, "output_shape": list(result.shape), - "output_dtype": str(result.dtype), "average_us": _average_event_us(call, args.repetitions, device), } ) @@ -134,7 +122,6 @@ def call() -> torch.Tensor: "hip": torch.version.hip, "torch": torch.__version__, "quant_type": "Q4_0", - "output_fp32": args.output_fp32, "warmup": args.warmup, "repetitions": args.repetitions, "measurements": measurements, diff --git a/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu b/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu index 53444e2d..d88960d5 100644 --- a/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu +++ b/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu @@ -95,20 +95,12 @@ torch::Tensor ggml_mul_mat_vec_a8( torch::Tensor W, // quant weight torch::Tensor X, // input int64_t type, - int64_t row, - bool output_fp32) { + int64_t row) { int col = X.sizes()[1]; int vecs = X.sizes()[0]; const int padded = (col + 512 - 1) / 512 * 512; const at::cuda::OptionalCUDAGuard device_guard(device_of(X)); - // Production callers keep their input dtype output. ``output_fp32`` is an - // explicitly opt-in, Q4_0-only benchmark probe used to measure whether the - // destination type changes HIP register allocation. It must not silently - // change model-layer numerics or the public GGUF layer contract. - const bool q4_fp32_probe = output_fp32 && type == 2; - auto options = torch::TensorOptions() - .dtype(q4_fp32_probe ? torch::kFloat32 : X.scalar_type()) - .device(W.device()); + auto options = torch::TensorOptions().dtype(X.dtype()).device(W.device()); at::Tensor Y = torch::empty({vecs, row}, options); cudaStream_t stream = at::cuda::getCurrentCUDAStream().stream(); options = torch::TensorOptions().dtype(torch::kInt32).device(W.device()); @@ -117,16 +109,8 @@ torch::Tensor ggml_mul_mat_vec_a8( quantize_row_q8_1_cuda((scalar_t*)X.data_ptr(), (void*)quant_X.data_ptr(), col, vecs, stream); switch (type) { case 2: - if (q4_fp32_probe) { - // The input still uses ``scalar_t`` during Q8_1 quantization. Only - // the final GEMV store changes type, isolating the code-generation - // question from all production inference behavior. - mul_mat_vec_q4_0_q8_1_cuda( - (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (float*)Y.data_ptr(), col, row, vecs, stream); - } else { - mul_mat_vec_q4_0_q8_1_cuda( - (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (scalar_t*)Y.data_ptr(), col, row, vecs, stream); - } + mul_mat_vec_q4_0_q8_1_cuda( + (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (scalar_t*)Y.data_ptr(), col, row, vecs, stream); break; case 3: mul_mat_vec_q4_1_q8_1_cuda( diff --git a/python/freetoken/kernel/gguf.py b/python/freetoken/kernel/gguf.py index b60f5909..60dbbb19 100644 --- a/python/freetoken/kernel/gguf.py +++ b/python/freetoken/kernel/gguf.py @@ -190,20 +190,10 @@ def ggml_dequantize( def ggml_mul_mat_vec_a8( - weight: torch.Tensor, - x: torch.Tensor, - quant_type: int, - row: int, - *, - output_fp32: bool = False, + weight: torch.Tensor, x: torch.Tensor, quant_type: int, row: int ) -> torch.Tensor: - """Run small-batch quantized GEMV; ``output_fp32`` is an isolated Q4 benchmark probe. - - Normal inference leaves ``output_fp32`` false, preserving the output dtype - expected by GGUF layers. The opt-in mode changes only Q4_0's destination - storage so LAN-223 profiling can compare register allocation with llama.cpp. - """ - return _module().ggml_mul_mat_vec_a8(weight, x, quant_type, row, output_fp32) + """MMVQ: small-batch GEMV with on-the-fly dequant. ``row`` = output features.""" + return _module().ggml_mul_mat_vec_a8(weight, x, quant_type, row) def ggml_mul_mat_a8( diff --git a/tests/kernels/test_gguf_dense_fp32_probe.py b/tests/kernels/test_gguf_dense_fp32_probe.py deleted file mode 100644 index f17bd4a2..00000000 --- a/tests/kernels/test_gguf_dense_fp32_probe.py +++ /dev/null @@ -1,26 +0,0 @@ -"""Guard the benchmark-only FP32 dense-output experiment's public boundary.""" - -from __future__ import annotations - -import ast -from pathlib import Path - - -REPOSITORY_ROOT = Path(__file__).resolve().parents[2] - - -def test_dense_probe_is_explicitly_opt_in_and_not_used_by_layers() -> None: - """Prevent an experiment flag from changing normal GGUF model execution.""" - wrapper_path = REPOSITORY_ROOT / "python" / "freetoken" / "kernel" / "gguf.py" - layer_path = REPOSITORY_ROOT / "python" / "freetoken" / "layers" / "gguf.py" - wrapper_module = ast.parse(wrapper_path.read_text(encoding="utf-8")) - function = next( - node - for node in wrapper_module.body - if isinstance(node, ast.FunctionDef) and node.name == "ggml_mul_mat_vec_a8" - ) - output_argument = next(arg for arg in function.args.kwonlyargs if arg.arg == "output_fp32") - default = function.args.kw_defaults[function.args.kwonlyargs.index(output_argument)] - - assert isinstance(default, ast.Constant) and default.value is False - assert "output_fp32" not in layer_path.read_text(encoding="utf-8") From 9dda5ef3384b8c3a8b9cf1fe26e178c0ccf88a50 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 16:54:04 -0700 Subject: [PATCH 47/72] docs(rocm): record rejected dense FP32 probe --- docs/lan223-rocm-validation-2026-08-28.md | 26 +++++++++++++++++++++++ 1 file changed, 26 insertions(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index 496d41e4..bc30563e 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -249,6 +249,32 @@ The retained raw evidence is: /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/ +``` + The best verified FreeToken command shape is: ```bash From b5975d8464c18f8febcdcb37060d1ec070105a87 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 16:54:45 -0700 Subject: [PATCH 48/72] perf(rocm): constrain dense Q4 GEMV occupancy --- python/freetoken/kernel/csrc/gguf/mmvq.cuh | 14 +++++++++++++- 1 file changed, 13 insertions(+), 1 deletion(-) diff --git a/python/freetoken/kernel/csrc/gguf/mmvq.cuh b/python/freetoken/kernel/csrc/gguf/mmvq.cuh index 7331731a..3c9c2bfe 100644 --- a/python/freetoken/kernel/csrc/gguf/mmvq.cuh +++ b/python/freetoken/kernel/csrc/gguf/mmvq.cuh @@ -1,8 +1,18 @@ // copied from // https://github.com/vllm-project/vllm/blob/4492e3a55428e161ca8db381edc28263e5da4c8d/csrc/quantization/gguf/mmvq.cuh // copied and adapted from https://github.com/ggerganov/llama.cpp/blob/b2899/ggml-cuda/mmvq.cu +// The LAN-223 HIP profiler reports materially higher VGPR use for this +// one-wave GEMV than the matched llama.cpp implementation. Constrain only the +// HIP compiler to one 32-lane wave per workgroup, matching that reference's +// launch contract. CUDA keeps its upstream scheduling and code generation. +#if defined(USE_ROCM) +#define FREETOKEN_DENSE_GEMV_LAUNCH_BOUNDS __launch_bounds__(WARP_SIZE, 1) +#else +#define FREETOKEN_DENSE_GEMV_LAUNCH_BOUNDS +#endif + template -static __global__ void mul_mat_vec_q( +static __global__ FREETOKEN_DENSE_GEMV_LAUNCH_BOUNDS void mul_mat_vec_q( const void* __restrict__ vx, const void* __restrict__ vy, scalar_t* __restrict__ dst, @@ -47,6 +57,8 @@ static __global__ void mul_mat_vec_q( } } +#undef FREETOKEN_DENSE_GEMV_LAUNCH_BOUNDS + template static void mul_mat_vec_q4_0_q8_1_cuda( const void* vx, From 5461b28604b1f1d1c4e4edd942c170a140d74e59 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 17:04:25 -0700 Subject: [PATCH 49/72] Revert "perf(rocm): constrain dense Q4 GEMV occupancy" This reverts commit c7009a97638f802dcafaa375729d7a6a5b3835d4. --- python/freetoken/kernel/csrc/gguf/mmvq.cuh | 14 +------------- 1 file changed, 1 insertion(+), 13 deletions(-) diff --git a/python/freetoken/kernel/csrc/gguf/mmvq.cuh b/python/freetoken/kernel/csrc/gguf/mmvq.cuh index 3c9c2bfe..7331731a 100644 --- a/python/freetoken/kernel/csrc/gguf/mmvq.cuh +++ b/python/freetoken/kernel/csrc/gguf/mmvq.cuh @@ -1,18 +1,8 @@ // copied from // https://github.com/vllm-project/vllm/blob/4492e3a55428e161ca8db381edc28263e5da4c8d/csrc/quantization/gguf/mmvq.cuh // copied and adapted from https://github.com/ggerganov/llama.cpp/blob/b2899/ggml-cuda/mmvq.cu -// The LAN-223 HIP profiler reports materially higher VGPR use for this -// one-wave GEMV than the matched llama.cpp implementation. Constrain only the -// HIP compiler to one 32-lane wave per workgroup, matching that reference's -// launch contract. CUDA keeps its upstream scheduling and code generation. -#if defined(USE_ROCM) -#define FREETOKEN_DENSE_GEMV_LAUNCH_BOUNDS __launch_bounds__(WARP_SIZE, 1) -#else -#define FREETOKEN_DENSE_GEMV_LAUNCH_BOUNDS -#endif - template -static __global__ FREETOKEN_DENSE_GEMV_LAUNCH_BOUNDS void mul_mat_vec_q( +static __global__ void mul_mat_vec_q( const void* __restrict__ vx, const void* __restrict__ vy, scalar_t* __restrict__ dst, @@ -57,8 +47,6 @@ static __global__ FREETOKEN_DENSE_GEMV_LAUNCH_BOUNDS void mul_mat_vec_q( } } -#undef FREETOKEN_DENSE_GEMV_LAUNCH_BOUNDS - template static void mul_mat_vec_q4_0_q8_1_cuda( const void* vx, From f1cd13a067f1b40ecc71238ad8e764ba9f0f4951 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 17:05:13 -0700 Subject: [PATCH 50/72] docs(rocm): record rejected dense launch bounds --- docs/lan223-rocm-validation-2026-08-28.md | 26 +++++++++++++++++++++++ 1 file changed, 26 insertions(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index bc30563e..e5b9d62b 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -275,6 +275,32 @@ The retained raw evidence is: /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/ +``` + The best verified FreeToken command shape is: ```bash From 05b13b695b4c05f101f2b95a767f4ce3e7ee23d9 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 17:08:46 -0700 Subject: [PATCH 51/72] perf(rocm): test indexed Q4 dense HIP kernel --- python/freetoken/kernel/csrc/gguf/mmvq.cuh | 79 ++++++++++++++++++++++ 1 file changed, 79 insertions(+) diff --git a/python/freetoken/kernel/csrc/gguf/mmvq.cuh b/python/freetoken/kernel/csrc/gguf/mmvq.cuh index 7331731a..5a9ab437 100644 --- a/python/freetoken/kernel/csrc/gguf/mmvq.cuh +++ b/python/freetoken/kernel/csrc/gguf/mmvq.cuh @@ -1,6 +1,76 @@ // copied from // https://github.com/vllm-project/vllm/blob/4492e3a55428e161ca8db381edc28263e5da4c8d/csrc/quantization/gguf/mmvq.cuh // copied and adapted from https://github.com/ggerganov/llama.cpp/blob/b2899/ggml-cuda/mmvq.cu + +#if defined(USE_ROCM) +// This is deliberately Q4_0-specific. Current llama.cpp's gfx1151 kernel +// keeps the base weight pointer and the Q4 block index separate until the +// vector-dot helper, whereas the inherited FreeToken template creates a +// per-iteration typed pointer. The two forms are numerically equivalent but +// can lead to different HIP register allocation. Keep it independent from +// CUDA and from all other GGUF quantization types while LAN-223 benchmarks +// establish whether the compiler actually benefits. +static __device__ __forceinline__ float vec_dot_q4_0_q8_1_hip_indexed( + const void* __restrict__ vx, + const block_q8_1* __restrict__ bq8_1, + const int& weight_block, + const int& iqs) { + const block_q4_0* bq4_0 = (const block_q4_0*)vx + weight_block; + int v[VDR_Q4_0_Q8_1_MMVQ]; + int u[2 * VDR_Q4_0_Q8_1_MMVQ]; + +#pragma unroll + for (int i = 0; i < VDR_Q4_0_Q8_1_MMVQ; ++i) { + v[i] = get_int_from_uint8(bq4_0->qs, iqs + i); + u[2 * i + 0] = get_int_from_int8_aligned(bq8_1->qs, iqs + i); + u[2 * i + 1] = get_int_from_int8_aligned(bq8_1->qs, iqs + i + QI4_0); + } + + return vec_dot_q4_0_q8_1_impl(v, u, __half2float(bq4_0->d), bq8_1->ds); +} + +template +static __global__ void mul_mat_vec_q4_0_hip_indexed( + const void* __restrict__ vx, + const void* __restrict__ vy, + scalar_t* __restrict__ dst, + const int ncols, + const int nrows, + const int nvecs) { + // Preserve the wrapper's row mapping if an operator changes + // ``GGML_CUDA_MMV_Y`` from its LAN-223 value of one. The experiment changes + // pointer/index representation, not the externally selected launch shape. + const int row = blockIdx.x * blockDim.y + threadIdx.y; + const int vec = blockIdx.y; + if (row >= nrows || vec >= nvecs) { + return; + } + + constexpr int blocks_per_iter = VDR_Q4_0_Q8_1_MMVQ * WARP_SIZE / QI4_0; + const int blocks_per_row = ncols / QK4_0; + const int quant_rows = (ncols + 512 - 1) / 512 * 512; + const int weight_block_base = row * blocks_per_row; + const block_q8_1* y = (const block_q8_1*)vy + vec * (quant_rows / QK8_1); + float sum = 0.0f; + + for (int weight_block = threadIdx.x / (QI4_0 / VDR_Q4_0_Q8_1_MMVQ); + weight_block < blocks_per_row; + weight_block += blocks_per_iter) { + const int quant_block = weight_block * (QK4_0 / QK8_1); + const int iqs = VDR_Q4_0_Q8_1_MMVQ * (threadIdx.x % (QI4_0 / VDR_Q4_0_Q8_1_MMVQ)); + sum += vec_dot_q4_0_q8_1_hip_indexed(vx, &y[quant_block], weight_block_base + weight_block, iqs); + } + +#pragma unroll + for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) { + sum += SGLANG_SHFL_XOR_SYNC(uint32_t(-1), sum, mask); + } + if (threadIdx.x == 0) { + dst[vec * nrows + row] = sum; + } +} +#endif + template static __global__ void mul_mat_vec_q( const void* __restrict__ vx, @@ -59,8 +129,17 @@ static void mul_mat_vec_q4_0_q8_1_cuda( const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y; const dim3 block_nums(block_num_y, nvecs, 1); const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1); +#if defined(USE_ROCM) + // HIP uses the indexed form above only for Q4_0. It has the same launch + // geometry and arithmetic as the generic path, but separates the weight + // base and block index to test the code-generation difference observed in + // the current llama.cpp source and profiler trace. + mul_mat_vec_q4_0_hip_indexed<<>>( + vx, vy, dst, ncols, nrows, nvecs); +#else mul_mat_vec_q <<>>(vx, vy, dst, ncols, nrows, nvecs); +#endif } template From 70ac13bb25ab0698b0dc825db3621284dc052157 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 17:13:40 -0700 Subject: [PATCH 52/72] Revert "perf(rocm): test indexed Q4 dense HIP kernel" This reverts commit a5e04d1f9c1f2763c52fe019aafa14558c819800. --- python/freetoken/kernel/csrc/gguf/mmvq.cuh | 79 ---------------------- 1 file changed, 79 deletions(-) diff --git a/python/freetoken/kernel/csrc/gguf/mmvq.cuh b/python/freetoken/kernel/csrc/gguf/mmvq.cuh index 5a9ab437..7331731a 100644 --- a/python/freetoken/kernel/csrc/gguf/mmvq.cuh +++ b/python/freetoken/kernel/csrc/gguf/mmvq.cuh @@ -1,76 +1,6 @@ // copied from // https://github.com/vllm-project/vllm/blob/4492e3a55428e161ca8db381edc28263e5da4c8d/csrc/quantization/gguf/mmvq.cuh // copied and adapted from https://github.com/ggerganov/llama.cpp/blob/b2899/ggml-cuda/mmvq.cu - -#if defined(USE_ROCM) -// This is deliberately Q4_0-specific. Current llama.cpp's gfx1151 kernel -// keeps the base weight pointer and the Q4 block index separate until the -// vector-dot helper, whereas the inherited FreeToken template creates a -// per-iteration typed pointer. The two forms are numerically equivalent but -// can lead to different HIP register allocation. Keep it independent from -// CUDA and from all other GGUF quantization types while LAN-223 benchmarks -// establish whether the compiler actually benefits. -static __device__ __forceinline__ float vec_dot_q4_0_q8_1_hip_indexed( - const void* __restrict__ vx, - const block_q8_1* __restrict__ bq8_1, - const int& weight_block, - const int& iqs) { - const block_q4_0* bq4_0 = (const block_q4_0*)vx + weight_block; - int v[VDR_Q4_0_Q8_1_MMVQ]; - int u[2 * VDR_Q4_0_Q8_1_MMVQ]; - -#pragma unroll - for (int i = 0; i < VDR_Q4_0_Q8_1_MMVQ; ++i) { - v[i] = get_int_from_uint8(bq4_0->qs, iqs + i); - u[2 * i + 0] = get_int_from_int8_aligned(bq8_1->qs, iqs + i); - u[2 * i + 1] = get_int_from_int8_aligned(bq8_1->qs, iqs + i + QI4_0); - } - - return vec_dot_q4_0_q8_1_impl(v, u, __half2float(bq4_0->d), bq8_1->ds); -} - -template -static __global__ void mul_mat_vec_q4_0_hip_indexed( - const void* __restrict__ vx, - const void* __restrict__ vy, - scalar_t* __restrict__ dst, - const int ncols, - const int nrows, - const int nvecs) { - // Preserve the wrapper's row mapping if an operator changes - // ``GGML_CUDA_MMV_Y`` from its LAN-223 value of one. The experiment changes - // pointer/index representation, not the externally selected launch shape. - const int row = blockIdx.x * blockDim.y + threadIdx.y; - const int vec = blockIdx.y; - if (row >= nrows || vec >= nvecs) { - return; - } - - constexpr int blocks_per_iter = VDR_Q4_0_Q8_1_MMVQ * WARP_SIZE / QI4_0; - const int blocks_per_row = ncols / QK4_0; - const int quant_rows = (ncols + 512 - 1) / 512 * 512; - const int weight_block_base = row * blocks_per_row; - const block_q8_1* y = (const block_q8_1*)vy + vec * (quant_rows / QK8_1); - float sum = 0.0f; - - for (int weight_block = threadIdx.x / (QI4_0 / VDR_Q4_0_Q8_1_MMVQ); - weight_block < blocks_per_row; - weight_block += blocks_per_iter) { - const int quant_block = weight_block * (QK4_0 / QK8_1); - const int iqs = VDR_Q4_0_Q8_1_MMVQ * (threadIdx.x % (QI4_0 / VDR_Q4_0_Q8_1_MMVQ)); - sum += vec_dot_q4_0_q8_1_hip_indexed(vx, &y[quant_block], weight_block_base + weight_block, iqs); - } - -#pragma unroll - for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) { - sum += SGLANG_SHFL_XOR_SYNC(uint32_t(-1), sum, mask); - } - if (threadIdx.x == 0) { - dst[vec * nrows + row] = sum; - } -} -#endif - template static __global__ void mul_mat_vec_q( const void* __restrict__ vx, @@ -129,17 +59,8 @@ static void mul_mat_vec_q4_0_q8_1_cuda( const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y; const dim3 block_nums(block_num_y, nvecs, 1); const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1); -#if defined(USE_ROCM) - // HIP uses the indexed form above only for Q4_0. It has the same launch - // geometry and arithmetic as the generic path, but separates the weight - // base and block index to test the code-generation difference observed in - // the current llama.cpp source and profiler trace. - mul_mat_vec_q4_0_hip_indexed<<>>( - vx, vy, dst, ncols, nrows, nvecs); -#else mul_mat_vec_q <<>>(vx, vy, dst, ncols, nrows, nvecs); -#endif } template From 2beadb0a83bed3952e5c52aba9a902ed368aaaaf Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 17:14:11 -0700 Subject: [PATCH 53/72] docs(rocm): record rejected indexed Q4 candidate --- docs/lan223-rocm-validation-2026-08-28.md | 28 +++++++++++++++++++++++ 1 file changed, 28 insertions(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index e5b9d62b..49e16944 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -301,6 +301,34 @@ The retained raw evidence is: /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/ +``` + The best verified FreeToken command shape is: ```bash From 50a8ae63b136502ba951cb1be8dfc373f516b683 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 17:17:46 -0700 Subject: [PATCH 54/72] perf(rocm): group Q4 MoE routes per workgroup --- python/freetoken/kernel/csrc/gguf/moe_vec.cuh | 100 +++++++++++++++++- 1 file changed, 96 insertions(+), 4 deletions(-) diff --git a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh index 3d12f1e9..a6c6f379 100644 --- a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh +++ b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh @@ -66,6 +66,19 @@ static void moe_vec_q4_0_q8_1_hip_two_rows_cuda( const int nrows, const int token_stride, cudaStream_t stream); + +template +static void moe_vec_q4_0_q8_1_hip_route_group8_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 @@ -81,10 +94,10 @@ static void moe_vec_q4_0_q8_1_cuda( 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( + // Route AMD builds through the grouped-route specialization. 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_route_group8_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; @@ -176,6 +189,85 @@ static void moe_vec_q4_0_q8_1_hip_two_rows_cuda( moe_vec_q4_0_hip_two_rows <<>>(vx, vy, dst, topk_ids, top_k, ncols, nrows, token_stride); } + +// Current llama.cpp's dedicated MMVQ MoE path groups the routed-token axis as +// one wave per route inside a multi-wave workgroup. FreeToken's original +// launcher put every route in a separate 32-thread workgroup. This HIP-only +// variant retains the accepted two-output-row arithmetic above, but groups up +// to eight routes so the gfx1151 scheduler sees the same route-axis shape as +// llama.cpp. Each wave remains independent, so no inter-route synchronization +// or shared-memory reduction is required. +template +__launch_bounds__(WARP_SIZE * 8, 1) +static __global__ void moe_vec_q4_0_hip_route_group8( + const void* __restrict__ vx, + const void* __restrict__ vy, + scalar_t* __restrict__ dst, + const int* __restrict__ topk_ids, + const int topk, + const int routes, + const int ncols, + const int nrows, + const int token_stride) { + const int row0 = 2 * blockIdx.x; + const int route = blockIdx.y * blockDim.y + threadIdx.y; + if (row0 >= nrows || route >= routes) { + 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); + 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); + } + } + + // ROCm's shuffle is wave scoped. This reduces only lanes in the current + // route wave even though the workgroup contains up to eight such waves. +#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_route_group8_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) { + constexpr int routes_per_block = 8; + const int routes = tokens * top_k; + const dim3 block_nums((nrows + 1) / 2, (routes + routes_per_block - 1) / routes_per_block, 1); + const dim3 block_dims(WARP_SIZE, routes_per_block, 1); + moe_vec_q4_0_hip_route_group8<<>>( + vx, vy, dst, topk_ids, top_k, routes, ncols, nrows, token_stride); +} #endif template From d934215c77495698ded6d632b6e77f15ce6294ba Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 17:20:29 -0700 Subject: [PATCH 55/72] Revert "perf(rocm): group Q4 MoE routes per workgroup" This reverts commit dc73e8e9578e5c8a14b694eb6acc7058e48da562. --- python/freetoken/kernel/csrc/gguf/moe_vec.cuh | 100 +----------------- 1 file changed, 4 insertions(+), 96 deletions(-) diff --git a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh index a6c6f379..3d12f1e9 100644 --- a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh +++ b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh @@ -66,19 +66,6 @@ static void moe_vec_q4_0_q8_1_hip_two_rows_cuda( const int nrows, const int token_stride, cudaStream_t stream); - -template -static void moe_vec_q4_0_q8_1_hip_route_group8_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 @@ -94,10 +81,10 @@ static void moe_vec_q4_0_q8_1_cuda( const int token_stride, cudaStream_t stream) { #if defined(USE_ROCM) - // Route AMD builds through the grouped-route specialization. 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_route_group8_cuda( + // 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; @@ -189,85 +176,6 @@ static void moe_vec_q4_0_q8_1_hip_two_rows_cuda( moe_vec_q4_0_hip_two_rows <<>>(vx, vy, dst, topk_ids, top_k, ncols, nrows, token_stride); } - -// Current llama.cpp's dedicated MMVQ MoE path groups the routed-token axis as -// one wave per route inside a multi-wave workgroup. FreeToken's original -// launcher put every route in a separate 32-thread workgroup. This HIP-only -// variant retains the accepted two-output-row arithmetic above, but groups up -// to eight routes so the gfx1151 scheduler sees the same route-axis shape as -// llama.cpp. Each wave remains independent, so no inter-route synchronization -// or shared-memory reduction is required. -template -__launch_bounds__(WARP_SIZE * 8, 1) -static __global__ void moe_vec_q4_0_hip_route_group8( - const void* __restrict__ vx, - const void* __restrict__ vy, - scalar_t* __restrict__ dst, - const int* __restrict__ topk_ids, - const int topk, - const int routes, - const int ncols, - const int nrows, - const int token_stride) { - const int row0 = 2 * blockIdx.x; - const int route = blockIdx.y * blockDim.y + threadIdx.y; - if (row0 >= nrows || route >= routes) { - 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); - 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); - } - } - - // ROCm's shuffle is wave scoped. This reduces only lanes in the current - // route wave even though the workgroup contains up to eight such waves. -#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_route_group8_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) { - constexpr int routes_per_block = 8; - const int routes = tokens * top_k; - const dim3 block_nums((nrows + 1) / 2, (routes + routes_per_block - 1) / routes_per_block, 1); - const dim3 block_dims(WARP_SIZE, routes_per_block, 1); - moe_vec_q4_0_hip_route_group8<<>>( - vx, vy, dst, topk_ids, top_k, routes, ncols, nrows, token_stride); -} #endif template From 4cd975a0d13c570981131454c9dbdc385040965f Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 17:20:46 -0700 Subject: [PATCH 56/72] docs(rocm): record rejected MoE route grouping --- docs/lan223-rocm-validation-2026-08-28.md | 23 +++++++++++++++++++++++ 1 file changed, 23 insertions(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index 49e16944..040715bf 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -329,6 +329,29 @@ The retained raw evidence is: /home/david/freetoken-amd/artifacts/amd-deep-investigation-2026-08-28/dense-indexed-pointer-api-20260829T001156Z/ ``` +### 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/ +``` + The best verified FreeToken command shape is: ```bash From d3a8c9b475eb433b478c17d4c10ce5a3ef12b56a Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 17:23:00 -0700 Subject: [PATCH 57/72] bench(rocm): probe Gemma GQA decode tile --- benchmarks/bench_rocm_gqa_attention.py | 98 +++++++++++++++++++++ python/freetoken/kernel/triton/attention.py | 17 +++- 2 files changed, 113 insertions(+), 2 deletions(-) create mode 100644 benchmarks/bench_rocm_gqa_attention.py diff --git a/benchmarks/bench_rocm_gqa_attention.py b/benchmarks/bench_rocm_gqa_attention.py new file mode 100644 index 00000000..2963938a --- /dev/null +++ b/benchmarks/bench_rocm_gqa_attention.py @@ -0,0 +1,98 @@ +"""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, default=2) + parser.add_argument("--sequence-length", type=int, default=1024) + 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, 8, 256, 8 + 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: 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=1024, rocm_block_h_probe=probe, + ) + + for _ in range(args.warmup): + default = call(None) + for _ in range(args.warmup): + candidate = call(args.probe_block_h) + 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, + "probe_block_h": args.probe_block_h, + "default_us": _event_us(lambda: call(None), args.repetitions), + "probe_us": _event_us(lambda: call(args.probe_block_h), 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/python/freetoken/kernel/triton/attention.py b/python/freetoken/kernel/triton/attention.py index bc1fd11a..661d2a08 100644 --- a/python/freetoken/kernel/triton/attention.py +++ b/python/freetoken/kernel/triton/attention.py @@ -364,8 +364,15 @@ 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, ) -> 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. + + ``rocm_block_h_probe`` is benchmark-only: it asks HIP Triton for an explicit + power-of-two query-head tile so LAN-223 can measure whether a smaller GQA + tile lowers correctly. Normal serving leaves it ``None`` and therefore + preserves the established ROCm 16-head padded tile. + """ 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,7 +403,13 @@ 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 torch.version.hip is not None: + 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 From 43d350e2bc68747d7932fdbbf5ea444fbec8a605 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 17:24:27 -0700 Subject: [PATCH 58/72] bench(rocm): parameterize GQA attention probes --- benchmarks/bench_rocm_gqa_attention.py | 22 ++++++++++------ python/freetoken/kernel/triton/attention.py | 28 ++++++++++++++++----- 2 files changed, 37 insertions(+), 13 deletions(-) diff --git a/benchmarks/bench_rocm_gqa_attention.py b/benchmarks/bench_rocm_gqa_attention.py index 2963938a..0a6efd68 100644 --- a/benchmarks/bench_rocm_gqa_attention.py +++ b/benchmarks/bench_rocm_gqa_attention.py @@ -23,7 +23,9 @@ 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, default=2) + 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("--warmup", type=int, default=20) parser.add_argument("--repetitions", type=int, default=200) @@ -64,16 +66,17 @@ def main() -> int: 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: int | None) -> torch.Tensor: + 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=1024, rocm_block_h_probe=probe, + splits, head_dim**-0.5, sliding_window=1024, 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) + default = call(None, None, None) for _ in range(args.warmup): - candidate = call(args.probe_block_h) + 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 = { @@ -82,8 +85,13 @@ def call(probe: int | None) -> torch.Tensor: "geometry": {"q_heads": query_heads, "kv_heads": kv_heads, "head_dim": head_dim}, "sequence_length": args.sequence_length, "probe_block_h": args.probe_block_h, - "default_us": _event_us(lambda: call(None), args.repetitions), - "probe_us": _event_us(lambda: call(args.probe_block_h), args.repetitions), + "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, } diff --git a/python/freetoken/kernel/triton/attention.py b/python/freetoken/kernel/triton/attention.py index 661d2a08..cbf36637 100644 --- a/python/freetoken/kernel/triton/attention.py +++ b/python/freetoken/kernel/triton/attention.py @@ -365,13 +365,15 @@ def decode_paged_attention( 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. - ``rocm_block_h_probe`` is benchmark-only: it asks HIP Triton for an explicit - power-of-two query-head tile so LAN-223 can measure whether a smaller GQA - tile lowers correctly. Normal serving leaves it ``None`` and therefore - preserves the established ROCm 16-head padded tile. + 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 @@ -418,6 +420,20 @@ def decode_paged_attention( 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) @@ -448,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)]( From 2b6a139d2c8e818dfe7b99d9b16def990a85a937 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 17:25:43 -0700 Subject: [PATCH 59/72] bench(rocm): validate Gemma attention warp candidate --- benchmarks/bench_rocm_gqa_attention.py | 12 ++++++++++-- python/freetoken/kernel/triton/attention.py | 13 +++++++++++++ 2 files changed, 23 insertions(+), 2 deletions(-) diff --git a/benchmarks/bench_rocm_gqa_attention.py b/benchmarks/bench_rocm_gqa_attention.py index 0a6efd68..503b64d3 100644 --- a/benchmarks/bench_rocm_gqa_attention.py +++ b/benchmarks/bench_rocm_gqa_attention.py @@ -27,6 +27,9 @@ def _parse_args() -> argparse.Namespace: 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) @@ -55,7 +58,9 @@ def main() -> int: 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, 8, 256, 8 + 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) @@ -69,7 +74,9 @@ def main() -> int: 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=1024, rocm_block_h_probe=probe_h, + 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, ) @@ -84,6 +91,7 @@ def call(probe_h: int | None, probe_n: int | None, probe_warps: int | None) -> t "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, diff --git a/python/freetoken/kernel/triton/attention.py b/python/freetoken/kernel/triton/attention.py index cbf36637..060649ce 100644 --- a/python/freetoken/kernel/triton/attention.py +++ b/python/freetoken/kernel/triton/attention.py @@ -1,6 +1,7 @@ from __future__ import annotations import functools +import os import torch import triton @@ -434,6 +435,18 @@ def decode_paged_attention( 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 + elif torch.version.hip is not None: + # Keep production behavior at four warps unless the LAN-223 experiment + # explicitly opts in. Parsing happens before Triton dispatch and has no + # device-side cost or graph-captured data dependency. + configured_warps = os.environ.get("FREETOKEN_ROCM_ATTENTION_WARPS") + if configured_warps is not None: + try: + num_warps = int(configured_warps) + except ValueError as error: + raise ValueError("FREETOKEN_ROCM_ATTENTION_WARPS must be an integer") from error + if num_warps not in (1, 2, 4, 8): + raise ValueError("FREETOKEN_ROCM_ATTENTION_WARPS must be one of 1, 2, 4, or 8") _decode_grouped_stage1_kernel[ (batch, triton.cdiv(num_q_heads, valid_block_h), max_kv_splits) From c07c45f412717c0bc0e08d7976b4ed5bb9678959 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 17:31:23 -0700 Subject: [PATCH 60/72] docs(rocm): reject GQA warp tuning candidate --- docs/lan223-rocm-validation-2026-08-28.md | 32 +++++++++++++++++++++ python/freetoken/kernel/triton/attention.py | 13 --------- 2 files changed, 32 insertions(+), 13 deletions(-) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index 040715bf..0cb156cb 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -352,6 +352,38 @@ The retained raw evidence is: /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 diff --git a/python/freetoken/kernel/triton/attention.py b/python/freetoken/kernel/triton/attention.py index 060649ce..cbf36637 100644 --- a/python/freetoken/kernel/triton/attention.py +++ b/python/freetoken/kernel/triton/attention.py @@ -1,7 +1,6 @@ from __future__ import annotations import functools -import os import torch import triton @@ -435,18 +434,6 @@ def decode_paged_attention( 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 - elif torch.version.hip is not None: - # Keep production behavior at four warps unless the LAN-223 experiment - # explicitly opts in. Parsing happens before Triton dispatch and has no - # device-side cost or graph-captured data dependency. - configured_warps = os.environ.get("FREETOKEN_ROCM_ATTENTION_WARPS") - if configured_warps is not None: - try: - num_warps = int(configured_warps) - except ValueError as error: - raise ValueError("FREETOKEN_ROCM_ATTENTION_WARPS must be an integer") from error - if num_warps not in (1, 2, 4, 8): - raise ValueError("FREETOKEN_ROCM_ATTENTION_WARPS must be one of 1, 2, 4, or 8") _decode_grouped_stage1_kernel[ (batch, triton.cdiv(num_q_heads, valid_block_h), max_kv_splits) From c49938e8c4cbea7e4da53bbcd4c0da7ba653f9da Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 17:37:02 -0700 Subject: [PATCH 61/72] perf(rocm): scalarize dense Q4 dot temporaries --- python/freetoken/kernel/csrc/gguf/mmvq.cuh | 8 +++++ python/freetoken/kernel/csrc/gguf/vecdotq.cuh | 31 +++++++++++++++++++ 2 files changed, 39 insertions(+) diff --git a/python/freetoken/kernel/csrc/gguf/mmvq.cuh b/python/freetoken/kernel/csrc/gguf/mmvq.cuh index 7331731a..6e6b703a 100644 --- a/python/freetoken/kernel/csrc/gguf/mmvq.cuh +++ b/python/freetoken/kernel/csrc/gguf/mmvq.cuh @@ -59,8 +59,16 @@ static void mul_mat_vec_q4_0_q8_1_cuda( const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y; const dim3 block_nums(block_num_y, nvecs, 1); const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1); +#if defined USE_ROCM + // The scalarized helper is HIP-only and keeps the CUDA code object stable. + // It implements the same Q4_0/Q8_1 dot-product sequence with shorter-lived + // temporaries for an isolated gfx1151 dense-decode experiment. + mul_mat_vec_q + <<>>(vx, vy, dst, ncols, nrows, nvecs); +#else mul_mat_vec_q <<>>(vx, vy, dst, ncols, nrows, nvecs); +#endif } template diff --git a/python/freetoken/kernel/csrc/gguf/vecdotq.cuh b/python/freetoken/kernel/csrc/gguf/vecdotq.cuh index 08b4cd26..e9312707 100644 --- a/python/freetoken/kernel/csrc/gguf/vecdotq.cuh +++ b/python/freetoken/kernel/csrc/gguf/vecdotq.cuh @@ -552,6 +552,37 @@ vec_dot_q4_0_q8_1(const void* __restrict__ vbq, const block_q8_1* __restrict__ b return vec_dot_q4_0_q8_1_impl(v, u, __half2float(bq4_0->d), bq8_1->ds); } +#if defined USE_ROCM +// HIP-only dense-GEMV variant of the Q4_0 x Q8_1 inner product. The generic +// helper above stages two packed Q4 words and four Q8 words in local arrays +// before forwarding them to a template helper. A decode lane always consumes +// exactly those six words, so this variant keeps them as named scalars and +// executes the same four DP4A instructions in the same order. It is wired +// only to the dense HIP launcher below, leaving CUDA and the separately tuned +// routed-MoE kernel on their established code paths. The intent is to give +// the AMD register allocator shorter temporary lifetimes without changing the +// numerical contract or packed GGUF layout. +static __device__ __forceinline__ float +vec_dot_q4_0_q8_1_hip_scalarized(const void* __restrict__ vbq, + const block_q8_1* __restrict__ bq8_1, + const int& iqs) { + const block_q4_0* bq4_0 = (const block_q4_0*)vbq; + const int v0 = get_int_from_uint8(bq4_0->qs, iqs); + const int v1 = get_int_from_uint8(bq4_0->qs, iqs + 1); + const int u0 = get_int_from_int8_aligned(bq8_1->qs, iqs); + const int u1 = get_int_from_int8_aligned(bq8_1->qs, iqs + QI4_0); + const int u2 = get_int_from_int8_aligned(bq8_1->qs, iqs + 1); + const int u3 = get_int_from_int8_aligned(bq8_1->qs, iqs + 1 + QI4_0); + int sumi = 0; + sumi = __dp4a(v0 & 0x0F0F0F0F, u0, sumi); + sumi = __dp4a((v0 >> 4) & 0x0F0F0F0F, u1, sumi); + sumi = __dp4a(v1 & 0x0F0F0F0F, u2, sumi); + sumi = __dp4a((v1 >> 4) & 0x0F0F0F0F, u3, sumi); + const float2 ds8f = __half22float2(bq8_1->ds); + return __half2float(bq4_0->d) * (sumi * ds8f.x - (8 * 2 / QI4_0) * ds8f.y); +} +#endif + template static __device__ __forceinline__ void allocate_tiles_q4_0(int** x_ql, half2** x_dm, int** x_qh, int** x_sc) { __shared__ int tile_x_qs[mmq_y * (WARP_SIZE_GGUF) + mmq_y]; From 20e6c1a3675972dd4c63425030d38dd9e0776d43 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 17:47:59 -0700 Subject: [PATCH 62/72] Revert "perf(rocm): scalarize dense Q4 dot temporaries" This reverts commit 4bffe92031fa6f0e41672dba6c6aaef3c3a12ddd. --- python/freetoken/kernel/csrc/gguf/mmvq.cuh | 8 ----- python/freetoken/kernel/csrc/gguf/vecdotq.cuh | 31 ------------------- 2 files changed, 39 deletions(-) diff --git a/python/freetoken/kernel/csrc/gguf/mmvq.cuh b/python/freetoken/kernel/csrc/gguf/mmvq.cuh index 6e6b703a..7331731a 100644 --- a/python/freetoken/kernel/csrc/gguf/mmvq.cuh +++ b/python/freetoken/kernel/csrc/gguf/mmvq.cuh @@ -59,16 +59,8 @@ static void mul_mat_vec_q4_0_q8_1_cuda( const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y; const dim3 block_nums(block_num_y, nvecs, 1); const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1); -#if defined USE_ROCM - // The scalarized helper is HIP-only and keeps the CUDA code object stable. - // It implements the same Q4_0/Q8_1 dot-product sequence with shorter-lived - // temporaries for an isolated gfx1151 dense-decode experiment. - mul_mat_vec_q - <<>>(vx, vy, dst, ncols, nrows, nvecs); -#else mul_mat_vec_q <<>>(vx, vy, dst, ncols, nrows, nvecs); -#endif } template diff --git a/python/freetoken/kernel/csrc/gguf/vecdotq.cuh b/python/freetoken/kernel/csrc/gguf/vecdotq.cuh index e9312707..08b4cd26 100644 --- a/python/freetoken/kernel/csrc/gguf/vecdotq.cuh +++ b/python/freetoken/kernel/csrc/gguf/vecdotq.cuh @@ -552,37 +552,6 @@ vec_dot_q4_0_q8_1(const void* __restrict__ vbq, const block_q8_1* __restrict__ b return vec_dot_q4_0_q8_1_impl(v, u, __half2float(bq4_0->d), bq8_1->ds); } -#if defined USE_ROCM -// HIP-only dense-GEMV variant of the Q4_0 x Q8_1 inner product. The generic -// helper above stages two packed Q4 words and four Q8 words in local arrays -// before forwarding them to a template helper. A decode lane always consumes -// exactly those six words, so this variant keeps them as named scalars and -// executes the same four DP4A instructions in the same order. It is wired -// only to the dense HIP launcher below, leaving CUDA and the separately tuned -// routed-MoE kernel on their established code paths. The intent is to give -// the AMD register allocator shorter temporary lifetimes without changing the -// numerical contract or packed GGUF layout. -static __device__ __forceinline__ float -vec_dot_q4_0_q8_1_hip_scalarized(const void* __restrict__ vbq, - const block_q8_1* __restrict__ bq8_1, - const int& iqs) { - const block_q4_0* bq4_0 = (const block_q4_0*)vbq; - const int v0 = get_int_from_uint8(bq4_0->qs, iqs); - const int v1 = get_int_from_uint8(bq4_0->qs, iqs + 1); - const int u0 = get_int_from_int8_aligned(bq8_1->qs, iqs); - const int u1 = get_int_from_int8_aligned(bq8_1->qs, iqs + QI4_0); - const int u2 = get_int_from_int8_aligned(bq8_1->qs, iqs + 1); - const int u3 = get_int_from_int8_aligned(bq8_1->qs, iqs + 1 + QI4_0); - int sumi = 0; - sumi = __dp4a(v0 & 0x0F0F0F0F, u0, sumi); - sumi = __dp4a((v0 >> 4) & 0x0F0F0F0F, u1, sumi); - sumi = __dp4a(v1 & 0x0F0F0F0F, u2, sumi); - sumi = __dp4a((v1 >> 4) & 0x0F0F0F0F, u3, sumi); - const float2 ds8f = __half22float2(bq8_1->ds); - return __half2float(bq4_0->d) * (sumi * ds8f.x - (8 * 2 / QI4_0) * ds8f.y); -} -#endif - template static __device__ __forceinline__ void allocate_tiles_q4_0(int** x_ql, half2** x_dm, int** x_qh, int** x_sc) { __shared__ int tile_x_qs[mmq_y * (WARP_SIZE_GGUF) + mmq_y]; From d8f5e99e63800c36c2aa566ac9625d416a53624e Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 17:48:23 -0700 Subject: [PATCH 63/72] docs(rocm): record scalarized Q4 candidate result --- docs/lan223-rocm-validation-2026-08-28.md | 35 +++++++++++++++++++++++ 1 file changed, 35 insertions(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index 0cb156cb..acc9165d 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -329,6 +329,41 @@ The retained raw evidence is: /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 From 18445f0e50d64fe6af9f53d1ef697ce1c8270992 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 17:53:18 -0700 Subject: [PATCH 64/72] perf(rocm): reuse Q8 words in two-row MoE --- python/freetoken/kernel/csrc/gguf/moe_vec.cuh | 39 ++++++++++++++++++- 1 file changed, 37 insertions(+), 2 deletions(-) diff --git a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh index 3d12f1e9..1287305f 100644 --- a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh +++ b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh @@ -52,6 +52,31 @@ static __global__ void moe_vec_q( } #if defined(USE_ROCM) +// Compute one Q4_0 row against four Q8_1 packed words that the caller has +// already loaded. The two-row HIP kernel invokes this twice per lane with two +// adjacent Q4_0 rows but one common activation block. Keeping the activation +// words in the caller avoids issuing the same four global Q8 loads for both +// rows. The DP4A order and Q4_0 scale correction match vec_dot_q4_0_q8_1, +// preserving the exact native GGUF arithmetic contract. +static __device__ __forceinline__ float moe_vec_q4_0_dot_shared_q8( + const block_q4_0* __restrict__ row, + const int u0, + const int u1, + const int u2, + const int u3, + const half2 ds8, + const int iqs) { + const int v0 = get_int_from_uint8(row->qs, iqs); + const int v1 = get_int_from_uint8(row->qs, iqs + 1); + int sumi = 0; + sumi = __dp4a(v0 & 0x0F0F0F0F, u0, sumi); + sumi = __dp4a((v0 >> 4) & 0x0F0F0F0F, u1, sumi); + sumi = __dp4a(v1 & 0x0F0F0F0F, u2, sumi); + sumi = __dp4a((v1 >> 4) & 0x0F0F0F0F, u3, sumi); + const float2 ds8f = __half22float2(ds8); + return __half2float(row->d) * (sumi * ds8f.x - (8 * 2 / QI4_0) * ds8f.y); +} + // The HIP launcher is defined after the CUDA-compatible wrapper so the // generic wrappers remain grouped by quantization format below. template @@ -136,9 +161,19 @@ static __global__ void moe_vec_q4_0_hip_two_rows( 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); + // Both output rows consume one identical Q8_1 activation block. Load its + // four packed words once per lane and pass them to the two row products. + // This is deliberately confined to the HIP two-row specialization: the + // generic CUDA and non-Q4 routes retain their established helper calls. + const block_q8_1* y_block = &y[iby]; + const int u0 = get_int_from_int8_aligned(y_block->qs, iqs); + const int u1 = get_int_from_int8_aligned(y_block->qs, iqs + QI4_0); + const int u2 = get_int_from_int8_aligned(y_block->qs, iqs + 1); + const int u3 = get_int_from_int8_aligned(y_block->qs, iqs + 1 + QI4_0); + tmp0 += moe_vec_q4_0_dot_shared_q8(&x[row0 * blocks_per_row + i], u0, u1, u2, u3, y_block->ds, iqs); if (row0 + 1 < nrows) { - tmp1 += vec_dot_q4_0_q8_1(&x[(row0 + 1) * blocks_per_row + i], &y[iby], iqs); + tmp1 += moe_vec_q4_0_dot_shared_q8( + &x[(row0 + 1) * blocks_per_row + i], u0, u1, u2, u3, y_block->ds, iqs); } } From 2a9c264cd778e3f2d3e8e0b52dc72fe1650da606 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 18:03:31 -0700 Subject: [PATCH 65/72] Revert "perf(rocm): reuse Q8 words in two-row MoE" This reverts commit 684148d65e0e1b4677d572b660adeeaa00ccbc66. --- python/freetoken/kernel/csrc/gguf/moe_vec.cuh | 39 +------------------ 1 file changed, 2 insertions(+), 37 deletions(-) diff --git a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh index 1287305f..3d12f1e9 100644 --- a/python/freetoken/kernel/csrc/gguf/moe_vec.cuh +++ b/python/freetoken/kernel/csrc/gguf/moe_vec.cuh @@ -52,31 +52,6 @@ static __global__ void moe_vec_q( } #if defined(USE_ROCM) -// Compute one Q4_0 row against four Q8_1 packed words that the caller has -// already loaded. The two-row HIP kernel invokes this twice per lane with two -// adjacent Q4_0 rows but one common activation block. Keeping the activation -// words in the caller avoids issuing the same four global Q8 loads for both -// rows. The DP4A order and Q4_0 scale correction match vec_dot_q4_0_q8_1, -// preserving the exact native GGUF arithmetic contract. -static __device__ __forceinline__ float moe_vec_q4_0_dot_shared_q8( - const block_q4_0* __restrict__ row, - const int u0, - const int u1, - const int u2, - const int u3, - const half2 ds8, - const int iqs) { - const int v0 = get_int_from_uint8(row->qs, iqs); - const int v1 = get_int_from_uint8(row->qs, iqs + 1); - int sumi = 0; - sumi = __dp4a(v0 & 0x0F0F0F0F, u0, sumi); - sumi = __dp4a((v0 >> 4) & 0x0F0F0F0F, u1, sumi); - sumi = __dp4a(v1 & 0x0F0F0F0F, u2, sumi); - sumi = __dp4a((v1 >> 4) & 0x0F0F0F0F, u3, sumi); - const float2 ds8f = __half22float2(ds8); - return __half2float(row->d) * (sumi * ds8f.x - (8 * 2 / QI4_0) * ds8f.y); -} - // The HIP launcher is defined after the CUDA-compatible wrapper so the // generic wrappers remain grouped by quantization format below. template @@ -161,19 +136,9 @@ static __global__ void moe_vec_q4_0_hip_two_rows( 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)); - // Both output rows consume one identical Q8_1 activation block. Load its - // four packed words once per lane and pass them to the two row products. - // This is deliberately confined to the HIP two-row specialization: the - // generic CUDA and non-Q4 routes retain their established helper calls. - const block_q8_1* y_block = &y[iby]; - const int u0 = get_int_from_int8_aligned(y_block->qs, iqs); - const int u1 = get_int_from_int8_aligned(y_block->qs, iqs + QI4_0); - const int u2 = get_int_from_int8_aligned(y_block->qs, iqs + 1); - const int u3 = get_int_from_int8_aligned(y_block->qs, iqs + 1 + QI4_0); - tmp0 += moe_vec_q4_0_dot_shared_q8(&x[row0 * blocks_per_row + i], u0, u1, u2, u3, y_block->ds, iqs); + tmp0 += vec_dot_q4_0_q8_1(&x[row0 * blocks_per_row + i], &y[iby], iqs); if (row0 + 1 < nrows) { - tmp1 += moe_vec_q4_0_dot_shared_q8( - &x[(row0 + 1) * blocks_per_row + i], u0, u1, u2, u3, y_block->ds, iqs); + tmp1 += vec_dot_q4_0_q8_1(&x[(row0 + 1) * blocks_per_row + i], &y[iby], iqs); } } From 42d5b7950ebb73b03ce46b5352ddd4f3a290fc1e Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 18:04:05 -0700 Subject: [PATCH 66/72] docs(rocm): reject Q8 reuse MoE candidate --- docs/lan223-rocm-validation-2026-08-28.md | 39 +++++++++++++++++++++++ 1 file changed, 39 insertions(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index acc9165d..eb4e6552 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -222,6 +222,45 @@ Artifacts are retained on LAN-223: /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 From e13eac5cecd88e78f8d9b443475b4cf56fa9b663 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 18:22:23 -0700 Subject: [PATCH 67/72] bench: expose explicit KV token capacity --- benchmarks/bench_decode_moe.py | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/benchmarks/bench_decode_moe.py b/benchmarks/bench_decode_moe.py index 56621792..ccbdcc94 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,11 @@ 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: + cmd += ["--num-token-override", str(args.num_token_override)] if args.cache > 0: cmd += ["--moe-cache-size", str(args.cache)] elif args.cache_rate is not None: From 6cb6f395e50fe9e732afae6c60080ff69589c83a Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 18:23:13 -0700 Subject: [PATCH 68/72] fix(bench): use server num-tokens flag --- benchmarks/bench_decode_moe.py | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/benchmarks/bench_decode_moe.py b/benchmarks/bench_decode_moe.py index ccbdcc94..313af333 100644 --- a/benchmarks/bench_decode_moe.py +++ b/benchmarks/bench_decode_moe.py @@ -200,7 +200,9 @@ def serve_cmd(args: argparse.Namespace, backend: str, port: int) -> list[str]: # 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: - cmd += ["--num-token-override", str(args.num_token_override)] + # 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: From 1a0033fea03299e515486ba2dcc95e2cdc26c1bb Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 18:34:20 -0700 Subject: [PATCH 69/72] docs(rocm): record full cache validation --- benchmarks/README.md | 5 ++ docs/lan223-rocm-validation-2026-08-28.md | 70 +++++++++++++++++++++-- 2 files changed, 71 insertions(+), 4 deletions(-) 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/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index eb4e6552..9492b78d 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -166,13 +166,75 @@ prompt tokens and llama.cpp's reused 58. | 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 | 60.42 client, 58.88 internal | 128.6 ms | Matched reference | The graph configuration removes approximately 1.5 percent of the eager decode -cost, but FreeToken still trails llama.cpp by 7.8 percent using client TPS and -by approximately 5.4 percent compared with llama.cpp's internal decode timing. -The requested criterion of meeting or exceeding llama.cpp is therefore **not -met** by the first configuration pass. +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/ +``` ### Accepted HIP Q4_0 one-wave/two-row MoE specialization From 0745a53abc317be0305803a013dcada9007c35ac Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 18:39:56 -0700 Subject: [PATCH 70/72] docs(rocm): add current llama control --- docs/lan223-rocm-validation-2026-08-28.md | 45 ++++++++++++++++++++++- 1 file changed, 44 insertions(+), 1 deletion(-) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index 9492b78d..ee413e80 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -167,7 +167,8 @@ prompt tokens and llama.cpp's reused 58. | 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 | 60.42 client, 58.88 internal | 128.6 ms | Matched reference | +| 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 @@ -236,6 +237,48 @@ The retained raw evidence is: /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/ +``` + ### Accepted HIP Q4_0 one-wave/two-row MoE specialization The first two-row experiment did not reproduce llama.cpp's execution shape: it From b8e48bd494e46879131e19c989e7eb86ad8b1928 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 18:43:05 -0700 Subject: [PATCH 71/72] docs(rocm): reconcile current operating guidance --- docs/lan223-rocm-validation-2026-08-28.md | 40 +++++++++++++---------- 1 file changed, 23 insertions(+), 17 deletions(-) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index ee413e80..d024b06f 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -571,9 +571,9 @@ 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-auto \ - --memory-ratio 0.50 --max-running-requests 1 --max-seq-len-override 8320 \ - --cuda-graph-max-bs 1 + --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 @@ -583,20 +583,26 @@ 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 automatic -cache selected all 3,840 routed-expert slots. Closing the gap requires a -profile-guided improvement to the HIP GGUF decode kernels or another proven -ROCm attention or quantized-linear implementation. The available ROCm 10 -`rocprofv3` installation could not yet provide that kernel breakdown: attach -mode reports that the PyTorch process has no `rocp-bg-attach` registration -thread even when launched with `ROCP_TOOL_ATTACH=1`, while launch mode aborts -before FreeToken starts with LLVM's duplicate `spirv-expand-step` option. The -full error evidence is retained in `rocprof-gfx1151*/` and -`rocprof-launch-gfx1151-v2/` under the raw artifact directory. This is a -toolchain issue, not a FreeToken performance result, so no profiler-derived -optimization claim is made here. 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. +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. + +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: From 22bcf7fd91b0f783961943377ecbcd845a019a79 Mon Sep 17 00:00:00 2001 From: David Date: Fri, 28 Aug 2026 18:51:24 -0700 Subject: [PATCH 72/72] docs(rocm): record current-main API revalidation --- docs/lan223-rocm-validation-2026-08-28.md | 42 +++++++++++++++++++++++ 1 file changed, 42 insertions(+) diff --git a/docs/lan223-rocm-validation-2026-08-28.md b/docs/lan223-rocm-validation-2026-08-28.md index d024b06f..54ce44ea 100644 --- a/docs/lan223-rocm-validation-2026-08-28.md +++ b/docs/lan223-rocm-validation-2026-08-28.md @@ -279,6 +279,48 @@ retained on LAN-223 at: /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/ +``` + ### Accepted HIP Q4_0 one-wave/two-row MoE specialization The first two-row experiment did not reproduce llama.cpp's execution shape: it