From e3ae139d874e33f700d0fe74687e22a26d0cf4aa Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Tue, 23 Jun 2026 17:21:52 -0600 Subject: [PATCH 01/17] [Bench] Add dense hybrid_w4a16 GEMM CLI benchmark Add benchmarks/kernels/benchmark_hybrid_w4a16.py, a CLI entry point for the dense W4A16 prefill GEMM that reuses the perf regression test's measurement rig (packed-int4 weights via pack_skinny_int4, cold-weight MALL rotation, do_bench_cudagraph median timing). The reported time/call matches a real prefill forward pass and the torch profiler trace. Defaults reproduce the Qwen3-VL-4B-Instruct-AWQ-4bit fused gate_up_proj GEMM at the 1322-token prefill shape (M=1322, N=19456, K=2560, group_size=32, symmetric, bf16): 23.5 TFLOP/s / 5.6 ms on gfx1151, matching the in-model profile within bench variance. Changes: - --model qwen3vl-4b-awq preset sweeps all four decoder GEMMs. - Per-shape overrides (--k/--n/--group-size/--provider/--batches) for ad-hoc shapes; provider follows the test's hybrid-w4a16[-zp][-bf16] convention. Co-authored-by: Claude Signed-off-by: Matthias Gehre --- benchmarks/kernels/benchmark_hybrid_w4a16.py | 130 +++++++++++++++++++ 1 file changed, 130 insertions(+) create mode 100644 benchmarks/kernels/benchmark_hybrid_w4a16.py diff --git a/benchmarks/kernels/benchmark_hybrid_w4a16.py b/benchmarks/kernels/benchmark_hybrid_w4a16.py new file mode 100644 index 000000000000..b5f46d4f50a5 --- /dev/null +++ b/benchmarks/kernels/benchmark_hybrid_w4a16.py @@ -0,0 +1,130 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""CLI benchmark for the dense hybrid_w4a16 GEMM kernel. + +Drives the same measurement rig as the perf regression test +(``tests/kernels/quantization/test_hybrid_w4a16_perf.py``): packed int4 +weights via the production ``pack_skinny_int4`` helper, cold-weight MALL +rotation, and ``do_bench_cudagraph`` median timing -- so the reported +time/call matches what a real prefill forward pass sees (and what shows up +in a torch profiler trace). + +Defaults reproduce the Qwen3-VL-4B-Instruct-AWQ-4bit fused gate_up_proj +GEMM at the 1322-token prefill shape from the gfx1151 roofline analysis: +M=1322, N=19456, K=2560, group_size=32, symmetric (no zero-point), bf16. + +Usage:: + + # Default: the up_proj shape at M=1322 + python benchmarks/kernels/benchmark_hybrid_w4a16.py + + # All Qwen3-VL-4B AWQ decoder GEMMs across a batch sweep + python benchmarks/kernels/benchmark_hybrid_w4a16.py --model qwen3vl-4b-awq + + # A custom shape / quant config + python benchmarks/kernels/benchmark_hybrid_w4a16.py \\ + --k 2560 --n 19456 --group-size 32 --provider hybrid-w4a16-bf16 \\ + --batches 1024 1322 2048 + +The provider name follows the test convention +``hybrid-w4a16[-zp][-bf16]``: ``-zp`` selects the asymmetric per-group +zero-point dequant path; ``-bf16`` runs with bfloat16 activations/scales +instead of float16. Symmetric fp16 is the bare ``hybrid-w4a16``. +""" + +from __future__ import annotations + +import argparse +import sys +from pathlib import Path + +import torch + +# Make the in-tree tests helper importable from a stock checkout, mirroring +# bench_hybrid_w4a16_moe.py. +REPO_ROOT = Path(__file__).resolve().parents[2] +sys.path.insert(0, str(REPO_ROOT)) + +from tests.kernels.quantization.test_hybrid_w4a16_perf import ( # noqa: E402 + PROVIDERS, + _provider_dtype, + measure_tflops, + prepare_hybrid_weights, +) + +# Named (K, N) decoder GEMM presets. K=in_features, N=out_features. +# qwen3vl-4b-awq: group_size=32, symmetric, bf16 (see model config). +MODEL_SHAPES: dict[str, list[dict]] = { + "qwen3vl-4b-awq": [ + {"k": 2560, "n": 19456, "group_size": 32, "comment": "gate_up_proj"}, + {"k": 9728, "n": 2560, "group_size": 32, "comment": "down_proj"}, + {"k": 2560, "n": 6144, "group_size": 32, "comment": "qkv_proj"}, + {"k": 2560, "n": 4096, "group_size": 32, "comment": "o_proj"}, + ], +} + + +def _bench_one( + m: int, k: int, n: int, group_size: int, provider: str +) -> tuple[str, float, float]: + """Return (kernel, tflops, ms) for one shape/provider/batch.""" + weights = prepare_hybrid_weights(k, n, group_size, dtype=_provider_dtype(provider)) + measure_tflops(m, weights, k, n, group_size, provider) # warmup + kernel, tflops = measure_tflops(m, weights, k, n, group_size, provider) + ms = (2 * m * n * k) * 1e-12 / (tflops * 1e-3) + return kernel, tflops, ms + + +def main() -> None: + p = argparse.ArgumentParser( + description=__doc__, + formatter_class=argparse.RawDescriptionHelpFormatter, + ) + p.add_argument( + "--model", + choices=sorted(MODEL_SHAPES), + help="Benchmark all decoder GEMMs for a named model preset.", + ) + p.add_argument("--k", type=int, default=2560, help="in_features (default 2560)") + p.add_argument("--n", type=int, default=19456, help="out_features (default 19456)") + p.add_argument("--group-size", type=int, default=32, help="quant group size") + p.add_argument( + "--provider", + default="hybrid-w4a16-bf16", + choices=PROVIDERS, + help="quant variant (default hybrid-w4a16-bf16: bf16, symmetric)", + ) + p.add_argument( + "--batches", + type=int, + nargs="+", + default=[1322], + help="M (token) batch sizes to sweep (default: 1322)", + ) + args = p.parse_args() + + if args.model: + shapes = MODEL_SHAPES[args.model] + else: + shapes = [ + {"k": args.k, "n": args.n, "group_size": args.group_size, "comment": ""} + ] + + hdr = f"{'shape (MxNxK)':>22} {'g':>4} {'kernel':>22} {'TFLOP/s':>9} {'ms/call':>9}" + print(hdr) + print("-" * len(hdr)) + for s in shapes: + k, n, gs = s["k"], s["n"], s["group_size"] + for m in args.batches: + kernel, tflops, ms = _bench_one(m, k, n, gs, args.provider) + tag = f" # {s['comment']}" if s["comment"] else "" + shape_str = f"{m}x{n}x{k}" + print( + f"{shape_str:>22} {gs:>4} {kernel:>22} {tflops:>9.2f} {ms:>9.3f}{tag}" + ) + torch.accelerator.synchronize() + + +if __name__ == "__main__": + main() From 51820ce8237dd015ac69a88a91a832fad648182c Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Wed, 24 Jun 2026 00:46:50 -0600 Subject: [PATCH 02/17] [ROCm] Hand-editable AMDGCN path for the W4A16 skinny GEMM (gfx1151) Add an opt-in workflow to directly hand-tune the ISA of the W4A16 skinny prefill GEMM instead of going through the Triton JIT. The .amdgcn lives in the source tree, is assembled to a .hsaco at vLLM build time, and is dispatched to at runtime via a direct HIP module launch when the selected tile/config exactly matches what the ISA was built for. Pinned config (Qwen3-VL-4B-AWQ up/gate_proj class on gfx1151): bf16, symmetric, group_size=32, BLOCK_M=128/BLOCK_N=64/BLOCK_K=32, 8 warps. Verified numerically correct (cos 0.999996 vs torch reference) at M=1322,N=19456,K=2560. Changes: - asm/hybrid_w4a16_skinny_gfx1151.amdgcn: the hand-editable ISA (dumped from the Triton kernel) and .meta.json pinning the kernarg ABI + constexpr config. - asm_w4a16.py: ctypes/libamdhip64 loader that packs the 88-byte kernarg buffer per the pinned layout and launches via hipModuleLaunchKernel; assembles the .amdgcn on demand (clang) when newer than the cached .hsaco, so the edit->rerun loop needs no full rebuild. - hybrid_w4a16.py: gate the gfx11 dispatch on VLLM_W4A16_HANDASM=1; route to the hand-asm kernel only when config_matches(), asserting every pinned field, else fall through to Triton unchanged. - CMakeLists.txt: assemble .amdgcn -> .hsaco at build time on gfx1151 HIP builds. - pyproject.toml: exclude *.amdgcn from the typos hook (asm mnemonics like "offen" are not prose). .gitignore: the generated .hsaco is a build artifact. Co-authored-by: Claude Signed-off-by: Matthias Gehre --- .gitignore | 3 + CMakeLists.txt | 19 + pyproject.toml | 3 +- .../asm/hybrid_w4a16_skinny_gfx1151.amdgcn | 1237 +++++++++++++++++ .../asm/hybrid_w4a16_skinny_gfx1151.meta.json | 32 + .../linear/mixed_precision/asm_w4a16.py | 268 ++++ .../linear/mixed_precision/hybrid_w4a16.py | 37 + 7 files changed, 1598 insertions(+), 1 deletion(-) create mode 100644 vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn create mode 100644 vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.meta.json create mode 100644 vllm/model_executor/kernels/linear/mixed_precision/asm_w4a16.py diff --git a/.gitignore b/.gitignore index 2c4e135e58dc..cf5ff4a763a7 100644 --- a/.gitignore +++ b/.gitignore @@ -249,3 +249,6 @@ vllm/grpc/vllm_engine_pb2.pyi # Ignore generated cpu headers csrc/cpu/cpu_attn_dispatch_generated.h + +# Hand-tuned AMDGCN build artifact (assembled from .amdgcn at build time) +*.hsaco diff --git a/CMakeLists.txt b/CMakeLists.txt index aae50f1a6dac..626fb5bad75b 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1376,6 +1376,25 @@ if(VLLM_GPU_LANG STREQUAL "HIP") endif() endif() +# Hand-editable AMDGCN kernels (gfx1151): assemble .amdgcn -> .hsaco at build +# time so editable installs (and wheels) ship a ready binary next to the source +# .amdgcn. The runtime loader (asm_w4a16.py) re-assembles on demand if the +# .amdgcn is edited without a rebuild, so this is a convenience, not a hard dep. +if (VLLM_GPU_LANG STREQUAL "HIP" AND VLLM_GPU_ARCHES MATCHES "gfx1151") + set(_HANDASM_DIR + ${CMAKE_CURRENT_SOURCE_DIR}/vllm/model_executor/kernels/linear/mixed_precision/asm) + set(_HANDASM_SRC ${_HANDASM_DIR}/hybrid_w4a16_skinny_gfx1151.amdgcn) + set(_HANDASM_OUT ${_HANDASM_DIR}/hybrid_w4a16_skinny_gfx1151.hsaco) + add_custom_command( + OUTPUT ${_HANDASM_OUT} + COMMAND ${CMAKE_HIP_COMPILER} -target amdgcn-amd-amdhsa -mcpu=gfx1151 + -x assembler ${_HANDASM_SRC} -o ${_HANDASM_OUT} + DEPENDS ${_HANDASM_SRC} + COMMENT "Assembling hand-tuned W4A16 AMDGCN -> hsaco (gfx1151)" + VERBATIM) + add_custom_target(vllm_handasm_w4a16 ALL DEPENDS ${_HANDASM_OUT}) +endif() + # Must run after the last HIP `define_extension_target` so every extension # has registered its sources. if (VLLM_GPU_LANG STREQUAL "HIP") diff --git a/pyproject.toml b/pyproject.toml index 117bc925e9dd..2399eb5b44a3 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -131,7 +131,8 @@ extend-exclude = ["tests/models/fixtures/*", "tests/prompts/*", "tests/tokenizer "tests/v1/engine/test_fast_incdec_prefix_err.py", ".git/*", "csrc/cpu/sgl-kernels/*", "rust/src/chat/src/renderer/deepseek_v32/fixtures/*", "rust/src/tool-parser/src/gemma4.rs", "rust/src/text/src/output/decoded.rs", - "rust/src/tokenizer/src/incremental.rs", "rust/src/reasoning-parser/src/tests.rs"] + "rust/src/tokenizer/src/incremental.rs", "rust/src/reasoning-parser/src/tests.rs", + "*.amdgcn"] ignore-hidden = false [tool.typos.default] diff --git a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn new file mode 100644 index 000000000000..cc169cbda7db --- /dev/null +++ b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn @@ -0,0 +1,1237 @@ + .amdgcn_target "amdgcn-amd-amdhsa--gfx1151" + .amdhsa_code_object_version 5 + .text + .globl _triton_w4a16_skinny_fmt_kernel ; -- Begin function _triton_w4a16_skinny_fmt_kernel + .p2align 8 + .type _triton_w4a16_skinny_fmt_kernel,@function +_triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel +.Lfunc_begin0: + .file 1 "/scratch/mgehre/vllm/vllm/model_executor/kernels/linear/mixed_precision" "hybrid_w4a16.py" + .loc 1 81 0 ; hybrid_w4a16.py:81:0 + .cfi_sections .debug_frame + .cfi_startproc +; %bb.0: + s_clause 0x1 + s_load_b32 s26, s[0:1], 0x30 + s_load_b128 s[20:23], s[0:1], 0x20 + v_dual_mov_b32 v8, 0 :: v_dual_and_b32 v33, 15, v0 +.Ltmp0: + .loc 1 125 44 prologue_end ; hybrid_w4a16.py:125:44 + v_and_b32_e32 v34, 0xc0, v0 + .loc 1 126 44 ; hybrid_w4a16.py:126:44 + v_and_b32_e32 v35, 32, v0 + .loc 1 125 21 ; hybrid_w4a16.py:125:21 + s_lshl_b32 s24, s2, 7 + s_delay_alu instid0(VALU_DEP_3) + v_mov_b32_e32 v7, v8 + v_mov_b32_e32 v6, v8 + v_mov_b32_e32 v5, v8 + v_mov_b32_e32 v4, v8 + v_mov_b32_e32 v3, v8 + v_mov_b32_e32 v2, v8 + v_mov_b32_e32 v1, v8 + v_mov_b32_e32 v16, v8 + v_mov_b32_e32 v15, v8 + v_mov_b32_e32 v14, v8 + v_mov_b32_e32 v13, v8 + v_mov_b32_e32 v12, v8 + v_mov_b32_e32 v11, v8 + v_mov_b32_e32 v10, v8 + v_mov_b32_e32 v9, v8 + v_mov_b32_e32 v24, v8 + v_mov_b32_e32 v23, v8 + v_mov_b32_e32 v22, v8 + v_mov_b32_e32 v21, v8 + v_mov_b32_e32 v20, v8 + v_mov_b32_e32 v19, v8 + v_mov_b32_e32 v18, v8 + v_mov_b32_e32 v17, v8 + v_mov_b32_e32 v32, v8 + v_mov_b32_e32 v31, v8 + v_mov_b32_e32 v30, v8 + v_mov_b32_e32 v29, v8 + v_mov_b32_e32 v28, v8 + v_mov_b32_e32 v27, v8 + v_mov_b32_e32 v26, v8 + v_mov_b32_e32 v25, v8 +.Ltmp1: + .file 2 "/scratch/mgehre/vllm/.venv/lib/python3.12/site-packages/triton/language" "standard.py" + .loc 2 43 17 ; standard.py:43:17 @[ hybrid_w4a16.py:141:39 ] + s_waitcnt lgkmcnt(0) + s_add_i32 s2, s26, 31 +.Ltmp2: + .loc 1 126 21 ; hybrid_w4a16.py:126:21 + s_lshl_b32 s25, s3, 6 + .loc 1 141 28 ; hybrid_w4a16.py:141:28 + s_cmp_lt_i32 s2, 32 + s_cbranch_scc1 .LBB0_3 +; %bb.1: ; %.lr.ph + .loc 1 0 28 is_stmt 0 ; hybrid_w4a16.py:0:28 + s_clause 0x2 + s_load_b128 s[4:7], s[0:1], 0x34 + s_load_b128 s[16:19], s[0:1], 0x0 + s_load_b64 s[8:9], s[0:1], 0x10 + .loc 1 125 44 is_stmt 1 ; hybrid_w4a16.py:125:44 + v_lshrrev_b32_e32 v1, 2, v0 +.Ltmp3: + .loc 2 43 30 ; standard.py:43:30 @[ hybrid_w4a16.py:141:39 ] + s_ashr_i32 s0, s2, 31 + v_dual_mov_b32 v25, 0 :: v_dual_and_b32 v36, 3, v0 + s_lshr_b32 s0, s0, 27 + s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2) +.Ltmp4: + .loc 1 125 44 ; hybrid_w4a16.py:125:44 + v_or_b32_e32 v5, s24, v1 +.Ltmp5: + .loc 2 43 30 ; standard.py:43:30 @[ hybrid_w4a16.py:141:39 ] + s_add_i32 s2, s2, s0 + v_dual_mov_b32 v29, v25 :: v_dual_lshlrev_b32 v4, 3, v0 + v_dual_mov_b32 v26, v25 :: v_dual_lshlrev_b32 v3, 1, v0 + s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4) +.Ltmp6: + .loc 1 125 31 ; hybrid_w4a16.py:125:31 + v_or_b32_e32 v8, 64, v5 + .loc 1 147 36 ; hybrid_w4a16.py:147:36 + v_cmp_gt_i32_e32 vcc_lo, s22, v5 + v_dual_mov_b32 v31, v25 :: v_dual_and_b32 v4, 48, v4 + v_dual_mov_b32 v27, v25 :: v_dual_lshlrev_b32 v2, 4, v0 + s_waitcnt lgkmcnt(0) + s_abs_i32 s27, s7 + .loc 1 126 31 ; hybrid_w4a16.py:126:31 + v_or_b32_e32 v6, s25, v1 + s_cvt_f32_u32 s0, s27 + v_dual_mov_b32 v30, v25 :: v_dual_and_b32 v3, 48, v3 + v_dual_mov_b32 v32, v25 :: v_dual_add_nc_u32 v7, s25, v1 + s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_4) | instid1(VALU_DEP_3) + v_rcp_iflag_f32_e32 v5, s0 + v_dual_mov_b32 v18, v25 :: v_dual_add_nc_u32 v1, s24, v1 + v_lshl_or_b32 v4, v33, 6, v4 +.Ltmp7: + .loc 2 43 30 ; standard.py:43:30 @[ hybrid_w4a16.py:141:39 ] + s_ashr_i32 s28, s2, 5 + v_xor_b32_e32 v2, v2, v3 + v_dual_mov_b32 v20, v25 :: v_dual_add_nc_u32 v3, 64, v1 +.Ltmp8: + .loc 1 147 36 ; hybrid_w4a16.py:147:36 + v_cmp_gt_i32_e64 s1, s22, v8 + v_lshl_or_b32 v8, v34, 4, v4 + s_delay_alu instid0(TRANS32_DEP_1) + v_readfirstlane_b32 s2, v5 + v_lshl_or_b32 v4, v35, 5, v4 + .loc 1 141 28 ; hybrid_w4a16.py:141:28 + v_mul_lo_u32 v7, s5, v7 + v_mul_lo_u32 v1, s26, v1 + v_mul_lo_u32 v3, s26, v3 + s_mul_f32 s2, s2, 0x4f7ffffe + .loc 1 153 36 ; hybrid_w4a16.py:153:36 + v_cmp_gt_i32_e64 s0, s23, v6 + v_mul_lo_u32 v38, s6, v6 + v_xor_b32_e32 v5, 16, v8 + v_xor_b32_e32 v6, 32, v8 + v_xor_b32_e32 v9, 48, v8 + v_xor_b32_e32 v10, 16, v4 + v_xor_b32_e32 v11, 32, v4 + v_xor_b32_e32 v12, 48, v4 + s_cvt_u32_f32 s2, s2 + v_dual_mov_b32 v22, v25 :: v_dual_lshlrev_b32 v13, 2, v36 + v_dual_mov_b32 v17, v25 :: v_dual_lshlrev_b32 v14, 4, v36 + s_sub_i32 s3, 0, s27 + v_dual_mov_b32 v28, v25 :: v_dual_lshlrev_b32 v37, 3, v36 + s_mul_i32 s3, s3, s2 + .loc 1 141 28 ; hybrid_w4a16.py:141:28 + v_lshl_add_u32 v39, v7, 2, v13 + v_lshl_add_u32 v40, v1, 1, v14 + v_lshl_add_u32 v41, v3, 1, v14 + v_dual_mov_b32 v19, v25 :: v_dual_add_nc_u32 v42, 0, v2 + v_dual_mov_b32 v24, v25 :: v_dual_add_nc_u32 v43, 0, v8 + v_dual_mov_b32 v21, v25 :: v_dual_add_nc_u32 v44, 0, v5 + v_add_nc_u32_e32 v45, 0, v6 + v_dual_mov_b32 v23, v25 :: v_dual_add_nc_u32 v46, 0, v9 + v_add_nc_u32_e32 v47, 0, v4 + v_dual_mov_b32 v9, v25 :: v_dual_add_nc_u32 v48, 0, v10 + v_dual_mov_b32 v14, v25 :: v_dual_add_nc_u32 v49, 0, v11 + v_dual_mov_b32 v11, v25 :: v_dual_add_nc_u32 v50, 0, v12 + v_mov_b32_e32 v10, v25 + v_mov_b32_e32 v12, v25 + v_mov_b32_e32 v13, v25 + v_mov_b32_e32 v15, v25 + v_mov_b32_e32 v16, v25 + v_mov_b32_e32 v1, v25 + v_mov_b32_e32 v2, v25 + v_mov_b32_e32 v3, v25 + v_mov_b32_e32 v4, v25 + v_mov_b32_e32 v5, v25 + v_mov_b32_e32 v6, v25 + v_mov_b32_e32 v7, v25 + v_mov_b32_e32 v8, v25 + s_mov_b32 s15, 0x31027000 + s_mov_b32 s14, 0x7ffffffe + s_and_b32 s17, s17, 0xffff + s_and_b32 s19, s19, 0xffff + s_mul_hi_u32 s3, s2, s3 + s_and_b32 s9, s9, 0xffff + s_ashr_i32 s6, s7, 31 + s_mov_b32 s5, 0 + s_add_i32 s7, s2, s3 + s_mov_b32 s12, s16 + s_mov_b32 s13, s17 + s_mov_b32 s16, s18 + s_mov_b32 s17, s19 + s_mov_b32 s18, s14 + s_mov_b32 s19, s15 + s_mov_b32 s10, s14 + s_mov_b32 s11, s15 +.LBB0_2: ; =>This Inner Loop Header: Depth=1 + .loc 1 143 26 ; hybrid_w4a16.py:143:26 + v_add_nc_u32_e32 v51, s5, v37 + .loc 1 153 61 ; hybrid_w4a16.py:153:61 + v_cmp_gt_i32_e64 s2, s4, v36 + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_ashr_i32 s3, s5, 31 + s_abs_i32 s29, s5 + s_xor_b32 s31, s3, s6 + .loc 1 143 26 ; hybrid_w4a16.py:143:26 + v_cmp_gt_i32_e64 s3, s26, v51 + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_mul_hi_u32 s30, s29, s7 + .loc 1 153 42 ; hybrid_w4a16.py:153:42 + s_and_b32 s2, s0, s2 + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_mul_i32 s33, s30, s27 + .loc 1 154 27 ; hybrid_w4a16.py:154:27 + v_cndmask_b32_e64 v51, 0x80000000, v39, s2 + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_sub_i32 s29, s29, s33 + .loc 1 147 41 ; hybrid_w4a16.py:147:41 + s_and_b32 s2, vcc_lo, s3 + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_add_i32 s34, s30, 1 + s_sub_i32 s33, s29, s27 + .loc 1 148 20 ; hybrid_w4a16.py:148:20 + v_cndmask_b32_e64 v52, 0x80000000, v40, s2 + .loc 1 147 41 ; hybrid_w4a16.py:147:41 + s_and_b32 s2, s1, s3 + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_cmp_ge_u32 s29, s27 + .loc 1 148 20 ; hybrid_w4a16.py:148:20 + v_cndmask_b32_e64 v55, 0x80000000, v41, s2 + .loc 1 154 27 ; hybrid_w4a16.py:154:27 + buffer_load_b32 v59, v51, s[16:19], 0 offen + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_cselect_b32 s2, s34, s30 + s_cselect_b32 s3, s33, s29 + s_add_i32 s29, s2, 1 + s_cmp_ge_u32 s3, s27 + .loc 1 148 20 ; hybrid_w4a16.py:148:20 + buffer_load_b128 v[51:54], v52, s[12:15], 0 offen + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_cselect_b32 s2, s29, s2 + .loc 1 141 28 ; hybrid_w4a16.py:141:28 + v_add_nc_u32_e32 v36, 4, v36 + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_xor_b32 s2, s2, s31 + .loc 1 141 28 ; hybrid_w4a16.py:141:28 + v_add_nc_u32_e32 v39, 16, v39 + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_sub_i32 s2, s2, s31 + .loc 1 141 28 ; hybrid_w4a16.py:141:28 + v_add_nc_u32_e32 v40, 64, v40 + .loc 1 224 16 ; hybrid_w4a16.py:224:16 + v_add_lshl_u32 v56, s2, v38, 1 + .loc 1 141 28 ; hybrid_w4a16.py:141:28 + v_add_nc_u32_e32 v41, 64, v41 + s_add_i32 s28, s28, -1 + s_add_i32 s5, s5, 32 + s_cmp_lg_u32 s28, 0 + .loc 1 224 16 ; hybrid_w4a16.py:224:16 + v_cndmask_b32_e64 v56, 0x80000000, v56, s0 + buffer_load_u16 v83, v56, s[8:11], 0 offen + .loc 1 148 20 ; hybrid_w4a16.py:148:20 + buffer_load_b128 v[55:58], v55, s[12:15], 0 offen + s_waitcnt vmcnt(0) + s_barrier + buffer_gl0_inv + ds_store_b128 v42, v[51:54] + ds_store_b128 v42, v[55:58] offset:4096 + .loc 1 190 22 ; hybrid_w4a16.py:190:22 + v_lshrrev_b32_e32 v60, 28, v59 + .loc 1 190 37 is_stmt 0 ; hybrid_w4a16.py:190:37 + v_and_b32_e32 v61, 15, v59 + v_bfe_u32 v62, v59, 16, 4 + v_bfe_u32 v63, v59, 4, 4 + v_bfe_u32 v64, v59, 20, 4 + v_bfe_u32 v65, v59, 8, 4 + v_bfe_u32 v66, v59, 24, 4 + v_bfe_u32 v59, v59, 12, 4 + .loc 1 234 28 is_stmt 1 ; hybrid_w4a16.py:234:28 + v_add_nc_u32_e32 v84, -8, v61 + v_add_nc_u32_e32 v85, -8, v62 + v_add_nc_u32_e32 v86, -8, v63 + v_add_nc_u32_e32 v87, -8, v64 + v_add_nc_u32_e32 v88, -8, v65 + v_add_nc_u32_e32 v89, -8, v66 + v_add_nc_u32_e32 v90, -8, v59 + v_add_nc_u32_e32 v91, -8, v60 + .loc 1 234 34 is_stmt 0 ; hybrid_w4a16.py:234:34 + v_cvt_f32_i32_e32 v84, v84 + v_cvt_f32_i32_e32 v85, v85 + v_cvt_f32_i32_e32 v86, v86 + v_cvt_f32_i32_e32 v87, v87 + v_cvt_f32_i32_e32 v88, v88 + v_cvt_f32_i32_e32 v89, v89 + v_cvt_f32_i32_e32 v90, v90 + v_cvt_f32_i32_e32 v91, v91 + v_bfe_u32 v92, v84, 16, 1 + v_bfe_u32 v93, v85, 16, 1 + v_bfe_u32 v94, v86, 16, 1 + v_bfe_u32 v95, v87, 16, 1 + v_bfe_u32 v96, v88, 16, 1 + v_bfe_u32 v97, v89, 16, 1 + v_bfe_u32 v98, v90, 16, 1 + v_bfe_u32 v99, v91, 16, 1 + v_add3_u32 v84, v84, v92, 0x7fff + v_add3_u32 v85, v85, v93, 0x7fff + v_add3_u32 v86, v86, v94, 0x7fff + v_add3_u32 v87, v87, v95, 0x7fff + v_add3_u32 v88, v88, v96, 0x7fff + v_add3_u32 v89, v89, v97, 0x7fff + v_add3_u32 v90, v90, v98, 0x7fff + v_add3_u32 v91, v91, v99, 0x7fff + .loc 1 224 16 is_stmt 1 ; hybrid_w4a16.py:224:16 + v_cndmask_b32_e64 v83, 0x3f80, v83, s0 + .loc 1 234 50 ; hybrid_w4a16.py:234:50 + v_alignbit_b32 v84, 0, v84, 16 + v_alignbit_b32 v85, 0, v85, 16 + v_alignbit_b32 v86, 0, v86, 16 + v_alignbit_b32 v87, 0, v87, 16 + v_alignbit_b32 v88, 0, v88, 16 + v_alignbit_b32 v89, 0, v89, 16 + v_alignbit_b32 v90, 0, v90, 16 + v_alignbit_b32 v91, 0, v91, 16 + v_perm_b32 v83, 0, v83, 0x5040100 + .loc 1 148 20 ; hybrid_w4a16.py:148:20 + s_waitcnt lgkmcnt(0) + s_barrier + buffer_gl0_inv + ds_load_b128 v[51:54], v43 + ds_load_b128 v[59:62], v43 offset:4096 + ds_load_b128 v[55:58], v44 + ds_load_b128 v[63:66], v44 offset:4096 + ds_load_b128 v[67:70], v45 + ds_load_b128 v[75:78], v45 offset:4096 + ds_load_b128 v[71:74], v46 + ds_load_b128 v[79:82], v46 offset:4096 + .loc 1 234 50 ; hybrid_w4a16.py:234:50 + v_dot2_bf16_bf16 v92, v86, v83, 0 + v_dot2_bf16_bf16 v88, v88, v83, 0 + v_dot2_bf16_bf16 v86, v90, v83, 0 + v_dot2_bf16_bf16 v90, v91, v83, 0 + v_dot2_bf16_bf16 v89, v89, v83, 0 + v_dot2_bf16_bf16 v87, v87, v83, 0 + v_dot2_bf16_bf16 v91, v84, v83, 0 + v_dot2_bf16_bf16 v83, v85, v83, 0 + .loc 1 240 33 ; hybrid_w4a16.py:240:33 + v_perm_b32 v86, v90, v86, 0x5040100 + v_perm_b32 v85, v89, v88, 0x5040100 + v_perm_b32 v84, v87, v92, 0x5040100 + s_waitcnt lgkmcnt(0) + v_perm_b32 v83, v83, v91, 0x5040100 + s_barrier + buffer_gl0_inv + ds_store_b128 v42, v[83:86] + s_waitcnt lgkmcnt(0) + s_barrier + buffer_gl0_inv + ds_load_b128 v[83:86], v47 + ds_load_b128 v[91:94], v47 offset:2048 + ds_load_b128 v[87:90], v48 + ds_load_b128 v[95:98], v48 offset:2048 + ds_load_b128 v[99:102], v49 + ds_load_b128 v[107:110], v49 offset:2048 + ds_load_b128 v[103:106], v50 + ds_load_b128 v[111:114], v50 offset:2048 + s_waitcnt lgkmcnt(5) + v_wmma_f32_16x16x16_bf16 v[25:32], v[83:90], v[51:58], v[25:32] + s_waitcnt lgkmcnt(4) + v_wmma_f32_16x16x16_bf16 v[17:24], v[91:98], v[51:58], v[17:24] + v_wmma_f32_16x16x16_bf16 v[9:16], v[83:90], v[59:66], v[9:16] + v_wmma_f32_16x16x16_bf16 v[1:8], v[91:98], v[59:66], v[1:8] + s_waitcnt lgkmcnt(1) + v_wmma_f32_16x16x16_bf16 v[25:32], v[99:106], v[67:74], v[25:32] + s_waitcnt lgkmcnt(0) + v_wmma_f32_16x16x16_bf16 v[17:24], v[107:114], v[67:74], v[17:24] + v_wmma_f32_16x16x16_bf16 v[9:16], v[99:106], v[75:82], v[9:16] + v_wmma_f32_16x16x16_bf16 v[1:8], v[107:114], v[75:82], v[1:8] + .loc 1 141 28 ; hybrid_w4a16.py:141:28 + s_cbranch_scc1 .LBB0_2 +.LBB0_3: ; %._crit_edge + .loc 1 125 44 ; hybrid_w4a16.py:125:44 + v_lshrrev_b32_e32 v34, 2, v34 + .loc 1 126 44 ; hybrid_w4a16.py:126:44 + v_lshrrev_b32_e32 v0, 4, v0 + v_lshrrev_b32_e32 v35, 1, v35 + .loc 1 244 21 ; hybrid_w4a16.py:244:21 + s_mul_i32 s0, s23, s24 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v63, v26, 16, 1 + .loc 1 125 44 ; hybrid_w4a16.py:125:44 + v_or_b32_e32 v33, v34, v33 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v34, v25, 16, 1 + .loc 1 126 44 ; hybrid_w4a16.py:126:44 + v_and_or_b32 v39, v0, 1, v35 + .loc 1 244 43 ; hybrid_w4a16.py:244:43 + s_add_i32 s0, s0, s25 + s_lshl_b32 s1, s23, 6 + .loc 1 125 44 ; hybrid_w4a16.py:125:44 + v_or_b32_e32 v66, s24, v33 + .loc 1 244 43 ; hybrid_w4a16.py:244:43 + v_mul_lo_u32 v33, s23, v33 + .loc 1 126 44 ; hybrid_w4a16.py:126:44 + v_or_b32_e32 v0, 46, v39 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v34, v25, v34, 0x7fff + .loc 1 126 31 ; hybrid_w4a16.py:126:31 + v_or_b32_e32 v68, s25, v39 + .loc 1 126 44 is_stmt 0 ; hybrid_w4a16.py:126:44 + v_or_b32_e32 v62, 2, v39 + .loc 1 245 32 is_stmt 1 ; hybrid_w4a16.py:245:32 + v_cmp_gt_i32_e64 s16, s22, v66 + .loc 1 126 31 ; hybrid_w4a16.py:126:31 + v_or_b32_e32 v43, s25, v0 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v34, 16, v34 + .loc 1 244 43 ; hybrid_w4a16.py:244:43 + v_add_nc_u32_e32 v99, s0, v33 + .loc 1 245 56 ; hybrid_w4a16.py:245:56 + v_cmp_gt_i32_e64 s15, s23, v68 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s17, v25, v25 + .loc 1 126 31 ; hybrid_w4a16.py:126:31 + v_or_b32_e32 v65, s25, v62 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v63, v26, v63, 0x7fff + .loc 1 244 43 ; hybrid_w4a16.py:244:43 + v_add3_u32 v33, s0, s1, v33 + .loc 1 245 56 ; hybrid_w4a16.py:245:56 + v_cmp_gt_i32_e64 s0, s23, v43 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v43, v99, v39, 1 + .loc 1 126 44 ; hybrid_w4a16.py:126:44 + v_or_b32_e32 v61, 4, v39 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v67, v27, 16, 1 + v_cndmask_b32_e64 v25, 0x7fff, v34, s17 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s17, s16, s15 + .loc 1 126 44 ; hybrid_w4a16.py:126:44 + v_or_b32_e32 v35, 44, v39 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v63, 16, v63 + .loc 1 245 56 ; hybrid_w4a16.py:245:56 + v_cmp_gt_i32_e64 s14, s23, v65 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v34, 0x80000000, v43, s17 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s17, v26, v26 + .loc 1 126 31 ; hybrid_w4a16.py:126:31 + v_or_b32_e32 v64, s25, v61 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v67, v27, v67, 0x7fff + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v43, v99, v62, 1 + .loc 1 126 31 ; hybrid_w4a16.py:126:31 + v_or_b32_e32 v44, s25, v35 + .loc 1 126 44 is_stmt 0 ; hybrid_w4a16.py:126:44 + v_or_b32_e32 v56, 6, v39 + .loc 1 243 23 is_stmt 1 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v70, v28, 16, 1 + v_cndmask_b32_e64 v26, 0x7fff, v63, s17 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s17, s16, s14 + .loc 1 126 44 ; hybrid_w4a16.py:126:44 + v_or_b32_e32 v36, 42, v39 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v67, 16, v67 + .loc 1 245 56 ; hybrid_w4a16.py:245:56 + v_cmp_gt_i32_e64 s13, s23, v64 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v43, 0x80000000, v43, s17 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s17, v27, v27 + .loc 1 126 31 ; hybrid_w4a16.py:126:31 + v_or_b32_e32 v60, s25, v56 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v70, v28, v70, 0x7fff + .loc 1 245 56 ; hybrid_w4a16.py:245:56 + v_cmp_gt_i32_e64 s1, s23, v44 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v44, v99, v61, 1 + .loc 1 126 31 ; hybrid_w4a16.py:126:31 + v_or_b32_e32 v45, s25, v36 + .loc 1 126 44 is_stmt 0 ; hybrid_w4a16.py:126:44 + v_or_b32_e32 v55, 8, v39 + .loc 1 243 23 is_stmt 1 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v71, v29, 16, 1 + v_cndmask_b32_e64 v27, 0x7fff, v67, s17 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s17, s16, s13 + .loc 1 126 44 ; hybrid_w4a16.py:126:44 + v_or_b32_e32 v37, 40, v39 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v70, 16, v70 + .loc 1 245 56 ; hybrid_w4a16.py:245:56 + v_cmp_gt_i32_e64 s12, s23, v60 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v44, 0x80000000, v44, s17 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s17, v28, v28 + .loc 1 126 31 ; hybrid_w4a16.py:126:31 + v_or_b32_e32 v59, s25, v55 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v71, v29, v71, 0x7fff + .loc 1 245 56 ; hybrid_w4a16.py:245:56 + v_cmp_gt_i32_e64 s2, s23, v45 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v45, v99, v56, 1 + .loc 1 126 44 ; hybrid_w4a16.py:126:44 + v_or_b32_e32 v38, 38, v39 + v_or_b32_e32 v40, 36, v39 + v_or_b32_e32 v41, 34, v39 + v_or_b32_e32 v42, 32, v39 + .loc 1 126 31 is_stmt 0 ; hybrid_w4a16.py:126:31 + v_or_b32_e32 v46, s25, v37 + .loc 1 126 44 ; hybrid_w4a16.py:126:44 + v_or_b32_e32 v49, 14, v39 + v_or_b32_e32 v52, 12, v39 + v_or_b32_e32 v53, 10, v39 + .loc 1 243 23 is_stmt 1 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v72, v30, 16, 1 + v_cndmask_b32_e64 v28, 0x7fff, v70, s17 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s17, s16, s12 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v71, 16, v71 + .loc 1 245 56 ; hybrid_w4a16.py:245:56 + v_cmp_gt_i32_e64 s11, s23, v59 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v45, 0x80000000, v45, s17 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s17, v29, v29 + .loc 1 126 31 ; hybrid_w4a16.py:126:31 + v_or_b32_e32 v47, s25, v38 + v_or_b32_e32 v48, s25, v40 + v_or_b32_e32 v50, s25, v41 + v_or_b32_e32 v51, s25, v42 + v_or_b32_e32 v54, s25, v49 + v_or_b32_e32 v57, s25, v52 + v_or_b32_e32 v58, s25, v53 + .loc 1 125 31 ; hybrid_w4a16.py:125:31 + v_or_b32_e32 v69, 64, v66 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v72, v30, v72, 0x7fff + .loc 1 245 56 ; hybrid_w4a16.py:245:56 + v_cmp_gt_i32_e64 s3, s23, v46 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v46, v99, v55, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v73, v31, 16, 1 + v_cndmask_b32_e64 v29, 0x7fff, v71, s17 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s17, s16, s11 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v72, 16, v72 + .loc 1 245 32 ; hybrid_w4a16.py:245:32 + v_cmp_gt_i32_e32 vcc_lo, s22, v69 + .loc 1 245 56 is_stmt 0 ; hybrid_w4a16.py:245:56 + v_cmp_gt_i32_e64 s10, s23, v58 + v_cmp_gt_i32_e64 s9, s23, v57 + v_cmp_gt_i32_e64 s8, s23, v54 + v_cmp_gt_i32_e64 s7, s23, v51 + v_cmp_gt_i32_e64 s6, s23, v50 + v_cmp_gt_i32_e64 s5, s23, v48 + v_cmp_gt_i32_e64 s4, s23, v47 + .loc 1 246 21 is_stmt 1 ; hybrid_w4a16.py:246:21 + s_and_b32 s21, s21, 0xffff + s_mov_b32 s23, 0x31027000 + s_mov_b32 s22, 0x7ffffffe + v_cndmask_b32_e64 v46, 0x80000000, v46, s17 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s17, v30, v30 + v_add3_u32 v73, v31, v73, 0x7fff + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + s_clause 0x4 + buffer_store_b16 v25, v34, s[20:23], 0 offen + buffer_store_b16 v26, v43, s[20:23], 0 offen + buffer_store_b16 v27, v44, s[20:23], 0 offen + buffer_store_b16 v28, v45, s[20:23], 0 offen + buffer_store_b16 v29, v46, s[20:23], 0 offen + v_add_lshl_u32 v25, v99, v53, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v74, v32, 16, 1 + v_cndmask_b32_e64 v26, 0x7fff, v72, s17 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s17, s16, s10 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v73, 16, v73 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v25, 0x80000000, v25, s17 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s17, v31, v31 + v_add3_u32 v74, v32, v74, 0x7fff + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v27, v99, v52, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v75, v17, 16, 1 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v29, v99, v49, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cndmask_b32_e64 v28, 0x7fff, v73, s17 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s17, s16, s9 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v74, 16, v74 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v27, 0x80000000, v27, s17 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s17, v32, v32 + v_add3_u32 v75, v17, v75, 0x7fff + v_bfe_u32 v76, v18, 16, 1 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v31, v99, v42, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v77, v19, 16, 1 + v_cndmask_b32_e64 v30, 0x7fff, v74, s17 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s17, s16, s8 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v75, 16, v75 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v29, 0x80000000, v29, s17 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s17, v17, v17 + v_add3_u32 v76, v18, v76, 0x7fff + v_add3_u32 v77, v19, v77, 0x7fff + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + s_clause 0x2 + buffer_store_b16 v26, v25, s[20:23], 0 offen + buffer_store_b16 v28, v27, s[20:23], 0 offen + buffer_store_b16 v30, v29, s[20:23], 0 offen + v_add_lshl_u32 v25, v99, v41, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cndmask_b32_e64 v17, 0x7fff, v75, s17 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s17, s16, s7 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v76, 16, v76 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v31, 0x80000000, v31, s17 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s17, v18, v18 + v_bfe_u32 v78, v20, 16, 1 + v_lshrrev_b32_e32 v77, 16, v77 + v_bfe_u32 v79, v21, 16, 1 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + buffer_store_b16 v17, v31, s[20:23], 0 offen + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cndmask_b32_e64 v17, 0x7fff, v76, s17 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s17, s16, s6 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v78, v20, v78, 0x7fff + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v18, 0x80000000, v25, s17 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s17, v19, v19 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v25, v99, v40, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v79, v21, v79, 0x7fff + v_lshrrev_b32_e32 v78, 16, v78 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v26, v99, v38, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cndmask_b32_e64 v19, 0x7fff, v77, s17 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s17, s16, s5 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v80, v22, 16, 1 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v25, 0x80000000, v25, s17 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s17, v20, v20 + v_lshrrev_b32_e32 v79, 16, v79 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v27, v99, v37, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v80, v22, v80, 0x7fff + v_bfe_u32 v81, v23, 16, 1 + v_cndmask_b32_e64 v20, 0x7fff, v78, s17 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s17, s16, s4 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v28, v99, v36, 1 + v_cndmask_b32_e64 v26, 0x80000000, v26, s17 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s17, v21, v21 + v_lshrrev_b32_e32 v80, 16, v80 + v_add3_u32 v81, v23, v81, 0x7fff + v_bfe_u32 v82, v24, 16, 1 + v_bfe_u32 v83, v9, 16, 1 + v_cndmask_b32_e64 v21, 0x7fff, v79, s17 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s17, s16, s3 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v81, 16, v81 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v27, 0x80000000, v27, s17 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s17, v22, v22 + v_add3_u32 v82, v24, v82, 0x7fff + v_add3_u32 v83, v9, v83, 0x7fff + v_bfe_u32 v84, v10, 16, 1 + v_bfe_u32 v85, v11, 16, 1 + v_cndmask_b32_e64 v22, 0x7fff, v80, s17 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s17, s16, s2 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v82, 16, v82 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v28, 0x80000000, v28, s17 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s17, v23, v23 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + s_clause 0x4 + buffer_store_b16 v17, v18, s[20:23], 0 offen + buffer_store_b16 v19, v25, s[20:23], 0 offen + buffer_store_b16 v20, v26, s[20:23], 0 offen + buffer_store_b16 v21, v27, s[20:23], 0 offen + buffer_store_b16 v22, v28, s[20:23], 0 offen + v_add_lshl_u32 v17, v99, v35, 1 + v_add_lshl_u32 v19, v99, v0, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v83, 16, v83 + v_cndmask_b32_e64 v18, 0x7fff, v81, s17 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s17, s16, s1 + s_and_b32 s16, s16, s0 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v17, 0x80000000, v17, s17 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s17, v24, v24 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v19, 0x80000000, v19, s16 + v_add_lshl_u32 v21, v33, v39, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s16, v9, v9 + v_add3_u32 v84, v10, v84, 0x7fff + v_bfe_u32 v86, v12, 16, 1 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s15, vcc_lo, s15 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v87, v13, 16, 1 + v_cndmask_b32_e64 v20, 0x7fff, v82, s17 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v22, v33, v62, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v88, v14, 16, 1 + v_cndmask_b32_e64 v9, 0x7fff, v83, s16 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v21, 0x80000000, v21, s15 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v89, v15, 16, 1 + v_lshrrev_b32_e32 v84, 16, v84 + v_add3_u32 v85, v11, v85, 0x7fff + v_cmp_o_f32_e64 s15, v10, v10 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s14, vcc_lo, s14 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v86, v12, v86, 0x7fff + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + s_clause 0x2 + buffer_store_b16 v18, v17, s[20:23], 0 offen + buffer_store_b16 v20, v19, s[20:23], 0 offen + buffer_store_b16 v9, v21, s[20:23], 0 offen + v_add_lshl_u32 v9, v33, v61, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v87, v13, v87, 0x7fff + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v22, 0x80000000, v22, s14 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s14, v11, v11 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v11, v33, v56, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v88, v14, v88, 0x7fff + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v17, v33, v55, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v89, v15, v89, 0x7fff + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s13, vcc_lo, s13 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v18, v33, v53, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v85, 16, v85 + v_cndmask_b32_e64 v10, 0x7fff, v84, s15 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s12, vcc_lo, s12 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v86, 16, v86 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v9, 0x80000000, v9, s13 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s13, v12, v12 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s11, vcc_lo, s11 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v87, 16, v87 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v11, 0x80000000, v11, s12 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s12, v13, v13 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s10, vcc_lo, s10 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v88, 16, v88 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v17, 0x80000000, v17, s11 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s11, v14, v14 + v_bfe_u32 v90, v16, 16, 1 + v_lshrrev_b32_e32 v89, 16, v89 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v19, v33, v52, 1 + v_cndmask_b32_e64 v18, 0x80000000, v18, s10 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s10, v15, v15 + v_bfe_u32 v91, v1, 16, 1 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + buffer_store_b16 v10, v22, s[20:23], 0 offen + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cndmask_b32_e64 v10, 0x7fff, v85, s14 + v_bfe_u32 v92, v2, 16, 1 + v_cndmask_b32_e64 v12, 0x7fff, v86, s13 + v_cndmask_b32_e64 v13, 0x7fff, v87, s12 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s9, vcc_lo, s9 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v93, v3, 16, 1 + v_cndmask_b32_e64 v14, 0x7fff, v88, s11 + v_add3_u32 v90, v16, v90, 0x7fff + v_cndmask_b32_e64 v15, 0x7fff, v89, s10 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v19, 0x80000000, v19, s9 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v91, v1, v91, 0x7fff + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + s_clause 0x4 + buffer_store_b16 v10, v9, s[20:23], 0 offen + buffer_store_b16 v12, v11, s[20:23], 0 offen + buffer_store_b16 v13, v17, s[20:23], 0 offen + buffer_store_b16 v14, v18, s[20:23], 0 offen + buffer_store_b16 v15, v19, s[20:23], 0 offen + v_add_lshl_u32 v9, v33, v49, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v92, v2, v92, 0x7fff + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v11, v33, v42, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v93, v3, v93, 0x7fff + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s8, vcc_lo, s8 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v12, v33, v41, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v90, 16, v90 + v_cmp_o_f32_e64 s9, v16, v16 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s7, vcc_lo, s7 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v91, 16, v91 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v9, 0x80000000, v9, s8 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s8, v1, v1 + v_lshrrev_b32_e32 v92, 16, v92 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v11, 0x80000000, v11, s7 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s7, v2, v2 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s6, vcc_lo, s6 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v93, 16, v93 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v13, v33, v40, 1 + v_cndmask_b32_e64 v12, 0x80000000, v12, s6 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s6, v3, v3 + v_cndmask_b32_e64 v10, 0x7fff, v90, s9 + v_cndmask_b32_e64 v1, 0x7fff, v91, s8 + v_bfe_u32 v94, v4, 16, 1 + v_cndmask_b32_e64 v2, 0x7fff, v92, s7 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s5, vcc_lo, s5 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v95, v5, 16, 1 + v_bfe_u32 v96, v6, 16, 1 + v_cndmask_b32_e64 v3, 0x7fff, v93, s6 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v13, 0x80000000, v13, s5 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v97, v7, 16, 1 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + s_clause 0x2 + buffer_store_b16 v10, v9, s[20:23], 0 offen + buffer_store_b16 v1, v11, s[20:23], 0 offen + buffer_store_b16 v2, v12, s[20:23], 0 offen + v_add_lshl_u32 v1, v33, v38, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_bfe_u32 v98, v8, 16, 1 + v_add3_u32 v94, v4, v94, 0x7fff + v_add3_u32 v95, v5, v95, 0x7fff + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s4, vcc_lo, s4 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v96, v6, v96, 0x7fff + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + buffer_store_b16 v3, v13, s[20:23], 0 offen + v_add_lshl_u32 v3, v33, v37, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v97, v7, v97, 0x7fff + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v1, 0x80000000, v1, s4 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s4, v5, v5 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v5, v33, v36, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_add3_u32 v98, v8, v98, 0x7fff + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v9, v33, v35, 1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v94, 16, v94 + v_cmp_o_f32_e64 s5, v4, v4 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s3, vcc_lo, s3 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v95, 16, v95 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s2, vcc_lo, s2 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v96, 16, v96 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v3, 0x80000000, v3, s3 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s3, v6, v6 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 s1, vcc_lo, s1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_lshrrev_b32_e32 v97, 16, v97 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e64 v5, 0x80000000, v5, s2 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s2, v7, v7 + v_lshrrev_b32_e32 v98, 16, v98 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_add_lshl_u32 v0, v33, v0, 1 + v_cndmask_b32_e64 v9, 0x80000000, v9, s1 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cmp_o_f32_e64 s1, v8, v8 + v_cndmask_b32_e64 v2, 0x7fff, v94, s5 + v_cndmask_b32_e64 v4, 0x7fff, v95, s4 + v_cndmask_b32_e64 v6, 0x7fff, v96, s3 + .loc 1 245 38 ; hybrid_w4a16.py:245:38 + s_and_b32 vcc_lo, vcc_lo, s0 + .loc 1 243 23 ; hybrid_w4a16.py:243:23 + v_cndmask_b32_e64 v7, 0x7fff, v97, s2 + v_cndmask_b32_e64 v8, 0x7fff, v98, s1 + .loc 1 246 21 ; hybrid_w4a16.py:246:21 + v_cndmask_b32_e32 v0, 0x80000000, v0, vcc_lo + s_clause 0x4 + buffer_store_b16 v2, v1, s[20:23], 0 offen + buffer_store_b16 v4, v3, s[20:23], 0 offen + buffer_store_b16 v6, v5, s[20:23], 0 offen + buffer_store_b16 v7, v9, s[20:23], 0 offen + buffer_store_b16 v8, v0, s[20:23], 0 offen + .loc 1 246 4 is_stmt 0 ; hybrid_w4a16.py:246:4 + s_nop 0 + s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) + s_endpgm +.Ltmp9: + .section .rodata,"a",@progbits + .p2align 6, 0x0 + .amdhsa_kernel _triton_w4a16_skinny_fmt_kernel + .amdhsa_group_segment_fixed_size 0 + .amdhsa_private_segment_fixed_size 0 + .amdhsa_kernarg_size 88 + .amdhsa_user_sgpr_count 2 + .amdhsa_user_sgpr_dispatch_ptr 0 + .amdhsa_user_sgpr_queue_ptr 0 + .amdhsa_user_sgpr_kernarg_segment_ptr 1 + .amdhsa_user_sgpr_dispatch_id 0 + .amdhsa_user_sgpr_private_segment_size 0 + .amdhsa_wavefront_size32 1 + .amdhsa_uses_dynamic_stack 0 + .amdhsa_enable_private_segment 0 + .amdhsa_system_sgpr_workgroup_id_x 1 + .amdhsa_system_sgpr_workgroup_id_y 1 + .amdhsa_system_sgpr_workgroup_id_z 0 + .amdhsa_system_sgpr_workgroup_info 0 + .amdhsa_system_vgpr_workitem_id 0 + .amdhsa_next_free_vgpr 115 + .amdhsa_next_free_sgpr 35 + .amdhsa_reserve_vcc 1 + .amdhsa_float_round_mode_32 0 + .amdhsa_float_round_mode_16_64 0 + .amdhsa_float_denorm_mode_32 3 + .amdhsa_float_denorm_mode_16_64 3 + .amdhsa_dx10_clamp 1 + .amdhsa_ieee_mode 1 + .amdhsa_fp16_overflow 0 + .amdhsa_workgroup_processor_mode 1 + .amdhsa_memory_ordered 1 + .amdhsa_forward_progress 1 + .amdhsa_shared_vgpr_count 0 + .amdhsa_inst_pref_size 36 + .amdhsa_exception_fp_ieee_invalid_op 0 + .amdhsa_exception_fp_denorm_src 0 + .amdhsa_exception_fp_ieee_div_zero 0 + .amdhsa_exception_fp_ieee_overflow 0 + .amdhsa_exception_fp_ieee_underflow 0 + .amdhsa_exception_fp_ieee_inexact 0 + .amdhsa_exception_int_div_zero 0 + .end_amdhsa_kernel + .text +.Lfunc_end0: + .size _triton_w4a16_skinny_fmt_kernel, .Lfunc_end0-_triton_w4a16_skinny_fmt_kernel + .cfi_endproc + ; -- End function + .set _triton_w4a16_skinny_fmt_kernel.num_vgpr, 115 + .set _triton_w4a16_skinny_fmt_kernel.num_agpr, 0 + .set _triton_w4a16_skinny_fmt_kernel.numbered_sgpr, 35 + .set _triton_w4a16_skinny_fmt_kernel.num_named_barrier, 0 + .set _triton_w4a16_skinny_fmt_kernel.private_seg_size, 0 + .set _triton_w4a16_skinny_fmt_kernel.uses_vcc, 1 + .set _triton_w4a16_skinny_fmt_kernel.uses_flat_scratch, 0 + .set _triton_w4a16_skinny_fmt_kernel.has_dyn_sized_stack, 0 + .set _triton_w4a16_skinny_fmt_kernel.has_recursion, 0 + .set _triton_w4a16_skinny_fmt_kernel.has_indirect_call, 0 + .section .AMDGPU.csdata,"",@progbits +; Kernel info: +; codeLenInByte = 4548 +; TotalNumSgprs: 37 +; NumVgprs: 115 +; ScratchSize: 0 +; MemoryBound: 0 +; FloatMode: 240 +; IeeeMode: 1 +; LDSByteSize: 0 bytes/workgroup (compile time only) +; SGPRBlocks: 0 +; VGPRBlocks: 14 +; NumSGPRsForWavesPerEU: 37 +; NumVGPRsForWavesPerEU: 115 +; Occupancy: 12 +; WaveLimiterHint : 0 +; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0 +; COMPUTE_PGM_RSRC2:USER_SGPR: 2 +; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0 +; COMPUTE_PGM_RSRC2:TGID_X_EN: 1 +; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1 +; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0 +; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0 + .text + .p2alignl 7, 3214868480 + .fill 96, 4, 3214868480 + .section .AMDGPU.gpr_maximums,"",@progbits + .set amdgpu.max_num_vgpr, 0 + .set amdgpu.max_num_agpr, 0 + .set amdgpu.max_num_sgpr, 0 + .text + .section .debug_abbrev,"",@progbits + .byte 1 ; Abbreviation Code + .byte 17 ; DW_TAG_compile_unit + .byte 1 ; DW_CHILDREN_yes + .byte 37 ; DW_AT_producer + .byte 14 ; DW_FORM_strp + .byte 19 ; DW_AT_language + .byte 5 ; DW_FORM_data2 + .byte 3 ; DW_AT_name + .byte 14 ; DW_FORM_strp + .byte 16 ; DW_AT_stmt_list + .byte 23 ; DW_FORM_sec_offset + .byte 27 ; DW_AT_comp_dir + .byte 14 ; DW_FORM_strp + .byte 17 ; DW_AT_low_pc + .byte 1 ; DW_FORM_addr + .byte 18 ; DW_AT_high_pc + .byte 6 ; DW_FORM_data4 + .byte 0 ; EOM(1) + .byte 0 ; EOM(2) + .byte 2 ; Abbreviation Code + .byte 46 ; DW_TAG_subprogram + .byte 0 ; DW_CHILDREN_no + .byte 3 ; DW_AT_name + .byte 14 ; DW_FORM_strp + .byte 32 ; DW_AT_inline + .byte 11 ; DW_FORM_data1 + .byte 0 ; EOM(1) + .byte 0 ; EOM(2) + .byte 3 ; Abbreviation Code + .byte 46 ; DW_TAG_subprogram + .byte 1 ; DW_CHILDREN_yes + .byte 17 ; DW_AT_low_pc + .byte 1 ; DW_FORM_addr + .byte 18 ; DW_AT_high_pc + .byte 6 ; DW_FORM_data4 + .byte 49 ; DW_AT_abstract_origin + .byte 19 ; DW_FORM_ref4 + .byte 0 ; EOM(1) + .byte 0 ; EOM(2) + .byte 4 ; Abbreviation Code + .byte 29 ; DW_TAG_inlined_subroutine + .byte 0 ; DW_CHILDREN_no + .byte 49 ; DW_AT_abstract_origin + .byte 19 ; DW_FORM_ref4 + .byte 85 ; DW_AT_ranges + .byte 23 ; DW_FORM_sec_offset + .byte 88 ; DW_AT_call_file + .byte 11 ; DW_FORM_data1 + .byte 89 ; DW_AT_call_line + .byte 11 ; DW_FORM_data1 + .byte 87 ; DW_AT_call_column + .byte 11 ; DW_FORM_data1 + .byte 0 ; EOM(1) + .byte 0 ; EOM(2) + .byte 0 ; EOM(3) + .section .debug_info,"",@progbits +.Lcu_begin0: + .long .Ldebug_info_end0-.Ldebug_info_start0 ; Length of Unit +.Ldebug_info_start0: + .short 4 ; DWARF version number + .long .debug_abbrev ; Offset Into Abbrev. Section + .byte 8 ; Address Size (in bytes) + .byte 1 ; Abbrev [1] 0xb:0x44 DW_TAG_compile_unit + .long .Linfo_string0 ; DW_AT_producer + .short 2 ; DW_AT_language + .long .Linfo_string1 ; DW_AT_name + .long .Lline_table_start0 ; DW_AT_stmt_list + .long .Linfo_string2 ; DW_AT_comp_dir + .quad .Lfunc_begin0 ; DW_AT_low_pc + .long .Lfunc_end0-.Lfunc_begin0 ; DW_AT_high_pc + .byte 2 ; Abbrev [2] 0x2a:0x6 DW_TAG_subprogram + .long .Linfo_string3 ; DW_AT_name + .byte 1 ; DW_AT_inline + .byte 3 ; Abbrev [3] 0x30:0x1e DW_TAG_subprogram + .quad .Lfunc_begin0 ; DW_AT_low_pc + .long .Lfunc_end0-.Lfunc_begin0 ; DW_AT_high_pc + .long 42 ; DW_AT_abstract_origin + .byte 4 ; Abbrev [4] 0x41:0xc DW_TAG_inlined_subroutine + .long 42 ; DW_AT_abstract_origin + .long .Ldebug_ranges0 ; DW_AT_ranges + .byte 1 ; DW_AT_call_file + .byte 141 ; DW_AT_call_line + .byte 39 ; DW_AT_call_column + .byte 0 ; End Of Children Mark + .byte 0 ; End Of Children Mark +.Ldebug_info_end0: + .section .debug_ranges,"",@progbits +.Ldebug_ranges0: + .quad .Ltmp1-.Lfunc_begin0 + .quad .Ltmp2-.Lfunc_begin0 + .quad .Ltmp3-.Lfunc_begin0 + .quad .Ltmp4-.Lfunc_begin0 + .quad .Ltmp5-.Lfunc_begin0 + .quad .Ltmp6-.Lfunc_begin0 + .quad .Ltmp7-.Lfunc_begin0 + .quad .Ltmp8-.Lfunc_begin0 + .quad 0 + .quad 0 + .section .debug_str,"MS",@progbits,1 +.Linfo_string0: + .asciz "triton" ; string offset=0 +.Linfo_string1: + .asciz "hybrid_w4a16.py" ; string offset=7 +.Linfo_string2: + .asciz "/scratch/mgehre/vllm/vllm/model_executor/kernels/linear/mixed_precision" ; string offset=23 +.Linfo_string3: + .asciz "_triton_w4a16_skinny_fmt_kernel" ; string offset=95 + .section ".note.GNU-stack","",@progbits + .amdgpu_metadata +--- +amdhsa.kernels: + - .args: + - .address_space: global + .offset: 0 + .size: 8 + .value_kind: global_buffer + - .address_space: global + .offset: 8 + .size: 8 + .value_kind: global_buffer + - .address_space: global + .offset: 16 + .size: 8 + .value_kind: global_buffer + - .address_space: global + .offset: 24 + .size: 8 + .value_kind: global_buffer + - .address_space: global + .offset: 32 + .size: 8 + .value_kind: global_buffer + - .offset: 40 + .size: 4 + .value_kind: by_value + - .offset: 44 + .size: 4 + .value_kind: by_value + - .offset: 48 + .size: 4 + .value_kind: by_value + - .offset: 52 + .size: 4 + .value_kind: by_value + - .offset: 56 + .size: 4 + .value_kind: by_value + - .offset: 60 + .size: 4 + .value_kind: by_value + - .offset: 64 + .size: 4 + .value_kind: by_value + - .address_space: global + .offset: 72 + .size: 8 + .value_kind: global_buffer + - .address_space: global + .offset: 80 + .size: 8 + .value_kind: global_buffer + .group_segment_fixed_size: 0 + .kernarg_segment_align: 8 + .kernarg_segment_size: 88 + .max_flat_workgroup_size: 256 + .name: _triton_w4a16_skinny_fmt_kernel + .private_segment_fixed_size: 0 + .sgpr_count: 37 + .sgpr_spill_count: 0 + .symbol: _triton_w4a16_skinny_fmt_kernel.kd + .uniform_work_group_size: 1 + .uses_dynamic_stack: false + .vgpr_count: 115 + .vgpr_spill_count: 0 + .wavefront_size: 32 + .workgroup_processor_mode: 1 +amdhsa.target: amdgcn-amd-amdhsa--gfx1151 +amdhsa.version: + - 1 + - 2 +... + + .end_amdgpu_metadata + .section .debug_line,"",@progbits +.Lline_table_start0: diff --git a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.meta.json b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.meta.json new file mode 100644 index 000000000000..ef40574ef0bf --- /dev/null +++ b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.meta.json @@ -0,0 +1,32 @@ +{ + "_comment": "Pinned ABI + constexpr config for the hand-editable AMDGCN of the W4A16 skinny GEMM. Dumped from Triton 3.6.0 _triton_w4a16_skinny_fmt_kernel for the Qwen3-VL-4B-AWQ up/gate_proj-class shapes (M>1024, N<2K, group_size=32) on gfx1151. Edit the .amdgcn next to this file; the loader asserts the launch matches every field here.", + "kernel_name": "_triton_w4a16_skinny_fmt_kernel", + "arch": "gfx1151", + "dtype": "bfloat16", + "has_zp": false, + "group_size": 32, + "block_m": 128, + "block_n": 64, + "block_k": 32, + "num_warps": 8, + "warp_size": 32, + "shared_bytes": 8192, + "kernarg_size": 88, + "kernarg_align": 8, + "kernarg_layout": [ + {"name": "a_ptr", "offset": 0, "kind": "ptr"}, + {"name": "b_ptr", "offset": 8, "kind": "ptr"}, + {"name": "scales_ptr", "offset": 16, "kind": "ptr"}, + {"name": "packed_scale_zp_ptr","offset": 24, "kind": "ptr"}, + {"name": "c_ptr", "offset": 32, "kind": "ptr"}, + {"name": "M", "offset": 40, "kind": "i32"}, + {"name": "N", "offset": 44, "kind": "i32"}, + {"name": "K", "offset": 48, "kind": "i32"}, + {"name": "K8", "offset": 52, "kind": "i32"}, + {"name": "stride_bn", "offset": 56, "kind": "i32"}, + {"name": "num_groups", "offset": 60, "kind": "i32"}, + {"name": "group_size", "offset": 64, "kind": "i32"}, + {"name": "_hidden_scratch0", "offset": 72, "kind": "ptr_null"}, + {"name": "_hidden_scratch1", "offset": 80, "kind": "ptr_null"} + ] +} diff --git a/vllm/model_executor/kernels/linear/mixed_precision/asm_w4a16.py b/vllm/model_executor/kernels/linear/mixed_precision/asm_w4a16.py new file mode 100644 index 000000000000..99a2a6b0b14a --- /dev/null +++ b/vllm/model_executor/kernels/linear/mixed_precision/asm_w4a16.py @@ -0,0 +1,268 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Direct-HIP loader/launcher for the hand-editable AMDGCN W4A16 skinny GEMM. + +This bypasses the Triton JIT entirely: it loads a prebuilt ``.hsaco`` (assembled +from a hand-editable ``.amdgcn`` under ``asm/``) with ``libamdhip64`` and launches +it via ``hipModuleLaunchKernel`` using the kernarg ABI pinned in +``asm/hybrid_w4a16_skinny_gfx1151.meta.json``. + +Workflow: edit ``asm/hybrid_w4a16_skinny_gfx1151.amdgcn`` -> rebuild vLLM (CMake +assembles it) OR just rerun (this module re-assembles on demand when the +``.amdgcn`` is newer than the cached ``.hsaco``). The dispatcher in +``hybrid_w4a16.py`` routes the matching launch here and ``assert``s that every +constexpr/config field in the meta matches the request. + +Enabled only when ``VLLM_W4A16_HANDASM=1``. +""" + +from __future__ import annotations + +import ctypes +import functools +import glob +import json +import math +import os +import struct +import subprocess +from pathlib import Path + +import torch + +_ASM_DIR = Path(__file__).parent / "asm" +_STEM = "hybrid_w4a16_skinny_gfx1151" +_AMDGCN = _ASM_DIR / f"{_STEM}.amdgcn" +_HSACO = _ASM_DIR / f"{_STEM}.hsaco" +_META = _ASM_DIR / f"{_STEM}.meta.json" + +# hipModuleLaunchKernel "extra" sentinels. +_HIP_LAUNCH_PARAM_BUFFER_POINTER = ctypes.c_void_p(0x01) +_HIP_LAUNCH_PARAM_BUFFER_SIZE = ctypes.c_void_p(0x02) +_HIP_LAUNCH_PARAM_END = ctypes.c_void_p(0x03) + + +@functools.lru_cache(maxsize=1) +def load_meta() -> dict: + return json.loads(_META.read_text()) + + +def _find_clang() -> str: + """Locate an amdgcn-capable clang for on-demand (re)assembly.""" + cand = os.environ.get("VLLM_ROCM_CLANG") + if cand and Path(cand).exists(): + return cand + # ROCm SDK wheel shipped in the active venv (next to torch in + # site-packages), then system ROCm. + site = os.path.dirname(os.path.dirname(torch.__file__)) # site-packages + patterns = [ + os.path.join(site, "_rocm_sdk_devel", "lib", "llvm", "bin", "clang"), + os.path.join(site, "_rocm_sdk_core", "lib", "llvm", "bin", "clang"), + "/opt/rocm*/llvm/bin/clang", + ] + for pat in patterns: + for hit in glob.glob(pat): + if Path(hit).exists(): + return str(Path(hit).resolve()) + raise FileNotFoundError( + "No amdgcn clang found to assemble the W4A16 .amdgcn. Set " + "VLLM_ROCM_CLANG to a ROCm clang, or build vLLM (CMake assembles it)." + ) + + +def assemble(force: bool = False) -> Path: + """Assemble the .amdgcn -> .hsaco if missing/stale. Returns the .hsaco path. + + This is the dev-loop fast path; the CMake build performs the same step at + build time so installed wheels ship a ready .hsaco. + """ + if ( + not force + and _HSACO.exists() + and _HSACO.stat().st_mtime >= _AMDGCN.stat().st_mtime + ): + return _HSACO + clang = _find_clang() + arch = load_meta()["arch"] + cmd = [ + clang, + "-target", + "amdgcn-amd-amdhsa", + f"-mcpu={arch}", + "-x", + "assembler", + str(_AMDGCN), + "-o", + str(_HSACO), + ] + subprocess.run(cmd, check=True, capture_output=True, text=True) + return _HSACO + + +def _open_hip() -> ctypes.CDLL: + """Open libamdhip64 (already in-process via torch's ROCm SDK) by full path. + + The bare soname is not on the loader path in the SDK-wheel layout, so resolve + the versioned file directly; dlopen-ing the same path just shares torch's + handle and HIP context. + """ + site = os.path.dirname(os.path.dirname(torch.__file__)) + for pat in ( + os.path.join(site, "_rocm_sdk_core", "lib", "libamdhip64.so*"), + os.path.join(os.path.dirname(torch.__file__), "lib", "libamdhip64.so*"), + "libamdhip64.so", + ): + for hit in sorted(glob.glob(pat)) or ([pat] if "*" not in pat else []): + try: + return ctypes.CDLL(hit) + except OSError: + continue + raise OSError("could not locate libamdhip64 for the hand-asm W4A16 loader") + + +class _HandAsmKernel: + """Loads the hsaco once and launches it via hipModuleLaunchKernel.""" + + def __init__(self) -> None: + self.meta = load_meta() + self.hip = _open_hip() + self.hip.hipModuleLoadData.argtypes = [ + ctypes.POINTER(ctypes.c_void_p), + ctypes.c_void_p, + ] + self.hip.hipModuleGetFunction.argtypes = [ + ctypes.POINTER(ctypes.c_void_p), + ctypes.c_void_p, + ctypes.c_char_p, + ] + self.hip.hipModuleLaunchKernel.argtypes = [ + ctypes.c_void_p, + ctypes.c_uint, + ctypes.c_uint, + ctypes.c_uint, # grid + ctypes.c_uint, + ctypes.c_uint, + ctypes.c_uint, # block + ctypes.c_uint, # shared + ctypes.c_void_p, # stream + ctypes.POINTER(ctypes.c_void_p), # kernelParams + ctypes.POINTER(ctypes.c_void_p), # extra + ] + image = bytearray(assemble().read_bytes()) + buf = (ctypes.c_char * len(image)).from_buffer(image) + self._image_keepalive = image # hipModuleLoadData copies, but be safe + module = ctypes.c_void_p() + self._check(self.hip.hipModuleLoadData(ctypes.byref(module), buf)) + func = ctypes.c_void_p() + self._check( + self.hip.hipModuleGetFunction( + ctypes.byref(func), module, self.meta["kernel_name"].encode() + ) + ) + self.module = module + self.func = func + + def _check(self, err: int) -> None: + if err != 0: + raise RuntimeError(f"HIP error {err} in hand-asm W4A16 loader") + + def launch( + self, a, b_q, scales, c, M, N, K, K8, stride_bn, num_groups, group_size, stream + ) -> None: + m = self.meta + ksize = m["kernarg_size"] + kbuf = bytearray(ksize) + # Pack per the pinned kernarg_layout. ptr_null args stay zero. + vals = { + "a_ptr": a.data_ptr(), + "b_ptr": b_q.data_ptr(), + "scales_ptr": scales.data_ptr(), + # symmetric: carrier unused, mirror production's dummy (scales). + "packed_scale_zp_ptr": scales.data_ptr(), + "c_ptr": c.data_ptr(), + "M": M, + "N": N, + "K": K, + "K8": K8, + "stride_bn": stride_bn, + "num_groups": num_groups, + "group_size": group_size, + } + for arg in m["kernarg_layout"]: + off, kind = arg["offset"], arg["kind"] + if kind == "ptr": + struct.pack_into(" _HandAsmKernel: + return _HandAsmKernel() + + +def config_matches( + dtype, has_zp, group_size, block_m, block_n, block_k, num_warps +) -> bool: + """True iff this launch is exactly what the pinned .amdgcn was built for.""" + m = load_meta() + return ( + str(dtype) == f"torch.{m['dtype']}" + and bool(has_zp) == m["has_zp"] + and group_size == m["group_size"] + and block_m == m["block_m"] + and block_n == m["block_n"] + and block_k == m["block_k"] + and num_warps == m["num_warps"] + ) + + +def launch_skinny_w4a16( + a, b_q, scales, c, M, N, K, K8, stride_bn, num_groups, group_size +) -> None: + """Assert-checked entry point used by the hybrid_w4a16 dispatcher.""" + m = load_meta() + assert a.dtype == getattr(torch, m["dtype"]), ( + f"hand-asm W4A16 built for {m['dtype']}, got {a.dtype}" + ) + assert group_size == m["group_size"], ( + f"hand-asm W4A16 built for group_size={m['group_size']}, got {group_size}" + ) + assert num_groups == K // group_size + assert K8 == K // 8 + stream = torch.cuda.current_stream().cuda_stream + _kernel().launch( + a, b_q, scales, c, M, N, K, K8, stride_bn, num_groups, group_size, stream + ) diff --git a/vllm/model_executor/kernels/linear/mixed_precision/hybrid_w4a16.py b/vllm/model_executor/kernels/linear/mixed_precision/hybrid_w4a16.py index 108e2111c5b8..97f67aad1ea1 100644 --- a/vllm/model_executor/kernels/linear/mixed_precision/hybrid_w4a16.py +++ b/vllm/model_executor/kernels/linear/mixed_precision/hybrid_w4a16.py @@ -12,6 +12,7 @@ weight copy. The triton kernel transposes tiles in-register. """ +import os from contextlib import nullcontext import torch @@ -43,6 +44,11 @@ LDS_CAPACITY_ELEMENTS = 64 * 1024 // 2 # 32768 fp16 elements +def _handasm_w4a16_enabled() -> bool: + """Opt-in hand-tuned AMDGCN override for the skinny W4A16 GEMM (gfx1151).""" + return os.environ.get("VLLM_W4A16_HANDASM", "0") == "1" and on_gfx1151() + + # --------------------------------------------------------------------------- # Triton kernel for the prefill path (reads skinny-format weights [N, K//8]) # --------------------------------------------------------------------------- @@ -447,6 +453,37 @@ def triton_w4a16_skinny_fmt_gemm( M, N, K, group_size, a.dtype ) grid = (triton.cdiv(M, block_m), triton.cdiv(N, block_n)) + # Opt-in hand-tuned AMDGCN override (VLLM_W4A16_HANDASM=1): dispatch to a + # prebuilt .hsaco assembled from a hand-editable .amdgcn when the selected + # tile/config exactly matches what that ISA was built for. The loader + # asserts every pinned constexpr/ABI field; otherwise we fall through to + # the Triton kernel below. See asm_w4a16.py / asm/. + if _handasm_w4a16_enabled(): + from .asm_w4a16 import config_matches, launch_skinny_w4a16 + + if config_matches( + a.dtype, + has_zp, + group_size, + block_m, + block_n, + block_k, + num_warps, + ): + launch_skinny_w4a16( + a, + b_q, + scales, + c, + M, + N, + K, + K8, + stride_bn, + num_groups, + group_size, + ) + return c # The kernel picks the packed (fp16/gfx1x) vs scalar unpack itself. _triton_w4a16_skinny_fmt_kernel[grid]( a, From becd074c9737336c1229f2bc51d9d1dd8f382e63 Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Wed, 24 Jun 2026 01:02:21 -0600 Subject: [PATCH 03/17] [ROCm] Enable hand-asm W4A16 skinny GEMM by default on gfx1151 Flip VLLM_W4A16_HANDASM to default-on so the hand-tuned AMDGCN kernel is used for the matching config (bf16, symmetric, group_size=32, tile 128x64x32) without needing the env var. Set VLLM_W4A16_HANDASM=0 to fall back to the Triton kernel. The dispatcher still asserts every pinned constexpr/ABI field and falls through to Triton for any non-matching shape. Co-authored-by: Claude Signed-off-by: Matthias Gehre --- .../kernels/linear/mixed_precision/hybrid_w4a16.py | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/vllm/model_executor/kernels/linear/mixed_precision/hybrid_w4a16.py b/vllm/model_executor/kernels/linear/mixed_precision/hybrid_w4a16.py index 97f67aad1ea1..e6f74cdf1b79 100644 --- a/vllm/model_executor/kernels/linear/mixed_precision/hybrid_w4a16.py +++ b/vllm/model_executor/kernels/linear/mixed_precision/hybrid_w4a16.py @@ -45,8 +45,11 @@ def _handasm_w4a16_enabled() -> bool: - """Opt-in hand-tuned AMDGCN override for the skinny W4A16 GEMM (gfx1151).""" - return os.environ.get("VLLM_W4A16_HANDASM", "0") == "1" and on_gfx1151() + """Hand-tuned AMDGCN override for the skinny W4A16 GEMM (gfx1151). + + On by default; set VLLM_W4A16_HANDASM=0 to fall back to the Triton kernel. + """ + return os.environ.get("VLLM_W4A16_HANDASM", "1") == "1" and on_gfx1151() # --------------------------------------------------------------------------- From caf54a81623d4dcb4f34f800a500a0144a6aa747 Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Wed, 24 Jun 2026 01:09:45 -0600 Subject: [PATCH 04/17] [Bench] Add --check correctness gate to benchmark_hybrid_w4a16 The benchmark only timed the kernel, so a bad hand-asm ISA edit (now the default path on gfx1151) could report a fast-but-wrong result. Add --check: run the active kernel and the Triton reference on identical inputs (toggling VLLM_W4A16_HANDASM) and assert cosine similarity >= --check-tol (default 0.999), exiting non-zero on mismatch. Using Triton as the reference keeps it general across providers with no dequant-math reconstruction, and it reports whether the hand-asm path was actually exercised for the shape. Co-authored-by: Claude Signed-off-by: Matthias Gehre --- benchmarks/kernels/benchmark_hybrid_w4a16.py | 110 +++++++++++++++++++ 1 file changed, 110 insertions(+) diff --git a/benchmarks/kernels/benchmark_hybrid_w4a16.py b/benchmarks/kernels/benchmark_hybrid_w4a16.py index b5f46d4f50a5..24ec2f6cf95d 100644 --- a/benchmarks/kernels/benchmark_hybrid_w4a16.py +++ b/benchmarks/kernels/benchmark_hybrid_w4a16.py @@ -36,6 +36,7 @@ from __future__ import annotations import argparse +import os import sys from pathlib import Path @@ -48,7 +49,9 @@ from tests.kernels.quantization.test_hybrid_w4a16_perf import ( # noqa: E402 PROVIDERS, + _compute_packed_scale_zp, _provider_dtype, + _provider_use_zp, measure_tflops, prepare_hybrid_weights, ) @@ -76,6 +79,82 @@ def _bench_one( return kernel, tflops, ms +def _apply(a, weights, group_size, provider): + """Run the production apply path once (respects VLLM_W4A16_HANDASM).""" + from vllm.model_executor.kernels.linear.mixed_precision.hybrid_w4a16 import ( + _hybrid_w4a16_apply_impl, + ) + from vllm.utils.platform_utils import num_compute_units + + use_zp = _provider_use_zp(provider) + packed_scale_zp = _compute_packed_scale_zp( + weights["w_s_skinny"], weights["w_zp"] if use_zp else None, a.dtype + ) + return _hybrid_w4a16_apply_impl( + a, + weights["w_q_skinny"], + weights["w_s_skinny"], + weights["w_q_skinny_i32"], + weights["w_zp"] if use_zp else None, + None, # bias + num_compute_units(), + group_size, + packed_scale_zp, + ) + + +def _check_one( + m: int, k: int, n: int, group_size: int, provider: str +) -> tuple[bool, float]: + """Compare the active (hand-asm) kernel against the Triton reference on the + SAME inputs. Returns (handasm_actually_used, cosine_similarity). + + Uses Triton as the trusted reference: run once with VLLM_W4A16_HANDASM=0 + (force Triton) and once with =1 (hand-asm if its config matches this shape), + then compare. When the hand-asm config does not match, both runs are Triton + and the check is trivially exact (reported as handasm_used=False). + """ + from vllm.model_executor.kernels.linear.mixed_precision import asm_w4a16 + from vllm.platforms.rocm import on_gfx1151 + + dtype = _provider_dtype(provider) + weights = prepare_hybrid_weights(k, n, group_size, dtype=dtype) + a = torch.randn((m, k), device="cuda", dtype=dtype) + + # Did the dispatcher's pinned config match this launch? + block = _select_skinny_config(m, n, k, group_size, dtype) + handasm_used = ( + on_gfx1151() + and block is not None + and asm_w4a16.config_matches( + dtype, _provider_use_zp(provider), group_size, *block + ) + ) + + os.environ["VLLM_W4A16_HANDASM"] = "0" + c_ref = _apply(a, weights, group_size, provider).float() + os.environ["VLLM_W4A16_HANDASM"] = "1" + c = _apply(a, weights, group_size, provider).float() + + cos = torch.nn.functional.cosine_similarity( + c.flatten(), c_ref.flatten(), dim=0 + ).item() + return handasm_used, cos + + +def _select_skinny_config(m, n, k, group_size, dtype): + """The tile (block_m, block_n, block_k, num_warps) the dispatcher would pick, + or None off gfx11.""" + from vllm.model_executor.kernels.linear.mixed_precision.hybrid_w4a16 import ( + _select_skinny_gfx11_config, + ) + from vllm.platforms.rocm import on_gfx1x + + if not on_gfx1x(): + return None + return _select_skinny_gfx11_config(m, n, k, group_size, dtype) + + def main() -> None: p = argparse.ArgumentParser( description=__doc__, @@ -102,6 +181,19 @@ def main() -> None: default=[1322], help="M (token) batch sizes to sweep (default: 1322)", ) + p.add_argument( + "--check", + action="store_true", + help="Verify the active kernel matches the Triton reference (cos > " + "0.999) on identical inputs before/instead of timing. Exits non-zero " + "on mismatch. Use this after editing the hand-asm .amdgcn.", + ) + p.add_argument( + "--check-tol", + type=float, + default=0.999, + help="Minimum cosine similarity for --check (default 0.999)", + ) args = p.parse_args() if args.model: @@ -111,6 +203,24 @@ def main() -> None: {"k": args.k, "n": args.n, "group_size": args.group_size, "comment": ""} ] + if args.check: + failures = 0 + for s in shapes: + k, n, gs = s["k"], s["n"], s["group_size"] + for m in args.batches: + used, cos = _check_one(m, k, n, gs, args.provider) + ok = cos >= args.check_tol + failures += not ok + where = "hand-asm" if used else "triton-only (no hand-asm)" + print( + f" CHECK {m}x{n}x{k} g{gs} {args.provider}: " + f"cos={cos:.6f} [{where}] -> {'PASS' if ok else 'FAIL'}" + ) + torch.accelerator.synchronize() + if failures: + raise SystemExit(f"{failures} correctness check(s) FAILED") + return + hdr = f"{'shape (MxNxK)':>22} {'g':>4} {'kernel':>22} {'TFLOP/s':>9} {'ms/call':>9}" print(hdr) print("-" * len(hdr)) From eea831ea3a88d2632c5b504c01cdc21449731356 Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Wed, 24 Jun 2026 09:43:09 -0600 Subject: [PATCH 05/17] [ROCm] W4A16 skinny GEMM: magic-constant bf16 dequant + CU mode (gfx1151) Hand-optimize the gfx1151 W4A16 skinny GEMM AMDGCN: - bf16 magic-constant dequant: replace the int4->fp32->bf16 path (v_cvt_f32_i32 + software RNE v_bfe/v_add3/v_alignbit) with `v_and_or_b32 ...,0xc3084300`, producing bf16(128+nibble) in the low half and bf16(-136) in the high half in one op per nibble. The existing v_dot2_bf16_bf16 then computes b_raw*scale + (-136)*scale = (nibble-8)*scale with scale broadcast to both halves -- folding the -8 offset for free, no int->float conversion and no float->bf16 cast. (bf16 128.0 = 0x4300; nibbles 0..15 fit the 7-bit mantissa at 2^7 exactly.) - CU mode + drop buffer_gl0_inv around the LDS-transpose barriers. End-to-end at M=1322,N=19456,K=2560 (Qwen3-VL-4B-AWQ gate_up): 5.47 -> 5.22 ms (~4.6% faster) vs the Triton path, cos 0.999995. Per-iter dequant VALU drops (v_cvt_f32_i32 8->0, v_alignbit 8->0). Co-authored-by: Claude Signed-off-by: Matthias Gehre --- .../asm/hybrid_w4a16_skinny_gfx1151.amdgcn | 79 +++++-------------- 1 file changed, 19 insertions(+), 60 deletions(-) diff --git a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn index cc169cbda7db..cdb450fcc603 100644 --- a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn +++ b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn @@ -250,69 +250,30 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel buffer_load_b128 v[55:58], v55, s[12:15], 0 offen s_waitcnt vmcnt(0) s_barrier - buffer_gl0_inv ds_store_b128 v42, v[51:54] ds_store_b128 v42, v[55:58] offset:4096 .loc 1 190 22 ; hybrid_w4a16.py:190:22 - v_lshrrev_b32_e32 v60, 28, v59 - .loc 1 190 37 is_stmt 0 ; hybrid_w4a16.py:190:37 - v_and_b32_e32 v61, 15, v59 - v_bfe_u32 v62, v59, 16, 4 - v_bfe_u32 v63, v59, 4, 4 - v_bfe_u32 v64, v59, 20, 4 - v_bfe_u32 v65, v59, 8, 4 - v_bfe_u32 v66, v59, 24, 4 - v_bfe_u32 v59, v59, 12, 4 - .loc 1 234 28 is_stmt 1 ; hybrid_w4a16.py:234:28 - v_add_nc_u32_e32 v84, -8, v61 - v_add_nc_u32_e32 v85, -8, v62 - v_add_nc_u32_e32 v86, -8, v63 - v_add_nc_u32_e32 v87, -8, v64 - v_add_nc_u32_e32 v88, -8, v65 - v_add_nc_u32_e32 v89, -8, v66 - v_add_nc_u32_e32 v90, -8, v59 - v_add_nc_u32_e32 v91, -8, v60 - .loc 1 234 34 is_stmt 0 ; hybrid_w4a16.py:234:34 - v_cvt_f32_i32_e32 v84, v84 - v_cvt_f32_i32_e32 v85, v85 - v_cvt_f32_i32_e32 v86, v86 - v_cvt_f32_i32_e32 v87, v87 - v_cvt_f32_i32_e32 v88, v88 - v_cvt_f32_i32_e32 v89, v89 - v_cvt_f32_i32_e32 v90, v90 - v_cvt_f32_i32_e32 v91, v91 - v_bfe_u32 v92, v84, 16, 1 - v_bfe_u32 v93, v85, 16, 1 - v_bfe_u32 v94, v86, 16, 1 - v_bfe_u32 v95, v87, 16, 1 - v_bfe_u32 v96, v88, 16, 1 - v_bfe_u32 v97, v89, 16, 1 - v_bfe_u32 v98, v90, 16, 1 - v_bfe_u32 v99, v91, 16, 1 - v_add3_u32 v84, v84, v92, 0x7fff - v_add3_u32 v85, v85, v93, 0x7fff - v_add3_u32 v86, v86, v94, 0x7fff - v_add3_u32 v87, v87, v95, 0x7fff - v_add3_u32 v88, v88, v96, 0x7fff - v_add3_u32 v89, v89, v97, 0x7fff - v_add3_u32 v90, v90, v98, 0x7fff - v_add3_u32 v91, v91, v99, 0x7fff - .loc 1 224 16 is_stmt 1 ; hybrid_w4a16.py:224:16 + v_mov_b32 v92, 15 + v_and_or_b32 v84, v59, v92, 0xc3084300 + v_lshrrev_b32_e32 v85, 16, v59 + v_and_or_b32 v85, v85, v92, 0xc3084300 + v_lshrrev_b32_e32 v86, 4, v59 + v_and_or_b32 v86, v86, v92, 0xc3084300 + v_lshrrev_b32_e32 v87, 20, v59 + v_and_or_b32 v87, v87, v92, 0xc3084300 + v_lshrrev_b32_e32 v88, 8, v59 + v_and_or_b32 v88, v88, v92, 0xc3084300 + v_lshrrev_b32_e32 v89, 24, v59 + v_and_or_b32 v89, v89, v92, 0xc3084300 + v_lshrrev_b32_e32 v90, 12, v59 + v_and_or_b32 v90, v90, v92, 0xc3084300 + v_lshrrev_b32_e32 v91, 28, v59 + v_and_or_b32 v91, v91, v92, 0xc3084300 v_cndmask_b32_e64 v83, 0x3f80, v83, s0 - .loc 1 234 50 ; hybrid_w4a16.py:234:50 - v_alignbit_b32 v84, 0, v84, 16 - v_alignbit_b32 v85, 0, v85, 16 - v_alignbit_b32 v86, 0, v86, 16 - v_alignbit_b32 v87, 0, v87, 16 - v_alignbit_b32 v88, 0, v88, 16 - v_alignbit_b32 v89, 0, v89, 16 - v_alignbit_b32 v90, 0, v90, 16 - v_alignbit_b32 v91, 0, v91, 16 - v_perm_b32 v83, 0, v83, 0x5040100 + v_lshl_or_b32 v83, v83, 16, v83 .loc 1 148 20 ; hybrid_w4a16.py:148:20 s_waitcnt lgkmcnt(0) s_barrier - buffer_gl0_inv ds_load_b128 v[51:54], v43 ds_load_b128 v[59:62], v43 offset:4096 ds_load_b128 v[55:58], v44 @@ -337,11 +298,9 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel s_waitcnt lgkmcnt(0) v_perm_b32 v83, v83, v91, 0x5040100 s_barrier - buffer_gl0_inv ds_store_b128 v42, v[83:86] s_waitcnt lgkmcnt(0) s_barrier - buffer_gl0_inv ds_load_b128 v[83:86], v47 ds_load_b128 v[91:94], v47 offset:2048 ds_load_b128 v[87:90], v48 @@ -990,7 +949,7 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel .amdhsa_dx10_clamp 1 .amdhsa_ieee_mode 1 .amdhsa_fp16_overflow 0 - .amdhsa_workgroup_processor_mode 1 + .amdhsa_workgroup_processor_mode 0 .amdhsa_memory_ordered 1 .amdhsa_forward_progress 1 .amdhsa_shared_vgpr_count 0 @@ -1225,7 +1184,7 @@ amdhsa.kernels: .vgpr_count: 115 .vgpr_spill_count: 0 .wavefront_size: 32 - .workgroup_processor_mode: 1 + .workgroup_processor_mode: 0 amdhsa.target: amdgcn-amd-amdhsa--gfx1151 amdhsa.version: - 1 From 8d213666b8f04b1e7a162d8ee17cc875f0527125 Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Wed, 24 Jun 2026 14:38:06 -0600 Subject: [PATCH 06/17] [ROCm] W4A16 skinny GEMM: op_sel dest-packing (drop v_perm) Pack the v_dot2_bf16_bf16 dequant outputs directly with op_sel:[0,0,0,1] (dest high-half write, preserving low) -- even-K to low, odd-K to high of consecutive VGPRs -- so the four v_perm_b32 that repacked them before the ds_store are eliminated. Correct (cos 0.999994); perf-neutral at M=1322 (the perms were not on the critical path), but removes redundant ops and is the WMMA-native packing. Co-authored-by: Claude Signed-off-by: Matthias Gehre --- .../asm/hybrid_w4a16_skinny_gfx1151.amdgcn | 23 ++++++++----------- 1 file changed, 9 insertions(+), 14 deletions(-) diff --git a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn index cdb450fcc603..b13b266ac610 100644 --- a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn +++ b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn @@ -283,22 +283,17 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel ds_load_b128 v[71:74], v46 ds_load_b128 v[79:82], v46 offset:4096 .loc 1 234 50 ; hybrid_w4a16.py:234:50 - v_dot2_bf16_bf16 v92, v86, v83, 0 - v_dot2_bf16_bf16 v88, v88, v83, 0 - v_dot2_bf16_bf16 v86, v90, v83, 0 - v_dot2_bf16_bf16 v90, v91, v83, 0 - v_dot2_bf16_bf16 v89, v89, v83, 0 - v_dot2_bf16_bf16 v87, v87, v83, 0 - v_dot2_bf16_bf16 v91, v84, v83, 0 - v_dot2_bf16_bf16 v83, v85, v83, 0 - .loc 1 240 33 ; hybrid_w4a16.py:240:33 - v_perm_b32 v86, v90, v86, 0x5040100 - v_perm_b32 v85, v89, v88, 0x5040100 - v_perm_b32 v84, v87, v92, 0x5040100 + v_dot2_bf16_bf16 v94, v84, v83, 0 + v_dot2_bf16_bf16 v94, v85, v83, 0 op_sel:[0,0,0,1] + v_dot2_bf16_bf16 v95, v86, v83, 0 + v_dot2_bf16_bf16 v95, v87, v83, 0 op_sel:[0,0,0,1] + v_dot2_bf16_bf16 v96, v88, v83, 0 + v_dot2_bf16_bf16 v96, v89, v83, 0 op_sel:[0,0,0,1] + v_dot2_bf16_bf16 v97, v90, v83, 0 + v_dot2_bf16_bf16 v97, v91, v83, 0 op_sel:[0,0,0,1] s_waitcnt lgkmcnt(0) - v_perm_b32 v83, v83, v91, 0x5040100 s_barrier - ds_store_b128 v42, v[83:86] + ds_store_b128 v42, v[94:97] s_waitcnt lgkmcnt(0) s_barrier ds_load_b128 v[83:86], v47 From 7a60e666d4622ffebb9b199f8724c1e796f7b87d Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Wed, 24 Jun 2026 14:39:37 -0600 Subject: [PATCH 07/17] [ROCm] W4A16 skinny GEMM: factor dequant into assembler macros Refactor the hand-written bf16 dequant into GAS/llvm-mc .macro definitions (DQ_EXT0/DQ_EXT for the magic-constant nibble extract, DQ_PACK for the op_sel dest-packed v_dot2 pair), invoked per nibble/pair. Assembler macros are pure text expansion, so this is zero runtime cost: the assembled .hsaco is byte-identical to the pre-refactor kernel (verified: same sha256). Purely a readability/maintainability improvement -- the quant logic is defined once. Verified: identical .hsaco, cos 0.999995 via the runtime loader. Co-authored-by: Claude Signed-off-by: Matthias Gehre --- .../asm/hybrid_w4a16_skinny_gfx1151.amdgcn | 49 ++++++++++--------- 1 file changed, 26 insertions(+), 23 deletions(-) diff --git a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn index b13b266ac610..d207a11764e4 100644 --- a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn +++ b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn @@ -1,5 +1,19 @@ .amdgcn_target "amdgcn-amd-amdhsa--gfx1151" .amdhsa_code_object_version 5 + +; --- W4A16 dequant macros (assembler text-expansion; zero runtime cost) --- +.macro DQ_EXT0 dst ; nibble@bit0 -> (bf16(128+n) low, bf16(-136) high) + v_and_or_b32 \dst, v59, v92, 0xc3084300 +.endm +.macro DQ_EXT dst, sh ; nibble@bit -> (bf16(128+n) low, bf16(-136) high) + v_lshrrev_b32_e32 \dst, \sh, v59 + v_and_or_b32 \dst, \dst, v92, 0xc3084300 +.endm +.macro DQ_PACK dst, klo, khi ; two dot2: klo->low half, khi->high half of dst + v_dot2_bf16_bf16 \dst, \klo, v83, 0 + v_dot2_bf16_bf16 \dst, \khi, v83, 0 op_sel:[0,0,0,1] +.endm + .text .globl _triton_w4a16_skinny_fmt_kernel ; -- Begin function _triton_w4a16_skinny_fmt_kernel .p2align 8 @@ -254,21 +268,14 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel ds_store_b128 v42, v[55:58] offset:4096 .loc 1 190 22 ; hybrid_w4a16.py:190:22 v_mov_b32 v92, 15 - v_and_or_b32 v84, v59, v92, 0xc3084300 - v_lshrrev_b32_e32 v85, 16, v59 - v_and_or_b32 v85, v85, v92, 0xc3084300 - v_lshrrev_b32_e32 v86, 4, v59 - v_and_or_b32 v86, v86, v92, 0xc3084300 - v_lshrrev_b32_e32 v87, 20, v59 - v_and_or_b32 v87, v87, v92, 0xc3084300 - v_lshrrev_b32_e32 v88, 8, v59 - v_and_or_b32 v88, v88, v92, 0xc3084300 - v_lshrrev_b32_e32 v89, 24, v59 - v_and_or_b32 v89, v89, v92, 0xc3084300 - v_lshrrev_b32_e32 v90, 12, v59 - v_and_or_b32 v90, v90, v92, 0xc3084300 - v_lshrrev_b32_e32 v91, 28, v59 - v_and_or_b32 v91, v91, v92, 0xc3084300 + DQ_EXT0 v84 + DQ_EXT v85, 16 + DQ_EXT v86, 4 + DQ_EXT v87, 20 + DQ_EXT v88, 8 + DQ_EXT v89, 24 + DQ_EXT v90, 12 + DQ_EXT v91, 28 v_cndmask_b32_e64 v83, 0x3f80, v83, s0 v_lshl_or_b32 v83, v83, 16, v83 .loc 1 148 20 ; hybrid_w4a16.py:148:20 @@ -283,14 +290,10 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel ds_load_b128 v[71:74], v46 ds_load_b128 v[79:82], v46 offset:4096 .loc 1 234 50 ; hybrid_w4a16.py:234:50 - v_dot2_bf16_bf16 v94, v84, v83, 0 - v_dot2_bf16_bf16 v94, v85, v83, 0 op_sel:[0,0,0,1] - v_dot2_bf16_bf16 v95, v86, v83, 0 - v_dot2_bf16_bf16 v95, v87, v83, 0 op_sel:[0,0,0,1] - v_dot2_bf16_bf16 v96, v88, v83, 0 - v_dot2_bf16_bf16 v96, v89, v83, 0 op_sel:[0,0,0,1] - v_dot2_bf16_bf16 v97, v90, v83, 0 - v_dot2_bf16_bf16 v97, v91, v83, 0 op_sel:[0,0,0,1] + DQ_PACK v94, v84, v85 + DQ_PACK v95, v86, v87 + DQ_PACK v96, v88, v89 + DQ_PACK v97, v90, v91 s_waitcnt lgkmcnt(0) s_barrier ds_store_b128 v42, v[94:97] From f65e797dfc15a6ecbd53714f92eb365f35555ba7 Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Wed, 24 Jun 2026 15:01:50 -0600 Subject: [PATCH 08/17] [ROCm] W4A16 skinny GEMM: 3-shift dequant extract (high-half materialization) Cut the per-iteration nibble-extract shifts from 7 to 3 by materializing the odd-K nibbles' bf16 directly in the high half of the VGPR. ExLlama packs nibble pair (K_even, K_odd) at bits (4p, 16+4p); bits 16-19 are the high-half bf16 mantissa LSBs, so `(view & 0x000F0000) | 0x4300c308` yields (bf16(-136) low, bf16(128+n) high) with NO shift, and v_dot2_bf16_bf16 consumes b_raw via its high lane (S0.hi*S1.hi) -- same (n-8)*scale. One shift per view (>>4,>>8,>>12) now covers both nibbles of the pair (DQ_EXT_LO/DQ_EXT_HI macros). Perf-neutral at M=1322 (~5.24 ms, same as the 7-shift form -- the extra shifts were already hidden under WMMA/LDS latency, not on the critical path), cos 0.999995. Kept for the leaner op count (the recurring dequant body is now 3 shifts + 8 v_and_or + 8 v_dot2 for 8 elements). Co-authored-by: Claude Signed-off-by: Matthias Gehre --- .../asm/hybrid_w4a16_skinny_gfx1151.amdgcn | 29 ++++++++++--------- 1 file changed, 16 insertions(+), 13 deletions(-) diff --git a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn index d207a11764e4..c4bf527a7f22 100644 --- a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn +++ b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn @@ -2,12 +2,11 @@ .amdhsa_code_object_version 5 ; --- W4A16 dequant macros (assembler text-expansion; zero runtime cost) --- -.macro DQ_EXT0 dst ; nibble@bit0 -> (bf16(128+n) low, bf16(-136) high) - v_and_or_b32 \dst, v59, v92, 0xc3084300 +.macro DQ_EXT_LO dst, src ; even-K: (bf16(128+n) low, bf16(-136) high) + v_and_or_b32 \dst, \src, v92, 0xc3084300 .endm -.macro DQ_EXT dst, sh ; nibble@bit -> (bf16(128+n) low, bf16(-136) high) - v_lshrrev_b32_e32 \dst, \sh, v59 - v_and_or_b32 \dst, \dst, v92, 0xc3084300 +.macro DQ_EXT_HI dst, src ; odd-K: (bf16(-136) low, bf16(128+n) HIGH); nibble already @bit16 + v_and_or_b32 \dst, \src, v93, 0x4300c308 .endm .macro DQ_PACK dst, klo, khi ; two dot2: klo->low half, khi->high half of dst v_dot2_bf16_bf16 \dst, \klo, v83, 0 @@ -268,14 +267,18 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel ds_store_b128 v42, v[55:58] offset:4096 .loc 1 190 22 ; hybrid_w4a16.py:190:22 v_mov_b32 v92, 15 - DQ_EXT0 v84 - DQ_EXT v85, 16 - DQ_EXT v86, 4 - DQ_EXT v87, 20 - DQ_EXT v88, 8 - DQ_EXT v89, 24 - DQ_EXT v90, 12 - DQ_EXT v91, 28 + v_mov_b32 v93, 0xf0000 + DQ_EXT_LO v84, v59 + DQ_EXT_HI v85, v59 + v_lshrrev_b32_e32 v87, 4, v59 + DQ_EXT_LO v86, v87 + DQ_EXT_HI v87, v87 + v_lshrrev_b32_e32 v89, 8, v59 + DQ_EXT_LO v88, v89 + DQ_EXT_HI v89, v89 + v_lshrrev_b32_e32 v91, 12, v59 + DQ_EXT_LO v90, v91 + DQ_EXT_HI v91, v91 v_cndmask_b32_e64 v83, 0x3f80, v83, s0 v_lshl_or_b32 v83, v83, 16, v83 .loc 1 148 20 ; hybrid_w4a16.py:148:20 From 022d996efebc0e5cfc48e93c8571909e2eff440f Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Wed, 24 Jun 2026 15:38:17 -0600 Subject: [PATCH 09/17] [ROCm] W4A16 skinny GEMM: hoist dequant masks to SGPRs Move the two dequant nibble masks (0xF / 0xF0000) from per-iteration VGPR v_mov (v92/v93) to SGPRs (s35/s36) set once in the loop prologue. v_and_or_b32 accepts an SGPR mask + literal OR within the constant-bus limit, and SGPRs survive the loop (ds_load only writes VGPRs) -- unlike v92/v93, which the WMMA B-fragment load `ds_load_b128 v[91:94]` clobbers every iteration, forcing the re-materialization the VGPR form needed. Bumps .amdhsa_next_free_sgpr 35->37 (sgpr_count 37->39) to allocate s35/s36. Removes 2 VALU v_mov from the loop body; perf-neutral at M=1322 (~5.24 ms), cos 0.999995. Frees v92/v93 in the loop and leaves the prologue SGPR setup loop-invariant. Co-authored-by: Claude Signed-off-by: Matthias Gehre --- .../asm/hybrid_w4a16_skinny_gfx1151.amdgcn | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn index c4bf527a7f22..f2d05c03dba9 100644 --- a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn +++ b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn @@ -3,10 +3,10 @@ ; --- W4A16 dequant macros (assembler text-expansion; zero runtime cost) --- .macro DQ_EXT_LO dst, src ; even-K: (bf16(128+n) low, bf16(-136) high) - v_and_or_b32 \dst, \src, v92, 0xc3084300 + v_and_or_b32 \dst, \src, s35, 0xc3084300 .endm .macro DQ_EXT_HI dst, src ; odd-K: (bf16(-136) low, bf16(128+n) HIGH); nibble already @bit16 - v_and_or_b32 \dst, \src, v93, 0x4300c308 + v_and_or_b32 \dst, \src, s36, 0x4300c308 .endm .macro DQ_PACK dst, klo, khi ; two dot2: klo->low half, khi->high half of dst v_dot2_bf16_bf16 \dst, \klo, v83, 0 @@ -194,6 +194,8 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel s_mov_b32 s19, s15 s_mov_b32 s10, s14 s_mov_b32 s11, s15 + s_mov_b32 s35, 15 ; mask_lo (SGPR; survives the loop, ds_load only writes VGPRs) + s_mov_b32 s36, 0xf0000 ; mask_hi .LBB0_2: ; =>This Inner Loop Header: Depth=1 .loc 1 143 26 ; hybrid_w4a16.py:143:26 v_add_nc_u32_e32 v51, s5, v37 @@ -266,8 +268,6 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel ds_store_b128 v42, v[51:54] ds_store_b128 v42, v[55:58] offset:4096 .loc 1 190 22 ; hybrid_w4a16.py:190:22 - v_mov_b32 v92, 15 - v_mov_b32 v93, 0xf0000 DQ_EXT_LO v84, v59 DQ_EXT_HI v85, v59 v_lshrrev_b32_e32 v87, 4, v59 @@ -941,7 +941,7 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel .amdhsa_system_sgpr_workgroup_info 0 .amdhsa_system_vgpr_workitem_id 0 .amdhsa_next_free_vgpr 115 - .amdhsa_next_free_sgpr 35 + .amdhsa_next_free_sgpr 37 .amdhsa_reserve_vcc 1 .amdhsa_float_round_mode_32 0 .amdhsa_float_round_mode_16_64 0 @@ -1177,7 +1177,7 @@ amdhsa.kernels: .max_flat_workgroup_size: 256 .name: _triton_w4a16_skinny_fmt_kernel .private_segment_fixed_size: 0 - .sgpr_count: 37 + .sgpr_count: 39 .sgpr_spill_count: 0 .symbol: _triton_w4a16_skinny_fmt_kernel.kd .uniform_work_group_size: 1 From 4ef173ca8f0fdf9e388f67bbed51d864014991bb Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Wed, 24 Jun 2026 15:45:14 -0600 Subject: [PATCH 10/17] [ROCm] W4A16 skinny GEMM: drop redundant OOB scale cndmask Remove the `v_cndmask v83, 0x3f80, v83, s0` that applied the other=1.0 default to out-of-bounds (offs_n >= N) scale lanes. It is redundant: the scale load address is already forced out-of-bounds for those lanes (`v_cndmask v56, 0x80000000, v56, s0`), so the buffer load returns 0 for them, and the output store masks those columns out entirely -- the OOB column value (0 vs 1.0 scaled) is discarded either way. The u16 load zero-extends, so the following `v_lshl_or v83,v83,16,v83` still broadcasts to (scale,scale). Validated correct at the aligned shape (N=19456, cos 0.999994) AND a non-64-aligned shape (N=19457, last tile = 1 valid + 63 OOB columns, cos 0.999994) -- the latter exercises the OOB path directly. One fewer VALU op per loop iteration. Co-authored-by: Claude Signed-off-by: Matthias Gehre --- .../mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn | 1 - 1 file changed, 1 deletion(-) diff --git a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn index f2d05c03dba9..713aa8781145 100644 --- a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn +++ b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn @@ -279,7 +279,6 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel v_lshrrev_b32_e32 v91, 12, v59 DQ_EXT_LO v90, v91 DQ_EXT_HI v91, v91 - v_cndmask_b32_e64 v83, 0x3f80, v83, s0 v_lshl_or_b32 v83, v83, 16, v83 .loc 1 148 20 ; hybrid_w4a16.py:148:20 s_waitcnt lgkmcnt(0) From e52939807ecfe5a76e2a591178c8c0725b1fce6d Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Wed, 24 Jun 2026 15:50:30 -0600 Subject: [PATCH 11/17] [ROCm] W4A16 skinny GEMM: hoist dequant shifts, break WAW/WAR chains Move the three nibble-extract shifts to the top of the dequant with distinct destination registers (reusing the v_dot output regs v95-v97, dead until the DQ_PACK dot2), then issue all 8 v_and_or against read-only sources into distinct operand regs. This removes the in-place WAR (the old `DQ_EXT_HI v87,v87` read+ wrote the same reg the LO extract had just read) and the shift->and_or serialization, so the 3 shifts are independent and the 8 extracts have no mutual WAW/WAR. Correct at aligned (N=19456) and OOB (N=19457) shapes, cos 0.999994; ~1% faster at M=1322 (5.24 -> ~5.18 ms, near the box noise floor). Co-authored-by: Claude Signed-off-by: Matthias Gehre --- .../asm/hybrid_w4a16_skinny_gfx1151.amdgcn | 24 +++++++++---------- 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn index 713aa8781145..bc245c384c34 100644 --- a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn +++ b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn @@ -268,18 +268,18 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel ds_store_b128 v42, v[51:54] ds_store_b128 v42, v[55:58] offset:4096 .loc 1 190 22 ; hybrid_w4a16.py:190:22 - DQ_EXT_LO v84, v59 - DQ_EXT_HI v85, v59 - v_lshrrev_b32_e32 v87, 4, v59 - DQ_EXT_LO v86, v87 - DQ_EXT_HI v87, v87 - v_lshrrev_b32_e32 v89, 8, v59 - DQ_EXT_LO v88, v89 - DQ_EXT_HI v89, v89 - v_lshrrev_b32_e32 v91, 12, v59 - DQ_EXT_LO v90, v91 - DQ_EXT_HI v91, v91 - v_lshl_or_b32 v83, v83, 16, v83 + v_lshrrev_b32_e32 v95, 4, v59 ; view>>4 + v_lshrrev_b32_e32 v96, 8, v59 ; view>>8 + v_lshrrev_b32_e32 v97, 12, v59 ; view>>12 + DQ_EXT_LO v84, v59 ; K0 + DQ_EXT_HI v85, v59 ; K1 + DQ_EXT_LO v86, v95 ; K2 + DQ_EXT_HI v87, v95 ; K3 + DQ_EXT_LO v88, v96 ; K4 + DQ_EXT_HI v89, v96 ; K5 + DQ_EXT_LO v90, v97 ; K6 + DQ_EXT_HI v91, v97 ; K7 + v_lshl_or_b32 v83, v83, 16, v83 ; (scale, scale) .loc 1 148 20 ; hybrid_w4a16.py:148:20 s_waitcnt lgkmcnt(0) s_barrier From bd8721fc648afc6eeaffddf675ab0efaef36fcfb Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Wed, 24 Jun 2026 16:45:53 -0600 Subject: [PATCH 12/17] [ROCm] W4A16 skinny GEMM: software-pipeline the global loads Prefetch iteration k+1's weight/scale/activation tiles while iteration k's LDS stores and WMMA execute, so the loop-top s_waitcnt vmcnt(0) no longer stalls on a just-issued global load. ATT showed that wait was ~40% of the loop's stall cycles; the activation load alone has ~250-cycle latency that was fully exposed each iteration. Changes: - Decouple the four global loads into a dedicated register window (W->v123, scale->v124, A0/A1->v115-122) so a load destination is never also a downstream consumer's source mid-iteration. - Emit the address/load block once in a prologue (loads iter 0) and once more after the dequant dot2 (prefetches iter k+1), overlapping it with the B-fragment ds_store/ds_load and the WMMA. The loop counter decrement + branch condition move into the relocated in-loop copy. - The last iteration's prefetch reads one K-tile past the end; addresses are already OOB-masked to 0, so it is harmless (correctness unchanged). - next_free_vgpr/vgpr_count 115->128. 115 already rounded up to the 128 VGPR granule, so occupancy is unchanged (12 waves/SIMD). Measured on gfx1151, shape 1322x19456x2560 g32 bf16 (up_proj), interleaved A/B in one process: median 5.239 -> 4.849 ms/call (~7.4% faster, 25.1 -> 27.1 TFLOP/s). Correctness: cos 0.999994 vs Triton reference. Co-authored-by: Claude Signed-off-by: Matthias Gehre --- .../asm/hybrid_w4a16_skinny_gfx1151.amdgcn | 118 ++++++++++++++---- 1 file changed, 92 insertions(+), 26 deletions(-) diff --git a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn index bc245c384c34..9e03073235bb 100644 --- a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn +++ b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn @@ -9,8 +9,8 @@ v_and_or_b32 \dst, \src, s36, 0x4300c308 .endm .macro DQ_PACK dst, klo, khi ; two dot2: klo->low half, khi->high half of dst - v_dot2_bf16_bf16 \dst, \klo, v83, 0 - v_dot2_bf16_bf16 \dst, \khi, v83, 0 op_sel:[0,0,0,1] + v_dot2_bf16_bf16 \dst, \klo, v124, 0 + v_dot2_bf16_bf16 \dst, \khi, v124, 0 op_sel:[0,0,0,1] .endm .text @@ -196,9 +196,9 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel s_mov_b32 s11, s15 s_mov_b32 s35, 15 ; mask_lo (SGPR; survives the loop, ds_load only writes VGPRs) s_mov_b32 s36, 0xf0000 ; mask_hi -.LBB0_2: ; =>This Inner Loop Header: Depth=1 + ; --- prologue: prefetch iter 0 --- .loc 1 143 26 ; hybrid_w4a16.py:143:26 - v_add_nc_u32_e32 v51, s5, v37 + v_add_nc_u32_e32 v125, s5, v37 .loc 1 153 61 ; hybrid_w4a16.py:153:61 v_cmp_gt_i32_e64 s2, s4, v36 .loc 1 193 39 ; hybrid_w4a16.py:193:39 @@ -206,7 +206,7 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel s_abs_i32 s29, s5 s_xor_b32 s31, s3, s6 .loc 1 143 26 ; hybrid_w4a16.py:143:26 - v_cmp_gt_i32_e64 s3, s26, v51 + v_cmp_gt_i32_e64 s3, s26, v125 .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_mul_hi_u32 s30, s29, s7 .loc 1 153 42 ; hybrid_w4a16.py:153:42 @@ -214,7 +214,7 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_mul_i32 s33, s30, s27 .loc 1 154 27 ; hybrid_w4a16.py:154:27 - v_cndmask_b32_e64 v51, 0x80000000, v39, s2 + v_cndmask_b32_e64 v125, 0x80000000, v39, s2 .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_sub_i32 s29, s29, s33 .loc 1 147 41 ; hybrid_w4a16.py:147:41 @@ -223,22 +223,22 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel s_add_i32 s34, s30, 1 s_sub_i32 s33, s29, s27 .loc 1 148 20 ; hybrid_w4a16.py:148:20 - v_cndmask_b32_e64 v52, 0x80000000, v40, s2 + v_cndmask_b32_e64 v126, 0x80000000, v40, s2 .loc 1 147 41 ; hybrid_w4a16.py:147:41 s_and_b32 s2, s1, s3 .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_cmp_ge_u32 s29, s27 .loc 1 148 20 ; hybrid_w4a16.py:148:20 - v_cndmask_b32_e64 v55, 0x80000000, v41, s2 + v_cndmask_b32_e64 v127, 0x80000000, v41, s2 .loc 1 154 27 ; hybrid_w4a16.py:154:27 - buffer_load_b32 v59, v51, s[16:19], 0 offen + buffer_load_b32 v123, v125, s[16:19], 0 offen .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_cselect_b32 s2, s34, s30 s_cselect_b32 s3, s33, s29 s_add_i32 s29, s2, 1 s_cmp_ge_u32 s3, s27 .loc 1 148 20 ; hybrid_w4a16.py:148:20 - buffer_load_b128 v[51:54], v52, s[12:15], 0 offen + buffer_load_b128 v[115:118], v126, s[12:15], 0 offen .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_cselect_b32 s2, s29, s2 .loc 1 141 28 ; hybrid_w4a16.py:141:28 @@ -252,34 +252,33 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel .loc 1 141 28 ; hybrid_w4a16.py:141:28 v_add_nc_u32_e32 v40, 64, v40 .loc 1 224 16 ; hybrid_w4a16.py:224:16 - v_add_lshl_u32 v56, s2, v38, 1 + v_add_lshl_u32 v125, s2, v38, 1 .loc 1 141 28 ; hybrid_w4a16.py:141:28 v_add_nc_u32_e32 v41, 64, v41 - s_add_i32 s28, s28, -1 s_add_i32 s5, s5, 32 - s_cmp_lg_u32 s28, 0 .loc 1 224 16 ; hybrid_w4a16.py:224:16 - v_cndmask_b32_e64 v56, 0x80000000, v56, s0 - buffer_load_u16 v83, v56, s[8:11], 0 offen + v_cndmask_b32_e64 v125, 0x80000000, v125, s0 + buffer_load_u16 v124, v125, s[8:11], 0 offen .loc 1 148 20 ; hybrid_w4a16.py:148:20 - buffer_load_b128 v[55:58], v55, s[12:15], 0 offen + buffer_load_b128 v[119:122], v127, s[12:15], 0 offen +.LBB0_2: ; =>This Inner Loop Header: Depth=1 s_waitcnt vmcnt(0) s_barrier - ds_store_b128 v42, v[51:54] - ds_store_b128 v42, v[55:58] offset:4096 + ds_store_b128 v42, v[115:118] + ds_store_b128 v42, v[119:122] offset:4096 .loc 1 190 22 ; hybrid_w4a16.py:190:22 - v_lshrrev_b32_e32 v95, 4, v59 ; view>>4 - v_lshrrev_b32_e32 v96, 8, v59 ; view>>8 - v_lshrrev_b32_e32 v97, 12, v59 ; view>>12 - DQ_EXT_LO v84, v59 ; K0 - DQ_EXT_HI v85, v59 ; K1 + v_lshrrev_b32_e32 v95, 4, v123 ; view>>4 + v_lshrrev_b32_e32 v96, 8, v123 ; view>>8 + v_lshrrev_b32_e32 v97, 12, v123 ; view>>12 + DQ_EXT_LO v84, v123 ; K0 + DQ_EXT_HI v85, v123 ; K1 DQ_EXT_LO v86, v95 ; K2 DQ_EXT_HI v87, v95 ; K3 DQ_EXT_LO v88, v96 ; K4 DQ_EXT_HI v89, v96 ; K5 DQ_EXT_LO v90, v97 ; K6 DQ_EXT_HI v91, v97 ; K7 - v_lshl_or_b32 v83, v83, 16, v83 ; (scale, scale) + v_lshl_or_b32 v124, v124, 16, v124 ; (scale, scale) .loc 1 148 20 ; hybrid_w4a16.py:148:20 s_waitcnt lgkmcnt(0) s_barrier @@ -296,6 +295,73 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel DQ_PACK v95, v86, v87 DQ_PACK v96, v88, v89 DQ_PACK v97, v90, v91 + ; --- prefetch iter k+1 (overlaps ds_store B + ds_load B + WMMA) --- + .loc 1 143 26 ; hybrid_w4a16.py:143:26 + v_add_nc_u32_e32 v125, s5, v37 + .loc 1 153 61 ; hybrid_w4a16.py:153:61 + v_cmp_gt_i32_e64 s2, s4, v36 + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_ashr_i32 s3, s5, 31 + s_abs_i32 s29, s5 + s_xor_b32 s31, s3, s6 + .loc 1 143 26 ; hybrid_w4a16.py:143:26 + v_cmp_gt_i32_e64 s3, s26, v125 + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_mul_hi_u32 s30, s29, s7 + .loc 1 153 42 ; hybrid_w4a16.py:153:42 + s_and_b32 s2, s0, s2 + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_mul_i32 s33, s30, s27 + .loc 1 154 27 ; hybrid_w4a16.py:154:27 + v_cndmask_b32_e64 v125, 0x80000000, v39, s2 + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_sub_i32 s29, s29, s33 + .loc 1 147 41 ; hybrid_w4a16.py:147:41 + s_and_b32 s2, vcc_lo, s3 + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_add_i32 s34, s30, 1 + s_sub_i32 s33, s29, s27 + .loc 1 148 20 ; hybrid_w4a16.py:148:20 + v_cndmask_b32_e64 v126, 0x80000000, v40, s2 + .loc 1 147 41 ; hybrid_w4a16.py:147:41 + s_and_b32 s2, s1, s3 + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_cmp_ge_u32 s29, s27 + .loc 1 148 20 ; hybrid_w4a16.py:148:20 + v_cndmask_b32_e64 v127, 0x80000000, v41, s2 + .loc 1 154 27 ; hybrid_w4a16.py:154:27 + buffer_load_b32 v123, v125, s[16:19], 0 offen + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_cselect_b32 s2, s34, s30 + s_cselect_b32 s3, s33, s29 + s_add_i32 s29, s2, 1 + s_cmp_ge_u32 s3, s27 + .loc 1 148 20 ; hybrid_w4a16.py:148:20 + buffer_load_b128 v[115:118], v126, s[12:15], 0 offen + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_cselect_b32 s2, s29, s2 + .loc 1 141 28 ; hybrid_w4a16.py:141:28 + v_add_nc_u32_e32 v36, 4, v36 + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_xor_b32 s2, s2, s31 + .loc 1 141 28 ; hybrid_w4a16.py:141:28 + v_add_nc_u32_e32 v39, 16, v39 + .loc 1 193 39 ; hybrid_w4a16.py:193:39 + s_sub_i32 s2, s2, s31 + .loc 1 141 28 ; hybrid_w4a16.py:141:28 + v_add_nc_u32_e32 v40, 64, v40 + .loc 1 224 16 ; hybrid_w4a16.py:224:16 + v_add_lshl_u32 v125, s2, v38, 1 + .loc 1 141 28 ; hybrid_w4a16.py:141:28 + v_add_nc_u32_e32 v41, 64, v41 + s_add_i32 s28, s28, -1 + s_add_i32 s5, s5, 32 + s_cmp_lg_u32 s28, 0 + .loc 1 224 16 ; hybrid_w4a16.py:224:16 + v_cndmask_b32_e64 v125, 0x80000000, v125, s0 + buffer_load_u16 v124, v125, s[8:11], 0 offen + .loc 1 148 20 ; hybrid_w4a16.py:148:20 + buffer_load_b128 v[119:122], v127, s[12:15], 0 offen s_waitcnt lgkmcnt(0) s_barrier ds_store_b128 v42, v[94:97] @@ -939,7 +1005,7 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel .amdhsa_system_sgpr_workgroup_id_z 0 .amdhsa_system_sgpr_workgroup_info 0 .amdhsa_system_vgpr_workitem_id 0 - .amdhsa_next_free_vgpr 115 + .amdhsa_next_free_vgpr 128 .amdhsa_next_free_sgpr 37 .amdhsa_reserve_vcc 1 .amdhsa_float_round_mode_32 0 @@ -1181,7 +1247,7 @@ amdhsa.kernels: .symbol: _triton_w4a16_skinny_fmt_kernel.kd .uniform_work_group_size: 1 .uses_dynamic_stack: false - .vgpr_count: 115 + .vgpr_count: 128 .vgpr_spill_count: 0 .wavefront_size: 32 .workgroup_processor_mode: 0 From 8f95f19ac70d4bf157b126d41588fa7315a58797 Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Wed, 24 Jun 2026 23:59:03 -0600 Subject: [PATCH 13/17] [ROCm] W4A16 skinny GEMM: strip .loc debug lines, frame the main loop Pure readability cleanup of the hand-asm kernel, no codegen change. Changes: - Remove all 311 .loc debug-line directives (they tracked the original Triton source lines and only added noise to the hand-edited ISA; the .file directive stays so the assembler is happy). - Surround the K-reduction loop with blank lines and a header comment describing what it iterates over (ceil(K / BLOCK_K) tiles: prefetched loads -> int4 dequant -> bf16 WMMA accumulate), and tag the backedge. Correctness unchanged: cos 0.999995 vs Triton reference. Co-authored-by: Claude Signed-off-by: Matthias Gehre --- .../asm/hybrid_w4a16_skinny_gfx1151.amdgcn | 325 +----------------- 1 file changed, 13 insertions(+), 312 deletions(-) diff --git a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn index 9e03073235bb..957eb75714da 100644 --- a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn +++ b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn @@ -20,7 +20,6 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel .Lfunc_begin0: .file 1 "/scratch/mgehre/vllm/vllm/model_executor/kernels/linear/mixed_precision" "hybrid_w4a16.py" - .loc 1 81 0 ; hybrid_w4a16.py:81:0 .cfi_sections .debug_frame .cfi_startproc ; %bb.0: @@ -29,11 +28,8 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel s_load_b128 s[20:23], s[0:1], 0x20 v_dual_mov_b32 v8, 0 :: v_dual_and_b32 v33, 15, v0 .Ltmp0: - .loc 1 125 44 prologue_end ; hybrid_w4a16.py:125:44 v_and_b32_e32 v34, 0xc0, v0 - .loc 1 126 44 ; hybrid_w4a16.py:126:44 v_and_b32_e32 v35, 32, v0 - .loc 1 125 21 ; hybrid_w4a16.py:125:21 s_lshl_b32 s24, s2, 7 s_delay_alu instid0(VALU_DEP_3) v_mov_b32_e32 v7, v8 @@ -69,48 +65,37 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel v_mov_b32_e32 v25, v8 .Ltmp1: .file 2 "/scratch/mgehre/vllm/.venv/lib/python3.12/site-packages/triton/language" "standard.py" - .loc 2 43 17 ; standard.py:43:17 @[ hybrid_w4a16.py:141:39 ] s_waitcnt lgkmcnt(0) s_add_i32 s2, s26, 31 .Ltmp2: - .loc 1 126 21 ; hybrid_w4a16.py:126:21 s_lshl_b32 s25, s3, 6 - .loc 1 141 28 ; hybrid_w4a16.py:141:28 s_cmp_lt_i32 s2, 32 s_cbranch_scc1 .LBB0_3 ; %bb.1: ; %.lr.ph - .loc 1 0 28 is_stmt 0 ; hybrid_w4a16.py:0:28 s_clause 0x2 s_load_b128 s[4:7], s[0:1], 0x34 s_load_b128 s[16:19], s[0:1], 0x0 s_load_b64 s[8:9], s[0:1], 0x10 - .loc 1 125 44 is_stmt 1 ; hybrid_w4a16.py:125:44 v_lshrrev_b32_e32 v1, 2, v0 .Ltmp3: - .loc 2 43 30 ; standard.py:43:30 @[ hybrid_w4a16.py:141:39 ] s_ashr_i32 s0, s2, 31 v_dual_mov_b32 v25, 0 :: v_dual_and_b32 v36, 3, v0 s_lshr_b32 s0, s0, 27 s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2) .Ltmp4: - .loc 1 125 44 ; hybrid_w4a16.py:125:44 v_or_b32_e32 v5, s24, v1 .Ltmp5: - .loc 2 43 30 ; standard.py:43:30 @[ hybrid_w4a16.py:141:39 ] s_add_i32 s2, s2, s0 v_dual_mov_b32 v29, v25 :: v_dual_lshlrev_b32 v4, 3, v0 v_dual_mov_b32 v26, v25 :: v_dual_lshlrev_b32 v3, 1, v0 s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4) .Ltmp6: - .loc 1 125 31 ; hybrid_w4a16.py:125:31 v_or_b32_e32 v8, 64, v5 - .loc 1 147 36 ; hybrid_w4a16.py:147:36 v_cmp_gt_i32_e32 vcc_lo, s22, v5 v_dual_mov_b32 v31, v25 :: v_dual_and_b32 v4, 48, v4 v_dual_mov_b32 v27, v25 :: v_dual_lshlrev_b32 v2, 4, v0 s_waitcnt lgkmcnt(0) s_abs_i32 s27, s7 - .loc 1 126 31 ; hybrid_w4a16.py:126:31 v_or_b32_e32 v6, s25, v1 s_cvt_f32_u32 s0, s27 v_dual_mov_b32 v30, v25 :: v_dual_and_b32 v3, 48, v3 @@ -120,23 +105,19 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel v_dual_mov_b32 v18, v25 :: v_dual_add_nc_u32 v1, s24, v1 v_lshl_or_b32 v4, v33, 6, v4 .Ltmp7: - .loc 2 43 30 ; standard.py:43:30 @[ hybrid_w4a16.py:141:39 ] s_ashr_i32 s28, s2, 5 v_xor_b32_e32 v2, v2, v3 v_dual_mov_b32 v20, v25 :: v_dual_add_nc_u32 v3, 64, v1 .Ltmp8: - .loc 1 147 36 ; hybrid_w4a16.py:147:36 v_cmp_gt_i32_e64 s1, s22, v8 v_lshl_or_b32 v8, v34, 4, v4 s_delay_alu instid0(TRANS32_DEP_1) v_readfirstlane_b32 s2, v5 v_lshl_or_b32 v4, v35, 5, v4 - .loc 1 141 28 ; hybrid_w4a16.py:141:28 v_mul_lo_u32 v7, s5, v7 v_mul_lo_u32 v1, s26, v1 v_mul_lo_u32 v3, s26, v3 s_mul_f32 s2, s2, 0x4f7ffffe - .loc 1 153 36 ; hybrid_w4a16.py:153:36 v_cmp_gt_i32_e64 s0, s23, v6 v_mul_lo_u32 v38, s6, v6 v_xor_b32_e32 v5, 16, v8 @@ -151,7 +132,6 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel s_sub_i32 s3, 0, s27 v_dual_mov_b32 v28, v25 :: v_dual_lshlrev_b32 v37, 3, v36 s_mul_i32 s3, s3, s2 - .loc 1 141 28 ; hybrid_w4a16.py:141:28 v_lshl_add_u32 v39, v7, 2, v13 v_lshl_add_u32 v40, v1, 1, v14 v_lshl_add_u32 v41, v3, 1, v14 @@ -196,77 +176,56 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel s_mov_b32 s11, s15 s_mov_b32 s35, 15 ; mask_lo (SGPR; survives the loop, ds_load only writes VGPRs) s_mov_b32 s36, 0xf0000 ; mask_hi + + + + ; =========================================================================== + ; main loop: K-reduction over in_features (K), BLOCK_K=32 per iteration + ; iterates ceil(K / BLOCK_K) tiles; each tile: prefetched global loads -> + ; int4 dequant -> bf16 WMMA accumulate into the output fragment + ; =========================================================================== ; --- prologue: prefetch iter 0 --- - .loc 1 143 26 ; hybrid_w4a16.py:143:26 v_add_nc_u32_e32 v125, s5, v37 - .loc 1 153 61 ; hybrid_w4a16.py:153:61 v_cmp_gt_i32_e64 s2, s4, v36 - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_ashr_i32 s3, s5, 31 s_abs_i32 s29, s5 s_xor_b32 s31, s3, s6 - .loc 1 143 26 ; hybrid_w4a16.py:143:26 v_cmp_gt_i32_e64 s3, s26, v125 - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_mul_hi_u32 s30, s29, s7 - .loc 1 153 42 ; hybrid_w4a16.py:153:42 s_and_b32 s2, s0, s2 - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_mul_i32 s33, s30, s27 - .loc 1 154 27 ; hybrid_w4a16.py:154:27 v_cndmask_b32_e64 v125, 0x80000000, v39, s2 - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_sub_i32 s29, s29, s33 - .loc 1 147 41 ; hybrid_w4a16.py:147:41 s_and_b32 s2, vcc_lo, s3 - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_add_i32 s34, s30, 1 s_sub_i32 s33, s29, s27 - .loc 1 148 20 ; hybrid_w4a16.py:148:20 v_cndmask_b32_e64 v126, 0x80000000, v40, s2 - .loc 1 147 41 ; hybrid_w4a16.py:147:41 s_and_b32 s2, s1, s3 - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_cmp_ge_u32 s29, s27 - .loc 1 148 20 ; hybrid_w4a16.py:148:20 v_cndmask_b32_e64 v127, 0x80000000, v41, s2 - .loc 1 154 27 ; hybrid_w4a16.py:154:27 buffer_load_b32 v123, v125, s[16:19], 0 offen - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_cselect_b32 s2, s34, s30 s_cselect_b32 s3, s33, s29 s_add_i32 s29, s2, 1 s_cmp_ge_u32 s3, s27 - .loc 1 148 20 ; hybrid_w4a16.py:148:20 buffer_load_b128 v[115:118], v126, s[12:15], 0 offen - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_cselect_b32 s2, s29, s2 - .loc 1 141 28 ; hybrid_w4a16.py:141:28 v_add_nc_u32_e32 v36, 4, v36 - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_xor_b32 s2, s2, s31 - .loc 1 141 28 ; hybrid_w4a16.py:141:28 v_add_nc_u32_e32 v39, 16, v39 - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_sub_i32 s2, s2, s31 - .loc 1 141 28 ; hybrid_w4a16.py:141:28 v_add_nc_u32_e32 v40, 64, v40 - .loc 1 224 16 ; hybrid_w4a16.py:224:16 v_add_lshl_u32 v125, s2, v38, 1 - .loc 1 141 28 ; hybrid_w4a16.py:141:28 v_add_nc_u32_e32 v41, 64, v41 s_add_i32 s5, s5, 32 - .loc 1 224 16 ; hybrid_w4a16.py:224:16 v_cndmask_b32_e64 v125, 0x80000000, v125, s0 buffer_load_u16 v124, v125, s[8:11], 0 offen - .loc 1 148 20 ; hybrid_w4a16.py:148:20 buffer_load_b128 v[119:122], v127, s[12:15], 0 offen .LBB0_2: ; =>This Inner Loop Header: Depth=1 s_waitcnt vmcnt(0) s_barrier ds_store_b128 v42, v[115:118] ds_store_b128 v42, v[119:122] offset:4096 - .loc 1 190 22 ; hybrid_w4a16.py:190:22 v_lshrrev_b32_e32 v95, 4, v123 ; view>>4 v_lshrrev_b32_e32 v96, 8, v123 ; view>>8 v_lshrrev_b32_e32 v97, 12, v123 ; view>>12 @@ -279,7 +238,6 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel DQ_EXT_LO v90, v97 ; K6 DQ_EXT_HI v91, v97 ; K7 v_lshl_or_b32 v124, v124, 16, v124 ; (scale, scale) - .loc 1 148 20 ; hybrid_w4a16.py:148:20 s_waitcnt lgkmcnt(0) s_barrier ds_load_b128 v[51:54], v43 @@ -290,77 +248,48 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel ds_load_b128 v[75:78], v45 offset:4096 ds_load_b128 v[71:74], v46 ds_load_b128 v[79:82], v46 offset:4096 - .loc 1 234 50 ; hybrid_w4a16.py:234:50 DQ_PACK v94, v84, v85 DQ_PACK v95, v86, v87 DQ_PACK v96, v88, v89 DQ_PACK v97, v90, v91 ; --- prefetch iter k+1 (overlaps ds_store B + ds_load B + WMMA) --- - .loc 1 143 26 ; hybrid_w4a16.py:143:26 v_add_nc_u32_e32 v125, s5, v37 - .loc 1 153 61 ; hybrid_w4a16.py:153:61 v_cmp_gt_i32_e64 s2, s4, v36 - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_ashr_i32 s3, s5, 31 s_abs_i32 s29, s5 s_xor_b32 s31, s3, s6 - .loc 1 143 26 ; hybrid_w4a16.py:143:26 v_cmp_gt_i32_e64 s3, s26, v125 - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_mul_hi_u32 s30, s29, s7 - .loc 1 153 42 ; hybrid_w4a16.py:153:42 s_and_b32 s2, s0, s2 - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_mul_i32 s33, s30, s27 - .loc 1 154 27 ; hybrid_w4a16.py:154:27 v_cndmask_b32_e64 v125, 0x80000000, v39, s2 - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_sub_i32 s29, s29, s33 - .loc 1 147 41 ; hybrid_w4a16.py:147:41 s_and_b32 s2, vcc_lo, s3 - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_add_i32 s34, s30, 1 s_sub_i32 s33, s29, s27 - .loc 1 148 20 ; hybrid_w4a16.py:148:20 v_cndmask_b32_e64 v126, 0x80000000, v40, s2 - .loc 1 147 41 ; hybrid_w4a16.py:147:41 s_and_b32 s2, s1, s3 - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_cmp_ge_u32 s29, s27 - .loc 1 148 20 ; hybrid_w4a16.py:148:20 v_cndmask_b32_e64 v127, 0x80000000, v41, s2 - .loc 1 154 27 ; hybrid_w4a16.py:154:27 buffer_load_b32 v123, v125, s[16:19], 0 offen - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_cselect_b32 s2, s34, s30 s_cselect_b32 s3, s33, s29 s_add_i32 s29, s2, 1 s_cmp_ge_u32 s3, s27 - .loc 1 148 20 ; hybrid_w4a16.py:148:20 buffer_load_b128 v[115:118], v126, s[12:15], 0 offen - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_cselect_b32 s2, s29, s2 - .loc 1 141 28 ; hybrid_w4a16.py:141:28 v_add_nc_u32_e32 v36, 4, v36 - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_xor_b32 s2, s2, s31 - .loc 1 141 28 ; hybrid_w4a16.py:141:28 v_add_nc_u32_e32 v39, 16, v39 - .loc 1 193 39 ; hybrid_w4a16.py:193:39 s_sub_i32 s2, s2, s31 - .loc 1 141 28 ; hybrid_w4a16.py:141:28 v_add_nc_u32_e32 v40, 64, v40 - .loc 1 224 16 ; hybrid_w4a16.py:224:16 v_add_lshl_u32 v125, s2, v38, 1 - .loc 1 141 28 ; hybrid_w4a16.py:141:28 v_add_nc_u32_e32 v41, 64, v41 s_add_i32 s28, s28, -1 s_add_i32 s5, s5, 32 s_cmp_lg_u32 s28, 0 - .loc 1 224 16 ; hybrid_w4a16.py:224:16 v_cndmask_b32_e64 v125, 0x80000000, v125, s0 buffer_load_u16 v124, v125, s[8:11], 0 offen - .loc 1 148 20 ; hybrid_w4a16.py:148:20 buffer_load_b128 v[119:122], v127, s[12:15], 0 offen s_waitcnt lgkmcnt(0) s_barrier @@ -387,163 +316,94 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel v_wmma_f32_16x16x16_bf16 v[17:24], v[107:114], v[67:74], v[17:24] v_wmma_f32_16x16x16_bf16 v[9:16], v[99:106], v[75:82], v[9:16] v_wmma_f32_16x16x16_bf16 v[1:8], v[107:114], v[75:82], v[1:8] - .loc 1 141 28 ; hybrid_w4a16.py:141:28 - s_cbranch_scc1 .LBB0_2 + s_cbranch_scc1 .LBB0_2 ; backedge: branch to next K-tile + ; --- end of main loop (fall through to epilogue) --- + + + .LBB0_3: ; %._crit_edge - .loc 1 125 44 ; hybrid_w4a16.py:125:44 v_lshrrev_b32_e32 v34, 2, v34 - .loc 1 126 44 ; hybrid_w4a16.py:126:44 v_lshrrev_b32_e32 v0, 4, v0 v_lshrrev_b32_e32 v35, 1, v35 - .loc 1 244 21 ; hybrid_w4a16.py:244:21 s_mul_i32 s0, s23, s24 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_bfe_u32 v63, v26, 16, 1 - .loc 1 125 44 ; hybrid_w4a16.py:125:44 v_or_b32_e32 v33, v34, v33 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_bfe_u32 v34, v25, 16, 1 - .loc 1 126 44 ; hybrid_w4a16.py:126:44 v_and_or_b32 v39, v0, 1, v35 - .loc 1 244 43 ; hybrid_w4a16.py:244:43 s_add_i32 s0, s0, s25 s_lshl_b32 s1, s23, 6 - .loc 1 125 44 ; hybrid_w4a16.py:125:44 v_or_b32_e32 v66, s24, v33 - .loc 1 244 43 ; hybrid_w4a16.py:244:43 v_mul_lo_u32 v33, s23, v33 - .loc 1 126 44 ; hybrid_w4a16.py:126:44 v_or_b32_e32 v0, 46, v39 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v34, v25, v34, 0x7fff - .loc 1 126 31 ; hybrid_w4a16.py:126:31 v_or_b32_e32 v68, s25, v39 - .loc 1 126 44 is_stmt 0 ; hybrid_w4a16.py:126:44 v_or_b32_e32 v62, 2, v39 - .loc 1 245 32 is_stmt 1 ; hybrid_w4a16.py:245:32 v_cmp_gt_i32_e64 s16, s22, v66 - .loc 1 126 31 ; hybrid_w4a16.py:126:31 v_or_b32_e32 v43, s25, v0 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v34, 16, v34 - .loc 1 244 43 ; hybrid_w4a16.py:244:43 v_add_nc_u32_e32 v99, s0, v33 - .loc 1 245 56 ; hybrid_w4a16.py:245:56 v_cmp_gt_i32_e64 s15, s23, v68 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s17, v25, v25 - .loc 1 126 31 ; hybrid_w4a16.py:126:31 v_or_b32_e32 v65, s25, v62 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v63, v26, v63, 0x7fff - .loc 1 244 43 ; hybrid_w4a16.py:244:43 v_add3_u32 v33, s0, s1, v33 - .loc 1 245 56 ; hybrid_w4a16.py:245:56 v_cmp_gt_i32_e64 s0, s23, v43 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v43, v99, v39, 1 - .loc 1 126 44 ; hybrid_w4a16.py:126:44 v_or_b32_e32 v61, 4, v39 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_bfe_u32 v67, v27, 16, 1 v_cndmask_b32_e64 v25, 0x7fff, v34, s17 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s17, s16, s15 - .loc 1 126 44 ; hybrid_w4a16.py:126:44 v_or_b32_e32 v35, 44, v39 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v63, 16, v63 - .loc 1 245 56 ; hybrid_w4a16.py:245:56 v_cmp_gt_i32_e64 s14, s23, v65 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v34, 0x80000000, v43, s17 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s17, v26, v26 - .loc 1 126 31 ; hybrid_w4a16.py:126:31 v_or_b32_e32 v64, s25, v61 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v67, v27, v67, 0x7fff - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v43, v99, v62, 1 - .loc 1 126 31 ; hybrid_w4a16.py:126:31 v_or_b32_e32 v44, s25, v35 - .loc 1 126 44 is_stmt 0 ; hybrid_w4a16.py:126:44 v_or_b32_e32 v56, 6, v39 - .loc 1 243 23 is_stmt 1 ; hybrid_w4a16.py:243:23 v_bfe_u32 v70, v28, 16, 1 v_cndmask_b32_e64 v26, 0x7fff, v63, s17 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s17, s16, s14 - .loc 1 126 44 ; hybrid_w4a16.py:126:44 v_or_b32_e32 v36, 42, v39 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v67, 16, v67 - .loc 1 245 56 ; hybrid_w4a16.py:245:56 v_cmp_gt_i32_e64 s13, s23, v64 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v43, 0x80000000, v43, s17 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s17, v27, v27 - .loc 1 126 31 ; hybrid_w4a16.py:126:31 v_or_b32_e32 v60, s25, v56 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v70, v28, v70, 0x7fff - .loc 1 245 56 ; hybrid_w4a16.py:245:56 v_cmp_gt_i32_e64 s1, s23, v44 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v44, v99, v61, 1 - .loc 1 126 31 ; hybrid_w4a16.py:126:31 v_or_b32_e32 v45, s25, v36 - .loc 1 126 44 is_stmt 0 ; hybrid_w4a16.py:126:44 v_or_b32_e32 v55, 8, v39 - .loc 1 243 23 is_stmt 1 ; hybrid_w4a16.py:243:23 v_bfe_u32 v71, v29, 16, 1 v_cndmask_b32_e64 v27, 0x7fff, v67, s17 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s17, s16, s13 - .loc 1 126 44 ; hybrid_w4a16.py:126:44 v_or_b32_e32 v37, 40, v39 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v70, 16, v70 - .loc 1 245 56 ; hybrid_w4a16.py:245:56 v_cmp_gt_i32_e64 s12, s23, v60 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v44, 0x80000000, v44, s17 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s17, v28, v28 - .loc 1 126 31 ; hybrid_w4a16.py:126:31 v_or_b32_e32 v59, s25, v55 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v71, v29, v71, 0x7fff - .loc 1 245 56 ; hybrid_w4a16.py:245:56 v_cmp_gt_i32_e64 s2, s23, v45 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v45, v99, v56, 1 - .loc 1 126 44 ; hybrid_w4a16.py:126:44 v_or_b32_e32 v38, 38, v39 v_or_b32_e32 v40, 36, v39 v_or_b32_e32 v41, 34, v39 v_or_b32_e32 v42, 32, v39 - .loc 1 126 31 is_stmt 0 ; hybrid_w4a16.py:126:31 v_or_b32_e32 v46, s25, v37 - .loc 1 126 44 ; hybrid_w4a16.py:126:44 v_or_b32_e32 v49, 14, v39 v_or_b32_e32 v52, 12, v39 v_or_b32_e32 v53, 10, v39 - .loc 1 243 23 is_stmt 1 ; hybrid_w4a16.py:243:23 v_bfe_u32 v72, v30, 16, 1 v_cndmask_b32_e64 v28, 0x7fff, v70, s17 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s17, s16, s12 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v71, 16, v71 - .loc 1 245 56 ; hybrid_w4a16.py:245:56 v_cmp_gt_i32_e64 s11, s23, v59 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v45, 0x80000000, v45, s17 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s17, v29, v29 - .loc 1 126 31 ; hybrid_w4a16.py:126:31 v_or_b32_e32 v47, s25, v38 v_or_b32_e32 v48, s25, v40 v_or_b32_e32 v50, s25, v41 @@ -551,24 +411,15 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel v_or_b32_e32 v54, s25, v49 v_or_b32_e32 v57, s25, v52 v_or_b32_e32 v58, s25, v53 - .loc 1 125 31 ; hybrid_w4a16.py:125:31 v_or_b32_e32 v69, 64, v66 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v72, v30, v72, 0x7fff - .loc 1 245 56 ; hybrid_w4a16.py:245:56 v_cmp_gt_i32_e64 s3, s23, v46 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v46, v99, v55, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_bfe_u32 v73, v31, 16, 1 v_cndmask_b32_e64 v29, 0x7fff, v71, s17 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s17, s16, s11 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v72, 16, v72 - .loc 1 245 32 ; hybrid_w4a16.py:245:32 v_cmp_gt_i32_e32 vcc_lo, s22, v69 - .loc 1 245 56 is_stmt 0 ; hybrid_w4a16.py:245:56 v_cmp_gt_i32_e64 s10, s23, v58 v_cmp_gt_i32_e64 s9, s23, v57 v_cmp_gt_i32_e64 s8, s23, v54 @@ -576,15 +427,12 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel v_cmp_gt_i32_e64 s6, s23, v50 v_cmp_gt_i32_e64 s5, s23, v48 v_cmp_gt_i32_e64 s4, s23, v47 - .loc 1 246 21 is_stmt 1 ; hybrid_w4a16.py:246:21 s_and_b32 s21, s21, 0xffff s_mov_b32 s23, 0x31027000 s_mov_b32 s22, 0x7ffffffe v_cndmask_b32_e64 v46, 0x80000000, v46, s17 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s17, v30, v30 v_add3_u32 v73, v31, v73, 0x7fff - .loc 1 246 21 ; hybrid_w4a16.py:246:21 s_clause 0x4 buffer_store_b16 v25, v34, s[20:23], 0 offen buffer_store_b16 v26, v43, s[20:23], 0 offen @@ -592,140 +440,87 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel buffer_store_b16 v28, v45, s[20:23], 0 offen buffer_store_b16 v29, v46, s[20:23], 0 offen v_add_lshl_u32 v25, v99, v53, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_bfe_u32 v74, v32, 16, 1 v_cndmask_b32_e64 v26, 0x7fff, v72, s17 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s17, s16, s10 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v73, 16, v73 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v25, 0x80000000, v25, s17 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s17, v31, v31 v_add3_u32 v74, v32, v74, 0x7fff - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v27, v99, v52, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_bfe_u32 v75, v17, 16, 1 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v29, v99, v49, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cndmask_b32_e64 v28, 0x7fff, v73, s17 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s17, s16, s9 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v74, 16, v74 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v27, 0x80000000, v27, s17 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s17, v32, v32 v_add3_u32 v75, v17, v75, 0x7fff v_bfe_u32 v76, v18, 16, 1 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v31, v99, v42, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_bfe_u32 v77, v19, 16, 1 v_cndmask_b32_e64 v30, 0x7fff, v74, s17 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s17, s16, s8 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v75, 16, v75 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v29, 0x80000000, v29, s17 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s17, v17, v17 v_add3_u32 v76, v18, v76, 0x7fff v_add3_u32 v77, v19, v77, 0x7fff - .loc 1 246 21 ; hybrid_w4a16.py:246:21 s_clause 0x2 buffer_store_b16 v26, v25, s[20:23], 0 offen buffer_store_b16 v28, v27, s[20:23], 0 offen buffer_store_b16 v30, v29, s[20:23], 0 offen v_add_lshl_u32 v25, v99, v41, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cndmask_b32_e64 v17, 0x7fff, v75, s17 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s17, s16, s7 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v76, 16, v76 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v31, 0x80000000, v31, s17 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s17, v18, v18 v_bfe_u32 v78, v20, 16, 1 v_lshrrev_b32_e32 v77, 16, v77 v_bfe_u32 v79, v21, 16, 1 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 buffer_store_b16 v17, v31, s[20:23], 0 offen - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cndmask_b32_e64 v17, 0x7fff, v76, s17 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s17, s16, s6 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v78, v20, v78, 0x7fff - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v18, 0x80000000, v25, s17 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s17, v19, v19 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v25, v99, v40, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v79, v21, v79, 0x7fff v_lshrrev_b32_e32 v78, 16, v78 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v26, v99, v38, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cndmask_b32_e64 v19, 0x7fff, v77, s17 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s17, s16, s5 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_bfe_u32 v80, v22, 16, 1 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v25, 0x80000000, v25, s17 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s17, v20, v20 v_lshrrev_b32_e32 v79, 16, v79 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v27, v99, v37, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v80, v22, v80, 0x7fff v_bfe_u32 v81, v23, 16, 1 v_cndmask_b32_e64 v20, 0x7fff, v78, s17 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s17, s16, s4 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v28, v99, v36, 1 v_cndmask_b32_e64 v26, 0x80000000, v26, s17 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s17, v21, v21 v_lshrrev_b32_e32 v80, 16, v80 v_add3_u32 v81, v23, v81, 0x7fff v_bfe_u32 v82, v24, 16, 1 v_bfe_u32 v83, v9, 16, 1 v_cndmask_b32_e64 v21, 0x7fff, v79, s17 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s17, s16, s3 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v81, 16, v81 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v27, 0x80000000, v27, s17 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s17, v22, v22 v_add3_u32 v82, v24, v82, 0x7fff v_add3_u32 v83, v9, v83, 0x7fff v_bfe_u32 v84, v10, 16, 1 v_bfe_u32 v85, v11, 16, 1 v_cndmask_b32_e64 v22, 0x7fff, v80, s17 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s17, s16, s2 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v82, 16, v82 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v28, 0x80000000, v28, s17 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s17, v23, v23 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 s_clause 0x4 buffer_store_b16 v17, v18, s[20:23], 0 offen buffer_store_b16 v19, v25, s[20:23], 0 offen @@ -734,122 +529,76 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel buffer_store_b16 v22, v28, s[20:23], 0 offen v_add_lshl_u32 v17, v99, v35, 1 v_add_lshl_u32 v19, v99, v0, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v83, 16, v83 v_cndmask_b32_e64 v18, 0x7fff, v81, s17 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s17, s16, s1 s_and_b32 s16, s16, s0 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v17, 0x80000000, v17, s17 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s17, v24, v24 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v19, 0x80000000, v19, s16 v_add_lshl_u32 v21, v33, v39, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s16, v9, v9 v_add3_u32 v84, v10, v84, 0x7fff v_bfe_u32 v86, v12, 16, 1 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s15, vcc_lo, s15 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_bfe_u32 v87, v13, 16, 1 v_cndmask_b32_e64 v20, 0x7fff, v82, s17 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v22, v33, v62, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_bfe_u32 v88, v14, 16, 1 v_cndmask_b32_e64 v9, 0x7fff, v83, s16 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v21, 0x80000000, v21, s15 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_bfe_u32 v89, v15, 16, 1 v_lshrrev_b32_e32 v84, 16, v84 v_add3_u32 v85, v11, v85, 0x7fff v_cmp_o_f32_e64 s15, v10, v10 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s14, vcc_lo, s14 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v86, v12, v86, 0x7fff - .loc 1 246 21 ; hybrid_w4a16.py:246:21 s_clause 0x2 buffer_store_b16 v18, v17, s[20:23], 0 offen buffer_store_b16 v20, v19, s[20:23], 0 offen buffer_store_b16 v9, v21, s[20:23], 0 offen v_add_lshl_u32 v9, v33, v61, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v87, v13, v87, 0x7fff - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v22, 0x80000000, v22, s14 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s14, v11, v11 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v11, v33, v56, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v88, v14, v88, 0x7fff - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v17, v33, v55, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v89, v15, v89, 0x7fff - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s13, vcc_lo, s13 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v18, v33, v53, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v85, 16, v85 v_cndmask_b32_e64 v10, 0x7fff, v84, s15 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s12, vcc_lo, s12 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v86, 16, v86 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v9, 0x80000000, v9, s13 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s13, v12, v12 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s11, vcc_lo, s11 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v87, 16, v87 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v11, 0x80000000, v11, s12 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s12, v13, v13 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s10, vcc_lo, s10 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v88, 16, v88 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v17, 0x80000000, v17, s11 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s11, v14, v14 v_bfe_u32 v90, v16, 16, 1 v_lshrrev_b32_e32 v89, 16, v89 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v19, v33, v52, 1 v_cndmask_b32_e64 v18, 0x80000000, v18, s10 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s10, v15, v15 v_bfe_u32 v91, v1, 16, 1 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 buffer_store_b16 v10, v22, s[20:23], 0 offen - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cndmask_b32_e64 v10, 0x7fff, v85, s14 v_bfe_u32 v92, v2, 16, 1 v_cndmask_b32_e64 v12, 0x7fff, v86, s13 v_cndmask_b32_e64 v13, 0x7fff, v87, s12 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s9, vcc_lo, s9 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_bfe_u32 v93, v3, 16, 1 v_cndmask_b32_e64 v14, 0x7fff, v88, s11 v_add3_u32 v90, v16, v90, 0x7fff v_cndmask_b32_e64 v15, 0x7fff, v89, s10 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v19, 0x80000000, v19, s9 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v91, v1, v91, 0x7fff - .loc 1 246 21 ; hybrid_w4a16.py:246:21 s_clause 0x4 buffer_store_b16 v10, v9, s[20:23], 0 offen buffer_store_b16 v12, v11, s[20:23], 0 offen @@ -857,122 +606,75 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel buffer_store_b16 v14, v18, s[20:23], 0 offen buffer_store_b16 v15, v19, s[20:23], 0 offen v_add_lshl_u32 v9, v33, v49, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v92, v2, v92, 0x7fff - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v11, v33, v42, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v93, v3, v93, 0x7fff - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s8, vcc_lo, s8 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v12, v33, v41, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v90, 16, v90 v_cmp_o_f32_e64 s9, v16, v16 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s7, vcc_lo, s7 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v91, 16, v91 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v9, 0x80000000, v9, s8 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s8, v1, v1 v_lshrrev_b32_e32 v92, 16, v92 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v11, 0x80000000, v11, s7 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s7, v2, v2 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s6, vcc_lo, s6 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v93, 16, v93 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v13, v33, v40, 1 v_cndmask_b32_e64 v12, 0x80000000, v12, s6 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s6, v3, v3 v_cndmask_b32_e64 v10, 0x7fff, v90, s9 v_cndmask_b32_e64 v1, 0x7fff, v91, s8 v_bfe_u32 v94, v4, 16, 1 v_cndmask_b32_e64 v2, 0x7fff, v92, s7 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s5, vcc_lo, s5 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_bfe_u32 v95, v5, 16, 1 v_bfe_u32 v96, v6, 16, 1 v_cndmask_b32_e64 v3, 0x7fff, v93, s6 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v13, 0x80000000, v13, s5 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_bfe_u32 v97, v7, 16, 1 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 s_clause 0x2 buffer_store_b16 v10, v9, s[20:23], 0 offen buffer_store_b16 v1, v11, s[20:23], 0 offen buffer_store_b16 v2, v12, s[20:23], 0 offen v_add_lshl_u32 v1, v33, v38, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_bfe_u32 v98, v8, 16, 1 v_add3_u32 v94, v4, v94, 0x7fff v_add3_u32 v95, v5, v95, 0x7fff - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s4, vcc_lo, s4 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v96, v6, v96, 0x7fff - .loc 1 246 21 ; hybrid_w4a16.py:246:21 buffer_store_b16 v3, v13, s[20:23], 0 offen v_add_lshl_u32 v3, v33, v37, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v97, v7, v97, 0x7fff - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v1, 0x80000000, v1, s4 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s4, v5, v5 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v5, v33, v36, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_add3_u32 v98, v8, v98, 0x7fff - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v9, v33, v35, 1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v94, 16, v94 v_cmp_o_f32_e64 s5, v4, v4 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s3, vcc_lo, s3 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v95, 16, v95 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s2, vcc_lo, s2 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v96, 16, v96 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v3, 0x80000000, v3, s3 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s3, v6, v6 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 s1, vcc_lo, s1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_lshrrev_b32_e32 v97, 16, v97 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e64 v5, 0x80000000, v5, s2 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s2, v7, v7 v_lshrrev_b32_e32 v98, 16, v98 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_add_lshl_u32 v0, v33, v0, 1 v_cndmask_b32_e64 v9, 0x80000000, v9, s1 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cmp_o_f32_e64 s1, v8, v8 v_cndmask_b32_e64 v2, 0x7fff, v94, s5 v_cndmask_b32_e64 v4, 0x7fff, v95, s4 v_cndmask_b32_e64 v6, 0x7fff, v96, s3 - .loc 1 245 38 ; hybrid_w4a16.py:245:38 s_and_b32 vcc_lo, vcc_lo, s0 - .loc 1 243 23 ; hybrid_w4a16.py:243:23 v_cndmask_b32_e64 v7, 0x7fff, v97, s2 v_cndmask_b32_e64 v8, 0x7fff, v98, s1 - .loc 1 246 21 ; hybrid_w4a16.py:246:21 v_cndmask_b32_e32 v0, 0x80000000, v0, vcc_lo s_clause 0x4 buffer_store_b16 v2, v1, s[20:23], 0 offen @@ -980,7 +682,6 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel buffer_store_b16 v6, v5, s[20:23], 0 offen buffer_store_b16 v7, v9, s[20:23], 0 offen buffer_store_b16 v8, v0, s[20:23], 0 offen - .loc 1 246 4 is_stmt 0 ; hybrid_w4a16.py:246:4 s_nop 0 s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) s_endpgm From b0ab3617c0b64adc2731c9da55d8fc4885342c95 Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Thu, 25 Jun 2026 00:38:02 -0600 Subject: [PATCH 14/17] [ROCm] W4A16 skinny GEMM: drop in-loop OOB guards, size descriptors exactly Remove the four in-loop v_cndmask OOB address guards (and their now-dead predicate chain) from the software-pipelined prefetch, and instead bound each read buffer's descriptor to its real size so the hardware clamps any out-of-bounds lane to 0 without issuing a memory access. Why the guards are removable: grouped quantization requires group_size to divide K, and group_size == BLOCK_K == 32, so K is always a multiple of BLOCK_K -- no K-tail can occur. The only OOB accesses are M-tail rows (output rows beyond M, masked out by the existing output-store guards) and the discarded speculative last prefetch. Why num_records must be set: the descriptors previously used the Triton "unbounded" sentinel num_records = 0x7ffffffe (~2GB), which only returns 0 when paired with the cndmask sentinel address 0x80000000. Dropping the cndmasks alone left OOB lanes reading their real (sub-2GB) addresses -- both a latent page-fault risk (it only stayed safe by landing in the allocator's mapped slack) and a measured slowdown, because those spurious reads to adjacent memory inflated the loop's s_waitcnt vmcnt(0) wait. Sizing each descriptor to its exact byte extent fixes both: OOB offsets now exceed num_records, so the hardware returns 0 with no memory access. Changes: - Drop the 4 cndmask guards + dead chain (2 K-bound v_cmp, 3 s_and, dead v125=s5+v37 and v36 increment); load direct from the incrementing bases (v41 advance reordered after the A1 load to keep the pre-inc address). - Compute per-buffer num_records in the prologue from dims already live there (M=s22, N=s23, K=s26, stride_bn=s5, num_groups=s6): activation M*K*2, weight N*stride_bn*4, scale N*num_groups*2. No kernarg ABI change. +2 scratch SGPRs (s37/s38), no occupancy change (still 12 waves/SIMD). - Output-store descriptor and its guards are left untouched. ATT (s_waitcnt vmcnt(0) stall): guarded 6.35M, naive no-guard 9.58M, exact-num_records 7.94M. Wall-clock at M=1322 N=19456 K=2560 g32 bf16 is parity with guarded (4.852 vs 4.854 ms/call, interleaved A/B). Correctness: cos 0.999994 vs Triton reference. Co-authored-by: Claude Signed-off-by: Matthias Gehre --- .../asm/hybrid_w4a16_skinny_gfx1151.amdgcn | 42 +++++++++---------- 1 file changed, 20 insertions(+), 22 deletions(-) diff --git a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn index 957eb75714da..170671592219 100644 --- a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn +++ b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn @@ -158,7 +158,14 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel v_mov_b32_e32 v7, v25 v_mov_b32_e32 v8, v25 s_mov_b32 s15, 0x31027000 - s_mov_b32 s14, 0x7ffffffe + ; --- per-buffer num_records (bytes): hardware clamps OOB loads to 0 (no fault). + ; --- dims still live here: M=s22 N=s23 K=s26 stride_bn=s5 num_groups=s6. + s_mul_i32 s14, s22, s26 ; activation: M*K + s_lshl_b32 s14, s14, 1 ; *2 (bf16) -> s[12:15] dword2 + s_mul_i32 s37, s23, s5 ; weight: N*stride_bn + s_lshl_b32 s37, s37, 2 ; *4 (int32) -> s[16:19] dword2 (temp) + s_mul_i32 s38, s23, s6 ; scale: N*num_groups + s_lshl_b32 s38, s38, 1 ; *2 (bf16) -> s[8:11] dword2 (temp) s_and_b32 s17, s17, 0xffff s_and_b32 s19, s19, 0xffff s_mul_hi_u32 s3, s2, s3 @@ -170,9 +177,9 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel s_mov_b32 s13, s17 s_mov_b32 s16, s18 s_mov_b32 s17, s19 - s_mov_b32 s18, s14 + s_mov_b32 s18, s37 ; weight num_records = N*stride_bn*4 s_mov_b32 s19, s15 - s_mov_b32 s10, s14 + s_mov_b32 s10, s38 ; scale num_records = N*num_groups*2 s_mov_b32 s11, s15 s_mov_b32 s35, 15 ; mask_lo (SGPR; survives the loop, ds_load only writes VGPRs) s_mov_b32 s36, 0xf0000 ; mask_hi @@ -253,44 +260,35 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel DQ_PACK v96, v88, v89 DQ_PACK v97, v90, v91 ; --- prefetch iter k+1 (overlaps ds_store B + ds_load B + WMMA) --- - v_add_nc_u32_e32 v125, s5, v37 - v_cmp_gt_i32_e64 s2, s4, v36 + ; OOB guards removed: group_size(32)==BLOCK_K divides K => no K-tail ever; + ; M-tail + speculative last prefetch are clamped by the buffer descriptor. s_ashr_i32 s3, s5, 31 s_abs_i32 s29, s5 s_xor_b32 s31, s3, s6 - v_cmp_gt_i32_e64 s3, s26, v125 s_mul_hi_u32 s30, s29, s7 - s_and_b32 s2, s0, s2 s_mul_i32 s33, s30, s27 - v_cndmask_b32_e64 v125, 0x80000000, v39, s2 s_sub_i32 s29, s29, s33 - s_and_b32 s2, vcc_lo, s3 s_add_i32 s34, s30, 1 s_sub_i32 s33, s29, s27 - v_cndmask_b32_e64 v126, 0x80000000, v40, s2 - s_and_b32 s2, s1, s3 s_cmp_ge_u32 s29, s27 - v_cndmask_b32_e64 v127, 0x80000000, v41, s2 - buffer_load_b32 v123, v125, s[16:19], 0 offen + buffer_load_b32 v123, v39, s[16:19], 0 offen ; weight (pre-inc base) s_cselect_b32 s2, s34, s30 s_cselect_b32 s3, s33, s29 s_add_i32 s29, s2, 1 s_cmp_ge_u32 s3, s27 - buffer_load_b128 v[115:118], v126, s[12:15], 0 offen + buffer_load_b128 v[115:118], v40, s[12:15], 0 offen ; act A0 (pre-inc base) s_cselect_b32 s2, s29, s2 - v_add_nc_u32_e32 v36, 4, v36 s_xor_b32 s2, s2, s31 v_add_nc_u32_e32 v39, 16, v39 s_sub_i32 s2, s2, s31 v_add_nc_u32_e32 v40, 64, v40 - v_add_lshl_u32 v125, s2, v38, 1 - v_add_nc_u32_e32 v41, 64, v41 + v_add_lshl_u32 v125, s2, v38, 1 ; scale group addr s_add_i32 s28, s28, -1 s_add_i32 s5, s5, 32 s_cmp_lg_u32 s28, 0 - v_cndmask_b32_e64 v125, 0x80000000, v125, s0 - buffer_load_u16 v124, v125, s[8:11], 0 offen - buffer_load_b128 v[119:122], v127, s[12:15], 0 offen + buffer_load_u16 v124, v125, s[8:11], 0 offen ; scale + buffer_load_b128 v[119:122], v41, s[12:15], 0 offen ; act A1 (pre-inc base) + v_add_nc_u32_e32 v41, 64, v41 s_waitcnt lgkmcnt(0) s_barrier ds_store_b128 v42, v[94:97] @@ -707,7 +705,7 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel .amdhsa_system_sgpr_workgroup_info 0 .amdhsa_system_vgpr_workitem_id 0 .amdhsa_next_free_vgpr 128 - .amdhsa_next_free_sgpr 37 + .amdhsa_next_free_sgpr 39 .amdhsa_reserve_vcc 1 .amdhsa_float_round_mode_32 0 .amdhsa_float_round_mode_16_64 0 @@ -943,7 +941,7 @@ amdhsa.kernels: .max_flat_workgroup_size: 256 .name: _triton_w4a16_skinny_fmt_kernel .private_segment_fixed_size: 0 - .sgpr_count: 39 + .sgpr_count: 41 .sgpr_spill_count: 0 .symbol: _triton_w4a16_skinny_fmt_kernel.kd .uniform_work_group_size: 1 From 8c0435b0e29ad74bf36b826637aacc84548177d3 Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Thu, 25 Jun 2026 01:19:23 -0600 Subject: [PATCH 15/17] [ROCm] W4A16 skinny GEMM: fold A1 activation pointer into a scalar offset The two activation b128 loads address the two 64-row halves of the BLOCK_M=128 tile. Their base addresses differ by a loop-invariant constant (v41 = v40 + 128*K, i.e. 64 rows), and both advanced by the same +64/iter, so the second pointer (v41) and its per-iteration increment were redundant. Load A1 by reusing the A0 voffset (v40) with the constant 128*K supplied as the buffer_load scalar-offset operand (precomputed once as s39 = K<<7). This removes the v41 prologue init, both v41 increments, and the now-dead v3 multiply; the loop's lone v40 advance moves after the A1 load so both loads read the pre-increment base. Performance is unchanged (4.813 vs 4.818 ms/call, interleaved A/B; the loop is latency-bound). The benefit is one fewer VALU op per iteration and a simpler address scheme. Correctness: cos 0.999994 vs Triton reference. Co-authored-by: Claude Signed-off-by: Matthias Gehre --- .../asm/hybrid_w4a16_skinny_gfx1151.amdgcn | 17 +++++++---------- 1 file changed, 7 insertions(+), 10 deletions(-) diff --git a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn index 170671592219..0dac6eb1f06b 100644 --- a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn +++ b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn @@ -116,7 +116,6 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel v_lshl_or_b32 v4, v35, 5, v4 v_mul_lo_u32 v7, s5, v7 v_mul_lo_u32 v1, s26, v1 - v_mul_lo_u32 v3, s26, v3 s_mul_f32 s2, s2, 0x4f7ffffe v_cmp_gt_i32_e64 s0, s23, v6 v_mul_lo_u32 v38, s6, v6 @@ -134,7 +133,6 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel s_mul_i32 s3, s3, s2 v_lshl_add_u32 v39, v7, 2, v13 v_lshl_add_u32 v40, v1, 1, v14 - v_lshl_add_u32 v41, v3, 1, v14 v_dual_mov_b32 v19, v25 :: v_dual_add_nc_u32 v42, 0, v2 v_dual_mov_b32 v24, v25 :: v_dual_add_nc_u32 v43, 0, v8 v_dual_mov_b32 v21, v25 :: v_dual_add_nc_u32 v44, 0, v5 @@ -162,6 +160,7 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel ; --- dims still live here: M=s22 N=s23 K=s26 stride_bn=s5 num_groups=s6. s_mul_i32 s14, s22, s26 ; activation: M*K s_lshl_b32 s14, s14, 1 ; *2 (bf16) -> s[12:15] dword2 + s_lshl_b32 s39, s26, 7 ; A1 voffset delta = 128*K (64 rows); share base+soffset s_mul_i32 s37, s23, s5 ; weight: N*stride_bn s_lshl_b32 s37, s37, 2 ; *4 (int32) -> s[16:19] dword2 (temp) s_mul_i32 s38, s23, s6 ; scale: N*num_groups @@ -209,7 +208,7 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel v_cndmask_b32_e64 v126, 0x80000000, v40, s2 s_and_b32 s2, s1, s3 s_cmp_ge_u32 s29, s27 - v_cndmask_b32_e64 v127, 0x80000000, v41, s2 + v_cndmask_b32_e64 v127, 0x80000000, v40, s2 ; A1 mask on shared v40 base buffer_load_b32 v123, v125, s[16:19], 0 offen s_cselect_b32 s2, s34, s30 s_cselect_b32 s3, s33, s29 @@ -223,11 +222,10 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel s_sub_i32 s2, s2, s31 v_add_nc_u32_e32 v40, 64, v40 v_add_lshl_u32 v125, s2, v38, 1 - v_add_nc_u32_e32 v41, 64, v41 s_add_i32 s5, s5, 32 v_cndmask_b32_e64 v125, 0x80000000, v125, s0 buffer_load_u16 v124, v125, s[8:11], 0 offen - buffer_load_b128 v[119:122], v127, s[12:15], 0 offen + buffer_load_b128 v[119:122], v127, s[12:15], s39 offen ; act A1 = v40+128*K .LBB0_2: ; =>This Inner Loop Header: Depth=1 s_waitcnt vmcnt(0) s_barrier @@ -281,14 +279,13 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel s_xor_b32 s2, s2, s31 v_add_nc_u32_e32 v39, 16, v39 s_sub_i32 s2, s2, s31 - v_add_nc_u32_e32 v40, 64, v40 v_add_lshl_u32 v125, s2, v38, 1 ; scale group addr s_add_i32 s28, s28, -1 s_add_i32 s5, s5, 32 s_cmp_lg_u32 s28, 0 buffer_load_u16 v124, v125, s[8:11], 0 offen ; scale - buffer_load_b128 v[119:122], v41, s[12:15], 0 offen ; act A1 (pre-inc base) - v_add_nc_u32_e32 v41, 64, v41 + buffer_load_b128 v[119:122], v40, s[12:15], s39 offen ; act A1 = v40+128*K + v_add_nc_u32_e32 v40, 64, v40 ; advance shared act base (after A0+A1) s_waitcnt lgkmcnt(0) s_barrier ds_store_b128 v42, v[94:97] @@ -705,7 +702,7 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel .amdhsa_system_sgpr_workgroup_info 0 .amdhsa_system_vgpr_workitem_id 0 .amdhsa_next_free_vgpr 128 - .amdhsa_next_free_sgpr 39 + .amdhsa_next_free_sgpr 40 .amdhsa_reserve_vcc 1 .amdhsa_float_round_mode_32 0 .amdhsa_float_round_mode_16_64 0 @@ -941,7 +938,7 @@ amdhsa.kernels: .max_flat_workgroup_size: 256 .name: _triton_w4a16_skinny_fmt_kernel .private_segment_fixed_size: 0 - .sgpr_count: 41 + .sgpr_count: 42 .sgpr_spill_count: 0 .symbol: _triton_w4a16_skinny_fmt_kernel.kd .uniform_work_group_size: 1 From c9760cdf31c0506bea974ad4226941f5215f9e3a Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Thu, 25 Jun 2026 01:36:51 -0600 Subject: [PATCH 16/17] [ROCm] W4A16 skinny GEMM: split WMMA SRC0/SRC1 into distinct VGPR banks All 8 WMMAs had SRC0 (weight frag) and SRC1 (activation frag) starting in the same VGPR bank (reg mod 4 == 3), which costs extra cycles per WMMA on gfx11. The register file is fully packed (128/128 VGPRs, at the 12-wave occupancy boundary), so rather than add registers, swap within the file: - Shift the activation SRC1 fragments down by one (v51-82 -> v50-81) so their first registers land in bank 2, distinct from SRC0 (bank 3) and the SRC2 accumulators (bank 1). - Relocate the freed-up collision: the weight LDS load address moves from v50 to the now-free v82. VGPR count and occupancy are unchanged. Performance is within noise (4.820 vs 4.830 ms/call, interleaved A/B) because the loop is currently global-load-latency-bound, so the WMMA bank savings are mostly hidden behind the s_waitcnt vmcnt(0) wait; the benefit will surface once the prefetch latency is hidden (deeper pipelining). Correctness: cos 0.999994. Co-authored-by: Claude Signed-off-by: Matthias Gehre --- .../asm/hybrid_w4a16_skinny_gfx1151.amdgcn | 38 +++++++++---------- 1 file changed, 19 insertions(+), 19 deletions(-) diff --git a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn index 0dac6eb1f06b..bfb2786c1c7b 100644 --- a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn +++ b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn @@ -141,7 +141,7 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel v_add_nc_u32_e32 v47, 0, v4 v_dual_mov_b32 v9, v25 :: v_dual_add_nc_u32 v48, 0, v10 v_dual_mov_b32 v14, v25 :: v_dual_add_nc_u32 v49, 0, v11 - v_dual_mov_b32 v11, v25 :: v_dual_add_nc_u32 v50, 0, v12 + v_dual_mov_b32 v11, v25 :: v_dual_add_nc_u32 v82, 0, v12 v_mov_b32_e32 v10, v25 v_mov_b32_e32 v12, v25 v_mov_b32_e32 v13, v25 @@ -245,14 +245,14 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel v_lshl_or_b32 v124, v124, 16, v124 ; (scale, scale) s_waitcnt lgkmcnt(0) s_barrier - ds_load_b128 v[51:54], v43 - ds_load_b128 v[59:62], v43 offset:4096 - ds_load_b128 v[55:58], v44 - ds_load_b128 v[63:66], v44 offset:4096 - ds_load_b128 v[67:70], v45 - ds_load_b128 v[75:78], v45 offset:4096 - ds_load_b128 v[71:74], v46 - ds_load_b128 v[79:82], v46 offset:4096 + ds_load_b128 v[50:53], v43 + ds_load_b128 v[58:61], v43 offset:4096 + ds_load_b128 v[54:57], v44 + ds_load_b128 v[62:65], v44 offset:4096 + ds_load_b128 v[66:69], v45 + ds_load_b128 v[74:77], v45 offset:4096 + ds_load_b128 v[70:73], v46 + ds_load_b128 v[78:81], v46 offset:4096 DQ_PACK v94, v84, v85 DQ_PACK v95, v86, v87 DQ_PACK v96, v88, v89 @@ -297,20 +297,20 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel ds_load_b128 v[95:98], v48 offset:2048 ds_load_b128 v[99:102], v49 ds_load_b128 v[107:110], v49 offset:2048 - ds_load_b128 v[103:106], v50 - ds_load_b128 v[111:114], v50 offset:2048 + ds_load_b128 v[103:106], v82 + ds_load_b128 v[111:114], v82 offset:2048 s_waitcnt lgkmcnt(5) - v_wmma_f32_16x16x16_bf16 v[25:32], v[83:90], v[51:58], v[25:32] + v_wmma_f32_16x16x16_bf16 v[25:32], v[83:90], v[50:57], v[25:32] s_waitcnt lgkmcnt(4) - v_wmma_f32_16x16x16_bf16 v[17:24], v[91:98], v[51:58], v[17:24] - v_wmma_f32_16x16x16_bf16 v[9:16], v[83:90], v[59:66], v[9:16] - v_wmma_f32_16x16x16_bf16 v[1:8], v[91:98], v[59:66], v[1:8] + v_wmma_f32_16x16x16_bf16 v[17:24], v[91:98], v[50:57], v[17:24] + v_wmma_f32_16x16x16_bf16 v[9:16], v[83:90], v[58:65], v[9:16] + v_wmma_f32_16x16x16_bf16 v[1:8], v[91:98], v[58:65], v[1:8] s_waitcnt lgkmcnt(1) - v_wmma_f32_16x16x16_bf16 v[25:32], v[99:106], v[67:74], v[25:32] + v_wmma_f32_16x16x16_bf16 v[25:32], v[99:106], v[66:73], v[25:32] s_waitcnt lgkmcnt(0) - v_wmma_f32_16x16x16_bf16 v[17:24], v[107:114], v[67:74], v[17:24] - v_wmma_f32_16x16x16_bf16 v[9:16], v[99:106], v[75:82], v[9:16] - v_wmma_f32_16x16x16_bf16 v[1:8], v[107:114], v[75:82], v[1:8] + v_wmma_f32_16x16x16_bf16 v[17:24], v[107:114], v[66:73], v[17:24] + v_wmma_f32_16x16x16_bf16 v[9:16], v[99:106], v[74:81], v[9:16] + v_wmma_f32_16x16x16_bf16 v[1:8], v[107:114], v[74:81], v[1:8] s_cbranch_scc1 .LBB0_2 ; backedge: branch to next K-tile ; --- end of main loop (fall through to epilogue) --- From bffbf90f8864d004c1291fd241cf3b92aecf7b43 Mon Sep 17 00:00:00 2001 From: Matthias Gehre Date: Thu, 25 Jun 2026 01:43:41 -0600 Subject: [PATCH 17/17] [ROCm] W4A16 skinny GEMM: drive load offsets from a scalar, not VGPR induction The weight and activation prefetch pointers were advanced every iteration with per-lane VGPR adds (v39 += 16, v40 += 64). The increments are uniform across lanes, so they belong in a scalar SOFFSET rather than the VALU. Both byte offsets are exact functions of the existing K-element counter s5 (s5 += 32/iter): weight = s5/2 (4-bit packed), activation = s5*2 (bf16), A1 = activation + 128*K. Compute these three scalars at the top of the prefetch block (while s5 is still the current iteration's value) and pass them as the buffer_load scalar offset; v39/v40 now stay fixed at their iter-0 per-lane base. Changes: - Loop: replace the two v_add_nc_u32 pointer increments with s_lshr/s_lshl/ s_add into s37/s38/s40, used as the weight/A0/A1 soffsets (2 VALU -> 3 SALU, moving induction off the dequant-heavy VALU path). - Prologue: drop the now-dead v36/v39/v40 increments (the iter-0 loads keep their literal soffsets, which equal the s5=0 offsets). - +1 scratch SGPR (s40); occupancy unchanged. Performance within noise (4.837 vs 4.845 ms/call, interleaved A/B) -- the loop is global-load-latency-bound, so freeing VALU slots is groundwork that pays off once the prefetch latency is hidden. Correctness: cos 0.999994. Co-authored-by: Claude Signed-off-by: Matthias Gehre --- .../asm/hybrid_w4a16_skinny_gfx1151.amdgcn | 18 ++++++++---------- 1 file changed, 8 insertions(+), 10 deletions(-) diff --git a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn index bfb2786c1c7b..22d2431461c1 100644 --- a/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn +++ b/vllm/model_executor/kernels/linear/mixed_precision/asm/hybrid_w4a16_skinny_gfx1151.amdgcn @@ -216,11 +216,8 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel s_cmp_ge_u32 s3, s27 buffer_load_b128 v[115:118], v126, s[12:15], 0 offen s_cselect_b32 s2, s29, s2 - v_add_nc_u32_e32 v36, 4, v36 s_xor_b32 s2, s2, s31 - v_add_nc_u32_e32 v39, 16, v39 s_sub_i32 s2, s2, s31 - v_add_nc_u32_e32 v40, 64, v40 v_add_lshl_u32 v125, s2, v38, 1 s_add_i32 s5, s5, 32 v_cndmask_b32_e64 v125, 0x80000000, v125, s0 @@ -260,6 +257,9 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel ; --- prefetch iter k+1 (overlaps ds_store B + ds_load B + WMMA) --- ; OOB guards removed: group_size(32)==BLOCK_K divides K => no K-tail ever; ; M-tail + speculative last prefetch are clamped by the buffer descriptor. + s_lshr_b32 s37, s5, 1 ; weight K byte offset = s5/2 (4-bit packed) + s_lshl_b32 s38, s5, 1 ; act A0 K byte offset = s5*2 (bf16) + s_add_i32 s40, s38, s39 ; act A1 K byte offset = A0 + 128*K s_ashr_i32 s3, s5, 31 s_abs_i32 s29, s5 s_xor_b32 s31, s3, s6 @@ -269,23 +269,21 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel s_add_i32 s34, s30, 1 s_sub_i32 s33, s29, s27 s_cmp_ge_u32 s29, s27 - buffer_load_b32 v123, v39, s[16:19], 0 offen ; weight (pre-inc base) + buffer_load_b32 v123, v39, s[16:19], s37 offen ; weight (base + s5/2) s_cselect_b32 s2, s34, s30 s_cselect_b32 s3, s33, s29 s_add_i32 s29, s2, 1 s_cmp_ge_u32 s3, s27 - buffer_load_b128 v[115:118], v40, s[12:15], 0 offen ; act A0 (pre-inc base) + buffer_load_b128 v[115:118], v40, s[12:15], s38 offen ; act A0 (base + s5*2) s_cselect_b32 s2, s29, s2 s_xor_b32 s2, s2, s31 - v_add_nc_u32_e32 v39, 16, v39 s_sub_i32 s2, s2, s31 v_add_lshl_u32 v125, s2, v38, 1 ; scale group addr s_add_i32 s28, s28, -1 s_add_i32 s5, s5, 32 s_cmp_lg_u32 s28, 0 buffer_load_u16 v124, v125, s[8:11], 0 offen ; scale - buffer_load_b128 v[119:122], v40, s[12:15], s39 offen ; act A1 = v40+128*K - v_add_nc_u32_e32 v40, 64, v40 ; advance shared act base (after A0+A1) + buffer_load_b128 v[119:122], v40, s[12:15], s40 offen ; act A1 (base + s5*2 + 128*K) s_waitcnt lgkmcnt(0) s_barrier ds_store_b128 v42, v[94:97] @@ -702,7 +700,7 @@ _triton_w4a16_skinny_fmt_kernel: ; @_triton_w4a16_skinny_fmt_kernel .amdhsa_system_sgpr_workgroup_info 0 .amdhsa_system_vgpr_workitem_id 0 .amdhsa_next_free_vgpr 128 - .amdhsa_next_free_sgpr 40 + .amdhsa_next_free_sgpr 41 .amdhsa_reserve_vcc 1 .amdhsa_float_round_mode_32 0 .amdhsa_float_round_mode_16_64 0 @@ -938,7 +936,7 @@ amdhsa.kernels: .max_flat_workgroup_size: 256 .name: _triton_w4a16_skinny_fmt_kernel .private_segment_fixed_size: 0 - .sgpr_count: 42 + .sgpr_count: 43 .sgpr_spill_count: 0 .symbol: _triton_w4a16_skinny_fmt_kernel.kd .uniform_work_group_size: 1