Blockscaled GEMV drops the K tail when floor(K/tile) is a multiple of the stage count: example 91 fails at --k=1152

Description

The blockscaled GEMV kernel (include/cutlass/gemm/kernel/gemv_blockscaled.h) silently drops up to a tile's worth of K elements for accepted problem sizes, producing wrong output on real hardware.

Mechanism:

1. The mainloop loads whole tiles only: total_tiles = gemm_k / tileA_k_local (:320), and load_stages_gmem_to_smem advances unroll_col_k by tileA_k_local for every stage unconditionally, including stages whose cp.async was predicated off because global_k >= gemm_k (:557; the predicate itself is only valid_tile at :551/:555).
2. After the mainloop, the remainder is handled by a single guarded tail call gated on the inflated counter:

// gemv_blockscaled.h:488
if (unroll_col_k + idx_col_k * kPackedElementsA < gemm_k) {
  accum += process_tail_elements(unroll_col_k, ...);
}

When floor(gemm_k / tileA_k_local) is a multiple of the stage count, unroll_col_k after the loop is rounded up past gemm_k, the guard is false for every thread, and every remaining K element is skipped.

Reproduction (hardware)

Example 91 (91_fp4_gemv, which instantiates this kernel) built for sm_120a with CUDA 13.0.3 and run on an RTX 5060 Ti:

$ 91_fp4_gemv --k=1024 --m=4096 --batch=1   -> profiling completed
$ 91_fp4_gemv --k=1152 --m=4096 --batch=1   -> tensorD mismatch
                                              compare_reference() failed
                                              test fail
$ 91_fp4_gemv --k=1280 --m=4096 --batch=1   -> profiling completed
$ 91_fp4_gemv --k=2048 --m=4096 --batch=1   -> profiling completed

Only the predicted window fails (floor(K/tileA_k_local) % kStageCount == 0 && K % tileA_k_local != 0; here tileA_k_local == 256, so K=1152 drops the 128 elements of the fifth partial tile).

Secondary observation (same pipeline)

The same load path can also issue global reads beyond the current row/batch's K extent: the first load group before the mainloop is issued with a literal-true predicate (:346-367 region) and later groups use the per-tile valid_tile above, so cp.async sources ptr_A + unroll_col_k / kPackedElementsA up to three whole tiles past the row end in some schedules. The results feed the MMA and are currently masked out by the accumulator logic, but the out-of-range global reads themselves are not gated by anything stronger than the tile predicate.

Suggested fix

Track the actually-consumed K extent separately from the issued-load counter (or clamp the tail-guard operand to min(unroll_col_k, ...)) so the tail fires whenever any thread still has elements below gemm_k.

Start with include/cutlass/gemm/kernel/gemv_blockscaled.h, especially the mainloop and tail handling around lines 320, 488, and 551-557. Reproduce the failure with example 91_fp4_gemv at --k=1152 and compare it with the passing sizes. Done means the partial K tile is processed without dropping elements, and the reported out-of-range load path is addressed or clearly scoped.