diff --git a/benchmarks/single_node/agentic/apply_k3_container_patches.sh b/benchmarks/single_node/agentic/apply_k3_container_patches.sh new file mode 100755 index 000000000..cdfd1d33f --- /dev/null +++ b/benchmarks/single_node/agentic/apply_k3_container_patches.sh @@ -0,0 +1,186 @@ +#!/usr/bin/env bash +# ============================================================================= +# apply_k3_container_patches.sh +# +# Turn the pinned base ROCm vLLM nightly image into the EXACT container the +# Kimi-K3 FP4 MI355X DSpark agentic benchmark runs in. Run this INSIDE a +# container started from: +# +# vllm/vllm-openai-rocm:nightly-cb8104839c141609d99f1254459ef3a4f1bd4263 +# +# (method borrowed from InferenceX #2508: fetch/build the deltas the base image +# lacks, apply them into the installed dist-packages + a node-local aiter, then +# verify by anchor grep). Self-contained and idempotent — no bind mounts, no +# host paths. Everything it needs ships in ./k3_patches/. +# +# docker run -d --name k3-dspark-benchmark \ +# --ipc=host --network=host --shm-size=137438953472 \ +# --device=/dev/kfd --device=/dev/dri --group-add video --group-add render \ +# --security-opt seccomp=unconfined --security-opt label=disable \ +# --cap-add=SYS_PTRACE -e GPU_ARCHS=gfx950 \ +# --entrypoint sleep \ +# vllm/vllm-openai-rocm:nightly-cb8104839c141609d99f1254459ef3a4f1bd4263 infinity +# docker cp benchmarks/single_node/agentic k3-dspark-benchmark:/opt/k3-recipe +# docker exec k3-dspark-benchmark bash /opt/k3-recipe/apply_k3_container_patches.sh +# +# Result matches `setup_benchmark.sh setup-dspark` from the source tree exactly: +# - aiter rebuilt at pin 55dbc4f47 (#4579 d3ddaabf9 + #4575 22beb1caa) +# - bundled tuned K3 GEMM CSV installed + merged -> merged_bf16_tuned_gemm.csv +# - triton 3.7.0 + tabulate (nightly ships 3.6.0) +# - 5 vLLM ASM base patches (decode #50578, fp8 prefill PR-A, PS metadata16, +# skip-k3-fp8-ps, wvSplitK #50618) +# - DSpark fp8-asm enablement layer (apply_dspark_fp8asm.sh) +# - FlyDSL->torch decode-GEMM reroute (patch_flydsl_decode_to_torch.sh) +# +# Overridable knobs (env): +# AITER_PIN aiter commit to build (default 55dbc4f47...) +# AITER_SRC pre-cloned aiter checkout to stage instead of git clone +# LOCAL_AITER install location (default /opt/aiter-local; the serve scripts +# reference this path for the merged GEMM CSV) +# SKIP_TRITON=1 skip the triton 3.7.0 upgrade (if the image already has it) +# ============================================================================= +set -euo pipefail + +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +PATCHES="$SCRIPT_DIR/k3_patches" + +AITER_PIN="${AITER_PIN:-55dbc4f475da26c23cdaf73ce6ed38342a2d7f83}" +AITER_4579="${AITER_4579:-d3ddaabf9}" # int-32 K offset fix +AITER_4575="${AITER_4575:-22beb1caa}" # int-32 V offset fix +LOCAL_AITER="${LOCAL_AITER:-/opt/aiter-local}" +AITER_REPO="${AITER_REPO:-https://github.com/ROCm/aiter.git}" +DIST="${DIST:-/usr/local/lib/python3.12/dist-packages}" +MLA="$DIST/vllm/v1/attention/backends/mla/rocm_aiter_mla.py" +UTILS="$DIST/vllm/model_executor/layers/utils.py" +TUNED_CSV="$PATCHES/kimik3_bf16_tuned_gemm.csv" + +say() { echo; echo "############### $* ###############"; } + +[ -d "$PATCHES" ] || { echo "!! missing $PATCHES bundle next to this script" >&2; exit 1; } +[ -f "$TUNED_CSV" ] || { echo "!! missing tuned GEMM CSV $TUNED_CSV" >&2; exit 1; } +[ -f "$MLA" ] || { echo "!! $MLA not found — is this the pinned base image?" >&2; exit 1; } + +# --------------------------------------------------------------------------- +say "1/6 build node-local aiter @ $AITER_PIN (#4579 + #4575 K/V int-32 offsets)" +# Stage from AITER_SRC if provided, else clone. JIT-compiles on demand against +# the container torch + system triton (PREBUILD_KERNELS=0, AITER_USE_SYSTEM_TRITON=1). +export PREBUILD_KERNELS=0 AITER_USE_SYSTEM_TRITON=1 +if [ -d "$LOCAL_AITER/.git" ]; then + echo " reusing existing $LOCAL_AITER checkout" +elif [ -n "${AITER_SRC:-}" ] && [ -d "$AITER_SRC/.git" ]; then + echo " staging aiter from $AITER_SRC" + rm -rf "$LOCAL_AITER"; cp -a "$AITER_SRC" "$LOCAL_AITER" +else + echo " cloning $AITER_REPO -> $LOCAL_AITER" + rm -rf "$LOCAL_AITER"; git clone "$AITER_REPO" "$LOCAL_AITER" +fi +git config --global --add safe.directory "$LOCAL_AITER" +git -C "$LOCAL_AITER" fetch --tags --depth=1 origin "$AITER_PIN" 2>/dev/null \ + || git -C "$LOCAL_AITER" fetch --tags origin 2>/dev/null || true +git -C "$LOCAL_AITER" reset --hard "$AITER_PIN" +git -C "$LOCAL_AITER" submodule update --init 3rdparty/composable_kernel +[ -d "$LOCAL_AITER/3rdparty/composable_kernel/include" ] \ + || { echo "!! composable_kernel submodule not populated" >&2; exit 1; } +git -C "$LOCAL_AITER" merge-base --is-ancestor "$AITER_4579" HEAD \ + || { echo "!! aiter missing #4579 ($AITER_4579) after checkout $AITER_PIN" >&2; exit 1; } +git -C "$LOCAL_AITER" merge-base --is-ancestor "$AITER_4575" HEAD \ + || { echo "!! aiter missing #4575 ($AITER_4575) after checkout $AITER_PIN" >&2; exit 1; } +echo " aiter HEAD: $(git -C "$LOCAL_AITER" log --oneline -1)" +# Never inherit stale JIT batons/build from a prior tree (blocks rank 0 in RCCL). +rm -rf "$LOCAL_AITER/aiter/jit/build" +find "$LOCAL_AITER/aiter/jit" -maxdepth 1 -name "module_*.so" -delete 2>/dev/null || true +pip uninstall -y aiter amd-aiter >/dev/null 2>&1 || true +( cd "$LOCAL_AITER" && pip install -e . --no-build-isolation --no-deps ) +rm -rf /root/aiter; ln -s "$LOCAL_AITER" /root/aiter +python3 -c "import aiter; assert '/opt/aiter-local' in aiter.__file__ or '/root/aiter' in aiter.__file__, aiter.__file__; print(' aiter:', aiter.__file__)" + +# --------------------------------------------------------------------------- +say "2/6 install + merge tuned K3 BF16 GEMM CSV" +CONFIGS="$LOCAL_AITER/aiter/configs" +mkdir -p "$CONFIGS/model_configs" +cp "$TUNED_CSV" "$CONFIGS/model_configs/kimik3_bf16_tuned_gemm.csv" +cmp -s "$TUNED_CSV" "$CONFIGS/model_configs/kimik3_bf16_tuned_gemm.csv" \ + || { echo "!! tuned GEMM CSV copy verification failed" >&2; exit 1; } +python3 - "$CONFIGS" <<'PY' +import os, shutil, sys +from pathlib import Path +from aiter.jit.core import AITER_CONFIGS +configs = Path(sys.argv[1]) +sources = [configs / "bf16_tuned_gemm.csv"] +sources.extend( + p for p in sorted((configs / "model_configs").glob("*bf16_tuned_gemm*.csv")) + if "untuned" not in p.name +) +source_list = os.pathsep.join(str(p) for p in sources if p.is_file()) +if not source_list: + raise SystemExit("ERROR: no BF16 tuned GEMM CSVs found") +try: + merged = AITER_CONFIGS.update_config_files(source_list, "bf16_tuned_gemm") +except RuntimeError as exc: + # aiter raises once after resolving cross-file dupes in place; second pass is clean. + if "Auto-resolved by keeping best performing" not in str(exc): + raise + merged = AITER_CONFIGS.update_config_files(source_list, "bf16_tuned_gemm") +dest = configs / "merged_bf16_tuned_gemm.csv" +shutil.copyfile(merged, dest) +print(f" merged BF16 GEMM CSV -> {dest}") +PY + +# --------------------------------------------------------------------------- +if [ "${SKIP_TRITON:-0}" = "1" ]; then + say "3/6 triton upgrade SKIPPED (SKIP_TRITON=1)" +else + say "3/6 triton 3.7.0 + tabulate (nightly ships 3.6.0)" + pip install -q --extra-index-url https://pypi.amd.com/triton/release/rocm-7.2.0/simple/ \ + triton==3.7.0 tabulate +fi +python3 -c "import triton; print(' triton', triton.__version__)" + +# --------------------------------------------------------------------------- +say "4/6 vLLM ASM base patches (decode #50578, fp8 prefill PR-A, PS16, skip-k3-fp8-ps, wvSplitK #50618)" +if grep -q "PATCH(fp8-asm)" "$MLA" && grep -q "PATCH(fp8-prefill-pad)" "$MLA" \ + && grep -q "num_head_k = max(16, self.num_heads)" "$MLA" \ + && grep -q "PATCH(skip-k3-fp8-ps)" "$MLA" \ + && grep -q "PATCH(vLLM #50618)" "$UTILS"; then + echo " all 5 ASM patches already present" +else + for p in patch_fp8asm.py patch_fp8_prefill.py patch_ps_metadata16.py patch_skip_k3_fp8_ps.py patch_wvsplitk.py; do + echo " applying $p ..." + python3 "$PATCHES/$p" + done +fi + +# --------------------------------------------------------------------------- +say "5/6 DSpark fp8-asm enablement layer" +bash "$PATCHES/apply_dspark_fp8asm.sh" + +# --------------------------------------------------------------------------- +say "6/6 FlyDSL -> torch decode-GEMM reroute (cudagraph-capturable dense GEMMs)" +CSV="$CONFIGS/merged_bf16_tuned_gemm.csv" bash "$PATCHES/patch_flydsl_decode_to_torch.sh" + +# --------------------------------------------------------------------------- +say "VERIFY (matches setup_benchmark.sh verify-dspark-patches)" +AITER_MLA="$LOCAL_AITER/aiter/mla.py" +KDA="$DIST/vllm/models/kimi_k3/amd/ops/third_party/kda/fused_recurrent.py" +ok=1 +chk() { local n; n=$(grep -c "$2" "$1" 2>/dev/null || echo 0); \ + if [ "$n" -ge "$3" ]; then echo " OK $4 ($n)"; else echo " FAIL $4 ($n < $3)"; ok=0; fi; } +chk "$MLA" "PATCH(fp8-asm)" 1 "decode pad-to-16 (#50578)" +chk "$MLA" "PATCH(fp8-prefill-pad)" 1 "fp8 prefill pad (PR-A)" +chk "$MLA" "num_head_k = max(16, self.num_heads)" 1 "PS metadata16 (PR-A)" +chk "$MLA" "PATCH(skip-k3-fp8-ps)" 1 "skip K3 fp8 PS" +chk "$UTILS" "PATCH(vLLM #50618)" 1 "wvSplitK (#50618)" +chk "$MLA" "_mtp_decode_qlen" 1 "DSpark _mtp_decode_qlen" +chk "$MLA" 'method == "dspark"' 1 "dspark verify qlen branch" +chk "$MLA" "uses_asm_decode" 2 "persistent-metadata gate" +chk "$AITER_MLA" "80: 64" 1 "aiter get_block_n_fp8 key 80" +chk "$AITER_MLA" "get_block_n_fp8.get(" 1 "aiter get_block_n_fp8.get()" +chk "$KDA" "stride_indices_seq" 5 "KDA PR#27 stride fix" +python3 -c "import vllm.v1.attention.backends.mla.rocm_aiter_mla; print(' IMPORT_OK')" +[ "$ok" = 1 ] || { echo; echo "!! one or more anchors missing — see FAIL lines above" >&2; exit 1; } + +echo +echo "DONE — container matches k3-dspark-benchmark. Serve with:" +echo " export VLLM_ROCM_AITER_MLA_ASM_PADDING=asm" +echo " NUM_SPEC=2 PORT=8890 GPU_MEM=0.95 MAX_NUM_SEQS=64 MNBT=16384 \\" +echo " SYNTHETIC_ACCEPT_LEN=2.51 bash " diff --git a/benchmarks/single_node/agentic/k3_patches/apply_dspark_fp8asm.sh b/benchmarks/single_node/agentic/k3_patches/apply_dspark_fp8asm.sh new file mode 100755 index 000000000..ecc295b58 --- /dev/null +++ b/benchmarks/single_node/agentic/k3_patches/apply_dspark_fp8asm.sh @@ -0,0 +1,380 @@ +#!/usr/bin/env bash +# ============================================================================= +# apply_dspark_fp8asm.sh (portable — no bind mounts / host paths) +# DSpark + native fp8-asm KV enablement layer for Kimi-K3 on MI355X. Applied by +# apply_k3_container_patches.sh AFTER the 5 base vLLM ASM patches and the aiter +# rebuild. Sits on the SHIPPED recipe rocm_aiter_mla.py (main + #51011 + #51040 +# + #51606), which the base ASM patches install. +# +# Idempotent: every step guards on its anchor / target text and takes a .bak. +# The sibling patch_*.py scripts are resolved relative to this file's directory. +# ============================================================================= +set -uo pipefail +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +D=/usr/local/lib/python3.12/dist-packages +R="$D/vllm/v1/attention/backends/mla/rocm_aiter_mla.py" +# RESEED_SHIPPED_MLA is off by default; only set SHIPPED_MLA_REF if you reseed. +SHIPPED_MLA_REF="${SHIPPED_MLA_REF:-}" +if [ "${RESEED_SHIPPED_MLA:-0}" = "1" ] && [ -n "$SHIPPED_MLA_REF" ] && [ -f "$SHIPPED_MLA_REF" ]; then + cp -a "$R" "$R.pre_reseed.bak" 2>/dev/null || true + cp "$SHIPPED_MLA_REF" "$R" + echo "reseeded rocm_aiter_mla.py from $SHIPPED_MLA_REF" +fi +if [ -f /opt/aiter-local/aiter/mla.py ]; then + AITER=/opt/aiter-local/aiter/mla.py +elif [ -f /root/aiter/aiter/mla.py ]; then + AITER=/root/aiter/aiter/mla.py +else + echo "!! aiter mla.py not found under /opt/aiter-local or /root/aiter — run setup first" + exit 1 +fi +KDA="$D/vllm/models/kimi_k3/amd/ops/third_party/kda/fused_recurrent.py" +export R AITER KDA SHIPPED_MLA_REF +say(){ echo; echo "=================== $* ==================="; } + +say "0/7 InferenceX fp8-asm base patches (decode, prefill, PS skip, wvSplitK)" +cd "$SCRIPT_DIR" +need_asm=0 +grep -q 'PATCH(fp8-asm)' "$R" || need_asm=1 +grep -q 'PATCH(skip-k3-fp8-ps)' "$R" || need_asm=1 +grep -q 'PATCH(vLLM #50618)' "$D/vllm/model_executor/layers/utils.py" || need_asm=1 +if [ "$need_asm" = 1 ]; then + for p in patch_fp8asm.py patch_fp8_prefill.py patch_ps_metadata16.py patch_skip_k3_fp8_ps.py patch_wvsplitk.py; do + echo " applying $p ..." + python3 "$SCRIPT_DIR/$p" + done +else + echo " all 5 InferenceX ASM patches already present" +fi + +say "1/7 sanity: shipped recipe backend present (_mtp_decode_qlen)" +grep -q _mtp_decode_qlen "$R" || { echo "!! recipe rocm_aiter_mla.py not installed — apply the base recipe first"; exit 1; } + +say "2/7 force DSpark draft causal (dflash_config.causal=true)" +DCACHE=/dev/shm/hf-cache/models--Inferact--Kimi-K3-DSpark/snapshots +CFG="$(ls -d "$DCACHE"/*/ 2>/dev/null | head -1)config.json" +if [ -f "$CFG" ]; then + cp -n "$CFG" "$CFG.orig.bak" + python3 - "$CFG" <<'PY' +import json, sys +f = sys.argv[1] +c = json.load(open(f)) +d = c.setdefault("dflash_config", {}) +if d.get("causal") is True: + print(" draft already forced causal") +else: + d["causal"] = True + json.dump(c, open(f, "w"), indent=2) + print(" forced causal:", f) +PY +else + echo " !! draft config not found at $DCACHE — stage the draft or set it manually" +fi + +say "3/7 aiter get_block_n_fp8 — add DSpark verify-width key (16*5=80)" +python3 - <<'PY' +import re, py_compile, shutil, os +F = os.environ.get("AITER", "/root/aiter/aiter/mla.py") +s = open(F).read() +shutil.copy2(F, F + ".pre_dspark.bak") +if "80: 64" not in s: + s = re.sub(r"(get_block_n_fp8\s*=\s*\{)", r"\g<1>\n 80: 64, 96: 64, 112: 64,", s, count=1) +s = s.replace( + "min_block_n = get_block_n_fp8[int(nhead * max_seqlen_q)]", + "min_block_n = get_block_n_fp8.get(int(nhead * max_seqlen_q), 64)", +) +open(F, "w").write(s); py_compile.compile(F, doraise=True) +print(" 80-key:", "80: 64" in s, " get():", "get_block_n_fp8.get(" in s) +PY + +say "4/7 recipe backend — DSpark verify width (_mtp_decode_qlen = 2*num_spec+1)" +python3 - <<'PY' +import py_compile, shutil, os +F = os.environ["R"] +s = open(F).read() +shutil.copy2(F, F + ".pre_dsparkqlen.bak") +if 'speculative_config.method == "dspark"' in s: + print(" dspark qlen branch already present") +else: + anchor = " else:\n self._mtp_decode_qlen = 1" + branch = ( + " elif (\n" + " speculative_config is not None\n" + ' and speculative_config.method == "dspark"\n' + " and speculative_config.num_speculative_tokens is not None\n" + " ):\n" + " # DSpark drafts a semi-autoregressive block; the target verifies\n" + " # 1 anchor + 2*num_spec positions.\n" + " self._mtp_decode_qlen = 2 * int(speculative_config.num_speculative_tokens) + 1\n" + ) + if s.count(anchor) == 1: + s = s.replace(anchor, branch + anchor, 1) + open(F, "w").write(s); py_compile.compile(F, doraise=True) + print(" dspark qlen branch added") + else: + print(f" !! anchor count {s.count(anchor)} — add the dspark branch manually") +PY + +say "4b/7 recipe backend — _aiter_mla_small_head_mode (shipped helper + os.environ fallback)" +python3 - <<'PY' +import os, py_compile, re, shutil + +F = os.environ["R"] +# Optional shipped-recipe reference for restoring the Gluon helpers; empty by +# default (the base ASM patches already install a recipe that carries them). +REF = os.environ.get("SHIPPED_MLA_REF", "") or "" +s = open(F).read() + +# Some patch orderings drop Gluon support helpers; restore from the shipped file. +if "_gluon_mla_decode_supported" not in s and os.path.isfile(REF): + ref = open(REF).read() + m = re.search( + r"@functools\.lru_cache\(maxsize=1\)\ndef _gluon_mla_decode_supported\(\)[\s\S]*?return on_gfx950\(\)\n", + ref, + ) + if m: + block = m.group(0) + if "def _aiter_mla_small_head_mode" in s: + s = s.replace("\ndef _aiter_mla_small_head_mode", "\n" + block + "\n\ndef _aiter_mla_small_head_mode", 1) + else: + s = s.replace("\nclass AiterMLABackend", "\n" + block + "\n\ndef _aiter_mla_small_head_mode() -> str:\n import os\n\n return (os.environ.get(\"VLLM_ROCM_AITER_MLA_ASM_PADDING\") or \"auto\").lower()\n\n\nclass AiterMLABackend", 1) + open(F, "w").write(s) + print(" restored _gluon_mla_decode_supported from shipped ref") + s = open(F).read() + +# Nightly cb8104839 may lack envs.VLLM_ROCM_AITER_MLA_ASM_PADDING — read the env var directly. +if "envs.VLLM_ROCM_AITER_MLA_ASM_PADDING" in s: + s = s.replace( + " import vllm.envs as envs\n\n mode = (envs.VLLM_ROCM_AITER_MLA_ASM_PADDING or \"auto\").lower()", + " import os\n\n mode = (os.environ.get(\"VLLM_ROCM_AITER_MLA_ASM_PADDING\") or \"auto\").lower()", + ) + open(F, "w").write(s) + py_compile.compile(F, doraise=True) + print(" full small_head_mode helper: envs -> os.environ") +elif 'os.environ.get("VLLM_ROCM_AITER_MLA_ASM_PADDING")' in s: + if "Small-head (<16) MLA decode" in s: + print(" full small_head_mode helper with os.environ already present") + elif os.path.isfile(REF): + ref = open(REF).read() + m = re.search( + r"def _aiter_mla_small_head_mode\(\)[\s\S]*?return mode\n", + ref, + ) + if m: + block = m.group(0).replace( + " import vllm.envs as envs\n\n mode = (envs.VLLM_ROCM_AITER_MLA_ASM_PADDING or \"auto\").lower()", + " import os\n\n mode = (os.environ.get(\"VLLM_ROCM_AITER_MLA_ASM_PADDING\") or \"auto\").lower()", + ) + s = re.sub( + r"\ndef _aiter_mla_small_head_mode\(\)[\s\S]*?(?=\n\nclass AiterMLABackend)", + "\n" + block, + s, + count=1, + ) + open(F, "w").write(s) + py_compile.compile(F, doraise=True) + print(" upgraded stub -> full shipped small_head_mode helper") + else: + print(" stub small_head_mode helper present (nightly gap)") + else: + print(" stub small_head_mode helper present (nightly gap)") +elif "def _aiter_mla_small_head_mode" not in s: + helper = ( + "\n\n" + "def _aiter_mla_small_head_mode() -> str:\n" + ' """Return VLLM_ROCM_AITER_MLA_ASM_PADDING mode (auto/gluon/asm)."""\n' + " import os\n\n" + ' return (os.environ.get("VLLM_ROCM_AITER_MLA_ASM_PADDING") or "auto").lower()\n\n' + ) + anchor = "class AiterMLABackend(MLACommonBackend):" + if anchor not in s: + print(" !! AiterMLABackend anchor not found — add small_head_mode manually") + else: + s = s.replace(anchor, helper + anchor, 1) + open(F, "w").write(s) + py_compile.compile(F, doraise=True) + print(" stub small_head_mode helper added") +else: + print(" !! _aiter_mla_small_head_mode present but unreadable — inspect manually") +PY + +say "4c/7 recipe backend — skip gluon verify flatten when ASM_PADDING=asm" +python3 - <<'PY' +import py_compile, os +F = os.environ["R"] +s = open(F).read() +old = ( + " if (\n" + " self.num_heads < AiterMLAHelper._AITER_MIN_MLA_HEADS\n" + " and int(decode.max_qo_len) > 1\n" + " ):\n" +) +new = ( + " if (\n" + " self.num_heads < AiterMLAHelper._AITER_MIN_MLA_HEADS\n" + " and int(decode.max_qo_len) > 1\n" + ' and _aiter_mla_small_head_mode() != "asm"\n' + " ):\n" +) +if 'and _aiter_mla_small_head_mode() != "asm"' in s and "multi-token verify (DSpark)" in s: + print(" gluon verify flatten already gated") +elif old in s: + s = s.replace(old, new, 1) + open(F, "w").write(s) + py_compile.compile(F, doraise=True) + print(" gluon verify flatten gated on asm mode") +else: + print(" !! gluon verify flatten anchor not found — inspect decode path manually") +PY + +say "5/7 recipe backend — broaden persistent-metadata gate" +python3 - <<'PY' +import py_compile, shutil, os +F = os.environ["R"] +s = open(F).read() +shutil.copy2(F, F + ".pre_gate.bak") +old = ( + " use_persistent_metadata = (\n" + " self.num_heads >= AiterMLAHelper._AITER_MIN_MLA_HEADS\n" + " and max_qo_len >= 1\n" + " and max_qo_len <= self._mtp_decode_qlen\n" + " )" +) +new = ( + " uses_asm_decode = (\n" + " self.num_heads >= AiterMLAHelper._AITER_MIN_MLA_HEADS\n" + ' or _aiter_mla_small_head_mode() == "asm"\n' + " or max_qo_len > 1\n" + " )\n" + " use_persistent_metadata = (\n" + " uses_asm_decode\n" + " and max_qo_len >= 1\n" + " and max_qo_len <= self._mtp_decode_qlen\n" + " )" +) +if "uses_asm_decode" in s: + print(" gate already broadened") +elif old in s: + s = s.replace(old, new, 1); open(F, "w").write(s); py_compile.compile(F, doraise=True) + print(" gate broadened") +else: + print(" !! gate anchor not found — inspect use_persistent_metadata manually") +PY + +say "6/7 KDA stride fix (Fangzhou-Ai/vllm PR #27; correctness)" +python3 - <<'PY' +import py_compile, re, shutil, os +F = os.environ["KDA"] +if not os.path.exists(F): + print(" !! KDA file not found — skip"); raise SystemExit(0) +s = open(F).read() +shutil.copy2(F, F + ".pre_pr27.bak") +changed = False + +def sub_once(pattern, repl, text, label): + global changed + text2, n = re.subn(pattern, repl, text, count=1, flags=re.DOTALL) + if n: + changed = True + print(f" + {label}") + return text2 + +# fused_recurrent_kda_fwd_kernel — sibling already takes stride_indices_seq upstream. +s = sub_once( + r"(def fused_recurrent_kda_fwd_kernel\([\s\S]*? stride_state_token: tl\.constexpr,\n)" + r"(?! stride_indices_seq: tl\.constexpr,\n)" + r"( IS_SPEC_DECODING: tl\.constexpr,)", + r"\1 stride_indices_seq: tl.constexpr,\n\2", + s, + "fwd kernel stride_indices_seq param", +) + +# fused_recurrent_kda_packed_decode_kernel — the cudagraph/eager path that failed at 10/15. +s = sub_once( + r"(def fused_recurrent_kda_packed_decode_kernel\([\s\S]*? stride_state_token: tl\.constexpr,\n)" + r"(?! stride_indices_seq: tl\.constexpr,\n)" + r"( H: tl\.constexpr,)", + r"\1 stride_indices_seq: tl.constexpr,\n\2", + s, + "packed_decode kernel stride_indices_seq param", +) + +s2 = s.replace( + "state_idx = tl.load(state_indices + i_n).to(tl.int64)", + "state_idx = tl.load(state_indices + i_n * stride_indices_seq).to(tl.int64)", +) +if s2 != s: + changed = True + print(" + packed_decode state_idx stride load") +s = s2 + +s2 = s.replace( + " if state_indices.ndim != 1 or state_indices.stride(0) != 1:\n" + " state_indices = state_indices.reshape(-1).contiguous()", + " if state_indices.ndim != 1:\n" + ' raise ValueError("`state_indices` must be one-dimensional.")', +) +if s2 != s: + changed = True + print(" + drop unit-stride requirement on state_indices") +s = s2 + +s2 = s.replace( + " if state_indices.ndim != 1 or state_indices.stride(0) != 1:\n" + ' raise ValueError("`state_indices` must be contiguous and one-dimensional.")', + " if state_indices.ndim != 1:\n" + ' raise ValueError("`state_indices` must be one-dimensional.")', +) +if s2 != s: + changed = True + print(" + drop contiguous check on state_indices") +s = s2 + +if not re.search( + r"def fused_recurrent_kda_fwd\([\s\S]*?stride_indices_seq=ssm_state_indices\.stride\(0\)", + s, +): + s2 = s.replace( + " stride_state_token=initial_state.stride(0),\n" + " IS_SPEC_DECODING=", + " stride_state_token=initial_state.stride(0),\n" + " stride_indices_seq=ssm_state_indices.stride(0),\n" + " IS_SPEC_DECODING=", + 1, + ) + if s2 != s: + changed = True + print(" + fwd kernel launch stride_indices_seq") + s = s2 + +if not re.search( + r"fused_recurrent_kda_packed_decode_kernel\[grid\][\s\S]*?" + r"stride_indices_seq=state_indices\.stride\(0\)", + s, +): + s = sub_once( + r"(fused_recurrent_kda_packed_decode_kernel\[grid\]\([\s\S]*?" + r" stride_state_token=initial_state\.stride\(0\),\n)" + r"( H=H,)", + r"\1 stride_indices_seq=state_indices.stride(0),\n\2", + s, + "packed_decode kernel launch stride_indices_seq", + ) + +if changed: + open(F, "w").write(s) +py_compile.compile(F, doraise=True) +print(" KDA PR#27 applied" if changed else " KDA PR#27 already present") +PY + +say "7/7 verify" +echo "aiter 80-key = $(grep -c '80: 64' "$AITER") (expect >=1)" +echo "aiter get() = $(grep -c 'get_block_n_fp8.get(' "$AITER") (expect 1)" +echo "dspark qlen = $(grep -c 'method == \"dspark\"' "$R") (expect >=1)" +echo "asm gate = $(grep -c 'uses_asm_decode' "$R") (expect >=2)" +echo "kda stride = $(grep -c 'stride_indices_seq' "$KDA") (expect >=5)" +python -c "import vllm.v1.attention.backends.mla.rocm_aiter_mla; print('IMPORT_OK')" +echo +echo "DONE. Serve with:" +echo " export VLLM_ROCM_AITER_MLA_ASM_PADDING=asm" +echo " NUM_SPEC=2 PORT=8890 GPU_MEM=0.88 MAX_NUM_SEQS=16 bash _serve_k3_bench_spec.sh" diff --git a/benchmarks/single_node/agentic/k3_patches/kimik3_bf16_tuned_gemm.csv b/benchmarks/single_node/agentic/k3_patches/kimik3_bf16_tuned_gemm.csv new file mode 100644 index 000000000..03c7666a8 --- /dev/null +++ b/benchmarks/single_node/agentic/k3_patches/kimik3_bf16_tuned_gemm.csv @@ -0,0 +1,501 @@ +gfx,cu_num,M,N,K,bias,dtype,outdtype,scaleAB,bpreshuffle,libtype,solidx,splitK,us,kernelName,err_ratio,tflops,bw +gfx950,256,1,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439841,0,4.7958,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x16x512_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_1_MO40_NTn1_NTA6_NTB1_NTC3_NTD0_NTM0_NEPBS14_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_4,0.0,0.66,657.0 +gfx950,256,1,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4023,8,5.7227,flydsl_gemm7_abf16_wbf16_bf16_t16x64x64_split_k8_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0234,1.47,1467.64 +gfx950,256,2,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2237,4,4.7475,flydsl_gemm3_abf16_wbf16_bf16_t16x64x128_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0073,1.33,664.76 +gfx950,256,2,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3659,8,5.7806,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k8_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0239,2.9,1454.71 +gfx950,256,4,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439633,0,4.8546,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x16x512_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_1_MO40_NTn1_NTA7_NTB2_NTC2_NTD0_NTM0_NEPBS4_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS0_SPO1_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_4,0.0,2.59,652.21 +gfx950,256,4,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4023,8,5.9363,flydsl_gemm7_abf16_wbf16_bf16_t16x64x64_split_k8_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.019,5.65,1420.0 +gfx950,256,8,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440906,0,4.9667,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x16x512_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_1_MO40_NTn1_NTA4_NTB0_NTC2_NTD1_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS0_SPO1_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_4,0.0,5.07,641.61 +gfx950,256,8,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4271,8,6.2762,flydsl_gemm8_abf16_wbf16_bf16_t16x64x64_split_k8_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0203,10.69,1349.63 +gfx950,256,16,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4993,4,4.6455,flydsl_gemm5_abf16_wbf16_bf16_t16x64x64_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0073,10.83,694.79 +gfx950,256,16,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3659,8,6.2278,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k8_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0188,21.55,1373.27 +gfx950,256,24,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5801,4,5.1879,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0079,14.55,630.04 +gfx950,256,32,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5801,4,5.2517,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0077,19.17,630.19 +gfx950,256,40,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3793,3,5.1982,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k3_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0059,24.21,644.56 +gfx950,256,40,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2463,4,7.1374,flydsl_gemm4_abf16_wbf16_bf16_t16x64x256_split_k4_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0188,47.01,1232.69 +gfx950,256,48,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,6423,3,5.2304,flydsl_gemm7_abf16_wbf16_bf16_t16x64x64_split_k3_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0055,28.87,648.42 +gfx950,256,56,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441420,0,5.4513,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x16x512_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_1_MO40_NTn1_NTA0_NTB0_NTC2_NTD1_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS0_SPO1_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_4,0.0,32.32,629.66 +gfx950,256,64,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440289,0,5.6641,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x32x512_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB2048_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT1_2_MO40_NTn1_NTA5_NTB0_NTC3_NTD2_NTM0_NEPBS14_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB2_WSGRA0_WSGRB0_WS64_WG16_4_4,0.0,35.54,613.23 +gfx950,256,96,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,6741,2,6.1574,flydsl_gemm5_abf16_wbf16_bf16_t16x64x128_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0025,49.05,590.71 +gfx950,256,96,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4291,4,8.9554,flydsl_gemm5_abf16_wbf16_bf16_t32x64x128_split_k4_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0117,89.92,1046.48 +gfx950,256,128,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440904,0,6.4316,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x32x256_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_2_MO40_NTn1_NTA4_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_4,0.0,62.61,591.0 +gfx950,256,256,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5892,1,6.9283,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0031,116.23,643.22 +gfx950,256,512,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,6282,1,8.1995,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,196.43,703.36 +gfx950,256,1024,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,6734,1,9.6467,flydsl_gemm4_abf16_wbf16_bf16_t64x64x128_split_k1_block_m_warp2_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,333.92,869.58 +gfx950,256,1024,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3156,1,18.7876,flydsl_gemm3_abf16_wbf16_bf16_t64x64x128_split_k1_block_m_warp2_block_n_warp2_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.005,457.21,1004.62 +gfx950,256,2048,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5724,1,13.2811,flydsl_gemm4_abf16_wbf16_bf16_t128x64x64_split_k1_block_m_warp4_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,485.08,1026.38 +gfx950,256,4096,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440595,0,18.5682,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x128x128_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_4_MO40_NTn1_NTA0_NTB6_NTC0_NTD6_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW4_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,693.92,1298.85 +gfx950,256,8192,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440366,0,26.6355,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_8_MO40_NTn1_NTA3_NTB7_NTC6_NTD6_NTM0_NEPBS12_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA4_VWB8_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,967.5,1692.81 +gfx950,256,16384,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440529,0,41.6977,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA3_NTB0_NTC6_NTD5_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1236.03,2087.21 +gfx950,256,16384,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439464,0,104.1127,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA3_NTB2_NTC4_NTD4_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1320.1,1692.02 +gfx950,256,24576,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440191,0,71.0111,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x192x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT16_3_MO40_NTn1_NTA2_NTB1_NTC7_NTD6_NTM0_NEPBS12_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM4_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,1088.69,1816.26 +gfx950,256,32768,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441024,0,87.1549,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA0_NTB4_NTC0_NTD4_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1182.71,1961.08 +gfx950,256,49152,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441458,0,128.2288,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA3_NTB4_NTC4_NTD5_NTM0_NEPBS14_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO4_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1205.8,1987.1 +gfx950,256,65536,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439465,0,167.4537,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA0_NTB4_NTC7_NTD7_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1231.14,2022.59 +gfx950,256,65536,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441458,0,378.8682,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA3_NTB4_NTC4_NTD5_NTM0_NEPBS14_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO4_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1451.05,1793.44 +gfx950,256,67268,1024,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440195,0,187.8122,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x224x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_7_MO40_NTn1_NTA2_NTB3_NTC7_NTD5_NTM0_NEPBS14_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1126.69,1850.56 +gfx950,256,4,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,29,1,3.557,flydsl_gemm2_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,0.44,114.29 +gfx950,256,8,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4950,8,7.5352,flydsl_gemm7_abf16_wbf16_bf16_t16x64x64_split_k8_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0208,23.38,2940.78 +gfx950,256,32,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,29,1,3.5318,flydsl_gemm2_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,3.56,141.49 +gfx950,256,40,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3204,4,10.4365,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k4_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0142,84.4,2176.63 +gfx950,256,64,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,189,1,3.5574,flydsl_gemm2_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,7.07,170.41 +gfx950,256,96,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4663,4,12.8527,flydsl_gemm4_abf16_wbf16_bf16_t64x64x128_split_k4_block_m_warp2_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0139,164.47,1843.29 +gfx950,256,512,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,545,1,3.8906,flydsl_gemm2_abf16_wbf16_bf16_t64x64x64_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,51.75,539.03 +gfx950,256,1024,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440060,0,35.685,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT96x128x128_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT6_2_MO40_NTn1_NTA3_NTB3_NTC5_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK3_SKFTR0_SKXCCM8_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA2_VWB2_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,631.88,1116.6 +gfx950,256,8192,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439530,0,9.5998,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x128x64_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT4_4_MO40_NTn1_NTA1_NTB1_NTC7_NTD7_NTM0_NEPBS10_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW4_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA4_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,335.55,2880.91 +gfx950,256,1,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,627,7,8.0835,flydsl_gemm3_abf16_wbf16_bf16_t16x64x128_split_k7_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0147,3.75,3747.91 +gfx950,256,16,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,1582,7,9.2773,flydsl_gemm5_abf16_wbf16_bf16_t16x64x64_split_k7_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0201,52.22,3295.63 +gfx950,256,48,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,1701,7,12.1468,flydsl_gemm5_abf16_wbf16_bf16_t48x64x64_split_k7_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0266,119.65,2565.99 +gfx950,256,128,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439702,0,18.4489,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x64x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB2_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT4_2_MO40_NTn1_NTA1_NTB1_NTC6_NTD1_NTM0_NEPBS10_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA4_VWB2_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,210.07,1769.93 +gfx950,256,2048,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440189,0,71.545,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT192x192x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA256_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT6_6_MO40_NTn1_NTA3_NTB0_NTC7_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA2_VWB2_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,866.71,954.48 +gfx950,256,2,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5801,4,4.9418,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0078,2.86,1435.36 +gfx950,256,24,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439978,0,6.0292,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x32x512_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB2048_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_2_MO40_NTn1_NTA4_NTB2_NTC1_NTD3_NTM0_NEPBS10_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB2_WSGRA0_WSGRB0_WS64_WG16_4_4,0.0,28.17,1204.51 +gfx950,256,56,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3786,1,6.9831,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.003,56.76,1075.16 +gfx950,256,256,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4320,1,9.2119,flydsl_gemm3_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.003,196.7,981.77 +gfx950,256,4096,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440625,0,29.4956,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x160x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_5_MO40_NTn1_NTA3_NTB2_NTC1_NTD6_NTM0_NEPBS10_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,982.89,1306.47 +gfx950,256,4,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3132,1,3.9246,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,3.21,808.85 +gfx950,256,32,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3132,1,4.6231,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,21.77,730.05 +gfx950,256,64,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,1845,1,4.2913,flydsl_gemm3_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,46.92,839.95 +gfx950,256,512,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3701,1,6.851,flydsl_gemm4_abf16_wbf16_bf16_t64x64x64_split_k1_block_m_warp2_block_n_warp1_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,235.09,994.85 +gfx950,256,4608,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440488,0,18.8897,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x224x64_MI16x16x1_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_7_MO40_NTn1_NTA1_NTB1_NTC5_NTD7_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,767.38,1915.11 +gfx950,256,9216,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439738,0,34.2747,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x224x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_7_MO40_NTn1_NTA2_NTB3_NTC7_NTD5_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO4_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,845.84,2019.16 +gfx950,256,15360,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441461,0,53.9537,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,895.55,2098.95 +gfx950,256,19219,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,72.2691,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,836.56,1949.76 +gfx950,256,21678,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441462,0,75.3908,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,904.53,2102.82 +gfx950,256,25600,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,92.4865,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,870.73,2018.09 +gfx950,256,29554,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441462,0,106.8728,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,869.9,2011.63 +gfx950,256,32768,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,112.6565,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,914.99,2112.85 +gfx950,256,35488,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441462,0,124.4909,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,896.74,2068.61 +gfx950,256,41472,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,141.7294,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,920.48,2119.65 +gfx950,256,47616,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,163.8358,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,914.25,2102.45 +gfx950,256,52224,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,175.0168,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,938.67,2156.86 +gfx950,256,58368,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,192.7636,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,952.51,2186.76 +gfx950,256,62279,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,206.2715,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,949.78,2179.47 +gfx950,256,66048,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,222.6049,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,933.35,2140.91 +gfx950,256,70709,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,234.737,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,947.58,2172.59 +gfx950,256,73928,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,240.7166,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,966.1,2214.48 +gfx950,256,79872,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,265.2275,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,947.32,2170.47 +gfx950,256,86016,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,278.3358,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,972.15,2226.48 +gfx950,256,89088,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,287.4484,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,974.95,2232.5 +gfx950,256,95232,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,305.203,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,981.56,2246.93 +gfx950,256,2,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5227,7,11.2344,flydsl_gemm8_abf16_wbf16_bf16_t16x128x64_split_k7_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0186,9.15,4577.3 +gfx950,256,24,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439842,0,14.792,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT32x32x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT1_1_MO40_NTn1_NTA7_NTB1_NTC0_NTD3_NTM0_NEPBS14_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS0_SPO1_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM4_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,83.36,3508.4 +gfx950,256,56,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2005,4,17.1785,flydsl_gemm3_abf16_wbf16_bf16_t32x64x128_split_k4_block_m_warp1_block_n_warp2_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0212,167.49,3061.06 +gfx950,256,256,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441293,0,28.3056,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT224x128x64_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA128_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT7_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,464.69,2009.68 +gfx950,256,4096,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441281,0,161.0823,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA0_NTB0_NTC0_NTD4_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1306.5,865.77 +gfx950,256,2,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,1484,4,8.6107,flydsl_gemm3_abf16_wbf16_bf16_t16x64x128_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0111,7.79,3900.64 +gfx950,256,4,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439633,0,5.4097,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x16x512_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_1_MO40_NTn1_NTA7_NTB2_NTC2_NTD0_NTM0_NEPBS4_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS0_SPO1_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_4,0.0,6.2,1558.23 +gfx950,256,24,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439911,0,10.7925,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x32x512_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB2048_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_2_MO40_NTn1_NTA4_NTB2_NTC3_NTD2_NTM0_NEPBS10_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB2_WSGRA0_WSGRB0_WS64_WG16_4_4,0.0,74.62,3145.49 +gfx950,256,32,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3079,1,5.6431,flydsl_gemm5_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,47.57,1544.59 +gfx950,256,56,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4325,2,12.3324,flydsl_gemm8_abf16_wbf16_bf16_t32x64x64_split_k2_block_m_warp2_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0049,152.37,2795.23 +gfx950,256,64,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2158,1,5.9344,flydsl_gemm3_abf16_wbf16_bf16_t16x64x256_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,90.47,1523.99 +gfx950,256,256,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4888,1,19.3188,flydsl_gemm5_abf16_wbf16_bf16_t64x64x128_split_k1_block_m_warp2_block_n_warp2_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0049,444.64,1953.99 +gfx950,256,512,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3521,1,10.5588,flydsl_gemm4_abf16_wbf16_bf16_t128x64x64_split_k1_block_m_warp2_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,406.77,1291.01 +gfx950,256,4096,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440505,0,102.8572,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA0_NTB1_NTC4_NTD7_NTM0_NEPBS12_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1336.21,978.67 +gfx950,256,8192,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441462,0,59.9041,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,1147.16,1540.37 +gfx950,256,49152,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441461,0,340.0353,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,1212.57,1504.86 +gfx950,256,1,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3574,2,5.3402,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0011,1.77,1769.31 +gfx950,256,16,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2370,1,5.8714,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,25.72,1638.01 +gfx950,256,48,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4183,1,5.8103,flydsl_gemm8_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,77.96,1717.27 +gfx950,256,128,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2672,1,7.8867,flydsl_gemm3_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,153.16,1379.41 +gfx950,256,2048,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440436,0,21.8094,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x160x64_MI16x16x1_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_5_MO40_NTn1_NTA0_NTB3_NTC6_NTD3_NTM0_NEPBS12_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,886.19,1490.45 +gfx950,256,24576,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440614,0,208.0274,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA1_NTB2_NTC5_NTD5_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1114.89,1376.07 +gfx950,256,67268,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,519.4038,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,1222.21,1476.97 +gfx950,256,8,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440306,0,18.3141,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT32x16x512_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_1_MO40_NTn1_NTA7_NTB3_NTC1_NTD0_NTM0_NEPBS10_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG32_4_2,0.0,39.38,4933.91 +gfx950,256,40,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440599,0,23.2429,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x48x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_3_MO40_NTn1_NTA1_NTB1_NTC1_NTD1_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW1_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,155.14,3924.69 +gfx950,256,96,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440051,0,27.1443,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x96x128_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_3_MO40_NTn1_NTA3_NTB1_NTC6_NTD3_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,318.81,3416.12 +gfx950,256,1024,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440800,0,101.9095,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x224x64_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB128_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_7_MO40_NTn1_NTA4_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,905.79,1154.97 +gfx950,256,1,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5329,1,4.7814,flydsl_gemm8_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,2.3,2306.0 +gfx950,256,2,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,6420,2,7.1904,flydsl_gemm7_abf16_wbf16_bf16_t16x64x64_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0021,6.12,3067.27 +gfx950,256,4,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4047,2,11.1513,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0054,18.43,4615.27 +gfx950,256,4,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2916,2,12.2681,flydsl_gemm5_abf16_wbf16_bf16_t16x64x64_split_k2_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0114,19.74,4943.42 +gfx950,256,8,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,1484,4,12.3301,flydsl_gemm3_abf16_wbf16_bf16_t16x64x128_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.012,38.1,4776.97 +gfx950,256,16,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3770,1,5.0853,flydsl_gemm5_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,34.64,2215.01 +gfx950,256,24,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3787,1,7.9033,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,66.87,2839.05 +gfx950,256,32,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4119,2,12.5225,flydsl_gemm6_abf16_wbf16_bf16_t32x64x128_split_k2_block_m_warp2_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0048,131.3,4157.98 +gfx950,256,32,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3022,2,15.0051,flydsl_gemm5_abf16_wbf16_bf16_t32x64x64_split_k2_block_m_warp2_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0048,129.14,4084.23 +gfx950,256,40,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3700,1,15.408,flydsl_gemm6_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.005,152.44,3869.51 +gfx950,256,48,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3036,1,5.3724,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp2_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,98.37,2191.18 +gfx950,256,56,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5077,1,8.7281,flydsl_gemm5_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,141.28,2634.59 +gfx950,256,64,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2617,1,15.1478,flydsl_gemm3_abf16_wbf16_bf16_t32x64x256_split_k1_block_m_warp1_block_n_warp2_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0048,217.08,3482.78 +gfx950,256,64,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3288,1,17.9141,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,216.34,3461.71 +gfx950,256,96,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2373,1,18.7467,flydsl_gemm3_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0049,300.7,3247.66 +gfx950,256,128,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5206,1,7.0673,flydsl_gemm4_abf16_wbf16_bf16_t64x64x128_split_k1_block_m_warp4_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,199.41,1845.35 +gfx950,256,256,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3519,1,13.0052,flydsl_gemm3_abf16_wbf16_bf16_t128x64x128_split_k1_block_m_warp4_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,433.45,2035.84 +gfx950,256,512,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440199,0,34.3668,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x128x128_MI32x32x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA8_LPB8_LPM0_LRVW8_LWPMn1_MIAV0_MIWT2_2_MO40_NTn1_NTA2_NTB3_NTC7_NTD7_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW2_SK3_SKFTR0_SKXCCM8_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA2_VWB2_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,765.47,1815.42 +gfx950,256,512,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440199,0,39.6501,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x128x128_MI32x32x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA8_LPB8_LPM0_LRVW8_LWPMn1_MIAV0_MIWT2_2_MO40_NTn1_NTA2_NTB3_NTC7_NTD7_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW2_SK3_SKFTR0_SKXCCM8_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA2_VWB2_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,781.95,1821.45 +gfx950,256,1024,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441313,0,60.4613,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x256x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_8_MO40_NTn1_NTA0_NTB0_NTC4_NTD4_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW4_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA4_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,994.51,1352.75 +gfx950,256,2048,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441462,0,24.1266,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,934.59,1803.65 +gfx950,256,4096,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440587,0,83.1428,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA0_NTB1_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1084.81,1122.45 +gfx950,256,8192,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,asm,10,1,315.322,_ZN5aiter24bf16gemm_bf16_tn_256x256E,0.0,1334.85,721.61 +gfx950,256,8192,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439609,0,360.7631,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA2_NTB2_NTC2_NTD5_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO4_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1375.05,685.22 +gfx950,256,16384,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440271,0,632.1955,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA0_NTB3_NTC4_NTD5_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK0_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1521.8,676.72 +gfx950,256,8,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2975,7,24.5052,flydsl_gemm4_abf16_wbf16_bf16_t16x256x64_split_k7_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0191,39.54,4952.43 +gfx950,256,40,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440176,0,29.0145,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT80x64x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT5_1_MO40_NTn1_NTA3_NTB0_NTC1_NTD1_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,166.97,4217.2 +gfx950,256,96,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441310,0,33.045,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT160x96x64_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA128_LBSPPB128_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT5_3_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,351.84,3755.75 +gfx950,256,1024,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440011,0,122.1848,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x160x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_5_MO40_NTn1_NTA2_NTB3_NTC4_NTD6_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1015.0,1252.96 +gfx950,256,1,16160,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440907,0,36.562,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT32x16x512_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT2_1_MO40_NTn1_NTA4_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_4,0.0,6.34,6337.63 +gfx950,256,16,16160,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440891,0,37.4289,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x32x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT2_1_MO40_NTn1_NTA4_NTB0_NTC3_NTD1_NTM0_NEPBS10_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK3_SKFTR0_SKXCCM8_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,99.03,6209.54 +gfx950,256,48,16160,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440929,0,39.8612,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x48x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_3_MO40_NTn1_NTA4_NTB0_NTC1_NTD5_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,278.97,5868.09 +gfx950,256,128,16160,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440209,0,47.7987,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x128x128_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT4_4_MO40_NTn1_NTA7_NTB2_NTC2_NTD2_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,620.39,4971.72 +gfx950,256,1,20480,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440602,0,45.8447,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x16x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT1_1_MO40_NTn1_NTA7_NTB1_NTC1_NTD1_NTM0_NEPBS10_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS0_SPO1_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_4,0.0,6.4,6405.46 +gfx950,256,16,20480,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440308,0,46.9462,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT96x16x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT3_1_MO40_NTn1_NTA7_NTB0_NTC3_NTD2_NTM0_NEPBS10_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS0_SPO0_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG32_4_1,0.0,100.06,6272.84 +gfx950,256,48,20480,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439915,0,50.2078,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT80x64x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT5_1_MO40_NTn1_NTA5_NTB2_NTC1_NTD3_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS0_SPO1_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,280.69,5900.59 +gfx950,256,128,20480,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439644,0,62.7421,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT96x128x128_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT6_2_MO40_NTn1_NTA7_NTB0_NTC4_NTD1_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK3_SKFTR0_SKXCCM4_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA2_VWB2_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,598.98,4792.3 +gfx950,256,48,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2463,4,7.2689,flydsl_gemm4_abf16_wbf16_bf16_t16x64x256_split_k4_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0193,55.39,1221.66 +gfx950,256,128,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440531,0,9.2171,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x32x512_MI16x16x1_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT1_1_MO40_NTn1_NTA7_NTB2_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_8_2,0.0,116.49,1052.32 +gfx950,256,2048,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439992,0,27.9873,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT96x96x128_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA256_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT3_3_MO40_NTn1_NTA3_NTB2_NTC0_NTD4_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,613.85,1049.05 +gfx950,256,24576,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,opus,6400,0,159.375,opus_gemm_mono_tile_512x192x256x64_2x4_16x16x32_0x0x0_nooob_4g_safe,0.0,1293.54,1631.67 +gfx950,256,67268,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440488,0,420.6321,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x224x64_MI16x16x1_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_7_MO40_NTn1_NTA1_NTB1_NTC5_NTD7_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1341.52,1657.54 +gfx950,256,1,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5257,7,7.0665,flydsl_gemm8_abf16_wbf16_bf16_t16x64x64_split_k7_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0163,3.12,3118.59 +gfx950,256,8,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,29,1,3.5135,flydsl_gemm2_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,0.9,119.49 +gfx950,256,16,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4495,8,7.9527,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k8_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0225,44.3,2803.91 +gfx950,256,40,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,29,1,3.573,flydsl_gemm2_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,4.4,147.31 +gfx950,256,48,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3204,4,10.793,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k4_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0137,97.93,2117.64 +gfx950,256,96,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,189,1,3.5675,flydsl_gemm2_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,10.58,199.78 +gfx950,256,128,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3375,4,15.5803,flydsl_gemm3_abf16_wbf16_bf16_t32x64x128_split_k4_block_m_warp1_block_n_warp2_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0216,180.91,1556.34 +gfx950,256,1024,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,365,1,3.7674,flydsl_gemm2_abf16_wbf16_bf16_t32x64x64_split_k1_block_m_warp2_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,106.88,1008.94 +gfx950,256,2048,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439726,0,55.632,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT192x256x64_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT12_4_MO40_NTn1_NTA1_NTB1_NTC4_NTD2_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW4_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB4_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,810.63,1036.66 +gfx950,256,2,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,1142,7,8.1456,flydsl_gemm4_abf16_wbf16_bf16_t16x64x64_split_k7_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0154,7.43,3721.61 +gfx950,256,24,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,1137,4,12.2555,flydsl_gemm4_abf16_wbf16_bf16_t16x64x64_split_k4_block_m_warp1_block_n_warp1_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0141,59.29,2506.88 +gfx950,256,56,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,asm,5,7,14.5813,_ZN5aiter39bf16gemm_fp32bf16_tn_64x64_splitk_cleanE,0.0198,116.28,2147.75 +gfx950,256,256,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440339,0,22.7924,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x128x64_MI32x32x1_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA256_LBSPPB256_LBSPPM0_LPA8_LPB8_LPM0_LRVW8_LWPMn1_MIAV1_MIWT2_2_MO40_NTn1_NTA0_NTB3_NTC1_NTD3_NTM0_NEPBS2_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA2_VWB2_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,340.07,1536.87 +gfx950,256,4096,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440189,0,120.7251,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT192x192x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA256_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT6_6_MO40_NTn1_NTA3_NTB0_NTC7_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA2_VWB2_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1027.27,880.51 +gfx950,256,4,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4993,4,5.1098,flydsl_gemm5_abf16_wbf16_bf16_t16x64x64_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.008,5.54,1391.17 +gfx950,256,32,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440288,0,6.1231,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x32x512_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT1_1_MO40_NTn1_NTA7_NTB0_NTC0_NTD0_NTM0_NEPBS14_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_8_2,0.0,36.99,1196.07 +gfx950,256,64,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5279,1,6.7931,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,66.68,1114.28 +gfx950,256,512,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5427,1,11.3147,flydsl_gemm3_abf16_wbf16_bf16_t96x64x128_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,320.28,973.07 +gfx950,256,8192,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440644,0,55.5526,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x320x64_MI32x32x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB128_LBSPPM0_LPA8_LPB8_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_5_MO40_NTn1_NTA0_NTB0_NTC0_NTD4_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW4_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB1_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,1043.73,1259.93 +gfx950,256,8,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3484,1,4.0325,flydsl_gemm7_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,6.24,794.31 +gfx950,256,40,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,1182,1,4.2641,flydsl_gemm3_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,29.51,804.96 +gfx950,256,96,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,681,1,4.7247,flydsl_gemm2_abf16_wbf16_bf16_t32x64x256_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,63.92,811.45 +gfx950,256,1024,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439860,0,8.5843,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT96x128x128_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT6_2_MO40_NTn1_NTA1_NTB1_NTC1_NTD5_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK0_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA2_VWB2_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,375.25,1221.5 +gfx950,256,6144,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440649,0,25.0184,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT384x192x64_MI32x32x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA256_LBSPPB128_LBSPPM0_LPA8_LPB8_LPM0_LRVW8_LWPMn1_MIAV0_MIWT6_3_MO40_NTn1_NTA0_NTB0_NTC4_NTD4_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO4_SVW2_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,772.53,1886.05 +gfx950,256,10752,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441461,0,37.0619,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,912.6,2164.38 +gfx950,256,16384,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439589,0,54.1052,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA3_NTB6_NTC5_NTD3_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,952.58,2228.74 +gfx950,256,19968,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,72.2095,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,869.88,2025.72 +gfx950,256,22161,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440144,0,84.9509,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x224x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_7_MO40_NTn1_NTA1_NTB3_NTC5_NTD5_NTM0_NEPBS14_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,820.62,1906.93 +gfx950,256,26112,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440622,0,91.4236,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA1_NTB1_NTC6_NTD6_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,898.47,2081.7 +gfx950,256,30720,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441462,0,109.4361,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,883.04,2040.89 +gfx950,256,33792,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441461,0,124.757,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,852.06,1966.76 +gfx950,256,36864,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,123.5056,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,938.94,2164.98 +gfx950,256,43008,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,147.1409,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,919.47,2116.52 +gfx950,256,49152,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,163.5048,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,945.65,2174.05 +gfx950,256,53760,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,178.242,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,948.79,2179.61 +gfx950,256,59904,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,196.2044,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,960.44,2204.53 +gfx950,256,62976,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,211.7489,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,935.57,2146.68 +gfx950,256,67584,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,221.3087,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,960.65,2203.2 +gfx950,256,71610,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,230.5183,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,977.21,2240.37 +gfx950,256,75264,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,242.7539,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,975.31,2235.34 +gfx950,256,81408,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,261.3043,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,980.04,2245.19 +gfx950,256,87552,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,282.1448,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,976.15,2235.44 +gfx950,256,90624,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,290.2721,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,982.11,2248.71 +gfx950,256,96768,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,309.7207,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,982.84,2249.7 +gfx950,256,4,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4489,4,11.4151,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0134,18.0,4508.61 +gfx950,256,32,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4611,4,14.0537,flydsl_gemm6_abf16_wbf16_bf16_t32x64x64_split_k4_block_m_warp2_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0215,116.99,3704.96 +gfx950,256,64,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4310,2,18.0418,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k2_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.006,182.26,2924.13 +gfx950,256,512,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440779,0,39.4911,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT224x128x64_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA128_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT7_4_MO40_NTn1_NTA4_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,666.14,1579.86 +gfx950,256,8192,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441462,0,303.7587,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,1385.66,749.08 +gfx950,256,4,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,1483,4,8.5035,flydsl_gemm3_abf16_wbf16_bf16_t16x64x128_split_k4_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0215,15.78,3953.66 +gfx950,256,8,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440305,0,5.3835,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x16x512_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_1_MO40_NTn1_NTA7_NTB3_NTC3_NTD1_NTM0_NEPBS2_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS0_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_4,0.0,12.47,1573.42 +gfx950,256,32,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,1493,2,10.2083,flydsl_gemm3_abf16_wbf16_bf16_t16x64x256_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0049,105.18,3338.33 +gfx950,256,40,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4183,1,6.1491,flydsl_gemm8_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,54.57,1430.81 +gfx950,256,64,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5442,2,12.7591,flydsl_gemm8_abf16_wbf16_bf16_t32x64x64_split_k2_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0049,168.31,2712.03 +gfx950,256,96,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2310,1,6.2534,flydsl_gemm3_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,128.78,1498.65 +gfx950,256,512,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440209,0,27.3245,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x128x128_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT4_4_MO40_NTn1_NTA7_NTB2_NTC2_NTD2_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,628.73,1535.0 +gfx950,256,1024,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3438,1,13.2672,flydsl_gemm4_abf16_wbf16_bf16_t128x128x64_split_k1_block_m_warp4_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,647.46,1422.63 +gfx950,256,8192,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,asm,10,1,195.4438,_ZN5aiter24bf16gemm_bf16_tn_256x256E,0.0,1406.43,858.42 +gfx950,256,16384,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439532,0,123.059,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA0_NTB1_NTC5_NTD4_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1116.85,1431.51 +gfx950,256,65536,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441459,0,440.7286,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA1_NTB1_NTC7_NTD5_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM8_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1247.38,1541.71 +gfx950,256,2,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3938,2,5.4042,flydsl_gemm7_abf16_wbf16_bf16_t16x64x64_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0017,3.49,1750.44 +gfx950,256,24,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2370,1,5.6754,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,39.91,1710.46 +gfx950,256,56,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2369,1,6.0049,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0022,88.01,1676.63 +gfx950,256,256,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4358,1,8.3638,flydsl_gemm4_abf16_wbf16_bf16_t96x64x64_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,288.85,1473.11 +gfx950,256,4096,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440644,0,40.3847,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x320x64_MI32x32x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB128_LBSPPM0_LPA8_LPB8_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_5_MO40_NTn1_NTA0_NTB0_NTC0_NTD4_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW4_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB1_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,957.16,1376.13 +gfx950,256,32768,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440624,0,259.1705,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA1_NTB3_NTC4_NTD5_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1193.18,1460.57 +gfx950,256,1,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440306,0,17.1174,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT32x16x512_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_1_MO40_NTn1_NTA7_NTB3_NTC1_NTD0_NTM0_NEPBS10_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG32_4_2,0.0,5.27,5267.84 +gfx950,256,16,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440306,0,18.3622,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT32x16x512_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_1_MO40_NTn1_NTA7_NTB3_NTC1_NTD0_NTM0_NEPBS10_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG32_4_2,0.0,78.55,4932.71 +gfx950,256,48,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440599,0,23.2367,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x48x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_3_MO40_NTn1_NTA1_NTB1_NTC1_NTD1_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW1_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,186.21,3935.01 +gfx950,256,128,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440676,0,28.9546,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x128x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_4_MO40_NTn1_NTA0_NTB0_NTC4_NTD0_NTM0_NEPBS2_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA4_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,398.5,3232.28 +gfx950,256,2048,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439532,0,162.057,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA0_NTB1_NTC5_NTD4_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1139.21,896.35 +gfx950,256,1,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2450,4,12.1912,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k4_block_m_warp1_block_n_warp1_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0153,4.82,4818.46 +gfx950,256,2,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5329,1,4.8262,flydsl_gemm8_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,4.56,2287.89 +gfx950,256,4,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4987,2,6.7614,flydsl_gemm5_abf16_wbf16_bf16_t16x64x64_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0025,13.03,3267.03 +gfx950,256,8,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2665,4,11.8274,flydsl_gemm4_abf16_wbf16_bf16_t16x64x64_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.012,34.75,4358.71 +gfx950,256,8,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3453,2,12.4243,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0043,38.99,4888.61 +gfx950,256,16,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439704,0,13.2082,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT32x16x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT2_1_MO40_NTn1_NTA7_NTB1_NTC0_NTD2_NTM0_NEPBS8_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK3_SKFTR0_SKXCCM4_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_2,0.0,71.13,4473.03 +gfx950,256,24,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4426,1,4.9985,flydsl_gemm6_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,52.86,2278.88 +gfx950,256,32,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4965,1,8.1295,flydsl_gemm5_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0028,86.68,2777.19 +gfx950,256,40,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4856,1,13.9618,flydsl_gemm8_abf16_wbf16_bf16_t32x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0047,147.2,3741.67 +gfx950,256,40,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2310,1,15.6981,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0047,154.3,3915.55 +gfx950,256,48,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4262,1,15.933,flydsl_gemm8_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,176.9,3753.32 +gfx950,256,56,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5044,1,5.4542,flydsl_gemm7_abf16_wbf16_bf16_t32x64x64_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,113.04,2181.6 +gfx950,256,64,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5549,1,8.8813,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,158.68,2604.82 +gfx950,256,96,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2580,1,16.7457,flydsl_gemm3_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,294.55,3191.54 +gfx950,256,96,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2132,1,19.256,flydsl_gemm3_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0048,301.9,3258.34 +gfx950,256,128,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4888,1,20.8001,flydsl_gemm5_abf16_wbf16_bf16_t64x64x128_split_k1_block_m_warp2_block_n_warp2_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0049,361.35,2961.71 +gfx950,256,256,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2639,1,8.9951,flydsl_gemm3_abf16_wbf16_bf16_t128x64x128_split_k1_block_m_warp2_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,313.35,1675.72 +gfx950,256,512,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5575,1,17.9038,flydsl_gemm4_abf16_wbf16_bf16_t128x128x64_split_k1_block_m_warp4_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,629.71,1727.73 +gfx950,256,1024,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440412,0,54.5307,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_4_MO40_NTn1_NTA2_NTB3_NTC5_NTD5_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA8_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,964.84,1346.04 +gfx950,256,1024,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441293,0,64.4159,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT224x128x64_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA128_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT7_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,962.63,1302.26 +gfx950,256,2048,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440282,0,97.4815,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA3_NTB3_NTC5_NTD5_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1233.66,1075.67 +gfx950,256,4096,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441462,0,45.4354,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,992.56,1673.18 +gfx950,256,8192,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441462,0,162.5715,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,1109.6,1012.64 +gfx950,256,16817,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440271,0,702.3755,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA0_NTB3_NTC4_NTD5_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK0_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1405.94,622.99 +gfx950,256,1,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,1777,7,22.8172,flydsl_gemm3_abf16_wbf16_bf16_t16x256x64_split_k7_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0183,5.31,5309.23 +gfx950,256,16,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439577,0,26.5909,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x16x128_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS1_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_1_MO40_NTn1_NTA3_NTB1_NTC0_NTD1_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR1_PLR1_PKA1_SIA3_SS0_SPO1_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_1,0.0,72.87,4573.38 +gfx950,256,48,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440176,0,29.5913,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT80x64x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT5_1_MO40_NTn1_NTA3_NTB0_NTC1_NTD1_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,196.45,4143.44 +gfx950,256,128,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440177,0,35.5119,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT80x128x128_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA256_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT5_2_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB2_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,436.53,3522.99 +gfx950,256,2048,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440649,0,219.9997,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT384x192x64_MI32x32x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA256_LBSPPB128_LBSPPM0_LPA8_LPB8_LPM0_LRVW8_LWPMn1_MIAV0_MIWT6_3_MO40_NTn1_NTA0_NTB0_NTC4_NTD4_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO4_SVW2_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,1127.43,841.25 +gfx950,256,2,16160,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439704,0,36.9768,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT32x16x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT2_1_MO40_NTn1_NTA7_NTB1_NTC0_NTD2_NTM0_NEPBS8_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK3_SKFTR0_SKXCCM4_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_2,0.0,12.53,6267.8 +gfx950,256,24,16160,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439637,0,37.9611,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x32x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT2_1_MO40_NTn1_NTA7_NTB1_NTC0_NTD6_NTM0_NEPBS12_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK3_SKFTR0_SKXCCM8_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,146.47,6132.32 +gfx950,256,56,16160,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440552,0,41.5177,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x64x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA2048_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_1_MO40_NTn1_NTA5_NTB3_NTC5_NTD2_NTM0_NEPBS14_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO4_SVW4_SK3_SKFTR0_SKXCCM8_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,312.48,5642.95 +gfx950,256,256,16160,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440799,0,72.8463,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x256x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_8_MO40_NTn1_NTA4_NTB0_NTC4_NTD4_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW4_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA4_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,814.15,3344.22 +gfx950,256,2,20480,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440889,0,45.9204,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT96x32x256_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT3_1_MO40_NTn1_NTA4_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,12.79,6396.11 +gfx950,256,24,20480,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440889,0,47.1916,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT96x32x256_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT3_1_MO40_NTn1_NTA4_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,149.32,6249.59 +gfx950,256,56,20480,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439915,0,50.6412,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT80x64x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT5_1_MO40_NTn1_NTA5_NTB2_NTC1_NTD3_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS0_SPO1_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,324.67,5858.82 +gfx950,256,24,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2452,4,6.7999,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0116,29.61,1269.78 +gfx950,256,56,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2463,4,7.7053,flydsl_gemm4_abf16_wbf16_bf16_t16x64x256_split_k4_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0192,60.97,1163.1 +gfx950,256,256,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3946,2,12.0071,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k2_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0049,178.85,916.96 +gfx950,256,4096,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440677,0,39.6483,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x128x128_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB2048_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT2_8_MO40_NTn1_NTA0_NTB4_NTC0_NTD4_NTM0_NEPBS8_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA2_VWB8_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,866.61,1269.45 +gfx950,256,32768,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440451,0,199.0305,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA0_NTB3_NTC4_NTD7_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1381.08,1728.04 +gfx950,256,1,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,29,1,3.5102,flydsl_gemm2_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,0.11,112.97 +gfx950,256,2,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5257,7,7.2455,flydsl_gemm8_abf16_wbf16_bf16_t16x64x64_split_k7_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.015,6.08,3043.95 +gfx950,256,16,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,29,1,3.5742,flydsl_gemm2_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,1.76,124.91 +gfx950,256,24,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2987,4,9.418,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0142,56.11,2382.45 +gfx950,256,48,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,29,1,3.593,flydsl_gemm2_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,5.25,153.9 +gfx950,256,56,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3204,4,11.0712,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k4_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.013,111.38,2077.01 +gfx950,256,128,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,189,1,2.9663,flydsl_gemm2_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,16.97,276.17 +gfx950,256,256,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441318,0,19.2485,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x96x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_3_MO40_NTn1_NTA0_NTB0_NTC7_NTD4_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA4_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,292.86,1375.51 +gfx950,256,2048,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,553,1,4.6652,flydsl_gemm2_abf16_wbf16_bf16_t64x64x64_split_k1_block_m_warp4_block_n_warp1_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,172.62,1545.26 +gfx950,256,4096,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439726,0,91.1902,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT192x256x64_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT12_4_MO40_NTn1_NTA1_NTB1_NTC4_NTD2_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW4_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB4_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,989.08,1023.39 +gfx950,256,4,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,1115,7,8.42,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k7_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0187,14.38,3604.74 +gfx950,256,32,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2043,7,10.7108,flydsl_gemm6_abf16_wbf16_bf16_t32x64x64_split_k7_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0196,90.46,2882.28 +gfx950,256,64,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2783,7,14.5533,flydsl_gemm6_abf16_wbf16_bf16_t64x64x64_split_k7_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0194,133.15,2162.09 +gfx950,256,512,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441293,0,30.5838,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT224x128x64_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA128_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT7_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,506.87,1300.7 +gfx950,256,8192,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441355,0,220.8372,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT176x256x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA4_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA128_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT11_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB4_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,1123.15,825.59 +gfx950,256,8,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4993,4,5.2247,flydsl_gemm5_abf16_wbf16_bf16_t16x64x64_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0076,10.84,1366.46 +gfx950,256,40,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3790,2,6.6023,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0027,42.88,1118.56 +gfx950,256,96,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5278,1,7.0003,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.003,97.06,1116.4 +gfx950,256,1024,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439427,0,15.1625,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x80x128_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT2_5_MO40_NTn1_NTA3_NTB3_NTC1_NTD6_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,478.01,985.47 +gfx950,256,1,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3132,1,3.9325,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,0.8,801.75 +gfx950,256,16,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3484,1,4.0831,flydsl_gemm7_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,12.33,798.51 +gfx950,256,48,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2027,1,4.1748,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,36.17,835.92 +gfx950,256,128,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4592,1,4.3285,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,93.02,938.72 +gfx950,256,2048,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439931,0,11.3739,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT192x128x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT12_2_MO40_NTn1_NTA3_NTB3_NTC6_NTD4_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK0_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA4_VWB2_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,566.42,1567.25 +gfx950,256,7680,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440571,0,29.9853,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x192x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT16_3_MO40_NTn1_NTA0_NTB1_NTC6_NTD5_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,805.7,1940.82 +gfx950,256,12288,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440571,0,43.7197,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x192x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT16_3_MO40_NTn1_NTA0_NTB1_NTC6_NTD5_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,884.15,2086.61 +gfx950,256,16896,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,64.8801,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,819.21,1915.17 +gfx950,256,20969,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440575,0,73.6329,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA0_NTB2_NTC7_NTD6_NTM0_NEPBS12_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,895.83,2084.01 +gfx950,256,23040,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,85.9306,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,843.44,1958.52 +gfx950,256,27648,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440606,0,103.9709,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x224x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_7_MO40_NTn1_NTA1_NTB0_NTC7_NTD4_NTM0_NEPBS14_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,836.51,1936.37 +gfx950,256,31232,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,111.1431,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,883.97,2042.56 +gfx950,256,34087,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441461,0,122.3874,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,876.14,2022.11 +gfx950,256,38400,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,139.935,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,863.23,1989.47 +gfx950,256,44544,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440267,0,162.0713,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA1_NTB3_NTC2_NTD2_NTM0_NEPBS10_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK0_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,864.58,1989.48 +gfx950,256,49190,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441461,0,172.4651,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,897.22,2062.68 +gfx950,256,55296,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,189.1058,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,919.84,2112.61 +gfx950,256,60928,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,208.5895,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,918.85,2108.82 +gfx950,256,64512,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,214.2631,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,947.14,2172.88 +gfx950,256,69120,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,226.2852,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,960.88,2203.4 +gfx950,256,72192,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,238.5843,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,951.85,2182.11 +gfx950,256,76800,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441461,0,255.0874,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,947.09,2170.43 +gfx950,256,82944,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,269.7568,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,967.24,2215.66 +gfx950,256,87883,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,285.6736,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,967.73,2216.13 +gfx950,256,92160,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,294.8084,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,983.39,2251.46 +gfx950,256,98304,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,311.0214,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,994.26,2275.69 +gfx950,256,8,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4944,4,11.5249,flydsl_gemm7_abf16_wbf16_bf16_t16x64x64_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0138,35.67,4473.12 +gfx950,256,40,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3194,2,15.6396,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k2_block_m_warp1_block_n_warp2_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0138,131.41,3340.26 +gfx950,256,96,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439436,0,20.4963,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x96x128_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT2_3_MO40_NTn1_NTA1_NTB2_NTC1_NTD1_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW2_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,240.65,2607.52 +gfx950,256,1024,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440679,0,64.3874,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT224x256x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA128_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT7_8_MO40_NTn1_NTA0_NTB0_NTC4_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,817.14,1139.98 +gfx950,256,1,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3938,2,5.1103,flydsl_gemm7_abf16_wbf16_bf16_t16x64x64_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.001,1.64,1643.51 +gfx950,256,8,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4020,4,8.9622,flydsl_gemm7_abf16_wbf16_bf16_t16x64x64_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0125,29.95,3758.62 +gfx950,256,16,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440303,0,5.4088,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x16x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_1_MO40_NTn1_NTA4_NTB1_NTC3_NTD0_NTM0_NEPBS8_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW1_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_2,0.0,24.81,1581.21 +gfx950,256,40,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2622,2,11.4887,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k2_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.005,116.83,2977.69 +gfx950,256,48,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3079,1,5.6986,flydsl_gemm5_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,70.66,1558.3 +gfx950,256,96,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439635,0,15.2922,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x48x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_3_MO40_NTn1_NTA7_NTB2_NTC1_NTD3_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS0_SPO0_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM4_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,210.65,2297.07 +gfx950,256,128,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3861,1,6.222,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp2_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,172.57,1558.88 +gfx950,256,1024,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440209,0,38.1803,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x128x128_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT4_4_MO40_NTn1_NTA7_NTB2_NTC2_NTD2_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,899.93,1318.26 +gfx950,256,2048,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440655,0,19.8682,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_4_MO40_NTn1_NTA0_NTB0_NTC4_NTD4_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA8_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,864.69,1477.74 +gfx950,256,16384,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440282,0,373.0705,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA3_NTB3_NTC5_NTD5_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1473.6,809.47 +gfx950,256,24576,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440588,0,178.8387,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA1_NTB0_NTC7_NTD5_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1152.76,1454.09 +gfx950,256,67268,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441462,0,468.0484,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,1205.61,1489.62 +gfx950,256,4,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3574,2,5.7611,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0024,6.55,1645.91 +gfx950,256,32,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3079,1,6.3664,flydsl_gemm5_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,47.43,1538.96 +gfx950,256,64,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3261,1,6.6995,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0022,90.15,1516.24 +gfx950,256,512,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3209,1,11.2497,flydsl_gemm3_abf16_wbf16_bf16_t80x128x128_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,429.51,1351.53 +gfx950,256,8192,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441461,0,79.8646,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,968.01,1273.55 +gfx950,256,49152,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,378.6002,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,1225.19,1487.28 +gfx950,256,2,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440907,0,17.9379,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT32x16x512_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT2_1_MO40_NTn1_NTA4_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_4,0.0,10.05,5028.38 +gfx950,256,24,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440651,0,21.184,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x32x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_2_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS12_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB2_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,102.13,4285.81 +gfx950,256,56,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439974,0,24.8945,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x64x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA2_GRVWB2_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA256_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT4_2_MO40_NTn1_NTA1_NTB3_NTC6_NTD1_NTM0_NEPBS12_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM7_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB2_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,202.78,3681.61 +gfx950,256,256,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441293,0,39.7968,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT224x128x64_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA128_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT7_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,579.87,2438.24 +gfx950,256,4096,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440543,0,297.0309,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA2_NTB1_NTC7_NTD4_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1243.08,674.6 +gfx950,256,1,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3453,4,11.3674,flydsl_gemm5_abf16_wbf16_bf16_t16x128x64_split_k4_block_m_warp1_block_n_warp1_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0133,4.52,4521.85 +gfx950,256,1,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4188,2,12.1281,flydsl_gemm8_abf16_wbf16_bf16_t16x64x64_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0049,4.99,4994.85 +gfx950,256,2,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3656,4,12.2926,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0104,9.55,4780.54 +gfx950,256,4,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5329,1,4.8262,flydsl_gemm8_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,9.13,2294.46 +gfx950,256,8,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,6851,2,7.1991,flydsl_gemm8_abf16_wbf16_bf16_t16x64x64_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0026,24.47,3078.07 +gfx950,256,16,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4047,2,12.0581,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0047,68.18,4289.59 +gfx950,256,16,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2916,2,12.3525,flydsl_gemm5_abf16_wbf16_bf16_t16x64x64_split_k2_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.012,78.44,4931.78 +gfx950,256,24,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440111,0,14.4802,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT32x32x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT1_1_MO40_NTn1_NTA7_NTB2_NTC2_NTD1_NTM0_NEPBS14_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,97.33,4092.55 +gfx950,256,32,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4441,1,5.0821,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,69.33,2266.38 +gfx950,256,40,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4057,1,7.6168,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,115.64,2982.41 +gfx950,256,48,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2638,1,14.24,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,173.19,3680.65 +gfx950,256,48,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2146,1,16.2132,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,179.28,3802.39 +gfx950,256,56,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,1683,1,16.357,flydsl_gemm3_abf16_wbf16_bf16_t32x64x256_split_k1_block_m_warp2_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0047,201.04,3667.04 +gfx950,256,64,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5297,1,6.1278,flydsl_gemm5_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,114.99,1962.51 +gfx950,256,96,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,6001,1,9.5942,flydsl_gemm4_abf16_wbf16_bf16_t64x64x128_split_k1_block_m_warp2_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,220.33,2469.33 +gfx950,256,128,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5517,1,18.8788,flydsl_gemm5_abf16_wbf16_bf16_t64x64x128_split_k1_block_m_warp2_block_n_warp2_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0049,348.36,2867.38 +gfx950,256,128,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4097,1,21.9988,flydsl_gemm4_abf16_wbf16_bf16_t64x64x128_split_k1_block_m_warp4_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,352.34,2885.23 +gfx950,256,256,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440680,0,28.1272,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x128x128_MI32x32x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB256_LBSPPM0_LPA8_LPB8_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_1_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,534.44,2292.71 +gfx950,256,512,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4266,1,11.9413,flydsl_gemm4_abf16_wbf16_bf16_t128x128x64_split_k1_block_m_warp2_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,472.07,1602.55 +gfx950,256,1024,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440204,0,25.8436,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_4_MO40_NTn1_NTA1_NTB0_NTC2_NTD7_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM4_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,872.5,1541.81 +gfx950,256,2048,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441462,0,85.5936,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,1229.38,1114.81 +gfx950,256,2048,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440335,0,99.428,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA1_NTB3_NTC5_NTD4_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO4_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1247.31,1078.34 +gfx950,256,4096,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440489,0,185.83,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA0_NTB0_NTC4_NTD3_NTM0_NEPBS12_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1294.29,812.54 +gfx950,256,8192,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,95.526,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,944.19,1476.39 +gfx950,256,2,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2975,7,22.7894,flydsl_gemm4_abf16_wbf16_bf16_t16x256x64_split_k7_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0186,10.63,5317.08 +gfx950,256,24,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441381,0,28.5034,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT96x32x128_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA256_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT3_1_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,101.98,4275.28 +gfx950,256,56,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440176,0,30.1355,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT80x64x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT5_1_MO40_NTn1_NTA3_NTB0_NTC1_NTD1_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,225.06,4076.9 +gfx950,256,256,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439447,0,48.9717,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT192x256x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT12_4_MO40_NTn1_NTA1_NTB0_NTC5_NTD2_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB4_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,633.11,2636.34 +gfx950,256,4096,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440570,0,387.4138,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x192x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_6_MO40_NTn1_NTA3_NTB2_NTC6_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB2_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1280.46,642.82 +gfx950,256,4,16160,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439704,0,36.1371,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT32x16x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT2_1_MO40_NTn1_NTA7_NTB1_NTC0_NTD2_NTM0_NEPBS8_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK3_SKFTR0_SKXCCM4_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_2,0.0,25.64,6416.02 +gfx950,256,32,16160,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440891,0,37.7139,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x32x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT2_1_MO40_NTn1_NTA4_NTB0_NTC3_NTD1_NTM0_NEPBS10_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK3_SKFTR0_SKXCCM8_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,196.57,6182.41 +gfx950,256,64,16160,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440890,0,41.6415,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x64x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT2_2_MO40_NTn1_NTA4_NTB0_NTC3_NTD3_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA2_VWB2_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,356.06,5635.14 +gfx950,256,512,16160,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440205,0,116.198,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI32x32x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB256_LBSPPM0_LPA8_LPB8_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_2_MO40_NTn1_NTA3_NTB1_NTC5_NTD7_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW4_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB2_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,1020.8,2199.33 +gfx950,256,4,20480,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440180,0,45.8857,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x16x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT1_1_MO40_NTn1_NTA4_NTB2_NTC2_NTD1_NTM0_NEPBS8_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_4,0.0,25.59,6403.36 +gfx950,256,32,20480,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440889,0,47.7656,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT96x32x256_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT3_1_MO40_NTn1_NTA4_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,196.69,6183.75 +gfx950,256,64,20480,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439915,0,52.6489,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT80x64x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT5_1_MO40_NTn1_NTA5_NTB2_NTC1_NTD3_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS0_SPO1_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,356.9,5643.81 +gfx950,256,32,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4092,8,7.2351,flydsl_gemm7_abf16_wbf16_bf16_t32x64x64_split_k8_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0184,37.1,1204.72 +gfx950,256,64,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3344,4,7.9179,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0126,67.8,1142.22 +gfx950,256,512,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3937,1,14.5053,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,296.1,939.76 +gfx950,256,8192,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440204,0,62.4518,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_4_MO40_NTn1_NTA1_NTB0_NTC2_NTD7_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM4_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1100.36,1477.53 +gfx950,256,49152,1024,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440394,0,288.9558,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA2_NTB7_NTC5_NTD2_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1426.92,1770.88 +gfx950,256,2,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,29,1,2.9072,flydsl_gemm2_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,0.27,137.55 +gfx950,256,4,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5257,7,7.2285,flydsl_gemm8_abf16_wbf16_bf16_t16x64x64_split_k7_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.02,12.19,3055.92 +gfx950,256,24,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,29,1,2.9696,flydsl_gemm2_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,3.18,159.31 +gfx950,256,32,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3213,7,9.04,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k7_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0199,77.95,2497.47 +gfx950,256,56,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,29,1,2.9479,flydsl_gemm2_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,7.47,196.61 +gfx950,256,64,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2920,4,10.7759,flydsl_gemm3_abf16_wbf16_bf16_t32x64x128_split_k4_block_m_warp1_block_n_warp2_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0216,130.78,2146.85 +gfx950,256,256,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,362,1,3.1267,flydsl_gemm2_abf16_wbf16_bf16_t32x64x64_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,32.19,398.24 +gfx950,256,512,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440108,0,24.9535,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x128x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT4_4_MO40_NTn1_NTA2_NTB1_NTC6_NTD3_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA4_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,451.81,1239.63 +gfx950,256,4096,1536,128,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,140,1,6.787,flydsl_gemm2_abf16_wbf16_bf16_t128x64x64_split_k1_block_m_warp4_block_n_warp1_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,237.31,2066.41 +gfx950,256,8192,1536,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440606,0,147.5026,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x224x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_7_MO40_NTn1_NTA1_NTB0_NTC7_NTD4_NTM0_NEPBS14_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1222.95,1116.09 +gfx950,256,8,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2242,7,8.5852,flydsl_gemm7_abf16_wbf16_bf16_t16x64x64_split_k7_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0188,28.21,3544.02 +gfx950,256,40,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439948,0,14.6927,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x48x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB2_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT4_3_MO40_NTn1_NTA5_NTB3_NTC3_NTD3_NTM0_NEPBS8_NLCA1_NLCB1_ONLL1_PGR2_PLR0_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA4_VWB1_WSGRA0_WSGRB0_WS64_WG32_4_2,0.0,82.43,2111.25 +gfx950,256,96,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,asm,8,7,16.3282,_ZN5aiter39bf16gemm_fp32bf16_tn_96x64_splitk_cleanE,0.0195,178.01,1963.44 +gfx950,256,1024,2112,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440405,0,44.867,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x160x64_MI16x16x1_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_10_MO40_NTn1_NTA0_NTB3_NTC3_NTD3_NTM0_NEPBS2_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB2_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,691.03,1098.43 +gfx950,256,1,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5801,4,5.0083,flydsl_gemm6_abf16_wbf16_bf16_t16x64x64_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0095,1.41,1414.77 +gfx950,256,16,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439841,0,5.5452,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x16x512_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_1_MO40_NTn1_NTA6_NTB1_NTC3_NTD0_NTM0_NEPBS14_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_4,0.0,20.42,1298.56 +gfx950,256,48,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4969,2,6.7411,flydsl_gemm5_abf16_wbf16_bf16_t16x64x128_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0027,50.4,1104.65 +gfx950,256,128,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5892,1,7.5355,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.003,120.23,1069.73 +gfx950,256,2048,2304,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439408,0,20.8367,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x160x128_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT4_5_MO40_NTn1_NTA1_NTB1_NTC7_NTD7_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW4_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,695.67,1094.54 +gfx950,256,2,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439772,0,4.3116,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT16x16x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_1_MO40_NTn1_NTA7_NTB1_NTC3_NTD0_NTM0_NEPBS14_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_2,0.0,1.46,732.92 +gfx950,256,24,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3484,1,4.0117,flydsl_gemm7_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,18.82,827.02 +gfx950,256,56,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3484,1,4.6225,flydsl_gemm7_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,38.11,767.36 +gfx950,256,256,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4809,1,5.0567,flydsl_gemm6_abf16_wbf16_bf16_t64x64x64_split_k1_block_m_warp4_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,159.26,984.98 +gfx950,256,4096,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439726,0,17.4384,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT192x256x64_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT12_4_MO40_NTn1_NTA1_NTB1_NTC4_NTD2_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW4_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB4_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,738.88,1864.04 +gfx950,256,8192,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439726,0,31.5209,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT192x256x64_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT12_4_MO40_NTn1_NTA1_NTB1_NTC4_NTD2_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW4_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB4_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,817.55,1962.7 +gfx950,256,13824,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440144,0,51.1432,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x224x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_7_MO40_NTn1_NTA1_NTB3_NTC5_NTD5_NTM0_NEPBS14_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,850.29,1999.02 +gfx950,256,18432,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441462,0,68.4397,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,847.2,1976.43 +gfx950,256,21504,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441459,0,74.5083,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA1_NTB1_NTC7_NTD5_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM8_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,907.9,2110.99 +gfx950,256,24576,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439531,0,88.8562,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x192x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT16_3_MO40_NTn1_NTA0_NTB6_NTC4_NTD2_NTM0_NEPBS12_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,870.05,2017.94 +gfx950,256,29184,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441462,0,107.9862,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,850.15,1966.33 +gfx950,256,32256,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,114.1838,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,888.64,2052.45 +gfx950,256,35328,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,121.6092,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,913.85,2108.2 +gfx950,256,39936,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,137.8542,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,911.31,2099.37 +gfx950,256,46080,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441462,0,155.0495,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,934.9,2150.59 +gfx950,256,50688,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,172.0072,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,927.0,2130.59 +gfx950,256,56832,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,192.5446,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,928.5,2132.06 +gfx950,256,61440,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,204.0453,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,947.21,2173.77 +gfx950,256,65536,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,212.5045,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,970.14,2225.4 +gfx950,256,70656,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,233.2789,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,952.78,2184.54 +gfx950,256,73728,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,238.2667,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,973.4,2231.23 +gfx950,256,78336,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,257.8939,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,955.52,2189.5 +gfx950,256,84480,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,273.2157,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,972.68,2227.9 +gfx950,256,88290,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441461,0,293.8028,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,945.32,2164.75 +gfx950,256,93696,3072,512,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,301.2282,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,978.47,2240.02 +gfx950,256,1,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2958,7,11.4073,flydsl_gemm4_abf16_wbf16_bf16_t16x128x64_split_k7_block_m_warp1_block_n_warp1_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0187,4.5,4506.04 +gfx950,256,16,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4473,4,11.8967,flydsl_gemm6_abf16_wbf16_bf16_t16x64x128_split_k4_block_m_warp1_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0214,69.1,4347.78 +gfx950,256,48,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4654,4,16.2013,flydsl_gemm6_abf16_wbf16_bf16_t48x64x64_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0137,152.23,3235.07 +gfx950,256,128,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439880,0,20.8963,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT112x128x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA2_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA256_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT7_2_MO40_NTn1_NTA1_NTB2_NTC1_NTD5_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM8_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB2_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,314.73,2590.54 +gfx950,256,2048,3584,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440411,0,104.6213,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_4_MO40_NTn1_NTA3_NTB1_NTC6_NTD4_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1005.79,912.06 +gfx950,256,1,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,1484,4,8.3177,flydsl_gemm3_abf16_wbf16_bf16_t16x64x128_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0137,4.03,4036.07 +gfx950,256,2,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2373,2,5.2443,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0017,3.2,1603.47 +gfx950,256,16,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3176,4,9.4724,flydsl_gemm5_abf16_wbf16_bf16_t16x64x64_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0119,56.68,3570.01 +gfx950,256,24,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4183,1,5.5699,flydsl_gemm8_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,36.15,1550.18 +gfx950,256,48,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2621,2,12.0069,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k2_block_m_warp1_block_n_warp2_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.012,134.14,2860.09 +gfx950,256,56,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4183,1,6.1966,flydsl_gemm8_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,75.81,1446.28 +gfx950,256,128,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3937,1,15.4787,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,277.48,2303.27 +gfx950,256,256,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4122,1,7.5615,flydsl_gemm4_abf16_wbf16_bf16_t64x64x128_split_k1_block_m_warp2_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,284.0,1456.07 +gfx950,256,2048,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440411,0,62.2314,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_4_MO40_NTn1_NTA3_NTB1_NTC6_NTD4_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1104.26,1078.38 +gfx950,256,4096,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441462,0,30.1261,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,1140.53,1670.7 +gfx950,256,16817,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440397,0,414.8081,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x224x64_MI16x16x1_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_7_MO40_NTn1_NTA0_NTB1_NTC6_NTD7_NTM0_NEPBS14_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO4_SVW8_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1360.35,745.13 +gfx950,256,32768,4096,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440622,0,235.7907,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA1_NTB1_NTC6_NTD6_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1165.77,1458.64 +gfx950,256,8,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2370,1,5.6614,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,13.34,1682.85 +gfx950,256,40,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4183,1,6.2892,flydsl_gemm8_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,60.02,1572.18 +gfx950,256,96,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3426,1,6.277,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,144.33,1675.73 +gfx950,256,1024,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439441,0,15.9983,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x160x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB128_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_5_MO40_NTn1_NTA3_NTB0_NTC1_NTD7_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW4_SK3_SKFTR0_SKXCCM4_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,604.04,1310.86 +gfx950,256,16384,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441461,0,146.7146,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950,0.0,1053.87,1322.2 +gfx950,256,65536,4608,1024,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440614,0,496.5217,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA1_NTB2_NTC5_NTD5_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1245.62,1505.74 +gfx950,256,4,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440306,0,18.1922,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT32x16x512_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_1_MO40_NTn1_NTA7_NTB3_NTC1_NTD0_NTM0_NEPBS10_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG32_4_2,0.0,19.82,4961.05 +gfx950,256,32,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440651,0,21.7912,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x32x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_2_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS12_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB2_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,132.38,4176.27 +gfx950,256,64,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439837,0,24.8874,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x64x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB2_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT4_2_MO40_NTn1_NTA3_NTB1_NTC4_NTD1_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA4_VWB2_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,231.81,3691.31 +gfx950,256,512,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439391,0,61.1145,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT224x256x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA128_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT7_8_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO4_SVW1_SK3_SKFTR0_SKXCCM4_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,755.21,1700.48 +gfx950,256,8192,6288,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440656,0,566.6903,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT320x224x64_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA256_LBSPPB128_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT10_7_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW2_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1303.12,548.11 +gfx950,256,1,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4987,2,6.7793,flydsl_gemm5_abf16_wbf16_bf16_t16x64x64_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0021,3.25,3250.7 +gfx950,256,2,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4490,2,11.356,flydsl_gemm7_abf16_wbf16_bf16_t16x64x64_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0056,9.05,4528.29 +gfx950,256,2,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4188,2,12.0741,flydsl_gemm8_abf16_wbf16_bf16_t16x64x64_split_k2_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0049,10.03,5019.08 +gfx950,256,4,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,1484,4,11.6416,flydsl_gemm3_abf16_wbf16_bf16_t16x64x128_split_k4_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.013,20.18,5051.74 +gfx950,256,8,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,5329,1,5.4733,flydsl_gemm8_abf16_wbf16_bf16_t16x64x64_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,16.09,2034.79 +gfx950,256,16,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3787,1,7.583,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,46.46,2940.61 +gfx950,256,24,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440111,0,13.1209,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT32x32x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT1_1_MO40_NTn1_NTA7_NTB2_NTC2_NTD1_NTM0_NEPBS14_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,93.98,3955.24 +gfx950,256,24,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2905,1,15.7054,flydsl_gemm5_abf16_wbf16_bf16_t16x64x128_split_k1_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,92.54,3890.51 +gfx950,256,32,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439705,0,15.1107,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT32x32x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT1_1_MO40_NTn1_NTA4_NTB0_NTC0_NTD2_NTM0_NEPBS10_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS0_SPO1_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM4_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,124.35,3933.71 +gfx950,256,40,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3857,1,5.7057,flydsl_gemm5_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,77.19,2040.93 +gfx950,256,48,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4057,1,7.8795,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,134.14,2900.65 +gfx950,256,56,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4859,1,15.3349,flydsl_gemm8_abf16_wbf16_bf16_t32x64x64_split_k1_block_m_warp2_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0048,187.63,3429.07 +gfx950,256,56,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2314,1,17.2804,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k1_block_m_warp2_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,196.24,3578.11 +gfx950,256,64,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2415,1,16.5992,flydsl_gemm3_abf16_wbf16_bf16_t32x64x256_split_k1_block_m_warp2_block_n_warp1_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0047,226.4,3624.39 +gfx950,256,96,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4684,1,6.6815,flydsl_gemm4_abf16_wbf16_bf16_t64x64x64_split_k1_block_m_warp2_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,158.19,1875.89 +gfx950,256,128,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,6737,1,9.9229,flydsl_gemm4_abf16_wbf16_bf16_t64x64x128_split_k1_block_m_warp4_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,284.05,2443.67 +gfx950,256,256,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2441,1,25.7079,flydsl_gemm3_abf16_wbf16_bf16_t128x64x128_split_k1_block_m_warp2_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0,511.65,2212.75 +gfx950,256,256,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440054,0,29.8257,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x128x128_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT4_2_MO40_NTn1_NTA2_NTB2_NTC4_NTD5_NTM0_NEPBS12_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK0_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA4_VWB2_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,519.76,2225.86 +gfx950,256,512,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440199,0,37.7962,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x128x128_MI32x32x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA8_LPB8_LPM0_LRVW8_LWPMn1_MIAV0_MIWT2_2_MO40_NTn1_NTA2_NTB3_NTC7_NTD7_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW2_SK3_SKFTR0_SKXCCM8_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA2_VWB2_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,795.44,1858.77 +gfx950,256,1024,7168,768,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440203,0,16.7421,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_4_MO40_NTn1_NTA3_NTB3_NTC4_NTD7_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,673.41,1628.41 +gfx950,256,2048,7168,1536,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440543,0,39.7373,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA2_NTB1_NTC7_NTD4_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1134.88,1451.32 +gfx950,256,4096,7168,3584,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440614,0,166.8382,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA1_NTB2_NTC5_NTD5_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1261.42,835.9 +gfx950,256,4096,7168,4224,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440588,0,190.4156,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA1_NTB0_NTC7_NTD5_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1302.59,808.12 +gfx950,256,8192,7168,4096,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441460,0,352.1007,Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950,0.0,1366.19,690.91 +gfx950,256,4,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,1777,7,23.3609,flydsl_gemm3_abf16_wbf16_bf16_t16x256x64_split_k7_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0193,20.74,5189.67 +gfx950,256,32,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439918,0,29.115,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT96x64x128_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT6_1_MO40_NTn1_NTA0_NTB2_NTC0_NTD2_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK3_SKFTR0_SKXCCM4_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,133.11,4194.06 +gfx950,256,64,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439915,0,28.458,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT80x64x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT5_1_MO40_NTn1_NTA5_NTB2_NTC1_NTD3_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS0_SPO1_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,272.37,4326.0 +gfx950,256,512,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440294,0,78.5565,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT192x256x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA256_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT6_8_MO40_NTn1_NTA3_NTB3_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO1_SVW2_SK3_SKFTR0_SKXCCM4_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA2_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,789.35,1745.26 +gfx950,256,8192,8448,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440547,0,710.522,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x224x64_MI16x16x1_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_7_MO40_NTn1_NTA3_NTB1_NTC5_NTD5_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO4_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1396.35,530.54 +gfx950,256,8,16160,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440907,0,36.273,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT32x16x512_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT2_1_MO40_NTn1_NTA4_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM8_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_4,0.0,51.09,6397.13 +gfx950,256,40,16160,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440929,0,39.0933,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x48x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_3_MO40_NTn1_NTA4_NTB0_NTC1_NTD5_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG64_4_1,0.0,237.04,5973.81 +gfx950,256,96,16160,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440514,0,44.1248,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x96x128_MI16x16x1_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT4_3_MO40_NTn1_NTA4_NTB0_NTC1_NTD2_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM6_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA4_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,504.03,5351.84 +gfx950,256,1024,16160,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440527,0,176.7004,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_CMS_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_8_MO40_NTn1_NTA3_NTB1_NTC5_NTD2_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW8_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA8_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,1342.55,1581.47 +gfx950,256,8,20480,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440308,0,45.8147,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT96x16x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT3_1_MO40_NTn1_NTA7_NTB0_NTC3_NTD2_NTM0_NEPBS10_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS0_SPO0_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG32_4_1,0.0,51.27,6418.11 +gfx950,256,40,20480,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439915,0,49.3711,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT80x64x256_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT5_1_MO40_NTn1_NTA5_NTB2_NTC1_NTD3_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS0_SPO1_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1,0.0,237.87,5991.63 +gfx950,256,96,20480,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,439920,0,56.8313,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT96x96x128_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA256_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV1_MIWT3_3_MO40_NTn1_NTA4_NTB3_NTC3_NTD1_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM8_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,495.95,5259.6 +gfx950,256,1,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2303,8,6.3996,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k8_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.019,2.01,2009.69 +gfx950,256,2,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3645,8,6.4512,flydsl_gemm6_abf16_wbf16_bf16_t16x64x128_split_k8_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0184,3.98,1996.11 +gfx950,256,4,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2303,8,6.3371,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k8_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0204,8.11,2037.14 +gfx950,256,8,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3045,8,6.8618,flydsl_gemm5_abf16_wbf16_bf16_t16x64x128_split_k8_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0197,14.98,1890.77 +gfx950,256,16,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2303,8,6.763,flydsl_gemm4_abf16_wbf16_bf16_t16x64x128_split_k8_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0204,30.39,1937.47 +gfx950,256,24,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4112,7,8.1387,flydsl_gemm7_abf16_wbf16_bf16_t16x64x64_split_k7_block_m_warp1_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0193,37.88,1625.83 +gfx950,256,32,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3747,8,7.7027,flydsl_gemm6_abf16_wbf16_bf16_t32x64x128_split_k8_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0206,53.36,1734.61 +gfx950,256,40,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4193,7,8.6454,flydsl_gemm7_abf16_wbf16_bf16_t32x64x64_split_k7_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.02,59.43,1560.39 +gfx950,256,48,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,4193,7,9.0129,flydsl_gemm7_abf16_wbf16_bf16_t32x64x64_split_k7_block_m_warp1_block_n_warp4_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0193,68.41,1511.08 +gfx950,256,56,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2481,7,9.4115,flydsl_gemm4_abf16_wbf16_bf16_t32x64x128_split_k7_block_m_warp2_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.019,76.43,1460.79 +gfx950,256,64,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,2243,7,10.0918,flydsl_gemm3_abf16_wbf16_bf16_t32x64x128_split_k7_block_m_warp2_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0189,81.46,1375.1 +gfx950,256,96,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3458,4,11.561,flydsl_gemm4_abf16_wbf16_bf16_t32x64x64_split_k4_block_m_warp1_block_n_warp2_block_k_warp2_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.0211,106.66,1244.99 +gfx950,256,128,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,flydsl,3126,7,12.6667,flydsl_gemm3_abf16_wbf16_bf16_t64x64x128_split_k7_block_m_warp2_block_n_warp2_block_k_warp1_async_copyTrue_b_to_ldsTrue_b_preshuffleFalse_c_to_ldsFalse_gfx950,0.02,129.8,1177.06 +gfx950,256,256,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,opus,1208,4,17.2088,opus_gemm_flatmm_splitk_256x64x64x128_2x1_16x16x32_0x0x0_wgpcu1_nooob,0.0,191.08,986.35 +gfx950,256,512,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441318,0,20.4935,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x96x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA1_DTLB1_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_3_MO40_NTn1_NTA0_NTB0_NTC7_NTD4_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO1_VSn1_VWA4_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,320.91,1029.72 +gfx950,256,1024,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441293,0,27.1063,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT224x128x64_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA128_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT7_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,485.25,1083.15 +gfx950,256,2048,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441293,0,40.2166,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT224x128x64_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA128_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT7_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,654.12,1140.7 +gfx950,256,4096,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,440679,0,63.1627,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT224x256x64_MI16x16x1_SN_LDSB0_AFC0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA128_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT7_8_MO40_NTn1_NTA0_NTB0_NTC4_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS2_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB8_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,832.98,1249.24 +gfx950,256,8192,896,7168,False,torch.bfloat16,torch.bfloat16,False,False,hipblaslt,441293,0,106.4082,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT224x128x64_MI16x16x1_SN_LDSB1_AFC0_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA128_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT7_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK3_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGRO0_VSn1_VWA1_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0.0,988.9,1362.35 diff --git a/benchmarks/single_node/agentic/k3_patches/patch_flydsl_decode_to_torch.sh b/benchmarks/single_node/agentic/k3_patches/patch_flydsl_decode_to_torch.sh new file mode 100755 index 000000000..364d2fb27 --- /dev/null +++ b/benchmarks/single_node/agentic/k3_patches/patch_flydsl_decode_to_torch.sh @@ -0,0 +1,42 @@ +#!/usr/bin/env bash +# Reroute aiter FlyDSL decode-range GEMMs off FlyDSL for Kimi-K3 FP4 DSpark on MI355X/gfx950. +# +# Runs INSIDE the container (awk-edits the merged tuned-GEMM CSV in place). The +# orchestrator apply_k3_container_patches.sh invokes this as the last patch step. +# +# WHY: At decode M=72 (conc-24, M=3*conc), the K3 dense BF16 GEMMs +# (N,K) in {(1024,7168),(384,7168),(7168,512)} route to libtype=="flydsl" in +# merged_bf16_tuned_gemm.csv. FlyDSL's per-call Python launcher + split-K +# semaphore path is NOT vLLM-FULL-cudagraph-capturable, so the decode step +# falls OFF the replayed FULL graph and runs kernel-by-kernel from the host. +# That desyncs the 8 TP ranks -> the straggler starves -> the others spin in +# cross_device_reduce_2stage (the "all-reduce" symptom). conc-16 (M=48) and +# conc-48 (M=144) do NOT hit this because their FlyDSL kernelNames aren't +# catalog-valid, so they already fall back to `torch` (torch.matmul -> +# rocBLAS/hipBLASLt Cijk kernels, which ARE captured). +# +# FIX: convert the FlyDSL rows for those 3 dense (N,K) shapes at all decode-range +# M (<=192 = max_num_seqs 64 * decode_query_len 3) to `torch` -- exactly the +# capturable backend the working concurrencies already use. Prefill (M>192) +# keeps FlyDSL. +# +# Idempotent. Backs up once to .pre_flydsl_fix.bak. +set -euo pipefail + +CSV="${CSV:-/opt/aiter-local/aiter/configs/merged_bf16_tuned_gemm.csv}" +[ -f "$CSV" ] || { echo "!! tuned GEMM CSV not found at $CSV" >&2; exit 1; } + +cp -n "$CSV" "$CSV.pre_flydsl_fix.bak" || true +awk -F, -v OFS="," ' +NR==1 {print; next} +# NOTE: encode exactly like a native torch row -> splitK(col13)=0, kernelName(col15)=native. +# Leaving splitK EMPTY makes pandas.read_csv infer the whole column as float (NaN), +# which then feeds float splitK (e.g. 3.0) into aiter::_gemm_a16w16_asm (Optional[int] +# only) and crashes engine init. Keep the column integer-typed. +($11=="flydsl" && $3<=192 && ( (($4==1024||$4==384)&&$5==7168) || ($4==7168&&$5==512) )) { + $11="torch"; $12="0"; $13="0"; $15="native"; print; next +} +{print} +' "$CSV" > "$CSV.new" && mv "$CSV.new" "$CSV" +echo "flydsl rows remaining for K3 dense shapes at M<=192 (expect 0): $(awk -F, '$11=="flydsl" && $3<=192 && ( (($4==1024||$4==384)&&$5==7168) || ($4==7168&&$5==512) )' "$CSV" | wc -l)" +echo "flydsl->torch reroute applied to $CSV" diff --git a/benchmarks/single_node/agentic/k3_patches/patch_fp8_prefill.py b/benchmarks/single_node/agentic/k3_patches/patch_fp8_prefill.py new file mode 100644 index 000000000..2eae34936 --- /dev/null +++ b/benchmarks/single_node/agentic/k3_patches/patch_fp8_prefill.py @@ -0,0 +1,60 @@ +F="/usr/local/lib/python3.12/dist-packages/vllm/v1/attention/backends/mla/rocm_aiter_mla.py" +s=open(F).read(); orig=s + +# --- A) relax the gate (2 occurrences) --- +g_old="_fp8_mla_prefill_supported() and self.num_heads % 16 == 0" +g_new="_fp8_mla_prefill_supported() and (self.num_heads % 16 == 0 or 0 < self.num_heads < 16)" +n=s.count(g_old); assert n>=1, "gate anchor not found"; s=s.replace(g_old,g_new) + +# --- B) pad q/k/v to 16 in _mla_fp8_prefill_attn --- +b_old=""" fp8_dtype = current_platform.fp8_dtype() + total_q = q.shape[0] + nhead = self.num_heads + v_head_dim = self.v_head_dim + tile_q = _FP8_PREFILL_TILE_Q""" +b_new=""" fp8_dtype = current_platform.fp8_dtype() + total_q = q.shape[0] + # PATCH(fp8-prefill-pad): PS asm prefill + mla_reduce_v1 need 16-aligned + # heads and the metadata is built for _num_attention_heads=16. K3 has 12 + # heads/rank -> replicate-pad q/k/v to 16 (heads independent over shared + # KV, exact like decode); output sliced back to the real head count. + _real_nhead = self.num_heads + _pad16 = _real_nhead < 16 + if _pad16: + q = AiterMLAHelper.get_mla_padded_q(_real_nhead, q) + k = AiterMLAHelper.get_mla_padded_q(_real_nhead, k) + v = AiterMLAHelper.get_mla_padded_q(_real_nhead, v) + nhead = 16 if _pad16 else self.num_heads + v_head_dim = self.v_head_dim + tile_q = _FP8_PREFILL_TILE_Q""" +assert b_old in s, "nhead anchor not found"; s=s.replace(b_old,b_new,1) + +# --- C) out_3d: padded case needs its own buffer (can't alias real-head out) --- +c_old=" out_3d = out.view(total_q, nhead, v_head_dim)" +c_new=""" if _pad16: + out_3d = torch.empty( + total_q, nhead, v_head_dim, dtype=out.dtype, device=out.device + ) + else: + out_3d = out.view(total_q, nhead, v_head_dim)""" +assert c_old in s, "out_3d anchor not found"; s=s.replace(c_old,c_new,1) + +# --- D) copy padded output back into the caller's real-head buffer --- +d_old=""" 0, + out_3d, + final_lse, + )""" +d_new=""" 0, + out_3d, + final_lse, + ) + + if _pad16: + out.view(total_q, _real_nhead, v_head_dim).copy_( + out_3d[:, :_real_nhead, :] + )""" +assert d_old in s, "reduce anchor not found"; s=s.replace(d_old,d_new,1) + +assert s!=orig +open(F,"w").write(s) +print("fp8 prefill pad patch applied; gate replaced:",n,"; markers:",s.count("PATCH(fp8-prefill-pad)")) diff --git a/benchmarks/single_node/agentic/k3_patches/patch_fp8asm.py b/benchmarks/single_node/agentic/k3_patches/patch_fp8asm.py new file mode 100644 index 000000000..331955a0a --- /dev/null +++ b/benchmarks/single_node/agentic/k3_patches/patch_fp8asm.py @@ -0,0 +1,53 @@ +F="/usr/local/lib/python3.12/dist-packages/vllm/v1/attention/backends/mla/rocm_aiter_mla.py" +s=open(F).read(); orig=s +a=''' @staticmethod + def get_mla_padded_q(num_heads: int, q: torch.Tensor) -> torch.Tensor: + return ( + q + if num_heads >= AiterMLAHelper._AITER_MIN_MLA_HEADS + else q.repeat_interleave( + AiterMLAHelper._AITER_MIN_MLA_HEADS // num_heads, dim=1 + ) + )''' +b=''' @staticmethod + def get_mla_padded_q(num_heads: int, q: torch.Tensor) -> torch.Tensor: + MIN = AiterMLAHelper._AITER_MIN_MLA_HEADS + if num_heads >= MIN: + return q + if MIN % num_heads == 0: + return q.repeat_interleave(MIN // num_heads, dim=1) + # PATCH(fp8-asm): non-divisor head counts (K3=12) -> replicate-append + # real heads to 16 (zero-pad produced garbage output). + return torch.cat([q, q[:, : MIN - num_heads, :]], dim=1)''' +assert a in s, "padded_q anchor"; s=s.replace(a,b) +c=''' @staticmethod + def get_mla_unpadded_o(num_heads: int, o: torch.Tensor) -> torch.Tensor: + return ( + o + if num_heads >= AiterMLAHelper._AITER_MIN_MLA_HEADS + else o[:, :: AiterMLAHelper._AITER_MIN_MLA_HEADS // num_heads, :] + )''' +d=''' @staticmethod + def get_mla_unpadded_o(num_heads: int, o: torch.Tensor) -> torch.Tensor: + MIN = AiterMLAHelper._AITER_MIN_MLA_HEADS + if num_heads >= MIN: + return o + if MIN % num_heads == 0: + return o[:, :: MIN // num_heads, :] + # PATCH(fp8-asm): drop appended heads; first num_heads are the real ones. + return o[:, :num_heads, :]''' +assert c in s, "unpadded_o anchor"; s=s.replace(c,d) +e=''' @staticmethod + def use_gluon_decode(num_heads: int, max_qo_len: int) -> bool: + return num_heads < AiterMLAHelper._AITER_MIN_MLA_HEADS and max_qo_len == 1''' +f=''' @staticmethod + def use_gluon_decode(num_heads: int, max_qo_len: int) -> bool: + # PATCH(fp8-asm): route single-token decode to ASM 576/512 (append-pad + # to 16). Keep gluon only for multi-token verify (no asm kernel there). + if max_qo_len == 1: + return False + return num_heads < AiterMLAHelper._AITER_MIN_MLA_HEADS and max_qo_len == 1''' +assert e in s, "use_gluon_decode anchor"; s=s.replace(e,f) +assert s!=orig +open(F,"w").write(s) +print("patched OK; markers:", s.count("PATCH(fp8-asm)")) diff --git a/benchmarks/single_node/agentic/k3_patches/patch_ps_metadata16.py b/benchmarks/single_node/agentic/k3_patches/patch_ps_metadata16.py new file mode 100644 index 000000000..77b82664f --- /dev/null +++ b/benchmarks/single_node/agentic/k3_patches/patch_ps_metadata16.py @@ -0,0 +1,9 @@ +F="/usr/local/lib/python3.12/dist-packages/vllm/v1/attention/backends/mla/rocm_aiter_mla.py" +s=open(F).read(); orig=s +old=" num_head_k = self.num_heads" +new=" num_head_k = max(16, self.num_heads) # PATCH(fp8-prefill-pad): 16-head PS metadata to match padded q/k/v (gcd(16,256)=16 -> ~960 tiles, not 4032)" +n=s.count(old) +assert n==2, f"expected 2 occurrences, found {n}" +s=s.replace(old,new) +open(F,"w").write(s) +print("PS metadata num_head_k -> max(16,num_heads) at both sites; replaced:",n) diff --git a/benchmarks/single_node/agentic/k3_patches/patch_skip_k3_fp8_ps.py b/benchmarks/single_node/agentic/k3_patches/patch_skip_k3_fp8_ps.py new file mode 100644 index 000000000..044a389ce --- /dev/null +++ b/benchmarks/single_node/agentic/k3_patches/patch_skip_k3_fp8_ps.py @@ -0,0 +1,127 @@ +"""Skip FP8 ASM MLA prefill PS workspace on Kimi-K3 (fused FA prefill path).""" +import re + +F = "/usr/local/lib/python3.12/dist-packages/vllm/v1/attention/backends/mla/rocm_aiter_mla.py" +s = open(F).read() +orig = s + +MARKER = "PATCH(skip-k3-fp8-ps)" +if MARKER in s and "_model_uses_fused_mla_prefill" in s: + print(f"{MARKER} v2 already applied") + raise SystemExit(0) + +helper_v2 = ''' +# PATCH(skip-k3-fp8-ps): Kimi-K3 prefill uses fused FA, not forward_mha FP8 ASM. +_FUSED_MLA_PREFILL_MODEL_TYPES = frozenset({"kimi_k3", "kimi_k3_mtp"}) +_FUSED_MLA_PREFILL_ARCHITECTURES = frozenset({ + "KimiK3ForConditionalGeneration", + "KimiK3MTPModel", +}) + + +def _model_uses_fused_mla_prefill(vllm_config=None) -> bool: + if vllm_config is None: + from vllm.config import get_current_vllm_config_or_none + + vllm_config = get_current_vllm_config_or_none() + if vllm_config is None or vllm_config.model_config is None: + return False + + model_config = vllm_config.model_config + hf_config = model_config.hf_config + if getattr(hf_config, "model_type", None) in _FUSED_MLA_PREFILL_MODEL_TYPES: + return True + + architectures = getattr(hf_config, "architectures", None) or [] + if any(arch in _FUSED_MLA_PREFILL_ARCHITECTURES for arch in architectures): + return True + + hf_text_config = model_config.hf_text_config + if getattr(hf_text_config, "model_type", None) in _FUSED_MLA_PREFILL_MODEL_TYPES: + return True + + return False + + +def _uses_aiter_fp8_asm_mla_prefill( + *, + kv_cache_dtype: str, + num_heads: int, + vllm_config=None, +) -> bool: + from vllm.utils.torch_utils import is_quantized_kv_cache + + if not _fp8_mla_prefill_supported(): + return False + if not is_quantized_kv_cache(kv_cache_dtype): + return False + if not (num_heads % 16 == 0 or 0 < num_heads < 16): + return False + + if vllm_config is None: + from vllm.config import get_current_vllm_config_or_none + + vllm_config = get_current_vllm_config_or_none() + + if _model_uses_fused_mla_prefill(vllm_config): + hf_config = ( + vllm_config.model_config.hf_config + if vllm_config is not None and vllm_config.model_config is not None + else None + ) + model_label = getattr(hf_config, "model_type", None) or ( + getattr(hf_config, "architectures", ["Kimi-K3"])[0] + if hf_config is not None + else "Kimi-K3" + ) + logger.info_once( + "Skipping FP8 ASM MLA prefill workspace for %s: prefill uses " + "the fused MLA path (forward_mha is not called).", + model_label, + ) + return False + + return True +''' + +if MARKER in s: + s = re.sub( + rf"# {re.escape(MARKER)}:.*?^class AiterMLABackend", + helper_v2 + "\n\nclass AiterMLABackend", + s, + count=1, + flags=re.MULTILINE | re.DOTALL, + ) +else: + insert_after = " return True\n\n\nclass AiterMLABackend" + assert insert_after in s, "insert anchor (AiterMLABackend) not found" + s = s.replace(insert_after, " return True\n" + helper_v2 + "\n\nclass AiterMLABackend", 1) + + builder_old = """ self._fp8_prefill_enabled = ( + _fp8_mla_prefill_supported() and (self.num_heads % 16 == 0 or 0 < self.num_heads < 16) + )""" + builder_new = f""" self._fp8_prefill_enabled = _uses_aiter_fp8_asm_mla_prefill( + kv_cache_dtype=kv_cache_dtype_str, + num_heads=self.num_heads, + vllm_config=vllm_config, + ) # {MARKER}""" + assert builder_old in s, "builder gate not found" + s = s.replace(builder_old, builder_new, 1) + + impl_old = """ self._fp8_prefill_enabled = ( + _fp8_mla_prefill_supported() and (self.num_heads % 16 == 0 or 0 < self.num_heads < 16) + ) + if self._fp8_prefill_enabled: + from aiter import mla_prefill_ps_asm_fwd, mla_reduce_v1""" + impl_new = f""" self._fp8_prefill_enabled = _uses_aiter_fp8_asm_mla_prefill( + kv_cache_dtype=kv_cache_dtype, + num_heads=num_heads, + ) # {MARKER} + if self._fp8_prefill_enabled: + from aiter import mla_prefill_ps_asm_fwd, mla_reduce_v1""" + assert impl_old in s, "impl gate not found" + s = s.replace(impl_old, impl_new, 1) + +assert s != orig +open(F, "w").write(s) +print(f"{MARKER} v2 applied") diff --git a/benchmarks/single_node/agentic/k3_patches/patch_wvsplitk.py b/benchmarks/single_node/agentic/k3_patches/patch_wvsplitk.py new file mode 100644 index 000000000..c834c01c9 --- /dev/null +++ b/benchmarks/single_node/agentic/k3_patches/patch_wvsplitk.py @@ -0,0 +1,24 @@ +F="/usr/local/lib/python3.12/dist-packages/vllm/model_executor/layers/utils.py" +s=open(F).read(); orig=s +# guard 1: wvSplitKrc branch +a=''' if use_skinny_reduce_counting: + return ops.wvSplitKrc(x, weight, cu_count, bias)''' +b=''' if use_skinny_reduce_counting: + # PATCH(vLLM #50618): skinny GEMMs index operands linearly and do not + # accept strides -> OOB read on strided activations (the cudagraph + # capture fault). Force contiguous. + if not x.is_contiguous(): + x = x.contiguous() + return ops.wvSplitKrc(x, weight, cu_count, bias)''' +assert a in s, "wvSplitKrc anchor"; s=s.replace(a,b,1) +# guard 2: wvSplitK branch +c=''' if use_skinny: + x_view = x.reshape(-1, x.size(-1))''' +d=''' if use_skinny: + x_view = x.reshape(-1, x.size(-1)) + # PATCH(vLLM #50618): wvSplitK indexes linearly; strided activation -> OOB. + if not x_view.is_contiguous(): + x_view = x_view.contiguous()''' +assert c in s, "wvSplitK anchor"; s=s.replace(c,d,1) +assert s!=orig; open(F,"w").write(s) +print("wvSplitK OOB fix applied; markers:", s.count("PATCH(vLLM #50618)")) diff --git a/benchmarks/single_node/agentic/kimik3_fp4_mi355x_vllm.sh b/benchmarks/single_node/agentic/kimik3_fp4_mi355x_vllm.sh new file mode 100755 index 000000000..7e37f2578 --- /dev/null +++ b/benchmarks/single_node/agentic/kimik3_fp4_mi355x_vllm.sh @@ -0,0 +1,160 @@ +#!/usr/bin/env bash +set -eo pipefail +set -x + +# Agentic trace-replay benchmark for Kimi-K3 on AMD MI355X (gfx950) via vLLM. +# ROCm adaptation of benchmarks/single_node/agentic/kimik3_fp4_b300_vllm.sh: +# same AIPerf agentic harness + trace dataset, but the serve flags swap the +# NVIDIA-only bits (FlashInfer allreduce, fastsafetensors, Rust frontend, +# FLASHINFER MLA prefill) for the MI355X day-0 recipe (AITER). +# +# Attention: ROCM_AITER_MLA (asm persistent) — the fast path for K3's 12 heads/ +# rank (TP8). REQUIRES a vLLM image with the asm-MLA non-divisor head support: +# * vLLM #50578 (asm decode pad-to-16 for non-divisor small head counts) +# * vLLM PR-A (fp8 asm MLA *prefill* pad-to-16 + 16-head PS metadata) +# On stock vLLM these are absent and the asm path breaks for 12 heads. The asm +# prefill also needs ROCm/aiter #4452 (64-bit paged-KV offsets) for >4GB KV. +# +# Validated base image (carries the vLLM patches + AITER build above): +# vllm/vllm-openai-rocm:nightly-cb8104839c141609d99f1254459ef3a4f1bd4263 +# Pinned in configs/amd-master.yaml (kimik3-fp4-mi355x-vllm-agentic). +# +# Required env vars: MODEL, TP, CONC, KV_OFFLOADING, TOTAL_CPU_DRAM_GB, RESULT_DIR, DURATION +# Optional: MAX_MODEL_LEN (replay --max-context-length), MODEL_PATH (pre-staged weights) + +source "$(dirname "$0")/../../benchmark_lib.sh" + +check_env_vars MODEL TP CONC KV_OFFLOADING TOTAL_CPU_DRAM_GB RESULT_DIR DURATION + +if [ "$TP" -ne 8 ]; then + echo "Error: Kimi-K3 on MI355X requires TP=8 (~1.5 TB checkpoint), got TP='$TP'" >&2 + exit 1 +fi + +# ---- Bootstrap the container from the pinned base image ---------------------- +# The image pinned in configs/amd-master.yaml is the STOCK ROCm vLLM nightly +# (cb8104839c...); this idempotently applies the aiter rebuild + tuned GEMM CSV +# + triton 3.7.0 + vLLM ASM patches it needs (see apply_k3_container_patches.sh). +# No-op once markers are present. Set SKIP_K3_BOOTSTRAP=1 for a pre-baked image. +RECIPE_DIR="$(cd "$(dirname "$0")" && pwd)" +if [ "${SKIP_K3_BOOTSTRAP:-0}" != "1" ]; then + bash "$RECIPE_DIR/apply_k3_container_patches.sh" +fi + +# Resolve weights: MODEL_PATH (pre-staged) else HF cache. +if [[ -n "${MODEL_PATH:-}" ]]; then + if [[ ! -d "$MODEL_PATH" || -z "$(ls -A "$MODEL_PATH" 2>/dev/null)" ]]; then + hf download "$MODEL" --local-dir "$MODEL_PATH" + fi +else + if [[ "$MODEL" != /* ]]; then hf download "$MODEL"; fi + export MODEL_PATH="$MODEL" +fi +if [ -n "$ROCR_VISIBLE_DEVICES" ]; then export HIP_VISIBLE_DEVICES="$ROCR_VISIBLE_DEVICES"; fi + +# ---- MI355X day-0 serving environment (AITER) ------------------------------- +export VLLM_ROCM_USE_AITER=1 +export GPU_ARCHS=gfx950 +export AITER_SITUV2_A8W4=1 +export AITER_BF16_FP8_MOE_BOUND=0 +export VLLM_USE_BREAKABLE_CUDAGRAPH=0 +export SAFETENSORS_FAST_GPU=1 +export VLLM_ENGINE_READY_TIMEOUT_S=3600 +# Outlast the AIPerf keep-alive pool so an inter-turn idle gap can't race a +# socket close into a warmup-abort (see the B300 recipe's note). +export VLLM_HTTP_TIMEOUT_KEEP_ALIVE=900 +export AIPERF_HTTP_TCP_USER_TIMEOUT=900000 + +# ---- Resolve traces + install AIPerf (isolated venv) ------------------------ +resolve_trace_source +install_agentic_deps + +SERVER_LOG="$RESULT_DIR/server.log" +mkdir -p "$RESULT_DIR" + +# ---- KV offloading ---------------------------------------------------------- +OFFLOAD_ARGS=() +case "${KV_OFFLOAD_BACKEND:-}" in + "") + require_agentic_kv_offload_none + ;; + vllm-simple) + require_agentic_kv_offload_backend vllm-simple + CPU_BYTES_PER_RANK=$(( TOTAL_CPU_DRAM_GB * 1000 * 1000 * 1000 / TP )) + export PYTHONHASHSEED=42 + OFFLOAD_ARGS=( + --kv-transfer-config + "{\"kv_connector\":\"SimpleCPUOffloadConnector\",\"kv_role\":\"kv_both\",\"kv_connector_extra_config\":{\"cpu_bytes_to_use_per_rank\":${CPU_BYTES_PER_RANK},\"lazy_offload\":false}}" + ) + ;; + *) + echo "Error: unsupported KV_OFFLOAD_BACKEND='$KV_OFFLOAD_BACKEND'" >&2 + exit 1 + ;; +esac + +# Fixed max-num-seqs (not 2*CONC): 64 is the validated headroom for the ASM MLA +# fp8 prefill at uncapped (~1M) context — higher values re-introduce the +# activation-arena OOM that capped long context. Override via MAX_NUM_SEQS. +MAX_NUM_SEQS="${MAX_NUM_SEQS:-64}" + +# --- fp8 KV on K3 (dense MLA) via the ASM persistent MLA path ------------------ +# K3's attention is DENSE MLA with 12 heads/rank at TP8 (a non-divisor of 16). +# * The AITER *asm persistent* MLA does batched fp8 decode once the query +# heads are padded 12->16 (vLLM #50578) and, for prefill, the fp8 PS asm +# kernel + 16-head PS metadata (vLLM PR-A). This is faster than TRITON_MLA +# and, with ROCm/aiter #4452 (64-bit paged-KV offsets), serves the FULL +# native (~1M) context uncapped with no OOM (validated 470k/590k prefills). +# * Do NOT set VLLM_ROCM_SHUFFLE_KV_CACHE_LAYOUT (asm reads the native layout). +# KV_CACHE_DTYPE=auto falls back to bf16 KV on the same asm path. +KVDTYPE_ARGS=( + --kv-cache-dtype "${KV_CACHE_DTYPE:-fp8}" + --attention-backend "${ATTENTION_BACKEND:-ROCM_AITER_MLA}" +) + +export VLLM_ROCM_USE_AITER_MOE=1 + +echo "Starting vllm server (MI355X/AITER, DSV4-agentic-derived config)..." +VLLM_CMD=( + vllm serve "$MODEL_PATH" --served-model-name "$MODEL" + --host 0.0.0.0 --port "$PORT" + --tensor-parallel-size "$TP" + --async-scheduling + --distributed-executor-backend mp + --gpu-memory-utilization 0.95 + --max-num-seqs "$MAX_NUM_SEQS" + --max-num-batched-tokens 4096 + --trust-remote-code + --load-format auto + --moe-backend auto + --mm-encoder-tp-mode data + "${KVDTYPE_ARGS[@]}" + --compilation-config '{"mode":3,"cudagraph_mode":"FULL_AND_PIECEWISE"}' + --enable-prefix-caching + # native hybrid KV (MLA + KDA) — no padding; the fix for the capture fault. + --no-disable-hybrid-kv-cache-manager + --reasoning-parser kimi_k3 + --tool-call-parser kimi_k3 + --enable-auto-tool-choice + --disable-uvicorn-access-log + "${OFFLOAD_ARGS[@]}" +) +# Optional per-seq context cap (default: model native, like the DSV4 recipe). +if [ -n "${MAX_MODEL_LEN:-}" ]; then VLLM_CMD+=(--max-model-len "$MAX_MODEL_LEN"); fi +printf '%q ' "${VLLM_CMD[@]}" | tee "$RESULT_DIR/vllm_command.txt"; printf '\n' | tee -a "$RESULT_DIR/vllm_command.txt" +"${VLLM_CMD[@]}" > "$SERVER_LOG" 2>&1 & +SERVER_PID=$! +echo "Server PID: $SERVER_PID" + +wait_for_server_ready --port "$PORT" --server-log "$SERVER_LOG" --server-pid "$SERVER_PID" + +build_replay_cmd "$RESULT_DIR" +run_agentic_replay_and_write_outputs "$RESULT_DIR" + +# cleanup: free the GPU (orphaned TP workers otherwise hold VRAM) +[[ -n "${SERVER_PID:-}" ]] && kill "$SERVER_PID" 2>/dev/null || true +pkill -9 -f "/usr/local/bin/vll[m]" 2>/dev/null || true +pkill -9 -f "EngineCore" 2>/dev/null || true +pkill -9 -f "multiprocessing.spawn" 2>/dev/null || true +for _ in $(seq 1 30); do pgrep -f "EngineCore|multiprocessing.spawn" >/dev/null 2>&1 || break; sleep 2; done +set +x diff --git a/benchmarks/single_node/agentic/kimik3_fp4_mi355x_vllm_mtp.sh b/benchmarks/single_node/agentic/kimik3_fp4_mi355x_vllm_mtp.sh new file mode 100755 index 000000000..ebf31e732 --- /dev/null +++ b/benchmarks/single_node/agentic/kimik3_fp4_mi355x_vllm_mtp.sh @@ -0,0 +1,300 @@ +#!/usr/bin/env bash +set -eo pipefail +set -x + +# Agentic trace-replay benchmark for Kimi-K3 (FP4) on AMD MI355X (gfx950) via +# vLLM with DSpark speculative decoding at level 2, probabilistic drafting, and +# synthetic acceptance pinned to the committed golden AL. +# +# MTP sibling of kimik3_fp4_mi355x_vllm.sh and the ROCm counterpart of +# kimik3_fp4_b300_vllm_mtp.sh. The target-server config is identical to the base +# MI355X recipe; the deltas are the DSpark --speculative-config, the KV-cache pin +# + spec-aware cudagraph capture sizes, and the fp8 ASM-MLA env. This is the +# variant whose numbers are apples-to-apples with the B300-MTP baseline +# (synthetic acceptance vs synthetic acceptance -- see the SPEC_CONFIG block). +# +# Attention: ROCM_AITER_MLA (asm persistent) for BOTH target and draft. On ROCm, +# the DSpark draft's semi-autoregressive parallel drafting is forced CAUSAL +# (dflash_config.causal=true, applied by apply_k3_container_patches.sh) so the +# draft no longer needs TRITON_MLA's non-causal path and runs on the same fp8 asm +# path as the target. This requires the vLLM ASM patches + DSpark fp8-asm layer +# baked into the validated base image below. +# +# Validated base image (carries the vLLM patches + AITER build + DSpark layer): +# vllm/vllm-openai-rocm:nightly-cb8104839c141609d99f1254459ef3a4f1bd4263 +# reproduced from the pinned nightly by apply_k3_container_patches.sh. +# Pinned in configs/amd-master.yaml (kimik3-fp4-mi355x-vllm-agentic-mtp). +# +# Required env vars: MODEL, TP, CONC, KV_OFFLOADING, TOTAL_CPU_DRAM_GB, RESULT_DIR, DURATION +# Optional: MODEL_PATH (pre-staged target), DRAFT_MODEL_PATH (pre-staged draft), +# EVAL_ONLY (true -> real block verification for accuracy runs) + +source "$(dirname "$0")/../../benchmark_lib.sh" + +check_env_vars MODEL TP CONC KV_OFFLOADING TOTAL_CPU_DRAM_GB RESULT_DIR DURATION + +if [ "$TP" -ne 8 ]; then + echo "Error: Kimi-K3 on MI355X requires TP=8 (~1.5 TB checkpoint), got TP='$TP'" >&2 + exit 1 +fi + +# ---- Bootstrap the container from the pinned base image ---------------------- +# The image pinned in configs/amd-master.yaml is the STOCK ROCm vLLM nightly +# (cb8104839c...); this idempotently turns it into the exact k3-dspark-benchmark +# container (aiter @ 55dbc4f47 rebuild, tuned GEMM CSV, triton 3.7.0, 5 vLLM ASM +# patches, DSpark fp8-asm layer, FlyDSL->torch reroute). No-op once markers are +# present, so re-runs / pre-patched images cost only the grep verify. Set +# SKIP_K3_BOOTSTRAP=1 to skip (e.g. when serving a pre-baked patched image). +RECIPE_DIR="$(cd "$(dirname "$0")" && pwd)" +if [ "${SKIP_K3_BOOTSTRAP:-0}" != "1" ]; then + bash "$RECIPE_DIR/apply_k3_container_patches.sh" +fi + +DRAFT_MODEL="${DRAFT_MODEL:-Inferact/Kimi-K3-DSpark}" + +# Resolve target + draft weights (pre-staged else HF cache). +if [[ -n "${MODEL_PATH:-}" ]]; then + if [[ ! -d "$MODEL_PATH" || -z "$(ls -A "$MODEL_PATH" 2>/dev/null)" ]]; then + hf download "$MODEL" --local-dir "$MODEL_PATH" + fi + DRAFT_MODEL_PATH="${DRAFT_MODEL_PATH:-${WRITABLE_MODELS_DIR:-/data/models}/${DRAFT_MODEL##*/}}" + if [[ ! -d "$DRAFT_MODEL_PATH" || -z "$(ls -A "$DRAFT_MODEL_PATH" 2>/dev/null)" ]]; then + hf download "$DRAFT_MODEL" --local-dir "$DRAFT_MODEL_PATH" + fi +else + if [[ "$MODEL" != /* ]]; then hf download "$MODEL"; fi + export MODEL_PATH="$MODEL" + DRAFT_MODEL_PATH="${DRAFT_MODEL_PATH:-$DRAFT_MODEL}" + if [[ "$DRAFT_MODEL_PATH" != /* ]]; then hf download "$DRAFT_MODEL"; fi +fi +if [ -n "$ROCR_VISIBLE_DEVICES" ]; then export HIP_VISIBLE_DEVICES="$ROCR_VISIBLE_DEVICES"; fi + +# Force the DSpark draft causal (non-causal parallel drafting is a cudagraph OOB +# source on ROCm; with causal=true the draft runs the fp8 asm path like the +# target). Idempotent write — the draft may live anywhere the harness staged it, +# so set it here rather than relying on the bootstrap's /dev/shm-scoped step. +if [ -f "$DRAFT_MODEL_PATH/config.json" ]; then + python3 - "$DRAFT_MODEL_PATH/config.json" <<'PY' +import json, sys +f = sys.argv[1] +c = json.load(open(f)) +d = c.setdefault("dflash_config", {}) +if d.get("causal") is True: + print("draft already forced causal") +else: + d["causal"] = True + json.dump(c, open(f, "w"), indent=2) + print("forced draft dflash_config.causal=true:", f) +PY +fi + +# ---- MI355X day-0 serving environment (AITER + fp8 ASM MLA) ------------------ +export VLLM_ROCM_USE_AITER=1 +export VLLM_ROCM_USE_AITER_MOE=1 +export GPU_ARCHS=gfx950 +export AITER_SITUV2_A8W4=1 +export VLLM_ROCM_USE_AITER_MOE_SITUV2_A8W4=1 +export AITER_BF16_FP8_MOE_BOUND=0 +export VLLM_USE_BREAKABLE_CUDAGRAPH=0 +export SAFETENSORS_FAST_GPU=1 +export VLLM_ENGINE_READY_TIMEOUT_S=3600 +export VLLM_HTTP_TIMEOUT_KEEP_ALIVE=900 +export AIPERF_HTTP_TCP_USER_TIMEOUT=900000 +# fp8 ASM persistent MLA on the native (unshuffled) KV layout. Do NOT set +# VLLM_ROCM_SHUFFLE_KV_CACHE_LAYOUT. Baked DSpark layer reads this to route the +# 12-head fp8 spec verify to the asm q-row-fold. +export VLLM_ROCM_AITER_MLA_ASM_PADDING=asm +# Merged tuned BF16 GEMM table installed by apply_k3_container_patches.sh. +MERGED_GEMM_CSV="${MERGED_GEMM_CSV:-/opt/aiter-local/aiter/configs/merged_bf16_tuned_gemm.csv}" +if [ -z "${AITER_CONFIG_GEMM_BF16:-}" ] && [ -f "$MERGED_GEMM_CSV" ]; then + export AITER_CONFIG_GEMM_BF16="$MERGED_GEMM_CSV" +fi +# Guard: the tuned BF16 GEMM table is a large perf lever. Without it the dense +# GEMMs fall back to an untuned path and throughput drops materially -- and it does +# so SILENTLY. Warn loudly if the CSV is missing (e.g. a pre-baked image that never +# ran apply_k3_container_patches.sh) so an untuned run is never mistaken for a valid +# number. Set REQUIRE_TUNED_GEMM=1 to make a missing CSV fatal instead of a warning. +if [ -z "${AITER_CONFIG_GEMM_BF16:-}" ] || [ ! -f "${AITER_CONFIG_GEMM_BF16:-/nonexistent}" ]; then + echo "############################################################################" >&2 + echo "WARNING: tuned BF16 GEMM CSV not found (looked for '$MERGED_GEMM_CSV')." >&2 + echo " AITER_CONFIG_GEMM_BF16 is unset -> dense GEMMs run UNTUNED and" >&2 + echo " throughput will be materially lower than the published recipe." >&2 + echo " Fix: ensure apply_k3_container_patches.sh ran (it installs" >&2 + echo " k3_patches/kimik3_bf16_tuned_gemm.csv), or point AITER_CONFIG_GEMM_BF16">&2 + echo " at the merged CSV. Set REQUIRE_TUNED_GEMM=1 to make this fatal." >&2 + echo "############################################################################" >&2 + if [ "${REQUIRE_TUNED_GEMM:-0}" = "1" ]; then + echo "REQUIRE_TUNED_GEMM=1 -> aborting to avoid an untuned (invalid) run." >&2 + exit 1 + fi +else + echo "tuned BF16 GEMM CSV active: $AITER_CONFIG_GEMM_BF16 ($(wc -l < "$AITER_CONFIG_GEMM_BF16" 2>/dev/null) rows)" +fi + +# ---- Resolve traces + install AIPerf (isolated venv) ------------------------ +resolve_trace_source +install_agentic_deps + +SERVER_LOG="$RESULT_DIR/server.log" +mkdir -p "$RESULT_DIR" + +# ---- KV offloading ---------------------------------------------------------- +# MI355X analogue of the B300-MTP offload_mode="on" arm. On the long-ISL agentic +# corpus the on-device KV pin saturates at conc>=8 (peak ~97-100%), the shared +# prefix is evicted, prefix-cache hit collapses and TTFT/throughput fall off a +# cliff. Spilling evicted KV to host keeps the prefix resident. Enabled only when +# the harness sets KV_OFFLOAD_BACKEND (no-op otherwise, so low-conc is unaffected). +OFFLOAD_ARGS=() +case "${KV_OFFLOAD_BACKEND:-}" in + "") + require_agentic_kv_offload_none + ;; + native) + require_agentic_kv_offload_backend native + OFFLOAD_ARGS=( + --kv-offloading-size "${KV_OFFLOADING_SIZE:-$TOTAL_CPU_DRAM_GB}" + --kv-offloading-backend native + ) + ;; + vllm-simple) + require_agentic_kv_offload_backend vllm-simple + CPU_BYTES_PER_RANK=$(( TOTAL_CPU_DRAM_GB * 1000 * 1000 * 1000 / TP )) + export PYTHONHASHSEED=42 + OFFLOAD_ARGS=( + --kv-transfer-config + "{\"kv_connector\":\"SimpleCPUOffloadConnector\",\"kv_role\":\"kv_both\",\"kv_connector_extra_config\":{\"cpu_bytes_to_use_per_rank\":${CPU_BYTES_PER_RANK},\"lazy_offload\":false}}" + ) + ;; + *) + echo "Error: unsupported KV_OFFLOAD_BACKEND='$KV_OFFLOAD_BACKEND'" >&2 + exit 1 + ;; +esac + +# ---- DSpark speculative decoding ------------------------------------------- +# Probabilistic drafting at DSpark level 2 (K=2), synthetic acceptance pinned to +# the committed golden AL per the AgentX policy: "InferenceX is evaluating +# inference-system performance, not the ability to fine-tune a benchmark-specific +# speculative head" (golden_al_distribution/README.md). Synthetic-vs-synthetic is +# the only apples-to-apples comparison with the B300-MTP baseline, which pins the +# same golden AL. K=2 rather than 7: verify width is a real cost synthetic +# acceptance does not remove; K=2 clears break-even at both ends of the conc range. +NUM_SPEC_TOKENS="${NUM_SPEC_TOKENS:-2}" +TOKENS_PER_SEQ=$((1 + NUM_SPEC_TOKENS)) +# Committed golden AL at K=2 on the probabilistic curve +# (kimik3_dspark_probabilistic_sample_method_block_rejection_sample_method.yaml: +# thinking_on 2 -> 2.51). Do not mix with the greedy curve's 2.45. +SYNTHETIC_ACCEPT_LEN="${SYNTHETIC_ACCEPT_LEN:-2.51}" + +# Draft attention backend: ROCM_AITER_MLA (draft forced causal -> fp8 asm path, +# same as target). Throughput runs pin synthetic acceptance; EVAL_ONLY accuracy +# runs use real block verification (synthetic commits drafted tokens regardless +# of the target's logits, so generated text would be wrong -> eval scores 0). +DRAFT_BACKEND="${DRAFT_BACKEND:-ROCM_AITER_MLA}" +if [ "${EVAL_ONLY:-false}" = "true" ]; then + SPEC_CONFIG="{\"method\": \"dspark\", \"model\": \"$DRAFT_MODEL_PATH\", \"num_speculative_tokens\": $NUM_SPEC_TOKENS, \"attention_backend\": \"$DRAFT_BACKEND\", \"draft_sample_method\": \"probabilistic\", \"rejection_sample_method\": \"block\"}" +else + SPEC_CONFIG="{\"method\": \"dspark\", \"model\": \"$DRAFT_MODEL_PATH\", \"num_speculative_tokens\": $NUM_SPEC_TOKENS, \"attention_backend\": \"$DRAFT_BACKEND\", \"draft_sample_method\": \"probabilistic\", \"rejection_sample_method\": \"synthetic\", \"synthetic_acceptance_length\": $SYNTHETIC_ACCEPT_LEN}" +fi + +# ---- Mandated DSpark serving knobs (do not change) -------------------------- +# gpu-mem 0.95 / max-num-seqs 64 / MNBT 16384 / FULL_AND_PIECEWISE are mandated. +MAX_NUM_SEQS="${MAX_NUM_SEQS:-64}" +MNBT="${MNBT:-16384}" +GPU_MEM="${GPU_MEM:-0.95}" +KVDTYPE_ARGS=( + --kv-cache-dtype "${KV_CACHE_DTYPE:-fp8}" + --attention-backend "${ATTENTION_BACKEND:-ROCM_AITER_MLA}" +) + +# KV-cache memory pin. At gpu-mem 0.95 with TP8 K3-fp4, weights are ~201 GiB of +# the 273.6 GiB budget; vLLM's auto-sizer under-estimates the real prefill/verify +# activation peak (MNBT chunks x up-to-64 concurrent 68k-token reqs + DSpark +# verify buffers) and OOMs to "0 MB free". Prefix caching stores the ~64k prefix +# ONCE (~350k KV tokens, ~6 GiB needed); pin KV to 32 GiB (~2M tokens, >5x) so a +# large physical headroom remains for the activation peak. Do NOT touch gpu-mem. +KV_CACHE_MEMORY="${KV_CACHE_MEMORY:-34359738368}" +KVMEM_ARG=(); [ -n "$KV_CACHE_MEMORY" ] && KVMEM_ARG=(--kv-cache-memory "$KV_CACHE_MEMORY") + +# cudagraph capture sizes — pin explicitly so DSpark decode (M = TOKENS_PER_SEQ * +# conc, uniform_decode_query_len = 1 + num_spec = 3) lands on FULL decode graphs +# at every benchmark concurrency. vLLM derives the FULL/decode graph set as +# round_up(size, 3) over this ladder; the AUTO ladder leaves gaps at 12 and 36 +# (conc-4 -> 12 tok, conc-12 -> 36 tok fall to a PIECEWISE graph -> attention runs +# eager every step -> get_mla_metadata_v1 host bubble, ~75 ms ITL). Adding 12 and +# 36 gives exact FULL decode graphs. Rule: for a new conc C, add 3*C. +CAPTURE_SIZES="${CAPTURE_SIZES:-1,2,4,8,12,16,24,32,36,40,48,56,64,72,80,88,96,104,112,120,128,136,144,152,160,168,176,184,192,200,208,216,224,232,240,248,256,272,288,304,320,336,352,368,384}" +CUDAGRAPH_MODE="${CUDAGRAPH_MODE:-FULL_AND_PIECEWISE}" +COMPILATION_CONFIG="{\"mode\":3,\"cudagraph_mode\":\"$CUDAGRAPH_MODE\",\"custom_ops\":[\"+fused_rms_norm_gated\"],\"cudagraph_capture_sizes\":[$CAPTURE_SIZES]}" + +# Banner must reflect the ACTUAL spec arm: EVAL_ONLY uses block/real verify (an +# accuracy run, NOT throughput-comparable), the default uses synthetic golden AL. +# Printing "synthetic AL" unconditionally would let a block run be mistaken for a +# throughput run when checking the log -- report the real mode. +if [ "${EVAL_ONLY:-false}" = "true" ]; then + SPEC_MODE_DESC="block/REAL verify (EVAL_ONLY accuracy run -- NOT throughput-comparable)" + # Loud guard: an EVAL_ONLY run uses real block verification, so its acceptance + # length is the corpus-measured ~1.2-1.7 (not the synthetic golden 2.51) and its + # tok/s is correspondingly lower. Those numbers are for ACCURACY, not throughput. + # Warn unmistakably so an eval run's tok/s can't be reported as a recipe result. + echo "############################################################################" >&2 + echo "WARNING: EVAL_ONLY=true -> ACCURACY run (rejection_sample_method=block)." >&2 + echo " Acceptance is real-verify (~1.2-1.7), NOT the synthetic golden 2.51," >&2 + echo " so throughput / interactivity from this run are LOWER and are NOT" >&2 + echo " comparable to the published recipe. Do NOT report tok/s from an" >&2 + echo " EVAL_ONLY run. For throughput, run with EVAL_ONLY unset (default)." >&2 + echo "############################################################################" >&2 +else + SPEC_MODE_DESC="synthetic AL=$SYNTHETIC_ACCEPT_LEN (AgentX golden, throughput methodology)" +fi +echo "Starting vllm server (MI355X/AITER DSpark fp8-asm, spec=$SPEC_MODE_DESC)..." +{ set +x; } 2>/dev/null +VLLM_CMD=( + vllm serve "$MODEL_PATH" --served-model-name "$MODEL" + --host 0.0.0.0 --port "$PORT" + --tensor-parallel-size "$TP" + --async-scheduling + --distributed-executor-backend mp + --gpu-memory-utilization "$GPU_MEM" + --max-num-seqs "$MAX_NUM_SEQS" + --max-model-len 1048576 + --max-num-batched-tokens "$MNBT" + --trust-remote-code + --load-format auto + --moe-backend aiter + --mm-encoder-tp-mode data + "${KVDTYPE_ARGS[@]}" + "${KVMEM_ARG[@]}" + --speculative-config "$SPEC_CONFIG" + --compilation-config "$COMPILATION_CONFIG" + --enable-prefix-caching + --enable-prompt-tokens-details + # native hybrid KV (MLA + KDA) — no padding; the fix for the capture fault. + --no-disable-hybrid-kv-cache-manager + --reasoning-parser kimi_k3 + --tool-call-parser kimi_k3 + --enable-auto-tool-choice + --disable-uvicorn-access-log + "${OFFLOAD_ARGS[@]}" +) +printf '%q ' "${VLLM_CMD[@]}" | tee "$RESULT_DIR/vllm_command.txt"; printf '\n' | tee -a "$RESULT_DIR/vllm_command.txt" +"${VLLM_CMD[@]}" > "$SERVER_LOG" 2>&1 & +SERVER_PID=$! +echo "Server PID: $SERVER_PID" + +wait_for_server_ready --port "$PORT" --server-log "$SERVER_LOG" --server-pid "$SERVER_PID" + +if [ "${EVAL_ONLY:-false}" = "true" ]; then + run_eval --port "$PORT" +else + build_replay_cmd "$RESULT_DIR" + run_agentic_replay_and_write_outputs "$RESULT_DIR" +fi + +# cleanup: free the GPU (orphaned TP workers otherwise hold VRAM) +[[ -n "${SERVER_PID:-}" ]] && kill "$SERVER_PID" 2>/dev/null || true +pkill -9 -f "/usr/local/bin/vll[m]" 2>/dev/null || true +pkill -9 -f "EngineCore" 2>/dev/null || true +pkill -9 -f "multiprocessing.spawn" 2>/dev/null || true +for _ in $(seq 1 30); do pgrep -f "EngineCore|multiprocessing.spawn" >/dev/null 2>&1 || break; sleep 2; done +set +x diff --git a/configs/amd-master.yaml b/configs/amd-master.yaml index 2ce535a42..1e8eaa716 100644 --- a/configs/amd-master.yaml +++ b/configs/amd-master.yaml @@ -1250,6 +1250,59 @@ dsv4-fp4-mi355x-vllm-agentic-mtp: # while MTP creates two. Restore these points after the upstream hybrid # KV recovery fix lands: https://github.com/vllm-project/vllm/pull/45497 +kimik3-fp4-mi355x-vllm-agentic: + # Day-zero Kimi-K3 recipe for MI355X (gfx950), the ROCm counterpart of + # kimik3-fp4-b300-vllm-agentic. The pinned image is the STOCK ROCm vLLM + # nightly; the recipe idempotently patches it into the validated container at + # startup (aiter @ 55dbc4f47 rebuild + tuned GEMM CSV + triton 3.7.0 + the 5 + # vLLM ASM-MLA patches K3's 12-head/rank fp8 attention needs) via + # benchmarks/single_node/agentic/apply_k3_container_patches.sh. K3 ships as the + # out-of-tree vllm.models.kimi_k3 plugin, present in this nightly. + image: vllm/vllm-openai-rocm:nightly-cb8104839c141609d99f1254459ef3a4f1bd4263 + model: moonshotai/Kimi-K3 + model-prefix: kimik3 + runner: cluster:mi355x-amds + precision: fp4 + framework: vllm + multinode: false + scenarios: + # Agentic-coding only, mirroring the B300 sister's ladder so GPU-resident + # and host-DRAM-offload are directly comparable at equal concurrency. + agentic-coding: + - dram-utilization: 0.60 + search-space: + # TP8-only: a ~1.5 TB FP4 checkpoint does not fit below 8 GPUs. + - { tp: 8, ep: 1, kv-offloading: none, conc-list: [1, 2, 4, 8, 16, 24] } + - { tp: 8, ep: 1, kv-offloading: dram, kv-offload-backend: { name: native }, conc-list: [8, 16, 24] } + +kimik3-fp4-mi355x-vllm-agentic-dspark: + # DSpark speculative decoding (Inferact/Kimi-K3-DSpark draft) at level 2 with + # probabilistic drafting and synthetic acceptance pinned to the committed golden + # AL 2.51 (golden_al_distribution/ + # kimik3_dspark_probabilistic_sample_method_block_rejection_sample_method.yaml, + # thinking_on 2 -> 2.51), per the AgentX policy: draft length is the + # submission's choice, the acceptance target is not. Synthetic-vs-synthetic is + # the only apples-to-apples comparison with kimik3-fp4-b300-vllm-agentic-dspark. + # EVAL_ONLY runs switch to real block verification. spec-decoding: mtp routes to + # kimik3_fp4_mi355x_vllm_mtp.sh, which forces the draft causal (the ROCm fp8-asm + # requirement) and serves the mandated DSpark config (gpu-mem 0.95 / seqs 64 / + # MNBT 16384 / FULL_AND_PIECEWISE / KV pin 32 GiB). + image: vllm/vllm-openai-rocm:nightly-cb8104839c141609d99f1254459ef3a4f1bd4263 + model: moonshotai/Kimi-K3 + model-prefix: kimik3 + runner: cluster:mi355x-amds + precision: fp4 + framework: vllm + multinode: false + scenarios: + agentic-coding: + - dram-utilization: 0.60 + search-space: + # GPU-resident through the ladder; host-DRAM offload for conc>=8 where the + # on-device KV pin saturates and the shared prefix would otherwise evict. + - { tp: 8, ep: 1, spec-decoding: mtp, kv-offloading: none, conc-list: [1, 2, 4, 8, 16, 24] } + - { tp: 8, ep: 1, spec-decoding: mtp, kv-offloading: dram, kv-offload-backend: { name: native }, conc-list: [8, 16, 24] } + dsr1-fp4-mi355x-sglang-disagg-mtp: image: lmsysorg/sglang-rocm:v0.5.12-rocm720-mi35x-20260519 model: amd/DeepSeek-R1-0528-MXFP4-v2