fix(tensilelite): bound custom kernel stores at 0xfffff000 by tony-davis · Pull Request #12482 · ROCm/rocm-libraries

tony-davis · 2026-09-23T16:00:11Z

JIRA ID : ROCM-31258

Motivation

hipBLASLt chooses a GPU kernel for each matrix multiplication from a tuned list, and some entries are hand-written kernels stored under tensilelite/Tensile/CustomKernels. A kernel writes its output matrix D through a buffer resource descriptor, a small structure that bounds the memory its store instructions may touch. The hardware silently discards any store at or past the descriptor's length field (num_records), and it raises no fault.

Code-generated kernels set that length to 0xfffff000 (commit ae0ca64a9c8), and the host admits shapes up to that bound. 120 hand-written kernels still used 0x80000000, or 2 GiB, so any admitted shape whose store range passed 2 GiB returned success with part of D never written. On gfx950 this is reachable through an ordinary fp32 matmul on the xf32 fast path: at M = 3,000,000, N = 256, K = 32, the kernel the heuristic ranks first leaves 14,976 elements of D unwritten.

Technical Details

Each affected kernel uses the old value in two roles. It is the length of the D, C and workspace descriptors, and it is also the store offset given to out-of-bounds lanes so the hardware drops their writes. Both roles must change together: raising only the length would bring those lanes' offsets back in bounds, so they would write past the end of D.

Readable kernels (22 files; 12 gfx950, 10 gfx942): .set BufferOOB, 0x80000000 becomes 0xfffff000. Nine of the gfx942 files also wrote 0x80000000 into the descriptor directly; those lines now use BufferOOB.
Raw-encoded kernels (98 gfx942 files, stored as .long instruction words): only the 32-bit literal of each descriptor-length s_mov_b32 and each out-of-bounds v_mov_b32 changes to 0xFFFFF000.
Reassembling all 120 files before and after shows the same instruction count and addresses. The only differences are 277 s_mov_b32 and 3,383 v_mov_b32 literals going from 0x80000000 to 0xfffff000.
The new value leaves at least 4 KiB of headroom below 2^32. Every out-of-bounds store in these kernels uses a scalar offset of 0 and an instruction offset of at most 512 bytes, so a discarded lane's address cannot wrap back into bounds. Workspace accesses never receive the out-of-bounds offset.
No host code or tuned logic changes. Two new tests are added (see Test Plan).

A build note for reviewers: the .s files are not tracked as build dependencies, so an incremental build does not reassemble them. Delete device-library/tensilelite-device-libraries.stamp in the build tree to force it.

Test Plan

Tensile/Tests/unit/test_CustomKernels_buffer_oob.py, in the TensileLite -m unit lane, with no GPU needed. It reads the value KernelWriterAssembly.py emits and checks every bundled custom kernel against it. For readable kernels it checks the BufferOOB definition and the D, C and workspace descriptor lengths. For raw-encoded kernels it decodes descriptor-length and out-of-bounds moves from the .long words. Its decoder counts matched an independent classification from disassembly exactly.
matmul_integer_exact_custom_kernel_store_extent_xf32_ROCM31258 in clients/tests/data/matmul_gtest.yaml, category pre_checkin, gfx950 only. It runs fp32 NN on the xf32 path at M = 2,097,152, 2,097,153 and 3,000,000 with N = 256 and K = 32. These are the sizes just at, one past, and well past 2 GiB of checked store range. It uses integer_exact inputs, so the check is bit-exact, and c_equal_d, so an unwritten element keeps its C value and fails the check.
Hardware evidence on an MI350X (gfx950), not automated: the reproducer attached to the ticket fills D with NaN (not-a-number), runs each custom kernel the heuristic offers, and counts elements still NaN afterward.

Test Result

Unpatched and patched builds of the same develop commit (fdec6c677df), standalone hipBLASLt for gfx950, ROCm 10.2.0a20260922, MI350X:

Unit test: unpatched, 120 failed, which is every affected kernel, and 17 passed. Patched, 137 passed.
gtest: unpatched, 2,097,152 passed, and 2,097,153 and 3,000,000 failed. Patched, all 3 passed, in about 38 seconds total.
Full TensileLite unit suite, patched: 8,506 passed, 47 skipped, 2 xfailed, 1 xpassed, 0 failed; 765 saved expected results unchanged.
The hipblaslt-bench command from the ticket: norm_error 8.9e+32, failed, becomes 3.75e-06, passed, with the same custom kernel selected.

Unwritten elements of D per custom kernel offered, fp32 xf32 NN, K = 32:

`M` × `N`	Unpatched	Patched
2,097,152 × 256	0	0
2,097,153 × 256	1	0
3,000,000 × 256	14,976	0
6,000,000 × 128	231,129,088	0
12,000,000 × 64	231,129,088	0
4,194,303 × 256	32,640	255

The last row's store range falls between 0xfffff000 and 2^32. The host still admits it because its threshold on develop is 2^32, and #12281 lowers that threshold to 0xfffff000. The two changes are independent and together close the gap.

