From b8754e0af930b6e38937d815d4cc80f64b4b2ac7 Mon Sep 17 00:00:00 2001 From: Robert Esclapez Garcia Date: Fri, 2 Oct 2026 11:34:40 -0700 Subject: [PATCH 1/2] [Benchmarks] attention: honour the dtype, sliding window, KV out of cache The benchmark built its model config with dtype="auto" and ignored the requested dtype, so it ran the stand-in model's native bfloat16 whatever was asked: everything labelled fp16 was bf16. --dtype (float16 or bfloat16) now sets Q/K/V and the output, and the bfloat16 -> float16 cast warning that then fires once per cell is filtered (that message only, so a backend's fallback warning stays visible). BenchmarkConfig.sliding_window: the impl and the KV cache spec get the window, and the block table holds only the blocks some query's window reaches (the rest name the null block), as vLLM's sliding-window manager keeps a windowed layer, instead of the whole sequence. benchmark_fn walks one KV cache per layer and divides by the layer count, so whether a cell reads DRAM or a resident last-level cache depends on num_layers * kv_bytes_per_layer. On parts with a large one (Strix Halo's 32 MiB MALL, never flushed between CUDA-graph replays) a small model at a short context stays resident: TRITON_ATTN decode cells under 32 MiB read 35 % fast (median, worst 64 %), those over it 0 %. --min-working-set-mb raises --num-layers until the working set (the window, for a windowed layer) reaches the given size; it only raises, so cells already past the cache do not move. The README documents it. AI assistance was used for this patch. Co-Authored-By: Claude Signed-off-by: Robert Esclapez Garcia --- benchmarks/attention_benchmarks/README.md | 35 +++++++ benchmarks/attention_benchmarks/benchmark.py | 34 +++++++ benchmarks/attention_benchmarks/common.py | 31 ++++++ benchmarks/attention_benchmarks/runner.py | 100 ++++++++++++++++--- 4 files changed, 187 insertions(+), 13 deletions(-) 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( From d7b04026a294006a3297eb3a333abbf5d232e4d8 Mon Sep 17 00:00:00 2001 From: Robert Esclapez Garcia Date: Fri, 2 Oct 2026 11:34:40 -0700 Subject: [PATCH 2/2] [ROCm][gfx1151] ROCM_ATTN: HIP decode attention on gfx1151 On gfx1151 (Strix Halo, RDNA3.5) ROCM_ATTN serves decode, speculative decode (M <= 8 query tokens per sequence) and the decodes of mixed batches with a HIP WMMA kernel, and everything else with TRITON_ATTN's kernels, on Triton's packed KV cache in the HND layout. It becomes the default attention backend there, ahead of TRITON_ATTN. On every other device ROCM_ATTN is unchanged. Kernel (csrc/rocm/rdna35_decode_attn.cu, csrc/rocm/rdna35_attn/): one workgroup per KV head, GQA-packed WMMA for Q@K and P@V in fp16 and bf16, D in {64, 128, 256, 512}, full and sliding-window attention, split-KV with the cross-workgroup merge done in-kernel, a dot-product mode for short sequences, V staged in LDS so the next tile's loads overlap, batched decode. Each configuration (heads, kv heads, D, M, window) is tuned and compiled into _rocm_C as a variant (variants.def, written by vllm/v1/attention/ops/rdna35_hip_decode.py from its tables; cmake/rdna35_attn.cmake); VLLM_RDNA35_ATTN_JIT builds them from source for kernel development. RocmAttentionBackend, on gfx1151 only, returns the packed shape and stride order of TRITON_ATTN, requires HND, adds head size 512, and builds RocmAttentionRdna35MetadataBuilder (Triton's, with the batch ordered decodes first) and RocmAttentionRdna35Impl. The impl subclasses TritonAttentionImpl without modifying it: its forward launches the kernel when the call is a uniform decode it can serve and hands anything else to Triton's forward (prefill, alibi, softcap, non-causal, quantized KV or output, shapes not built); in a mixed batch the decodes go to the kernel and the prefills to Triton on a slice of the metadata. The branches check the class, so ROCM_AITER_UNIFIED_ATTN, a subclass, keeps its own layout. What ROCM_ATTN does not declare (sinks, KV connectors, encoder-decoder) is served by TRITON_ATTN through the normal selection. Gemma 4's forced single backend is ROCM_ATTN on gfx1151. benchmarks/kernels/gfx1151_decode_attn/tools: the tuning, sweep, roofline and batch tools that produced variants.def and the measurements, on top of attention_benchmarks. Tests: tests/kernels/attention/test_rdna35_hip_decode.py. AI assistance was used for this patch. Co-Authored-By: Claude Signed-off-by: Robert Esclapez Garcia --- CMakeLists.txt | 14 + .../gfx1151_decode_attn/tools/batch.py | 149 ++ .../gfx1151_decode_attn/tools/check.py | 170 ++ .../gfx1151_decode_attn/tools/dump_asm.sh | 30 + .../gfx1151_decode_attn/tools/floor.py | 193 ++ .../gfx1151_decode_attn/tools/matrix.py | 300 +++ .../gfx1151_decode_attn/tools/roofline.py | 278 +++ .../gfx1151_decode_attn/tools/shapes.csv | 60 + .../gfx1151_decode_attn/tools/shapeset.py | 184 ++ .../gfx1151_decode_attn/tools/sweep.py | 276 +++ .../kernels/gfx1151_decode_attn/tools/tune.py | 364 ++++ cmake/rdna35_attn.cmake | 55 + csrc/rocm/ops.h | 9 + csrc/rocm/rdna35_attn/instance.h | 77 + csrc/rocm/rdna35_attn/launch.h | 35 + csrc/rocm/rdna35_attn/macros.h | 60 + csrc/rocm/rdna35_attn/prelude.h | 15 + csrc/rocm/rdna35_attn/registry.hip | 112 ++ csrc/rocm/rdna35_attn/variants.def | 1284 ++++++++++++ csrc/rocm/rdna35_attn/variants.h | 30 + csrc/rocm/rdna35_decode_attn.cu | 1784 +++++++++++++++++ csrc/rocm/torch_bindings.cpp | 11 + .../attention/test_rdna35_hip_decode.py | 634 ++++++ vllm/envs.py | 6 + vllm/model_executor/models/config.py | 16 +- vllm/platforms/rocm.py | 8 +- vllm/v1/attention/backends/rocm_attn.py | 424 +++- vllm/v1/attention/ops/rdna35_hip_decode.py | 899 +++++++++ vllm/v1/spec_decode/gemma4.py | 5 +- 29 files changed, 7471 insertions(+), 11 deletions(-) create mode 100644 benchmarks/kernels/gfx1151_decode_attn/tools/batch.py create mode 100644 benchmarks/kernels/gfx1151_decode_attn/tools/check.py create mode 100755 benchmarks/kernels/gfx1151_decode_attn/tools/dump_asm.sh create mode 100644 benchmarks/kernels/gfx1151_decode_attn/tools/floor.py create mode 100644 benchmarks/kernels/gfx1151_decode_attn/tools/matrix.py create mode 100644 benchmarks/kernels/gfx1151_decode_attn/tools/roofline.py create mode 100644 benchmarks/kernels/gfx1151_decode_attn/tools/shapes.csv create mode 100644 benchmarks/kernels/gfx1151_decode_attn/tools/shapeset.py create mode 100644 benchmarks/kernels/gfx1151_decode_attn/tools/sweep.py create mode 100644 benchmarks/kernels/gfx1151_decode_attn/tools/tune.py create mode 100644 cmake/rdna35_attn.cmake create mode 100644 csrc/rocm/rdna35_attn/instance.h create mode 100644 csrc/rocm/rdna35_attn/launch.h create mode 100644 csrc/rocm/rdna35_attn/macros.h create mode 100644 csrc/rocm/rdna35_attn/prelude.h create mode 100644 csrc/rocm/rdna35_attn/registry.hip create mode 100644 csrc/rocm/rdna35_attn/variants.def create mode 100644 csrc/rocm/rdna35_attn/variants.h create mode 100644 csrc/rocm/rdna35_decode_attn.cu create mode 100644 tests/kernels/attention/test_rdna35_hip_decode.py create mode 100644 vllm/v1/attention/ops/rdna35_hip_decode.py 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/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