Fix blockwise scale indexing under split-K (#3537)

[amacharla15]

Fixes #3537

Under serial split-K, the blockwise scale lookup used the threadblock-local
K-tile counter (k_iter_idx) with no contribution from the slice's position in
the global K range, so every slice applied slice-0's scale factors to its own
partial products.

Reproduction

M=N=256, K=512, scale block 128 (4 K-blocks), FP8 e4m3 with all-ones A and B,
kSplitKSerial = true, and powers-of-two scales per K-block:

scale_A row 0: 1, 2, 4, 8 scale_B row 0: 1, 2, 4, 8
scale_A row 1: 2, 4, 8, 16 scale_B row 1: 2, 4, 8, 16

For output element (0,0) the correct value is
128 × (1·1 + 2·2 + 4·4 + 8·8) = 128 × 85 = 10880.

split_k_slices	before	after
1	10880	10880
2	1280	10880
4	not measured	10880

The observed 1280 is exactly 2 × 128 × (1·1 + 2·2) — slice 1 re-applied
K-blocks 0 and 1 to data from blocks 2 and 3. All 65536 output elements were
affected. can_implement, initialize and the kernel launch all returned
kSuccess, so the mis-scaled result was reported as success.

Change

Thread the slice's K-tile base
(threadblock_tile_offset.k() * params.gemm_k_size / Mma::Shape::kK) from
gemm_universal_blockwise.h down through operator() → gemm_iters →
mac_loop_iter, and add it to k_block_idx before the existing clamp. The
clamp is retained as a bounds guard. With split_k_slices == 1 the offset is
zero, so existing behaviour is unchanged.

Coverage gap

The existing coverage on this path, gemm_f8t_f8n_bf16t_tensor_op_f32_blockwise_sm89.cu,
used K=128 with a single scale block and no split-K, so it could not observe the
mis-indexing. This PR extends that file to use K=512 with 4 scale blocks and
distinct per-block scales, and exercises split_k_slices of 1, 2 and 4 under
both threadblock swizzles.

Testing

RTX 4090 (SM89), CUDA 12.4.
cutlass_test_unit_gemm_device_tensorop_sm89_blockwise passes with
split_k_slices of 1, 2 and 4.

[amacharla15]

Hi @hwu36 — could you take a look at this when you have a chance? This fixes the split-K scale indexing issue in the SM89 blockwise GEMM path and extends the regression coverage introduced in #3394. Thanks!

[hwu36]

@jackkosaian , could you please take a look?