diff --git a/CMakeLists.txt b/CMakeLists.txt index 25010ebbcda1..d6843d5c4b8d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1591,6 +1591,17 @@ if(VLLM_GPU_LANG STREQUAL "HIP") "csrc/rocm/gdn_chunked.cu") endif() + # RDNA3.5 decode attention: the variants of csrc/rocm/rdna35_attn/variants.def. + # Their code is emitted for gfx1151 only, where they were tuned; other shapes + # and devices are served by Triton. + set(VLLM_ROCM_HAS_RDNA35_ATTN OFF) + if(VLLM_GPU_ARCHES MATCHES "gfx1151") + set(VLLM_ROCM_HAS_RDNA35_ATTN ON) + include(cmake/rdna35_attn.cmake) + rdna35_attn_sources(VLLM_RDNA35_ATTN_SRC) + list(APPEND VLLM_ROCM_EXT_SRC ${VLLM_RDNA35_ATTN_SRC}) + endif() + define_extension_target( _rocm_C DESTINATION vllm @@ -1639,6 +1650,9 @@ if(VLLM_GPU_LANG STREQUAL "HIP") if(VLLM_ROCM_HAS_GDN_CHUNKED) target_compile_definitions(_rocm_C PRIVATE VLLM_ROCM_GDN_CHUNKED) endif() + if(VLLM_ROCM_HAS_RDNA35_ATTN) + target_compile_definitions(_rocm_C PRIVATE VLLM_ROCM_RDNA35_ATTN) + endif() endif() # Must run after the last HIP `define_extension_target` so every extension diff --git a/benchmarks/attention_benchmarks/README.md b/benchmarks/attention_benchmarks/README.md index 944ceb91af91..867dc68b1fb5 100644 --- a/benchmarks/attention_benchmarks/README.md +++ b/benchmarks/attention_benchmarks/README.md @@ -21,6 +21,41 @@ python benchmark.py \ --output-csv results.csv ``` +## Cache residency + +`benchmark_fn` walks one KV cache per layer and divides the result by the layer +count, so the layer loop doubles as a rotation over disjoint KV regions. Whether +a measurement reads DRAM or a resident last-level cache therefore depends on +`num_layers * kv_bytes_per_layer`, not on one layer's cache. + +The default of 10 layers is enough for most shapes. It stops being enough when +the product falls below the cache, which on a Strix Halo means a 32 MiB MALL +that CUDA-graph replay never flushes. Measured on that board, TRITON_ATTN +decode, with and without the option: + +| working set (10 layers) | cells | median bias | worst | +| --- | --- | --- | --- | +| under 32 MiB | 6 | +35% | +64% | +| over 32 MiB | 18 | +0% | +3% | + +The split is clean: every cell that fits the MALL was reporting a time too low +to be sustainable, and every cell that does not was already fine. Small models +at short contexts are the ones affected — `32x8x128` at 512 tokens reads 20 MB +across ten layers and came out 64% fast. + +`--min-working-set-mb` raises `--num-layers` until the KV working set reaches +the given size. It only ever raises, so shapes already past the cache are +untouched and their timings do not move. + +```bash +# 20 MB across the default 10 layers: inside the MALL. +python benchmark.py --backends TRITON_ATTN --batch-specs q1s512 \ + --num-q-heads 32 --num-kv-heads 8 --head-dim 128 # 11 us, cached +python benchmark.py --backends TRITON_ATTN --batch-specs q1s512 \ + --num-q-heads 32 --num-kv-heads 8 --head-dim 128 \ + --min-working-set-mb 96 # 18 us, DRAM +``` + ## Simplified Batch Specification Grammar Express workloads concisely using query length and sequence length: diff --git a/benchmarks/attention_benchmarks/benchmark.py b/benchmarks/attention_benchmarks/benchmark.py index a21999217e39..43fa9b3b68e1 100644 --- a/benchmarks/attention_benchmarks/benchmark.py +++ b/benchmarks/attention_benchmarks/benchmark.py @@ -35,6 +35,7 @@ from dataclasses import replace from pathlib import Path +import torch import yaml from rich.console import Console from tqdm import tqdm @@ -688,6 +689,25 @@ def main(): # Model config parser.add_argument("--num-layers", type=int, default=10, help="Number of layers") + parser.add_argument( + "--dtype", + choices=["float16", "bfloat16"], + default="float16", + help="Q/K/V and output dtype.", + ) + parser.add_argument( + "--min-working-set-mb", + type=float, + default=None, + help=( + "Grow --num-layers until the KV working set reaches this many MB. " + "The layer loop walks one KV cache per layer, so it doubles as a " + "rotation over disjoint KV; without this a small model at a short " + "context can stay resident in a large last-level cache and report " + "bandwidth the memory system cannot sustain. Suggested value on " + "Strix Halo (32 MiB MALL): 96." + ), + ) parser.add_argument("--head-dim", type=int, default=128, help="Head dimension") parser.add_argument("--num-q-heads", type=int, default=32, help="Query heads") parser.add_argument("--num-kv-heads", type=int, default=8, help="KV heads") @@ -1087,6 +1107,8 @@ def main(): backend=decode_backend, batch_spec=spec, num_layers=args.num_layers, + dtype=getattr(torch, args.dtype), + min_working_set_mb=args.min_working_set_mb, head_dim=args.head_dim, num_q_heads=args.num_q_heads, num_kv_heads=args.num_kv_heads, @@ -1157,6 +1179,8 @@ def main(): backend=backend, # Will be overridden later batch_spec=batch_spec, num_layers=args.num_layers, + dtype=getattr(torch, args.dtype), + min_working_set_mb=args.min_working_set_mb, head_dim=args.head_dim, num_q_heads=args.num_q_heads, num_kv_heads=args.num_kv_heads, @@ -1426,6 +1450,8 @@ def main(): backend=f"{backend}_{variant_label}", batch_spec=spec, num_layers=args.num_layers, + dtype=getattr(torch, args.dtype), + min_working_set_mb=args.min_working_set_mb, head_dim=args.head_dim, num_q_heads=args.num_q_heads, num_kv_heads=args.num_kv_heads, @@ -1500,6 +1526,8 @@ def main(): # Model parameter sweep base_config_args = { "num_layers": args.num_layers, + "dtype": getattr(torch, args.dtype), + "min_working_set_mb": args.min_working_set_mb, "head_dim": args.head_dim, "v_head_dim": args.v_head_dim, "num_q_heads": args.num_q_heads, @@ -1532,6 +1560,8 @@ def main(): # Unified parameter sweep base_config_args = { "num_layers": args.num_layers, + "dtype": getattr(torch, args.dtype), + "min_working_set_mb": args.min_working_set_mb, "head_dim": args.head_dim, "v_head_dim": args.v_head_dim, "num_q_heads": args.num_q_heads, @@ -1573,6 +1603,8 @@ def main(): backend=backend, batch_spec=spec, num_layers=args.num_layers, + dtype=getattr(torch, args.dtype), + min_working_set_mb=args.min_working_set_mb, head_dim=args.head_dim, v_head_dim=getattr(args, "v_head_dim", None), num_q_heads=args.num_q_heads, @@ -1649,6 +1681,8 @@ def main(): backend=decode_backend, batch_spec=spec, num_layers=args.num_layers, + dtype=getattr(torch, args.dtype), + min_working_set_mb=args.min_working_set_mb, head_dim=args.head_dim, num_q_heads=args.num_q_heads, num_kv_heads=args.num_kv_heads, diff --git a/benchmarks/attention_benchmarks/common.py b/benchmarks/attention_benchmarks/common.py index 0bf8a46184f1..ae64df4bc169 100644 --- a/benchmarks/attention_benchmarks/common.py +++ b/benchmarks/attention_benchmarks/common.py @@ -38,6 +38,30 @@ def batch_spec_sort_key(spec: str) -> tuple[int, int, int]: return (0, 0, 0) +def layers_for_working_set( + num_layers: int, + kv_bytes_per_layer: int, + min_working_set_mb: float | None, +) -> int: + """Raise the layer count until the KV working set outgrows the cache. + + ``benchmark_fn`` walks one KV cache per layer and the result is divided by + the layer count, so the layer loop doubles as a rotation over disjoint KV + regions. Whether a measurement reads DRAM or a last-level cache therefore + depends on ``num_layers * kv_bytes_per_layer``, not on the size of a single + layer's cache. + + Left at the default this is a no-op. Set it on parts with a large + last-level cache (for example the 32 MiB MALL on Strix Halo), where a small + model at a short context otherwise stays resident across replays and + reports bandwidth the memory system cannot sustain. + """ + if not min_working_set_mb or kv_bytes_per_layer <= 0: + return num_layers + needed = int(min_working_set_mb * 1024**2) + return max(num_layers, -(-needed // kv_bytes_per_layer)) + + def run_do_bench( benchmark_fn, use_cuda_graphs: bool, @@ -284,10 +308,17 @@ class BenchmarkConfig: torch_profile_dir: str | None = None torch_profile_iters: int = 3 warmup_ms: int | None = None + # Grow num_layers until the KV working set reaches this, so the timed + # kernel reads DRAM instead of a resident last-level cache. + min_working_set_mb: float | None = None # "auto" or "fp8" kv_cache_dtype: str = "auto" + # Sliding window in keys (None: full attention). Set on the impl and the + # KV cache spec, so every backend sees the same windowed layer. + sliding_window: int | None = None + # MLA-specific prefill_backend: str | None = None kv_lora_rank: int | None = None diff --git a/benchmarks/attention_benchmarks/runner.py b/benchmarks/attention_benchmarks/runner.py index c51e62727cfb..73d949c58cf9 100644 --- a/benchmarks/attention_benchmarks/runner.py +++ b/benchmarks/attention_benchmarks/runner.py @@ -12,6 +12,7 @@ import statistics import types from contextlib import contextmanager +from dataclasses import replace import torch from batch_spec import parse_batch_spec, reorder_for_flashinfer @@ -20,6 +21,7 @@ BenchmarkResult, MockLayer, get_attention_scale, + layers_for_working_set, run_do_bench, run_ncu_profile, ) @@ -89,11 +91,41 @@ def log_warnings_and_errors_only(): # ============================================================================ +def _block_rows( + q_lens: list[int], + kv_lens: list[int], + block_size: int, + sliding_window: int | None, +) -> tuple[list[list[int]], int]: + """Block table rows and the number of KV blocks to allocate. + + A windowed layer, as vLLM's sliding-window manager keeps it, holds only + the blocks some query's window reaches; the earlier entries name the null + block 0. Allocating the whole sequence instead would spread the few + blocks read over a buffer many times larger than what is touched. + """ + max_blocks = (max(kv_lens) + block_size - 1) // block_size + if sliding_window is None: + rows = [ + list(range(r * max_blocks, (r + 1) * max_blocks)) + for r in range(len(kv_lens)) + ] + return rows, len(kv_lens) * max_blocks + rows, nxt = [], 1 + for q, kv in zip(q_lens, kv_lens): + first = max(0, kv - q - (sliding_window - 1)) // block_size + live = max_blocks - first + rows.append([0] * first + list(range(nxt, nxt + live))) + nxt += live + return rows, nxt + + def _build_common_attn_metadata( q_lens: list[int], kv_lens: list[int], block_size: int, device: torch.device, + sliding_window: int | None = None, ) -> CommonAttentionMetadata: """Build CommonAttentionMetadata from query/kv lengths.""" batch_size = len(q_lens) @@ -108,11 +140,8 @@ def _build_common_attn_metadata( seq_lens = torch.tensor(kv_lens, dtype=torch.int32, device=device) max_seq_len = int(seq_lens.max().item()) - max_blocks = (max(kv_lens) + block_size - 1) // block_size - num_blocks = batch_size * max_blocks - block_table_tensor = torch.arange( - num_blocks, dtype=torch.int32, device=device - ).view(batch_size, max_blocks) + rows, _ = _block_rows(q_lens, kv_lens, block_size, sliding_window) + block_table_tensor = torch.tensor(rows, dtype=torch.int32, device=device) slot_mapping = torch.arange(total_tokens, dtype=torch.int64, device=device) max_query_len = max(q_lens) @@ -131,17 +160,38 @@ def _build_common_attn_metadata( ) +class _DropCastWarning(logging.Filter): + """Drop vLLM's bfloat16 -> float16 cast warning. + + The stand-in model's config is bfloat16 and this harness asks for float16 + on purpose (see the dtype argument below), so the cast is expected and the + warning fires once per ModelConfig -- once per benchmarked cell, thousands + of lines in a sweep. Scoped to that one message rather than to the logger + or the level, because the warning that must stay visible is the backend + announcing it fell back to Triton. + """ + + def filter(self, record: logging.LogRecord) -> bool: + return not record.getMessage().startswith("Casting ") + + def _create_vllm_config( config: BenchmarkConfig, max_num_blocks: int, ) -> VllmConfig: """Create a VllmConfig for benchmarking with mock model methods.""" + cast_logger = logging.getLogger("vllm.config.model") + if not any(isinstance(f, _DropCastWarning) for f in cast_logger.filters): + cast_logger.addFilter(_DropCastWarning()) model_config = ModelConfig( model="HuggingFaceTB/SmolLM2-135M", # Use public model to avoid login issues tokenizer="HuggingFaceTB/SmolLM2-135M", skip_tokenizer_init=True, trust_remote_code=False, - dtype="auto", # Use model's native dtype + # Honour the configured dtype. With "auto" this silently became the + # stand-in model's native bfloat16 while BenchmarkConfig.dtype claimed + # float16, so every result was labelled with a dtype it did not use. + dtype=str(config.dtype).removeprefix("torch."), seed=0, max_model_len=1024, ) @@ -229,7 +279,7 @@ def _create_backend_impl( scale=scale, num_kv_heads=config.num_kv_heads, alibi_slopes=None, - sliding_window=None, + sliding_window=config.sliding_window, kv_cache_dtype=config.kv_cache_dtype, ) @@ -238,6 +288,7 @@ def _create_backend_impl( num_kv_heads=config.num_kv_heads, head_size=config.head_dim, dtype=dtype, + sliding_window=config.sliding_window, ) layer = MockLayer(device, kv_cache_spec=kv_cache_spec) @@ -335,6 +386,11 @@ def _create_input_tensors( return q_list, k_list, v_list +def dtype_size(config) -> int: + """Bytes per KV element for the configured cache dtype.""" + return 1 if config.kv_cache_dtype.startswith("fp8") else config.dtype.itemsize + + def _create_kv_cache( config: BenchmarkConfig, max_num_blocks: int, @@ -481,12 +537,29 @@ def run_attention_benchmark(config: BenchmarkConfig) -> BenchmarkResult: q_lens = [r.q_len for r in requests] kv_lens = [r.kv_len for r in requests] total_q = sum(q_lens) - max_kv = max(kv_lens) - batch_size = len(q_lens) - # Calculate total blocks needed: batch_size * max_blocks_per_request - max_blocks_per_request = (max_kv + config.block_size - 1) // config.block_size - max_num_blocks = batch_size * max_blocks_per_request + # KV blocks to allocate: every request's whole sequence, or its window. + _, max_num_blocks = _block_rows( + q_lens, kv_lens, config.block_size, config.sliding_window + ) + + # One KV cache per layer, so the layer loop is also a rotation over + # disjoint KV; grow it when asked, to push the working set out of cache. + # A windowed layer reads only its window, so that is what must outgrow the + # cache; counting the whole sequence would leave the read part resident. + read_lens = [ + kv if config.sliding_window is None else min(kv, config.sliding_window + q - 1) + for kv, q in zip(kv_lens, q_lens) + ] + kv_bytes_per_layer = ( + 2 * sum(read_lens) * config.num_kv_heads * config.head_dim * dtype_size(config) + ) + config = replace( + config, + num_layers=layers_for_working_set( + config.num_layers, kv_bytes_per_layer, config.min_working_set_mb + ), + ) # Suppress vLLM logs during setup to reduce spam with log_warnings_and_errors_only(): @@ -509,7 +582,7 @@ def run_attention_benchmark(config: BenchmarkConfig) -> BenchmarkResult: get_kv_cache_layout.cache_clear() common_metadata = _build_common_attn_metadata( - q_lens, kv_lens, config.block_size, device + q_lens, kv_lens, config.block_size, device, config.sliding_window ) kv_cache_spec = FullAttentionSpec( @@ -517,6 +590,7 @@ def run_attention_benchmark(config: BenchmarkConfig) -> BenchmarkResult: num_kv_heads=config.num_kv_heads, head_size=config.head_dim, dtype=dtype, + sliding_window=config.sliding_window, ) builder = _create_metadata_builder( diff --git a/benchmarks/kernels/gfx1151_decode_attn/tools/batch.py b/benchmarks/kernels/gfx1151_decode_attn/tools/batch.py new file mode 100644 index 000000000000..c04cd715e1dc --- /dev/null +++ b/benchmarks/kernels/gfx1151_decode_attn/tools/batch.py @@ -0,0 +1,149 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""ROCM_ATTN (gfx1151) against TRITON_ATTN on batches: many decodes, and mixed. + +`matrix.py` measures one sequence. This measures what a server runs: batches +of sequences decoding together (one launch, grid.y per sequence) and batches +mixing decodes with a prefill or a chunked-prefill extend (decodes on the +kernel, the rest on Triton). + +Decode batches are timed under CUDA graphs, as vLLM replays them. Mixed +batches are timed eagerly: vLLM does not capture them whole, and the backend +only splits outside a capture. Decodes are listed first in every mixed spec, +the order vLLM's batch reordering gives the backend. + +Each cell names the path the backend took: `kernel` (every call on the kernel), +`split` (decodes on the kernel, the rest on Triton) or `triton` (fallback). + + cd && VLLM_KV_CACHE_LAYOUT=HND amd-gpu-lock \\ + /bin/python benchmarks/kernels/gfx1151_decode_attn/tools/batch.py +""" + +import argparse +import pathlib +import statistics +import sys +from typing import Any + +sys.path.insert(0, str(pathlib.Path(__file__).resolve().parent)) +import shapeset # noqa: E402 + +DECODE = ["2q1s4k", "8q1s4k", "8q4s4k", "32q1s1k", "64q1s1k", "32q4s1k"] +MIXED = ["4q1s8k_q512", "16q1s4k_2q1k", "16q1s2k_q64s2k", "32q1s1k_q64s1k"] + + +def roofline_us(spec: str, hq: int, hkv: int, d: int, window: int, itemsize: int): + """One dispatch plus every decode sequence's bytes; None if any request + is a prefill (its floor is compute, not bytes).""" + from batch_spec import parse_batch_spec + + reqs = parse_batch_spec(spec) + if any(r.q_len > 8 for r in reqs): + return None + moved = sum( + shapeset.roofline_us( + hq, hkv, d, r.q_len, r.kv_len, itemsize=itemsize, window=window + ) + - shapeset.DISPATCH_US + for r in reqs + ) + return shapeset.DISPATCH_US + moved + + +def main() -> None: + p = argparse.ArgumentParser() + shapeset.add_arguments(p) + shapeset.add_dtype_argument(p) + p.add_argument("--decode", nargs="*", default=DECODE) + p.add_argument("--mixed", nargs="*", default=MIXED) + p.add_argument("--block-size", type=int, default=16) + p.add_argument("--reps", type=int, default=1) + args = p.parse_args() + dtype = shapeset.torch_dtype(args.dtype) + itemsize = 2 + + bench = pathlib.Path(__file__).resolve().parents[3] / "attention_benchmarks" + sys.path.insert(0, str(bench)) + from common import BenchmarkConfig + from runner import run_attention_benchmark + + import vllm.v1.attention.backends.rocm_attn as backend_mod + + shapes, _ = shapeset.load(args) + groups: dict[tuple[int, int, int, int], list[str]] = {} + for sh in shapes: + if sh.d == 96: # three elements per lane: never served + continue + groups.setdefault((sh.hq, sh.hkv, sh.d, sh.window), []).append(sh.model) + + def run(backend, spec, hq, hkv, d, window, graphs): + cfg = BenchmarkConfig( + backend=backend, + batch_spec=spec, + num_layers=10, + min_working_set_mb=96, + head_dim=d, + num_q_heads=hq, + num_kv_heads=hkv, + block_size=args.block_size, + device="cuda:0", + dtype=dtype, + sliding_window=window or None, + use_cuda_graphs=graphs, + ) + rs = [run_attention_benchmark(cfg) for _ in range(args.reps)] + return statistics.median(r.median_time for r in rs) * 1e6 + + def ours(spec, hq, hkv, d, window, graphs): + impls: list[Any] = [] + orig = backend_mod.RocmAttentionRdna35Impl.__init__ + + def spy(self, *a, _init=orig, _seen=impls, **kw): + _init(self, *a, **kw) + _seen.append(self) + + backend_mod.RocmAttentionRdna35Impl.__init__ = spy + try: + # First call builds whatever variant the batch needs; a build + # inside the timed run would corrupt its median. + run("ROCM_ATTN", spec, hq, hkv, d, window, graphs) + impls.clear() + us = run("ROCM_ATTN", spec, hq, hkv, d, window, graphs) + finally: + backend_mod.RocmAttentionRdna35Impl.__init__ = orig + k = sum(i.kernel_calls for i in impls) + f = sum(i.fallback_calls for i in impls) + split = sum(getattr(i, "split_calls", 0) for i in impls) + path = "split" if split else ("kernel" if k and not f else "triton") + return us, path + + print( + f"" + ) + print( + "| models | Hq | Hkv | D | window | batch | graphs | path | roofline | " + "Triton | ours | vs | %roof |" + ) + print("| --- " * 13 + "|") + for (hq, hkv, d, window), models in sorted(groups.items()): + label = models[0] if len(models) == 1 else f"{models[0]} +{len(models) - 1}" + for spec, graphs in [(s, True) for s in args.decode] + [ + (s, False) for s in args.mixed + ]: + roof = roofline_us(spec, hq, hkv, d, window, itemsize) + us, path = ours(spec, hq, hkv, d, window, graphs) + tri = run("TRITON_ATTN", spec, hq, hkv, d, window, graphs) + pct = "-" if roof is None else f"{roof / us * 100:.1f} %" + rf = "-" if roof is None else f"{roof:.1f}" + print( + f"| {label} | {hq} | {hkv} | {d} | {window or '-'} | {spec} | " + f"{'yes' if graphs else 'no'} | {path} | {rf} | {tri:.1f} | " + f"{us:.1f} | {tri / us:.2f}x | {pct} |", + flush=True, + ) + + +if __name__ == "__main__": + main() diff --git a/benchmarks/kernels/gfx1151_decode_attn/tools/check.py b/benchmarks/kernels/gfx1151_decode_attn/tools/check.py new file mode 100644 index 000000000000..4117bf4e834f --- /dev/null +++ b/benchmarks/kernels/gfx1151_decode_attn/tools/check.py @@ -0,0 +1,170 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Validate a kernel variant against a float reference. + +Run this before trusting any timing: several bugs this session were silent, +producing finite but wrong numbers (an empty workgroup whose gmax was -inf, a +wave reaching the LDS stores with exec = 0, bf16 read as fp16). + + cd && PYTHONPATH=$PWD amd-gpu-lock /bin/python \\ + benchmarks/kernels/gfx1151_decode_attn/tools/check.py --hq 8 --hkv 4 + +S=48 matters: a causal off-by-one moves the softmax by ~1/S while the tolerance +is fixed, so long contexts hide exactly the bug most likely to be present. +""" + +import argparse +import itertools +import sys + +import shapeset +import torch + + +def reference(q, kv, s, hq, hkv, head_dim, m): + flat = kv.transpose(1, 2).flatten(0, 1)[:s] + k, v = flat[..., :head_dim], flat[..., head_dim:] + gqa = hq // hkv + qf = q.float().permute(1, 0, 2) + kf = k.float().permute(1, 0, 2).repeat_interleave(gqa, 0) + vf = v.float().permute(1, 0, 2).repeat_interleave(gqa, 0) + scores = torch.bmm(qf, kf.transpose(1, 2)) * (head_dim**-0.5) + pos = torch.arange(s, device=q.device).view(1, s) + lim = (s - m + torch.arange(m, device=q.device)).view(m, 1) + scores = scores.masked_fill((pos > lim).view(1, m, s), float("-inf")) + return torch.bmm(torch.softmax(scores, -1), vf).permute(1, 0, 2) + + +def main() -> None: + p = argparse.ArgumentParser() + p.add_argument("--hq", type=int, default=32) + p.add_argument("--hkv", type=int, default=16) + p.add_argument("--head-dim", type=int, default=256) + p.add_argument("--m", type=int, default=4) + p.add_argument("--block-size", type=int, default=16) + p.add_argument("--contexts", type=int, nargs="+", default=[48, 1024]) + p.add_argument("--layouts", type=int, nargs="+", default=[0, 1]) + p.add_argument("--nseg", type=int, nargs="+", default=[None]) + p.add_argument("--rg", type=int, nargs="+", default=[None]) + p.add_argument("--minb", type=int, nargs="+", default=[None]) + p.add_argument("--nw", type=int, default=None) + p.add_argument("--dspl", type=int, default=None) + p.add_argument("--rspl", type=int, default=None) + p.add_argument("--ablate", type=int, default=None) + p.add_argument( + "--repeat", + type=int, + default=1, + help="invoke N times and check every result; catches state a kernel " + "leaves behind between launches, such as an arrival counter that is " + "not reset", + ) + p.add_argument( + "--mutate", + type=int, + default=0, + help="1 admits one key too many; the check MUST then fail", + ) + shapeset.add_dtype_argument(p) + args = p.parse_args() + dtype = shapeset.torch_dtype(args.dtype) + rtol = shapeset.RTOL[args.dtype] + + from vllm.v1.attention.ops.rdna35_hip_decode import ( + KernelVariant, + load, + make_scratch, + precompile, + ) + + dev = torch.device("cuda") + failures = 0 + + # Build every variant the matrix needs before checking any of them. The + # variant does not depend on S, so the product below is usually a handful + # of builds at ~19.6 s each -- serially that is minutes of nothing before + # the first result. Failures are left for the loop, which reports them + # against the shape that asked for them. + def variant_for(layout, nseg, rg, minb): + from vllm.v1.attention.ops.rdna35_hip_decode import knobs_for as _knobs_for + + kw = dict(_knobs_for(args.hq, args.hkv, args.head_dim, args.m, 0, dtype)) + kw["mutate"] = args.mutate + kw["dtype"] = dtype + for name, val in ( + ("nseg", nseg), + ("rg", rg), + ("minb", minb), + ("nw", args.nw), + ("dspl", args.dspl), + ("rspl", args.rspl), + ("ablate", args.ablate), + ): + if val is not None: + kw[name] = val + return KernelVariant( + args.head_dim, args.hq, args.hkv, args.m, args.block_size, layout, **kw + ) + + precompile( + [ + variant_for(*combo) + for combo in itertools.product(args.layouts, args.nseg, args.rg, args.minb) + ] + ) + for s, layout, nseg, rg, minb in itertools.product( + args.contexts, args.layouts, args.nseg, args.rg, args.minb + ): + torch.manual_seed(0) + # Round up, so S need not be a multiple of the page: a tile that runs + # off the end of the sequence is exactly the case where a kernel may + # read uninitialised slots, and an exact-division harness never builds + # one. The trailing slots are filled with NaN below for that reason. + blocks = -(-s // args.block_size) + shape = ( + (blocks, args.block_size, args.hkv, 2 * args.head_dim) + if layout == 0 + else (blocks, args.hkv, args.block_size, 2 * args.head_dim) + ) + kv = torch.randn(shape, device=dev, dtype=dtype) * 0.5 + if layout == 0: + kv = kv.transpose(1, 2) + # Slots past the sequence carry NaN, not plausible data. A kernel that + # addresses them is fine; one that lets them reach the accumulator is + # not, and 0 * NaN = NaN makes that failure total rather than subtle. + tail = blocks * args.block_size - s + if tail: + kv[-1, :, args.block_size - tail :, :] = float("nan") + q = torch.randn(args.m, args.hq, args.head_dim, device=dev, dtype=dtype) * 0.5 + bt = torch.arange(blocks, device=dev, dtype=torch.int32) + seq_lens = torch.tensor([s], device=dev, dtype=torch.int32) + + variant = variant_for(layout, nseg, rg, minb) + module = load(variant) + acc, smax, ssum, arrivals = make_scratch(variant, dev) + out = torch.empty_like(q) + ref = reference(q, kv, s, args.hq, args.hkv, args.head_dim, args.m) + max_rel = 0.0 + ok = True + for _ in range(args.repeat): + out.zero_() + module.decode_attn( + q, kv, bt, out, acc, smax, ssum, arrivals, seq_lens, args.head_dim**-0.5 + ) + torch.accelerator.synchronize() + got = out.float() + rel = ((got - ref).abs() / ref.abs().clamp_min(1e-3)).max().item() + max_rel = max(max_rel, rel) + ok = ok and rel <= rtol and torch.isfinite(got).all() + if args.mutate: + ok = not ok # the negative control must be detected + failures += not ok + label = f"S={s} l={layout} nseg={nseg} rg={rg} minb={minb}" + print(f"{label:<44} max_rel={max_rel:.3e} {'PASS' if ok else 'FAIL'}") + + sys.exit(1 if failures else 0) + + +if __name__ == "__main__": + main() diff --git a/benchmarks/kernels/gfx1151_decode_attn/tools/dump_asm.sh b/benchmarks/kernels/gfx1151_decode_attn/tools/dump_asm.sh new file mode 100755 index 000000000000..0cd8ff74514f --- /dev/null +++ b/benchmarks/kernels/gfx1151_decode_attn/tools/dump_asm.sh @@ -0,0 +1,30 @@ +#!/usr/bin/env bash +# Emit gfx1151 ISA for one kernel variant. +# +# tools/dump_asm.sh [out.s] [-DNSEG=4 -DNUM_Q_HEADS=8 -DKV_BF16=1 ...] +# +# Built without RDNA35_TORCH_EXT the source instantiates the kernel itself, so +# it compiles as is. Two findings this session came from reading the output: +# four redundant block-table loads per tile, and eleven instructions spent on +# one IEEE division. +set -euo pipefail + +VENV=${VENV:-/scratch/rogarcia/vllm/.venv} +HIPCC=${HIPCC:-$VENV/lib/python3.12/site-packages/_rocm_sdk_devel/bin/hipcc} +HERE="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +CSRC="$(cd "$HERE/../../../../csrc/rocm" && pwd)" +SRC=${SRC:-rdna35_decode_attn.cu} +OUT=${1:-/tmp/kernel.s}; shift || true + +"$HIPCC" -O3 --offload-arch=gfx1151 --cuda-device-only -S -o "$OUT" \ + -DHEAD_DIM=256 -DNUM_Q_HEADS=32 -DNUM_KV_HEADS=16 -DMAXM=4 \ + -DBS=16 -DLAYOUT=1 -DNSEG=1 "$@" "$CSRC/$SRC" 2> >(grep -v hip-link >&2) + +echo "$OUT" +printf 'instructions %s\n' "$(grep -cE '^\s+[a-z]' "$OUT")" +for op in global_load_b128 global_load_b64 v_wmma s_waitcnt v_dual \ + v_perm_b32 s_barrier; do + printf '%-18s %s\n' "$op" "$(grep -c "$op" "$OUT" || true)" +done +printf '%-18s %s\n' vgpr_count "$(grep -m1 '\.vgpr_count' "$OUT" | tr -dc 0-9)" +printf '%-18s %s\n' vgpr_spill "$(grep -m1 '\.vgpr_spill_count' "$OUT" | tr -dc 0-9)" diff --git a/benchmarks/kernels/gfx1151_decode_attn/tools/floor.py b/benchmarks/kernels/gfx1151_decode_attn/tools/floor.py new file mode 100644 index 000000000000..7a8cc6d867a2 --- /dev/null +++ b/benchmarks/kernels/gfx1151_decode_attn/tools/floor.py @@ -0,0 +1,193 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""The most any decode-attention kernel can reach against shapeset's roofline. + +The roofline is a floor for the problem, not for a kernel: one dispatch plus +the bytes at peak bandwidth. Real kernels pay latency on top -- a kernel +launch in a graph, a page-table read before the first KV address is known, a +DRAM round trip -- and at short context that is most of the time. This +measures that floor directly: a kernel that does nothing but stream N bytes, +after one dependent page-table read, cold, rotated through the same >= 96 MiB +working set and timed under a HIP graph exactly as the attention harness is. +No attention kernel can beat it, so its %roof at a configuration's byte count +is that configuration's ceiling. + + cd && amd-gpu-lock /bin/python \\ + benchmarks/kernels/gfx1151_decode_attn/tools/floor.py + +Prints the stream floor per size, then per (Hq, Hkv, D, M) the geomean %roof +the floor would score over the seven contexts matrix.py uses. +""" + +import argparse +import bisect +import math +import pathlib +import sys + +import torch +from torch.utils.cpp_extension import load_inline + +from vllm.triton_utils import triton + +sys.path.insert(0, str(pathlib.Path(__file__).resolve().parent)) +import shapeset # noqa: E402 + +_SRC = r""" +#include +#include +#include +typedef unsigned int u4 __attribute__((ext_vector_type(4))); +// Each workgroup first reads its page index -- the dependent load a paged +// kernel cannot avoid -- then streams its share of the buffer in b128. +template +__global__ __launch_bounds__(256) void stream(const u4* __restrict__ p, + const int* __restrict__ bt, + long n16, unsigned* out) { + const long stride = (long)gridDim.x * 256 * UNR; + u4 acc = {0, 0, 0, 0}; + const long base = __builtin_amdgcn_readfirstlane(bt[blockIdx.x & 7]); + for (long i = (long)blockIdx.x * 256 * UNR + threadIdx.x; i < n16; + i += stride) { + u4 v[UNR]; +#pragma unroll + for (int u = 0; u < UNR; ++u) { + const long j = i + (long)u * 256 + base; + v[u] = j < n16 ? p[j] : u4{0, 0, 0, 0}; + } +#pragma unroll + for (int u = 0; u < UNR; ++u) acc ^= v[u]; + } + const unsigned r = acc.x ^ acc.y ^ acc.z ^ acc.w; + if (r == 0x12345678u) out[blockIdx.x] = r; +} +__global__ void empty(unsigned* out) { + if (threadIdx.x == 999) out[0] = 1; +} +void run(torch::Tensor buf, torch::Tensor bt, int64_t n16, int64_t grid, + int64_t unr, torch::Tensor out) { + auto s = at::cuda::getCurrentCUDAStream(); + auto* p = (const u4*)buf.data_ptr(); + auto* o = (unsigned*)out.data_ptr(); + const int* b = bt.data_ptr(); + if (unr < 0) { + hipLaunchKernelGGL(empty, dim3(1), dim3(32), 0, s, o); + return; + } +#define L(U) hipLaunchKernelGGL((stream), dim3(grid), dim3(256), 0, s, p, b, n16, o) + if (unr == 1) L(1); else if (unr == 2) L(2); else if (unr == 4) L(4); else L(8); +} +""" + +CONTEXTS = [128, 512, 1024, 4096, 8192, 16384, 32768] + + +def _module(): + sdk = pathlib.Path(torch.__file__).resolve().parents[1] / "_rocm_sdk_devel" + return load_inline( + "rdna35_floor_probe", + cpp_sources="void run(torch::Tensor, torch::Tensor, int64_t, int64_t, " + "int64_t, torch::Tensor);", + cuda_sources=_SRC, + functions=["run"], + extra_cuda_cflags=["-O3"], + extra_ldflags=[f"-L{sdk / 'lib'}"], + ) + + +def _bench(mod, bufs, bts, n16, grid, unr, out): + def fn(): + for b, t in zip(bufs, bts): + mod.run(b, t, n16, grid, unr, out) + + fn() + torch.accelerator.synchronize() + ms = triton.testing.do_bench_cudagraph(fn, return_mode="median") + return ms * 1000 / len(bufs) + + +def measure(mod, sizes_kib, working_set_mb=96): + """Best-of-grid stream time, in us, for each size.""" + out = torch.zeros(4096, dtype=torch.int32, device="cuda") + floor = {} + for kib in sizes_kib: + nbytes = kib * 1024 + layers = max(10, -(-working_set_mb * 1024**2 // nbytes)) + bufs = [ + torch.randint(0, 255, (nbytes,), dtype=torch.uint8, device="cuda") + for _ in range(layers) + ] + bts = [torch.zeros(8, dtype=torch.int32, device="cuda") for _ in bufs] + n16 = nbytes // 16 + best = math.inf + for grid in (8, 20, 40, 80, 160, 320, 640, 1280): + for unr in (1, 2, 4, 8): + if grid * 256 * unr > 4 * n16 and grid > 8: + continue + best = min(best, _bench(mod, bufs, bts, n16, grid, unr, out)) + floor[kib] = best + del bufs, bts + torch.accelerator.empty_cache() + return floor + + +def floor_us(floor, nbytes): + """Log-log interpolation of the measured floor at nbytes.""" + pts = sorted(floor.items()) + xs = [math.log(k) for k, _ in pts] + x = math.log(max(nbytes / 1024, pts[0][0])) + i = min(max(bisect.bisect(xs, x), 1), len(xs) - 1) + (_, t0), (_, t1) = pts[i - 1], pts[i] + f = (x - xs[i - 1]) / (xs[i] - xs[i - 1]) + return math.exp(math.log(t0) + f * (math.log(t1) - math.log(t0))) + + +def main() -> None: + p = argparse.ArgumentParser() + shapeset.add_arguments(p) + p.add_argument( + "--sizes", + type=int, + nargs="+", + default=[32, 64, 128, 256, 512, 1024, 2048, 4096, 8192, 16384, 65536, 262144], + help="stream sizes in KiB; they must bracket every configuration's KV", + ) + p.add_argument("--m", type=int, nargs="+", default=[1, 4]) + args = p.parse_args() + + mod = _module() + out = torch.zeros(16, dtype=torch.int32, device="cuda") + tiny = [torch.zeros(16, dtype=torch.uint8, device="cuda") for _ in range(100)] + bts = [torch.zeros(8, dtype=torch.int32, device="cuda") for _ in tiny] + print(f"empty kernel: {_bench(mod, tiny, bts, 1, 1, -1, out):.2f} us") + + floor = measure(mod, args.sizes) + print("\n| KiB | floor us | roofline us | %roof |") + print("| --- | --- | --- | --- |") + for kib, us in floor.items(): + roof = shapeset.DISPATCH_US + kib * 1024 / (shapeset.PEAK_GIBS * 1024**3) * 1e6 + print(f"| {kib} | {us:.2f} | {roof:.2f} | {roof / us * 100:.1f} % |") + + shapes, _ = shapeset.load(args) + configs = sorted({(s.d, s.hq, s.hkv, s.window) for s in shapes if s.d != 96}) + win = "window | " if args.windowed else "" + print(f"\n| D | Hq | Hkv | {win}M | ceiling geomean | ceiling @128 |") + print("| --- " * (6 + bool(win)) + "|") + for d, hq, hkv, w in configs: + for m in args.m: + r = [] + for s in CONTEXTS: + roof = shapeset.roofline_us(hq, hkv, d, m, s, window=w) + nbytes = ( + (roof - shapeset.DISPATCH_US) * 1e-6 * shapeset.PEAK_GIBS * 1024**3 + ) + r.append(roof / floor_us(floor, nbytes)) + g = math.exp(sum(map(math.log, r)) / len(r)) * 100 + first = r[0] * 100 + wc = f"{w} | " if win else "" + print(f"| {d} | {hq} | {hkv} | {wc}{m} | {g:.1f} % | {first:.1f} % |") + + +if __name__ == "__main__": + main() diff --git a/benchmarks/kernels/gfx1151_decode_attn/tools/matrix.py b/benchmarks/kernels/gfx1151_decode_attn/tools/matrix.py new file mode 100644 index 000000000000..15e37f335756 --- /dev/null +++ b/benchmarks/kernels/gfx1151_decode_attn/tools/matrix.py @@ -0,0 +1,300 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Every shipped shape against Triton across the context range. + +roofline.py answers "where should I look next" in minutes by measuring one +context. This answers "what do we ship" and takes much longer, so it writes +markdown to stdout rather than pretending to be a gate. + + cd && VLLM_KV_CACHE_LAYOUT=HND PYTHONPATH=$PWD \\ + amd-gpu-lock /bin/python \\ + benchmarks/kernels/gfx1151_decode_attn/tools/matrix.py > table.md +""" + +import argparse +import contextlib +import statistics +import sys +import time +from pathlib import Path +from typing import Any + +import torch + +_HERE = Path(__file__).resolve() +_ROOT = _HERE.parents[4] +sys.path.insert(0, str(_ROOT / "benchmarks" / "attention_benchmarks")) +sys.path.insert(0, str(_HERE.parent)) + +import shapeset # noqa: E402 + + +def main() -> None: + p = argparse.ArgumentParser() + shapeset.add_arguments(p) + p.add_argument( + "--m", + type=int, + nargs="+", + default=[1, 4], + help="query tokens per sequence; 1 is plain decode, 4 is speculative", + ) + p.add_argument("--block-size", type=int, default=16) + p.add_argument( + "--reps", + type=int, + default=1, + help="whole re-setups per cell; do_bench medians many iterations " + "inside one already", + ) + p.add_argument( + "--contexts", + type=int, + nargs="+", + default=[128, 512, 1024, 4096, 8192, 16384, 32768], + ) + p.add_argument( + "--nseg", + type=int, + default=None, + help="force NSEG on every configuration, overriding the heuristics", + ) + p.add_argument( + "--rg", + type=int, + default=None, + help="force RG on every configuration, overriding the heuristics", + ) + p.add_argument( + "--minb", + type=int, + default=None, + help="force MINB on every configuration, overriding the heuristics", + ) + p.add_argument( + "--nw", + type=int, + default=None, + help="force NW on every configuration, overriding the heuristics", + ) + p.add_argument( + "--dspl", + type=int, + default=None, + help="force DSPL on every configuration, overriding the heuristics", + ) + p.add_argument( + "--rspl", + type=int, + default=None, + help="force RSPL on every configuration, overriding the heuristics", + ) + p.add_argument( + "--dot", + type=int, + default=None, + help="force DOT (the per-q-head dot decomposition) on every configuration", + ) + p.add_argument( + "--bfly", + type=int, + default=None, + help="force BFLY on every configuration (dot decomposition only)", + ) + p.add_argument( + "--vinlds", + type=int, + default=None, + help="force VINLDS on every configuration, overriding the heuristics", + ) + p.add_argument( + "--cpub", + type=int, + default=None, + help="force CPUB on every configuration, overriding the heuristics", + ) + p.add_argument( + "--no-triton", + dest="triton", + action="store_false", + help="skip the Triton column; %%roof does not need it and it is half the run", + ) + shapeset.add_dtype_argument(p) + args = p.parse_args() + dtype = shapeset.torch_dtype(args.dtype) + + from common import BenchmarkConfig + from runner import run_attention_benchmark + + import vllm.v1.attention.backends.rocm_attn as backend_mod + import vllm.v1.attention.ops.rdna35_hip_decode as decode_mod + from vllm.v1.attention.ops.rdna35_hip_decode import ( + KernelVariant, + precompile, + seal, + ) + from vllm.v1.attention.ops.rdna35_hip_decode import knobs_for as _knobs_for + + # Patch the module's own binding too: variant_for resolves knobs_for + # through the ops module, so rebinding only the local name here would + # precompile one set of variants and measure another. + forced = { + k: v + for k, v in ( + ("nseg", args.nseg), + ("rg", args.rg), + ("minb", args.minb), + ("nw", args.nw), + ("dspl", args.dspl), + ("rspl", args.rspl), + ("cpub", args.cpub), + ("vinlds", args.vinlds), + ("dot", args.dot), + ("bfly", args.bfly), + ) + if v is not None + } + if forced: + _heuristic = _knobs_for + + def _knobs_for(*args, **kwargs): # noqa: F811 + return {**_heuristic(*args, **kwargs), **forced} + + decode_mod.knobs_for = _knobs_for + + # Grouped by configuration, not listed by model. The table names 50 + # models but holds only 27 distinct (Hq, Hkv, D) tuples, and the kernel + # cannot tell two models with the same tuple apart -- measuring each is + # measuring the same thing twice. Ungrouped, 46% of the run was redundant. + # + # That redundancy did buy one thing before it was removed: 161 repeated + # cells put the harness noise floor at 0.2% median, 1.4% p90, 3.3% worst. + # The window is part of the key, not just the shape: gemma interleaves + # windowed and full-context layers, so one model contributes two rows with + # identical Hq/Hkv/D. Keyed on shape alone they would collapse into one + # group and the table would claim a windowed measurement it never took. + shapes, windowed = shapeset.load(args) + if not shapes: + print(f"") + return + groups: dict[tuple[int, int, int, int], list[str]] = {} + for sh in shapes: + groups.setdefault((sh.hq, sh.hkv, sh.d, sh.window), []).append(sh.model) + seen = set(groups) + rows = [(models, *key) for key, models in groups.items()] + + scratch: dict[Any, Any] = {} + orig_make = backend_mod.make_scratch + + def cached(variant, device): + if variant not in scratch: + scratch[variant] = orig_make(variant, device) + return scratch[variant] + + backend_mod.make_scratch = cached + + wanted = [] + for hq, hkv, d, window in seen: + for m in args.m: + for layout in (0, 1): + with contextlib.suppress(Exception): + wanted.append( + KernelVariant( + d, + hq, + hkv, + m, + args.block_size, + layout, + **_knobs_for(hq, hkv, d, m, window, dtype), + dtype=dtype, + window=window, + ) + ) + precompile(wanted) + # Nothing may build from here on: a build inside do_bench corrupts the + # median it reports (258% measured), which is why every cell used to pay a + # discarded warm-up call. Sealing makes the same hazard a loud failure and + # halves the run. Scratch is realised now for the same reason -- its memset + # would otherwise land in the first timed iteration. + for variant in wanted: + with contextlib.suppress(Exception): + cached(variant, torch.device("cuda:0")) + seal() + time.sleep(5) # clocks settle after a parallel build; see 00-protocol.md + + def timeit(backend, hq, hkv, d, s, m, window): + cfg = BenchmarkConfig( + backend=backend, + batch_spec=f"q{m}s{s}", + num_layers=10, + min_working_set_mb=96, + head_dim=d, + num_q_heads=hq, + num_kv_heads=hkv, + block_size=args.block_size, + device="cuda:0", + dtype=dtype, + sliding_window=window or None, + ) + rs = [run_attention_benchmark(cfg) for _ in range(args.reps)] + spread = max(r.std_time / r.median_time for r in rs) * 100 + return statistics.median(r.median_time for r in rs) * 1e6, spread + + active = shapeset.describe(args) + print( + f"" + ) + print( + "| modelos | S | M | Hq | Hkv | D | roofline | Triton | nuestro | vs " + "| %roof | spread |" + ) + print("| --- | --- | --- | --- | --- | --- | --- | --- | --- | --- | --- | --- |") + # Ordered by (S, M) first and configuration last, so every configuration + # sits side by side at the same context and token count -- the comparison + # the table exists to support. It also walks the contexts upward across + # the whole run, which is what warms the allocator: a pass that starts at a + # large context reads its first cell per configuration up to 10x slow. + for s in args.contexts: + for m in args.m: + for models, hq, hkv, d, window in rows: + roof = shapeset.roofline_us( + hq, hkv, d, m, s, block_size=args.block_size, window=window + ) + impls: list[Any] = [] + orig_init = backend_mod.RocmAttentionRdna35Impl.__init__ + + def spy(self, *a, _init=orig_init, _seen=impls, **kw): + _init(self, *a, **kw) + _seen.append(self) + + backend_mod.RocmAttentionRdna35Impl.__init__ = spy + try: + ours, spread = timeit("ROCM_ATTN", hq, hkv, d, s, m, window) + finally: + backend_mod.RocmAttentionRdna35Impl.__init__ = orig_init + ran = any(i.kernel_calls for i in impls) and not any( + i.fallback_calls for i in impls + ) + tri = None + if args.triton: + tri, _ = timeit("TRITON_ATTN", hq, hkv, d, s, m, window) + mark = "" if ran else " (fallback)" + label = ( + models[0] if len(models) == 1 else f"{models[0]} +{len(models) - 1}" + ) + (f" w={window}" if window else "") + print( + f"| {label}{mark} | {s} | {m} | {hq} | {hkv} | {d} | " + f"{roof:.2f} | {'-' if tri is None else f'{tri:.2f}'} | " + f"{ours:.2f} | {'-' if tri is None else f'{tri / ours:.2f}x'} | " + f"{roof / ours * 100:.1f} % | {spread:.1f} % |", + flush=True, + ) + + +if __name__ == "__main__": + main() diff --git a/benchmarks/kernels/gfx1151_decode_attn/tools/roofline.py b/benchmarks/kernels/gfx1151_decode_attn/tools/roofline.py new file mode 100644 index 000000000000..633ae145458e --- /dev/null +++ b/benchmarks/kernels/gfx1151_decode_attn/tools/roofline.py @@ -0,0 +1,278 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Where every shipped shape sits against its KV roofline, worst first. + +This is both halves of the optimisation loop in one command: it is the +regression gate, because it checks every shape against a float reference, and +it is the opportunity finder, because it sorts by distance from the roofline. + + cd && VLLM_KV_CACHE_LAYOUT=HND PYTHONPATH=$PWD \\ + amd-gpu-lock /bin/python \\ + benchmarks/kernels/gfx1151_decode_attn/tools/roofline.py + +S=128 only, deliberately. The fixed cost is the same at every context and the +roofline budget is smallest here, so this is where a shape's distance from the +bus is most visible -- and one context keeps a full 50-shape pass to minutes. + +`ran` is the column to read first. The backend falls back to Triton silently +for any shape it cannot serve, and a harness that does not check reports +Triton's number as if it were ours; that has happened twice on this project. +""" + +import argparse +import contextlib +import statistics +import sys +import time +from pathlib import Path +from typing import Any + +import torch + +_HERE = Path(__file__).resolve() +_ROOT = _HERE.parents[4] +sys.path.insert(0, str(_ROOT / "benchmarks" / "attention_benchmarks")) +sys.path.insert(0, str(_HERE.parent)) + +import shapeset # noqa: E402 + +RTOL = 1e-3 + + +def reference(q, kv, s, hq, hkv, head_dim, m): + flat = kv.transpose(1, 2).flatten(0, 1)[:s] + k, v = flat[..., :head_dim], flat[..., head_dim:] + gqa = hq // hkv + qf = q.float().permute(1, 0, 2) + kf = k.float().permute(1, 0, 2).repeat_interleave(gqa, 0) + vf = v.float().permute(1, 0, 2).repeat_interleave(gqa, 0) + scores = torch.bmm(qf, kf.transpose(1, 2)) * (head_dim**-0.5) + pos = torch.arange(s, device=q.device).view(1, s) + lim = (s - m + torch.arange(m, device=q.device)).view(m, 1) + scores = scores.masked_fill((pos > lim).view(1, m, s), float("-inf")) + return torch.bmm(torch.softmax(scores, -1), vf).permute(1, 0, 2) + + +def correctness(hq, hkv, d, m, block_size, dtype, s=48): + """max_rel against a float reference, or None if the shape does not build. + + S=48 rather than 128: a causal off-by-one moves the softmax by ~1/S while + the tolerance is fixed, so a short context is the one that catches it. + """ + from vllm.v1.attention.ops.rdna35_hip_decode import ( + KernelVariant, + UnexpectedBuildError, + load, + make_scratch, + ) + from vllm.v1.attention.ops.rdna35_hip_decode import knobs_for as _knobs_for + + dev = torch.device("cuda") + torch.manual_seed(0) + blocks = -(-s // block_size) + kv = torch.randn(blocks, hkv, block_size, 2 * d, device=dev, dtype=dtype) + kv *= 0.5 + q = torch.randn(m, hq, d, device=dev, dtype=dtype) * 0.5 + bt = torch.arange(blocks, device=dev, dtype=torch.int32) + try: + variant = KernelVariant( + d, + hq, + hkv, + m, + block_size, + 1, + **_knobs_for(hq, hkv, d, m), + dtype=dtype, + ) + module = load(variant) + acc, smax, ssum, arrivals = make_scratch(variant, dev) + except UnexpectedBuildError: + # Never swallowed: the seal exists to be heard. + raise + except Exception: + return None + out = torch.empty_like(q) + seq_lens = torch.tensor([s], device=dev, dtype=torch.int32) + module.decode_attn(q, kv, bt, out, acc, smax, ssum, arrivals, seq_lens, d**-0.5) + torch.accelerator.synchronize() + ref = reference(q, kv, s, hq, hkv, d, m) + got = out.float() + if not torch.isfinite(got).all(): + return float("inf") + return ((got - ref).abs() / ref.abs().clamp_min(1e-3)).max().item() + + +def main() -> None: + p = argparse.ArgumentParser() + shapeset.add_arguments(p) + p.add_argument("--s", type=int, default=128) + p.add_argument("--m", type=int, default=4) + p.add_argument("--block-size", type=int, default=16) + p.add_argument( + "--reps", + type=int, + default=1, + help="whole re-setups per cell; do_bench medians many iterations " + "inside one already", + ) + p.add_argument("--triton", action="store_true", help="also time TRITON_ATTN") + p.add_argument("--skip-check", action="store_true", help="timings only") + shapeset.add_dtype_argument(p) + args = p.parse_args() + dtype = shapeset.torch_dtype(args.dtype) + + from common import BenchmarkConfig + from runner import run_attention_benchmark + + import vllm.v1.attention.backends.rocm_attn as backend_mod + + rows, windowed = shapeset.load(args) + if not rows: + print(f"no shapes match {shapeset.describe(args) or 'the given filters'}") + return + + # Scratch is memoised for the same reason sweep.py memoises it: a rebuild + # per forward drags a device memset into the timed region. + scratch: dict[Any, Any] = {} + orig_make = backend_mod.make_scratch + + def cached(variant, device): + if variant not in scratch: + scratch[variant] = orig_make(variant, device) + return scratch[variant] + + backend_mod.make_scratch = cached + + # Build every variant the pass needs before timing anything. A cold build + # is ~19.6 s and they are independent, so serially the 27 distinct variants + # of this table cost nine minutes before the first measurement; spread over + # the machine it is under a minute. Nothing here is timed, so a build that + # lands late cannot contaminate a number. + from vllm.v1.attention.ops.rdna35_hip_decode import ( + KernelVariant, + precompile, + seal, + ) + from vllm.v1.attention.ops.rdna35_hip_decode import knobs_for as _knobs_for + + wanted = [] + for _, hq, hkv, d, _window in rows: + for layout in (0, 1): + with contextlib.suppress(Exception): + wanted.append( + KernelVariant( + d, + hq, + hkv, + args.m, + args.block_size, + layout, + **_knobs_for(hq, hkv, d, args.m), + dtype=dtype, + ) + ) + precompile(wanted) + # Nothing may build from here on: a build inside do_bench corrupts the + # median it reports (258% measured), which is why every cell used to pay a + # discarded warm-up call. Sealing makes the same hazard a loud failure and + # halves the run. Scratch is realised now for the same reason -- its memset + # would otherwise land in the first timed iteration. + for variant in wanted: + with contextlib.suppress(Exception): + cached(variant, torch.device("cuda:0")) + seal() + # Let the machine settle before timing anything. Compiling saturates the + # cores and moves the SoC clock, which is the contamination 00-protocol.md + # invented quiet-lock to avoid; doing it in-process moments before the + # first measurement is the same mistake by another route. A pass taken + # straight after a 50-way rebuild reported three shapes falling back that + # were served on every quiet run before and after. + time.sleep(5) + + active = shapeset.describe(args) + if active: + print(f"filtered to {len(rows)} shapes: {active}") + hdr = f"{'model':<36} {'Hq':>3} {'Hkv':>4} {'D':>4} {'ran':>4} {'max_rel':>9}" + hdr += f" {'us':>9} {'%roof':>6} {'spread':>7}" + if args.triton: + hdr += f" {'triton':>9} {'vs':>6}" + print(hdr) + + out_rows = [] + for model, hq, hkv, d, _window in rows: + rel = ( + None + if args.skip_check + else correctness(hq, hkv, d, args.m, args.block_size, dtype, s=48) + ) + roof = shapeset.roofline_us( + hq, hkv, d, args.m, args.s, block_size=args.block_size + ) + + def timeit(backend, d=d, hq=hq, hkv=hkv): + cfg = BenchmarkConfig( + backend=backend, + batch_spec=f"q{args.m}s{args.s}", + num_layers=10, + min_working_set_mb=96, + head_dim=d, + num_q_heads=hq, + num_kv_heads=hkv, + block_size=args.block_size, + device="cuda:0", + dtype=dtype, + ) + rs = [run_attention_benchmark(cfg) for _ in range(args.reps)] + spread = max(r.std_time / r.median_time for r in rs) * 100 + return statistics.median(r.median_time for r in rs) * 1e6, spread + + impls: list[Any] = [] + orig_init = backend_mod.RocmAttentionRdna35Impl.__init__ + + def spy(self, *a, _init=orig_init, _seen=impls, **kw): + _init(self, *a, **kw) + _seen.append(self) + + backend_mod.RocmAttentionRdna35Impl.__init__ = spy + try: + us, spread = timeit("ROCM_ATTN") + finally: + backend_mod.RocmAttentionRdna35Impl.__init__ = orig_init + ran = any(i.kernel_calls for i in impls) and not any( + i.fallback_calls for i in impls + ) + # Why, not just whether. warning_once suppresses the backend's own + # message after the first shape hits a given reason, so a pass over + # many shapes hides every fallback but the first. + why = next((i._rejected for i in impls if i._rejected), None) + + line = f"{model:<36} {hq:>3} {hkv:>4} {d:>4} {'yes' if ran else 'NO':>4}" + line += f" {'-' if rel is None else f'{rel:.2e}':>9}" + line += f" {us:>9.2f} {roof / us * 100:>5.1f}% {spread:>6.1f}%" + if args.triton: + t, _ = timeit("TRITON_ATTN") + line += f" {t:>9.2f} {t / us:>5.2f}x" + if not ran and why: + line += f" <- {why}" + print(line, flush=True) + out_rows.append((roof / us * 100, ran, model, hq, hkv, d, us)) + + print("\nworst first (only shapes the kernel actually serves):") + served = [r for r in out_rows if r[1]] + for pct, _, model, hq, hkv, d, us in sorted(served): + print(f" {pct:>5.1f}% {model:<36} Hq={hq} Hkv={hkv} D={d} {us:.2f} us") + missing = [r for r in out_rows if not r[1]] + if missing: + by_d: dict[int, int] = {} + for _, _, _, _, _, d, _ in missing: + by_d[d] = by_d.get(d, 0) + 1 + spread = ", ".join(f"D={k}: {v}" for k, v in sorted(by_d.items())) + print(f"\nnot served, {len(missing)} of {len(out_rows)} shapes -> {spread}") + if windowed: + print(f"skipped {windowed} sliding-window shapes: the kernel is full-context") + + +if __name__ == "__main__": + main() diff --git a/benchmarks/kernels/gfx1151_decode_attn/tools/shapes.csv b/benchmarks/kernels/gfx1151_decode_attn/tools/shapes.csv new file mode 100644 index 000000000000..7894ba8751ea --- /dev/null +++ b/benchmarks/kernels/gfx1151_decode_attn/tools/shapes.csv @@ -0,0 +1,60 @@ +model,Hq,Hkv,D,window +Qwen2.5-0.5B-Instruct,14,2,64,0 +Llama-3.2-1B-Instruct,32,8,64,0 +SmolLM2-1.7B-Instruct,32,32,64,0 +MiniCPM-V-0.53B-bosch,16,2,64,0 +geely-vlm-0320-ckpt14000,16,2,64,0 +Phi-3.5-vision-instruct,32,32,96,0 +Qwen2.5-3B-Instruct,16,2,128,0 +Qwen2.5-7B-Instruct,28,4,128,0 +DeepSeek-R1-Distill-Qwen-7B,28,4,128,0 +Qwen3-1.7B,16,8,128,0 +Qwen3-4B,32,8,128,0 +Qwen3-8B,32,8,128,0 +Qwen3-30B-A3B-Instruct-2507,32,4,128,0 +Llama-3.1-8B-Instruct,32,8,128,0 +Llama-2-7B,32,32,128,0 +Llama-3.2-3B-Instruct,24,8,128,0 +granite-4.1-8b,32,8,128,0 +Ministral-3-3B-Instruct-2512,32,8,128,0 +Ministral-3-8B-Instruct-2512,32,8,128,0 +deepseek-vl2-tiny,10,10,128,0 +Phi-4-multimodal-instruct,24,8,128,0 +cogagent-chat-hf,32,32,128,0 +MiniCPM-o-2_6,28,4,128,0 +MiniCPM-V-2_6,28,4,128,0 +MiniCPM-V-custom-checkpoint-2400,32,2,128,0 +Qwen2.5-VL-3B-Instruct,16,2,128,0 +Qwen2.5-VL-7B-Instruct,28,4,128,0 +Qwen3-VL-4B-Instruct,32,8,128,0 +Cosmos-Reason2-2B-LLM,16,8,128,0 +Cosmos-Reason2-8B-LLM,32,8,128,0 +Qwen3-Omni-30B-A3B-thinker,32,4,128,0 +Qwen3-Omni-30B-A3B-talker,16,2,128,0 +Qwen3-Omni-30B-A3B-code-predictor,16,8,128,0 +NVIDIA-Nemotron-3-Nano-4B,40,8,128,0 +gemma-2b-it,8,1,256,0 +paligemma2-3b-mix-448-LLM,8,4,256,0 +gemma-3-4b-it,8,4,256,0 +gemma-3-12b-it,16,8,256,0 +Qwen3.5-0.8B,8,2,256,0 +Qwen3.5-2B,8,2,256,0 +Qwen3.5-9B,16,4,256,0 +Qwen3.6-27B,24,4,256,0 +Qwen3.5-35B-A3B,16,2,256,0 +Qwen3.6-35B-A3B,16,2,256,0 +gemma-4-E2B-it,8,1,256,512 +gemma-4-E4B-it,8,2,256,512 +gemma-3-4b-it,8,4,256,1024 +gemma-3-12b-it,16,8,256,1024 +gemma-4-12b-it,16,8,256,1024 +gemma-4-26B-A4B-it,16,8,256,1024 +gemma-4-31B-it-AWQ,32,16,256,1024 +gemma-4-31B-it-assistant,32,16,256,1024 +paligemma2-3b-mix-448-LLM,8,4,256,4096 +gemma-4-E2B-it,8,1,512,0 +gemma-4-E4B-it,8,2,512,0 +gemma-4-12b-it,16,1,512,0 +gemma-4-26B-A4B-it,16,2,512,0 +gemma-4-31B-it-AWQ,32,4,512,0 +gemma-4-31B-it-assistant,32,4,512,0 \ No newline at end of file diff --git a/benchmarks/kernels/gfx1151_decode_attn/tools/shapeset.py b/benchmarks/kernels/gfx1151_decode_attn/tools/shapeset.py new file mode 100644 index 000000000000..5b3e94373ac3 --- /dev/null +++ b/benchmarks/kernels/gfx1151_decode_attn/tools/shapeset.py @@ -0,0 +1,184 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Loading and filtering shapes.csv. + +Shared by roofline.py, matrix.py and tune.py so that `--hkv 8` selects the same +rows whichever tool is asked. Selecting by hand instead -- editing the csv, or +passing a model substring that happens to cover the shapes you meant -- is how +a pass ends up measuring a different set than the one it reports. +""" + +import argparse +import csv +import os +from pathlib import Path +from typing import NamedTuple + +# The tools measure variants outside variants.def -- every point a search +# visits -- so they JIT-build all of them rather than use _rocm_C's. +os.environ.setdefault("VLLM_RDNA35_ATTN_JIT", "1") + +_SHAPES = Path(__file__).resolve().parent / "shapes.csv" + +# Measured peak DRAM bandwidth, LPDDR5X-8000 on Strix Halo: 256 GiB/s +# theoretical, 230 GiB/s achieved by a tuned streaming kernel. +PEAK_GIBS = 230.0 + +# Per-kernel dispatch overhead on the 8060S, measured under a HIP graph with a +# single-wave spin, so it is the floor even when the launch is already graphed +# and not an artifact of eager dispatch. Attention cannot be computed without +# dispatching at least one kernel, so it belongs in an algorithmic floor. +DISPATCH_US = 1.48 + + +def roofline_us( + hq: int, + hkv: int, + d: int, + m: int, + s: int, + itemsize: int = 2, + block_size: int = 16, + window: int = 0, +) -> float: + """The floor for one decode-attention call, in microseconds. + + Deliberately algorithmic. It counts what the problem requires -- one kernel + dispatch, the query in, the KV cache in, the output out, and the block + table that addresses the pages -- and nothing an implementation chose. Our + own split-KV partials go through global memory and can exceed the KV + traffic at short context, but they are excluded on purpose: a denominator + that grew with a kernel's own overhead would let a wasteful kernel report + itself near the roof. One dispatch is counted rather than NSEG of them, for + the same reason -- the decomposition is ours, the single launch is not. + + The dispatch term dominates at short context: at S=128 it is more than the + bytes for every shape in the table, which is why %roof there used to read + absurdly low. So this is a floor for any implementation, not a model of + ours. + + Args: + hq: Query heads. + hkv: KV heads. + d: Head dimension. + m: Query tokens per sequence. + s: Context length. + itemsize: Bytes per element of Q, K, V and the output. + block_size: KV cache page size, for the block table. + window: Sliding window in keys, 0 for full attention. A windowed layer + needs only the keys some query can see: min(s, window + m - 1). + + Returns: + One dispatch plus the time those bytes take at peak bandwidth, in + microseconds. + """ + keys = s if not window else min(s, window + m - 1) + query = m * hq * d * itemsize + kv = keys * hkv * 2 * d * itemsize + out = m * hq * d * itemsize + table = -(-keys // block_size) * 4 + moved = (query + kv + out + table) / (PEAK_GIBS * 1024**3) * 1e6 + return DISPATCH_US + moved + + +class Shape(NamedTuple): + model: str + hq: int + hkv: int + d: int + window: int + + +def add_arguments(p: argparse.ArgumentParser) -> None: + p.add_argument("--shapes", default=str(_SHAPES)) + p.add_argument("--filter", default="", help="substring match on the model name") + p.add_argument("--hq", type=int, nargs="+", help="keep only these Hq") + p.add_argument("--hkv", type=int, nargs="+", help="keep only these Hkv") + p.add_argument("--head-dim", type=int, nargs="+", help="keep only these D") + p.add_argument("--gqa", type=int, nargs="+", help="keep only these Hq/Hkv") + p.add_argument( + "--windowed", + action="store_true", + help="the sliding-window rows instead of the full-attention ones", + ) + + +def add_dtype_argument(p: argparse.ArgumentParser) -> None: + p.add_argument( + "--dtype", + choices=("fp16", "bf16"), + default="fp16", + help="element type of Q, the KV cache and the output", + ) + + +def torch_dtype(name: str): + """The torch dtype for a `--dtype` value.""" + import torch + + return {"fp16": torch.float16, "bf16": torch.bfloat16}[name] + + +# Relative error bound check.py applies per dtype; bf16 output rounding alone +# is up to 2^-8 relative. +RTOL = {"fp16": 1e-3, "bf16": 8e-3} + + +def load(args: argparse.Namespace) -> tuple[list[Shape], int]: + """The shapes the kernel can serve, after the filters in `args`. + + Returns the surviving shapes and how many sliding-window rows were dropped. + The window is checked last so that the count describes the selection rather + than the whole file. + + Args: + args: Namespace populated by `add_arguments`. + + Returns: + A `(shapes, windowed)` pair. + """ + shapes: list[Shape] = [] + windowed = 0 + with open(args.shapes) as fh: + for r in csv.DictReader(fh): + s = Shape( + r["model"], + int(r["Hq"]), + int(r["Hkv"]), + int(r["D"]), + int(r["window"]), + ) + if args.filter and args.filter not in s.model: + continue + if args.hq and s.hq not in args.hq: + continue + if args.hkv and s.hkv not in args.hkv: + continue + if args.head_dim and s.d not in args.head_dim: + continue + # A ratio that is not a whole number matches no integer --gqa, which + # is the wanted answer rather than a rounding decision. + if args.gqa and (s.hq % s.hkv or s.hq // s.hkv not in args.gqa): + continue + if bool(s.window) != getattr(args, "windowed", False): + windowed += bool(s.window) + continue + shapes.append(s) + return shapes, windowed + + +def describe(args: argparse.Namespace) -> str: + """The active filters, for tools that print what they measured.""" + bits = [] + for name, val in ( + ("model~", args.filter), + ("Hq", args.hq), + ("Hkv", args.hkv), + ("D", args.head_dim), + ("Hq/Hkv", args.gqa), + ): + if val: + joined = val if isinstance(val, str) else ",".join(str(v) for v in val) + bits.append(f"{name}={joined}") + return " ".join(bits) diff --git a/benchmarks/kernels/gfx1151_decode_attn/tools/sweep.py b/benchmarks/kernels/gfx1151_decode_attn/tools/sweep.py new file mode 100644 index 000000000000..8784a8d852ee --- /dev/null +++ b/benchmarks/kernels/gfx1151_decode_attn/tools/sweep.py @@ -0,0 +1,276 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Compare ROCM_ATTN (gfx1151) against TRITON_ATTN through vLLM's own harness. + +Both backends go through run_attention_benchmark, so the KV tensors, the +metadata and the timer are identical; only the kernel launch differs. + + cd && VLLM_KV_CACHE_LAYOUT=HND \\ + PYTHONPATH=$PWD amd-gpu-lock /bin/python \\ + benchmarks/kernels/gfx1151_decode_attn/tools/sweep.py --hq 8 --hkv 4 + +Pass --nseg/--rg/--minb/--nw/--dspl/--rspl to override what the backend would pick, +which is how the tuning sweeps are run. +""" + +import argparse +import contextlib +import itertools +import statistics +import sys +import time +from dataclasses import replace +from pathlib import Path +from typing import Any + +import torch + +_HERE = Path(__file__).resolve() +_ROOT = _HERE.parents[4] +sys.path.insert(0, str(_ROOT / "benchmarks" / "attention_benchmarks")) +sys.path.insert(0, str(_HERE.parent)) + +import shapeset # noqa: E402 + + +def main() -> None: + p = argparse.ArgumentParser() + p.add_argument("--hq", type=int, default=32) + p.add_argument("--hkv", type=int, default=16) + p.add_argument("--head-dim", type=int, default=256) + p.add_argument("--m", type=int, default=4, help="query tokens per sequence") + p.add_argument("--block-size", type=int, default=16) + p.add_argument( + "--reps", + type=int, + default=1, + help="whole re-setups per cell; do_bench already medians many " + "iterations inside one, so this only resamples allocation and " + "graph-capture placement", + ) + p.add_argument( + "--contexts", + type=int, + nargs="+", + default=[128, 16384, 32768], + help="three points, not a curve: the shortest context where fixed cost " + "dominates and the two longest where bandwidth does. A knob has to " + "serve the whole range with one value, so what matters is the ends", + ) + p.add_argument( + "--nseg", type=int, nargs="+", default=[None], help="most KV segments" + ) + p.add_argument( + "--rg", type=int, nargs="+", default=[None], help="row groups per kv head" + ) + p.add_argument( + "--minb", + type=int, + nargs="+", + default=[None], + help="least KV blocks per active segment", + ) + p.add_argument( + "--nw", type=int, nargs="+", default=[None], help="waves per workgroup" + ) + p.add_argument( + "--dspl", + type=int, + nargs="+", + default=[None], + help="waves sharing a key tile, each owning 1/dspl of the head dim", + ) + p.add_argument( + "--rspl", + type=int, + nargs="+", + default=[None], + help="waves sharing a key tile, each carrying 1/rspl of the row tiles", + ) + p.add_argument( + "--ablate", + type=int, + nargs="+", + default=[None], + help="MEASUREMENT ONLY, wrong numbers; see ABLATE in the kernel", + ) + p.add_argument("--triton", action="store_true", help="also measure Triton") + shapeset.add_dtype_argument(p) + args = p.parse_args() + dtype = shapeset.torch_dtype(args.dtype) + + from common import BenchmarkConfig + from runner import run_attention_benchmark + + import vllm.v1.attention.backends.rocm_attn as backend_mod + from vllm.v1.attention.ops.rdna35_hip_decode import ( + KernelVariant, + load, + make_scratch, + precompile, + seal, + ) + + knobs = ("nseg", "rg", "minb", "nw", "dspl", "rspl", "ablate") + combos = [ + {k: v for k, v in zip(knobs, vals, strict=True) if v is not None} + for vals in itertools.product(*(getattr(args, k) for k in knobs)) + ] + + # One row per combination, one column per context. Passing several values + # per knob builds the matrix here rather than in a shell loop: a loop pays + # interpreter startup and a serial build for every cell, while this one + # compiles the whole product at once and measures in a single process. + original = backend_mod.RocmAttentionRdna35Impl._prepare + scratch: dict[Any, Any] = {} + current: dict[str, Any] = {} + + def cached_scratch(variant, device): + if variant not in scratch: + scratch[variant] = make_scratch(variant, device) + return scratch[variant] + + # An override the backend would not have chosen leaves self._variant + # permanently disagreeing with what _prepare computes, so it rebuilds every + # forward; make_scratch zeroes the arrival counters, which is a device + # memset, i.e. an extra kernel launch inside the timed region. + backend_mod.make_scratch = cached_scratch + + def patched(self, kv_cache, **kw): + built = original(self, kv_cache, **kw) + if built is not None and current: + # Rebuild from the heuristic knobs, not from self._variant, so the + # override lands on the same base precompile() started from. + want = replace(self._variant, **{**base, **current}) + if self._variant != want: + self._variant = want + self._built = (load(want), cached_scratch(want, kw["q"].device)) + built = self._built + return built + + backend_mod.RocmAttentionRdna35Impl._prepare = patched + + # The override in `patched` starts from what _prepare computed -- the + # tuned knobs -- and replaces only the swept ones, so precompilation has to + # start from the same base. Building from the dataclass defaults instead + # produced a different variant than the one measured, which the discarded + # warm-up call used to hide by absorbing the rebuild. + from vllm.v1.attention.ops.rdna35_hip_decode import knobs_for as _knobs_for + + base = dict(_knobs_for(args.hq, args.hkv, args.head_dim, args.m)) + wanted = [] + for combo in combos: + for layout in (0, 1): + with contextlib.suppress(Exception): + wanted.append( + KernelVariant( + args.head_dim, + args.hq, + args.hkv, + args.m, + args.block_size, + layout, + **{**base, **combo}, + dtype=dtype, + ) + ) + # The base variant too: `patched` calls the unpatched _prepare first, which + # builds whatever the backend would have chosen before the override + # replaces it. Precompiling only the overridden variants leaves that first + # build to land inside the timed region -- which is what the seal now + # refuses. + for layout in (0, 1): + with contextlib.suppress(Exception): + wanted.append( + KernelVariant( + args.head_dim, + args.hq, + args.hkv, + args.m, + args.block_size, + layout, + **base, + dtype=dtype, + ) + ) + precompile(wanted) + # Nothing may build from here on: a build inside do_bench corrupts the + # median it reports (258% measured), which is why every cell used to pay a + # discarded warm-up call. Sealing makes the same hazard a loud failure and + # halves the run. Scratch is realised now for the same reason -- its memset + # would otherwise land in the first timed iteration. + for variant in wanted: + with contextlib.suppress(Exception): + cached_scratch(variant, torch.device("cuda:0")) + seal() + time.sleep(5) # clocks settle after a parallel build; see 00-protocol.md + + def timeit(backend, s): + cfg = BenchmarkConfig( + backend=backend, + batch_spec=f"q{args.m}s{s}", + num_layers=10, + min_working_set_mb=96, + head_dim=args.head_dim, + num_q_heads=args.hq, + num_kv_heads=args.hkv, + block_size=args.block_size, + device="cuda:0", + dtype=dtype, + ) + rs = [run_attention_benchmark(cfg) for _ in range(args.reps)] + # do_bench's own dispersion over its iterations, worst of the reps. A + # wide cell flags itself here instead of being averaged into silence. + spread = max(r.std_time / r.median_time for r in rs) * 100 + return statistics.median(r.median_time for r in rs) * 1e6, spread + + label_w = max(24, *(len(_label(c)) for c in combos)) + print(f"Hq={args.hq} Hkv={args.hkv} D={args.head_dim} M={args.m} {args.dtype}") + print( + f"{'':<{label_w}} " + + " ".join(f"{s:>9}" for s in args.contexts) + + f" {'spread':>7}" + ) + roofs = [ + shapeset.roofline_us( + args.hq, args.hkv, args.head_dim, args.m, s, block_size=args.block_size + ) + for s in args.contexts + ] + print(f"{'roofline':<{label_w}} " + " ".join(f"{r:>9.2f}" for r in roofs)) + if args.triton: + current.clear() + cells = [timeit("TRITON_ATTN", s) for s in args.contexts] + print( + f"{'TRITON_ATTN':<{label_w}} " + + " ".join(f"{t:>9.2f}" for t, _ in cells) + + f" {max(d for _, d in cells):>6.1f}%" + ) + for combo in combos: + current.clear() + current.update(combo) + try: + cells = [timeit("ROCM_ATTN", s) for s in args.contexts] + except Exception as exc: # noqa: BLE001 - a combination may not build + # The message, not just the class: a combination that legitimately + # fails a static_assert and one the seal caught because it was + # never precompiled are both exceptions here, and only the text + # tells them apart. + print(f"{_label(combo):<{label_w}} {exc.__class__.__name__}: {exc}") + continue + print( + f"{_label(combo):<{label_w}} " + + " ".join(f"{u:>9.2f}" for u, _ in cells) + + f" {max(d for _, d in cells):>6.1f}%" + ) + + +def _label(combo) -> str: + if not combo: + return "backend defaults" + return " ".join(f"{k}={v}" for k, v in combo.items()) + + +if __name__ == "__main__": + main() diff --git a/benchmarks/kernels/gfx1151_decode_attn/tools/tune.py b/benchmarks/kernels/gfx1151_decode_attn/tools/tune.py new file mode 100644 index 000000000000..0948155162f6 --- /dev/null +++ b/benchmarks/kernels/gfx1151_decode_attn/tools/tune.py @@ -0,0 +1,364 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Find the one set of launch knobs per configuration that serves every context. + +The constraint that shapes this whole search: a configuration may not depend +on S. The grid is fixed when the CUDA graph is captured and S is a runtime +argument, so one choice has to serve 128 and 32768 alike. Tuning per context +would be easy and unshippable. + +So the objective is the whole range: the geomean of the fraction of roofline +over the seven contexts matrix.py reports -- after the worst long-context cell +(S >= 16384) has reached 90 % of roof, which comes first. +A candidate that wins the long contexts by starving the short ones scores +badly on both. + +The search is coordinate descent -- one knob at a time over its values, the +rest held, keeping the best -- starting from the backend's own choice. The +knobs interact, which is why --rounds exists; an exhaustive grid is ~100 +builds per configuration. + + cd && VLLM_KV_CACHE_LAYOUT=HND PYTHONPATH=$PWD \\ + amd-gpu-lock /bin/python \\ + benchmarks/kernels/gfx1151_decode_attn/tools/tune.py --hq 32 --hkv 8 \\ + --head-dim 128 --m 1 + +Prints one `_TUNED` row per configuration, ready to paste. +""" + +import argparse +import contextlib +import math +import sys +import time +from pathlib import Path +from typing import Any + +_HERE = Path(__file__).resolve() +_ROOT = _HERE.parents[4] +sys.path.insert(0, str(_ROOT / "benchmarks" / "attention_benchmarks")) +sys.path.insert(0, str(_HERE.parent)) + +import shapeset # noqa: E402 + +CONTEXTS = [128, 512, 1024, 4096, 8192, 16384, 32768] +# Long-context workgroup counts worth trying; see _TARGET_WORKGROUPS in the +# backend for why the answer is small. +TARGETS = (8, 16, 32) + + +def geomean(xs): + return math.exp(sum(math.log(x) for x in xs) / len(xs)) + + +# Long contexts must reach this fraction of roof before anything else counts. +LONG_S, LONG_TARGET = 16384, 0.90 + + +def rank(ctxs, entry): + """How a measured point compares: first its worst long-context cell, up to + the target, then the geomean over every context. A point that trades + long-context throughput below the target for short-context gains loses.""" + g, cells = entry + if not cells: # did not build or run + return (-1.0, 0.0) + long_ = [c for s, c in zip(ctxs, cells) if s >= LONG_S] + floor_ = min(min(long_), LONG_TARGET) if long_ else LONG_TARGET + return (floor_, g) + + +def space(hq, hkv, d, start): + """The values each knob is searched over, for one configuration.""" + gqa = hq // hkv + rg0 = start["rg"] + default_dspl = d // 128 if d >= 256 else 1 + rgs = {1, rg0, 2 * rg0, 4 * rg0, max(1, rg0 // 2)} + return { + "nw": [2, 4, 8], + "dspl": [0, 2 * default_dspl] if d >= 128 else [0], + "rg": sorted(r for r in rgs if gqa % r == 0), + "rspl": [1, 2, 4], + "pf": [0, 1], + "cpub": [0, 1], + "vinlds": [0, 1], + "target": list(TARGETS), + "minb": [1, 2, 4], + # The dot decomposition's own knobs. + "dnseg": [1, 2, 4, 8], + "bfly": [0, 2, 4], + "gt": [0, 1], + } + + +def knobs_of(hkv, cand): + """KernelVariant keyword arguments for a search point.""" + if cand.get("dot"): + k = {"dot": 1, "nw": cand["nw"], "nseg": cand["dnseg"]} + if cand["bfly"]: + k["bfly"] = cand["bfly"] + if cand["gt"]: + k["gt"] = 1 + return k + k = { + "nw": cand["nw"], + "rg": cand["rg"], + "minb": cand["minb"], + "nseg": max(1, cand["target"] // (hkv * cand["rg"])), + } + if cand["dspl"]: + k["dspl"] = cand["dspl"] + if cand["rspl"] > 1: + k["rspl"] = cand["rspl"] + if cand["pf"]: + k["pf"] = 1 + if cand.get("cpub"): + k["cpub"] = 1 + if cand.get("vinlds"): + k["vinlds"] = 1 + return k + + +def main() -> None: + p = argparse.ArgumentParser() + shapeset.add_arguments(p) + p.add_argument("--m", type=int, nargs="+", default=[1, 4]) + p.add_argument("--block-size", type=int, default=16) + p.add_argument("--rounds", type=int, default=1) + p.add_argument("--contexts", type=int, nargs="+", default=CONTEXTS) + shapeset.add_dtype_argument(p) + args = p.parse_args() + dtype = shapeset.torch_dtype(args.dtype) + + from common import BenchmarkConfig + from runner import run_attention_benchmark + + import vllm.v1.attention.backends.rocm_attn as backend_mod + from vllm.v1.attention.ops.rdna35_hip_decode import ( + KernelVariant, + load, + precompile, + seal, + ) + from vllm.v1.attention.ops.rdna35_hip_decode import knobs_for as _knobs_for + + shapes, _windowed = shapeset.load(args) + configs = sorted({(s.hq, s.hkv, s.d, s.window) for s in shapes if s.d != 96}) + if not configs: + print(f"no shapes match {shapeset.describe(args) or 'the given filters'}") + return + + scratch: dict[Any, Any] = {} + orig_make = backend_mod.make_scratch + + def cached(variant, device): + if variant not in scratch: + scratch[variant] = orig_make(variant, device) + return scratch[variant] + + backend_mod.make_scratch = cached + + original = backend_mod.RocmAttentionRdna35Impl._prepare + override: dict[str, int] = {} + + # The override replaces the whole knob set rather than patching what + # _prepare chose, so nothing of the base leaks into a point that did not + # ask for it -- dspl in particular, where absent means the kernel's rule. + def patched(self, kv_cache, **kw): + built = original(self, kv_cache, **kw) + if built is not None and override: + v = self._variant + want = KernelVariant( + v.head_size, + v.num_q_heads, + v.num_kv_heads, + v.max_m, + v.block_size, + v.layout, + **override, + dtype=v.dtype, + window=v.window, + ) + if self._variant != want: + self._variant = want + self._built = (load(want), cached(want, kw["q"].device)) + built = self._built + return built + + backend_mod.RocmAttentionRdna35Impl._prepare = patched + + def timeit(hq, hkv, d, m, s, window): + cfg = BenchmarkConfig( + backend="ROCM_ATTN", + batch_spec=f"q{m}s{s}", + num_layers=10, + min_working_set_mb=96, + head_dim=d, + num_q_heads=hq, + num_kv_heads=hkv, + block_size=args.block_size, + device="cuda:0", + dtype=dtype, + sliding_window=window or None, + ) + return run_attention_benchmark(cfg).median_time * 1e6 + + def key(c): + return tuple(sorted(c.items())) + + for hq, hkv, d, window in configs: + for m in args.m: + base = dict(_knobs_for(hq, hkv, d, m, window, dtype)) + start = { + "nw": base.get("nw", 8), + "dspl": base.get("dspl", 0), + "rg": base.get("rg", 1), + "rspl": base.get("rspl", 1), + "pf": base.get("pf", 0), + "cpub": base.get("cpub", 0), + "vinlds": base.get("vinlds", 0), + "target": 16, + "minb": base.get("minb", 1), + } + values = space(hq, hkv, d, start) + roof = { + s: shapeset.roofline_us( + hq, hkv, d, m, s, block_size=args.block_size, window=window + ) + for s in args.contexts + } + scores: dict[tuple, tuple[float, list[float]]] = {} + + def evaluate( + cands, + ctxs, + hq=hq, + hkv=hkv, + d=d, + m=m, + window=window, + base=base, + roof=roof, + scores=scores, + ): + todo = [c for c in cands if (ctxs, key(c)) not in scores] + variants = {} + for c in todo: + with contextlib.suppress(Exception): + variants[key(c)] = KernelVariant( + d, + hq, + hkv, + m, + args.block_size, + 1, + **knobs_of(hkv, c), + dtype=dtype, + window=window, + ) + # The unpatched _prepare builds the backend's own choice before + # the override replaces it, so that variant is needed too. + seal(False) + precompile( + [ + *variants.values(), + KernelVariant( + d, + hq, + hkv, + m, + args.block_size, + 1, + **base, + dtype=dtype, + window=window, + ), + ] + ) + seal() + time.sleep(5) # clocks settle after a parallel build + for c in todo: + if key(c) not in variants: + scores[(ctxs, key(c))] = (0.0, []) + continue + override.clear() + override.update(knobs_of(hkv, c)) + try: + r = [roof[s] / timeit(hq, hkv, d, m, s, window) for s in ctxs] + except Exception: # noqa: BLE001 - a point may not build + scores[(ctxs, key(c))] = (0.0, []) + continue + scores[(ctxs, key(c))] = (geomean(r), r) + override.clear() + + # A second start with every row tile split inside one workgroup: + # reaching it from the backend's rg needs several knobs to move + # at once, which one-knob descent does not do. + gqa_m = hq // hkv * m + rt = -(-gqa_m // 16) + rs_max = max(r for r in (1, 2, 4) if rt % r == 0) + starts = [start] + if rs_max > 1: + starts.append(dict(start, rg=1, rspl=rs_max, nw=4, target=32)) + + # The dot decomposition, a start of its own: it shares no knob + # with WMMA but the waves. Scored like any other point, over the + # whole context range. + dot_start = dict(start, dot=1, nw=8, dnseg=2, bfly=0, gt=0) + + def descend(ctxs, starts=starts, dot_start=dot_start, dot=False): + found = [descend_from(ctxs, s) for s in starts] + if dot: + found.append(descend_from(ctxs, dot_start)) + return max(found, key=lambda f: rank(ctxs, f[1])) + + def descend_from(ctxs, start, values=values, scores=scores): + best = dict(start) + evaluate([best], ctxs) + for _ in range(args.rounds): + knobs_ = ( + ("nw", "dnseg", "bfly", "gt") + if best.get("dot") + else ( + "nw", + "dspl", + "rg", + "rspl", + "pf", + "cpub", + "vinlds", + "target", + "minb", + ) + ) + for knob in knobs_: + cands = [dict(best, **{knob: v}) for v in values[knob]] + if knob == "rspl": + # rspl splits rows inside the workgroup that rg + # splits across workgroups; trade one for the other. + cands += [ + dict(best, rspl=v, rg=best["rg"] // v) + for v in values["rspl"] + if v > 1 and best["rg"] % v == 0 + ] + evaluate(cands, ctxs) + for c in cands: + if rank(ctxs, scores[(ctxs, key(c))]) > rank( + ctxs, scores[(ctxs, key(best))] + ): + best = c + return best, scores[(ctxs, key(best))] + + best, (g, cells) = descend(tuple(args.contexts), dot=True) + row = ", ".join(f'"{k}": {v}' for k, v in knobs_of(hkv, best).items()) + worst = min(cells) * 100 if cells else 0.0 + key_ = f"{hq}, {hkv}, {d}, {m}" + (f", {window}" if window else "") + print( + f" ({key_}): {{{row}}}, " + f"# {g * 100:.1f} % geomean, worst {worst:.1f} %", + flush=True, + ) + + +if __name__ == "__main__": + main() diff --git a/cmake/rdna35_attn.cmake b/cmake/rdna35_attn.cmake new file mode 100644 index 000000000000..b568cd0967ad --- /dev/null +++ b/cmake/rdna35_attn.cmake @@ -0,0 +1,55 @@ +# +# Sources of the RDNA3.5 decode-attention variants built into _rocm_C. +# +# csrc/rocm/rdna35_attn/variants.def is the one list of variants. At +# configure time it is split by group into one generated unit per group, +# each rewritten only when its lines change, and CMake re-runs whenever the +# list changes. So editing one group's lines rebuilds that group's unit and +# the registry (which includes the whole list), and leaves every other unit +# untouched. +# +# Sets ${OUT_SRCS} to the registry plus the generated units. +# +function(rdna35_attn_sources OUT_SRCS) + set(_DIR ${CMAKE_SOURCE_DIR}/csrc/rocm/rdna35_attn) + set(_DEF ${_DIR}/variants.def) + set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS ${_DEF}) + + file(STRINGS ${_DEF} _LINES REGEX "^RDNA35_VARIANT\\(") + set(_GROUPS) + foreach(_LINE IN LISTS _LINES) + if(NOT _LINE MATCHES "^RDNA35_VARIANT\\(([A-Za-z0-9_]+),") + message(FATAL_ERROR "variants.def: no group in '${_LINE}'") + endif() + set(_G ${CMAKE_MATCH_1}) + list(APPEND _GROUPS ${_G}) + string(APPEND _BODY_${_G} + "#define RDNA35_V ${_LINE}\n" + "#include \"rdna35_attn/instance.h\"\n" + "#undef RDNA35_V\n") + endforeach() + list(REMOVE_DUPLICATES _GROUPS) + list(LENGTH _LINES _NV) + list(LENGTH _GROUPS _NG) + message(STATUS "RDNA35 decode attention: ${_NV} variants in ${_NG} units") + + set(_OUT ${CMAKE_CURRENT_BINARY_DIR}/rdna35_attn) + set(_SRCS ${_DIR}/registry.hip) + foreach(_G IN LISTS _GROUPS) + # file(CONFIGURE) leaves the file, and its timestamp, alone when the + # content is unchanged: that is what keeps the other groups built. + file(CONFIGURE OUTPUT ${_OUT}/${_G}.hip + CONTENT "// Generated from variants.def, group ${_G}. Do not edit.\n${_BODY_${_G}}" + @ONLY) + list(APPEND _SRCS ${_OUT}/${_G}.hip) + endforeach() + + set_source_files_properties(${_SRCS} PROPERTIES + INCLUDE_DIRECTORIES ${CMAKE_SOURCE_DIR}/csrc/rocm + # The optimisation level the kernels were tuned at, whatever the build + # type; -g0 keeps RelWithDebInfo from carrying ~1000 kernels' debug info. + # The kernel keeps values some variants leave unused rather than #if them + # out per knob, which _rocm_C's -Werror=unused-variable would refuse. + COMPILE_OPTIONS "-O3;-g0;-Wno-unused-variable") + set(${OUT_SRCS} ${_SRCS} PARENT_SCOPE) +endfunction() diff --git a/csrc/rocm/ops.h b/csrc/rocm/ops.h index bbd1c621119f..a05d59769f11 100644 --- a/csrc/rocm/ops.h +++ b/csrc/rocm/ops.h @@ -142,6 +142,15 @@ void gdn_chunked(torch::Tensor& q, torch::Tensor& k, torch::Tensor& v, torch::Tensor& cu_seqlens, torch::Tensor& out, torch::Tensor& final_state, double scale); +// RDNA3.5 decode attention, the variants of rdna35_attn/variants.def: the +// index of a variant's key (-1 if not built), and a launch by that index. +int64_t rdna35_decode_variant(torch::IntArrayRef key); +void rdna35_decode_attn(int64_t variant, torch::Tensor& q, + torch::Tensor& kv_cache, torch::Tensor& block_table, + torch::Tensor& out, torch::Tensor& acc, + torch::Tensor& m, torch::Tensor& l, torch::Tensor& cnt, + torch::Tensor& seq_lens, double scale); + void paged_attention( torch::Tensor& out, torch::Tensor& exp_sums, torch::Tensor& max_logits, torch::Tensor& tmp_out, torch::Tensor& query, torch::Tensor& key_cache, diff --git a/csrc/rocm/rdna35_attn/instance.h b/csrc/rocm/rdna35_attn/instance.h new file mode 100644 index 000000000000..86e7f3c590d3 --- /dev/null +++ b/csrc/rocm/rdna35_attn/instance.h @@ -0,0 +1,77 @@ +// One decode-attention variant, built into _rocm_C. +// +// Included once per variant by the generated group units, after +// +// #define RDNA35_V RDNA35_VARIANT() +// +// No include guard: every inclusion is another variant. The kernel's defines +// read the entry's fields, its code goes into a namespace named after the +// entry, and rdna35::launch_ -- declared for the registry by variants.h +// -- launches it. Everything defined here is undefined at the end, the +// kernel's own definitions by the kernel (RDNA35_VLLM). +#include "rdna35_attn/prelude.h" + +#define RDNA35_VARIANT(...) (__VA_ARGS__) + +#define HEAD_DIM RDNA35_FIELD(1, RDNA35_V) +#define NUM_Q_HEADS RDNA35_FIELD(2, RDNA35_V) +#define NUM_KV_HEADS RDNA35_FIELD(3, RDNA35_V) +#define MAXM RDNA35_FIELD(4, RDNA35_V) +#define BS RDNA35_FIELD(5, RDNA35_V) +#define LAYOUT RDNA35_FIELD(6, RDNA35_V) +#define NSEG RDNA35_FIELD(7, RDNA35_V) +#define RG RDNA35_FIELD(8, RDNA35_V) +#define MINB RDNA35_FIELD(9, RDNA35_V) +#define NW RDNA35_FIELD(10, RDNA35_V) +#define DSPL RDNA35_FIELD(11, RDNA35_V) +#define RSPL RDNA35_FIELD(12, RDNA35_V) +#define PF RDNA35_FIELD(13, RDNA35_V) +#define CPUB RDNA35_FIELD(14, RDNA35_V) +#define VINLDS RDNA35_FIELD(15, RDNA35_V) +#define BATCH RDNA35_FIELD(16, RDNA35_V) +#define DOT RDNA35_FIELD(17, RDNA35_V) +#define BFLY RDNA35_FIELD(18, RDNA35_V) +#define GT RDNA35_FIELD(19, RDNA35_V) +#define WIN RDNA35_FIELD(20, RDNA35_V) +#define KV_BF16 RDNA35_FIELD(21, RDNA35_V) +#define MUTATE 0 +#define ABLATE 0 +#define TIMING 0 + +namespace rdna35 { +namespace RDNA35_CALL(RDNA35_ID, RDNA35_PREPEND(v, RDNA35_V)) { +#include "rdna35_decode_attn.cu" +} // namespace RDNA35_CALL(RDNA35_ID,RDNA35_PREPEND(v,RDNA35_V)) + +void RDNA35_CALL(RDNA35_ID, + RDNA35_PREPEND(launch, RDNA35_V))(const LaunchArgs& a) { + RDNA35_CALL(RDNA35_ID, RDNA35_PREPEND(v, RDNA35_V))::launch(a); +} +} // namespace rdna35 + +#undef RDNA35_VARIANT +#undef HEAD_DIM +#undef NUM_Q_HEADS +#undef NUM_KV_HEADS +#undef MAXM +#undef BS +#undef LAYOUT +#undef NSEG +#undef RG +#undef MINB +#undef NW +#undef DSPL +#undef RSPL +#undef PF +#undef CPUB +#undef VINLDS +#undef BATCH +#undef DOT +#undef BFLY +#undef GT +#undef WIN +#undef KV_BF16 +#undef MUTATE +#undef ABLATE +#undef TIMING +#undef BTARGET diff --git a/csrc/rocm/rdna35_attn/launch.h b/csrc/rocm/rdna35_attn/launch.h new file mode 100644 index 000000000000..94d84fc83034 --- /dev/null +++ b/csrc/rocm/rdna35_attn/launch.h @@ -0,0 +1,35 @@ +// Launch ABI of the RDNA3.5 decode-attention variants. +// +// Every variant (rdna35_decode_attn.cu, specialised by compile-time defines) +// exposes one host function of type LaunchFn. The kernel translation units +// see only this header, not the list of variants, so a change to the list +// rebuilds only the units whose variants changed. +#pragma once + +#include + +namespace rdna35 { + +// Pointers are raw so the kernel units need no torch headers: a unit then +// compiles in about the time of its kernels, not of torch/extension.h. +struct LaunchArgs { + const void* q; + const void* kv; + const int* bt; + float* acc; + float* m; + float* l; + int* cnt; + void* out; + const int* seq_lens; + int bt_width; + // Per-sequence strides of a batch launch; 0 for one sequence. + int bt_stride, acc_stride, ml_stride, cnt_stride; + int nseq; + float scale; + hipStream_t stream; +}; + +using LaunchFn = void (*)(const LaunchArgs&); + +} // namespace rdna35 diff --git a/csrc/rocm/rdna35_attn/macros.h b/csrc/rocm/rdna35_attn/macros.h new file mode 100644 index 000000000000..79b15aa73b20 --- /dev/null +++ b/csrc/rocm/rdna35_attn/macros.h @@ -0,0 +1,60 @@ +// Helpers for the X-macro list of decode-attention variants (variants.def). +// +// An entry is +// +// RDNA35_VARIANT(group, HEAD_DIM, NUM_Q_HEADS, NUM_KV_HEADS, MAXM, BS, +// LAYOUT, NSEG, RG, MINB, NW, DSPL, RSPL, PF, CPUB, VINLDS, +// BATCH, DOT, BFLY, GT, WIN, KV_BF16) +// +// with every field an integer literal, exactly the defines +// rdna35_decode_attn.cu is built with (DSPL 0 is the kernel's own rule). +// `group` names the translation unit the variant is compiled in. +#pragma once + +#define RDNA35_STRIP(...) __VA_ARGS__ +// m applied to the fields of a parenthesised entry: RDNA35_CALL(m, (a, b)) +// is m(a, b), with the entry expanded before m collects its arguments. +#define RDNA35_CALL(m, args) m args +#define RDNA35_PREPEND(p, entry) (p, RDNA35_STRIP entry) + +// A C identifier unique to the variant, from every field but the group: +// RDNA35_ID(launch, g, 128, 16, 2, ...) is launch_d128_q16_kv2_... +#define RDNA35_ID(p, g, d, hq, hkv, m, bs, l, n, rg, mb, nw, ds, rs, pf, cp, \ + vl, b, dot, bf, gt, win, t) \ + p##_d##d##_q##hq##_kv##hkv##_m##m##_bs##bs##_l##l##_n##n##_rg##rg##_mb##mb##_w##nw##_ds##ds##_rs##rs##_pf##pf##_cp##cp##_vl##vl##_b##b##_dot##dot##_bf##bf##_gt##gt##_win##win##_t##t + +// The key a variant is looked up by: every field but the group, in order. +#define RDNA35_KEY(g, ...) __VA_ARGS__ +#define RDNA35_NFIELDS 21 + +// Field i (1-based, after the group) of a parenthesised entry. +#define RDNA35_FIELD(i, entry) RDNA35_CALL(RDNA35_F##i, entry) +#define RDNA35_F1(g, a, ...) a +#define RDNA35_F2(g, a, b, ...) b +#define RDNA35_F3(g, a, b, c, ...) c +#define RDNA35_F4(g, a, b, c, d, ...) d +#define RDNA35_F5(g, a, b, c, d, e, ...) e +#define RDNA35_F6(g, a, b, c, d, e, f, ...) f +#define RDNA35_F7(g, a, b, c, d, e, f, h, ...) h +#define RDNA35_F8(g, a, b, c, d, e, f, h, i, ...) i +#define RDNA35_F9(g, a, b, c, d, e, f, h, i, j, ...) j +#define RDNA35_F10(g, a, b, c, d, e, f, h, i, j, k, ...) k +#define RDNA35_F11(g, a, b, c, d, e, f, h, i, j, k, l, ...) l +#define RDNA35_F12(g, a, b, c, d, e, f, h, i, j, k, l, m, ...) m +#define RDNA35_F13(g, a, b, c, d, e, f, h, i, j, k, l, m, n, ...) n +#define RDNA35_F14(g, a, b, c, d, e, f, h, i, j, k, l, m, n, o, ...) o +#define RDNA35_F15(g, a, b, c, d, e, f, h, i, j, k, l, m, n, o, p, ...) p +#define RDNA35_F16(g, a, b, c, d, e, f, h, i, j, k, l, m, n, o, p, q, ...) q +#define RDNA35_F17(g, a, b, c, d, e, f, h, i, j, k, l, m, n, o, p, q, r, ...) r +#define RDNA35_F18(g, a, b, c, d, e, f, h, i, j, k, l, m, n, o, p, q, r, s, \ + ...) \ + s +#define RDNA35_F19(g, a, b, c, d, e, f, h, i, j, k, l, m, n, o, p, q, r, s, t, \ + ...) \ + t +#define RDNA35_F20(g, a, b, c, d, e, f, h, i, j, k, l, m, n, o, p, q, r, s, t, \ + u, ...) \ + u +#define RDNA35_F21(g, a, b, c, d, e, f, h, i, j, k, l, m, n, o, p, q, r, s, t, \ + u, v) \ + v diff --git a/csrc/rocm/rdna35_attn/prelude.h b/csrc/rocm/rdna35_attn/prelude.h new file mode 100644 index 000000000000..2f51890e8142 --- /dev/null +++ b/csrc/rocm/rdna35_attn/prelude.h @@ -0,0 +1,15 @@ +// What every variant of a group unit shares, included before the first one +// so that the kernel's own includes find their guards set and add nothing to +// the variant namespaces. +#pragma once + +#include + +#include +#include +#include + +#include "rdna35_attn/launch.h" +#include "rdna35_attn/macros.h" + +#define RDNA35_VLLM 1 diff --git a/csrc/rocm/rdna35_attn/registry.hip b/csrc/rocm/rdna35_attn/registry.hip new file mode 100644 index 000000000000..c5ff38dfc257 --- /dev/null +++ b/csrc/rocm/rdna35_attn/registry.hip @@ -0,0 +1,112 @@ +// Torch ops over the decode-attention variants built into _rocm_C. +// +// rdna35_decode_variant maps a variant's key (the fields of its variants.def +// line) to an index, once, when the backend picks a variant; +// rdna35_decode_attn launches by that index. The checks are the JIT op's +// (decode_attn_op in rdna35_decode_attn.cu), read from the key instead of +// from the defines. +#include +#include +#include + +#include +#include + +#include "rdna35_attn/variants.h" + +namespace { + +enum Field { + kHeadDim, + kNumQHeads, + kNumKvHeads, + kMaxM, + kBlockSize, + kLayout, + kNseg, + kRg, + kMinb, + kNw, + kDspl, + kRspl, + kPf, + kCpub, + kVinlds, + kBatch, + kDot, + kBfly, + kGt, + kWin, + kBf16, +}; +static_assert(kBf16 + 1 == RDNA35_NFIELDS, "one Field per key field"); + +constexpr int64_t kNumVariants = + sizeof(rdna35::kVariants) / sizeof(rdna35::kVariants[0]); + +} // namespace + +int64_t rdna35_decode_variant(torch::IntArrayRef key) { + TORCH_CHECK(key.size() == RDNA35_NFIELDS, "a variant key has ", + RDNA35_NFIELDS, " fields, got ", key.size()); + // The kernels carry code for gfx1151 only (rdna35_decode_attn.cu). + int dev; + hipDeviceProp_t prop; + if (hipGetDevice(&dev) != hipSuccess || + hipGetDeviceProperties(&prop, dev) != hipSuccess || + !strstr(prop.gcnArchName, "gfx1151")) + return -1; + for (int64_t i = 0; i < kNumVariants; ++i) { + const int* k = rdna35::kVariants[i].key; + if (std::equal(key.begin(), key.end(), k)) return i; + } + return -1; +} + +void rdna35_decode_attn(int64_t variant, torch::Tensor& q, + torch::Tensor& kv_cache, torch::Tensor& block_table, + torch::Tensor& out, torch::Tensor& acc, + torch::Tensor& m, torch::Tensor& l, torch::Tensor& cnt, + torch::Tensor& seq_lens, double scale) { + TORCH_CHECK(variant >= 0 && variant < kNumVariants, "no variant ", variant); + const rdna35::Variant& v = rdna35::kVariants[variant]; + const int maxm = v.key[kMaxM]; + TORCH_CHECK(q.is_contiguous() && out.is_contiguous(), + "q and out must be contiguous"); + TORCH_CHECK(block_table.scalar_type() == torch::kInt32 && + (block_table.dim() == 1 || block_table.dim() == 2) && + block_table.stride(-1) == 1, + "block table must be int32 rows, contiguous within a row"); + TORCH_CHECK(seq_lens.scalar_type() == torch::kInt32 && seq_lens.numel() >= 1, + "seq_lens must be an int32 tensor holding the sequence length"); + TORCH_CHECK(q.size(0) % maxm == 0 && q.size(0) > 0, "q has ", q.size(0), + " tokens, not a multiple of MAXM=", maxm); + const int nseq = q.size(0) / maxm; + TORCH_CHECK(seq_lens.numel() >= nseq, "seq_lens has ", seq_lens.numel(), + " entries for ", nseq, " sequences"); + TORCH_CHECK(block_table.dim() == 2 ? block_table.size(0) >= nseq : nseq == 1, + "block table has too few rows for ", nseq, " sequences"); + const bool batched = acc.dim() == 3; + TORCH_CHECK(batched ? acc.size(0) >= nseq && m.size(0) >= nseq && + l.size(0) >= nseq && cnt.size(0) >= nseq + : nseq == 1, + "scratch holds fewer sequences than the ", nseq, " launched"); + TORCH_CHECK(v.key[kBatch] || nseq == 1, "one-sequence build launched for ", + nseq, " sequences"); + TORCH_CHECK(q.size(1) == v.key[kNumQHeads] && q.size(2) == v.key[kHeadDim], + "q shape does not match the compiled variant"); + const auto dtype = v.key[kBf16] ? at::kBFloat16 : at::kHalf; + TORCH_CHECK(q.scalar_type() == dtype && kv_cache.scalar_type() == dtype && + out.scalar_type() == dtype, + "kernel built for ", dtype, ", got q=", q.scalar_type(), + " kv_cache=", kv_cache.scalar_type(), " out=", out.scalar_type()); + const at::cuda::OptionalCUDAGuard device_guard(device_of(q)); + v.fn({q.data_ptr(), kv_cache.data_ptr(), block_table.data_ptr(), + acc.data_ptr(), m.data_ptr(), l.data_ptr(), + cnt.data_ptr(), out.data_ptr(), seq_lens.data_ptr(), + (int)block_table.size(-1), + block_table.dim() == 2 ? (int)block_table.stride(0) : 0, + batched ? (int)acc.stride(0) : 0, batched ? (int)m.stride(0) : 0, + batched ? (int)cnt.stride(0) : 0, nseq, (float)scale, + at::cuda::getCurrentCUDAStream()}); +} diff --git a/csrc/rocm/rdna35_attn/variants.def b/csrc/rocm/rdna35_attn/variants.def new file mode 100644 index 000000000000..bdd348c01f73 --- /dev/null +++ b/csrc/rocm/rdna35_attn/variants.def @@ -0,0 +1,1284 @@ +// Decode-attention variants compiled into _rocm_C (see macros.h for +// the fields). Generated by write_variants_def from the backend's +// tables; regenerate rather than edit: +// python -m vllm.v1.attention.ops.rdna35_hip_decode +RDNA35_VARIANT(q10_kv10_d128_w0_bf16, 128, 10, 10, 1, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q10_kv10_d128_w0_bf16, 128, 10, 10, 1, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q10_kv10_d128_w0_bf16, 128, 10, 10, 2, 16, 1, 1, 1, 4, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q10_kv10_d128_w0_bf16, 128, 10, 10, 2, 16, 1, 1, 1, 4, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q10_kv10_d128_w0_bf16, 128, 10, 10, 3, 16, 1, 1, 1, 2, 8, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q10_kv10_d128_w0_bf16, 128, 10, 10, 3, 16, 1, 1, 1, 2, 8, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q10_kv10_d128_w0_bf16, 128, 10, 10, 4, 16, 1, 1, 1, 2, 2, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q10_kv10_d128_w0_bf16, 128, 10, 10, 4, 16, 1, 1, 1, 2, 2, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q10_kv10_d128_w0_bf16, 128, 10, 10, 5, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q10_kv10_d128_w0_bf16, 128, 10, 10, 5, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q10_kv10_d128_w0_bf16, 128, 10, 10, 6, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q10_kv10_d128_w0_bf16, 128, 10, 10, 6, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q10_kv10_d128_w0_bf16, 128, 10, 10, 7, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q10_kv10_d128_w0_bf16, 128, 10, 10, 7, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q10_kv10_d128_w0_bf16, 128, 10, 10, 8, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q10_kv10_d128_w0_bf16, 128, 10, 10, 8, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q10_kv10_d128_w0_fp16, 128, 10, 10, 1, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q10_kv10_d128_w0_fp16, 128, 10, 10, 1, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q10_kv10_d128_w0_fp16, 128, 10, 10, 2, 16, 1, 1, 1, 4, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q10_kv10_d128_w0_fp16, 128, 10, 10, 2, 16, 1, 1, 1, 4, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q10_kv10_d128_w0_fp16, 128, 10, 10, 3, 16, 1, 1, 1, 2, 8, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q10_kv10_d128_w0_fp16, 128, 10, 10, 3, 16, 1, 1, 1, 2, 8, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q10_kv10_d128_w0_fp16, 128, 10, 10, 4, 16, 1, 1, 1, 2, 2, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q10_kv10_d128_w0_fp16, 128, 10, 10, 4, 16, 1, 1, 1, 2, 2, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q10_kv10_d128_w0_fp16, 128, 10, 10, 5, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q10_kv10_d128_w0_fp16, 128, 10, 10, 5, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q10_kv10_d128_w0_fp16, 128, 10, 10, 6, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q10_kv10_d128_w0_fp16, 128, 10, 10, 6, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q10_kv10_d128_w0_fp16, 128, 10, 10, 7, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q10_kv10_d128_w0_fp16, 128, 10, 10, 7, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q10_kv10_d128_w0_fp16, 128, 10, 10, 8, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q10_kv10_d128_w0_fp16, 128, 10, 10, 8, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q14_kv2_d64_w0_bf16, 64, 14, 2, 1, 16, 1, 4, 1, 2, 4, 0, 1, 1, 1, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q14_kv2_d64_w0_bf16, 64, 14, 2, 1, 16, 1, 4, 1, 2, 4, 0, 1, 1, 1, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q14_kv2_d64_w0_bf16, 64, 14, 2, 2, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q14_kv2_d64_w0_bf16, 64, 14, 2, 2, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q14_kv2_d64_w0_bf16, 64, 14, 2, 3, 16, 1, 8, 1, 2, 4, 0, 2, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q14_kv2_d64_w0_bf16, 64, 14, 2, 3, 16, 1, 8, 1, 2, 4, 0, 2, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q14_kv2_d64_w0_bf16, 64, 14, 2, 4, 16, 1, 4, 1, 2, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q14_kv2_d64_w0_bf16, 64, 14, 2, 4, 16, 1, 4, 1, 2, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q14_kv2_d64_w0_bf16, 64, 14, 2, 5, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q14_kv2_d64_w0_bf16, 64, 14, 2, 5, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q14_kv2_d64_w0_bf16, 64, 14, 2, 6, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q14_kv2_d64_w0_bf16, 64, 14, 2, 6, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q14_kv2_d64_w0_bf16, 64, 14, 2, 7, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q14_kv2_d64_w0_bf16, 64, 14, 2, 7, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q14_kv2_d64_w0_bf16, 64, 14, 2, 8, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q14_kv2_d64_w0_bf16, 64, 14, 2, 8, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q14_kv2_d64_w0_fp16, 64, 14, 2, 1, 16, 1, 4, 1, 2, 4, 0, 1, 1, 1, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q14_kv2_d64_w0_fp16, 64, 14, 2, 1, 16, 1, 4, 1, 2, 4, 0, 1, 1, 1, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q14_kv2_d64_w0_fp16, 64, 14, 2, 2, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q14_kv2_d64_w0_fp16, 64, 14, 2, 2, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q14_kv2_d64_w0_fp16, 64, 14, 2, 3, 16, 1, 8, 1, 2, 4, 0, 2, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q14_kv2_d64_w0_fp16, 64, 14, 2, 3, 16, 1, 8, 1, 2, 4, 0, 2, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q14_kv2_d64_w0_fp16, 64, 14, 2, 4, 16, 1, 4, 1, 2, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q14_kv2_d64_w0_fp16, 64, 14, 2, 4, 16, 1, 4, 1, 2, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q14_kv2_d64_w0_fp16, 64, 14, 2, 5, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q14_kv2_d64_w0_fp16, 64, 14, 2, 5, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q14_kv2_d64_w0_fp16, 64, 14, 2, 6, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q14_kv2_d64_w0_fp16, 64, 14, 2, 6, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q14_kv2_d64_w0_fp16, 64, 14, 2, 7, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q14_kv2_d64_w0_fp16, 64, 14, 2, 7, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q14_kv2_d64_w0_fp16, 64, 14, 2, 8, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q14_kv2_d64_w0_fp16, 64, 14, 2, 8, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv1_d512_w0_bf16, 512, 16, 1, 1, 16, 1, 16, 1, 1, 8, 8, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv1_d512_w0_bf16, 512, 16, 1, 1, 16, 1, 16, 1, 1, 8, 8, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv1_d512_w0_bf16, 512, 16, 1, 2, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv1_d512_w0_bf16, 512, 16, 1, 2, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv1_d512_w0_bf16, 512, 16, 1, 3, 16, 1, 8, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv1_d512_w0_bf16, 512, 16, 1, 3, 16, 1, 8, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv1_d512_w0_bf16, 512, 16, 1, 4, 16, 1, 8, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv1_d512_w0_bf16, 512, 16, 1, 4, 16, 1, 8, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv1_d512_w0_bf16, 512, 16, 1, 5, 16, 1, 2, 8, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv1_d512_w0_bf16, 512, 16, 1, 5, 16, 1, 2, 8, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv1_d512_w0_bf16, 512, 16, 1, 6, 16, 1, 2, 8, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv1_d512_w0_bf16, 512, 16, 1, 6, 16, 1, 2, 8, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv1_d512_w0_bf16, 512, 16, 1, 7, 16, 1, 2, 8, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv1_d512_w0_bf16, 512, 16, 1, 7, 16, 1, 2, 8, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv1_d512_w0_bf16, 512, 16, 1, 8, 16, 1, 2, 8, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv1_d512_w0_bf16, 512, 16, 1, 8, 16, 1, 2, 8, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv1_d512_w0_fp16, 512, 16, 1, 1, 16, 1, 16, 1, 1, 8, 8, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv1_d512_w0_fp16, 512, 16, 1, 1, 16, 1, 16, 1, 1, 8, 8, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv1_d512_w0_fp16, 512, 16, 1, 2, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv1_d512_w0_fp16, 512, 16, 1, 2, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv1_d512_w0_fp16, 512, 16, 1, 3, 16, 1, 8, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv1_d512_w0_fp16, 512, 16, 1, 3, 16, 1, 8, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv1_d512_w0_fp16, 512, 16, 1, 4, 16, 1, 8, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv1_d512_w0_fp16, 512, 16, 1, 4, 16, 1, 8, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv1_d512_w0_fp16, 512, 16, 1, 5, 16, 1, 2, 8, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv1_d512_w0_fp16, 512, 16, 1, 5, 16, 1, 2, 8, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv1_d512_w0_fp16, 512, 16, 1, 6, 16, 1, 2, 8, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv1_d512_w0_fp16, 512, 16, 1, 6, 16, 1, 2, 8, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv1_d512_w0_fp16, 512, 16, 1, 7, 16, 1, 2, 8, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv1_d512_w0_fp16, 512, 16, 1, 7, 16, 1, 2, 8, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv1_d512_w0_fp16, 512, 16, 1, 8, 16, 1, 2, 8, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv1_d512_w0_fp16, 512, 16, 1, 8, 16, 1, 2, 8, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d128_w0_bf16, 128, 16, 2, 1, 16, 1, 4, 2, 2, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d128_w0_bf16, 128, 16, 2, 1, 16, 1, 4, 2, 2, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d128_w0_bf16, 128, 16, 2, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d128_w0_bf16, 128, 16, 2, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d128_w0_bf16, 128, 16, 2, 3, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d128_w0_bf16, 128, 16, 2, 3, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d128_w0_bf16, 128, 16, 2, 4, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d128_w0_bf16, 128, 16, 2, 4, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d128_w0_bf16, 128, 16, 2, 5, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d128_w0_bf16, 128, 16, 2, 5, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d128_w0_bf16, 128, 16, 2, 6, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d128_w0_bf16, 128, 16, 2, 6, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d128_w0_bf16, 128, 16, 2, 7, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d128_w0_bf16, 128, 16, 2, 7, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d128_w0_bf16, 128, 16, 2, 8, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d128_w0_bf16, 128, 16, 2, 8, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d128_w0_fp16, 128, 16, 2, 1, 16, 1, 4, 2, 2, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d128_w0_fp16, 128, 16, 2, 1, 16, 1, 4, 2, 2, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d128_w0_fp16, 128, 16, 2, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d128_w0_fp16, 128, 16, 2, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d128_w0_fp16, 128, 16, 2, 3, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d128_w0_fp16, 128, 16, 2, 3, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d128_w0_fp16, 128, 16, 2, 4, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d128_w0_fp16, 128, 16, 2, 4, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d128_w0_fp16, 128, 16, 2, 5, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d128_w0_fp16, 128, 16, 2, 5, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d128_w0_fp16, 128, 16, 2, 6, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d128_w0_fp16, 128, 16, 2, 6, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d128_w0_fp16, 128, 16, 2, 7, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d128_w0_fp16, 128, 16, 2, 7, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d128_w0_fp16, 128, 16, 2, 8, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d128_w0_fp16, 128, 16, 2, 8, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 1, 1056, 1, 8, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 1, 1056, 1, 8, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 1, 16, 1, 8, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 1, 16, 1, 8, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 2, 1056, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 2, 1056, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 3, 1056, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 3, 1056, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 3, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 3, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 4, 1056, 1, 8, 2, 1, 4, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 4, 1056, 1, 8, 2, 1, 4, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 4, 16, 1, 8, 2, 1, 4, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 4, 16, 1, 8, 2, 1, 4, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 5, 1056, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 5, 1056, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 5, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 5, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 6, 1056, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 6, 1056, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 6, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 6, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 7, 1056, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 7, 1056, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 7, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 7, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 8, 1056, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 8, 1056, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 8, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_bf16, 256, 16, 2, 8, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 1, 1056, 1, 8, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 1, 1056, 1, 8, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 1, 16, 1, 8, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 1, 16, 1, 8, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 2, 1056, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 2, 1056, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 3, 1056, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 3, 1056, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 3, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 3, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 4, 1056, 1, 8, 2, 1, 4, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 4, 1056, 1, 8, 2, 1, 4, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 4, 16, 1, 8, 2, 1, 4, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 4, 16, 1, 8, 2, 1, 4, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 5, 1056, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 5, 1056, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 5, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 5, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 6, 1056, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 6, 1056, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 6, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 6, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 7, 1056, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 7, 1056, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 7, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 7, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 8, 1056, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 8, 1056, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 8, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d256_w0_fp16, 256, 16, 2, 8, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 1, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 1, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 1, 32, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 1, 32, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 2, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 2, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 2, 32, 1, 4, 2, 1, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 2, 32, 1, 4, 2, 1, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 3, 16, 1, 4, 2, 1, 8, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 3, 16, 1, 4, 2, 1, 8, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 3, 32, 1, 4, 2, 1, 8, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 3, 32, 1, 4, 2, 1, 8, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 4, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 4, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 4, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 4, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 5, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 5, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 5, 32, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 5, 32, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 6, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 6, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 6, 32, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 6, 32, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 7, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 7, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 7, 32, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 7, 32, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 8, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 8, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 8, 32, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_bf16, 512, 16, 2, 8, 32, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 1, 16, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 1, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 1, 32, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 1, 32, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 2, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 2, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 2, 32, 1, 4, 2, 1, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 2, 32, 1, 4, 2, 1, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 3, 16, 1, 4, 2, 1, 8, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 3, 16, 1, 4, 2, 1, 8, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 3, 32, 1, 4, 2, 1, 8, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 3, 32, 1, 4, 2, 1, 8, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 4, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 4, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 4, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 4, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 5, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 5, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 5, 32, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 5, 32, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 6, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 6, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 6, 32, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 6, 32, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 7, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 7, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 7, 32, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 7, 32, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 8, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 8, 16, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 8, 32, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d512_w0_fp16, 512, 16, 2, 8, 32, 1, 2, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d64_w0_bf16, 64, 16, 2, 1, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d64_w0_bf16, 64, 16, 2, 1, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d64_w0_bf16, 64, 16, 2, 2, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d64_w0_bf16, 64, 16, 2, 2, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d64_w0_bf16, 64, 16, 2, 3, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d64_w0_bf16, 64, 16, 2, 3, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d64_w0_bf16, 64, 16, 2, 4, 16, 1, 4, 2, 2, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d64_w0_bf16, 64, 16, 2, 4, 16, 1, 4, 2, 2, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d64_w0_bf16, 64, 16, 2, 5, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d64_w0_bf16, 64, 16, 2, 5, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d64_w0_bf16, 64, 16, 2, 6, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d64_w0_bf16, 64, 16, 2, 6, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d64_w0_bf16, 64, 16, 2, 7, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d64_w0_bf16, 64, 16, 2, 7, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d64_w0_bf16, 64, 16, 2, 8, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d64_w0_bf16, 64, 16, 2, 8, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv2_d64_w0_fp16, 64, 16, 2, 1, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d64_w0_fp16, 64, 16, 2, 1, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d64_w0_fp16, 64, 16, 2, 2, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d64_w0_fp16, 64, 16, 2, 2, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d64_w0_fp16, 64, 16, 2, 3, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d64_w0_fp16, 64, 16, 2, 3, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d64_w0_fp16, 64, 16, 2, 4, 16, 1, 4, 2, 2, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d64_w0_fp16, 64, 16, 2, 4, 16, 1, 4, 2, 2, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d64_w0_fp16, 64, 16, 2, 5, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d64_w0_fp16, 64, 16, 2, 5, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d64_w0_fp16, 64, 16, 2, 6, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d64_w0_fp16, 64, 16, 2, 6, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d64_w0_fp16, 64, 16, 2, 7, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d64_w0_fp16, 64, 16, 2, 7, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d64_w0_fp16, 64, 16, 2, 8, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv2_d64_w0_fp16, 64, 16, 2, 8, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 1, 16, 1, 4, 1, 1, 2, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 1, 16, 1, 4, 1, 1, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 1, 528, 1, 4, 1, 1, 2, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 1, 528, 1, 4, 1, 1, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 2, 528, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 2, 528, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 3, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 3, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 3, 528, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 3, 528, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 4, 16, 1, 2, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 4, 16, 1, 2, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 4, 528, 1, 2, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 4, 528, 1, 2, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 5, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 5, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 5, 528, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 5, 528, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 6, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 6, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 6, 528, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 6, 528, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 7, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 7, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 7, 528, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 7, 528, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 8, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 8, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 8, 528, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_bf16, 256, 16, 4, 8, 528, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 1, 16, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 2, 1, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 1, 16, 1, 4, 1, 1, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 1, 528, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 2, 1, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 1, 528, 1, 4, 1, 1, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 2, 16, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 2, 1, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 2, 528, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 2, 1, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 2, 528, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 3, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 3, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 3, 528, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 3, 528, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 4, 16, 1, 2, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 4, 16, 1, 2, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 4, 528, 1, 2, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 4, 528, 1, 2, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 5, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 5, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 5, 528, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 5, 528, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 6, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 6, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 6, 528, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 6, 528, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 7, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 7, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 7, 528, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 7, 528, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 8, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 8, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 8, 528, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv4_d256_w0_fp16, 256, 16, 4, 8, 528, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d128_w0_bf16, 128, 16, 8, 1, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d128_w0_bf16, 128, 16, 8, 1, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d128_w0_bf16, 128, 16, 8, 2, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d128_w0_bf16, 128, 16, 8, 2, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d128_w0_bf16, 128, 16, 8, 3, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d128_w0_bf16, 128, 16, 8, 3, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d128_w0_bf16, 128, 16, 8, 4, 16, 1, 1, 2, 4, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d128_w0_bf16, 128, 16, 8, 4, 16, 1, 1, 2, 4, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d128_w0_bf16, 128, 16, 8, 5, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d128_w0_bf16, 128, 16, 8, 5, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d128_w0_bf16, 128, 16, 8, 6, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d128_w0_bf16, 128, 16, 8, 6, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d128_w0_bf16, 128, 16, 8, 7, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d128_w0_bf16, 128, 16, 8, 7, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d128_w0_bf16, 128, 16, 8, 8, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d128_w0_bf16, 128, 16, 8, 8, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d128_w0_fp16, 128, 16, 8, 1, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d128_w0_fp16, 128, 16, 8, 1, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d128_w0_fp16, 128, 16, 8, 2, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d128_w0_fp16, 128, 16, 8, 2, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d128_w0_fp16, 128, 16, 8, 3, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d128_w0_fp16, 128, 16, 8, 3, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d128_w0_fp16, 128, 16, 8, 4, 16, 1, 1, 2, 4, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d128_w0_fp16, 128, 16, 8, 4, 16, 1, 1, 2, 4, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d128_w0_fp16, 128, 16, 8, 5, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d128_w0_fp16, 128, 16, 8, 5, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d128_w0_fp16, 128, 16, 8, 6, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d128_w0_fp16, 128, 16, 8, 6, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d128_w0_fp16, 128, 16, 8, 7, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d128_w0_fp16, 128, 16, 8, 7, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d128_w0_fp16, 128, 16, 8, 8, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d128_w0_fp16, 128, 16, 8, 8, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d256_w0_bf16, 256, 16, 8, 1, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d256_w0_bf16, 256, 16, 8, 1, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d256_w0_bf16, 256, 16, 8, 2, 16, 1, 2, 1, 4, 2, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d256_w0_bf16, 256, 16, 8, 2, 16, 1, 2, 1, 4, 2, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d256_w0_bf16, 256, 16, 8, 3, 16, 1, 2, 1, 2, 2, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d256_w0_bf16, 256, 16, 8, 3, 16, 1, 2, 1, 2, 2, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d256_w0_bf16, 256, 16, 8, 4, 16, 1, 2, 1, 4, 4, 4, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d256_w0_bf16, 256, 16, 8, 4, 16, 1, 2, 1, 4, 4, 4, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d256_w0_bf16, 256, 16, 8, 5, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d256_w0_bf16, 256, 16, 8, 5, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d256_w0_bf16, 256, 16, 8, 6, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d256_w0_bf16, 256, 16, 8, 6, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d256_w0_bf16, 256, 16, 8, 7, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d256_w0_bf16, 256, 16, 8, 7, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d256_w0_bf16, 256, 16, 8, 8, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d256_w0_bf16, 256, 16, 8, 8, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q16_kv8_d256_w0_fp16, 256, 16, 8, 1, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0) +RDNA35_VARIANT(q16_kv8_d256_w0_fp16, 256, 16, 8, 1, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d256_w0_fp16, 256, 16, 8, 2, 16, 1, 2, 1, 4, 2, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d256_w0_fp16, 256, 16, 8, 2, 16, 1, 2, 1, 4, 2, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d256_w0_fp16, 256, 16, 8, 3, 16, 1, 2, 1, 2, 2, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d256_w0_fp16, 256, 16, 8, 3, 16, 1, 2, 1, 2, 2, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d256_w0_fp16, 256, 16, 8, 4, 16, 1, 2, 1, 4, 4, 4, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d256_w0_fp16, 256, 16, 8, 4, 16, 1, 2, 1, 4, 4, 4, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d256_w0_fp16, 256, 16, 8, 5, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d256_w0_fp16, 256, 16, 8, 5, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d256_w0_fp16, 256, 16, 8, 6, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d256_w0_fp16, 256, 16, 8, 6, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d256_w0_fp16, 256, 16, 8, 7, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d256_w0_fp16, 256, 16, 8, 7, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d256_w0_fp16, 256, 16, 8, 8, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d256_w0_fp16, 256, 16, 8, 8, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q16_kv8_d256_w1024_bf16, 256, 16, 8, 1, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 1, 4, 1, 1024, 1) +RDNA35_VARIANT(q16_kv8_d256_w1024_bf16, 256, 16, 8, 1, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q16_kv8_d256_w1024_bf16, 256, 16, 8, 2, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 1024, 1) +RDNA35_VARIANT(q16_kv8_d256_w1024_bf16, 256, 16, 8, 2, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q16_kv8_d256_w1024_bf16, 256, 16, 8, 3, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 1024, 1) +RDNA35_VARIANT(q16_kv8_d256_w1024_bf16, 256, 16, 8, 3, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q16_kv8_d256_w1024_bf16, 256, 16, 8, 4, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q16_kv8_d256_w1024_bf16, 256, 16, 8, 4, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q16_kv8_d256_w1024_bf16, 256, 16, 8, 5, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q16_kv8_d256_w1024_bf16, 256, 16, 8, 5, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q16_kv8_d256_w1024_bf16, 256, 16, 8, 6, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q16_kv8_d256_w1024_bf16, 256, 16, 8, 6, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q16_kv8_d256_w1024_bf16, 256, 16, 8, 7, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q16_kv8_d256_w1024_bf16, 256, 16, 8, 7, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q16_kv8_d256_w1024_bf16, 256, 16, 8, 8, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q16_kv8_d256_w1024_bf16, 256, 16, 8, 8, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q16_kv8_d256_w1024_fp16, 256, 16, 8, 1, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 1, 4, 1, 1024, 0) +RDNA35_VARIANT(q16_kv8_d256_w1024_fp16, 256, 16, 8, 1, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q16_kv8_d256_w1024_fp16, 256, 16, 8, 2, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 1024, 0) +RDNA35_VARIANT(q16_kv8_d256_w1024_fp16, 256, 16, 8, 2, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q16_kv8_d256_w1024_fp16, 256, 16, 8, 3, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 1024, 0) +RDNA35_VARIANT(q16_kv8_d256_w1024_fp16, 256, 16, 8, 3, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q16_kv8_d256_w1024_fp16, 256, 16, 8, 4, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q16_kv8_d256_w1024_fp16, 256, 16, 8, 4, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q16_kv8_d256_w1024_fp16, 256, 16, 8, 5, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q16_kv8_d256_w1024_fp16, 256, 16, 8, 5, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q16_kv8_d256_w1024_fp16, 256, 16, 8, 6, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q16_kv8_d256_w1024_fp16, 256, 16, 8, 6, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q16_kv8_d256_w1024_fp16, 256, 16, 8, 7, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q16_kv8_d256_w1024_fp16, 256, 16, 8, 7, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q16_kv8_d256_w1024_fp16, 256, 16, 8, 8, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q16_kv8_d256_w1024_fp16, 256, 16, 8, 8, 16, 1, 2, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 1, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 1, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 1, 784, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 1, 784, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 2, 784, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 2, 784, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 3, 16, 1, 4, 1, 2, 4, 0, 2, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 3, 16, 1, 4, 1, 2, 4, 0, 2, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 3, 784, 1, 4, 1, 2, 4, 0, 2, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 3, 784, 1, 4, 1, 2, 4, 0, 2, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 4, 16, 1, 4, 1, 1, 4, 0, 2, 1, 1, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 4, 16, 1, 4, 1, 1, 4, 0, 2, 1, 1, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 4, 784, 1, 4, 1, 1, 4, 0, 2, 1, 1, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 4, 784, 1, 4, 1, 1, 4, 0, 2, 1, 1, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 5, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 5, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 5, 784, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 5, 784, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 6, 16, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 6, 16, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 6, 784, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 6, 784, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 7, 16, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 7, 16, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 7, 784, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 7, 784, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 8, 16, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 8, 16, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 8, 784, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_bf16, 256, 24, 4, 8, 784, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 1, 16, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 2, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 1, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 1, 784, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 2, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 1, 784, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 2, 16, 1, 4, 1, 1, 4, 0, 1, 0, 0, 0, 0, 1, 2, 1, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 2, 784, 1, 4, 1, 1, 4, 0, 1, 0, 0, 0, 0, 1, 2, 1, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 2, 784, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 3, 16, 1, 4, 1, 2, 4, 0, 2, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 3, 16, 1, 4, 1, 2, 4, 0, 2, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 3, 784, 1, 4, 1, 2, 4, 0, 2, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 3, 784, 1, 4, 1, 2, 4, 0, 2, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 4, 16, 1, 4, 1, 1, 4, 0, 2, 1, 1, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 4, 16, 1, 4, 1, 1, 4, 0, 2, 1, 1, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 4, 784, 1, 4, 1, 1, 4, 0, 2, 1, 1, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 4, 784, 1, 4, 1, 1, 4, 0, 2, 1, 1, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 5, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 5, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 5, 784, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 5, 784, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 6, 16, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 6, 16, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 6, 784, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 6, 784, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 7, 16, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 7, 16, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 7, 784, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 7, 784, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 8, 16, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 8, 16, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 8, 784, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv4_d256_w0_fp16, 256, 24, 4, 8, 784, 1, 1, 3, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv8_d128_w0_bf16, 128, 24, 8, 1, 16, 1, 1, 1, 1, 4, 0, 1, 0, 1, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv8_d128_w0_bf16, 128, 24, 8, 1, 16, 1, 1, 1, 1, 4, 0, 1, 0, 1, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv8_d128_w0_bf16, 128, 24, 8, 2, 16, 1, 2, 1, 4, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv8_d128_w0_bf16, 128, 24, 8, 2, 16, 1, 2, 1, 4, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv8_d128_w0_bf16, 128, 24, 8, 3, 16, 1, 1, 1, 2, 4, 0, 1, 0, 1, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv8_d128_w0_bf16, 128, 24, 8, 3, 16, 1, 1, 1, 2, 4, 0, 1, 0, 1, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv8_d128_w0_bf16, 128, 24, 8, 4, 16, 1, 1, 1, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv8_d128_w0_bf16, 128, 24, 8, 4, 16, 1, 1, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv8_d128_w0_bf16, 128, 24, 8, 5, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv8_d128_w0_bf16, 128, 24, 8, 5, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv8_d128_w0_bf16, 128, 24, 8, 6, 16, 1, 2, 1, 2, 4, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv8_d128_w0_bf16, 128, 24, 8, 6, 16, 1, 2, 1, 2, 4, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv8_d128_w0_bf16, 128, 24, 8, 7, 16, 1, 2, 1, 2, 4, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv8_d128_w0_bf16, 128, 24, 8, 7, 16, 1, 2, 1, 2, 4, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv8_d128_w0_bf16, 128, 24, 8, 8, 16, 1, 1, 3, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv8_d128_w0_bf16, 128, 24, 8, 8, 16, 1, 1, 3, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q24_kv8_d128_w0_fp16, 128, 24, 8, 1, 16, 1, 1, 1, 1, 4, 0, 1, 0, 1, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv8_d128_w0_fp16, 128, 24, 8, 1, 16, 1, 1, 1, 1, 4, 0, 1, 0, 1, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv8_d128_w0_fp16, 128, 24, 8, 2, 16, 1, 2, 1, 4, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv8_d128_w0_fp16, 128, 24, 8, 2, 16, 1, 2, 1, 4, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv8_d128_w0_fp16, 128, 24, 8, 3, 16, 1, 1, 1, 2, 4, 0, 1, 0, 1, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv8_d128_w0_fp16, 128, 24, 8, 3, 16, 1, 1, 1, 2, 4, 0, 1, 0, 1, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv8_d128_w0_fp16, 128, 24, 8, 4, 16, 1, 1, 1, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv8_d128_w0_fp16, 128, 24, 8, 4, 16, 1, 1, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv8_d128_w0_fp16, 128, 24, 8, 5, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv8_d128_w0_fp16, 128, 24, 8, 5, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv8_d128_w0_fp16, 128, 24, 8, 6, 16, 1, 2, 1, 2, 4, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv8_d128_w0_fp16, 128, 24, 8, 6, 16, 1, 2, 1, 2, 4, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv8_d128_w0_fp16, 128, 24, 8, 7, 16, 1, 2, 1, 2, 4, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv8_d128_w0_fp16, 128, 24, 8, 7, 16, 1, 2, 1, 2, 4, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv8_d128_w0_fp16, 128, 24, 8, 8, 16, 1, 1, 3, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q24_kv8_d128_w0_fp16, 128, 24, 8, 8, 16, 1, 1, 3, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q28_kv4_d128_w0_bf16, 128, 28, 4, 1, 16, 1, 4, 1, 1, 4, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q28_kv4_d128_w0_bf16, 128, 28, 4, 1, 16, 1, 4, 1, 1, 4, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q28_kv4_d128_w0_bf16, 128, 28, 4, 2, 16, 1, 4, 1, 4, 8, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q28_kv4_d128_w0_bf16, 128, 28, 4, 2, 16, 1, 4, 1, 4, 8, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q28_kv4_d128_w0_bf16, 128, 28, 4, 3, 16, 1, 8, 1, 2, 4, 2, 2, 1, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q28_kv4_d128_w0_bf16, 128, 28, 4, 3, 16, 1, 8, 1, 2, 4, 2, 2, 1, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q28_kv4_d128_w0_bf16, 128, 28, 4, 4, 16, 1, 4, 1, 1, 8, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q28_kv4_d128_w0_bf16, 128, 28, 4, 4, 16, 1, 4, 1, 1, 8, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q28_kv4_d128_w0_bf16, 128, 28, 4, 5, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q28_kv4_d128_w0_bf16, 128, 28, 4, 5, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q28_kv4_d128_w0_bf16, 128, 28, 4, 6, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q28_kv4_d128_w0_bf16, 128, 28, 4, 6, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q28_kv4_d128_w0_bf16, 128, 28, 4, 7, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q28_kv4_d128_w0_bf16, 128, 28, 4, 7, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q28_kv4_d128_w0_bf16, 128, 28, 4, 8, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q28_kv4_d128_w0_bf16, 128, 28, 4, 8, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q28_kv4_d128_w0_fp16, 128, 28, 4, 1, 16, 1, 4, 1, 1, 4, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q28_kv4_d128_w0_fp16, 128, 28, 4, 1, 16, 1, 4, 1, 1, 4, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q28_kv4_d128_w0_fp16, 128, 28, 4, 2, 16, 1, 4, 1, 4, 8, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q28_kv4_d128_w0_fp16, 128, 28, 4, 2, 16, 1, 4, 1, 4, 8, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q28_kv4_d128_w0_fp16, 128, 28, 4, 3, 16, 1, 8, 1, 2, 4, 2, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q28_kv4_d128_w0_fp16, 128, 28, 4, 3, 16, 1, 8, 1, 2, 4, 2, 2, 1, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q28_kv4_d128_w0_fp16, 128, 28, 4, 4, 16, 1, 4, 1, 1, 8, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q28_kv4_d128_w0_fp16, 128, 28, 4, 4, 16, 1, 4, 1, 1, 8, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q28_kv4_d128_w0_fp16, 128, 28, 4, 5, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q28_kv4_d128_w0_fp16, 128, 28, 4, 5, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q28_kv4_d128_w0_fp16, 128, 28, 4, 6, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q28_kv4_d128_w0_fp16, 128, 28, 4, 6, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q28_kv4_d128_w0_fp16, 128, 28, 4, 7, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q28_kv4_d128_w0_fp16, 128, 28, 4, 7, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q28_kv4_d128_w0_fp16, 128, 28, 4, 8, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q28_kv4_d128_w0_fp16, 128, 28, 4, 8, 16, 1, 1, 7, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv16_d256_w1024_bf16, 256, 32, 16, 1, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 1, 4, 1, 1024, 1) +RDNA35_VARIANT(q32_kv16_d256_w1024_bf16, 256, 32, 16, 1, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q32_kv16_d256_w1024_bf16, 256, 32, 16, 2, 16, 1, 1, 1, 2, 4, 4, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q32_kv16_d256_w1024_bf16, 256, 32, 16, 2, 16, 1, 1, 1, 2, 4, 4, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q32_kv16_d256_w1024_bf16, 256, 32, 16, 3, 16, 1, 1, 1, 4, 2, 0, 1, 0, 0, 1, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q32_kv16_d256_w1024_bf16, 256, 32, 16, 3, 16, 1, 1, 1, 4, 2, 0, 1, 0, 0, 1, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q32_kv16_d256_w1024_bf16, 256, 32, 16, 4, 16, 1, 1, 1, 2, 2, 0, 1, 0, 0, 1, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q32_kv16_d256_w1024_bf16, 256, 32, 16, 4, 16, 1, 1, 1, 2, 2, 0, 1, 0, 0, 1, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q32_kv16_d256_w1024_bf16, 256, 32, 16, 5, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q32_kv16_d256_w1024_bf16, 256, 32, 16, 5, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q32_kv16_d256_w1024_bf16, 256, 32, 16, 6, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q32_kv16_d256_w1024_bf16, 256, 32, 16, 6, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q32_kv16_d256_w1024_bf16, 256, 32, 16, 7, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q32_kv16_d256_w1024_bf16, 256, 32, 16, 7, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q32_kv16_d256_w1024_bf16, 256, 32, 16, 8, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q32_kv16_d256_w1024_bf16, 256, 32, 16, 8, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q32_kv16_d256_w1024_fp16, 256, 32, 16, 1, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 1, 4, 1, 1024, 0) +RDNA35_VARIANT(q32_kv16_d256_w1024_fp16, 256, 32, 16, 1, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q32_kv16_d256_w1024_fp16, 256, 32, 16, 2, 16, 1, 1, 1, 2, 4, 4, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q32_kv16_d256_w1024_fp16, 256, 32, 16, 2, 16, 1, 1, 1, 2, 4, 4, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q32_kv16_d256_w1024_fp16, 256, 32, 16, 3, 16, 1, 1, 1, 4, 2, 0, 1, 0, 0, 1, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q32_kv16_d256_w1024_fp16, 256, 32, 16, 3, 16, 1, 1, 1, 4, 2, 0, 1, 0, 0, 1, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q32_kv16_d256_w1024_fp16, 256, 32, 16, 4, 16, 1, 1, 1, 2, 2, 0, 1, 0, 0, 1, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q32_kv16_d256_w1024_fp16, 256, 32, 16, 4, 16, 1, 1, 1, 2, 2, 0, 1, 0, 0, 1, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q32_kv16_d256_w1024_fp16, 256, 32, 16, 5, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q32_kv16_d256_w1024_fp16, 256, 32, 16, 5, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q32_kv16_d256_w1024_fp16, 256, 32, 16, 6, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q32_kv16_d256_w1024_fp16, 256, 32, 16, 6, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q32_kv16_d256_w1024_fp16, 256, 32, 16, 7, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q32_kv16_d256_w1024_fp16, 256, 32, 16, 7, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q32_kv16_d256_w1024_fp16, 256, 32, 16, 8, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q32_kv16_d256_w1024_fp16, 256, 32, 16, 8, 16, 1, 1, 1, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q32_kv2_d128_w0_bf16, 128, 32, 2, 1, 16, 1, 2, 4, 2, 4, 0, 1, 0, 1, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv2_d128_w0_bf16, 128, 32, 2, 1, 16, 1, 2, 4, 2, 4, 0, 1, 0, 1, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv2_d128_w0_bf16, 128, 32, 2, 2, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv2_d128_w0_bf16, 128, 32, 2, 2, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv2_d128_w0_bf16, 128, 32, 2, 3, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv2_d128_w0_bf16, 128, 32, 2, 3, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv2_d128_w0_bf16, 128, 32, 2, 4, 16, 1, 4, 4, 2, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv2_d128_w0_bf16, 128, 32, 2, 4, 16, 1, 4, 4, 2, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv2_d128_w0_bf16, 128, 32, 2, 5, 16, 1, 1, 8, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv2_d128_w0_bf16, 128, 32, 2, 5, 16, 1, 1, 8, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv2_d128_w0_bf16, 128, 32, 2, 6, 16, 1, 1, 8, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv2_d128_w0_bf16, 128, 32, 2, 6, 16, 1, 1, 8, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv2_d128_w0_bf16, 128, 32, 2, 7, 16, 1, 1, 8, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv2_d128_w0_bf16, 128, 32, 2, 7, 16, 1, 1, 8, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv2_d128_w0_bf16, 128, 32, 2, 8, 16, 1, 1, 8, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv2_d128_w0_bf16, 128, 32, 2, 8, 16, 1, 1, 8, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv2_d128_w0_fp16, 128, 32, 2, 1, 16, 1, 2, 4, 2, 4, 0, 1, 0, 1, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv2_d128_w0_fp16, 128, 32, 2, 1, 16, 1, 2, 4, 2, 4, 0, 1, 0, 1, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv2_d128_w0_fp16, 128, 32, 2, 2, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv2_d128_w0_fp16, 128, 32, 2, 2, 16, 1, 4, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv2_d128_w0_fp16, 128, 32, 2, 3, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv2_d128_w0_fp16, 128, 32, 2, 3, 16, 1, 2, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv2_d128_w0_fp16, 128, 32, 2, 4, 16, 1, 4, 4, 2, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv2_d128_w0_fp16, 128, 32, 2, 4, 16, 1, 4, 4, 2, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv2_d128_w0_fp16, 128, 32, 2, 5, 16, 1, 1, 8, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv2_d128_w0_fp16, 128, 32, 2, 5, 16, 1, 1, 8, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv2_d128_w0_fp16, 128, 32, 2, 6, 16, 1, 1, 8, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv2_d128_w0_fp16, 128, 32, 2, 6, 16, 1, 1, 8, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv2_d128_w0_fp16, 128, 32, 2, 7, 16, 1, 1, 8, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv2_d128_w0_fp16, 128, 32, 2, 7, 16, 1, 1, 8, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv2_d128_w0_fp16, 128, 32, 2, 8, 16, 1, 1, 8, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv2_d128_w0_fp16, 128, 32, 2, 8, 16, 1, 1, 8, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d128_w0_bf16, 128, 32, 32, 1, 16, 1, 1, 1, 2, 2, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d128_w0_bf16, 128, 32, 32, 1, 16, 1, 1, 1, 2, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d128_w0_bf16, 128, 32, 32, 2, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d128_w0_bf16, 128, 32, 32, 2, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d128_w0_bf16, 128, 32, 32, 3, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d128_w0_bf16, 128, 32, 32, 3, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d128_w0_bf16, 128, 32, 32, 4, 16, 1, 1, 1, 1, 2, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d128_w0_bf16, 128, 32, 32, 4, 16, 1, 1, 1, 1, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d128_w0_bf16, 128, 32, 32, 5, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d128_w0_bf16, 128, 32, 32, 5, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d128_w0_bf16, 128, 32, 32, 6, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d128_w0_bf16, 128, 32, 32, 6, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d128_w0_bf16, 128, 32, 32, 7, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d128_w0_bf16, 128, 32, 32, 7, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d128_w0_bf16, 128, 32, 32, 8, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d128_w0_bf16, 128, 32, 32, 8, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d128_w0_fp16, 128, 32, 32, 1, 16, 1, 1, 1, 2, 2, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d128_w0_fp16, 128, 32, 32, 1, 16, 1, 1, 1, 2, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d128_w0_fp16, 128, 32, 32, 2, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d128_w0_fp16, 128, 32, 32, 2, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d128_w0_fp16, 128, 32, 32, 3, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d128_w0_fp16, 128, 32, 32, 3, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d128_w0_fp16, 128, 32, 32, 4, 16, 1, 1, 1, 1, 2, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d128_w0_fp16, 128, 32, 32, 4, 16, 1, 1, 1, 1, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d128_w0_fp16, 128, 32, 32, 5, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d128_w0_fp16, 128, 32, 32, 5, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d128_w0_fp16, 128, 32, 32, 6, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d128_w0_fp16, 128, 32, 32, 6, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d128_w0_fp16, 128, 32, 32, 7, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d128_w0_fp16, 128, 32, 32, 7, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d128_w0_fp16, 128, 32, 32, 8, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d128_w0_fp16, 128, 32, 32, 8, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d64_w0_bf16, 64, 32, 32, 1, 16, 1, 1, 1, 2, 2, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d64_w0_bf16, 64, 32, 32, 1, 16, 1, 1, 1, 2, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d64_w0_bf16, 64, 32, 32, 2, 16, 1, 1, 1, 1, 2, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d64_w0_bf16, 64, 32, 32, 2, 16, 1, 1, 1, 1, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d64_w0_bf16, 64, 32, 32, 3, 16, 1, 1, 1, 1, 2, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d64_w0_bf16, 64, 32, 32, 3, 16, 1, 1, 1, 1, 2, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d64_w0_bf16, 64, 32, 32, 4, 16, 1, 1, 1, 4, 2, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d64_w0_bf16, 64, 32, 32, 4, 16, 1, 1, 1, 4, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d64_w0_bf16, 64, 32, 32, 5, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d64_w0_bf16, 64, 32, 32, 5, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d64_w0_bf16, 64, 32, 32, 6, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d64_w0_bf16, 64, 32, 32, 6, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d64_w0_bf16, 64, 32, 32, 7, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d64_w0_bf16, 64, 32, 32, 7, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d64_w0_bf16, 64, 32, 32, 8, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d64_w0_bf16, 64, 32, 32, 8, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv32_d64_w0_fp16, 64, 32, 32, 1, 16, 1, 1, 1, 2, 2, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d64_w0_fp16, 64, 32, 32, 1, 16, 1, 1, 1, 2, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d64_w0_fp16, 64, 32, 32, 2, 16, 1, 1, 1, 1, 2, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d64_w0_fp16, 64, 32, 32, 2, 16, 1, 1, 1, 1, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d64_w0_fp16, 64, 32, 32, 3, 16, 1, 1, 1, 1, 2, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d64_w0_fp16, 64, 32, 32, 3, 16, 1, 1, 1, 1, 2, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d64_w0_fp16, 64, 32, 32, 4, 16, 1, 1, 1, 4, 2, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d64_w0_fp16, 64, 32, 32, 4, 16, 1, 1, 1, 4, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d64_w0_fp16, 64, 32, 32, 5, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d64_w0_fp16, 64, 32, 32, 5, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d64_w0_fp16, 64, 32, 32, 6, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d64_w0_fp16, 64, 32, 32, 6, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d64_w0_fp16, 64, 32, 32, 7, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d64_w0_fp16, 64, 32, 32, 7, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d64_w0_fp16, 64, 32, 32, 8, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv32_d64_w0_fp16, 64, 32, 32, 8, 16, 1, 1, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d128_w0_bf16, 128, 32, 4, 1, 16, 1, 2, 2, 4, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d128_w0_bf16, 128, 32, 4, 1, 16, 1, 2, 2, 4, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d128_w0_bf16, 128, 32, 4, 2, 16, 1, 4, 2, 4, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d128_w0_bf16, 128, 32, 4, 2, 16, 1, 4, 2, 4, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d128_w0_bf16, 128, 32, 4, 3, 16, 1, 2, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d128_w0_bf16, 128, 32, 4, 3, 16, 1, 2, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d128_w0_bf16, 128, 32, 4, 4, 16, 1, 1, 4, 1, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d128_w0_bf16, 128, 32, 4, 4, 16, 1, 1, 4, 1, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d128_w0_bf16, 128, 32, 4, 5, 16, 1, 1, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d128_w0_bf16, 128, 32, 4, 5, 16, 1, 1, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d128_w0_bf16, 128, 32, 4, 6, 16, 1, 1, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d128_w0_bf16, 128, 32, 4, 6, 16, 1, 1, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d128_w0_bf16, 128, 32, 4, 7, 16, 1, 1, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d128_w0_bf16, 128, 32, 4, 7, 16, 1, 1, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d128_w0_bf16, 128, 32, 4, 8, 16, 1, 1, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d128_w0_bf16, 128, 32, 4, 8, 16, 1, 1, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d128_w0_fp16, 128, 32, 4, 1, 16, 1, 2, 2, 4, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d128_w0_fp16, 128, 32, 4, 1, 16, 1, 2, 2, 4, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d128_w0_fp16, 128, 32, 4, 2, 16, 1, 4, 2, 4, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d128_w0_fp16, 128, 32, 4, 2, 16, 1, 4, 2, 4, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d128_w0_fp16, 128, 32, 4, 3, 16, 1, 2, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d128_w0_fp16, 128, 32, 4, 3, 16, 1, 2, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d128_w0_fp16, 128, 32, 4, 4, 16, 1, 1, 4, 1, 4, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d128_w0_fp16, 128, 32, 4, 4, 16, 1, 1, 4, 1, 4, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d128_w0_fp16, 128, 32, 4, 5, 16, 1, 1, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d128_w0_fp16, 128, 32, 4, 5, 16, 1, 1, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d128_w0_fp16, 128, 32, 4, 6, 16, 1, 1, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d128_w0_fp16, 128, 32, 4, 6, 16, 1, 1, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d128_w0_fp16, 128, 32, 4, 7, 16, 1, 1, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d128_w0_fp16, 128, 32, 4, 7, 16, 1, 1, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d128_w0_fp16, 128, 32, 4, 8, 16, 1, 1, 4, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d128_w0_fp16, 128, 32, 4, 8, 16, 1, 1, 4, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 1, 16, 1, 4, 1, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 1, 16, 1, 4, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 1, 32, 1, 4, 1, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 1, 32, 1, 4, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 2, 32, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 2, 32, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 3, 16, 1, 2, 2, 2, 8, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 3, 16, 1, 2, 2, 2, 8, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 3, 32, 1, 2, 2, 2, 8, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 3, 32, 1, 2, 2, 2, 8, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 4, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 4, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 4, 32, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 4, 32, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 5, 16, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 5, 16, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 5, 32, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 5, 32, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 6, 16, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 6, 16, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 6, 32, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 6, 32, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 7, 16, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 7, 16, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 7, 32, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 7, 32, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 8, 16, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 8, 16, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 8, 32, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_bf16, 512, 32, 4, 8, 32, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 1, 16, 1, 2, 1, 1, 4, 0, 1, 0, 0, 0, 0, 1, 4, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 1, 16, 1, 4, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 1, 32, 1, 2, 1, 1, 4, 0, 1, 0, 0, 0, 0, 1, 4, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 1, 32, 1, 4, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 2, 16, 1, 2, 1, 1, 4, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 2, 32, 1, 2, 1, 1, 4, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 2, 32, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 3, 16, 1, 2, 2, 2, 8, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 3, 16, 1, 2, 2, 2, 8, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 3, 32, 1, 2, 2, 2, 8, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 3, 32, 1, 2, 2, 2, 8, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 4, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 4, 16, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 4, 32, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 4, 32, 1, 2, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 5, 16, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 5, 16, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 5, 32, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 5, 32, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 6, 16, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 6, 16, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 6, 32, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 6, 32, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 7, 16, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 7, 16, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 7, 32, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 7, 32, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 8, 16, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 8, 16, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 8, 32, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv4_d512_w0_fp16, 512, 32, 4, 8, 32, 1, 1, 4, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d128_w0_bf16, 128, 32, 8, 1, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d128_w0_bf16, 128, 32, 8, 1, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d128_w0_bf16, 128, 32, 8, 2, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d128_w0_bf16, 128, 32, 8, 2, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d128_w0_bf16, 128, 32, 8, 3, 16, 1, 1, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d128_w0_bf16, 128, 32, 8, 3, 16, 1, 1, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d128_w0_bf16, 128, 32, 8, 4, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d128_w0_bf16, 128, 32, 8, 4, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d128_w0_bf16, 128, 32, 8, 5, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d128_w0_bf16, 128, 32, 8, 5, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d128_w0_bf16, 128, 32, 8, 6, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d128_w0_bf16, 128, 32, 8, 6, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d128_w0_bf16, 128, 32, 8, 7, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d128_w0_bf16, 128, 32, 8, 7, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d128_w0_bf16, 128, 32, 8, 8, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d128_w0_bf16, 128, 32, 8, 8, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d128_w0_fp16, 128, 32, 8, 1, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d128_w0_fp16, 128, 32, 8, 1, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d128_w0_fp16, 128, 32, 8, 2, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d128_w0_fp16, 128, 32, 8, 2, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d128_w0_fp16, 128, 32, 8, 3, 16, 1, 1, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d128_w0_fp16, 128, 32, 8, 3, 16, 1, 1, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d128_w0_fp16, 128, 32, 8, 4, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d128_w0_fp16, 128, 32, 8, 4, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d128_w0_fp16, 128, 32, 8, 5, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d128_w0_fp16, 128, 32, 8, 5, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d128_w0_fp16, 128, 32, 8, 6, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d128_w0_fp16, 128, 32, 8, 6, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d128_w0_fp16, 128, 32, 8, 7, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d128_w0_fp16, 128, 32, 8, 7, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d128_w0_fp16, 128, 32, 8, 8, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d128_w0_fp16, 128, 32, 8, 8, 16, 1, 1, 2, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d64_w0_bf16, 64, 32, 8, 1, 16, 1, 1, 1, 2, 4, 0, 1, 1, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d64_w0_bf16, 64, 32, 8, 1, 16, 1, 1, 1, 2, 4, 0, 1, 1, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d64_w0_bf16, 64, 32, 8, 2, 16, 1, 1, 1, 1, 8, 0, 1, 1, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d64_w0_bf16, 64, 32, 8, 2, 16, 1, 1, 1, 1, 8, 0, 1, 1, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d64_w0_bf16, 64, 32, 8, 3, 16, 1, 1, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d64_w0_bf16, 64, 32, 8, 3, 16, 1, 1, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d64_w0_bf16, 64, 32, 8, 4, 16, 1, 1, 1, 1, 4, 0, 1, 1, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d64_w0_bf16, 64, 32, 8, 4, 16, 1, 1, 1, 1, 4, 0, 1, 1, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d64_w0_bf16, 64, 32, 8, 5, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d64_w0_bf16, 64, 32, 8, 5, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d64_w0_bf16, 64, 32, 8, 6, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d64_w0_bf16, 64, 32, 8, 6, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d64_w0_bf16, 64, 32, 8, 7, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d64_w0_bf16, 64, 32, 8, 7, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d64_w0_bf16, 64, 32, 8, 8, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d64_w0_bf16, 64, 32, 8, 8, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q32_kv8_d64_w0_fp16, 64, 32, 8, 1, 16, 1, 1, 1, 2, 4, 0, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d64_w0_fp16, 64, 32, 8, 1, 16, 1, 1, 1, 2, 4, 0, 1, 1, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d64_w0_fp16, 64, 32, 8, 2, 16, 1, 1, 1, 1, 8, 0, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d64_w0_fp16, 64, 32, 8, 2, 16, 1, 1, 1, 1, 8, 0, 1, 1, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d64_w0_fp16, 64, 32, 8, 3, 16, 1, 1, 2, 2, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d64_w0_fp16, 64, 32, 8, 3, 16, 1, 1, 2, 2, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d64_w0_fp16, 64, 32, 8, 4, 16, 1, 1, 1, 1, 4, 0, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d64_w0_fp16, 64, 32, 8, 4, 16, 1, 1, 1, 1, 4, 0, 1, 1, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d64_w0_fp16, 64, 32, 8, 5, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d64_w0_fp16, 64, 32, 8, 5, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d64_w0_fp16, 64, 32, 8, 6, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d64_w0_fp16, 64, 32, 8, 6, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d64_w0_fp16, 64, 32, 8, 7, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d64_w0_fp16, 64, 32, 8, 7, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d64_w0_fp16, 64, 32, 8, 8, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q32_kv8_d64_w0_fp16, 64, 32, 8, 8, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q40_kv8_d128_w0_bf16, 128, 40, 8, 1, 16, 1, 1, 1, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q40_kv8_d128_w0_bf16, 128, 40, 8, 1, 16, 1, 1, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q40_kv8_d128_w0_bf16, 128, 40, 8, 2, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q40_kv8_d128_w0_bf16, 128, 40, 8, 2, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q40_kv8_d128_w0_bf16, 128, 40, 8, 3, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q40_kv8_d128_w0_bf16, 128, 40, 8, 3, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q40_kv8_d128_w0_bf16, 128, 40, 8, 4, 16, 1, 2, 1, 4, 8, 2, 2, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q40_kv8_d128_w0_bf16, 128, 40, 8, 4, 16, 1, 2, 1, 4, 8, 2, 2, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q40_kv8_d128_w0_bf16, 128, 40, 8, 5, 16, 1, 2, 1, 2, 4, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q40_kv8_d128_w0_bf16, 128, 40, 8, 5, 16, 1, 2, 1, 2, 4, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q40_kv8_d128_w0_bf16, 128, 40, 8, 6, 16, 1, 2, 1, 2, 4, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q40_kv8_d128_w0_bf16, 128, 40, 8, 6, 16, 1, 2, 1, 2, 4, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q40_kv8_d128_w0_bf16, 128, 40, 8, 7, 16, 1, 1, 5, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q40_kv8_d128_w0_bf16, 128, 40, 8, 7, 16, 1, 1, 5, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q40_kv8_d128_w0_bf16, 128, 40, 8, 8, 16, 1, 1, 5, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q40_kv8_d128_w0_bf16, 128, 40, 8, 8, 16, 1, 1, 5, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q40_kv8_d128_w0_fp16, 128, 40, 8, 1, 16, 1, 1, 1, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q40_kv8_d128_w0_fp16, 128, 40, 8, 1, 16, 1, 1, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q40_kv8_d128_w0_fp16, 128, 40, 8, 2, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q40_kv8_d128_w0_fp16, 128, 40, 8, 2, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q40_kv8_d128_w0_fp16, 128, 40, 8, 3, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q40_kv8_d128_w0_fp16, 128, 40, 8, 3, 16, 1, 2, 1, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q40_kv8_d128_w0_fp16, 128, 40, 8, 4, 16, 1, 2, 1, 4, 8, 2, 2, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q40_kv8_d128_w0_fp16, 128, 40, 8, 4, 16, 1, 2, 1, 4, 8, 2, 2, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q40_kv8_d128_w0_fp16, 128, 40, 8, 5, 16, 1, 2, 1, 2, 4, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q40_kv8_d128_w0_fp16, 128, 40, 8, 5, 16, 1, 2, 1, 2, 4, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q40_kv8_d128_w0_fp16, 128, 40, 8, 6, 16, 1, 2, 1, 2, 4, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q40_kv8_d128_w0_fp16, 128, 40, 8, 6, 16, 1, 2, 1, 2, 4, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q40_kv8_d128_w0_fp16, 128, 40, 8, 7, 16, 1, 1, 5, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q40_kv8_d128_w0_fp16, 128, 40, 8, 7, 16, 1, 1, 5, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q40_kv8_d128_w0_fp16, 128, 40, 8, 8, 16, 1, 1, 5, 2, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q40_kv8_d128_w0_fp16, 128, 40, 8, 8, 16, 1, 1, 5, 2, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d256_w0_bf16, 256, 8, 1, 1, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d256_w0_bf16, 256, 8, 1, 1, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d256_w0_bf16, 256, 8, 1, 2, 16, 1, 16, 1, 1, 2, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d256_w0_bf16, 256, 8, 1, 2, 16, 1, 16, 1, 1, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d256_w0_bf16, 256, 8, 1, 3, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d256_w0_bf16, 256, 8, 1, 3, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d256_w0_bf16, 256, 8, 1, 4, 16, 1, 16, 2, 1, 4, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d256_w0_bf16, 256, 8, 1, 4, 16, 1, 16, 2, 1, 4, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d256_w0_bf16, 256, 8, 1, 5, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d256_w0_bf16, 256, 8, 1, 5, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d256_w0_bf16, 256, 8, 1, 6, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d256_w0_bf16, 256, 8, 1, 6, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d256_w0_bf16, 256, 8, 1, 7, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d256_w0_bf16, 256, 8, 1, 7, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d256_w0_bf16, 256, 8, 1, 8, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d256_w0_bf16, 256, 8, 1, 8, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d256_w0_fp16, 256, 8, 1, 1, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d256_w0_fp16, 256, 8, 1, 1, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d256_w0_fp16, 256, 8, 1, 2, 16, 1, 16, 1, 1, 2, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d256_w0_fp16, 256, 8, 1, 2, 16, 1, 16, 1, 1, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d256_w0_fp16, 256, 8, 1, 3, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d256_w0_fp16, 256, 8, 1, 3, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d256_w0_fp16, 256, 8, 1, 4, 16, 1, 16, 2, 1, 4, 2, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d256_w0_fp16, 256, 8, 1, 4, 16, 1, 16, 2, 1, 4, 2, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d256_w0_fp16, 256, 8, 1, 5, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d256_w0_fp16, 256, 8, 1, 5, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d256_w0_fp16, 256, 8, 1, 6, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d256_w0_fp16, 256, 8, 1, 6, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d256_w0_fp16, 256, 8, 1, 7, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d256_w0_fp16, 256, 8, 1, 7, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d256_w0_fp16, 256, 8, 1, 8, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d256_w0_fp16, 256, 8, 1, 8, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 1, 16, 1, 16, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 1, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 1, 32, 1, 16, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 1, 32, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 2, 16, 1, 16, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 2, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 2, 32, 1, 16, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 2, 32, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 2, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 3, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 3, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 3, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 3, 32, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 4, 16, 1, 16, 1, 1, 4, 0, 2, 1, 1, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 4, 16, 1, 16, 1, 1, 4, 0, 2, 1, 1, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 4, 32, 1, 16, 1, 1, 4, 0, 2, 1, 1, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 4, 32, 1, 16, 1, 1, 4, 0, 2, 1, 1, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 5, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 5, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 5, 32, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 5, 32, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 6, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 6, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 6, 32, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 6, 32, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 7, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 7, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 7, 32, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 7, 32, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 8, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 8, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 8, 32, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_bf16, 256, 8, 1, 8, 32, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 1, 16, 1, 16, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 1, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 1, 32, 1, 16, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 1, 32, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 2, 16, 1, 16, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 2, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 2, 32, 1, 16, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 2, 32, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 2, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 3, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 3, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 3, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 3, 32, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 4, 16, 1, 16, 1, 1, 4, 0, 2, 1, 1, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 4, 16, 1, 16, 1, 1, 4, 0, 2, 1, 1, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 4, 32, 1, 16, 1, 1, 4, 0, 2, 1, 1, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 4, 32, 1, 16, 1, 1, 4, 0, 2, 1, 1, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 5, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 5, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 5, 32, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 5, 32, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 6, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 6, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 6, 32, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 6, 32, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 7, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 7, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 7, 32, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 7, 32, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 8, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 8, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 8, 32, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d256_w512_fp16, 256, 8, 1, 8, 32, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv1_d512_w0_bf16, 512, 8, 1, 1, 16, 1, 16, 1, 1, 8, 8, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d512_w0_bf16, 512, 8, 1, 1, 16, 1, 16, 1, 1, 8, 8, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d512_w0_bf16, 512, 8, 1, 2, 16, 1, 16, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d512_w0_bf16, 512, 8, 1, 2, 16, 1, 16, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d512_w0_bf16, 512, 8, 1, 3, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d512_w0_bf16, 512, 8, 1, 3, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d512_w0_bf16, 512, 8, 1, 4, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d512_w0_bf16, 512, 8, 1, 4, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d512_w0_bf16, 512, 8, 1, 5, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d512_w0_bf16, 512, 8, 1, 5, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d512_w0_bf16, 512, 8, 1, 6, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d512_w0_bf16, 512, 8, 1, 6, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d512_w0_bf16, 512, 8, 1, 7, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d512_w0_bf16, 512, 8, 1, 7, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d512_w0_bf16, 512, 8, 1, 8, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d512_w0_bf16, 512, 8, 1, 8, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv1_d512_w0_fp16, 512, 8, 1, 1, 16, 1, 16, 1, 1, 8, 8, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d512_w0_fp16, 512, 8, 1, 1, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d512_w0_fp16, 512, 8, 1, 2, 16, 1, 16, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d512_w0_fp16, 512, 8, 1, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d512_w0_fp16, 512, 8, 1, 3, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d512_w0_fp16, 512, 8, 1, 3, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d512_w0_fp16, 512, 8, 1, 4, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d512_w0_fp16, 512, 8, 1, 4, 16, 1, 8, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d512_w0_fp16, 512, 8, 1, 5, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d512_w0_fp16, 512, 8, 1, 5, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d512_w0_fp16, 512, 8, 1, 6, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d512_w0_fp16, 512, 8, 1, 6, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d512_w0_fp16, 512, 8, 1, 7, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d512_w0_fp16, 512, 8, 1, 7, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d512_w0_fp16, 512, 8, 1, 8, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv1_d512_w0_fp16, 512, 8, 1, 8, 16, 1, 4, 4, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 1, 16, 1, 8, 1, 1, 4, 4, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 1, 16, 1, 8, 1, 1, 4, 4, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 1, 544, 1, 8, 1, 1, 4, 4, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 1, 544, 1, 8, 1, 1, 4, 4, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 2, 544, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 2, 544, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 3, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 3, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 3, 544, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 3, 544, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 4, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 4, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 4, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 4, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 5, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 5, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 5, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 5, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 6, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 6, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 6, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 6, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 7, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 7, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 7, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 7, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 8, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 8, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 8, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_bf16, 256, 8, 2, 8, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 1, 16, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 1, 16, 1, 8, 1, 1, 4, 4, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 1, 544, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 1, 544, 1, 8, 1, 1, 4, 4, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 2, 544, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 2, 544, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 3, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 3, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 3, 544, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 3, 544, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 4, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 4, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 4, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 4, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 5, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 5, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 5, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 5, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 6, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 6, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 6, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 6, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 7, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 7, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 7, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 7, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 8, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 8, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 8, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w0_fp16, 256, 8, 2, 8, 544, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 1, 16, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 1, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 1, 32, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 1, 32, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 2, 32, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 2, 32, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 3, 16, 1, 8, 1, 1, 4, 4, 1, 1, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 3, 16, 1, 8, 1, 1, 4, 4, 1, 1, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 3, 32, 1, 8, 1, 1, 4, 4, 1, 1, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 3, 32, 1, 8, 1, 1, 4, 4, 1, 1, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 4, 16, 1, 8, 1, 1, 4, 4, 1, 1, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 4, 16, 1, 8, 1, 1, 4, 4, 1, 1, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 4, 32, 1, 8, 1, 1, 4, 4, 1, 1, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 4, 32, 1, 8, 1, 1, 4, 4, 1, 1, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 5, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 5, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 5, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 5, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 6, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 6, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 6, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 6, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 7, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 7, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 7, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 7, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 8, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 8, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 8, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_bf16, 256, 8, 2, 8, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 1) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 1, 16, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 1, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 1, 32, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 1, 32, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 2, 32, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 2, 32, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 3, 16, 1, 8, 1, 1, 4, 4, 1, 1, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 3, 16, 1, 8, 1, 1, 4, 4, 1, 1, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 3, 32, 1, 8, 1, 1, 4, 4, 1, 1, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 3, 32, 1, 8, 1, 1, 4, 4, 1, 1, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 4, 16, 1, 8, 1, 1, 4, 4, 1, 1, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 4, 16, 1, 8, 1, 1, 4, 4, 1, 1, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 4, 32, 1, 8, 1, 1, 4, 4, 1, 1, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 4, 32, 1, 8, 1, 1, 4, 4, 1, 1, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 5, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 5, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 5, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 5, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 6, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 6, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 6, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 6, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 7, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 7, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 7, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 7, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 8, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 8, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 8, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d256_w512_fp16, 256, 8, 2, 8, 32, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 512, 0) +RDNA35_VARIANT(q8_kv2_d512_w0_bf16, 512, 8, 2, 1, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d512_w0_bf16, 512, 8, 2, 1, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d512_w0_bf16, 512, 8, 2, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d512_w0_bf16, 512, 8, 2, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d512_w0_bf16, 512, 8, 2, 3, 16, 1, 4, 1, 2, 4, 0, 1, 0, 1, 1, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d512_w0_bf16, 512, 8, 2, 3, 16, 1, 4, 1, 2, 4, 0, 1, 0, 1, 1, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d512_w0_bf16, 512, 8, 2, 4, 16, 1, 4, 2, 1, 8, 8, 1, 1, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d512_w0_bf16, 512, 8, 2, 4, 16, 1, 4, 2, 1, 8, 8, 1, 1, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d512_w0_bf16, 512, 8, 2, 5, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d512_w0_bf16, 512, 8, 2, 5, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d512_w0_bf16, 512, 8, 2, 6, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d512_w0_bf16, 512, 8, 2, 6, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d512_w0_bf16, 512, 8, 2, 7, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d512_w0_bf16, 512, 8, 2, 7, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d512_w0_bf16, 512, 8, 2, 8, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d512_w0_bf16, 512, 8, 2, 8, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv2_d512_w0_fp16, 512, 8, 2, 1, 16, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d512_w0_fp16, 512, 8, 2, 1, 16, 1, 8, 1, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d512_w0_fp16, 512, 8, 2, 2, 16, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d512_w0_fp16, 512, 8, 2, 2, 16, 1, 8, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d512_w0_fp16, 512, 8, 2, 3, 16, 1, 4, 1, 2, 4, 0, 1, 0, 1, 1, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d512_w0_fp16, 512, 8, 2, 3, 16, 1, 4, 1, 2, 4, 0, 1, 0, 1, 1, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d512_w0_fp16, 512, 8, 2, 4, 16, 1, 4, 2, 1, 8, 8, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d512_w0_fp16, 512, 8, 2, 4, 16, 1, 4, 2, 1, 8, 8, 1, 1, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d512_w0_fp16, 512, 8, 2, 5, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d512_w0_fp16, 512, 8, 2, 5, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d512_w0_fp16, 512, 8, 2, 6, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d512_w0_fp16, 512, 8, 2, 6, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d512_w0_fp16, 512, 8, 2, 7, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d512_w0_fp16, 512, 8, 2, 7, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d512_w0_fp16, 512, 8, 2, 8, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv2_d512_w0_fp16, 512, 8, 2, 8, 16, 1, 4, 2, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv4_d256_w0_bf16, 256, 8, 4, 1, 16, 1, 4, 1, 1, 2, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv4_d256_w0_bf16, 256, 8, 4, 1, 16, 1, 4, 1, 1, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv4_d256_w0_bf16, 256, 8, 4, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv4_d256_w0_bf16, 256, 8, 4, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv4_d256_w0_bf16, 256, 8, 4, 3, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv4_d256_w0_bf16, 256, 8, 4, 3, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv4_d256_w0_bf16, 256, 8, 4, 4, 16, 1, 2, 2, 2, 4, 4, 1, 1, 1, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv4_d256_w0_bf16, 256, 8, 4, 4, 16, 1, 2, 2, 2, 4, 4, 1, 1, 1, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv4_d256_w0_bf16, 256, 8, 4, 5, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv4_d256_w0_bf16, 256, 8, 4, 5, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv4_d256_w0_bf16, 256, 8, 4, 6, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv4_d256_w0_bf16, 256, 8, 4, 6, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv4_d256_w0_bf16, 256, 8, 4, 7, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv4_d256_w0_bf16, 256, 8, 4, 7, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv4_d256_w0_bf16, 256, 8, 4, 8, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv4_d256_w0_bf16, 256, 8, 4, 8, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1) +RDNA35_VARIANT(q8_kv4_d256_w0_fp16, 256, 8, 4, 1, 16, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 0, 0) +RDNA35_VARIANT(q8_kv4_d256_w0_fp16, 256, 8, 4, 1, 16, 1, 4, 1, 1, 2, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv4_d256_w0_fp16, 256, 8, 4, 2, 16, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 0, 0) +RDNA35_VARIANT(q8_kv4_d256_w0_fp16, 256, 8, 4, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv4_d256_w0_fp16, 256, 8, 4, 3, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv4_d256_w0_fp16, 256, 8, 4, 3, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv4_d256_w0_fp16, 256, 8, 4, 4, 16, 1, 2, 2, 2, 4, 4, 1, 1, 1, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv4_d256_w0_fp16, 256, 8, 4, 4, 16, 1, 2, 2, 2, 4, 4, 1, 1, 1, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv4_d256_w0_fp16, 256, 8, 4, 5, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv4_d256_w0_fp16, 256, 8, 4, 5, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv4_d256_w0_fp16, 256, 8, 4, 6, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv4_d256_w0_fp16, 256, 8, 4, 6, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv4_d256_w0_fp16, 256, 8, 4, 7, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv4_d256_w0_fp16, 256, 8, 4, 7, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv4_d256_w0_fp16, 256, 8, 4, 8, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv4_d256_w0_fp16, 256, 8, 4, 8, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0) +RDNA35_VARIANT(q8_kv4_d256_w1024_bf16, 256, 8, 4, 1, 16, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 1024, 1) +RDNA35_VARIANT(q8_kv4_d256_w1024_bf16, 256, 8, 4, 1, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q8_kv4_d256_w1024_bf16, 256, 8, 4, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 1024, 1) +RDNA35_VARIANT(q8_kv4_d256_w1024_bf16, 256, 8, 4, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q8_kv4_d256_w1024_bf16, 256, 8, 4, 3, 16, 1, 2, 1, 2, 4, 4, 1, 1, 0, 0, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q8_kv4_d256_w1024_bf16, 256, 8, 4, 3, 16, 1, 2, 1, 2, 4, 4, 1, 1, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q8_kv4_d256_w1024_bf16, 256, 8, 4, 4, 16, 1, 2, 1, 4, 4, 4, 1, 1, 1, 0, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q8_kv4_d256_w1024_bf16, 256, 8, 4, 4, 16, 1, 2, 1, 4, 4, 4, 1, 1, 1, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q8_kv4_d256_w1024_bf16, 256, 8, 4, 5, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q8_kv4_d256_w1024_bf16, 256, 8, 4, 5, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q8_kv4_d256_w1024_bf16, 256, 8, 4, 6, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q8_kv4_d256_w1024_bf16, 256, 8, 4, 6, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q8_kv4_d256_w1024_bf16, 256, 8, 4, 7, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q8_kv4_d256_w1024_bf16, 256, 8, 4, 7, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q8_kv4_d256_w1024_bf16, 256, 8, 4, 8, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q8_kv4_d256_w1024_bf16, 256, 8, 4, 8, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 1) +RDNA35_VARIANT(q8_kv4_d256_w1024_fp16, 256, 8, 4, 1, 16, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 1024, 0) +RDNA35_VARIANT(q8_kv4_d256_w1024_fp16, 256, 8, 4, 1, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q8_kv4_d256_w1024_fp16, 256, 8, 4, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 1024, 0) +RDNA35_VARIANT(q8_kv4_d256_w1024_fp16, 256, 8, 4, 2, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q8_kv4_d256_w1024_fp16, 256, 8, 4, 3, 16, 1, 2, 1, 2, 4, 4, 1, 1, 0, 0, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q8_kv4_d256_w1024_fp16, 256, 8, 4, 3, 16, 1, 2, 1, 2, 4, 4, 1, 1, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q8_kv4_d256_w1024_fp16, 256, 8, 4, 4, 16, 1, 2, 1, 4, 4, 4, 1, 1, 1, 0, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q8_kv4_d256_w1024_fp16, 256, 8, 4, 4, 16, 1, 2, 1, 4, 4, 4, 1, 1, 1, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q8_kv4_d256_w1024_fp16, 256, 8, 4, 5, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q8_kv4_d256_w1024_fp16, 256, 8, 4, 5, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q8_kv4_d256_w1024_fp16, 256, 8, 4, 6, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q8_kv4_d256_w1024_fp16, 256, 8, 4, 6, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q8_kv4_d256_w1024_fp16, 256, 8, 4, 7, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q8_kv4_d256_w1024_fp16, 256, 8, 4, 7, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q8_kv4_d256_w1024_fp16, 256, 8, 4, 8, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q8_kv4_d256_w1024_fp16, 256, 8, 4, 8, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1024, 0) +RDNA35_VARIANT(q8_kv4_d256_w4096_bf16, 256, 8, 4, 1, 16, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 4096, 1) +RDNA35_VARIANT(q8_kv4_d256_w4096_bf16, 256, 8, 4, 1, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 4096, 1) +RDNA35_VARIANT(q8_kv4_d256_w4096_bf16, 256, 8, 4, 2, 16, 1, 2, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 4096, 1) +RDNA35_VARIANT(q8_kv4_d256_w4096_bf16, 256, 8, 4, 2, 16, 1, 2, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 4096, 1) +RDNA35_VARIANT(q8_kv4_d256_w4096_bf16, 256, 8, 4, 3, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 4096, 1) +RDNA35_VARIANT(q8_kv4_d256_w4096_bf16, 256, 8, 4, 3, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 4096, 1) +RDNA35_VARIANT(q8_kv4_d256_w4096_bf16, 256, 8, 4, 4, 16, 1, 2, 1, 1, 4, 2, 1, 0, 0, 1, 0, 0, 0, 0, 4096, 1) +RDNA35_VARIANT(q8_kv4_d256_w4096_bf16, 256, 8, 4, 4, 16, 1, 2, 1, 1, 4, 2, 1, 0, 0, 1, 1, 0, 0, 0, 4096, 1) +RDNA35_VARIANT(q8_kv4_d256_w4096_bf16, 256, 8, 4, 5, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 4096, 1) +RDNA35_VARIANT(q8_kv4_d256_w4096_bf16, 256, 8, 4, 5, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 4096, 1) +RDNA35_VARIANT(q8_kv4_d256_w4096_bf16, 256, 8, 4, 6, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 4096, 1) +RDNA35_VARIANT(q8_kv4_d256_w4096_bf16, 256, 8, 4, 6, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 4096, 1) +RDNA35_VARIANT(q8_kv4_d256_w4096_bf16, 256, 8, 4, 7, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 4096, 1) +RDNA35_VARIANT(q8_kv4_d256_w4096_bf16, 256, 8, 4, 7, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 4096, 1) +RDNA35_VARIANT(q8_kv4_d256_w4096_bf16, 256, 8, 4, 8, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 4096, 1) +RDNA35_VARIANT(q8_kv4_d256_w4096_bf16, 256, 8, 4, 8, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 4096, 1) +RDNA35_VARIANT(q8_kv4_d256_w4096_fp16, 256, 8, 4, 1, 16, 1, 2, 1, 1, 8, 0, 1, 0, 0, 0, 0, 1, 4, 1, 4096, 0) +RDNA35_VARIANT(q8_kv4_d256_w4096_fp16, 256, 8, 4, 1, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 4096, 0) +RDNA35_VARIANT(q8_kv4_d256_w4096_fp16, 256, 8, 4, 2, 16, 1, 2, 2, 1, 4, 0, 1, 0, 0, 0, 0, 0, 0, 0, 4096, 0) +RDNA35_VARIANT(q8_kv4_d256_w4096_fp16, 256, 8, 4, 2, 16, 1, 2, 2, 1, 4, 0, 1, 0, 0, 0, 1, 0, 0, 0, 4096, 0) +RDNA35_VARIANT(q8_kv4_d256_w4096_fp16, 256, 8, 4, 3, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 4096, 0) +RDNA35_VARIANT(q8_kv4_d256_w4096_fp16, 256, 8, 4, 3, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 4096, 0) +RDNA35_VARIANT(q8_kv4_d256_w4096_fp16, 256, 8, 4, 4, 16, 1, 2, 1, 1, 4, 2, 1, 0, 0, 1, 0, 0, 0, 0, 4096, 0) +RDNA35_VARIANT(q8_kv4_d256_w4096_fp16, 256, 8, 4, 4, 16, 1, 2, 1, 1, 4, 2, 1, 0, 0, 1, 1, 0, 0, 0, 4096, 0) +RDNA35_VARIANT(q8_kv4_d256_w4096_fp16, 256, 8, 4, 5, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 4096, 0) +RDNA35_VARIANT(q8_kv4_d256_w4096_fp16, 256, 8, 4, 5, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 4096, 0) +RDNA35_VARIANT(q8_kv4_d256_w4096_fp16, 256, 8, 4, 6, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 4096, 0) +RDNA35_VARIANT(q8_kv4_d256_w4096_fp16, 256, 8, 4, 6, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 4096, 0) +RDNA35_VARIANT(q8_kv4_d256_w4096_fp16, 256, 8, 4, 7, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 4096, 0) +RDNA35_VARIANT(q8_kv4_d256_w4096_fp16, 256, 8, 4, 7, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 4096, 0) +RDNA35_VARIANT(q8_kv4_d256_w4096_fp16, 256, 8, 4, 8, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 0, 0, 0, 0, 4096, 0) +RDNA35_VARIANT(q8_kv4_d256_w4096_fp16, 256, 8, 4, 8, 16, 1, 4, 1, 1, 8, 0, 1, 0, 0, 0, 1, 0, 0, 0, 4096, 0) diff --git a/csrc/rocm/rdna35_attn/variants.h b/csrc/rocm/rdna35_attn/variants.h new file mode 100644 index 000000000000..0e7fd2f9d230 --- /dev/null +++ b/csrc/rocm/rdna35_attn/variants.h @@ -0,0 +1,30 @@ +// Every decode-attention variant the wheel carries, from variants.def. +// +// Only the registry includes this: a group unit sees its own entries alone +// (cmake/rdna35_attn.cmake), so editing one group's lines of variants.def +// rebuilds that unit and the registry, not the rest. +#pragma once + +#include "rdna35_attn/launch.h" +#include "rdna35_attn/macros.h" + +namespace rdna35 { + +#define RDNA35_VARIANT(...) \ + void RDNA35_CALL(RDNA35_ID, (launch, __VA_ARGS__))(const LaunchArgs&); +#include "rdna35_attn/variants.def" +#undef RDNA35_VARIANT + +struct Variant { + int key[RDNA35_NFIELDS]; + LaunchFn fn; +}; + +inline constexpr Variant kVariants[] = { +#define RDNA35_VARIANT(...) \ + {{RDNA35_KEY(__VA_ARGS__)}, &RDNA35_CALL(RDNA35_ID, (launch, __VA_ARGS__))}, +#include "rdna35_attn/variants.def" +#undef RDNA35_VARIANT +}; + +} // namespace rdna35 diff --git a/csrc/rocm/rdna35_decode_attn.cu b/csrc/rocm/rdna35_decode_attn.cu new file mode 100644 index 000000000000..2061248ec23d --- /dev/null +++ b/csrc/rocm/rdna35_decode_attn.cu @@ -0,0 +1,1784 @@ +// RDNA3.5 (gfx1151) decode attention over vLLM's paged KV cache, GQA-packed. +// +// Reached from ROCM_ATTN on gfx1151. The cache is vLLM's single paged +// tensor with K and V packed in the content dim, exactly as +// TritonAttentionBackend produces it, so no repacking is needed: +// +// logical (NB, HKV, BS, 2*D) +// NHD (NB, BS, HKV, 2*D) +// HND (NB, HKV, BS, 2*D) +// +// Every shape is a compile-time define, so one build serves one shape tuple; +// vllm/v1/attention/ops/rdna35_hip_decode.py drives the compilation. +// +// One workgroup per (kv head, row group, KV segment). Every query row that +// reads a kv head -- GQA heads times MAXM tokens -- is served by the same +// workgroup, so each KV byte is read from memory once, not once per q head. +// +// A block is TILES tiles of 16 keys. Each tile belongs to DSPL waves, each +// owning 1/DSPL of the head dim: they split Q@K over d and sum their partial +// scores through LDS, then each runs P@V for its own d. Every wave keeps its +// own online softmax, so the waves merge only once, at the end. Both +// products are WMMA 16x16x16 fp16 (or bf16, with KV_BF16=1) -> f32: +// +// S^T[key][row] = K[key][:] . Q[row][:] A = K tile, B = Q^T +// O^T[d][row] += V[key][d] . P[row][key] A = V^T, B = P^T +#include + +#include +#include +#include + +#include "rdna35_attn/launch.h" + +#ifndef RDNA35_BODY + #ifndef HEAD_DIM + #define HEAD_DIM 128 + #endif + #ifndef NUM_Q_HEADS + #define NUM_Q_HEADS 32 + #endif + #ifndef NUM_KV_HEADS + #define NUM_KV_HEADS 8 + #endif + #ifndef MAXM + #define MAXM 1 + #endif + #ifndef BS + #define BS 16 + #endif + #ifndef LAYOUT + #define LAYOUT 1 // 0 = NHD, 1 = HND + #endif + // Element type of Q, the KV cache and the output: 0 = fp16, 1 = bf16. + #ifndef KV_BF16 + #define KV_BF16 0 + #endif + // Most KV segments a (kv head, row group) is split over. The number active + // is clamp(nblocks / MINB, 1, NSEG), decided at run time from S: the grid is + // fixed when a CUDA graph is captured, S is not -- the kernel reads it from + // the device, since a host argument would be frozen into the graph too. + #ifndef NSEG + #define NSEG 8 + #endif + // Row groups per kv head. Each workgroup carries ROWS/RG of the kv head's + // rows and reads all of its KV, the RG workgroups sharing it through L2. + #ifndef RG + #define RG 1 + #endif + // Least KV blocks an active segment is given. + #ifndef MINB + #define MINB 1 + #endif + #ifndef MUTATE + #define MUTATE 0 + #endif + // Waves per workgroup. + #ifndef NW + #define NW 8 + #endif + // Waves sharing one key tile, each owning 1/DSPL of the head dim. Keeps a + // wave's accumulator, K and V slice within budget at large D. 0, like + // leaving it undefined, keeps this rule. + #if !defined(DSPL) || DSPL == 0 + #undef DSPL + #define DSPL (HEAD_DIM >= 256 ? HEAD_DIM / 128 : 1) + #endif + // Waves sharing one key tile, each carrying 1/RSPL of the workgroup's row + // tiles over its whole d part. The in-workgroup form of RG: the waves read + // the same tile at nearly the same time, so its KV comes from DRAM once + // without depending on how separate workgroups drift. + #ifndef RSPL + #define RSPL 1 + #endif + // 1 keeps a second tile's loads in flight per wave (unshared tiles). + #ifndef PF + #define PF 0 + #endif + // 1 runs the per-q-head dot-product decomposition instead of WMMA (see + // the dot body), with BFLY and GT its butterfly split and grid order. + #ifndef DOT + #define DOT 0 + #endif + #ifndef BFLY + #define BFLY 0 + #endif + #ifndef GT + #define GT 0 + #endif + + // Sliding window: a query at position p attends to keys p-WIN+1 .. p. 0 is + // full causal attention. Only the window's blocks are read, whatever S is. + #ifndef WIN + #define WIN 0 + #endif + + // Measurement only, wrong answers: 1 skips the V loads, 2 thins P@V to one + // WMMA per tile, 4 skips the K loads, 16 thins Q@K to one WMMA per tile, 32 + // returns at once, 64 loads each tile and does nothing else, 128 skips the + // DSPL waves' score exchange. + #ifndef ABLATE + #define ABLATE 0 + #endif + // 1 stages a split-KV partial in LDS before writing it, so that each store + // covers whole lines, where the merge's tree would otherwise end in one + // live tile writing its rows straight from registers, 32 bytes a row per + // store. + #ifndef CPUB + #define CPUB 0 + #endif + // 1 stages each unshared tile's V in LDS beside its K, so the tile's + // registers are free once staged and the next tile's loads go out before + // this one is computed -- PF's overlap without PF's second register set. + #ifndef VINLDS + #define VINLDS 0 + #endif + // 1 builds for a batch of sequences, one per grid.y; 0 for one sequence, + // without the per-sequence addressing. + #ifndef BATCH + #define BATCH 0 + #endif + // Workgroups a batch aims for when the host caps its segments. + #ifndef BTARGET + #define BTARGET 128 + #endif + // Measurement only: records the 100 MHz realtime counter at phase boundaries + // of thread 0 of workgroups 0..7 into g_ts, read back by the timings() op. + // 1 marks the phases, 3 also waits for the first tile's data to separate + // memory from compute. + #ifndef TIMING + #define TIMING 0 + #endif + + #define WAVE 32 + #define BLOCK (NW * WAVE) + #define GQA (NUM_Q_HEADS / NUM_KV_HEADS) + #define ROWS (GQA * MAXM) + #define NCHUNK (HEAD_DIM / 16) + #define KV_ROW (2 * HEAD_DIM) + #define PAGE_ELEMS (BS * NUM_KV_HEADS * KV_ROW) + #if LAYOUT == 0 + #define KEY_STRIDE (NUM_KV_HEADS * KV_ROW) + #else + #define KEY_STRIDE KV_ROW + #endif + #define LOG2E 1.44269504088896340736f + #define QPAD 8 + // Merge buffer: one slot per (live tile, d part) holds the real rows of a row + // tile, each DPART floats plus a pad. + #define PADM 4 + #define MBUDGET (36 * 1024) + +static_assert(NUM_Q_HEADS % NUM_KV_HEADS == 0, "GQA must be integral"); +static_assert(BS % 16 == 0, "a 16-key tile must sit inside one page"); + + #if TIMING + // Workgroups TS_WG0..TS_WG0+7 are recorded. + #ifndef TS_WG0 + #define TS_WG0 0 + #endif +__device__ unsigned long long g_ts[8 * 16]; + #define TS(i) \ + do { \ + const unsigned w_ = blockIdx.x - TS_WG0; \ + if (tid == 0 && w_ < 8) \ + g_ts[w_ * 16 + (i)] = __builtin_amdgcn_s_sendmsg_rtnl(0x83); \ + } while (0) + #else + #define TS(i) \ + do { \ + } while (0) + #endif + + #if KV_BF16 +typedef __bf16 elem_t; + #else +typedef _Float16 elem_t; + #endif +typedef elem_t e16 __attribute__((ext_vector_type(16))); +typedef elem_t e8 __attribute__((ext_vector_type(8))); +typedef elem_t e2 __attribute__((ext_vector_type(2))); +typedef short s16 __attribute__((ext_vector_type(16))); +typedef short s2v __attribute__((ext_vector_type(2))); +typedef float f8 __attribute__((ext_vector_type(8))); +typedef float f4 __attribute__((ext_vector_type(4))); +typedef unsigned u8v __attribute__((ext_vector_type(8))); +typedef unsigned u4v __attribute__((ext_vector_type(4))); +typedef unsigned u2v __attribute__((ext_vector_type(2))); + +// Element offset of (page, slot) for this kv head. +__device__ __forceinline__ size_t tile_off(int page, unsigned slot, int kvh) { + #if LAYOUT == 0 + return (size_t)page * PAGE_ELEMS + + ((size_t)slot * NUM_KV_HEADS + kvh) * KV_ROW; + #else + return (size_t)page * PAGE_ELEMS + ((size_t)kvh * BS + slot) * KV_ROW; + #endif +} + +// Every barrier in this kernel orders LDS only; global visibility, where the +// split-KV merge needs it, is __threadfence's job. __syncthreads() also +// waits for every outstanding global load (vmcnt(0)) and invalidates L0, so +// the barrier that publishes Q waited for the whole first KV tile to land. +__device__ __forceinline__ void lds_barrier() { + asm volatile("s_waitcnt lgkmcnt(0)\n\ts_barrier" ::: "memory"); +} + +// The cross-lane moves below all read lanes that are active, so fetch- +// inactive changes nothing -- except that with it set the compiler no longer +// ties the destination to a copy of the source: 66 of 207 v_mov gone. +// +// The other 16-lane half's value for this lane: lane l <-> l ^ 16. +__device__ __forceinline__ unsigned xhalf_u(unsigned v) { + return (unsigned)__builtin_amdgcn_permlanex16((int)v, (int)v, 0x76543210u, + 0xFEDCBA98u, true, false); +} + +__device__ __forceinline__ float xhalf(float v) { + return __builtin_bit_cast(float, xhalf_u(__builtin_bit_cast(unsigned, v))); +} + +// Within each 16-lane row: lane i takes lane i & 7's value, or lane i | 8's. +__device__ __forceinline__ unsigned lower8(unsigned v) { + return (unsigned)__builtin_amdgcn_permlane16((int)v, (int)v, 0x76543210u, + 0x76543210u, true, false); +} + +__device__ __forceinline__ float upper8f(float v) { + return __builtin_bit_cast( + float, __builtin_amdgcn_permlane16( + __builtin_bit_cast(int, v), __builtin_bit_cast(int, v), + 0xFEDCBA98u, 0xFEDCBA98u, true, false)); +} + +__device__ __forceinline__ float weight_of(float m, float gmax) { + return (gmax == -INFINITY) ? 0.f : __builtin_amdgcn_exp2f(m - gmax); +} + +// P's high half as an element, and as the float it stands for. bf16 +// truncates: gfx1151 has no f32 -> bf16 conversion, so rounding would be +// several VALU ops per value where truncation packs two in one v_perm, and +// the low half carries whatever truncation drops. +__device__ __forceinline__ unsigned pack2(float a, float b) { + #if KV_BF16 + return __builtin_amdgcn_perm(__builtin_bit_cast(unsigned, b), + __builtin_bit_cast(unsigned, a), 0x07060302u); + #else + return __builtin_bit_cast(unsigned, e2{(elem_t)a, (elem_t)b}); + #endif +} + +__device__ __forceinline__ float elem_part(float p) { + #if KV_BF16 + return __builtin_bit_cast(float, + __builtin_bit_cast(unsigned, p) & 0xFFFF0000u); + #else + return (float)(elem_t)p; + #endif +} + +// An output value, rounded to nearest even. The bf16 cast also keeps NaN a +// NaN, which costs a compare and a select per value; here a NaN may come out +// as Inf, still not finite. +__device__ __forceinline__ elem_t to_elem(float x) { + #if KV_BF16 + const unsigned u = __builtin_bit_cast(unsigned, x); + return __builtin_bit_cast( + elem_t, (unsigned short)((u + 0x7FFFu + ((u >> 16) & 1u)) >> 16)); + #else + return (elem_t)x; + #endif +} + +__device__ __forceinline__ f8 wmma(e16 a, e16 b, f8 c) { + #if KV_BF16 + return __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32( + __builtin_bit_cast(s16, a), __builtin_bit_cast(s16, b), c); + #else + return __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a, b, c); + #endif +} + +// WMMA's k index is free as long as A and B agree on it -- and each half of +// the wave may order it differently: measured on gfx1151, the lower half +// produces the even output rows from its own copy of A (only its even rows +// are read) and of B, the upper half the odd ones from its own. So both +// products order k as "this half's keys, then the other half's": the lower +// half puts keys 2i first, the upper half keys 2i+1, which is the order the +// S^T accumulator already hands each half its keys in. Every operand is +// then (own, other half's) in every lane, with no select on the half. +// +// B = P^T: lane l (row l16) has P for keys 2e + hi in p[e]. +__device__ __forceinline__ e16 p_frag(const float* p) { + u8v r; + #pragma unroll + for (int i = 0; i < 4; ++i) { + r[i] = pack2(p[2 * i], p[2 * i + 1]); + r[4 + i] = xhalf_u(r[i]); + } + return __builtin_bit_cast(e16, r); +} + + #define RDNA35_BODY 1 + #define M_NSEG NSEG + #define M_RG RG + #define M_MINB MINB + #define M_DSPL DSPL + #define M_RSPL RSPL + #define M_NW NW + #define M_PF PF + #define M_DOT DOT + #define M_BFLY BFLY + #define M_GT GT +namespace decomp { + #include __FILE_NAME__ +} // namespace decomp + +// Per-sequence strides of the batch: grid.y is the sequence, and every +// pointer is moved to that sequence's rows before the body runs, so the body +// sees one sequence exactly as it did before batching. +struct Batch { + int bt_stride, acc_stride, ml_stride, cnt_stride; +}; + +template +__global__ __launch_bounds__(BLOCK) void decode_attn( + const elem_t* __restrict__ q, const elem_t* __restrict__ kv, + const int* __restrict__ bt, float* __restrict__ p_acc, + float* __restrict__ p_m, float* __restrict__ p_l, int* __restrict__ p_cnt, + OutT* __restrict__ out, const int* __restrict__ seq_lens, int bt_width, + float scale, int coop + #if BATCH + , + Batch bs + #endif +) { + // _rocm_C is built for every gfx11 target; the code is gfx1151's alone, and + // the registry offers the variants on gfx1151 alone. + #if !defined(__HIP_DEVICE_COMPILE__) || defined(__gfx1151__) + __shared__ __attribute__((aligned(16))) char lds[decomp::kLds]; + #if BATCH + const int b = blockIdx.y; + constexpr int kTok = MAXM * NUM_Q_HEADS * HEAD_DIM; + q += (size_t)b * kTok; + out += (size_t)b * kTok; + bt += (size_t)b * bs.bt_stride; + p_acc += (size_t)b * bs.acc_stride; + p_m += (size_t)b * bs.ml_stride; + p_l += (size_t)b * bs.ml_stride; + p_cnt += (size_t)b * bs.cnt_stride; + seq_lens += b; + #endif + // S and the first page indices go out together: neither waits for the + // other. Every KV address depends on the page table. + const int S = __builtin_amdgcn_readfirstlane(*seq_lens); + const int pages = decomp::first_pages(bt, bt_width); + #if BATCH + // A CUDA-graph batch padded past its real sequences: S = 0, nothing to + // read or write. Only in batch builds: waiting for S before the body's + // first loads costs 1-3 % at one sequence. + if (S <= 0) return; + #endif + decomp::body(q, kv, bt, p_acc, p_m, p_l, p_cnt, out, S, pages, bt_width, + scale, coop, lds); + #endif +} + + #if !defined(RDNA35_TORCH_EXT) && !defined(RDNA35_VLLM) +// Instantiated for ISA inspection when built without the torch op. +template __global__ void decode_attn(const elem_t*, const elem_t*, + const int*, float*, float*, float*, + int*, elem_t*, const int*, int, + float, int + #if BATCH + , + Batch + #endif +); + #else +// One call, arguments already checked: by decode_attn_op below in a JIT +// build, by the registry (rdna35_attn/registry.hip) in _rocm_C's. +// Static: JIT builds of different variants share a process, and a global +// symbol would let one variant's op resolve to another's launch. +static void launch(const rdna35::LaunchArgs& a) { + // Row group fastest, then kv head, then segment: the row groups of one kv + // head read the same KV and are dispatched side by side to share it in L2. + dim3 grid(decomp::kGrid, a.nseq), block(BLOCK); + static const int resident = [&] { + if (!decomp::kSharedMerge) return 0; + int dev, wgps, per = 0; + (void)hipGetDevice(&dev); + (void)hipDeviceGetAttribute(&wgps, hipDeviceAttributeMultiprocessorCount, + dev); + (void)hipOccupancyMaxActiveBlocksPerMultiprocessor( + &per, decode_attn, BLOCK, 0); + // The runtime assumes 64 KiB of LDS where a gfx1151 WGP has 128, and + // reports one workgroup per WGP for anything over 32 KiB. Count what + // fits, rounding every way down: 1536 VGPRs per SIMD in blocks of 24, + // and no workgroup spanning the WGP's two CUs. + hipFuncAttributes fa; + hipDeviceProp_t prop; + if (hipFuncGetAttributes(&fa, reinterpret_cast( + decode_attn)) == hipSuccess && + hipGetDeviceProperties(&prop, dev) == hipSuccess && + strstr(prop.gcnArchName, "gfx1151")) { + const int wps = std::min(16, 1536 / ((fa.numRegs + 23) / 24 * 24)); + const int by_waves = 2 * (2 * wps / (BLOCK / WAVE)); + const int by_lds = + fa.sharedSizeBytes ? 128 * 1024 / (int)fa.sharedSizeBytes : by_waves; + per = std::max(per, std::min(by_waves, by_lds)); + } + return per * wgps; + }(); + // The shared merge waits across workgroups: only when all of them fit. + int coop = (int)(resident >= (int)(grid.x * grid.y)); + #if BATCH + // Segments per sequence, capped so the batch lands near BTARGET + // workgroups (the WMMA body reads it from coop's upper bits). + const int groups = decomp::kGrid / NSEG; + coop |= std::max(1, (BTARGET + a.nseq * groups - 1) / (a.nseq * groups)) << 1; + const Batch bs{a.bt_stride, a.acc_stride, a.ml_stride, a.cnt_stride}; + #endif + hipLaunchKernelGGL(decode_attn, grid, block, 0, a.stream, + static_cast(a.q), + static_cast(a.kv), a.bt, a.acc, a.m, a.l, + a.cnt, static_cast(a.out), a.seq_lens, a.bt_width, + a.scale, coop + #if BATCH + , + bs + #endif + ); +} + #endif + + #ifdef RDNA35_TORCH_EXT + #include + #include + #include + +// The scratch buffers (acc/m/l/cnt) are caller-allocated on purpose: this runs +// inside a CUDA-graph capture, and an allocation there would break it. +void decode_attn_op(torch::Tensor& q, torch::Tensor& kv_cache, + torch::Tensor& block_table, torch::Tensor& out, + torch::Tensor& acc, torch::Tensor& m, torch::Tensor& l, + torch::Tensor& cnt, torch::Tensor& seq_lens, double scale) { + TORCH_CHECK(q.is_contiguous() && out.is_contiguous(), + "q and out must be contiguous"); + TORCH_CHECK(block_table.scalar_type() == torch::kInt32 && + (block_table.dim() == 1 || block_table.dim() == 2) && + block_table.stride(-1) == 1, + "block table must be int32 rows, contiguous within a row"); + // A tensor, not an int: under CUDA-graph capture a host argument is frozen + // at its capture-time value, and every replay would attend over that many + // keys whatever the sequence has grown to. + TORCH_CHECK(seq_lens.scalar_type() == torch::kInt32 && seq_lens.numel() >= 1, + "seq_lens must be an int32 tensor holding the sequence length"); + // A batch of sequences with MAXM query tokens each. + TORCH_CHECK(q.size(0) % MAXM == 0 && q.size(0) > 0, "q has ", q.size(0), + " tokens, not a multiple of MAXM=", MAXM); + const int nseq = q.size(0) / MAXM; + TORCH_CHECK(seq_lens.numel() >= nseq, "seq_lens has ", seq_lens.numel(), + " entries for ", nseq, " sequences"); + TORCH_CHECK(block_table.dim() == 2 ? block_table.size(0) >= nseq : nseq == 1, + "block table has too few rows for ", nseq, " sequences"); + // Scratch is (sequences, ...) or, for one sequence, without that dim. + const bool batched = acc.dim() == 3; + TORCH_CHECK(batched ? acc.size(0) >= nseq && m.size(0) >= nseq && + l.size(0) >= nseq && cnt.size(0) >= nseq + : nseq == 1, + "scratch holds fewer sequences than the ", nseq, " launched"); + TORCH_CHECK(BATCH || nseq == 1, "one-sequence build launched for ", nseq, + " sequences"); + TORCH_CHECK(q.size(1) == NUM_Q_HEADS && q.size(2) == HEAD_DIM, + "q shape does not match the compiled variant"); + // The loads reinterpret every tensor as the element type the variant was + // built for. The other 16-bit type would read the same bits and return + // finite nonsense, so refuse it here rather than downstream. + constexpr auto dtype = KV_BF16 ? at::kBFloat16 : at::kHalf; + TORCH_CHECK(q.scalar_type() == dtype && kv_cache.scalar_type() == dtype && + out.scalar_type() == dtype, + "kernel built for ", dtype, ", got q=", q.scalar_type(), + " kv_cache=", kv_cache.scalar_type(), " out=", out.scalar_type()); + const at::cuda::OptionalCUDAGuard device_guard(device_of(q)); + launch({q.data_ptr(), kv_cache.data_ptr(), block_table.data_ptr(), + acc.data_ptr(), m.data_ptr(), l.data_ptr(), + cnt.data_ptr(), out.data_ptr(), seq_lens.data_ptr(), + (int)block_table.size(-1), + block_table.dim() == 2 ? (int)block_table.stride(0) : 0, + batched ? (int)acc.stride(0) : 0, batched ? (int)m.stride(0) : 0, + batched ? (int)cnt.stride(0) : 0, nseq, (float)scale, + at::cuda::getCurrentCUDAStream()}); +} + + #if TIMING +torch::Tensor timings() { + auto t = torch::zeros({8 * 16}, torch::kInt64); + hipMemcpyFromSymbol(t.data_ptr(), HIP_SYMBOL(g_ts), sizeof(g_ts)); + return t; +} + #endif + +PYBIND11_MODULE(TORCH_EXTENSION_NAME, mod) { + mod.def("decode_attn", &decode_attn_op, "RDNA3.5 GQA paged decode attention"); + #if TIMING + mod.def("timings", &timings); + #endif +} + #endif + + #ifdef RDNA35_VLLM + // _rocm_C builds many variants into one translation unit, each in a + // namespace of its own (rdna35_attn/instance.h), so the next one must + // start from none of this one's definitions. The body undefines its own. + #undef RDNA35_BODY + #undef M_NSEG + #undef M_RG + #undef M_MINB + #undef M_DSPL + #undef M_RSPL + #undef M_NW + #undef M_PF + #undef M_DOT + #undef M_BFLY + #undef M_GT + #undef WAVE + #undef BLOCK + #undef GQA + #undef ROWS + #undef NCHUNK + #undef KV_ROW + #undef PAGE_ELEMS + #undef KEY_STRIDE + #undef LOG2E + #undef QPAD + #undef PADM + #undef MBUDGET + #undef TS + #undef TS_WG0 + #undef DSPL + #endif + +#else + // ---------------------------------------------------------------- body ---- + // Everything below depends on the decomposition knobs, read through their + // M_ names; the outer pass instantiates it in namespace decomp. DOT picks + // the per-q-head dot-product decomposition instead of WMMA, for the whole + // configuration. + #if M_DOT + // ------------------------------------------------------- dot body ---- + // The per-q-head dot-product decomposition: one workgroup per + // (q head, segment), every wave a slice of + // KPW keys at a time with the whole head dim across its lanes, Q@K as fdot2 + // and a butterfly across the lanes of a row, P@V on the VALU. No WMMA + // padding and no score exchange, which is what wins at short context: few + // keys per kv head leave the WMMA path's 16-row tiles mostly empty and its + // cross-wave traffic exposed. Each kv head's KV is read once per q head, + // through L2. + #define D_DPL (HEAD_DIM == 128 ? 8 : HEAD_DIM / WAVE) + #define D_LPR (HEAD_DIM / D_DPL) + #define D_SUB (WAVE / D_LPR) + #define D_KPW 4 + #define D_KPWE ((D_KPW / D_SUB) > 0 ? (D_KPW / D_SUB) : 1) + #define D_FPL (D_DPL / 2) + #define D_SLICE (D_DPL + 1) + #define D_STRIDE (D_LPR * D_SLICE) + #define D_NPARTW (M_NW * D_SUB) + #define D_ACC_BYTES (MAXM * M_NW * D_SUB * D_STRIDE * 4) + #define D_ML_BYTES (MAXM * M_NW * D_SUB * 4) + +static_assert(D_LPR >= 1 && D_LPR <= WAVE && WAVE % D_LPR == 0, + "a row must be covered by a whole divisor of the wave"); +static_assert(BS % (D_SUB * D_KPWE) == 0, + "a tile of SUB*KPWE tokens must not straddle two blocks"); +static_assert(D_ACC_BYTES + 2 * D_ML_BYTES + 16 <= 65536, + "the dot mode's partials must fit LDS"); + +typedef float fvec __attribute__((ext_vector_type(D_FPL))); + +constexpr int kLds = D_ACC_BYTES + 2 * D_ML_BYTES + 16; +constexpr int kGrid = M_NSEG * NUM_Q_HEADS; +// One arrival counter per q head. +constexpr int kCnt = NUM_Q_HEADS; +constexpr bool kSharedMerge = false; + +__device__ __forceinline__ void dot_coords(int& seg, int& h) { + const int bid = blockIdx.x; + #if M_GT + // Head fastest: the q heads of a kv head, which read byte-identical KV, + // are dispatched side by side. + h = bid % NUM_Q_HEADS; + seg = bid / NUM_Q_HEADS; + #else + seg = bid % M_NSEG; + h = bid / M_NSEG; + #endif +} + +// The block index of wave's first tile, which goes out before S is known. +__device__ __forceinline__ int first_pages(const int* __restrict__ bt, + int bt_width) { + int seg, h; + dot_coords(seg, h); + const int gw = + __builtin_amdgcn_readfirstlane(seg * M_NW + threadIdx.x / WAVE); + return bt[min(gw * D_KPWE * D_SUB / BS, bt_width - 1)]; +} + +__device__ __forceinline__ float dot2(float a, float b, float c) { + #if KV_BF16 + return __builtin_amdgcn_fdot2_f32_bf16(__builtin_bit_cast(s2v, a), + __builtin_bit_cast(s2v, b), c, false); + #else + return __builtin_amdgcn_fdot2(__builtin_bit_cast(e2, a), + __builtin_bit_cast(e2, b), c, false); + #endif +} + +__device__ __forceinline__ int partial_row(int wb, int p) { + return wb * D_SUB + p; +} + + // BFLY routes each butterfly element to the LDS pipe (ds_bpermute) or the + // VALU (permlane), in quarters. + #define BFLY_ON_VALU(idx) (((idx) & 3) < M_BFLY) + #define BFLY_LO(st) \ + ((st) == 1 ? 0x67452301u \ + : (st) == 2 ? 0x54761032u \ + : (st) == 4 ? 0x32107654u \ + : 0xFEDCBA98u) + #define BFLY_HI(st) \ + ((st) == 1 ? 0xEFCDAB89u \ + : (st) == 2 ? 0xDCFE98BAu \ + : (st) == 4 ? 0xBA98FEDCu \ + : 0x76543210u) + #define BFLY_VALU(st, v) \ + ((st) >= 16 \ + ? __builtin_amdgcn_permlanex16( \ + __builtin_bit_cast(int, v), __builtin_bit_cast(int, v), \ + 0x76543210u, 0xFEDCBA98u, false, false) \ + : __builtin_amdgcn_permlane16( \ + __builtin_bit_cast(int, v), __builtin_bit_cast(int, v), \ + BFLY_LO(st), BFLY_HI(st), false, false)) + #define BFLY_XOR(st, addr, idx, v) \ + (BFLY_ON_VALU(idx) \ + ? BFLY_VALU(st, v) \ + : __builtin_amdgcn_ds_bpermute(addr, __builtin_bit_cast(int, v))) + +template +__device__ __forceinline__ void body( + const elem_t* __restrict__ q, const elem_t* __restrict__ kv, + const int* __restrict__ bt, float* __restrict__ p_acc, + float* __restrict__ p_m, float* __restrict__ p_l, int* __restrict__ p_cnt, + OutT* __restrict__ out, const int S, int pages, int bt_width, float scale, + int coop, char* __restrict__ lds) { + (void)coop; + if (blockIdx.x >= kGrid) return; + const int tid = threadIdx.x; + const int wave = __builtin_amdgcn_readfirstlane(tid / WAVE); + int seg, h; + dot_coords(seg, h); + const int lane = tid & (WAVE - 1); + const int lrow = lane % D_LPR; // this lane's slice of the row + const int grp = lane / D_LPR; // which of the SUB keys it carries + const int dl = lrow * D_DPL; + const int kvh = h / GQA; + + fvec qr[MAXM]; + #pragma unroll + for (int t = 0; t < MAXM; ++t) { + const unsigned qoff = + ((unsigned)t * NUM_Q_HEADS + (unsigned)h) * HEAD_DIM + (unsigned)dl; + qr[t] = *(const fvec*)(q + qoff); + } + + float acc[MAXM][D_DPL], mx[MAXM], ls[MAXM]; + #pragma unroll + for (int t = 0; t < MAXM; ++t) { + #pragma unroll + for (int i = 0; i < D_DPL; ++i) acc[t][i] = 0.f; + mx[t] = -INFINITY; + ls[t] = 0.f; + } + + __builtin_assume(S > 0); + const float scale2 = scale * LOG2E; + const int ctx = S - MAXM; + // readfirstlane: without it the compiler cannot see the start is + // wave-uniform and guards the loop with an exec mask, so a wave with no + // tiles reached the LDS stores with exec = 0 (reference kernel, NSEG > 1). + const int gw = __builtin_amdgcn_readfirstlane(seg * M_NW + wave); + #if WIN + // Tiles start at the first key some query can see, on a tile boundary so no + // tile straddles two pages; a window wider than S starts at 0. + const int kbase = + max(0, S - MAXM - (WIN - 1)) / (D_KPWE * D_SUB) * (D_KPWE * D_SUB); + #else + constexpr int kbase = 0; + #endif + const int jstart = kbase + gw * D_KPWE * D_SUB; + const int jstep = M_NSEG * M_NW * D_KPWE * D_SUB; + // The first tile's block came in with S; each later one is read a tile + // ahead, so no tile's loads wait on the block table. With a window the + // first tile depends on S, so its block is read now. + #if WIN + int blk = __builtin_amdgcn_readfirstlane(bt[min(jstart / BS, bt_width - 1)]); + (void)pages; + #else + int blk = __builtin_amdgcn_readfirstlane(pages); + #endif + + for (int jb = jstart; jb < S; jb += jstep) { + fvec kr[D_KPWE], vr[D_KPWE]; + __builtin_assume(blk >= 0); + #if LAYOUT == 1 + const size_t base = (size_t)blk * PAGE_ELEMS + + ((size_t)kvh * BS + (unsigned)jb % BS) * KV_ROW + + (size_t)grp * KV_ROW + dl; + #pragma unroll + for (int c = 0; c < D_KPWE; ++c) { + const size_t off = base + (size_t)c * D_SUB * KV_ROW; + kr[c] = *(const fvec*)(kv + off); + vr[c] = *(const fvec*)(kv + off + HEAD_DIM); + } + #else + #pragma unroll + for (int c = 0; c < D_KPWE; ++c) { + const unsigned jj = (unsigned)min(jb + c * D_SUB + grp, S - 1); + const size_t off = (size_t)blk * PAGE_ELEMS + + ((size_t)(jj % BS) * NUM_KV_HEADS + kvh) * KV_ROW + dl; + kr[c] = *(const fvec*)(kv + off); + vr[c] = *(const fvec*)(kv + off + HEAD_DIM); + } + #endif + blk = __builtin_amdgcn_readfirstlane( + bt[min((jb + jstep) / BS, bt_width - 1)]); + // Keys past S read whatever their page holds; K's garbage dies in the + // mask, V's must not reach the accumulator as 0 * NaN. + if (__builtin_expect(jb + D_KPWE * D_SUB > S, 0)) { + #pragma unroll + for (int c = 0; c < D_KPWE; ++c) + if (jb + c * D_SUB + grp >= S) vr[c] = fvec{}; + } + + float s[D_KPWE][MAXM]; + #pragma unroll + for (int c = 0; c < D_KPWE; ++c) + #pragma unroll + for (int t = 0; t < MAXM; ++t) { + float d = 0.f; + #pragma unroll + for (int e = 0; e < D_FPL; ++e) d = dot2(qr[t][e], kr[c][e], d); + s[c][t] = d; + } + #pragma unroll + for (int st = 1; st < D_LPR; st <<= 1) { + const int addr = (lane ^ st) << 2; + #pragma unroll + for (int c = 0; c < D_KPWE; ++c) + #pragma unroll + for (int t = 0; t < MAXM; ++t) + s[c][t] = __builtin_bit_cast( + float, BFLY_XOR(st, addr, c * MAXM + t, s[c][t])) + + s[c][t]; + } + #pragma unroll + for (int c = 0; c < D_KPWE; ++c) { + const int jj = jb + c * D_SUB + grp; + #pragma unroll + for (int t = 0; t < MAXM; ++t) { + #if MUTATE == 1 + const bool valid = (jj <= ctx + t + 1); + #else + const bool valid = (jj <= ctx + t); + #endif + #if WIN + s[c][t] = (jj < ctx + t - (WIN - 1)) ? -INFINITY : s[c][t]; + #endif + s[c][t] = valid ? s[c][t] * scale2 : -INFINITY; + } + } + #pragma unroll + for (int t = 0; t < MAXM; ++t) { + float mnew = mx[t]; + #pragma unroll + for (int c = 0; c < D_KPWE; ++c) mnew = fmaxf(mnew, s[c][t]); + const float alpha = + (mnew == -INFINITY) ? 0.f : __builtin_amdgcn_exp2f(mx[t] - mnew); + mx[t] = mnew; + ls[t] *= alpha; + #pragma unroll + for (int i = 0; i < D_DPL; ++i) acc[t][i] *= alpha; + float lsum = 0.f; + #pragma unroll + for (int c = 0; c < D_KPWE; ++c) { + const float p = + (mnew == -INFINITY) ? 0.f : __builtin_amdgcn_exp2f(s[c][t] - mnew); + s[c][t] = p; + lsum += p; + } + ls[t] += lsum; + } + #pragma unroll + for (int c = 0; c < D_KPWE; ++c) { + const elem_t* vv = (const elem_t*)&vr[c]; + #pragma unroll + for (int t = 0; t < MAXM; ++t) + #pragma unroll + for (int i = 0; i < D_DPL; ++i) acc[t][i] += s[c][t] * (float)vv[i]; + } + } + + // The waves' partials meet in LDS, then (with segments) in global memory. + float* const lds_acc = reinterpret_cast(lds); + float* const lds_m = reinterpret_cast(lds + D_ACC_BYTES); + float* const lds_l = lds_m + MAXM * M_NW * D_SUB; + int& lds_last = *reinterpret_cast(lds_l + MAXM * M_NW * D_SUB); + #pragma unroll + for (int t = 0; t < MAXM; ++t) { + if (lrow == 0) { + lds_m[(t * M_NW + wave) * D_SUB + grp] = mx[t]; + lds_l[(t * M_NW + wave) * D_SUB + grp] = ls[t]; + } + #pragma unroll + for (int i = 0; i < D_DPL; ++i) + lds_acc[((t * M_NW + wave) * D_SUB + grp) * D_STRIDE + lrow * D_SLICE + + i] = acc[t][i]; + } + lds_barrier(); + + #pragma unroll + for (int t = 0; t < MAXM; ++t) { + const int wb = t * M_NW; + float gmax = -INFINITY; + #pragma unroll + for (int p = 0; p < D_NPARTW; ++p) + gmax = fmaxf(gmax, lds_m[partial_row(wb, p)]); + float a[D_NPARTW], den = 0.f; + #pragma unroll + for (int p = 0; p < D_NPARTW; ++p) { + const int w = partial_row(wb, p); + a[p] = weight_of(lds_m[w], gmax); + den = fmaf(a[p], lds_l[w], den); + } + const size_t pb = ((size_t)h * M_NSEG + seg) * MAXM + t; + #if M_NSEG > 1 + if (tid == 0) { + p_m[pb] = gmax; + p_l[pb] = den; + } + #endif + for (int d = tid; d < HEAD_DIM; d += M_NW * WAVE) { + float num = 0.f; + #pragma unroll + for (int p = 0; p < D_NPARTW; ++p) + num = fmaf(a[p], + lds_acc[partial_row(wb, p) * D_STRIDE + + (d / D_DPL) * D_SLICE + d % D_DPL], + num); + #if M_NSEG == 1 + out[((size_t)t * NUM_Q_HEADS + h) * HEAD_DIM + d] = + to_elem(num * __builtin_amdgcn_rcpf(den)); + #else + p_acc[pb * HEAD_DIM + d] = num; + #endif + } + } + + #if M_NSEG > 1 + // Last arriver for this head merges; it resets the counter for the next + // launch, so a graph replay starts clean. + __threadfence(); + if (tid == 0) lds_last = (atomicAdd(&p_cnt[h], 1) == M_NSEG - 1); + lds_barrier(); + if (!lds_last) return; + if (tid == 0) p_cnt[h] = 0; + __threadfence(); + #pragma unroll + for (int t = 0; t < MAXM; ++t) { + const size_t rb = (size_t)h * M_NSEG * MAXM + t; + float gmax = -INFINITY; + #pragma unroll + for (int s = 0; s < M_NSEG; ++s) + gmax = fmaxf(gmax, p_m[rb + (size_t)s * MAXM]); + for (int d = tid; d < HEAD_DIM; d += M_NW * WAVE) { + float num = 0.f, den = 0.f; + #pragma unroll + for (int s = 0; s < M_NSEG; ++s) { + const size_t b = rb + (size_t)s * MAXM; + const float a = weight_of(p_m[b], gmax); + den = fmaf(a, p_l[b], den); + num = fmaf(a, p_acc[b * HEAD_DIM + d], num); + } + out[((size_t)t * NUM_Q_HEADS + h) * HEAD_DIM + d] = + to_elem(num * __builtin_amdgcn_rcpf(den)); + } + } + #endif +} + + #undef D_DPL + #undef D_LPR + #undef D_SUB + #undef D_KPW + #undef D_KPWE + #undef D_FPL + #undef D_SLICE + #undef D_STRIDE + #undef D_NPARTW + #undef D_ACC_BYTES + #undef D_ML_BYTES + #undef BFLY_ON_VALU + #undef BFLY_LO + #undef BFLY_HI + #undef BFLY_VALU + #undef BFLY_XOR + #else + #define ROWS_W (ROWS / M_RG) + #define RTILES ((ROWS_W + 15) / 16) + #define ROWPAD (RTILES * 16) + #define TILES (M_NW / (M_DSPL * M_RSPL)) + // Row tiles per wave; wave rs of a tile carries row tiles rs, rs + M_RSPL, + // ... + #define RTW (RTILES / M_RSPL) + // With M_RSPL > 1 the tile is shared: each of its M_RSPL waves loads EPW of + // the eight key pairs, stages them in LDS, and all read the whole tile + // back. + #define SHT (M_RSPL > 1) + #define EPW (8 / M_RSPL) + #define KBLK (16 * TILES) + #define DPART (HEAD_DIM / M_DSPL) + // d per lane of a K or V row: lane l16 owns DPART/16 consecutive d. + #define VD (DPART / 16) + #define NCP (NCHUNK / M_DSPL) + // P reaches P@V as a high plus a low half, ~22 bits rather than 11 in fp16 + // and ~16 rather than 8 in bf16: with one fp16 P the output misses the 1e-3 + // relative bound wherever it is near zero. When a row tile has at most 8 + // real rows the low half rides in the padding columns of the same WMMA; + // otherwise it takes a second one. + #define PPACK (ROWS_W <= 8) + + // Per-wave K tile: 16 keys of DPART halves. K is read row-wise, like V, + // and transposed to the lane-per-key WMMA operand through it; the 16-byte + // pad puts the sixteen keys one operand read touches on distinct banks. + #define KT_ROW (DPART * 2 + 16) + #define KT_BYTES (M_NW * (VINLDS ? 2 : 1) * 16 * KT_ROW) + #define QS_BYTES (ROWPAD * (HEAD_DIM + QPAD) * 2) + #define MROW (DPART + PADM) + #define RV (ROWS_W < 16 ? ROWS_W : 16) + #define MSLOT (RV * MROW) + #define SLOTS_FIT(t) ((t) * M_DSPL * M_RSPL * MSLOT * 4 <= MBUDGET) + // Live tiles after the merge's tree rounds: the most that fit the budget. + #define TFIN \ + (SLOTS_FIT(TILES) ? TILES \ + : SLOTS_FIT(TILES / 2) ? TILES / 2 \ + : SLOTS_FIT(TILES / 4) ? TILES / 4 \ + : SLOTS_FIT(TILES / 8) ? TILES / 8 \ + : 1) + #define MRG_SLOTS \ + (TFIN == TILES ? TILES : (TFIN > TILES / 2 ? TFIN : TILES / 2)) + #define MRG_BYTES \ + (TILES == 1 && M_RSPL == 1 && !CPUB \ + ? 0 \ + : MRG_SLOTS * M_DSPL * M_RSPL * MSLOT * 4) + #if SHT + // Per (tile, d part) a K and a V tile, double-buffered, and the score + // exchange apart from them: partners may still read the shared K. + #define TB_BYTES (2 * TILES * M_DSPL * 2 * 16 * KT_ROW) + #define SX_BYTES (M_DSPL > 1 ? M_NW * RTW * 1024 : 0) + #define LOOP_BYTES (QS_BYTES + TB_BYTES + SX_BYTES) + #else + #define LOOP_BYTES (QS_BYTES + KT_BYTES) + #endif + #define LDS_RAW (LOOP_BYTES > MRG_BYTES ? LOOP_BYTES : MRG_BYTES) + // Split KV: a single workgroup merging every segment's partials is bound by + // its own CU's bandwidth once they reach 64 KiB -- 2 us at 256 KiB. Past + // that the segments wait for each other and each merges a slice instead. + #define SHARED_MERGE (ROWS_W * HEAD_DIM * M_NSEG * 4 >= 64 * 1024) + #define M_BLOCK (M_NW * WAVE) + // m_s, l_s, gm_s, gl_s and the merge's go flag, after the raw buffer. + #define LDS_MS ((LDS_RAW + 15) / 16 * 16) + #define LDS_ALL (LDS_MS + (2 * M_NW * ROWPAD + 2 * ROWPAD) * 4 + 16) + +static_assert(GQA % M_RG == 0, "row groups split whole q heads"); +static_assert(M_NW % (M_DSPL * M_RSPL) == 0, "whole tiles per workgroup"); +static_assert(RTILES % M_RSPL == 0, "whole row tiles per wave"); +static_assert(8 % M_RSPL == 0, "whole key pairs per wave of a shared tile"); +static_assert(NCHUNK % M_DSPL == 0, "whole chunks per d part"); +static_assert(VD == 4 || VD == 8, "a lane's K and V slice is one b64 or b128"); +static_assert(TILES <= WAVE, "one lane per tile loads its page index"); +static_assert(M_DSPL == 1 || SHT || RTW * 1024 <= 16 * KT_ROW, + "the score exchange must fit a wave's K tile"); + + #if VD == 8 +typedef u4v vrow_t; + #else +typedef u2v vrow_t; + #endif + +// One wave's share of a tile: half hi of the wave holds keys 2e + hi, lane +// l16 VD contiguous d of each, for K and for V. WMMA wants the two halves of +// a wave to hold the same operand; loading different keys into them and +// exchanging halves with one permlanex16 per dword halves both the registers +// and the loads. +struct Tile { + vrow_t k[8 / M_RSPL]; + vrow_t v[8 / M_RSPL]; +}; + +// Page index of every 16-key tile of block b: lane t holds tile t's. Tiles +// past the sequence clamp to its last page, so every address stays inside a +// page vLLM allocated. +__device__ __forceinline__ int block_pages(const int* __restrict__ bt, int b, + int lane, int S) { + const int key = b * KBLK + (lane % TILES) * 16; + // One tile: a uniform index, which then comes from the scalar cache. + if (TILES == 1) + return bt[__builtin_amdgcn_readfirstlane(min(key, S - 1)) / BS]; + return bt[(unsigned)min(key, S - 1) / BS]; +} + +// The same before S is known: clamped to the block table's width instead. +// Every entry of the table names an allocated page, so the address stays in +// the cache, and the keys past S are masked like any others. +__device__ __forceinline__ int block_pages_w(const int* __restrict__ bt, int b, + int lane, int width) { + const int blk = (b * KBLK + (lane % TILES) * 16) / BS; + if (TILES == 1) + return bt[__builtin_amdgcn_readfirstlane(min(blk, width - 1))]; + return bt[min(blk, width - 1)]; +} + +// A = V^T for output element j of this lane's slice: row d = base + VD*l16 + +// j, keys in the order above. Half hi holds keys 2e + hi in v[e]; it packs +// its four dwords and receives the other half's four. +__device__ __forceinline__ e16 v_frag(const vrow_t* v, int j) { + const int w = j >> 1; + // v_perm_b32(hi_src, lo_src, sel): half j&1 of each source dword, lo_src + // into the low half of the result. + const unsigned sel = (j & 1) ? 0x07060302u : 0x05040100u; + u8v r; + #pragma unroll + for (int a = 0; a < 4; ++a) { + r[a] = __builtin_amdgcn_perm(v[2 * a + 1][w], v[2 * a][w], sel); + r[4 + a] = xhalf_u(r[a]); + } + return __builtin_bit_cast(e16, r); +} + +// A wave's real rows of one row tile, rows `stride` floats apart. Lane l16 +// is the row; for element e it owns VD consecutive d at VD*(2e + hi). +__device__ __forceinline__ void store_rows(float* dst, const f8* acc, int l16, + int hi, int stride = MROW) { + float* row = dst + l16 * stride; + #pragma unroll + for (int e = 0; e < 8; ++e) { + float v[VD]; + #pragma unroll + for (int j = 0; j < VD; ++j) v[j] = acc[j][e]; + #pragma unroll + for (int k = 0; k < VD; k += 4) + *(f4*)(row + VD * (2 * e + hi) + k) = *(const f4*)(v + k); + } +} + +__device__ __forceinline__ void add_rows(f8* acc, const float* src, int l16, + int hi) { + const float* row = src + l16 * MROW; + #pragma unroll + for (int e = 0; e < 8; ++e) + #pragma unroll + for (int k = 0; k < VD; k += 4) { + const f4 v = *(const f4*)(row + VD * (2 * e + hi) + k); + #pragma unroll + for (int j = 0; j < 4; ++j) acc[k + j][e] += v[j]; + } +} + +constexpr int kLds = LDS_ALL; +constexpr int kGrid = M_NSEG * NUM_KV_HEADS * M_RG; +// Arrival counters, then merge generations, one each per (kv head, row group). +constexpr int kCnt = 2 * NUM_KV_HEADS * M_RG; +constexpr bool kSharedMerge = SHARED_MERGE; + +// The first block's page indices, which go out before S is known. +__device__ __forceinline__ int first_pages(const int* __restrict__ bt, + int bt_width) { + const int seg = blockIdx.x / (M_RG * NUM_KV_HEADS); + return block_pages_w(bt, seg, threadIdx.x & (WAVE - 1), bt_width); +} + +template +__device__ __forceinline__ void body( + const elem_t* __restrict__ q, const elem_t* __restrict__ kv, + const int* __restrict__ bt, float* __restrict__ p_acc, + float* __restrict__ p_m, float* __restrict__ p_l, int* __restrict__ p_cnt, + OutT* __restrict__ out, const int S, int pages, int bt_width, float scale, + int coop, char* __restrict__ lds) { + const int tid = threadIdx.x; + const int lane = tid & (WAVE - 1); + const int wave = __builtin_amdgcn_readfirstlane(tid / WAVE); + const int l16 = lane & 15; + const int hi = lane >> 4; + const int tw = wave / (M_DSPL * M_RSPL); // this wave's tile within the block + const int rs = (wave / M_DSPL) % M_RSPL; // its share of the row tiles + const int dp = wave % M_DSPL; // and its part of the head dim + // Workgroup row tile of this wave's row tile rt. + #define RT(rt) ((rt) * M_RSPL + rs) + + const int bid = blockIdx.x; + const int rg = bid % M_RG; + const int kvh = (bid / M_RG) % NUM_KV_HEADS; + const int seg = bid / (M_RG * NUM_KV_HEADS); + + __builtin_assume(S > 0); + const int nblk = (S + KBLK - 1) / KBLK; + #if WIN + // The first block holding a key some query can see. Clamped at 0: a window + // wider than the sequence reads it all and nothing before it. + const int b0 = max(0, S - MAXM - (WIN - 1)) / KBLK; + #else + constexpr int b0 = 0; + #endif + #if BATCH + // A batch brings its own parallelism: the host caps the segments so the + // whole grid does not pay a merge per sequence it no longer needs. + const int ncap = coop >> 1; + coop &= 1; + const int nseg = max(1, min(min(M_NSEG, ncap), (nblk - b0) / M_MINB)); + #else + const int nseg = max(1, min(M_NSEG, (nblk - b0) / M_MINB)); + #endif + if (seg >= nseg) return; + #if WIN + // The page indices that came in with S were for block seg, not b0 + seg. + pages = block_pages(bt, b0 + seg, lane, S); + #endif + TS(0); + if (ABLATE & 32) { + if (S == -1) out[0] = (OutT)0; + return; + } + + // Q goes out before the KV. Load completion is counted in order, so a Q + // load issued behind the KV tiles would hold its LDS store, and the barrier + // behind it, until the last KV byte landed. + constexpr int QITER = (ROWPAD * HEAD_DIM + M_BLOCK * 8 - 1) / (M_BLOCK * 8); + e8 qx[QITER]; + #pragma unroll + for (int it = 0; it < QITER; ++it) { + const int i = (it * M_BLOCK + tid) * 8; + // Unconditional, from a clamped row: a load under a branch made the + // compiler wait vmcnt(0) -- for Q as well as the page table -- before it + // could issue the first KV load. + const int r = min(i / HEAD_DIM, ROWS_W - 1), d = i % HEAD_DIM; + const int h = kvh * GQA + rg * (GQA / M_RG) + r / MAXM; + const int m = r % MAXM; + qx[it] = *(const e8*)(q + ((size_t)m * NUM_Q_HEADS + h) * HEAD_DIM + d); + if (i / HEAD_DIM >= ROWS_W) qx[it] = e8{}; + } + // The group's merge generation cannot move before this segment arrives, so + // it is read now rather than on the way to the arrival. + int* const gen = p_cnt + NUM_KV_HEADS * M_RG + kvh * M_RG + rg; + const int g0 = SHARED_MERGE ? __hip_atomic_load(gen, __ATOMIC_RELAXED, + __HIP_MEMORY_SCOPE_AGENT) + : 0; + + // q_s and the K tiles are dead once the main loop ends; the merge reuses + // their storage. + char* const lds_raw = lds; + auto q_s = reinterpret_cast(lds_raw); + float* mrg_s = reinterpret_cast(lds_raw); + #if SHT + // This wave's tile group's K tile in buffer `buf`; its V tile follows. + char* const tb_s = lds_raw + QS_BYTES + (tw * M_DSPL + dp) * 2 * 16 * KT_ROW; + #define KT_BUF(buf) (tb_s + (buf) * (TB_BYTES / 2)) + float* sx_s = reinterpret_cast(lds_raw + QS_BYTES + TB_BYTES); + #define SX_SLOT(w) ((w) * RTW * 256) + #else + static_assert(!VINLDS || !M_PF, "VINLDS replaces PF's second tile"); + char* kt_s = lds_raw + QS_BYTES + wave * (VINLDS ? 2 : 1) * 16 * KT_ROW; + // Wave w's partial scores go in wave w's own K tile. It writes them only + // after its own Q@K has read that tile, its partners read them between the + // two barriers of the exchange, and it refills the tile with the next K + // only after the second barrier. + float* sx_s = reinterpret_cast(lds_raw + QS_BYTES); + #define SX_SLOT(w) ((w) * (VINLDS ? 2 : 1) * (16 * KT_ROW / 4)) + #endif + auto m_s = reinterpret_cast(lds + LDS_MS); + auto l_s = + reinterpret_cast(lds + LDS_MS + M_NW * ROWPAD * 4); + float* const gm_s = + reinterpret_cast(lds + LDS_MS + 2 * M_NW * ROWPAD * 4); + float* const gl_s = gm_s + ROWPAD; + + const elem_t* kvh_base = kv + tile_off(0, 0, kvh); + // This wave's slice of a K or V row, and its chunks of the head dim. + const int col = dp * DPART + VD * l16; + const int c0 = dp * NCP; + + // Issue block b's tile for this wave, and the page lookup of the block + // after it. Blocks are issued in order, so `pages` always holds b's. + auto issue = [&](Tile& t, int b) { + const int page = __builtin_amdgcn_readlane(pages, tw); + #if !SHT + if (b + nseg < nblk) pages = block_pages(bt, b + nseg, lane, S); + #endif + const elem_t* tp = kvh_base + (size_t)page * PAGE_ELEMS + + (size_t)((b * KBLK + tw * 16) % BS) * KEY_STRIDE + + (size_t)hi * KEY_STRIDE + col; + // All of V, then all of K. Interleaving them row by row measured 1.7 to + // 3.3 us slower to land at S=128 (32/32/128), same bytes, same addresses. + // A shared tile's wave loads only its EPW key pairs. + const elem_t* tw_p = tp + (size_t)(SHT ? 2 * rs * EPW : 0) * KEY_STRIDE; + #pragma unroll + for (int e = 0; e < 8 / M_RSPL; ++e) + t.v[e] = (ABLATE & 1) + ? vrow_t{} + : *(const vrow_t*)(tw_p + 2 * e * KEY_STRIDE + HEAD_DIM); + #pragma unroll + for (int e = 0; e < 8 / M_RSPL; ++e) + t.k[e] = + (ABLATE & 4) ? vrow_t{} : *(const vrow_t*)(tw_p + 2 * e * KEY_STRIDE); + // The scheduler must neither sink these to their first use nor hoist + // work above them: left alone under VGPR pressure it issued a tile's K + // two loads at a time with a wait after each. + asm volatile("" ::: "memory"); + #if SHT + // A shared tile's next page indices go out after its loads: ahead of + // them, a vector page-table load held the KV loads behind a vmcnt(0). + // Unconditional (block_pages clamps to S), so the waits can count it. + pages = block_pages(bt, b + nseg, lane, S); + #endif + }; + + Tile ta; + issue(ta, b0 + seg); + + #pragma unroll + for (int it = 0; it < QITER; ++it) { + const int i = (it * M_BLOCK + tid) * 8; + if (i < ROWPAD * HEAD_DIM) *(e8*)&q_s[i / HEAD_DIM][i % HEAD_DIM] = qx[it]; + } + + const float scale2 = scale * LOG2E; + const int ctx = S - MAXM; + + float m_run[RTW], l_run[RTW]; + f8 acc[RTW][VD]; + #pragma unroll + for (int rt = 0; rt < RTW; ++rt) { + m_run[rt] = -INFINITY; + l_run[rt] = 0.f; + #pragma unroll + for (int j = 0; j < VD; ++j) acc[rt][j] = f8{}; + } + // Token index of this lane's row in each row tile, for the causal mask. + int mrow[RTW]; + #pragma unroll + for (int rt = 0; rt < RTW; ++rt) { + const int r = RT(rt) * 16 + l16; + mrow[rt] = (r < ROWS_W) ? (r % MAXM) : (MAXM - 1); + } + + lds_barrier(); + TS(1); + + // Stage a tile's K for Q@K: rows in, keys out. LDS ops of one wave + // complete in order, so no barrier. + // Stored as integers and read back as halves, which type-based alias + // analysis would let pass each other without the asm. + #if SHT + // A shared tile: this wave's key pairs of K and V into buffer kb. + auto stage = [&](const Tile& t, char* kb) { + #pragma unroll + for (int i = 0; i < EPW; ++i) { + const int row = (2 * (rs * EPW + i) + hi) * KT_ROW + VD * l16 * 2; + *(vrow_t*)(kb + row) = t.k[i]; + *(vrow_t*)(kb + 16 * KT_ROW + row) = t.v[i]; + } + asm volatile("" ::: "memory"); + }; + #else + #if VINLDS + auto stage_kv = [&](const Tile& t) { + #pragma unroll + for (int e = 0; e < 8; ++e) { + const int row = (2 * e + hi) * KT_ROW + VD * l16 * 2; + *(vrow_t*)(kt_s + row) = t.k[e]; + *(vrow_t*)(kt_s + 16 * KT_ROW + row) = t.v[e]; + } + asm volatile("" ::: "memory"); + }; + #endif + auto stage_k = [&](const Tile& t) { + #pragma unroll + for (int e = 0; e < 8; ++e) + *(vrow_t*)(kt_s + (2 * e + hi) * KT_ROW + VD * l16 * 2) = t.k[e]; + asm volatile("" ::: "memory"); + }; + #endif + + // Q@K, the softmax update and P@V for one tile whose K is staged at kb. + // v is the tile's V in registers; a shared tile reads it from after K. + auto process = [&](Tile& t, vrow_t* v, const char* kb, int b) { + const int kt = b * KBLK + tw * 16; + #if ABLATE & 64 + unsigned x = 0; + #pragma unroll + for (int e = 0; e < 8 / M_RSPL; ++e) + #pragma unroll + for (int w = 0; w < VD / 2; ++w) x ^= t.v[e][w] ^ t.k[e][w]; + acc[0][0][0] += __builtin_bit_cast(float, x & 0x3f800000u); + return; + #endif + #if TIMING == 3 + if (b == b0 + seg) { + asm volatile("s_waitcnt vmcnt(0)" ::: "memory"); + TS(10); + } + #endif + f8 s[RTW]; + #pragma unroll + for (int rt = 0; rt < RTW; ++rt) s[rt] = f8{}; + #pragma unroll + for (int c = 0; c < NCP; ++c) { + const e16 a = *(const e16*)(kb + l16 * KT_ROW + c * 32); + #pragma unroll + for (int rt = 0; rt < RTW; ++rt) { + const e16 bq = *(const e16*)&q_s[RT(rt) * 16 + l16][(c0 + c) * 16]; + if (!(ABLATE & 16) || c == 0) s[rt] = wmma(a, bq, s[rt]); + } + } + #if M_DSPL > 1 && !(ABLATE & 128) + // The scores overwrite this wave's K tile, which the Q@K above read as + // halves; keep the float stores below those reads. Then sum the tile's + // M_DSPL partials in the same order in every wave, so they all see + // bit-identical S and agree on the softmax. + asm volatile("" ::: "memory"); + #pragma unroll + for (int rt = 0; rt < RTW; ++rt) + *(f8*)&sx_s[SX_SLOT(wave) + rt * 256 + lane * 8] = s[rt]; + lds_barrier(); + #pragma unroll + for (int rt = 0; rt < RTW; ++rt) { + f8 sum = *(const f8*)&sx_s[SX_SLOT(wave - dp) + rt * 256 + lane * 8]; + #pragma unroll + for (int q2 = 1; q2 < M_DSPL; ++q2) + sum += *(const f8*)&sx_s[SX_SLOT(wave - dp + q2) + rt * 256 + lane * 8]; + s[rt] = sum; + } + #if !SHT + // Partners are done reading before anyone refills its K tile. A shared + // tile's barrier before the next Q@K already orders it. + lds_barrier(); + #endif + #endif + if (b == b0 + seg) TS(2); + #if SHT || VINLDS + vrow_t vs[8]; + #pragma unroll + for (int e = 0; e < 8; ++e) + vs[e] = *(const vrow_t*)(kb + 16 * KT_ROW + (2 * e + hi) * KT_ROW + + VD * l16 * 2); + v = vs; + #endif + + // A tile that reaches past S: its keys are masked below, and their V is + // zeroed so that a NaN in an unused slot cannot reach P@V as 0 * NaN. + if (__builtin_expect(kt + 16 > S, 0)) { + #pragma unroll + for (int e = 0; e < 8; ++e) + if (kt + 2 * e + hi >= S) v[e] = vrow_t{}; + } + #if WIN + // Likewise keys before every query's window: vLLM may have freed their + // pages, and the table then names a block whose contents are anyone's. + if (__builtin_expect(kt < ctx - (WIN - 1), 0)) { + #pragma unroll + for (int e = 0; e < 8; ++e) + if (kt + 2 * e + hi < ctx - (WIN - 1)) v[e] = vrow_t{}; + } + #endif + + // Lane holds row l16, keys kt + 2e + hi. Only a tile that reaches past + // the first query token's keys needs the causal mask. The scale is + // applied inside the exponent: it is positive, so the max commutes. + const bool tail = kt + 16 > ctx + 1; + #if WIN + // And only a tile reaching below the last query's window needs its + // lower edge. + const bool head = kt < ctx + MAXM - WIN; + #else + constexpr bool head = false; + #endif + #pragma unroll + for (int rt = 0; rt < RTW; ++rt) { + if (__builtin_expect(tail || head, 0)) { + #pragma unroll + for (int e = 0; e < 8; ++e) { + const int key = kt + 2 * e + hi; + #if MUTATE == 1 + const bool valid = key <= ctx + mrow[rt] + 1 && key < S; + #else + const bool valid = key <= ctx + mrow[rt]; + #endif + #if WIN + if (key < ctx + mrow[rt] - (WIN - 1)) s[rt][e] = -INFINITY; + #endif + if (!valid) s[rt][e] = -INFINITY; + } + } + float mx = s[rt][0]; + #pragma unroll + for (int e = 1; e < 8; ++e) mx = fmaxf(mx, s[rt][e]); + mx = fmaxf(mx, xhalf(mx)) * scale2; + const float mnew = fmaxf(m_run[rt], mx); + const float alpha = + (mnew == -INFINITY) ? 1.f : __builtin_amdgcn_exp2f(m_run[rt] - mnew); + m_run[rt] = mnew; + float p[8], lo[8], sum = 0.f; + #pragma unroll + for (int e = 0; e < 8; ++e) { + p[e] = (mnew == -INFINITY) ? 0.f + : __builtin_amdgcn_exp2f(__builtin_fmaf( + s[rt][e], scale2, -mnew)); + lo[e] = p[e] - elem_part(p[e]); + sum += p[e]; + } + l_run[rt] = l_run[rt] * alpha + sum + xhalf(sum); + + #if PPACK + // Columns 0..7 carry P's high half for rows 0..7 and columns 8..15 its + // low half for the same rows, so one WMMA does the work of two. A lane + // of columns 8..15 therefore accumulates for row l16-8, and scales by + // that row's alpha. A mask, not a ?: -- the compiler turned the + // conditional into a branch that computed pl only in lanes 8..15, and + // those read it from 0..7. + const u8v ph = __builtin_bit_cast(u8v, p_frag(p)); + const u8v pl = __builtin_bit_cast(u8v, p_frag(lo)); + const unsigned upper = 0u - (unsigned)(l16 >> 3); + u8v pk; + #pragma unroll + for (int i = 0; i < 8; ++i) + pk[i] = (ph[i] & ~upper) | (lower8(pl[i]) & upper); + const e16 bp = __builtin_bit_cast(e16, pk); + const float ascale = __builtin_bit_cast( + float, (__builtin_bit_cast(unsigned, alpha) & ~upper) | + (lower8(__builtin_bit_cast(unsigned, alpha)) & upper)); + #else + const e16 bp = p_frag(p); + const e16 bl = p_frag(lo); + const float ascale = alpha; + #endif + #pragma unroll + for (int j = 0; j < VD; ++j) { + acc[rt][j] *= ascale; + if ((ABLATE & 2) && j) continue; + const e16 a = v_frag(v, j); + acc[rt][j] = wmma(a, bp, acc[rt][j]); + #if !PPACK + acc[rt][j] = wmma(a, bl, acc[rt][j]); + #endif + } + } + asm volatile("" ::: "memory"); + }; + + // One tile per wave in flight. Keeping the next tile's loads in flight + // as well (two register sets, or the next K and V held until the current + // P@V is done) needs 60-100 VGPRs more than a wave has without spilling, + // and spilled it measured 59 % against 77 % of roof on 32/4/128. + #if SHT + // Stage, put the next tile's loads in flight, then compute this one: the + // registers a shared tile frees pay for the prefetch. The LDS wait goes + // before the loads, so the barrier does not wait for a page-table s_load + // counted with LDS. Two buffers: a wave refills one only after the + // barrier that every wave reaches after reading it. + int buf = 0; + for (int b = b0 + seg; b < nblk; b += nseg) { + char* const kb = KT_BUF(buf); + stage(ta, kb); + asm volatile("s_waitcnt lgkmcnt(0)" ::: "memory"); + if (b + nseg < nblk) issue(ta, b + nseg); + asm volatile("s_barrier" ::: "memory"); + process(ta, nullptr, kb, b); + buf ^= 1; + } + #elif M_PF + // Two tiles per wave in flight: the next one's loads go out before this + // one is computed. Only where a second tile's registers fit (small D). + Tile tb; + for (int b = b0 + seg; b < nblk; b += 2 * nseg) { + stage_k(ta); + if (b + nseg < nblk) issue(tb, b + nseg); + process(ta, ta.v, kt_s, b); + if (b + nseg >= nblk) break; + stage_k(tb); + if (b + 2 * nseg < nblk) issue(ta, b + 2 * nseg); + process(tb, tb.v, kt_s, b + nseg); + } + #elif VINLDS + // The tile's registers are free once staged: the next tile's loads go out + // before this one is computed. Its own buffer, so no barrier. + for (int b = b0 + seg; b < nblk; b += nseg) { + stage_kv(ta); + if (b + nseg < nblk) issue(ta, b + nseg); + process(ta, nullptr, kt_s, b); + } + #else + for (int b = b0 + seg; b < nblk; b += nseg) { + stage_k(ta); + process(ta, ta.v, kt_s, b); + if (b + nseg < nblk) issue(ta, b + nseg); + } + #endif + TS(3); + #if PPACK + // Fold the low half's columns onto the rows they belong to. + #pragma unroll + for (int j = 0; j < VD; ++j) + #pragma unroll + for (int e = 0; e < 8; ++e) acc[0][j][e] += upper8f(acc[0][j][e]); + #endif + + // Merge the waves. Every wave rescales to the workgroup's max; then, if + // the partials of all tiles do not fit the budget, tree rounds halve the + // live tiles (upper half stores, lower half adds in registers) until they + // do; then the live tiles store and the whole workgroup sums them in one + // pass. Only real rows are stored -- at M=1 and GQA=4 that is 4 of the 16 + // a row tile carries. Lane holds row l16 and, for element e of slice j, + // output d = dp*DPART + VD*(2e + hi) + j. + #pragma unroll + for (int rt = 0; rt < RTW; ++rt) + if (hi == 0) { + m_s[wave][RT(rt) * 16 + l16] = m_run[rt]; + l_s[wave][RT(rt) * 16 + l16] = l_run[rt]; + } + lds_barrier(); + float gsum[RTW]; + #pragma unroll + for (int rt = 0; rt < RTW; ++rt) { + const int r = RT(rt) * 16 + l16; + float mx = -INFINITY; + #pragma unroll + for (int w = rs * M_DSPL + dp; w < M_NW; w += M_DSPL * M_RSPL) + mx = fmaxf(mx, m_s[w][r]); + float sm = 0.f; + #pragma unroll + for (int w = rs * M_DSPL + dp; w < M_NW; w += M_DSPL * M_RSPL) + sm += weight_of(m_s[w][r], mx) * l_s[w][r]; + gsum[rt] = sm; + if (tw == 0 && dp == 0 && hi == 0) { + gm_s[r] = mx; + gl_s[r] = sm; + } + const float f = weight_of(m_run[rt], mx); + #pragma unroll + for (int j = 0; j < VD; ++j) acc[rt][j] *= f; + } + + TS(4); + const int grp = kvh * M_RG + rg; + const size_t pb = ((size_t)grp * M_NSEG + seg) * ROWPAD; + const bool publish = nseg > 1; + + #pragma unroll + for (int rt = 0; rt < RTW; ++rt) { + // Real rows in this row tile. + const int rv = min(16, ROWS_W - RT(rt) * 16); + // The previous row tile's readers must be done with the region. + if (rt) lds_barrier(); + #pragma unroll + for (int st = TILES / 2; st >= TFIN; st >>= 1) { + if (tw >= st && tw < 2 * st && l16 < rv) + store_rows(mrg_s + (((tw - st) * M_RSPL + rs) * M_DSPL + dp) * MSLOT, + acc[rt], l16, hi); + lds_barrier(); + if (tw < st && l16 < rv) + add_rows(acc[rt], mrg_s + ((tw * M_RSPL + rs) * M_DSPL + dp) * MSLOT, + l16, hi); + lds_barrier(); + } + #if TFIN == 1 && !CPUB + // One live tile: its waves write straight from registers. + if (tw == 0 && l16 < rv) { + const int r = RT(rt) * 16 + l16; + if (!publish) { + const int h = kvh * GQA + rg * (GQA / M_RG) + r / MAXM; + const int m = r % MAXM; + const float inv = __builtin_amdgcn_rcpf(gsum[rt]); + OutT* o = out + ((size_t)m * NUM_Q_HEADS + h) * HEAD_DIM + dp * DPART; + #pragma unroll + for (int e = 0; e < 8; ++e) { + elem_t x[VD]; + #pragma unroll + for (int j = 0; j < VD; ++j) x[j] = to_elem(acc[rt][j][e] * inv); + *(vrow_t*)(o + VD * (2 * e + hi)) = *(const vrow_t*)x; + } + } else { + store_rows(p_acc + (pb + RT(rt) * 16) * HEAD_DIM + dp * DPART, acc[rt], + l16, hi, HEAD_DIM); + } + } + #else + if (tw < TFIN && l16 < rv) + store_rows(mrg_s + ((tw * M_RSPL + rs) * M_DSPL + dp) * MSLOT, acc[rt], + l16, hi); + lds_barrier(); + // Eight consecutive d of one real row per thread, over the M_RSPL row + // tiles this round's waves stored. + const int r0 = rt * M_RSPL * 16; + const int nr = min(M_RSPL * 16, ROWS_W - r0); + for (int i = tid * 8; i < nr * HEAD_DIM; i += M_BLOCK * 8) { + const int rl = i / HEAD_DIM, d = i % HEAD_DIM; + const int r = r0 + rl; + const float* src = mrg_s + ((rl / 16) * M_DSPL + d / DPART) * MSLOT + + (rl % 16) * MROW + d % DPART; + f4 lo4 = *(const f4*)src, hi4 = *(const f4*)(src + 4); + #pragma unroll + for (int t = 1; t < TFIN; ++t) { + lo4 += *(const f4*)(src + t * M_RSPL * M_DSPL * MSLOT); + hi4 += *(const f4*)(src + t * M_RSPL * M_DSPL * MSLOT + 4); + } + if (!publish) { + const float inv = __builtin_amdgcn_rcpf(gl_s[r]); + e8 o; + #pragma unroll + for (int k = 0; k < 4; ++k) { + o[k] = to_elem(lo4[k] * inv); + o[4 + k] = to_elem(hi4[k] * inv); + } + const int h = kvh * GQA + rg * (GQA / M_RG) + r / MAXM; + const int m = r % MAXM; + *(e8*)(out + ((size_t)m * NUM_Q_HEADS + h) * HEAD_DIM + d) = o; + } else { + float* dst = p_acc + (pb + r) * HEAD_DIM + d; + *(f4*)dst = lo4; + *(f4*)(dst + 4) = hi4; + } + } + #endif + } + TS(5); + if (!publish) return; + + // Split KV: publish (acc, m, l), and the last segment to arrive merges. + // Nobody waits: the late arriver is already resident when it does the + // atomic, and it resets the counter, so a CUDA-graph replay starts clean. + // With a shared merge (`coop`: the host has checked the whole grid can be + // resident at once, so waiting cannot deadlock) the others wait for it to + // bump the group's generation, and every segment merges its slice. + // Counting arrivals first and letting only the non-last segments fence + // their partials out measured 0.5-1.6 % worse: their fence then sits on + // the last one's path instead of ahead of its atomic. + #if M_RSPL > 1 + // gm_s holds every wave's rows; with one tile no merge barrier ordered it. + lds_barrier(); + #endif + if (tid < ROWS_W) { + p_m[pb + tid] = gm_s[tid]; + p_l[pb + tid] = gl_s[tid]; + } + __threadfence(); + coop = SHARED_MERGE && coop; + int& lds_go = *reinterpret_cast(gl_s + ROWPAD); + lds_barrier(); + if (tid == 0) { + const bool last = atomicAdd(&p_cnt[grp], 1) == nseg - 1; + if (last) { + // The generation first: the counter only has to be clean by the next + // launch, and a release store behind it would wait for its ack. + if (coop) + __hip_atomic_store(gen, g0 + 1, __ATOMIC_RELEASE, + __HIP_MEMORY_SCOPE_AGENT); + p_cnt[grp] = 0; + } else if (coop) { + while (__hip_atomic_load(gen, __ATOMIC_RELAXED, + __HIP_MEMORY_SCOPE_AGENT) == g0) + __builtin_amdgcn_s_sleep(1); + } + lds_go = last || coop; + } + lds_barrier(); + TS(6); + if (!lds_go) return; + __threadfence(); + constexpr int NEL = ROWS_W * HEAD_DIM; + const int chunk = coop ? (NEL / 4 + nseg - 1) / nseg * 4 : NEL; + const int i0 = coop ? seg * chunk : 0; + const int i1 = min(NEL, i0 + chunk); + + // Every load of the merge goes out at once -- each thread's rows' m and l + // and its slice of every segment's partial, M_NSEG unrolled -- so the last + // arriver pays one L2 round trip, not one for the running max, one for the + // weights and one for the partials. + for (int i = i0 + tid * 4; i < i1; i += M_BLOCK * 4) { + const int r = i / HEAD_DIM, d = i % HEAD_DIM; + float pm[M_NSEG], pl[M_NSEG]; + f4 pa[M_NSEG]; + #pragma unroll + for (int sg = 0; sg < M_NSEG; ++sg) { + const size_t row = ((size_t)grp * M_NSEG + sg) * ROWPAD + r; + pm[sg] = sg < nseg ? p_m[row] : -INFINITY; + pl[sg] = sg < nseg ? p_l[row] : 0.f; + pa[sg] = sg < nseg ? *(const f4*)(p_acc + row * HEAD_DIM + d) : f4{}; + } + float gm = -INFINITY; + #pragma unroll + for (int sg = 0; sg < M_NSEG; ++sg) gm = fmaxf(gm, pm[sg]); + f4 num = {}; + float den = 0.f; + #pragma unroll + for (int sg = 0; sg < M_NSEG; ++sg) { + const float a = weight_of(pm[sg], gm); + den = fmaf(a, pl[sg], den); + num += a * pa[sg]; + } + const float inv = __builtin_amdgcn_rcpf(den); + const int h = kvh * GQA + rg * (GQA / M_RG) + r / MAXM; + const int m = r % MAXM; + OutT* o = out + ((size_t)m * NUM_Q_HEADS + h) * HEAD_DIM + d; + #pragma unroll + for (int k = 0; k < 4; ++k) o[k] = to_elem(num[k] * inv); + } + TS(7); +} + + #undef RT + #undef KT_BUF + #undef SX_SLOT + #undef ROWS_W + #undef RTILES + #undef ROWPAD + #undef TILES + #undef RTW + #undef SHT + #undef EPW + #undef KBLK + #undef DPART + #undef VD + #undef NCP + #undef PPACK + #undef KT_ROW + #undef KT_BYTES + #undef QS_BYTES + #undef MROW + #undef RV + #undef MSLOT + #undef SLOTS_FIT + #undef TFIN + #undef MRG_SLOTS + #undef MRG_BYTES + #undef TB_BYTES + #undef SX_BYTES + #undef LOOP_BYTES + #undef LOOP_BYTES + #undef LDS_RAW + #undef SHARED_MERGE + #undef LDS_MS + #undef M_BLOCK + #undef LDS_ALL + + #endif // M_DOT +#endif // RDNA35_BODY diff --git a/csrc/rocm/torch_bindings.cpp b/csrc/rocm/torch_bindings.cpp index aa9a7951cd5e..19858c8454c1 100644 --- a/csrc/rocm/torch_bindings.cpp +++ b/csrc/rocm/torch_bindings.cpp @@ -179,6 +179,17 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, rocm_ops) { rocm_ops.impl("gdn_chunked", torch::kCUDA, &gdn_chunked); #endif +#ifdef VLLM_ROCM_RDNA35_ATTN + // RDNA3.5 decode attention (gfx1151): the variants built into the wheel. + rocm_ops.def("rdna35_decode_variant(int[] key) -> int", + &rdna35_decode_variant); + rocm_ops.def( + "rdna35_decode_attn(int variant, Tensor q, Tensor kv_cache, " + "Tensor block_table, Tensor! out, Tensor! acc, Tensor! m, Tensor! l, " + "Tensor! cnt, Tensor seq_lens, float scale) -> ()"); + rocm_ops.impl("rdna35_decode_attn", torch::kCUDA, &rdna35_decode_attn); +#endif + // W4A16 MoE prefill WMMA GEMM for AMD RDNA3 (gfx11). Always registered; the // kernel body is gfx11-only (stub elsewhere) and Python gates calls on // on_gfx1x(). Mutates c in place; an unsupported shape raises via TORCH_CHECK diff --git a/tests/kernels/attention/test_rdna35_hip_decode.py b/tests/kernels/attention/test_rdna35_hip_decode.py new file mode 100644 index 000000000000..c8fdf2e85626 --- /dev/null +++ b/tests/kernels/attention/test_rdna35_hip_decode.py @@ -0,0 +1,634 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""RDNA3.5 HIP decode attention: agreement with a reference, and honest fallback. + +The kernel walks the paged KV cache with its own address arithmetic rather than +reading the tensor's strides, so a layout it did not expect yields finite, +wrong numbers instead of an error. These tests therefore cover three things a +plain tolerance check would miss: + +- a short context, where a causal off-by-one is visible. The error from one + masked-off-by-one key falls off as ~1/S while the tolerance is fixed, so long + contexts hide the bug most likely to be present. +- a deliberately mutated kernel, which the comparison *must* fail. A test that + passes a known-broken kernel proves nothing about the working one. It also + exercises two variants in one process, which only works because each build + gets its own device-function symbols. +- that an unsupported shape falls back to Triton rather than being served + wrong. +""" + +import pytest +import torch + +from vllm.platforms import current_platform + +pytestmark = pytest.mark.skipif(not current_platform.is_rocm(), reason="ROCm only") + +rdna35 = pytest.importorskip("vllm.v1.attention.ops.rdna35_hip_decode") + +HQ, HKV, HEAD_DIM, M, BLOCK_SIZE = 32, 16, 256, 4, 16 +# Partials big enough (256 KiB per group) for the segments to share the merge, +# on a grid small enough (32 workgroups) to be resident at once, which the +# host requires before it lets them wait for each other. +SHARED = dict(hq=16, hkv=2, hd=512, rg=2, nseg=8) +# Four row tiles per kv head, split over the waves that share each key tile +# (RSPL). With nw=4 there is one tile per workgroup, so no merge round orders +# the waves' rows before they are published; nseg=16 makes the merge shared. +RSPL = { + "one tile, shared merge": dict(hq=32, hkv=2, hd=128, rspl=4, nw=4, nseg=16), + "two tiles": dict(hq=32, hkv=2, hd=128, rspl=4, nw=8, nseg=4), + "with dspl": dict(hq=32, hkv=2, hd=128, rspl=2, nw=8, dspl=2, rg=2, nseg=4), + # Not a row split: unshared tiles with the next tile's loads in flight. + "prefetch": dict(hq=16, hkv=2, hd=64, rg=2, nw=4, pf=1, nseg=8), + # One live tile after the merge's tree, its partials staged in LDS. + "staged publish": dict(hq=8, hkv=1, hd=256, rspl=2, nw=4, pf=1, cpub=1, nseg=16), + # Unshared tiles with V staged in LDS, the next tile's loads before this + # one's compute (VINLDS). + "v in lds": dict(hq=32, hkv=2, hd=128, rg=4, nw=4, vinlds=1, nseg=4), + # ... and with the head dim split over waves, whose score exchange writes + # into the K half of each wave's buffer. + "v in lds, dspl": dict(hq=16, hkv=8, hd=256, nw=4, dspl=2, vinlds=1, nseg=2), + # Not WMMA at all: the per-q-head dot-product decomposition. + "dot": dict(hq=16, hkv=2, hd=256, nw=8, dot=1, bfly=4, nseg=2), +} +# Sliding windows: the WMMA path with and without segments, and the dot path. +# 1000 keys starts the window off any tile and page boundary. +WINDOWED = { + "wmma": dict(hq=16, hkv=8, hd=256, nw=4, window=1000, nseg=1), + "wmma segments": dict(hq=8, hkv=1, hd=256, nw=4, rg=2, window=512, nseg=8), + "dot": dict(hq=8, hkv=2, hd=256, nw=8, dot=1, bfly=4, window=1000, nseg=2), +} +# Relative error bound per dtype. The reference sees the same rounded inputs, +# so what differs is the kernel's arithmetic and its output rounding -- the +# latter alone up to 2^-8 relative in bf16. +TOL = {torch.float16: 1e-3, torch.bfloat16: 8e-3} +DTYPES = list(TOL) + + +def _skip_unless_gfx1151(): + from vllm.platforms.rocm import on_gfx1151 + + if not on_gfx1151(): + pytest.skip("kernel is built for gfx1151") + + +@pytest.fixture(scope="session") +def _jit(): + """The kernel tests build variants of their own -- layouts, segment + counts, a mutated kernel -- which only the JIT loader does. The tests of + the variants built into _rocm_C turn it back off.""" + with pytest.MonkeyPatch.context() as mp: + mp.setenv("VLLM_RDNA35_ATTN_JIT", "1") + yield + + +@pytest.fixture(scope="session", autouse=True) +def _build_variants(_jit): + """Build the variants this module needs before any test runs. + + They are independent builds of ~19.6 s each, almost all of it torch's + headers rather than the kernel, so serially they were the whole runtime of + the suite on a cold cache. Built together it is one build's worth. + + Not a correctness concern: a build that fails here fails again in the test + that needs it, where it is reported against that test rather than as a + collection error. + """ + from vllm.platforms.rocm import on_gfx1151 + + if not on_gfx1151(): + return + rdna35.precompile( + v + for dtype in DTYPES + for v in ( + _variant(layout=0, dtype=dtype), + _variant(layout=1, dtype=dtype), + _variant(layout=0, nseg=4, dtype=dtype), + _variant(layout=1, nseg=4, dtype=dtype), + _variant(layout=1, mutate=1, dtype=dtype), + _variant(**SHARED, dtype=dtype), + *(_variant(**shape, dtype=dtype) for shape in RSPL.values()), + *( + _variant(**{k: v for k, v in shape.items()}, dtype=dtype) + for shape in WINDOWED.values() + ), + *( + _variant( + **{k: v for k, v in shape.items() if k != "nseg"}, + nseg=shape["nseg"], + dtype=dtype, + batch=1, + ) + for shape in BATCHED.values() + ), + ) + ) + + +def _variant( + layout: int = 1, + mutate: int = 0, + nseg: int = 1, + hq: int = HQ, + hkv: int = HKV, + hd: int = HEAD_DIM, + rg: int = 1, + dtype: torch.dtype = torch.float16, + rspl: int = 1, + nw: int = 8, + dspl: int = 0, + pf: int = 0, + cpub: int = 0, + vinlds: int = 0, + dot: int = 0, + bfly: int = 0, + window: int = 0, + batch: int = 0, +): + # nseg > 1 with minb = 1 splits even a short context over several + # workgroups, which is the only way to reach the cross-workgroup merge. + return rdna35.KernelVariant( + head_size=hd, + num_q_heads=hq, + num_kv_heads=hkv, + max_m=M, + block_size=BLOCK_SIZE, + layout=layout, + nseg=nseg, + rg=rg, + minb=1, + mutate=mutate, + dtype=dtype, + rspl=rspl, + nw=nw, + dspl=dspl, + pf=pf, + cpub=cpub, + vinlds=vinlds, + dot=dot, + bfly=bfly, + window=window, + batch=batch, + ) + + +def _paged_inputs( + seq_len: int, + layout: int, + dtype: torch.dtype, + seed: int = 0, + hq: int = HQ, + hkv: int = HKV, + hd: int = HEAD_DIM, + m: int = M, + bs: int = BLOCK_SIZE, +): + """Build a KV cache whose physical order matches the layout, then present + it in the logical (num_blocks, num_kv_heads, block_size, 2*hs) order the + backend passes down. The last page may be partly filled.""" + torch.manual_seed(seed) + dev = torch.device("cuda") + num_blocks = -(-seq_len // bs) + if layout == 0: # NHD + kv = torch.randn(num_blocks, bs, hkv, 2 * hd, device=dev, dtype=dtype) + kv = kv.transpose(1, 2) + else: # HND + kv = torch.randn(num_blocks, hkv, bs, 2 * hd, device=dev, dtype=dtype) + kv = kv * 0.5 + q = torch.randn(m, hq, hd, device=dev, dtype=dtype) * 0.5 + return q, kv, torch.arange(num_blocks, device=dev, dtype=torch.int32) + + +def _reference(q, kv, seq_len, window=0): + m, hq, hkv, hd = q.shape[0], q.shape[1], kv.shape[1], q.shape[2] + flat = kv.transpose(1, 2).reshape(-1, hkv, 2 * hd)[:seq_len] + k, v = flat[..., :hd], flat[..., hd:] + gqa = hq // hkv + qf = q.float().permute(1, 0, 2) + kf = k.float().permute(1, 0, 2).repeat_interleave(gqa, 0) + vf = v.float().permute(1, 0, 2).repeat_interleave(gqa, 0) + scores = torch.bmm(qf, kf.transpose(1, 2)) * (hd**-0.5) + pos = torch.arange(seq_len, device=q.device).view(1, seq_len) + lim = (seq_len - m + torch.arange(m, device=q.device)).view(m, 1) + masked = pos > lim + if window: + masked |= pos < lim - (window - 1) + scores = scores.masked_fill(masked.view(1, m, seq_len), float("-inf")) + return torch.bmm(torch.softmax(scores, -1), vf).permute(1, 0, 2) + + +def _run( + seq_len, + layout=1, + dtype=torch.float16, + mutate=0, + nseg=1, + repeat=1, + build_dtype=None, + free_before_window=False, + **shape, +): + _skip_unless_gfx1151() + dims = {k: v for k, v in shape.items() if k in ("hq", "hkv", "hd")} + q, kv, block_table = _paged_inputs(seq_len, layout, dtype, **dims) + variant = _variant(layout, mutate, nseg, dtype=build_dtype or dtype, **shape) + window = shape.get("window", 0) + ref = _reference(q, kv, seq_len, window) + if free_before_window: + # vLLM frees the pages no query's window reaches and points their + # table entries elsewhere: here, at a page of NaN the kernel must never + # let into the output. + assert layout == 1 + first = max(0, seq_len - M - (window - 1)) // BLOCK_SIZE + kv = torch.cat([kv, torch.full_like(kv[:1], float("nan"))]) + block_table = block_table.clone() + block_table[:first] = kv.shape[0] - 1 + module = rdna35.load(variant) + acc, m, ln, arrivals = rdna35.make_scratch(variant, q.device) + out = torch.empty_like(q) + seq_lens = torch.tensor([seq_len], device=q.device, dtype=torch.int32) + for _ in range(repeat): + out.zero_() + module.decode_attn( + q, kv, block_table, out, acc, m, ln, arrivals, seq_lens, q.shape[2] ** -0.5 + ) + torch.accelerator.synchronize() + return out.float(), ref + + +def _max_rel(got, ref) -> float: + """Relative error with a floor, which is what catches a causal off-by-one. + + max_abs alone does not: one key masked wrongly gives max_abs=1.2e-03 at S=2048 + and slips past a 2e-2 absolute threshold. + """ + return ((got - ref).abs() / ref.abs().clamp_min(1e-3)).max().item() + + +@pytest.mark.parametrize("dtype", DTYPES) +@pytest.mark.parametrize("seq_len", [48, 1024]) +@pytest.mark.parametrize("layout", [0, 1]) +@pytest.mark.parametrize("nseg", [1, 4]) +def test_matches_reference(seq_len, layout, nseg, dtype): + got, ref = _run(seq_len, layout=layout, nseg=nseg, dtype=dtype) + assert torch.isfinite(got).all() + assert _max_rel(got, ref) <= TOL[dtype] + + +@pytest.mark.parametrize("dtype", DTYPES) +def test_split_merge_survives_relaunch(dtype): + """The arrival counters must be back at zero after every launch. + + The last workgroup to arrive resets them; if it did not, the second launch + would elect no merger and leave the output unwritten. + """ + got, ref = _run(1024, nseg=4, repeat=3, dtype=dtype) + assert _max_rel(got, ref) <= TOL[dtype] + + +@pytest.mark.parametrize("dtype", DTYPES) +@pytest.mark.parametrize("seq_len", [48, 1024]) +def test_shared_split_merge(seq_len, dtype): + """Segments that wait for each other and merge a slice each. + + Relaunched, because the wait is on a generation that must keep advancing: + one stuck at the value a later launch reads first would release nobody. + """ + shape = {k: v for k, v in SHARED.items() if k != "nseg"} + got, ref = _run(seq_len, nseg=SHARED["nseg"], repeat=3, dtype=dtype, **shape) + assert torch.isfinite(got).all() + assert _max_rel(got, ref) <= TOL[dtype] + + +@pytest.mark.parametrize("dtype", DTYPES) +@pytest.mark.parametrize("seq_len", [48, 1024]) +@pytest.mark.parametrize("name", list(RSPL)) +def test_row_tiles_split_over_waves(name, seq_len, dtype): + """Waves that share a key tile through LDS and split its rows.""" + shape = dict(RSPL[name]) + nseg = shape.pop("nseg") + got, ref = _run(seq_len, nseg=nseg, repeat=2, dtype=dtype, **shape) + assert torch.isfinite(got).all() + assert _max_rel(got, ref) <= TOL[dtype] + + +@pytest.mark.parametrize("dtype", DTYPES) +@pytest.mark.parametrize("seq_len", [48, 512, 1024, 2000]) +@pytest.mark.parametrize("name", list(WINDOWED)) +def test_sliding_window(name, seq_len, dtype): + """Keys before the window are masked; a window wider than S (48, 512 at + 1000) reads everything and nothing before it.""" + shape = dict(WINDOWED[name]) + nseg = shape.pop("nseg") + got, ref = _run(seq_len, nseg=nseg, repeat=2, dtype=dtype, **shape) + assert torch.isfinite(got).all() + assert _max_rel(got, ref) <= TOL[dtype] + + +@pytest.mark.parametrize("name", list(WINDOWED)) +def test_sliding_window_never_reads_freed_pages(name): + """Pages wholly before the window may be freed; their contents must not + reach the output even as 0 * NaN. + + S=2016 puts the window's first key late in its WMMA block, so that block's + first tile is a whole freed page: with the V zeroing removed the WMMA + cases fail. (The dot path's tiles never leave the first key's page.) + """ + shape = dict(WINDOWED[name]) + nseg = shape.pop("nseg") + got, ref = _run(2016, nseg=nseg, free_before_window=True, **shape) + assert torch.isfinite(got).all() + assert _max_rel(got, ref) <= TOL[torch.float16] + + +def test_graph_replay_follows_seq_lens(): + """S is read on the device, so a captured graph serves a longer sequence. + + A host int is frozen into the graph at capture: every replay would attend + over the capture-time length, which is what full CUDA-graph decode does. + """ + _skip_unless_gfx1151() + long_s = 1024 + q, kv, block_table = _paged_inputs(long_s, 1, torch.float16) + variant = _variant(layout=1, nseg=4) + module = rdna35.load(variant) + scratch = rdna35.make_scratch(variant, q.device) + out = torch.empty_like(q) + seq_lens = torch.tensor([48], device=q.device, dtype=torch.int32) + + def launch(): + module.decode_attn( + q, kv, block_table, out, *scratch, seq_lens, q.shape[2] ** -0.5 + ) + + launch() # warm-up outside the capture + graph = torch.cuda.CUDAGraph() + with torch.cuda.graph(graph): + launch() + for s in (48, 208, long_s): + seq_lens.fill_(s) + out.zero_() + graph.replay() + torch.accelerator.synchronize() + n = s // BLOCK_SIZE + ref = _reference(q, kv[:n], s) + assert _max_rel(out.float(), ref) <= TOL[torch.float16], f"S={s}" + + +# One launch over several sequences (grid.y), each with its own length, table +# row, query rows, output rows and scratch. S=0 is a CUDA-graph batch padded +# past its real sequences. +BATCHED = { + "split, last arriver merges": dict(hq=HQ, hkv=HKV, hd=HEAD_DIM, nseg=4), + "shared merge": dict(SHARED), + "dot": dict(RSPL["dot"]), + "window": dict(WINDOWED["wmma segments"]), +} + + +@pytest.mark.parametrize("dtype", DTYPES) +@pytest.mark.parametrize("name", list(BATCHED)) +def test_batch_of_sequences(name, dtype): + """Every sequence of a batch matches its own reference, and a padded + sequence (S=0) reads nothing and disturbs none of them.""" + _skip_unless_gfx1151() + shape = dict(BATCHED[name]) + nseg = shape.pop("nseg") + _check_batch( + _variant(1, 0, nseg, dtype=dtype, batch=1, **shape), [1008, 48, 0, 2016] + ) + + +def _check_batch(variant, lens): + """Launch a batch build over sequences of these lengths and check each + against its own reference, twice, a padded one (S=0) left unwritten.""" + v = variant + dims = dict( + hq=v.num_q_heads, hkv=v.num_kv_heads, hd=v.head_size, m=v.max_m, bs=v.block_size + ) + width = max(lens) // v.block_size + 1 + parts = [ + _paged_inputs(max(s, 16), 1, v.dtype, seed=i, **dims) + for i, s in enumerate(lens) + ] + kv = torch.cat([p[1] for p in parts]) + first = torch.tensor([0] + [p[1].shape[0] for p in parts]).cumsum(0) + bt = torch.zeros(len(lens), width, device=kv.device, dtype=torch.int32) + for i, p in enumerate(parts): + bt[i, : p[1].shape[0]] = p[2] + int(first[i]) + q = torch.cat([p[0] for p in parts]) + module = rdna35.load(v) + scratch = rdna35.make_scratch(v, q.device, max_seqs=len(lens)) + seq_lens = torch.tensor(lens, device=q.device, dtype=torch.int32) + out = torch.full_like(q, 7.0) + for _ in range(2): + module.decode_attn(q, kv, bt, out, *scratch, seq_lens, q.shape[2] ** -0.5) + torch.accelerator.synchronize() + for i, s in enumerate(lens): + rows = out[i * v.max_m : (i + 1) * v.max_m].float() + if s == 0: + assert (rows == 7.0).all(), f"{v.name}: a padded sequence was written" + continue + ref = _reference(parts[i][0], parts[i][1], s, v.window) + assert _max_rel(rows, ref) <= TOL[v.dtype], f"{v.name}: sequence {i}, S={s}" + + +@pytest.mark.parametrize("dtype", DTYPES) +def test_other_dtype_is_refused_not_miscomputed(dtype): + """A build must refuse the other 16-bit type, not return nonsense. + + The loads reinterpret the tensors as the element type the variant was + built for, so fp16 read as bf16 or the reverse stays finite and is wrong + by orders of magnitude. + """ + other = next(d for d in DTYPES if d != dtype) + with pytest.raises(RuntimeError, match="kernel built for"): + _run(1024, dtype=dtype, build_dtype=other) + + +@pytest.mark.parametrize("dtype", DTYPES) +def test_negative_control_is_detected(dtype): + """A kernel mutated to admit one key too many must fail the comparison.""" + got, ref = _run(48, mutate=1, dtype=dtype) + assert _max_rel(got, ref) > TOL[dtype], ( + "the mutated kernel passed, so this comparison cannot detect a causal " + "off-by-one and proves nothing about the real kernel" + ) + + +def test_layouts_disagree_when_data_differs(): + """NHD and HND must be distinct code paths, not the same one twice.""" + assert rdna35.expected_kv_cache_strides( + _variant(layout=0) + ) != rdna35.expected_kv_cache_strides(_variant(layout=1)) + + +def test_unsupported_head_size_falls_back_to_triton(): + """An unbuilt shape must be served by Triton, not served wrong.""" + from vllm.v1.attention.backends.rocm_attn import RocmAttentionRdna35Impl + from vllm.v1.kv_cache_interface import KVQuantMode + + impl = RocmAttentionRdna35Impl.__new__(RocmAttentionRdna35Impl) + impl._rejected = None + # 96 is a real shipped head size (Phi-3.5-vision) and is deliberately not + # built: 96/32 is 3 fp16 per lane, which is not a power of two and so has + # no single load width. It must fall back, not be served wrong. + impl.head_size, impl.num_heads, impl.num_kv_heads = 96, HQ, HKV + + fits = impl._prepare( + kv_cache=torch.empty(0, dtype=torch.float16), + q=torch.empty(0, dtype=torch.float16), + alibi_slopes=None, + sinks=None, + softcap=0, + causal=True, + window_size=None, + kv_quant_mode=KVQuantMode.NONE, + seqused_k=torch.zeros(1), + max_seqlen_q=0, + ) + assert not fits + assert "head_size 96" in impl._rejected + + +@pytest.mark.parametrize( + "nseq, tokens, max_q, reason", + [ + (3, 5, 4, "unequal query length"), # a prefill mixed with decodes + (1, 64, 64, "more than"), # a prompt: one build per length otherwise + ], +) +def test_non_decode_batches_fall_back_to_triton(nseq, tokens, max_q, reason): + """Only batches of equal, decode-sized query lengths reach the kernel.""" + from vllm.v1.attention.backends.rocm_attn import RocmAttentionRdna35Impl + from vllm.v1.kv_cache_interface import KVQuantMode + + impl = RocmAttentionRdna35Impl.__new__(RocmAttentionRdna35Impl) + impl._rejected = None + impl.head_size, impl.num_heads, impl.num_kv_heads = HEAD_DIM, HQ, HKV + fits = impl._prepare( + kv_cache=torch.empty(0, dtype=torch.float16), + q=torch.empty(tokens, HQ, HEAD_DIM, dtype=torch.float16), + alibi_slopes=None, + sinks=None, + softcap=0, + causal=True, + window_size=None, + kv_quant_mode=KVQuantMode.NONE, + seqused_k=torch.zeros(nseq), + max_seqlen_q=max_q, + ) + assert not fits + assert reason in impl._rejected + + +def _built_variants(): + return rdna35.built_variants() + + +def test_variants_def_matches_the_tables(tmp_path): + """variants.def is what _rocm_C is built from. A table row changed without + regenerating it would leave the backend asking for a variant that was + never built, and every call of that shape silently on Triton.""" + path = tmp_path / "variants.def" + rdna35.write_variants_def(_built_variants(), path) + assert path.read_text() == rdna35.VARIANTS_DEF.read_text(), ( + "variants.def is stale: python -m vllm.v1.attention.ops.rdna35_hip_decode" + ) + + +def test_every_listed_variant_is_built_in(monkeypatch): + _skip_unless_gfx1151() + monkeypatch.delenv("VLLM_RDNA35_ATTN_JIT") + missing = [v.name for v in _built_variants() if rdna35._builtin_index(v) < 0] + assert not missing, f"{len(missing)} variants not in _rocm_C, e.g. {missing[0]}" + + +@pytest.mark.parametrize( + "group", sorted({rdna35.variant_group(v) for v in _built_variants()}) +) +def test_built_in_variants_match_reference(group, monkeypatch): + """Every variant built into _rocm_C, launched through the registry: one + sequence past a partial tile, and a batch with a padded sequence. Windows + run past the window, so that it masks; large pages end partly filled.""" + _skip_unless_gfx1151() + monkeypatch.delenv("VLLM_RDNA35_ATTN_JIT") + for v in _built_variants(): + if rdna35.variant_group(v) != group: + continue + s = max(1024, v.window + 64, v.block_size + 1000) + if v.batch: + _check_batch(v, [s, 48, 0]) + continue + dims = dict( + hq=v.num_q_heads, + hkv=v.num_kv_heads, + hd=v.head_size, + m=v.max_m, + bs=v.block_size, + ) + q, kv, block_table = _paged_inputs(s, 1, v.dtype, **dims) + module = rdna35.load(v) + assert isinstance(module, rdna35._Builtin), v.name + out = torch.empty_like(q) + module.decode_attn( + q, + kv, + block_table, + out, + *rdna35.make_scratch(v, q.device), + torch.tensor([s], device=q.device, dtype=torch.int32), + q.shape[2] ** -0.5, + ) + torch.accelerator.synchronize() + ref = _reference(q, kv, s, v.window) + assert _max_rel(out.float(), ref) <= TOL[v.dtype], v.name + + +@pytest.mark.parametrize("gfx1151", [True, False]) +def test_rocm_attn_ahead_of_triton_only_on_gfx1151(gfx1151, monkeypatch): + """On gfx1151 ROCM_ATTN is this kernel, ahead of Triton, which serves what + it does not; the other RDNA targets keep Triton first.""" + import vllm.platforms.rocm as rocm + from vllm.v1.attention.backends.registry import AttentionBackendEnum as B + + monkeypatch.setattr(rocm, "on_gfx1x", lambda: True) + monkeypatch.setattr(rocm, "on_gfx1151", lambda: gfx1151) + order = rocm._get_backend_priorities(use_mla=False, use_sparse=False) + assert (order.index(B.ROCM_ATTN) < order.index(B.TRITON_ATTN)) == gfx1151 + + +@pytest.mark.parametrize("gfx1151", [True, False]) +def test_rocm_attn_layout_and_impl_only_on_gfx1151(gfx1151, monkeypatch): + """On gfx1151 ROCM_ATTN runs this kernel on Triton's packed HND cache; + elsewhere it keeps its own cache and impl. ROCM_AITER_UNIFIED_ATTN, its + subclass, keeps its own on both.""" + import vllm.platforms.rocm as rocm + from vllm.v1.attention.backends import rocm_attn + from vllm.v1.attention.backends.rocm_aiter_unified_attn import ( + RocmAiterUnifiedAttentionBackend as Aiter, + ) + from vllm.v1.attention.backends.triton_attn import TritonAttentionBackend + + monkeypatch.setattr(rocm, "on_gfx1151", lambda: gfx1151) + backend = rocm_attn.RocmAttentionBackend + shape = backend.get_kv_cache_shape(8, 16, HKV, HEAD_DIM) + if gfx1151: + assert backend.get_impl_cls() is rocm_attn.RocmAttentionRdna35Impl + assert backend.get_builder_cls() is rocm_attn.RocmAttentionRdna35MetadataBuilder + assert shape == TritonAttentionBackend.get_kv_cache_shape(8, 16, HKV, HEAD_DIM) + assert backend.get_required_kv_cache_layout() == "HND" + assert 512 in backend.get_supported_head_sizes() + else: + assert backend.get_impl_cls() is rocm_attn.RocmAttentionImpl + assert shape == (2, 8, 16, HKV, HEAD_DIM) + assert backend.get_required_kv_cache_layout() is None + assert 512 not in backend.get_supported_head_sizes() + with pytest.raises(NotImplementedError): + backend.get_kv_cache_stride_order() + assert Aiter.get_required_kv_cache_layout() is None + assert 512 not in Aiter.get_supported_head_sizes() + with pytest.raises(NotImplementedError): + Aiter.get_kv_cache_stride_order() diff --git a/vllm/envs.py b/vllm/envs.py index b2d2977c71a7..0cfde4cd2e0f 100755 --- a/vllm/envs.py +++ b/vllm/envs.py @@ -154,6 +154,7 @@ VLLM_ROCM_USE_AITER_FUSION_SHARED_EXPERTS: bool = False VLLM_ROCM_USE_AITER_TRITON_GEMM: bool = True VLLM_ROCM_USE_SKINNY_GEMM: bool = True + VLLM_RDNA35_ATTN_JIT: bool = False VLLM_ROCM_FP8_PADDING: bool = True VLLM_ROCM_MOE_PADDING: bool = True VLLM_ROCM_SHUFFLE_KV_CACHE_LAYOUT: bool = False @@ -1369,6 +1370,11 @@ def _resolve_rust_cli_path() -> str | None: "VLLM_ROCM_USE_SKINNY_GEMM": lambda: ( os.getenv("VLLM_ROCM_USE_SKINNY_GEMM", "True").lower() in ("true", "1") ), + # ROCM_ATTN on gfx1151: JIT-build every decode kernel from the source checkout + # instead of using the variants compiled into _rocm_C (kernel development) + "VLLM_RDNA35_ATTN_JIT": lambda: ( + os.getenv("VLLM_RDNA35_ATTN_JIT", "False").lower() in ("true", "1") + ), # Pad the fp8 weights to 256 bytes for ROCm "VLLM_ROCM_FP8_PADDING": lambda: bool(int(os.getenv("VLLM_ROCM_FP8_PADDING", "1"))), # Pad the weights for the moe kernel diff --git a/vllm/model_executor/models/config.py b/vllm/model_executor/models/config.py index 12c9f5630ec5..840cc2a70d46 100644 --- a/vllm/model_executor/models/config.py +++ b/vllm/model_executor/models/config.py @@ -217,7 +217,7 @@ def verify_and_update_config(vllm_config: "VllmConfig") -> None: When FA4 is available we force it for ALL layers, giving a uniform kernel path and avoiding the mixed FA3+FA4 penalty. When FA4 is not available we fall back - to Triton. + to Triton. On gfx1151 every layer runs ROCM_ATTN, ahead of both. """ model_config = vllm_config.model_config arch_config = model_config.model_arch_config @@ -230,9 +230,23 @@ def verify_and_update_config(vllm_config: "VllmConfig") -> None: if len(set(head_dims.values())) <= 1: return + from vllm.platforms import current_platform from vllm.v1.attention.backends.fa_utils import is_fa_version_supported from vllm.v1.attention.backends.registry import AttentionBackendEnum + if current_platform.is_rocm() and vllm_config.attention_config.backend is None: + from vllm.platforms.rocm import on_gfx1151 + + # gfx1151: ROCM_ATTN, the RDNA3.5 kernel, built for both head sizes. + if on_gfx1151(): + vllm_config.attention_config.backend = AttentionBackendEnum.ROCM_ATTN + logger.info( + "Gemma4 model has heterogeneous head dimensions %s. Using " + "ROCM_ATTN for all layers.", + head_dims, + ) + return + max_head_dim = max(head_dims.values()) if is_fa_version_supported(4) and max_head_dim <= 512: diff --git a/vllm/platforms/rocm.py b/vllm/platforms/rocm.py index 5eb64c87f0d2..3528b7d1c239 100644 --- a/vllm/platforms/rocm.py +++ b/vllm/platforms/rocm.py @@ -515,7 +515,13 @@ def _get_backend_priorities( # Keep ROCM_ATTN disabled for KV connectors until connector transfer # semantics are validated for its asymmetric native K/V cache views. allow_rocm_attn = not use_kv_connector - if on_gfx1x(): + if on_gfx1151(): + # ROCM_ATTN on gfx1151 is the RDNA3.5 HIP decode kernel, with + # TRITON_ATTN's kernels for every call it does not cover. + if allow_rocm_attn: + backends.append(AttentionBackendEnum.ROCM_ATTN) + backends.append(AttentionBackendEnum.TRITON_ATTN) + elif on_gfx1x(): # On RDNA (gfx11/gfx12), TRITON_ATTN is faster than ROCM_ATTN # because ROCM_ATTN's custom paged attention kernel falls back # to Triton internally with extra overhead. diff --git a/vllm/v1/attention/backends/rocm_attn.py b/vllm/v1/attention/backends/rocm_attn.py index 07cd6a50d16c..c32ae32a7a52 100644 --- a/vllm/v1/attention/backends/rocm_attn.py +++ b/vllm/v1/attention/backends/rocm_attn.py @@ -1,9 +1,14 @@ # SPDX-License-Identifier: Apache-2.0 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project -"""Attention layer with PagedAttention and Triton prefix prefill.""" +"""Attention layer with PagedAttention and Triton prefix prefill. -from dataclasses import dataclass -from typing import ClassVar +On gfx1151 (Strix Halo) the backend serves decode and speculative decode with +the RDNA3.5 HIP kernel (csrc/rocm/rdna35_decode_attn.cu) on Triton's packed +HND KV cache, and everything else with TRITON_ATTN's kernels on that cache. +""" + +from dataclasses import dataclass, replace +from typing import Any, ClassVar import torch @@ -27,20 +32,50 @@ CommonAttentionMetadata, MultipleOf, ) +from vllm.v1.attention.backends.triton_attn import ( + TritonAttentionBackend, + TritonAttentionImpl, + TritonAttentionMetadata, + TritonAttentionMetadataBuilder, +) +from vllm.v1.attention.backends.utils import ( + KVCacheLayoutType, + split_decodes_and_prefills, +) from vllm.v1.attention.ops.chunked_prefill_paged_decode import ( chunked_prefill_paged_decode, has_native_kv_cache_layout, ) from vllm.v1.attention.ops.paged_attn import PagedAttention +from vllm.v1.attention.ops.rdna35_hip_decode import ( + MAX_M, + SUPPORTED_HEAD_SIZES, + KernelVariant, + VariantBuildError, + expected_kv_cache_strides, + load, + make_scratch, + scratch_bytes, + variant_for, +) from vllm.v1.attention.ops.triton_reshape_and_cache_flash import ( triton_reshape_and_cache_flash, ) -from vllm.v1.kv_cache_interface import AttentionSpec +from vllm.v1.kv_cache_interface import AttentionSpec, KVQuantMode from vllm.v1.utils import create_attention_profiler_scope logger = init_logger(__name__) +def _on_rdna35() -> bool: + """gfx1151: the RDNA3.5 HIP decode kernel on the packed HND cache.""" + if not current_platform.is_rocm(): + return False + from vllm.platforms.rocm import on_gfx1151 + + return on_gfx1151() + + @dataclass class RocmAttentionMetadata: # NOTE(sang): Definition of context_len, query_len, and seq_len. @@ -190,7 +225,10 @@ def get_supported_kernel_block_sizes() -> list[int | MultipleOf]: @classmethod def get_supported_head_sizes(cls) -> list[int]: - return [32, 64, 80, 96, 128, 160, 192, 224, 256] + sizes = [32, 64, 80, 96, 128, 160, 192, 224, 256] + if cls is RocmAttentionBackend and _on_rdna35(): + sizes.append(512) + return sizes @classmethod def supports_mm_prefix(cls) -> bool: @@ -224,7 +262,9 @@ def supports_sliding_window(cls) -> bool: return True @staticmethod - def get_impl_cls() -> type["RocmAttentionImpl"]: + def get_impl_cls() -> type[AttentionImpl]: + if _on_rdna35(): + return RocmAttentionRdna35Impl return RocmAttentionImpl @classmethod @@ -250,16 +290,42 @@ def get_kv_cache_shape( head_size: int, cache_dtype_str: str = "auto", ) -> tuple[int, ...]: + if _on_rdna35(): + return TritonAttentionBackend.get_kv_cache_shape( + num_blocks, block_size, num_kv_heads, head_size, cache_dtype_str + ) if block_size % 16 != 0: raise ValueError("Block size must be a multiple of 16.") return (2, num_blocks, block_size, num_kv_heads, head_size) + @classmethod + def get_kv_cache_stride_order( + cls, + include_num_layers_dimension: bool = False, + ) -> tuple[int, ...]: + if cls is RocmAttentionBackend and _on_rdna35(): + return TritonAttentionBackend.get_kv_cache_stride_order( + include_num_layers_dimension + ) + raise NotImplementedError + + @classmethod + def get_required_kv_cache_layout(cls) -> KVCacheLayoutType | None: + # Keys of one head contiguous in a page. The decode kernel is tuned + # on it, and the Triton paths it falls back to run faster on it too: + # 3-5 % prefill, 13-27 % batched decode. + if cls is RocmAttentionBackend and _on_rdna35(): + return "HND" + return None + @staticmethod def use_cascade_attention(*args, **kwargs) -> bool: return False @staticmethod - def get_builder_cls() -> type["RocmAttentionMetadataBuilder"]: + def get_builder_cls() -> type[AttentionMetadataBuilder]: + if _on_rdna35(): + return RocmAttentionRdna35MetadataBuilder return RocmAttentionMetadataBuilder @@ -581,3 +647,347 @@ def do_rope_and_kv_cache_update( flash_layout, is_fp8_kv_cache, ) + + +# The JIT-compiled module plus the scratch buffers sized for it. The module is +# a pybind extension built at runtime, so it has no static type. +_Built = tuple[Any, tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]] + +# Split-KV scratch, one set per variant and device, shared by every layer that +# runs it: layers launch one after another on one stream, and the counters are +# left clean by each launch. Sized for the scheduler's batch, up to a budget; +# a larger batch falls back to Triton. +_SCRATCH: dict[tuple[KernelVariant, torch.device], tuple[int, Any]] = {} +_SCRATCH_BUDGET = 64 * 1024**2 +# Mixed batches are split only from this head size up. +_SPLIT_MIN_HEAD_SIZE = 256 +_SPLIT_MIN_DECODES = 16 +_SPLIT_MIN_PREFILL = 256 +_SPLIT_NEEDS_LONG_PREFILL = {(2, 256)} +_SPLIT_SMALL_WINDOW = 512 + + +class RocmAttentionRdna35MetadataBuilder(TritonAttentionMetadataBuilder): + """Triton's metadata, with the batch ordered decodes first and the + number of leading uniform decodes counted, so that a batch mixing + prefills and decodes can send its decodes to the kernel.""" + + def __init__(self, *args, **kwargs): + super().__init__(*args, **kwargs) + self._init_reorder_batch_threshold(1, supports_spec_as_decode=True) + if self.reorder_batch_threshold is not None: + self.reorder_batch_threshold = min(self.reorder_batch_threshold, MAX_M) + + def build(self, common_prefix_len, common_attn_metadata, fast_build=False): + md = super().build(common_prefix_len, common_attn_metadata, fast_build) + num_decodes, num_decode_tokens = 0, 0 + if self.reorder_batch_threshold is not None: + num_decodes, _, num_decode_tokens, _ = split_decodes_and_prefills( + common_attn_metadata, + decode_threshold=self.reorder_batch_threshold, + require_uniform=True, + ) + md.num_decodes = num_decodes # type: ignore[attr-defined] + md.num_decode_tokens = num_decode_tokens # type: ignore[attr-defined] + return md + + +class RocmAttentionRdna35Impl(TritonAttentionImpl): + """ROCM_ATTN on gfx1151: the RDNA3.5 HIP kernel where it serves the call, + TRITON_ATTN's forward for everything else, on the same packed cache.""" + + def __init__(self, *args, **kwargs): + super().__init__(*args, **kwargs) + self._variant: KernelVariant | None = None + self._built: _Built | None = None + self._rejected: str | None = None + # Counters, so a test can assert the kernel really ran. A benchmark + # that silently falls back measures Triton and reports it as this + # backend, which is worse than an error. + self.kernel_calls = 0 + self.fallback_calls = 0 + # Mixed batches served as decodes on the kernel plus the rest on Triton. + self.split_calls = 0 + try: + from vllm.config import get_current_vllm_config + + max_seqs = get_current_vllm_config().scheduler_config.max_num_seqs + except Exception: + max_seqs = 1 + self._max_seqs = max(1, max_seqs) + self._cap = 0 + + def _reject(self, reason: str) -> None: + """Record why this shape falls back. Info, not a warning: as gfx1151's + default backend it falls back in normal operation (prefills, head + sizes the kernel does not build).""" + if self._rejected != reason: + self._rejected = reason + logger.info_once( + "ROCM_ATTN (gfx1151) falling back to Triton: %s", reason, scope="local" + ) + + def _prepare(self, kv_cache: torch.Tensor, **kwargs) -> _Built | None: + """Decide whether the kernel can serve this call, and build it if so. + + Every condition is checked rather than assumed. The kernel walks the + paged KV cache with its own address arithmetic instead of reading the + tensor's strides, so a layout it did not expect would not fault — it + would read the wrong addresses and return finite, wrong numbers. + + Returns the compiled module and its scratch buffers, or None to fall + back to Triton. + """ + if kwargs["alibi_slopes"] is not None or kwargs["sinks"] is not None: + self._reject("alibi/sinks unsupported") + return None + if kwargs["softcap"] or not kwargs["causal"]: + self._reject("softcap or non-causal unsupported") + return None + # vLLM passes a causal window of w keys as (w - 1, 0). + window = kwargs["window_size"] + win = 0 + if window is not None and window[0] >= 0: + if window[1] != 0: + self._reject(f"only causal sliding windows, got {window}") + return None + win = window[0] + 1 + # Features the kernel does not implement must not reach it silently. + if kwargs.get("mm_prefix_range") is not None: + self._reject("multimodal bidirectional prefix unsupported") + return None + if kwargs.get("rswa_prefix_lens") is not None: + self._reject("rswa unsupported") + return None + if kwargs.get("chunk_lookback", -1) >= 0: + self._reject("chunk lookback unsupported") + return None + # Not the descale tensors: on the unquantized path k_descale is still a + # broadcast of a 1.0 scale, so testing it for None never fires. + if kwargs["kv_quant_mode"] != KVQuantMode.NONE: + self._reject(f"KV quant mode {kwargs['kv_quant_mode']!r} unsupported") + return None + + # Every sequence with the same number of query tokens: decode, or + # speculative decode. Mixed batches (prefill in them) go to Triton. + nseq = kwargs["seqused_k"].shape[0] + max_m = kwargs["max_seqlen_q"] + if kwargs["q"].shape[0] != nseq * max_m: + self._reject("sequences of unequal query length") + return None + # A decode kernel: every distinct M is a build of its own, so a + # prompt served here would compile a variant per prompt length. + if max_m > MAX_M: + self._reject(f"{max_m} query tokens per sequence, more than {MAX_M}") + return None + dtype = kwargs["q"].dtype + if dtype not in (torch.float16, torch.bfloat16): + self._reject(f"kernel is fp16 or bf16 only, got {dtype}") + return None + if kv_cache.dtype != dtype: + self._reject(f"KV cache is {kv_cache.dtype}, query is {dtype}") + return None + if self.head_size not in SUPPORTED_HEAD_SIZES: + self._reject(f"head_size {self.head_size} not built") + return None + if self.num_heads % self.num_kv_heads: + self._reject("q heads must divide evenly over kv heads") + return None + if kv_cache.shape[2] % 16: + self._reject(f"block size {kv_cache.shape[2]} is not a multiple of 16") + return None + + # Logical KV cache order is (num_blocks, num_kv_heads, block_size, 2*hs). + q = kwargs["q"] + block_size = kv_cache.shape[2] + variant = variant_for( + self.num_heads, + self.num_kv_heads, + self.head_size, + max_m, + block_size, + 0 if kv_cache.stride(1) < kv_cache.stride(2) else 1, + win, + dtype, + nseq > 1, + ) + expected = expected_kv_cache_strides(variant) + actual = (kv_cache.stride(0), kv_cache.stride(1), kv_cache.stride(2)) + if actual != expected: + self._reject(f"KV strides {actual} != {expected} for this layout") + return None + + if self._variant != variant: + try: + module = load(variant) + except VariantBuildError as exc: + self._reject(str(exc).splitlines()[0]) + return None + key = (variant, q.device) + if key not in _SCRATCH: + if variant.batch: + per_seq = scratch_bytes(variant) + cap = min(self._max_seqs, max(1, _SCRATCH_BUDGET // per_seq)) + _SCRATCH[key] = (cap, make_scratch(variant, q.device, cap)) + else: + _SCRATCH[key] = (1, make_scratch(variant, q.device)) + self._cap, scratch = _SCRATCH[key] + self._built = (module, scratch) + self._variant = variant + if nseq > self._cap: + self._reject(f"{nseq} sequences, scratch holds {self._cap}") + return None + return self._built + + def forward( + self, + layer: torch.nn.Module, + query: torch.Tensor, + key: torch.Tensor, + value: torch.Tensor, + kv_cache: torch.Tensor, + attn_metadata: TritonAttentionMetadata, + output: torch.Tensor, + output_scale: torch.Tensor | None = None, + output_block_scale: torch.Tensor | None = None, + ) -> torch.Tensor: + """The kernel when it can serve the call, else Triton's forward.""" + if ( + attn_metadata is None + or self.attn_type in (AttentionType.ENCODER_ONLY, AttentionType.ENCODER) + or output_scale is not None + or output_block_scale is not None + ): + # Profiling, encoder attention and quantized output: Triton's. + return super().forward( + layer, + query, + key, + value, + kv_cache, + attn_metadata, + output, + output_scale, + output_block_scale, + ) + n = attn_metadata.num_actual_tokens + args = { + "q": query[:n], + "out": output[:n], + "seqused_k": attn_metadata.seq_lens, + "block_table": attn_metadata.block_table, + "max_seqlen_q": attn_metadata.max_query_len, + "softmax_scale": self.scale, + "causal": attn_metadata.causal, + "window_size": self.sliding_window, + "alibi_slopes": self.alibi_slopes, + "sinks": self.sinks, + "softcap": self.logits_soft_cap, + "mm_prefix_range": attn_metadata.mm_prefix_range_tensor, + "rswa_prefix_lens": attn_metadata.rswa_prefix_lens, + "chunk_lookback": self.chunk_lookback, + "kv_quant_mode": self._kv_quant_mode, + } + if self._split_mixed(layer, query, key, value, kv_cache, attn_metadata, args): + self.kernel_calls += 1 + self.split_calls += 1 + return output + built = self._prepare(kv_cache, **args) + if built is None: + self.fallback_calls += 1 + return super().forward( + layer, query, key, value, kv_cache, attn_metadata, output + ) + self.kernel_calls += 1 + self._launch(built, kv_cache, args) + return output + + def _split_mixed( + self, + layer: torch.nn.Module, + query: torch.Tensor, + key: torch.Tensor, + value: torch.Tensor, + kv_cache: torch.Tensor, + md: TritonAttentionMetadata, + kwargs: dict, + ) -> bool: + """Serve a batch of decodes followed by prefills in two launches: the + decodes on the kernel, the prefills on Triton. Returns False, having + done nothing, when the batch is not like that or the kernel cannot + take its decodes. + + Not under CUDA-graph capture: the split is host arithmetic on this + step's batch, and a captured graph would replay one step's split. + """ + nd = getattr(md, "num_decodes", 0) + ndt = getattr(md, "num_decode_tokens", 0) + nreq = kwargs["seqused_k"].shape[0] + if not 0 < nd < nreq or torch.cuda.is_current_stream_capturing(): + return False + # The prefills leave the launch they shared with the decodes, and a + # short one alone is latency-bound in Triton. Below D=256 the + # kernel's decode gain over Triton (1.0-1.1x) does not pay for that. + if self.head_size < _SPLIT_MIN_HEAD_SIZE: + return False + # A few decodes beside a short extend: the extend, alone, costs more + # than the kernel saves on the decodes (0.79-0.85x measured). + if nd < _SPLIT_MIN_DECODES and kwargs["q"].shape[0] - ndt < _SPLIT_MIN_PREFILL: + return False + # Two kv heads at D=256: Triton's short extend alone costs more than + # the kernel saves even on 16-32 decodes (0.73-0.93x measured). + prefill_tokens = kwargs["q"].shape[0] - ndt + pair = (self.num_kv_heads, self.head_size) + if pair in _SPLIT_NEEDS_LONG_PREFILL and prefill_tokens < _SPLIT_MIN_PREFILL: + return False + # A small window leaves a decode little KV to save on: a few of them + # do not pay for the prefill's own launch (0.95-0.98x at w512). + window = kwargs["window_size"] + if ( + window is not None + and 0 <= window[0] < _SPLIT_SMALL_WINDOW + and nd < _SPLIT_MIN_DECODES + ): + return False + dec = dict(kwargs) + for k in ("q", "out"): + dec[k] = kwargs[k][:ndt] + for k in ("seqused_k", "block_table"): + dec[k] = kwargs[k][:nd] + dec["max_seqlen_q"] = ndt // nd + built = self._prepare(kv_cache, **dec) + if built is None: + return False + self._launch(built, kv_cache, dec) + prefills = replace( + md, + num_actual_tokens=md.num_actual_tokens - ndt, + query_start_loc=md.query_start_loc[nd:] - ndt, + seq_lens=md.seq_lens[nd:], + block_table=md.block_table[nd:], + ) + super().forward( + layer, query[ndt:], key, value, kv_cache, prefills, kwargs["out"][ndt:] + ) + return True + + def _launch(self, built: _Built, kv_cache: torch.Tensor, kwargs: dict) -> None: + module, (acc, softmax_max, softmax_sum, arrivals) = built + # Called directly rather than through a registered custom op: this path + # is exercised under CUDA-graph capture, not torch.compile, so the op + # wrapper would only add indirection inside the region being measured. + module.decode_attn( + kwargs["q"], + kv_cache, + kwargs["block_table"], + kwargs["out"], + acc, + softmax_max, + softmax_sum, + arrivals, + # The device tensor, not max_seqlen_k: this runs under full + # CUDA-graph capture, where a host int would be frozen at its + # capture-time value for every replay. + kwargs["seqused_k"], + kwargs["softmax_scale"], + ) diff --git a/vllm/v1/attention/ops/rdna35_hip_decode.py b/vllm/v1/attention/ops/rdna35_hip_decode.py new file mode 100644 index 000000000000..a2a156c4e92d --- /dev/null +++ b/vllm/v1/attention/ops/rdna35_hip_decode.py @@ -0,0 +1,899 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Loader for the RDNA3.5 paged decode-attention kernel. + +The kernel in ``csrc/rocm/rdna35_decode_attn.cu`` takes its shapes as +compile-time defines, so one build serves exactly one shape tuple. The +variants the tables below name are compiled with the rest of vLLM into +``_rocm_C`` (the list is ``csrc/rocm/rdna35_attn/variants.def``, written by +``write_variants_def``); any other variant is served by Triton. + +The kernel tools (benchmarks/kernels/gfx1151_decode_attn) search variants +that are not in the list, so with ``VLLM_RDNA35_ATTN_JIT`` set every variant +is instead compiled here with ``torch.utils.cpp_extension.load``, from a +source checkout, in a few seconds rather than a vLLM rebuild. +""" + +import os +from collections.abc import Iterable +from concurrent.futures import ThreadPoolExecutor +from dataclasses import dataclass +from pathlib import Path +from typing import Any, TypedDict + +import torch + +from vllm.logger import init_logger + +logger = init_logger(__name__) + +# resolve() matters: ninja re-expands the path through a shell, so a symlinked +# checkout whose name contains '$' would be mangled into a missing file. +_CSRC = Path(__file__).resolve().parents[4] / "csrc" / "rocm" +_SOURCE = _CSRC / "rdna35_decode_attn.cu" +_LAUNCH_H = _CSRC / "rdna35_attn" / "launch.h" + + +@dataclass(frozen=True) +class KernelVariant: + """The compile-time shape tuple a single build is specialised for. + + One workgroup serves one (kv head, row group, KV segment): every query + row that reads a kv head -- GQA heads times max_m tokens -- is packed into + it, so the KV cache is read once per kv head rather than once per q head. + """ + + head_size: int + num_q_heads: int + num_kv_heads: int + max_m: int + block_size: int + layout: int # 0 = NHD, 1 = HND + # Most KV segments a (kv head, row group) is split over. The kernel + # activates clamp(nblocks // minb, 1, nseg) of them at run time, so short + # contexts do not pay for the cross-workgroup merge a long one needs. + nseg: int = 1 + # Row groups per kv head: its GQA*max_m rows split over rg workgroups that + # each read the whole of its KV, sharing it through L2. + rg: int = 1 + # Least KV blocks an active segment is given. + minb: int = 1 + # Waves per workgroup. + nw: int = 8 + # Waves sharing one key tile, each owning 1/dspl of the head dim. 0 keeps + # the kernel's rule. + dspl: int = 0 + # Waves sharing one key tile, each carrying 1/rspl of the row tiles: the + # in-workgroup form of rg, with the tile loaded once and shared in LDS. + rspl: int = 1 + # 1 keeps a second tile's loads in flight per wave (unshared tiles only; + # it needs a second tile's registers). + pf: int = 0 + # 1 stages split-KV partials in LDS and writes them a line at a time + # instead of from registers. + cpub: int = 0 + # 1 stages unshared tiles' V in LDS beside K, freeing the tile's + # registers so the next tile's loads go out before this one's compute. + vinlds: int = 0 + # 1 builds for a batch of sequences (grid.y), with scratch from + # make_scratch(max_seqs=...); 0 serves exactly one sequence. + batch: int = 0 + # 1 runs the per-q-head dot-product decomposition instead of WMMA: nseg + # segments per q head, nw waves, bfly its butterfly split, gt 1 to + # dispatch heads fastest. rg, minb, dspl, rspl, pf and cpub do not apply. + dot: int = 0 + bfly: int = 0 + gt: int = 0 + mutate: int = 0 + # Measurement only: skips blocks of work and returns wrong numbers. See + # the ABLATE comment in the kernel. + ablate: int = 0 + # Element type of Q, the KV cache and the output. + dtype: torch.dtype = torch.float16 + # Sliding window in keys (a query at p sees p-window+1 .. p); 0 is full + # causal attention. Only the window's pages are read. + window: int = 0 + + def __post_init__(self) -> None: + if self.dtype not in (torch.float16, torch.bfloat16): + raise ValueError(f"kernel is built for fp16 or bf16, not {self.dtype}") + + @property + def suffix(self) -> str: + return ( + f"d{self.head_size}_q{self.num_q_heads}_kv{self.num_kv_heads}" + f"_m{self.max_m}_bs{self.block_size}_l{self.layout}" + f"_n{self.nseg}_rg{self.rg}_mb{self.minb}_w{self.nw}" + f"{'' if not self.dspl else f'_ds{self.dspl}'}" + f"{'' if self.rspl == 1 else f'_rs{self.rspl}'}" + f"{'' if not self.pf else '_pf'}" + f"{'' if not self.cpub else '_cp'}" + f"{'' if not self.vinlds else '_vl'}" + f"{'' if not self.batch else '_batch'}" + f"{'' if not self.dot else f'_dot_bf{self.bfly}_gt{self.gt}'}" + f"_mut{self.mutate}" + f"{'' if not self.ablate else f'_ab{self.ablate}'}" + f"{'_bf16' if self.dtype == torch.bfloat16 else ''}" + f"{'' if not self.window else f'_win{self.window}'}" + ) + + @property + def name(self) -> str: + return f"rdna35_decode_{self.suffix}" + + @property + def rows_padded(self) -> int: + """Query rows one workgroup carries, rounded up to whole WMMA tiles.""" + gqa = self.num_q_heads // self.num_kv_heads + rows = gqa * self.max_m // self.rg + return -(-rows // 16) * 16 + + def scratch_shapes(self) -> tuple[tuple[int, ...], tuple[int, ...], int]: + """Shapes of the (acc, m/l) partials the split-KV merge goes through, + and the number of counters.""" + if self.dot: + rows = self.num_q_heads * self.nseg * self.max_m + cnt = self.num_q_heads + else: + rows = self.num_kv_heads * self.rg * self.nseg * self.rows_padded + cnt = 2 * self.num_kv_heads * self.rg + return (rows, self.head_size), (rows,), cnt + + +_loaded: dict[KernelVariant, Any] = {} + +VARIANTS_DEF = _CSRC / "rdna35_attn" / "variants.def" + + +def variant_key(variant: KernelVariant) -> list[int]: + """The fields of the variant's variants.def line, in its order.""" + v = variant + return [ + v.head_size, + v.num_q_heads, + v.num_kv_heads, + v.max_m, + v.block_size, + v.layout, + v.nseg, + v.rg, + v.minb, + v.nw, + v.dspl, + v.rspl, + v.pf, + v.cpub, + v.vinlds, + v.batch, + v.dot, + v.bfly, + v.gt, + v.window, + int(v.dtype == torch.bfloat16), + ] + + +def variant_group(variant: KernelVariant) -> str: + """The unit a variant is compiled in: one per configuration and dtype, so + re-tuning a configuration rebuilds only its own unit.""" + v = variant + dtype = "bf16" if v.dtype == torch.bfloat16 else "fp16" + return f"q{v.num_q_heads}_kv{v.num_kv_heads}_d{v.head_size}_w{v.window}_{dtype}" + + +def write_variants_def(variants: "Iterable[KernelVariant]", path: Path) -> None: + """Write the variants.def _rocm_C is built from.""" + lines = sorted( + { + f"RDNA35_VARIANT({variant_group(v)}, {', '.join(map(str, variant_key(v)))})" + for v in variants + if not (v.mutate or v.ablate) + } + ) + path.write_text( + "// Decode-attention variants compiled into _rocm_C (see macros.h for\n" + "// the fields). Generated by write_variants_def from the backend's\n" + "// tables; regenerate rather than edit:\n" + "// python -m vllm.v1.attention.ops.rdna35_hip_decode\n" + + "".join(f"{line}\n" for line in lines) + ) + + +class _Builtin: + """A variant compiled into _rocm_C, with the JIT module's interface.""" + + def __init__(self, index: int): + self.index = index + + def decode_attn(self, q, kv_cache, block_table, out, *scratch_and_args): + torch.ops._rocm_C.rdna35_decode_attn( + self.index, q, kv_cache, block_table, out, *scratch_and_args + ) + + +def _builtin_index(variant: KernelVariant) -> int: + """The variant's index in _rocm_C, or -1 if it was not built in.""" + if variant.mutate or variant.ablate: + return -1 + try: + op = torch.ops._rocm_C.rdna35_decode_variant + except (AttributeError, RuntimeError): + return -1 + return int(op(variant_key(variant))) + + +def _hip_runtime_ldflags() -> list[str]: + """Point the linker at libamdhip64. + + With the pip-installed ROCm SDK, torch sets ROCM_HOME to the venv root and + so searches ``/lib``, but the runtime actually ships inside + ``site-packages/_rocm_sdk_devel/lib``. + """ + sdk_lib = Path(torch.__file__).resolve().parents[1] / "_rocm_sdk_devel" / "lib" + return [f"-L{sdk_lib}"] if (sdk_lib / "libamdhip64.so").exists() else [] + + +def _staged_source(name: str) -> Path: + """Copy the kernel source into this variant's own build directory. + + On ROCm, torch hipifies each source before compiling it, and does so where + the source lives: it writes ``rdna35_decode_attn.hip`` next to the ``.cu`` + and records the result in ``hipify_python.HIPIFY_FINAL_RESULT``, a global + dict keyed by the absolute source path. Every variant compiles the same + file, so concurrent builds collide on that one key: the later one resets + the entry to a fresh record whose ``hipified_path`` is still the raw + ``.cu``, and the earlier build reads that back and compiles unhipified + CUDA, failing on ``cuda_runtime_api.h``. The generated ``.hip`` is shared + too, and is unlinked by whichever build's ``GeneratedFileCleaner`` exits + first, under the feet of the others. Giving each variant its own copy + gives it its own hipify key, its own ``.hip`` and its own cleanup. + + The header it includes is staged beside it for the same reason: hipify + writes a ``_hip`` copy of every header it follows next to that header, so + one found in csrc/rocm would be rewritten in the source tree. + """ + from torch.utils.cpp_extension import _get_build_directory + + build_dir = Path(_get_build_directory(name, False)) + for src, dst in ( + (_SOURCE, build_dir / _SOURCE.name), + (_LAUNCH_H, build_dir / "rdna35_attn" / _LAUNCH_H.name), + ): + text = src.read_bytes() + # Rewrite only on a real change: ninja keys off mtime, so touching + # this every time would force a full rebuild on every warm start. + if not dst.is_file() or dst.read_bytes() != text: + dst.parent.mkdir(exist_ok=True) + dst.write_bytes(text) + return build_dir / _SOURCE.name + + +class UnexpectedBuildError(RuntimeError): + """A variant was requested that precompilation did not cover.""" + + +class VariantBuildError(RuntimeError): + """A variant does not compile -- a static_assert refused its shape.""" + + +# Variants that already failed to build. A shape the kernel refuses fails the +# same way every time, and a rebuild costs ~20 s, so the first failure is +# remembered and every later request fails at once. +_failed: dict[KernelVariant, str] = {} + + +_sealed = False + + +def seal(on: bool = True) -> None: + """Refuse to JIT-build any further variant. + + For benchmark harnesses. A build landing inside ``do_bench`` does not just + add time: ``do_bench`` sizes its repeat count from a calibration run, so a + build during calibration corrupts the median it reports -- measured here as + 104.59 us against a true 29.99, a 258% error, with no warning. Harnesses + used to dodge that with a discarded warm-up call per cell, which doubled + the run. Sealing after ``precompile`` turns the same hazard into a loud + failure instead, and the warm-up call can go. + + Args: + on: True to refuse builds, False to allow them again. + """ + global _sealed + _sealed = on + + +def load(variant: KernelVariant) -> Any: + """The variant built into _rocm_C; with VLLM_RDNA35_ATTN_JIT, a JIT build. + + Raises: + VariantBuildError: The variant is not built into _rocm_C (JIT off), or + the JIT build failed: the kernel refuses the shape, or there is no + source checkout or compiler. + UnexpectedBuildError: The loader is sealed and the variant is not + built yet. + """ + import vllm.envs as envs + + if variant in _loaded: + return _loaded[variant] + if variant in _failed: + raise VariantBuildError(_failed[variant]) + if not envs.VLLM_RDNA35_ATTN_JIT: + index = _builtin_index(variant) + if index < 0: + _failed[variant] = f"{variant.name} is not built into _rocm_C" + raise VariantBuildError(_failed[variant]) + _loaded[variant] = _Builtin(index) + return _loaded[variant] + if _sealed: + raise UnexpectedBuildError( + f"{variant.name} was not precompiled, and building it now would " + "land inside the timed region. Precompile it or unseal." + ) + + if not _SOURCE.is_file(): + _failed[variant] = ( + f"{variant.name} is not built into _rocm_C, and there is no source " + f"checkout to build it from ({_SOURCE})" + ) + raise VariantBuildError(_failed[variant]) + + from torch.utils.cpp_extension import load as load_extension + + source = _staged_source(variant.name) + + flags = [ + "-O3", + "-DRDNA35_TORCH_EXT", + # HIP resolves device functions by name across the whole process, so two + # variants sharing the symbol `decode_attn` would both dispatch to + # whichever registered first — returning plausible, wrong numbers. Give + # each build its own symbols so variants can coexist. + f"-Ddecode_attn=decode_attn_{variant.suffix}", + f"-DHEAD_DIM={variant.head_size}", + f"-DNUM_Q_HEADS={variant.num_q_heads}", + f"-DNUM_KV_HEADS={variant.num_kv_heads}", + f"-DMAXM={variant.max_m}", + f"-DBS={variant.block_size}", + f"-DLAYOUT={variant.layout}", + f"-DNSEG={variant.nseg}", + f"-DRG={variant.rg}", + f"-DMINB={variant.minb}", + f"-DNW={variant.nw}", + *([f"-DDSPL={variant.dspl}"] if variant.dspl else []), + f"-DRSPL={variant.rspl}", + f"-DPF={variant.pf}", + f"-DCPUB={variant.cpub}", + f"-DVINLDS={variant.vinlds}", + f"-DBATCH={variant.batch}", + f"-DDOT={variant.dot}", + f"-DBFLY={variant.bfly}", + f"-DGT={variant.gt}", + f"-DWIN={variant.window}", + f"-DMUTATE={variant.mutate}", + f"-DABLATE={variant.ablate}", + f"-DKV_BF16={int(variant.dtype == torch.bfloat16)}", + ] + logger.info("Compiling %s", variant.name) + try: + module = load_extension( + name=variant.name, + sources=[str(source)], + extra_cuda_cflags=flags, + extra_ldflags=_hip_runtime_ldflags(), + ) + # Not only a refused shape: a machine without the ROCm compiler or ninja + # fails here too, and must fall back rather than fail the model. + except Exception as exc: + _failed[variant] = f"{variant.name} does not build: {exc}" + raise VariantBuildError(_failed[variant]) from exc + _loaded[variant] = module + return module + + +# Measured peak RSS of one build: 1.15 GB, almost all of it torch's headers. +# Budgeted at 2 GB because the measurement is a single sample and the cost of +# being wrong is an OOM that takes the machine with it, not a slow build. +_BUILD_RSS_BYTES = 2 << 30 + + +def _parallel_budget(n: int) -> int: + """How many builds fit in memory, not how many cores exist. + + A build peaks at 1.15 GB, so 32 of them want more than this machine has -- + running one per core filled 30 GB of RAM and had to be killed. Memory is + the binding constraint here, and it is read at call time because what is + free depends on what else is running. + """ + free = _available_bytes() + by_mem = max(1, int(free * 0.7) // _BUILD_RSS_BYTES) + return max(1, min(n, os.cpu_count() or 8, by_mem)) + + +def _available_bytes() -> int: + try: + with open("/proc/meminfo") as fh: + for line in fh: + if line.startswith("MemAvailable:"): + return int(line.split()[1]) * 1024 + except OSError: + pass + return 8 << 30 + + +def precompile(variants: "Iterable[KernelVariant]", workers: int | None = None) -> None: + """Build several variants at once. + + Each variant is already its own translation unit, its own ninja invocation + and -- since ``_staged_source`` gives it a private copy of the kernel -- + its own hipify output, so concurrent builds do not collide. + + This matters more than it looks. A cold build is ~19.6 s, of which the + kernel itself is 0.57 s -- the rest is torch/extension.h and the link. So + the cost is per-variant and fixed, and the 27 distinct variants of the + shape table cost nine minutes serially before a single measurement can be + taken. Spread over the machine that is well under a minute. + """ + # Deduplicated: two threads building the same variant would race on the + # same build directory and the same _loaded entry. Callers hand over + # whatever their loop produced -- roofline builds both layouts per shape + # and shapes repeat -- so duplicates are the normal case, not a mistake. + todo = list(dict.fromkeys(v for v in variants if v not in _loaded)) + if not todo: + return + n = workers or _parallel_budget(len(todo)) + logger.info( + "Compiling %d kernel variants across %d workers (%.1f GiB available)", + len(todo), + n, + _available_bytes() / (1 << 30), + ) + with ThreadPoolExecutor(max_workers=n) as pool: + futures = {pool.submit(load, v): v for v in todo} + for fut, variant in futures.items(): + try: + fut.result() + except Exception: + # A shape the kernel refuses to build is a normal outcome -- + # D=96 is three fp16 per lane and static_asserts. This is a + # warm-up, not a gate: let the real load report it where the + # caller already handles a fallback. + logger.debug( + "%s did not build; leaving it to the caller", + variant.name, + exc_info=True, + ) + + +def make_scratch( + variant: KernelVariant, device: torch.device, max_seqs: int = 0 +) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]: + """Allocate the partials once. + + They are passed into the op rather than allocated inside it because the op + runs under CUDA-graph capture, where an allocation would break the graph. + ``max_seqs`` > 0 gives every tensor a leading sequence dimension, for + batches of up to that many sequences; 0 is one sequence without it. + """ + acc_shape, ml_shape, cnt = variant.scratch_shapes() + lead = (max_seqs,) if max_seqs else () + opts = {"dtype": torch.float32, "device": device} + return ( + torch.empty(lead + acc_shape, **opts), + torch.empty(lead + ml_shape, **opts), + torch.empty(lead + ml_shape, **opts), + # Arrival counters, then merge generations, one each per (kv head, + # row group). Zeroed once: the kernel resets the counters as it + # consumes them and only compares generations for change, so every + # later launch starts clean without the host writing here -- which it + # could not do under graph capture anyway. + torch.zeros(lead + (cnt,), dtype=torch.int32, device=device), + ) + + +def scratch_bytes(variant: KernelVariant) -> int: + """Bytes of scratch one sequence needs.""" + acc_shape, ml_shape, cnt = variant.scratch_shapes() + return 4 * (acc_shape[0] * acc_shape[1] + 2 * ml_shape[0] + cnt) + + +def expected_kv_cache_strides(variant: KernelVariant) -> tuple[int, int, int]: + """Element strides of (block, kv head, token) the kernel indexes. + + The kernel walks the KV cache with its own arithmetic instead of reading + the tensor's strides, so the caller must confirm the tensor really is laid + out this way. Getting this wrong is silent: the kernel would read the wrong + addresses and still return finite numbers. + """ + kv_row = 2 * variant.head_size + page = variant.block_size * variant.num_kv_heads * kv_row + if variant.layout == 0: # NHD: (NB, BS, HKV, 2D) + return page, kv_row, variant.num_kv_heads * kv_row + return page, variant.block_size * kv_row, kv_row # HND: (NB, HKV, BS, 2D) + + +# Query tokens per sequence the kernel serves: decode and speculative decode. +MAX_M = 8 + +# The kernel's V and K slices are one b64 or b128 per lane, which the head +# dims below allow; its GQA packing assumes the q heads divide evenly over the +# kv heads. +SUPPORTED_HEAD_SIZES = (64, 128, 256, 512) + +# Workgroups the long-context split aims for. Fewer, longer streams keep +# LPDDR5X closer to its peak than many short ones: a tile-structured stream +# with this kernel's access order measured 97.9 % of peak at 16 workgroups, +# 96.4 % at 40 and 90-94 % at 64 and up. +_TARGET_WORKGROUPS = 16 + + +def _segments_for(num_kv_heads: int, rg: int) -> int: + """Most KV segments per (kv head, row group), for _TARGET_WORKGROUPS.""" + return max(1, _TARGET_WORKGROUPS // (num_kv_heads * rg)) + + +def _rows_split(gqa: int, max_m: int, head_size: int) -> tuple[int, int]: + """Row groups, and a d split if one is needed, for one kv head's rows. + + A wave's accumulator is its rows times the d it owns, and it has to stay + near 64 VGPRs: 16 rows at 128 d, 32 at 64. Row groups are the cheaper way + down -- they only re-read the kv head's KV from L2 -- but they split whole + q heads, so a GQA with no fitting divisor (7, 5) halves each wave's d + instead, which measured 63-73 % of roof against 48-69 % for rg = GQA. At + D=512 it is LDS rather than VGPRs: Q and the K tile leave room for one row + tile. + + Returns: + `(rg, dspl)`, dspl 0 meaning the kernel's own rule. + """ + rows = gqa * max_m + cap = 32 if head_size == 64 else 16 + for rg in range(1, gqa + 1): + if gqa % rg == 0 and rows // rg <= cap: + break + if rows // rg <= cap and (rg == 1 or rows // rg >= 8): + return rg, 0 + if head_size == 128 and rows <= 2 * cap: + return 1, 2 + return rg, 0 + + +# Best knobs measured per configuration, keyed on (Hq, Hkv, D, M). +# +# The heuristics in knobs_for are rules fitted to the whole shape table; this +# is the exceptions list, and it wins where both apply. Rows carry only the +# knobs actually measured -- a partial row is normal, and a configuration +# absent here behaves exactly as the heuristics say, so landing a row can only +# affect the configuration it names. +# +# Every row comes from a measurement (tools/tune.py); do not add one without +# it. +class _Knobs(TypedDict, total=False): + """Launch knobs a configuration may pin. Total=False: a row sets only what + was measured.""" + + nseg: int + rg: int + minb: int + nw: int + dspl: int + rspl: int + pf: int + cpub: int + vinlds: int + dot: int + bfly: int + gt: int + + +# Every shipped configuration, measured: coordinate descent over (nw, dspl, +# rg, workgroup target, minb) scored by geomean %roof over the seven contexts +# matrix.py reports (tools/tune.py). dspl 0 is the kernel's own rule. +_TUNED: dict[tuple[int, int, int, int], _Knobs] = { + # D=64 + (14, 2, 64, 1): {"nw": 4, "rg": 1, "minb": 2, "nseg": 4, "pf": 1, "cpub": 1}, + (14, 2, 64, 3): {"nw": 4, "rg": 1, "minb": 2, "nseg": 8, "rspl": 2}, + (14, 2, 64, 4): {"nseg": 4, "rg": 1, "minb": 2, "nw": 4, "vinlds": 1}, + (16, 2, 64, 1): {"nseg": 8, "rg": 1, "minb": 1, "nw": 4, "dspl": 0}, + (16, 2, 64, 4): {"nseg": 4, "rg": 2, "minb": 2, "nw": 4, "vinlds": 1}, + (32, 8, 64, 1): {"nw": 4, "rg": 1, "minb": 2, "nseg": 1, "pf": 1}, + (32, 8, 64, 2): {"nw": 8, "rg": 1, "minb": 1, "nseg": 1, "pf": 1}, + (32, 8, 64, 3): {"nw": 8, "rg": 2, "minb": 2, "nseg": 1}, + (32, 8, 64, 4): {"nw": 4, "rg": 1, "minb": 1, "nseg": 1, "pf": 1}, + (32, 32, 64, 1): {"nseg": 1, "rg": 1, "minb": 2, "nw": 2, "dspl": 0}, + (32, 32, 64, 2): {"nw": 2, "rg": 1, "minb": 1, "nseg": 1}, + (32, 32, 64, 3): {"nw": 2, "rg": 1, "minb": 1, "nseg": 1, "cpub": 1}, + (32, 32, 64, 4): {"nseg": 1, "rg": 1, "minb": 4, "nw": 2, "dspl": 0}, + # D=128 + (16, 2, 128, 1): {"nseg": 4, "rg": 2, "minb": 2, "nw": 4, "vinlds": 1}, + (16, 2, 128, 2): {"nw": 8, "rg": 1, "minb": 1, "nseg": 8}, + (16, 2, 128, 4): {"nw": 4, "rg": 4, "minb": 2, "nseg": 2, "vinlds": 1}, + (32, 2, 128, 1): {"nw": 4, "rg": 4, "minb": 2, "nseg": 2, "cpub": 1, "vinlds": 1}, + (32, 2, 128, 4): {"nseg": 4, "rg": 4, "minb": 2, "nw": 4, "vinlds": 1}, + (28, 4, 128, 1): {"nseg": 4, "rg": 1, "minb": 1, "nw": 4, "dspl": 2}, + (28, 4, 128, 2): {"nw": 8, "rg": 1, "minb": 4, "nseg": 4, "cpub": 1}, + (28, 4, 128, 3): { + "nw": 4, + "rg": 1, + "minb": 2, + "nseg": 8, + "dspl": 2, + "rspl": 2, + "pf": 1, + }, + (28, 4, 128, 4): {"nseg": 4, "rg": 1, "minb": 1, "nw": 8, "dspl": 2}, + (32, 4, 128, 1): {"nw": 4, "rg": 2, "minb": 4, "nseg": 2, "vinlds": 1}, + (32, 4, 128, 2): {"nw": 8, "rg": 2, "minb": 4, "nseg": 4}, + (32, 4, 128, 4): {"nw": 4, "rg": 4, "minb": 1, "nseg": 1, "vinlds": 1}, + (16, 8, 128, 1): {"nseg": 1, "rg": 2, "minb": 2, "nw": 4, "dspl": 0}, + (16, 8, 128, 4): {"nw": 4, "rg": 2, "minb": 4, "nseg": 1}, + (24, 8, 128, 1): {"nw": 4, "rg": 1, "minb": 1, "nseg": 1, "cpub": 1, "vinlds": 1}, + (24, 8, 128, 2): {"nw": 8, "rg": 1, "minb": 4, "nseg": 2}, + (24, 8, 128, 3): {"nw": 4, "rg": 1, "minb": 2, "nseg": 1, "cpub": 1, "vinlds": 1}, + (24, 8, 128, 4): {"nseg": 1, "rg": 1, "minb": 1, "nw": 4, "dspl": 0}, + (32, 8, 128, 1): {"nseg": 1, "rg": 2, "minb": 2, "nw": 4, "dspl": 0}, + (32, 8, 128, 3): {"nw": 8, "rg": 2, "minb": 2, "nseg": 1}, + (32, 8, 128, 4): {"nseg": 1, "rg": 2, "minb": 2, "nw": 4, "dspl": 0}, + (40, 8, 128, 1): {"nseg": 1, "rg": 1, "minb": 1, "nw": 4, "dspl": 0}, + (40, 8, 128, 4): {"nw": 8, "rg": 1, "minb": 4, "nseg": 2, "dspl": 2, "rspl": 2}, + (10, 10, 128, 1): {"nseg": 1, "rg": 1, "minb": 2, "nw": 4, "vinlds": 1}, + (10, 10, 128, 2): {"nw": 4, "rg": 1, "minb": 4, "nseg": 1, "vinlds": 1}, + (10, 10, 128, 3): {"nw": 8, "rg": 1, "minb": 2, "nseg": 1, "cpub": 1}, + (10, 10, 128, 4): {"nw": 2, "rg": 1, "minb": 2, "nseg": 1, "vinlds": 1}, + (32, 32, 128, 1): {"nseg": 1, "rg": 1, "minb": 2, "nw": 2, "dspl": 0}, + (32, 32, 128, 4): {"nseg": 1, "rg": 1, "minb": 1, "nw": 2, "dspl": 0}, + # D=256 + (8, 1, 256, 1): {"dot": 1, "nw": 8, "nseg": 4, "bfly": 4}, + (8, 1, 256, 2): {"nw": 2, "rg": 1, "minb": 1, "nseg": 16}, + (8, 1, 256, 4): {"nseg": 16, "rg": 2, "minb": 1, "nw": 4, "dspl": 2}, + (8, 2, 256, 1): {"dot": 1, "nw": 8, "nseg": 2, "bfly": 4}, + (8, 2, 256, 2): {"dot": 1, "nw": 8, "nseg": 4, "bfly": 4, "gt": 1}, + (8, 2, 256, 4): {"nseg": 4, "rg": 2, "minb": 1, "nw": 8, "dspl": 0}, + (16, 2, 256, 1): {"nseg": 8, "rg": 2, "minb": 2, "nw": 8, "dspl": 0}, + (16, 2, 256, 4): {"nseg": 8, "rg": 2, "minb": 1, "nw": 4, "dspl": 2}, + (8, 4, 256, 1): {"dot": 1, "nw": 8, "nseg": 2, "bfly": 4, "gt": 1}, + (8, 4, 256, 2): {"dot": 1, "nw": 8, "nseg": 2, "bfly": 4, "gt": 1}, + (8, 4, 256, 4): { + "nw": 4, + "rg": 2, + "minb": 2, + "nseg": 2, + "dspl": 4, + "pf": 1, + "cpub": 1, + }, + (16, 4, 256, 1): {"dot": 1, "nw": 8, "nseg": 2, "bfly": 2, "gt": 1}, + (16, 4, 256, 2): {"dot": 1, "nw": 8, "nseg": 2, "bfly": 2, "gt": 1}, + (16, 4, 256, 4): {"nw": 4, "rg": 2, "minb": 1, "nseg": 2}, + (24, 4, 256, 1): {"dot": 1, "nw": 8, "nseg": 2, "bfly": 2}, + (24, 4, 256, 2): {"dot": 1, "nw": 4, "nseg": 4, "bfly": 2, "gt": 1}, + (24, 4, 256, 3): {"nw": 4, "rg": 1, "minb": 2, "nseg": 4, "rspl": 2}, + (24, 4, 256, 4): { + "nw": 4, + "rg": 1, + "minb": 1, + "nseg": 4, + "rspl": 2, + "pf": 1, + "cpub": 1, + }, + (16, 8, 256, 1): {"dot": 1, "nw": 4, "nseg": 2, "gt": 1}, + (16, 8, 256, 2): {"nw": 2, "rg": 1, "minb": 4, "nseg": 2, "vinlds": 1}, + (16, 8, 256, 3): {"nw": 2, "rg": 1, "minb": 2, "nseg": 2, "vinlds": 1}, + (16, 8, 256, 4): {"nseg": 2, "rg": 1, "minb": 4, "nw": 4, "dspl": 4}, + # D=512 + (8, 1, 512, 1): {"dot": 1, "nw": 8, "nseg": 4, "bfly": 4}, + (8, 1, 512, 2): {"dot": 1, "nw": 8, "nseg": 4, "bfly": 4}, + (8, 1, 512, 4): {"nseg": 8, "rg": 2, "minb": 1, "nw": 8, "dspl": 0}, + (16, 1, 512, 1): {"nw": 8, "rg": 1, "minb": 1, "nseg": 16, "dspl": 8}, + (16, 1, 512, 3): {"nw": 8, "rg": 4, "minb": 1, "nseg": 8}, + (16, 1, 512, 4): {"nseg": 8, "rg": 4, "minb": 1, "nw": 8, "dspl": 0}, + (8, 2, 512, 1): {"dot": 1, "nw": 8, "nseg": 2}, + (8, 2, 512, 2): {"dot": 1, "nw": 8, "nseg": 2, "bfly": 4}, + (8, 2, 512, 3): {"nw": 4, "rg": 1, "minb": 2, "nseg": 4, "cpub": 1, "vinlds": 1}, + (8, 2, 512, 4): {"nw": 8, "rg": 2, "minb": 1, "nseg": 4, "dspl": 8, "pf": 1}, + (16, 2, 512, 1): {"dot": 1, "nw": 8, "nseg": 2, "bfly": 4}, + (16, 2, 512, 2): {"nw": 4, "rg": 2, "minb": 1, "nseg": 4, "vinlds": 1}, + (16, 2, 512, 3): {"nw": 8, "rg": 2, "minb": 1, "nseg": 4, "cpub": 1}, + (16, 2, 512, 4): {"nseg": 4, "rg": 2, "minb": 1, "nw": 8, "dspl": 0}, + (32, 4, 512, 1): {"dot": 1, "nw": 4, "nseg": 2, "bfly": 4}, + (32, 4, 512, 2): {"dot": 1, "nw": 4, "nseg": 2}, + (32, 4, 512, 3): {"nw": 8, "rg": 2, "minb": 2, "nseg": 2, "cpub": 1}, + (32, 4, 512, 4): {"nseg": 2, "rg": 2, "minb": 2, "nw": 8, "dspl": 0}, +} + +# bf16 rows that differ from _TUNED. The dot decomposition's bf16 products +# (v_dot2_f32_bf16) fall behind at long context -- 0.73-0.97x of the WMMA +# rows at 32k -- so in bf16 these shapes keep their WMMA configuration. +_TUNED_BF16: dict[tuple[int, int, int, int], _Knobs] = { + (8, 1, 256, 1): {"nseg": 8, "rg": 2, "minb": 1, "nw": 8, "dspl": 0}, + (8, 2, 256, 1): {"nseg": 8, "rg": 1, "minb": 1, "nw": 4, "dspl": 4}, + (8, 2, 256, 2): {"rg": 1, "nseg": 8, "nw": 8, "minb": 1}, + (8, 4, 256, 1): {"nseg": 4, "rg": 1, "minb": 1, "nw": 2, "dspl": 0}, + (8, 4, 256, 2): {"rg": 1, "nseg": 4, "nw": 8, "minb": 1}, + (16, 4, 256, 1): {"nseg": 4, "rg": 1, "minb": 1, "nw": 2, "dspl": 0}, + (16, 4, 256, 2): {"rg": 1, "nseg": 4, "nw": 8, "minb": 1}, + (24, 4, 256, 1): {"nseg": 4, "rg": 1, "minb": 1, "nw": 8, "dspl": 0}, + (24, 4, 256, 2): {"rg": 1, "nseg": 4, "nw": 8, "minb": 1}, + (16, 8, 256, 1): {"rg": 1, "nseg": 2, "nw": 8, "minb": 2}, + (8, 1, 512, 1): {"nseg": 16, "rg": 1, "minb": 1, "nw": 8, "dspl": 8}, + (8, 1, 512, 2): {"rg": 1, "nseg": 16, "nw": 8, "minb": 1}, + (8, 2, 512, 1): {"nseg": 8, "rg": 1, "minb": 1, "nw": 4, "dspl": 0}, + (8, 2, 512, 2): {"rg": 1, "nseg": 8, "nw": 8, "minb": 1}, + (16, 2, 512, 1): {"nseg": 8, "rg": 1, "minb": 1, "nw": 4, "dspl": 0}, + (32, 4, 512, 1): {"nseg": 4, "rg": 1, "minb": 1, "nw": 4, "dspl": 0}, + (32, 4, 512, 2): {"rg": 1, "nseg": 4, "nw": 8, "minb": 1}, +} + +# Sliding-window layers, keyed on (Hq, Hkv, D, M, window). They read only the +# window's keys whatever the sequence length, so they are tuned apart from the +# full-attention layers that share their shape (gemma interleaves the two). +_TUNED_SWA: dict[tuple[int, int, int, int, int], _Knobs] = { + (8, 1, 256, 1, 512): {"dot": 1, "nw": 8, "nseg": 4, "bfly": 4, "gt": 1}, + (8, 1, 256, 2, 512): {"dot": 1, "nw": 8, "nseg": 8, "bfly": 2}, + (8, 1, 256, 3, 512): {"dot": 1, "nw": 8, "nseg": 4, "bfly": 4}, + (8, 1, 256, 4, 512): { + "nw": 4, + "rg": 1, + "minb": 1, + "nseg": 16, + "rspl": 2, + "pf": 1, + "cpub": 1, + }, + (8, 2, 256, 1, 512): {"dot": 1, "nw": 8, "nseg": 2, "bfly": 4, "gt": 1}, + (8, 2, 256, 2, 512): {"dot": 1, "nw": 8, "nseg": 4, "bfly": 4, "gt": 1}, + (8, 2, 256, 3, 512): {"nw": 4, "rg": 1, "minb": 1, "nseg": 8, "dspl": 4, "pf": 1}, + (8, 2, 256, 4, 512): {"nw": 4, "rg": 1, "minb": 1, "nseg": 8, "dspl": 4, "pf": 1}, + (8, 4, 256, 1, 1024): {"dot": 1, "nw": 8, "nseg": 2, "bfly": 4, "gt": 1}, + (8, 4, 256, 2, 1024): {"dot": 1, "nw": 8, "nseg": 4, "bfly": 4, "gt": 1}, + (8, 4, 256, 3, 1024): {"nw": 4, "rg": 1, "minb": 2, "nseg": 2, "dspl": 4, "pf": 1}, + (8, 4, 256, 4, 1024): { + "nw": 4, + "rg": 1, + "minb": 4, + "nseg": 2, + "dspl": 4, + "pf": 1, + "cpub": 1, + }, + (8, 4, 256, 1, 4096): {"dot": 1, "nw": 8, "nseg": 2, "bfly": 4, "gt": 1}, + (8, 4, 256, 2, 4096): {"nw": 4, "rg": 2, "minb": 1, "nseg": 2}, + (8, 4, 256, 4, 4096): { + "nw": 4, + "rg": 1, + "minb": 1, + "nseg": 2, + "dspl": 2, + "vinlds": 1, + }, + (16, 8, 256, 1, 1024): {"dot": 1, "nw": 4, "nseg": 1, "bfly": 4, "gt": 1}, + (16, 8, 256, 2, 1024): {"dot": 1, "nw": 8, "nseg": 2, "bfly": 4, "gt": 1}, + (16, 8, 256, 3, 1024): {"dot": 1, "nw": 8, "nseg": 2, "bfly": 4, "gt": 1}, + (16, 8, 256, 4, 1024): {"nw": 4, "rg": 2, "minb": 2, "nseg": 1}, + (32, 16, 256, 1, 1024): {"dot": 1, "nw": 4, "nseg": 1, "bfly": 4, "gt": 1}, + (32, 16, 256, 2, 1024): {"nw": 4, "rg": 1, "minb": 2, "nseg": 1, "dspl": 4}, + (32, 16, 256, 3, 1024): {"nw": 2, "rg": 1, "minb": 4, "nseg": 1, "vinlds": 1}, + (32, 16, 256, 4, 1024): {"nw": 2, "rg": 1, "minb": 2, "nseg": 1, "vinlds": 1}, +} + + +def knobs_for( + num_q_heads: int, + num_kv_heads: int, + head_size: int, + max_m: int, + window: int = 0, + dtype: torch.dtype = torch.float16, + batch: bool = False, +) -> _Knobs: + """Launch knobs for one configuration: the heuristics, then the measured + overrides on top. + + Args: + num_q_heads: Query heads. + num_kv_heads: KV heads. + head_size: Head dimension. + max_m: Query tokens per sequence. + window: Sliding window in keys, 0 for full attention. + dtype: Element type; bf16 rows may differ (`_TUNED_BF16`). + batch: More than one sequence. The dot decomposition buys latency for + one sequence and loses to WMMA once a batch fills the machine, so + a dot row gives way to its WMMA row (`_TUNED_BF16`) or the + heuristics. + + Returns: + Keyword arguments for `KernelVariant`. + """ + if window: + tuned = _TUNED_SWA.get( + (num_q_heads, num_kv_heads, head_size, max_m, window), {} + ) + else: + key = (num_q_heads, num_kv_heads, head_size, max_m) + tuned = _TUNED.get(key, {}) + if dtype == torch.bfloat16 or (batch and tuned.get("dot")): + tuned = _TUNED_BF16.get(key, tuned) + if batch and tuned.get("dot"): + tuned = {} + rg, dspl = _rows_split(num_q_heads // num_kv_heads, max_m, head_size) + rg = tuned.get("rg", rg) + knobs: _Knobs = { + "rg": rg, + "nseg": _segments_for(num_kv_heads, rg), + # Four waves where a wave carries the whole head dim: the fewer tiles + # a workgroup holds, the more of them are in flight at short context. + "nw": 4 if head_size <= 128 else 8, + # Short contexts want one segment per block to fill the machine; long + # ones a few blocks per segment so a workgroup overlaps its own tiles. + # Only kv-head-rich configurations have the parallelism to spare. + "minb": 2 if num_kv_heads * rg >= 8 else 1, + } + if dspl: + knobs["dspl"] = dspl + knobs.update(tuned) + return knobs + + +def variant_for( + num_q_heads: int, + num_kv_heads: int, + head_size: int, + max_m: int, + block_size: int, + layout: int, + window: int, + dtype: torch.dtype, + batch: bool, +) -> KernelVariant: + """The kernel build that serves one call.""" + return KernelVariant( + **knobs_for(num_q_heads, num_kv_heads, head_size, max_m, window, dtype, batch), + head_size=head_size, + num_q_heads=num_q_heads, + num_kv_heads=num_kv_heads, + max_m=max_m, + block_size=block_size, + layout=layout, + dtype=dtype, + window=window, + batch=int(batch), + ) + + +# Page sizes _rocm_C carries kernels for, 16 everywhere plus what vLLM gives +# these layers when it evens out page sizes across a model's layers -- two +# head sizes in gemma-4, the GDN state in Qwen3.5/3.6 -- measured per model. +# Any other page size is served by Triton, which at these sizes is 2.5-6.5x +# slower than the kernel on the bs=16 knobs. +_BUILT_BLOCK_SIZES: dict[tuple[int, int, int, int], tuple[int, ...]] = { + (8, 1, 256, 512): (16, 32), # gemma-4-E2B sliding layers + (8, 2, 256, 512): (16, 32), # gemma-4-E4B sliding layers + (16, 2, 512, 0): (16, 32), # gemma-4-26B-A4B full layers + (32, 4, 512, 0): (16, 32), # gemma-4-31B full layers + (8, 2, 256, 0): (16, 544), # Qwen3.5-0.8B, -2B + (16, 4, 256, 0): (16, 528), # Qwen3.5-9B + (16, 2, 256, 0): (16, 1056), # Qwen3.5-35B-A3B, Qwen3.6-35B-A3B + (24, 4, 256, 0): (16, 784), # Qwen3.6-27B +} + + +def built_variants() -> list[KernelVariant]: + """Every variant built into _rocm_C (variants.def): each configuration + of the tables at every query length the kernel serves, in both dtypes, for + one sequence and for a batch, on the HND cache the backend requires.""" + configs = {(*k[:3], 0) for k in _TUNED} | {(*k[:3], k[4]) for k in _TUNED_SWA} + return [ + variant_for(hq, hkv, d, m, bs, 1, window, dtype, batch) + for hq, hkv, d, window in sorted(configs) + for m in range(1, MAX_M + 1) + for dtype in (torch.float16, torch.bfloat16) + for batch in (False, True) + for bs in _BUILT_BLOCK_SIZES.get((hq, hkv, d, window), (16,)) + ] + + +if __name__ == "__main__": + write_variants_def(built_variants(), VARIANTS_DEF) + print(f"wrote {VARIANTS_DEF}") diff --git a/vllm/v1/spec_decode/gemma4.py b/vllm/v1/spec_decode/gemma4.py index 07cfaa2f4db2..f4116156ffbc 100644 --- a/vllm/v1/spec_decode/gemma4.py +++ b/vllm/v1/spec_decode/gemma4.py @@ -146,9 +146,10 @@ def _setup_centroids_cuda_graphs(self) -> None: ) def _create_draft_vllm_config(self) -> VllmConfig: - """Preserve the target's forced TRITON_ATTN backend for draft layers. + """Preserve the target's forced backend for draft layers. - Gemma4 forces TRITON_ATTN due to heterogeneous head dimensions + Gemma4 forces TRITON_ATTN (ROCM_ATTN on gfx1151) due to + heterogeneous head dimensions (head_dim=256 sliding, global_head_dim=512 full). The base class resets attention_config.backend to None for draft models, causing sliding layers to fall back to FLASH_ATTN which cannot handle