Performance and output were compared on gfx950 for the six custom kernels that the harness could select directly (fp32 xf32 NN, NT and TN; bf16 TN, two sizes; fp16 TN), each at a size from the tuned table that selects it. Each ran alternately against the unpatched and patched kernel libraries, three runs apiece. Output D was bit-identical, and the median time changed by between -2.1% and +0.5%. The six fp8 custom kernels were not compared, because the harness hits a host-side error in the library's algorithm query for that type combination, identical on both libraries. In the rebuilt gfx950 kernel library, 21 of 315 code objects changed, and the other 294 were byte-identical.

gfx942 was verified on an MI300A; see this comment. In short, there's no regression on hardware, and on gfx942 the fix is defensive rather than reachable.

Automated GitHub checks are pending.

Submission Checklist

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

Risk level

Medium (3/5). This changes shipping assembly in 120 kernels on the default gfx942 and gfx950 selection paths. The reassembly comparison bounds the change to the two constants described above. Outside the 2 GiB to 0xfffff000 store range, output was bit-identical in every comparison run. Inside it, stores that were dropped are now written. gfx942 has been run on hardware, but the six gfx942 kernels whose stores could be affected can't currently be exercised past 2 GiB (see ROCM-31692), which is why this is not rated lower.

ROCM-31016 and fix(tensilelite): stop offering solutions whose output stores are dropped #12281, the matching host-threshold fix for generated kernels. The two PRs are independent and should land close together; either order is safe.
ROCM-31207, removing the 4 GiB ceiling
fix(tensilelite): validate custom kernel problem types in logic files #11385, draft build-time validation of custom kernel metadata; a per-kernel host threshold would build on it
AIHPBLAS-4650, custom kernels drifting from the metadata that selects them
ae0ca64a9c8 ([hipblaslt] adjust BufferOOB to 0xfffff000 #6664), the generator change this aligns with

Hand-written custom kernels set the length of their D, C and workspace buffer descriptors, and the store offset that discards out-of-bounds lanes, to 0x80000000. Generated kernels moved to 0xfffff000 in ae0ca64, and the host admits shapes up to that bound, so these kernels silently dropped every store past 2 GiB of output. The 22 readable kernels now set BufferOOB to 0xfffff000, and nine gfx942 files that wrote 0x80000000 into the descriptor directly now use BufferOOB. In the 98 gfx942 kernels stored as raw instruction words, the literal of each descriptor-length s_mov_b32 and each out-of-bounds v_mov_b32 changes to 0xFFFFF000. Reassembling all 120 files shows 277 s_mov_b32 and 3383 v_mov_b32 literals changed and nothing else. JIRA ID: ROCM-31258 Co-authored-by: Cursor <cursoragent@cursor.com>

test_CustomKernels_buffer_oob.py checks every bundled custom kernel, readable or raw-encoded, against the BufferOOB value that KernelWriterAssembly.py emits. It needs no GPU and fails for the 120 kernels fixed in the previous commit. matmul_integer_exact_custom_kernel_store_extent_xf32_ROCM31258 runs the gfx950 fp32 xf32 NN custom kernel at three sizes around 2 GiB of output with an exact check. Without the fix, the two sizes past 2 GiB fail. JIRA ID: ROCM-31258 Co-authored-by: Cursor <cursoragent@cursor.com>

therock-pr-bot · 2026-09-23T16:05:13Z

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	PR does not contain code files — Unit Test auto-passed
🔎 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 · 2026-09-23T16:05:16Z

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

codecov · 2026-09-23T16:59:51Z

Codecov Report

✅ All modified and coverable lines are covered by tests.

Additional details and impacted files

@@             Coverage Diff             @@
##           develop   #12482      +/-   ##
===========================================
+ Coverage    70.43%   70.77%   +0.34%     
===========================================
  Files         2933     2933              
  Lines       482485   482129     -356     
  Branches     69520    69513       -7     
===========================================
+ Hits        339819   341225    +1406     
+ Misses      118818   116974    -1844     
- Partials     23848    23930      +82

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	`48.80% <ø> (ø)`
TensileLite-Unit	`81.27% <ø> (ø)`
hipBLAS	`90.62% <ø> (ø)`		Carriedforward from ed0a44b
hipBLASLt	`35.28% <ø> (+0.01%)`	⬆️	Carriedforward from ed0a44b
hipCUB	`82.68% <ø> (ø)`		Carriedforward from ed0a44b
hipDNN	`86.94% <ø> (+0.30%)`	⬆️	Carriedforward from ed0a44b
hipFFT	`54.32% <ø> (ø)`		Carriedforward from ed0a44b
hipRAND	`76.12% <ø> (ø)`		Carriedforward from ed0a44b
hipSOLVER	`68.96% <ø> (ø)`		Carriedforward from ed0a44b
hipSPARSE	`86.99% <ø> (ø)`		Carriedforward from ed0a44b
rocBLAS	`48.31% <ø> (+0.08%)`	⬆️	Carriedforward from ed0a44b
rocFFT	`52.36% <ø> (+4.48%)`	⬆️	Carriedforward from ed0a44b
rocRAND	`57.42% <ø> (ø)`		Carriedforward from ed0a44b
rocSOLVER	`76.83% <ø> (ø)`		Carriedforward from ed0a44b
rocSPARSE	`74.76% <ø> (-<0.01%)`	⬇️	Carriedforward from ed0a44b
rocThrust	`78.34% <ø> (-<0.01%)`	⬇️	Carriedforward from ed0a44b

*This pull request uses carry forward flags. Click here to find out more.
see 72 files 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.

Copilot

Copilot review overview

🟡 Changes recommended

The host threshold remains unsynchronized, and the new validator can silently skip unrecognized kernel files.

Get a fresh assessment by requesting another Copilot review.

Review effort: Lite
Findings: 1

Open (1)

Host predicate admits extents beyond the kernel sentinel · New

What changed in this PR

This PR updates 120 TensileLite custom kernels to use the 0xfffff000 buffer-store bound, adds static validation, and introduces a large-output xf32 regression test.

Changes:

Updates descriptor and out-of-bounds store sentinels.
Adds validation for readable and raw-encoded kernels.
Adds gfx950 large-matrix regression coverage.

File	Change
`projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_CustomKernels_buffer_oob.py`	Validates custom-kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/CustomGSUs_Cijk_Ailk_Bljk_F8NH_HHS_BH_MT64x16x128_MI16x16x1_15_Equality_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/CustomGSUs_Cijk_Ailk_Bljk_F8NH_HHS_BH_MT128x16x128_MI16x16x1_16_Equality_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_S_MX_B_BIAS_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x32_MI16x16x1_shortname0_gfx950.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_HHS_BH_UserArgs_MT128x16x128_MI16x16x1_SN_K1_MIWT2_1_triple_buffer.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_HHS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_F8NBS_BH_BiasSB_AS_SABV_SAV_UserArgs_MT256x256x128_MI16x16x1_SN_K1_MIWT4_16_WSGRB2_DTVA.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_F8NBS_BH_BiasSB_AS_SABV_SAV_UserArgs_MT256x256x128_MI16x16x1_SN_K1_MIWT4_16_DTVA.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_F8NBS_BH_BiasSB_AS_SABV_SAV_UserArgs_MT256x224x128_MI16x16x1_SN_K1_MIWT4_14_WSGRB2_DTVA.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_F8NBS_BH_BiasSB_AS_SABV_SAV_UserArgs_MT256x224x128_MI16x16x1_SN_K1_MIWT4_14_DTVA.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_F8BS_BH_SAB_UserArgs_shortname0_gfx950.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_F8BS_BH_SAB_NTD_UserArgs_shortname1_gfx950.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_F8BS_BH_SAB_NTD_UserArgs_shortname0_gfx950.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_F8B8BS_BH_SAB_UserArgs_shortname0_gfx950.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_F8B8BS_BH_SAB_NTD_UserArgs_shortname1_gfx950.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_F8B8BS_BH_SAB_NTD_UserArgs_shortname0_gfx950.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname1_gfx950.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_BBS_BH_MT256x256x64_MI16x16x1_UserArgs_shortname0_gfx950.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_BBS_BH_MT16x16x128_MI16x16x1_SN_GSUM_K1_MIWT1_1_3.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_BBS_BH_MT16x16x128_MI16x16x1_SN_GSUM_K1_MIWT1_1_3_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_BBS_BH_MT16x16x128_MI16x16x1_SN_GSUM_K1_MIWT1_1_2.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_BBS_BH_MT16x16x128_MI16x16x1_SN_GSUM_K1_MIWT1_1_2_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_BBS_BH_MT128x16x128_MI16x16x1_SN_GSUM_K1_MIWT2_1_01.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_BBS_BH_MT128x16x128_MI16x16x1_SN_GSUM_K1_MIWT2_1_01_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_BBS_BH_Bias_AS_SAV_UserArgs_MT256x256x64_MI16x16x1_SN_K1_MIWT4_16_WSGRB2_DTVA.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_BBS_BH_Bias_AS_SAV_UserArgs_MT256x256x64_MI16x16x1_SN_K1_MIWT4_16_DTVA.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_BBS_BH_Bias_AS_SAV_UserArgs_MT256x224x64_MI16x16x1_SN_K1_MIWT4_14_WSGRB2_DTVA.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Alik_Bljk_BBS_BH_Bias_AS_SAV_UserArgs_MT256x224x64_MI16x16x1_SN_K1_MIWT4_14_DTVA.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_S_MX_B_BIAS_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x32_MI16x16x1_shortname0_gfx950.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_HSS_BH_Bias_GG_AS_SAV_UserArgs_shortname7_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_HSS_BH_Bias_GG_AS_SAV_UserArgs_shortname6_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_HSS_BH_Bias_GG_AS_SAV_UserArgs_shortname5_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_HSS_BH_Bias_GG_AS_SAV_UserArgs_shortname4_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_HHS_BH_Bias_GG_AS_SAV_UserArgs_shortname9_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_HHS_BH_Bias_GG_AS_SAV_UserArgs_shortname8_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_HHS_BH_Bias_GG_AS_SAV_UserArgs_shortname7_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_HHS_BH_Bias_GG_AS_SAV_UserArgs_shortname6_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_HHS_BH_Bias_GG_AS_SAV_UserArgs_shortname5_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_HHS_BH_Bias_GG_AS_SAV_UserArgs_shortname4_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_HHS_BH_Bias_GG_AS_SAV_UserArgs_shortname11_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_GG_AS_SAB_SAV_UserArgs_shortname7_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_GG_AS_SAB_SAV_UserArgs_shortname6_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_GG_AS_SAB_SAV_UserArgs_shortname5_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_GG_AS_SAB_SAV_UserArgs_shortname4_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_AS_SAB_SAV_shortname1_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_AS_SAB_SAV_shortname0_gfx942.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_AS_SAB_SAV_MT128x16x128_MI16x16x1_shortname2.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/Custom_Cijk_Ailk_Bjlk_S_MX_B_BIAS_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x32_MI16x16x1_shortname0_gfx950.s`	Updates kernel bounds.
`projects/hipblaslt/tensilelite/Tensile/CustomKernels/tensile/c_tensile_sk_Cijk_Ailk_Bjlk_HHS_BH_UserArgs_MT128x256x64_MI16x16x1.s`	Updates kernel bounds.
`projects/hipblaslt/clients/tests/data/matmul_gtest.yaml`	Adds large-output regression cases.

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

The store-bound check passed silently for kernels in which it found no BufferOOB definition or descriptor length. It now requires the set of such kernels to match an explicit list of six kernels from other generators (aiter, CK, rocRoller, Triton, Wave), each of which bounds its stores another way, so a new kernel in an unrecognized form fails instead of passing. JIRA ID: ROCM-31258 Co-authored-by: Cursor <cursoragent@cursor.com>

tony-davis · 2026-09-24T01:34:31Z

I ran this on an MI300A (gfx942, 228 CUs) with ROCm 10.2, comparing this branch against the unpatched 10.2.0a20260923 nightly kernel library.

The unit test passes here (138 passed). On the base commit it fails 120, and 94 of those are gfx942 kernels.

Of the changed gfx942 kernels, 85 move the D descriptor base forward one column at a time, so the old 0x80000000 length wasn't reachable for them with normal shapes. For those, I think this PR is alignment with the generator rather than a behavior change. All 38 changed grouped GEMM kernels pass the CPU reference when run by index, and one of them also passes at a 2.88 GB store range on both libraries.

The six kernels with the gfx950-style addressing are the fp8 x fp16 split-K ones (CustomGSUs_..._F8NH_HHS_BH_...). On current develop they don't write D at all, with or without this PR, because #11155 dropped the split-K factor from their launch grid. I filed that as ROCM-31692. With a one-line local host fix for the grid, all six give the same norm error on this branch as on the unpatched library, so I don't see a regression. I couldn't get them past 2 GiB, though, and I don't think a user can either. SynchronizerSizeCheck caps split-K output well under that when ldd = M, and with a large ldd the host passes that stride into a workspace it sizes from M x N. Both of those predate #11155, so testing an older commit wouldn't change the picture.

So for gfx942: no regression on hardware, and the fix is defensive there rather than reachable. I'd like to update the "gfx942 hardware verification is pending" line in the description to point at this comment.

…rnel-buffer-oob

tony-davis · 2026-09-25T20:51:24Z

@copilot resolve the merge conflicts in this pull request

…ng/rocm-31258-custom-kernel-buffer-oob # Conflicts: # projects/hipblaslt/clients/tests/data/matmul_gtest.yaml Co-authored-by: tony-davis <7938412+tony-davis@users.noreply.github.com>

Copilot · 2026-09-25T21:00:10Z

@copilot resolve the merge conflicts in this pull request

Resolved by merging the latest develop and keeping both test additions in projects/hipblaslt/clients/tests/data/matmul_gtest.yaml in commit fddfa37.

…pped (#12281) JIRA ID : ROCM-31016 ## Motivation hipBLASLt picks a GPU kernel for each matrix multiplication from a tuned list. Before a kernel is offered, host-side checks called predicates remove kernels that cannot run the requested problem. A kernel writes its output `D` through a buffer resource descriptor. Its 32-bit `num_records` field bounds every store, and the hardware silently discards any store at or past it. `allocPostLoopSrd` in `KernelWriterAssembly.py` sets that field to the `BufferOOB` sentinel, `0xfffff000`. `BufferStoreOffsetLimitCheck` is the predicate meant to keep selection away from shapes a kernel cannot write in full, but it compared against `2^32` instead of the sentinel. A shape whose store extent fell between the two was reported as supported and then only partly written, while `hipblasLtMatmul` returned success. This is the failure reported in [ROCm/hipBLASLt#2299](ROCm/hipBLASLt#2299), where a convolution in a video model returned a partly unwritten result with no error. ## Technical Details `BufferStoreOffsetLimitCheck` now compares against `0xfffff000`, held in a named constant with a comment pointing at where the generator emits it. The extent formula is unchanged: `ldd * min(MacroTile1, N) * elementBytes`. Each workgroup advances its descriptor base along `N` before storing (`computeStoreSrdStart`), so only one `MacroTile1` of columns has to fit, not all of `D`. That is why outputs larger than 4 GiB keep their solutions. Bounding the whole tensor instead would leave a 6.04 GiB bf16 matmul (`M` = 12,667,846, `N` = 256) with no solution at all, for a size that runs correctly today. The comparison stays strict, so an extent landing exactly on the sentinel is rejected even though its last store is still in range. That costs one `ldd` value per tile geometry and avoids reasoning about how wide a solution's final store is. A test records the choice. `debugEval` printed `strides()[1] * value`, skipping the `min()` against `N`, so for outputs narrower than `MacroTile1` it reported a larger extent than the one that decided the result. It now prints the compared value and formats the threshold from the constant. This changes host-side selection only. No kernel code changes, and the threshold is evaluated at run time rather than stored in the tuned library files, so a host-library rebuild applies it to libraries that have already shipped. Out of scope: the 21 hand-written kernels under `Tensile/CustomKernels` set `BufferOOB` to `0x80000000`, so this global threshold is still too permissive for them. That is tracked in ROCM-31258 and [#12482](#12482). ## Test Plan - `tensilelite-tests`, C++ predicate tests in `tensilelite/tests/Predicates_test.cpp`. Five cases: the measured broken shape is rejected, the shape one step below the sentinel is accepted, an extent exactly on the sentinel is rejected (the strict comparison), the shape from the original report stays selectable, and a tall output narrower than `MacroTile1` stays selectable. These run in the shared TensileLite lane without a GPU. - Hardware measurement on gfx950, as supporting evidence rather than an automated check: fill `D` with NaN, run each algorithm the heuristic offers, and count elements never written. A and B are finite and beta = 0, so any NaN left behind is a store that never happened. ## Test Result Predicate tests, local build on an MI350X host: - Full `tensilelite-tests` suite: 499 ran, 496 passed, 3 skipped. The skips are unrelated: `RangeLibraryTest.SpecificSizes/0` and `MXDataGenFP4Test.ZeroFrequencyWithinBounds/3` also skip on `develop`, and `DataInitializationReferenceRecompute.Gfx1250ReusesPreswizzledMXAcrossSolutions` is gfx1250-only. - With the threshold put back to `2^32`, 2 of the 5 fail: past-sentinel and exactly-at-sentinel. Those are the shapes whose extent falls between the sentinel and `2^32`. Hardware, MI350X, two host builds differing only in this threshold, against the tuned library from TheRock nightly `10.1.0a20260822`. The first 64 algorithms the heuristic offered were run for each shape: | Shape (bf16, `N` = 256, `K` = 48, `ldd` = `M`) | Threshold | Candidates offered | Left `D` partly unwritten | |---|---|---|---| | `M` = 8,388,607 | `2^32` | 1479 | 22 of 64 | | `M` = 8,388,607 | `0xfffff000` | 1124 | 0 of 64 | | `M` = 8,388,599 | `2^32` | 1479 | 0 of 64 | | `M` = 8,388,599 | `0xfffff000` | 1479 | 0 of 64 | All 22 had `MacroTile1` = 256, including the heuristic's first choice. The unwritten elements were the trailing partial tile of each solution's `MacroTile0` in the last column. The shape from the issue (`M` = 5,326,464) and the 6.04 GiB shape above were clean before and after and keep every candidate. A later full sweep on an MI355X against ROCm `10.2.0a20260918` ran all 2,207 offered solutions at `M` = 8,388,607. 605 left `D` partly unwritten and 1,601 ran clean. The change excludes the 606 with `MacroTile1` >= 256 and still admits exactly the 1,601 that ran clean. The one-solution gap between 606 and 605 is a counting difference between the two tools, not a working solution being excluded. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (2/5). Host-side selection only, with no kernel code change and no change to the extent formula. The predicate becomes stricter only for store extents in the 4,096 bytes between `0xfffff000` and `2^32`. Shapes outside that window get the same answer as before. Inside it, the measured solutions it removes were the ones leaving `D` partly unwritten, plus the one conservative exactly-at-sentinel extent, which would have run correctly. ## Related - [ROCm/hipBLASLt#2299](ROCm/hipBLASLt#2299), the report this fixes - ROCM-31258 and [#12482](#12482), the same mismatch for the hand-written custom kernels - ROCM-31207, removing the 4 GiB ceiling rather than checking against it --------- Co-authored-by: Cursor <cursoragent@cursor.com>

b-shi

Thanks for the fix, changes look good to me!

assistant-librarian · 2026-10-01T04:44:48Z

TheRock Submodule Bump Activity

Newest first

2026-10-01 — Included in TheRock via ROCm/TheRock#8607.

…GEMMs (#12876) JIRA ID : AIHPBLAS-4988 Stack, in merge order: **#12876 (this PR)**, #12960, #13005, #13013, #13015, #13021, #13022, #13025, #13031. Independent of the stack: #13027 (rocisa 64-bit add lowering) and #13060 (address-carry lint). ## Summary Recently we fixed some bugs where address, index or launch-size math was done in 32 bits when it needed 64. That kind of bug only shows up at big sizes or unlucky addresses, and it usually lives in a handful of solutions out of hundreds. Our tests have a hard time reaching it, because `unit_check` compares against an OpenBLAS GEMM on the host. That costs M·N·K and copies every padded buffer to the host, so the big-shape tests stay small and run only the heuristic's top pick, once. `integer_exact` already got rid of the tolerance, but not the cost. This PR adds fast_check, an opt-in check (`fast_check: 1`) for `integer_exact` tests. Instead of recomputing the GEMM, it compares a few random "fingerprints" of D with the same fingerprints computed from the inputs (Freivalds' algorithm, in exact modular arithmetic). That costs about as much as reading the matrices once, D never leaves the GPU, and a second set of fingerprints pins down which elements are wrong. Padding is filled with poison so stray reads show up, and D starts as a sentinel so missed writes show up. Every failure names the solution, its kernel and where each buffer sits in memory. Only test code changes. To make sure it catches what it's meant to, we reverted the ROCM-31016 fix (dropped stores past 2 GiB). fast_check flagged the affected solutions and pointed at the exact element where the stores stopped. It has also found real defects: on gfx950 the ROCM-31016 shape and the bf16 every-solution case fail today (ROCM-32152 and ROCM-32151), so those two cases are quarantined there, and running it next to `unit_check` found a bug in the CPU reference (ROCM-32062). **CI time added**, measured on one MI300X / one MI350X under ROCm 10.2 at the stack tip (the sum of gtest's per-test times, not counting test-data parsing at startup): - This PR: about 32 s / 15 s in pre_checkin (108 cases; the 51 self-tests take under a second), and about 3.5 min / 20 s in nightly (35 / 33 cases; gfx950 runs fewer because of the two quarantines above). - The whole stack: about 37 s / 58 s in pre_checkin, about 15 min / 9 min in nightly, and about 3.5 min / 1 min in the new `stress` tier (#13021), which needs large-memory runners. On MI300X the Stream-K cases skip unless `TENSILE_SOLUTION_SELECTION_METHOD=2` is set; with it they add about 5.5 min. ## Motivation hipBLASLt's gtest suite checks a matrix multiply (GEMM) result with `unit_check`, which compares the GPU output against a reference computed on the CPU. That reference costs time proportional to M·N·K, and the harness copies every input and output buffer, padding included, into pinned host memory. Tests for out-of-bounds indexing need large leading dimensions or large shapes, so this cost limits them. On gfx942, `matmul_size_t_index_overflow` with ldd = 1049100 spends 12 s of its 14.4 s allocating host memory. The existing overflow tests therefore keep K small and run only the heuristic's first solution, even though indexing defects such as ROCM-31016 and ROCM-31258 affect specific kernels. ## Technical Details This PR adds an opt-in YAML argument, `fast_check: 1`, for tests that use `initialization: integer_exact`. Under integer_exact, every input is a small integer, so the correct result is an exact integer as well. Only client test code and test data change. The library, TensileLite and kernel selection are untouched. - **Every element of D is checked without a CPU GEMM.** fast_check uses Freivalds' algorithm, which multiplies D by a random vector `r` and compares the result with `A·(B·r) + C·r + bias·sum(r)`. That costs time proportional to M·K + K·N + M·N instead of M·N·K. It does the same from the left with a second vector, to locate bad columns. The sums are taken modulo the prime 2^61 − 1, so the comparison is exact, and a wrong result passes one probe with probability below 2^−60. The input side is computed on the host once per test. D's side is computed on the GPU for each solution, so D is never copied to the host. When the probes disagree, the failure message lists the mismatched elements with their row, column and element offset, and marks offsets at or above 2^31 or 2^32. - **Failure context.** Each failure names the solution: its position in the test's list, its library index and its kernel. It also lists every operand's address range and flags any that crosses a 4 GiB boundary, since some address defects depend on placement rather than size. - **Out-of-bounds reads.** The padding of A, B and C (rows past the matrix inside the leading dimension, and gaps between batches) is filled with 40, a value outside the integer_exact range [−2, 2]. A kernel that reads padding therefore produces a wrong D. - **Out-of-bounds writes and missed stores.** Before each solution runs, D's whole allocation is filled with 0xFF bytes, which form a NaN in f16, bf16 and f32. Afterwards a device kernel checks that every element outside the M × N × batch region still holds that pattern, and that no element inside it does. A failed launch of any of these fills or scans fails the test; it is never read as clean. - **Results above the exact range.** bf16 holds integers exactly only up to 256. fast_check compares each larger result with its correctly rounded value (round to nearest, ties to even, pinned by a self-test), so f16 and bf16 tests can use K above the 256 limit that `unit_check` imposes on integer_exact. Those elements are recomputed exactly, in parallel, on the host. - **Supported configurations:** alpha, beta, bias, `scaleAlpha_vector`, strided batches, and every transpose, including conjugate transpose, which stores real data the same way as transpose. An output with zero rows, columns or batches needs no check, so it reads no operands and launches nothing. The test fails with a stated reason for timing runs, activations, scale factors, grouped GEMM and pointer-array batches; later PRs in the stack lift several of these. Output types run against real kernels here are f32, f16 and bf16; the stack adds int8 and fp8 cases (#13013, #13022). The code also accepts f64 and fp8 outputs, which only the self-tests exercise. New files are `clients/common/include/fast_check.hpp`, `clients/common/src/fast_check.cpp` and `clients/tests/src/fast_check_gtest.cpp`. `testing_matmul.hpp` calls them, and in fast_check-only runs it skips the CPU reference and the padded host copies. Running both checks on the same results found ROCM-32062, a separate defect in the CPU reference. When `scaleAlpha_vector` is combined with a leading dimension of A larger than its row count, `cast_mul` in `cblas_interface.cpp` scales the wrong rows of A, so `unit_check` fails a correct GPU result. No existing test combined the two. That case, `matmul_fast_check_epilogue_padded_lda_ROCM32062`, is quarantined in `known_bugs.yaml` until the fix PR removes the entry; #13015 adds a fast_check-only copy that keeps checking the GPU there. Related: ROCM-26455, ROCM-31016 (#12281), ROCM-31230 (#12332), ROCM-31258 (#12482), and the original overflow tests (#8231). ## Test Plan - `FastCheck` (host only) and `FastCheckDevice` (small device buffers, no hipBLASLt kernels) in `fast_check_gtest.cpp`. These build correct results, inject faults, and check that each fault is reported: a wrong element, a pair of errors that cancel in a plain row sum, a shifted row, a duplicated column, an unwritten element, an out-of-bounds write, wrong rounding of a large bf16 result (including a tie), an empty buffer, an empty output (zero M, N or batch count, with null unused operands), and an unsupported type. They also require the device and host passes to reach the same verdict. - pre_checkin, 126 cases: fast_check and `unit_check` on the same results across precisions and transposes, including all nine N/T/C pairs on two rectangular f32 shapes with padded leading dimensions and three batches; the size_t index-overflow shapes in f32, f16 and bf16; and the ROCM-26455 shape with bias and `scaleAlpha_vector`. - nightly, 35 cases: element offsets and a batch base offset above 2^32; shapes up to 16385 × 16383 × 4097 with partial tiles; and every solution at the output-overflow and ROCM-31016 shapes. - Regression replay: revert the ROCM-31016 threshold and run `matmul_fast_check_store_extent_ROCM31016`. - `standard` tier (smoke, quick, pre_checkin) to confirm existing tests are unchanged. ## Test Result **ROCm 10.2, at the stack tip (3364b79, which includes this PR).** Built and run in a ROCm 10.2.0 image (TheRock 10.2.0a20260928). Every fast_check case and self-test, with none failing: | GPU | Ran | Passed | Skipped | | --- | --- | --- | --- | | gfx90a (MI210) | 496 | 410 | 86 | | gfx942 (MI300X) | 539 | 473 | 66 | | gfx950 (MI350X) | 559 | 486 | 73 | **After newling's review.** With the empty-output fix, at stack tip 6296886715 (ROCm 10.2, before the rebase onto develop), the same filter ran 501 tests on gfx90a, 544 on gfx942 and 564 on gfx950, with 0 failures. newling's conjugate-transpose case failed 10 of its 18 cases before the fix (all those with a C operand) and passed all 18 after, on gfx1201 with ROCm 7.1. On the current tip (86d4ac3, with both fixes, rebased onto develop with #13119), the full fast_check filter on gfx942 (MI300X) ran 561 tests: 495 passed, 66 skipped for a stated reason, and none failed, including all 18 conjugate-transpose cases and 28 fp8 fnuz cases. The skips are quarantined known bugs, shapes the library declines, Stream-K cases that need `TENSILE_SOLUTION_SELECTION_METHOD=2`, and on the shared MI210, cases that need more free memory than it had. Earlier development runs, on gfx942 (MI300A): - Self-tests: each of six deliberate source mutations, four in the host code and two in the GPU kernels, made at least one self-test fail. - With the ROCM-31016 fix reverted, 170 of the 1,093 solutions failed. The first reported 1793 elements never written, starting at element offset 2147481599, which is the 2^32-byte boundary for bf16. With the fix in place, all 1,093 passed in 144 s. - With the fix reverted, each failure names the kernel, for example `Cijk_Ailk_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x448x32_…` (library index 75677). The placement record shows C and D, each about 4 GiB, crossing a 4 GiB boundary. - Same shapes, `unit_check` before and fast_check after: ldd = 1049100 f32 went from 14.4 s to 0.7 s, and lda = 1049100 f32 from 8.3 s to 1.2 s. - `standard` tier (smoke, quick, pre_checkin), run with the stack's second PR: 71,379 run, one failure, `matmul_fallback_compute_fp16`, which passes at the stack tip under ROCm 10.2 on an MI300X (see #12960). ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Medium (3/5): test code and test data only, with no library or device behavior change, and existing tests take the same paths unless they set `fast_check`. It is rated 3 rather than 2 because it adds shared machinery to `testing_matmul.hpp` and new GPU kernels that every architecture runs, and about half a minute of pre_checkin. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: James Newling <james.newling@gmail.com>

tony-davis and others added 2 commits September 23, 2026 15:59

github-actions Bot added project: hipblaslt project: hipsparselt ci:hipsparselt-fast labels Sep 23, 2026

assistant-librarian Bot added the organization: ROCm label Sep 23, 2026

tony-davis requested a lite review from Copilot September 23, 2026 19:11

Copilot started reviewing on behalf of tony-davis September 23, 2026 19:12

View session

Copilot AI reviewed Sep 23, 2026

View reviewed changes

Comment thread ...k_Alik_Bljk_S_MX_B_BIAS_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x32_MI16x16x1_shortname0_gfx950.s Resolved

tony-davis marked this pull request as ready for review September 23, 2026 19:39

tony-davis requested a review from a team as a code owner September 23, 2026 19:39

tony-davis requested a review from b-shi September 23, 2026 19:39

Merge branch 'develop' into users/todavis_amdeng/rocm-31258-custom-ke…

ed0a44b

…rnel-buffer-oob

tony-davis requested review from newling and talumbau September 24, 2026 21:55

Copilot started work on behalf of tony-davis September 25, 2026 20:52

View session

Merge remote-tracking branch 'origin/develop' into users/todavis_amde…

fddfa37

…ng/rocm-31258-custom-kernel-buffer-oob # Conflicts: # projects/hipblaslt/clients/tests/data/matmul_gtest.yaml Co-authored-by: tony-davis <7938412+tony-davis@users.noreply.github.com>

Copilot finished work on behalf of tony-davis September 25, 2026 21:00

tony-davis mentioned this pull request Sep 25, 2026

fix(tensilelite): stop offering solutions whose output stores are dropped #12281

Merged

1 task

b-shi approved these changes Sep 28, 2026

View reviewed changes

tony-davis merged commit 5517c00 into develop Sep 28, 2026
164 of 171 checks passed

tony-davis deleted the users/todavis_amdeng/rocm-31258-custom-kernel-buffer-oob branch September 28, 2026 15:10

tony-davis mentioned this pull request Oct 1, 2026

test(hipblaslt): Add fast_check verification for large integer_exact GEMMs #12876

Merged

1 task done

Navigation Menu

Repository navigation

fix(tensilelite): bound custom kernel stores at 0xfffff000 - #12482

fix(tensilelite): bound custom kernel stores at 0xfffff000#12482
tony-davis merged 5 commits into
developfrom
users/todavis_amdeng/rocm-31258-custom-kernel-buffer-oob

tony-davis commented Sep 23, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Sep 23, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Sep 23, 2026

Uh oh!

codecov Bot commented Sep 23, 2026 •

edited

Loading

Uh oh!

Copilot AI left a comment

Uh oh!

Uh oh!

tony-davis commented Sep 24, 2026

Uh oh!

tony-davis commented Sep 25, 2026

Uh oh!

Copilot AI commented Sep 25, 2026

Uh oh!

b-shi left a comment

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 1, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

Conversation

tony-davis commented Sep 23, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Risk level

Related

Uh oh!

therock-pr-bot Bot commented Sep 23, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

therock-pr-bot Bot commented Sep 23, 2026

Uh oh!

codecov Bot commented Sep 23, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Copilot review overview

🟡 Changes recommended

Uh oh!

Uh oh!

tony-davis commented Sep 24, 2026

Uh oh!

tony-davis commented Sep 25, 2026

Uh oh!

Copilot AI commented Sep 25, 2026

Uh oh!

b-shi left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 1, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

tony-davis commented Sep 23, 2026 •

edited

Loading

therock-pr-bot Bot commented Sep 23, 2026 •

edited

Loading

codecov Bot commented Sep 23, 2026 •

edited

Loading