[yenong-amd]

Resolves: #13361

Motivation

Batched TF32 GEMMs that select the handwritten SK3 gfx950 custom kernels fail to launch:

hipExtModuleLaunchKernel: work-item count exceeds the 32-bit globalWorkSize parameters
(numWorkItems (4558979072, 1, 1), from numWorkGroups (17808512, 1, 1) of workgroup size (256, 1, 1))
for kernel: Custom_Cijk_Ailk_Bljk_S_MX_B_BIAS_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x32_MI16x16x1_shortname0_gfx950
hipBLASLt status error: Expected HIPBLAS_STATUS_SUCCESS, received HIPBLAS_STATUS_INTERNAL_ERROR

Reproducer (MI350X): f32 in/out, compute xf32_r, NN, M=256 N=1856 K=320, batch_count=1492.

Root cause

These three custom kernels carry no explicit CustomKernel.grid, so CustomKernels.py infers one. For
StreamK kernels it picked the grid from ProblemType.Batched: Batched: True stamped StreamKWithBatch.
On the custom-kernel launch path (generateCustomCall), StreamKWithBatch launches sk.grid * tiles.z.

sk.grid already includes the batch: resolveStreamKSettings counts tiles as x * y * z, and Origami
sizes the grid against problem.batch. Multiplying by tiles.z applies the batch twice:

output tiles = 1 * 8 * 1492 = 11936 = sk.grid
launched     = 11936 * 1492 = 17,808,512 workgroups -> 4.56e9 work-items (> 2^32)

The kernels' assembly expects the batch-inclusive count. StreamKIdx = WorkGroup0, the SKGrid kernarg,
TotalIters, and the per-WG split all form one flat problem-wide space, and the batch index is recovered
from the tile ID via NumWorkGroups0 * NumWorkGroups1. The extra * tiles.z only inflates the launch.

Technical details

CustomKernels._buildCustomKernelFromMetadata now infers StreamKNoBatch for every StreamK custom kernel
that has no explicit grid, regardless of Batched. StreamKNoBatch launches sk.grid workgroups, which
matches what the assembly computes.

• Scope: the inference only applies when custom.config has no CustomKernel.grid. Today that is exactly
  the three SK3 gfx950 kernels:
  • Custom_Cijk_Ailk_Bljk_S_MX_B_BIAS_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x32_MI16x16x1_shortname0_gfx950.s
  • Custom_Cijk_Ailk_Bjlk_S_MX_B_BIAS_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x32_MI16x16x1_shortname0_gfx950.s
  • Custom_Cijk_Alik_Bljk_S_MX_B_BIAS_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x32_MI16x16x1_shortname0_gfx950.s
• Kernels that declare StreamKWithBatch explicitly (c_tensile_sk_Cijk_Ailk_Bjlk_HHS_BH_UserArgs_MT128x256x64_MI16x16x1.s,
  Tests/custom/custom_tensile_sk.yaml) keep it, and generateCustomCall is unchanged.
• The three .s files keep Batched: True, so their custom.config problem type still matches the batched
  problems they serve. No assembly, host code, or logic YAML changes.

An earlier revision of this PR set Batched: False in the three .s files instead. That broke the Tensile
benchmark flow: custom_xfp32.yaml failed the problem-type match against the kernels, and switching the test
to non-batched exposed an unrelated id2 codegen bug in KernelWriterConversion.py. The current revision
reverts both and makes the fix in the inference.

Tests added

• test_CustomKernelMetadata.py
  • test_build_from_metadata_streamk_grid_excludes_batch: replaces the test that expected StreamKWithBatch.
    It now expects StreamKNoBatch for Batched true and false.
  • test_sk3_tf32_custom_kernel_launches_batch_inclusive_grid: asserts the three bundled SK3 kernels infer
    StreamKNoBatch.
• clients/tests/data/matmul_gtest.yaml: matmul_tf32_streamk_custom_large_batch (gfx950, pre_checkin),
  the batch-1492 NN reproducer with norm_check.

Test plan

• tensilelite Python unit tests (tensilelite/Tests/unit).
• tox -e py3 -- -k custom_xfp32 on gfx950, which builds and benchmarks all three SK3 kernels.
• Rebuild hipBLASLt for gfx950, rerun the reproducer with hipblaslt-bench --verify, and run
  hipblaslt-test --gtest_filter='*tf32_streamk_custom_large_batch*'.

Test result

• Python unit tests: 10559 passed, 67 skipped, 794 snapshots passed.
• custom_xfp32.yaml: passed.
• Reproducer: launches 11936 workgroups instead of 17,808,512 and completes. norm_error 0.00570796 with
  --verify (hpl init, summed per-batch relative Frobenius error over 1492 batches, about 3.8e-6 per batch).
• matmul_tf32_streamk_custom_large_batch: passed.

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[therock-pr-bot]

🎉 All checks passed! This PR is ready for review.

[talumbau]

Please add a test that would fail on this situation in the future.

[yenong-amd]

Please add a test that would fail on this situation in the future.

Added root cause to the PR description.
Added the test case that failed to launch to hipblaslt gtest.
Added python test to ensure that the custom kernels keep Batched: False.

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.

Additional details and impacted files

@@             Coverage Diff             @@
##           develop   #13361      +/-   ##
===========================================
+ Coverage    71.86%   71.91%   +0.05%     
===========================================
  Files         2876     2876              
  Lines       475301   474734     -567     
  Branches     70235    70221      -14     
===========================================
- Hits        341564   341399     -165     
+ Misses      109956   109557     -399     
+ Partials     23781    23778       -3     

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	49.67% <ø> (ø)
TensileLite-Unit	83.44% <100.00%> (-<0.01%)	⬇️
hipBLAS	90.62% <ø> (ø)		Carriedforward from 2c69947
hipBLASLt	42.74% <ø> (+0.03%)	⬆️
hipCUB	82.68% <ø> (ø)		Carriedforward from 2c69947
hipDNN	86.66% <ø> (ø)		Carriedforward from 2c69947
hipFFT	42.58% <ø> (ø)		Carriedforward from 2c69947
hipRAND	76.12% <ø> (ø)		Carriedforward from 2c69947
hipSOLVER	68.96% <ø> (ø)		Carriedforward from 2c69947
hipSPARSE	86.99% <ø> (ø)		Carriedforward from 2c69947
rocBLAS	48.37% <ø> (+0.19%)	⬆️	Carriedforward from 2c69947
rocFFT	46.51% <ø> (ø)		Carriedforward from 2c69947
rocRAND	57.45% <ø> (ø)		Carriedforward from 2c69947
rocSOLVER	76.83% <ø> (ø)		Carriedforward from 2c69947
rocSPARSE	74.07% <ø> (ø)		Carriedforward from 2c69947
rocThrust	93.15% <ø> (ø)		Carriedforward from 2c69947

*This pull request uses carry forward flags. Click here to find out more.

Files with missing lines	Coverage Δ
...hipblaslt/tensilelite/tensilelite/CustomKernels.py	97.41% <100.00%> (-0.01%)	⬇️

... and 1 file with indirect coverage changes

🚀 New features to boost your workflow:

• ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
• 📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

[assistant-librarian]

TheRock Submodule Bump Activity

Newest first

• 2026-10-12 — Included in TheRock via ROCm/TheRock#8933.