[CK] blockscale split-K: pick the pipeline from the per-split loop count

[siliangchen-amd]

Summary

gemm_a8w8_blockscale_cktile picks the pipeline variant (has_hot_loop, tail_num) from the loop count of the full K. With split-K, each split only runs K / (k_batch * K_Tile) loops. When a split is short (two K loops), the selected variant does not match the loops it actually runs, and the output is silently wrong.

This PR computes the loop count per split, the same way the other CK-Tile launchers in aiter already do: ck_deepgemm, ck_tile_gemm_moe_2stages and cktile_gemm_a8w8_bpreshuffle all use k_grain = k_batch * K_Tile.

Results (MI325X, gfx942)

Public gemm_a8w8_blockscale_cktile with the default kernel (K_Tile 128), split-K output compared against split-K 0, fraction of mismatched elements:

M, N, K	splitK (splits)	loops per split	before	after
8, 256, 512	1 (2)	2	0.69	0
8, 256, 1024	2 (4)	2	0.87	0
32, 512, 2048	3 (8)	2	0.97	0

Every CK-Tile tuning candidate on gfx942 (80 instances × splitK 0–3 × 8 shapes, M from 1 to 16384, (N, K) = (128, 6144) and (256, 4096); 2560 runs), compared against a torch reference:

• Wrong results (errRatio > 0.05) drop from 80 to 0. All 80 were splits with two K loops.
• The 224 rejected candidates and every splitK = 0 result are unchanged.
• The remaining split-K results differ only by atomic-add ordering (errRatio at most 0.008).

The tuner already rejected these candidates through errRatio, so no tuned config was affected. Direct split-K calls with short splits were.

Test plan

• op_tests/test_gemm_a8w8_blockscale.py: test_splitk_correctness now asserts on the error ratio and covers three shapes with two K loops per split. Before this change the new cases fail (cktile_err=0.8716); after it, all cases pass. The existing cases pass either way.
• Candidate sweep and API check above.

Found by Hyperloom; reviewed and measured by hand.
It showed up while tuning the gfx942 table in #5839, which started from Hyperloom optimizing Qwen/Qwen3.5-397B-A17B-FP8.

[github-actions]

🏷️ CI Guide

Runs automatically on every PR:

• ✅ Pre-checks (submodule verification, code formatting)
• ✅ Aiter op tests (gfx942 + gfx950)
• ✅ Triton tests on MI35X (only when aiter/ops/triton/** or related paths are changed)

Extended tests (opt-in via labels):

Label	Tests
ci:gfx1250-ffm-triton	Run the five-shard gfx1250 FFM Triton test suite
ci:triton-300x	Run an additional Triton test job on MI300X in PRs; main branch always runs both MI35X and MI300X
multigpu	Aiter multi-GPU tests on the 8-GPU runner
ci:sglang	SGLang integration tests: DeepSeek-R1-MXFP4 accuracy, Qwen 3.5 accuracy
ci:atom	ATOM benchmark: DeepSeek-R1-0528, GPT-OSS-120B
ci:atom_full	ATOM accuracy suite for PR and main models from ATOM models_accuracy.json
ci:vllm	vLLM benchmark: GPT-OSS-120B, DeepSeek-R1-0528, Kimi-K2.5
ci:all	All standard extended tests (excludes ci:atom_full)

Only add ci:atom_full for FlyDSL or Triton upgrades.
Add labels via the sidebar or gh pr edit 5847 --add-label <label>

PR title tags & labels:
Component tags ([Triton/Gluon], [HIP], [CK], [ASM], ...) are added to the PR title and as PR labels automatically from the changed files and re-synced on every push — change-type tags like [fix]/[Perf], op tags like [MLA], and human labels (ci:*) are left untouched. Add the no-auto-title label to opt this PR out.

[yifehuan]

LGTM