fix(tensilelite): match 16-bit TN CMS prefetch order (AIHPBLAS-5228) by newling · Pull Request #14147 · ROCm/rocm-libraries

newling · 2026-10-21T20:14:40Z

JIRA ID : AIHPBLAS-5228

Motivation

hipBLASLt uses TensileLite to generate GPU matrix-multiplication kernels. #14129, tracked by AIHPBLAS-5228, reports intermittent wrong results on MI355X (gfx950) for BF16, a 16-bit floating-point input format, with A transposed and B untransposed (TN), at M=1024, N=4096, K=4096, batch=2, alpha=1, beta=0. The reported develop version is 1ae074b37e8; the reproducer and one-solution logic file generate a custom main-loop schedule (CMS) with a 256×160×64 block of matrix work.

RocJITsu identifies a missing dependency in that exact generated kernel: a read from local data share (LDS), the GPU's workgroup-local memory, can reach a prefetched B tile before its global-memory loads finish writing it. This PR aligns the prefetch order with the loop's existing waits, using the same repair mechanism as the earlier TensorFloat-32 (TF32) fix #12605.

Technical Details

The affected TN branch of _get_schedule_256x160x64_16bit uses direct-to-LDS loads for both inputs, two stages of global-read prefetching (PrefetchGlobalRead=2), and one stage of local-read prefetching (PrefetchLocalRead=1). This CMS is selected on gfx950 with TransposeLDS=1 and LDSTrInst=false. The schedule family was introduced in #3340.

BEFORE — the prefetch and loop disagree about load order. The final prefetch before loop entry issues 8 A loads followed by 5 B loads into the next LDS bank. s_waitcnt vmcnt(13) allows that entire set of 13 global-memory operations to remain outstanding. The loop then issues B before A. At the first read from the next B bank, 5 current-loop B loads and 2 current-loop A loads have been issued after the prefetched set. Its vmcnt(15) only requires the oldest 5 of those 20 loads to complete.

Those oldest 5 loads belong to A. All 5 prefetched B loads can still be outstanding when ds_read_b128 reads their destinations. The barrier after that wait, s_barrier, brings the waves, or groups of GPU threads, together; it does not replace the missing global-load completion wait.

AFTER — the same wait completes the needed B loads. Set the existing SwapGlobalReadOrder=1 property and exchange the GRA/GRB schedule keys. With swapping enabled, those keys name the first and second global-read streams, so exchanging them preserves the physical loop's B-before-A instruction order while changing the prefetch to B5 followed by A8.

The pending-operation accounting at the first next-bank B read becomes:

State, oldest operations first	BEFORE	AFTER
Last prefetched set	A8, B5	B5, A8
Younger current-loop loads	B5, A2	B5, A2
Operations retired by `vmcnt(15)`	First 5 prefetched A loads	All 5 prefetched B loads
Operations allowed to remain pending	Prefetched A3+B5, current B5+A2	Prefetched A8, current B5+A2
Next-bank B read	Can race with its producer	Its producers are complete

The later A wait completes the prefetched A loads before their reads. At K=128, the loop with no further global loads uses vmcnt(2): the original order leaves the final 2 B loads pending, while the reordered prefetch completes B before those reads. At K=64, the existing early-exit vmcnt(0) already completes the prefetch and gives no reports.

Both generated main-loop paths are identical before and after this fix. No wait or barrier is added or changed. In the reported solution, setting ScheduleGROverBarrier=0 produces assembly identical to the baseline; that parameter does not isolate this CMS ordering defect.

Generated assembly for the reported solution: before, after, and unified diff. Each assembly gist contains the complete file; use Raw to download it.

Performance implications. The repeated loop's instruction schedule, waits, barriers, and generated register/LDS resource metadata are unchanged for the reported kernel. The fix retains overlap with pending A loads instead of inserting a full global-memory drain. We expect long-K throughput to remain close to baseline, but the reordered prologue and tail loads can change startup and tail latency, particularly for short-K problems. No physical-GPU timing measurements were made; emulator wall times are not performance evidence. A timing comparison would measure that remaining effect separately from establishing the race.

This extends the TF32 prefetch-order fix #12605, AIHPBLAS-4866 to the 16-bit TN schedule and reuses its regression model. The pre-loop fixes #11174, ROCM-28430 and #12468, ROCM-31578 address the next dependency, from LDS reads to register consumers; their LDS waits do not complete the global-to-LDS writes involved here.

Test Plan

Run the emitted-assembly prefetch regressions and adjacent CMS/layout tests, including the strict YAML corpus check:

cd projects/hipblaslt/tensilelite
python -m pytest -q \
  tensilelite/Tests/unit/test_tf32_dtl_prefetch.py \
  tensilelite/Tests/unit/test_CustomSchedule.py \
  tensilelite/Tests/unit/test_CustomSchedule_LayoutAutoDetection.py \
  tensilelite/Tests/unit/test_input_yaml_corpus_clean.py

The 8 new cases cover BF16 and FP16, ordinary tile processing and StreamK, which distributes work along the reduction dimension, and both ScheduleGROverBarrier settings. They generate actual gfx950 assembly, seed the pending-load queue from its prologue, and check completion before both inputs' next-bank reads in each loop path. The existing steady-state CMS validator seeds from a previous loop iteration and misses this entry mismatch. These tests are collected by the shared TensileLite coverage lane and by Math CI's preliminary unit stage, including its gfx950 run, as documented in TensileLite testing.

Regenerate and assemble the reported one-solution logic with TensileCreateLibrary, then compare baseline and fixed code objects using RocJITsu's race plugin. Include prefetch-only, first-loop, repeated-loop, tail, partial-tile, and batched cases, plus controls that change only prefetch order or only the prologue wait. This establishes the dependency without relying on repeated hardware launches to produce wrong numbers.

Test Result

All 8 new regression cases fail on the unpatched source with prefetched B is still pending and pass with the fix. After rebasing onto develop at 7fe48480310, the command above passes 777 tests, including all 12 prefetch cases. The final regenerated assembly is identical to the previously validated source fix; the client rebuild completes without further work.

At M=256, N=160, K=192, batch=1, the baseline gives 5 reports. Changing only the last prefetch to B-before-A gives 0, as does changing only the prologue vmcnt(13) to vmcnt(0). Disabling CMS also gives 0. The two loop paths were exercised independently, with a controlled hardware-ID substitution selecting the second path; this does not model waves running on different physical execution units at the same time.

The baseline/reorder matrix completed 28 full-output numerical validations across K=64/128/192/256/512/4096, both loop paths, partial tiles, batching, and the original reported shape. Baseline produces 54 reports across that matrix; every reordered case produces 0. The regenerated source fix completed a further 24 validations with 0 reports, including K=65/129/193 tails and both paths at the reported shape. Three fresh runs against the final rebased code objects, at K=128, K=192, and the original shape, also pass without reports. There are no unrelated race reports in these retained runs.

Numerical validation passes in both arms under functional emulation, which does not establish whether the pending loads corrupt output on physical hardware. The independently tracked pending operations establish the race. Device code was assembled for gfx950 using ROCm 7.1/Clang 20, code-object version 5; the reporter used ROCm 7.2.3. RocJITsu version is 88ad60dee24848ffba2096d58c4711ba9fab906e.

Applicable root pre-commit hooks, whitespace checks, and Black/pyflakes checks on the modified test pass. CustomSchedule.py retains the same 65 pre-existing pyflakes diagnostics as the base, with no additions. Automated GitHub and Math CI results are pending.

Race log, BEFORE. At M=1024, N=4096, K=4096, batch=2, RocJITsu reports 5 races, one for each prefetched B load. The first complete record is shown below; the remaining records concern LDS bytes 103680, 108032, 112384, and 116736. The first arrow marks the pending direct-to-LDS write; the second marks its premature LDS read. The 112 instructions elided marker comes from the detector.

RACE kernel=Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x160x64_MI16x16x1_CMS_SN_LDSB0_AFC1_AG0_AGGSUA0_AGNTAB0_AFEM1_AFEM1_ASEM1_BL1_BS1_CD1_1_CLR1_CLS0_CADS0_DTLA1_DTLB1_DTLM0_DTVA0_DTVB0_DTVMXSA0_DTVMXSB0_DTVSM0_DPLB0_EPS0_ELFLR0_EMLLn1_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_HPLR0_ISA950_ICIW0_IU1_K1_LDSSI0_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPMXSA0_LBSPPMXSB0_LBSPPM0_LPA16_LPB32_LPMXSA0_LPMXSB0_LPMn1_LRVW8_LWPMn1_MIAV0_MIWT8_5_MXLIBL_MXSFNS_MO64_MGRIPM1_NTn1_NTA0_NTB4_NTC4_NTD4_NTE0_NTG0_NTMXSA0_NTMXSB0_NTM0_NTWS0_NVn1_NVA0_NVB0_NVC0_NVD0_NVE0_NVG0_NVMXSA0_NVMXSB0_NVM0_NVWS0_NEPBS10_NLCA1_NLCB1_ONLL1_PXCCM0_PAP0_PGL0_PGR2_PLR1_PKA1_RAP0_SGROB1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW8_SKFTR0_SNLL0_SIP1_SGRO0_TDMI0_TDMIM0_TDMLWS0_TDMPLB0_TDMS0_TIN0_THn1_THA0_THB0_THC0_THD0_THE0_THG0_THMXSA0_THMXSB0_THM0_THWS0_TPSSK_TLDS1_TLDSMn1_ULSGRO0_USL1_USLMX0_UDFMAC0_UIOFGRO0_UPLRP0_USFGRO0_USI0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WASG_WG32_8_1_WGMXCC1_WQS0 symbol=Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x160x64_MI16x16x1_CMS_SN_LDSB0_AFC1_AG0_AGGSUA0_AGNTAB0_AFEM1_AFEM1_ASEM1_BL1_BS1_CD1_1_CLR1_CLS0_CADS0_DTLA1_DTLB1_DTLM0_DTVA0_DTVB0_DTVMXSA0_DTVMXSB0_DTVSM0_DPLB0_EPS0_ELFLR0_EMLLn1_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_HPLR0_ISA950_ICIW0_IU1_K1_LDSSI0_LDSTI0_LBSPPA1024_LBSPPB1024_LBSPPMXSA0_LBSPPMXSB0_LBSPPM0_LPA16_LPB32_LPMXSA0_LPMXSB0_LPMn1_LRVW8_LWPMn1_MIAV0_MIWT8_5_MXLIBL_MXSFNS_MO64_MGRIPM1_NTn1_NTA0_NTB4_NTC4_NTD4_NTE0_NTG0_NTMXSA0_NTMXSB0_NTM0_NTWS0_NVn1_NVA0_NVB0_NVC0_NVD0_NVE0_NVG0_NVMXSA0_NVMXSB0_NVM0_NVWS0_NEPBS10_NLCA1_NLCB1_ONLL1_PXCCM0_PAP0_PGL0_PGR2_PLR1_PKA1_RAP0_SGROB1_SIA3_SS1_SPO1_SRVW0_SSO1_SVW8_SKFTR0_SNLL0_SIP1_SGRO0_TDMI0_TDMIM0_TDMLWS0_TDMPLB0_TDMS0_TIN0_THn1_THA0_THB0_THC0_THD0_THE0_THG0_THMXSA0_THMXSB0_THM0_THWS0_TPSSK_TLDS1_TLDSMn1_ULSGRO0_USL1_USLMX0_UDFMAC0_UIOFGRO0_UPLRP0_USFGRO0_USI0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WASG_WG32_8_1_WGMXCC1_WQS0 dispatch=10 type=LDS access=read reg=99328 wave=0 lane=0 wg=129,0,0 conflict=unknown
Race on LDS byte 99328 [workgroup (129, 0, 0), wave 0, lane 0]
  ==>  0x7d60adf0492c  buffer_load_dwordx4 v8, s[68:71], 0 offen nt lds  ; <-- wave 0
       0x7d60adf04934  s_add_u32 m0, m0, 0x1100
       0x7d60adf0493c  buffer_load_dwordx4 v9, s[68:71], 0 offen nt lds
       0x7d60adf04944  s_add_u32 m0, m0, 0x1100
       0x7d60adf0494c  buffer_load_dwordx4 v10, s[68:71], 0 offen nt lds
       0x7d60adf04954  s_add_u32 m0, m0, 0x1100
       0x7d60adf0495c  buffer_load_dwordx4 v11, s[68:71], 0 offen nt lds
       0x7d60adf04964  s_add_u32 m0, m0, 0x1100
       0x7d60adf0496c  buffer_load_dwordx4 v12, s[68:71], 0 offen nt lds
       0x7d60adf04974  s_xor_b32 s52, 0x10000, s52
       0x7d60adf0497c  s_xor_b32 s53, 0x10000, s53
       0x7d60adf04984  s_branch 1
       0x7d60adf0498c  s_waitcnt vmcnt(13) expcnt(7) lgkmcnt(15)
       0x7d60adf04990  s_barrier
       0x7d60adf04994  ds_read_b128 v[14:17], v118
       0x7d60adf0499c  ds_read_b128 v[18:21], v118 offset:128
       ... 112 instructions elided ...
       0x7d60adf04ce4  v_mfma_f32_16x16x32_bf16 acc[156:159], v[94:97], v[42:45], acc[156:159]
       0x7d60adf04cec  s_waitcnt vmcnt(63) expcnt(7) lgkmcnt(0)
       0x7d60adf04cf0  s_barrier
       0x7d60adf04cf4  v_mfma_f32_16x16x32_bf16 acc[0:3], v[98:101], v[46:49], acc[0:3]
       0x7d60adf04cfc  s_mov_b32 m0, s52
       0x7d60adf04d00  buffer_load_dwordx4 v0, s[64:67], 0 offen lds
       0x7d60adf04d08  v_mfma_f32_16x16x32_bf16 acc[4:7], v[98:101], v[50:53], acc[4:7]
       0x7d60adf04d10  v_mfma_f32_16x16x32_bf16 acc[8:11], v[98:101], v[54:57], acc[8:11]
       0x7d60adf04d18  v_mfma_f32_16x16x32_bf16 acc[12:15], v[98:101], v[58:61], acc[12:15]
       0x7d60adf04d20  s_add_u32 m0, m0, 0x1080
       0x7d60adf04d28  buffer_load_dwordx4 v1, s[64:67], 0 offen lds
       0x7d60adf04d30  v_mfma_f32_16x16x32_bf16 acc[16:19], v[98:101], v[62:65], acc[16:19]
       0x7d60adf04d38  v_mfma_f32_16x16x32_bf16 acc[20:23], v[98:101], v[66:69], acc[20:23]
       0x7d60adf04d40  s_waitcnt vmcnt(15) expcnt(7) lgkmcnt(15)
       0x7d60adf04d44  s_barrier
  ==>  0x7d60adf04d48  ds_read_b128 v[78:81], v13  ; <-- wave 0 lane 0
END_RACE

Race log, AFTER. The same reported shape with the regenerated source fix completes full-output validation with 0 race records on both independently selected loop paths.

Submission Checklist

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

Risk level

Medium (3/5): changes a shipping gfx950 kernel schedule for BF16/FP16 TN 256×160×64 direct-to-LDS kernels. Emitted-assembly regressions and capacity-aware detector controls cover the missing dependency; the reported kernel retains its loop schedule and resource metadata. Physical-GPU performance is unmeasured, and shared automated checks are pending.

The gfx950 MT256x160x64 TN custom schedule issues B before A, but its prologue prefetched A before B. Partial VMEM waits in the first loop and the no-global-load loop can leave required B DirectToLds writes pending at LDS reads. Set the existing SwapGlobalReadOrder property and exchange schedule stream keys so both prefetches issue B before A while retaining both SIMD loop bodies and every wait/barrier. Extend the emitted-assembly dependency regression to BF16 and FP16 with ordinary and StreamK scheduling. JIRA ID : AIHPBLAS-5228

therock-pr-bot · 2026-10-21T20:20:07Z

✅ 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 · 2026-10-21T20:20:10Z

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

nakajee

The code change and asm code are reasonable.
Looks good.
Please make sure no CI regression with this change.

codecov · 2026-10-21T23:44:03Z

Codecov Report

✅ All modified and coverable lines are covered by tests.

Additional details and impacted files

@@             Coverage Diff             @@
##           develop   #14147      +/-   ##
===========================================
+ Coverage    75.94%   75.94%   +0.01%     
===========================================
  Files         2873     2873              
  Lines       477013   477014       +1     
  Branches     71574    71574              
===========================================
+ Hits        362223   362253      +30     
+ Misses       90091    90068      -23     
+ Partials     24699    24693       -6

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	`48.93% <ø> (ø)`
TensileLite-Unit	`83.69% <100.00%> (+0.04%)`	⬆️
hipBLAS	`90.62% <ø> (ø)`		Carriedforward from 7fe4848
hipBLASLt	`64.43% <ø> (ø)`
hipCUB	`82.68% <ø> (ø)`		Carriedforward from 7fe4848
hipDNN	`87.03% <ø> (ø)`		Carriedforward from 7fe4848
hipFFT	`42.00% <ø> (ø)`		Carriedforward from 7fe4848
hipRAND	`76.12% <ø> (ø)`		Carriedforward from 7fe4848
hipSOLVER	`69.67% <ø> (ø)`		Carriedforward from 7fe4848
hipSPARSE	`86.99% <ø> (ø)`		Carriedforward from 7fe4848
rocBLAS	`68.73% <ø> (ø)`		Carriedforward from 7fe4848
rocFFT	`48.10% <ø> (ø)`		Carriedforward from 7fe4848
rocRAND	`57.30% <ø> (ø)`		Carriedforward from 7fe4848
rocSOLVER	`76.74% <ø> (ø)`		Carriedforward from 7fe4848
rocSPARSE	`76.92% <ø> (ø)`		Carriedforward from 7fe4848
rocThrust	`93.15% <ø> (ø)`		Carriedforward from 7fe4848

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

Files with missing lines	Coverage Δ
...nsilelite/tensilelite/Components/CustomSchedule.py	`96.30% <100.00%> (+<0.01%)`	⬆️

... and 3 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.

assistant-librarian · 2026-10-26T08:14:02Z

TheRock Submodule Bump Activity

Newest first

2026-10-26 — Included in TheRock via ROCm/TheRock#9478.

JIRA ID : AIHPBLAS-5076 ## Motivation hipBLASLt ships three hand-written gfx950 assembly kernels for FP32 GEMMs that use `F32XdlMathOp: X` (XF32): `Custom_Cijk_{Ailk_Bljk,Ailk_Bjlk,Alik_Bljk}_S_MX_B_BIAS_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x32_MI16x16x1_shortname0_gfx950`. They split each FP32 input into BF16 parts and feed BF16 matrix instructions (MFMA). Unlike TensileLite's generated kernels, their instruction schedule, including every wait, is written by hand in `tensilelite/tensilelite/CustomKernels/tensile/`. The defect scan prompted by [AIHPBLAS-5076](https://amd-hub.atlassian.net/browse/AIHPBLAS-5076) included a gfx950 race-detector sweep of hipBLASLt solution indices under the rocjitsu GPU emulator. It found that the TN kernel (`Alik_Bljk`) reads registers before the local data share (LDS) loads that fill them are guaranteed to complete. Both shipped solution indices that select this kernel report 24 register races and memory-wait warnings at M=N=K=256, while numerical validation still passes. The NN and NT siblings run clean. This PR moves one existing wait so the TN kernel waits for its B operand before converting it, as the siblings already do. ## Technical Details Each main-loop iteration reads FP32 tiles from LDS into temporary registers with `ds_read_b128`, converts them to BF16 with `v_cvt_pk_bf16_f32`, and computes the low-order remainder with `v_dot2c_f32_bf16`, which rewrites the same temporaries. LDS reads complete asynchronously and are tracked by the `lgkmcnt` counter, so every instruction that reads or overwrites those temporaries must first wait until the read has returned. **BEFORE — the B conversions precede their wait.** Near the end of the `MAINLOOP` macro, the kernel issues the eight next-iteration B reads, `LR B 0` through `LR B 7`, into `vgprValuB_T0_I0` (`v192`–`v223`). The last wait before them, `s_waitcnt lgkmcnt(5) // Wait for all LR A0 to finish`, retires only the A reads. After about 65 more instructions, the macro converts all 32 B temporaries, rewrites them with 32 `v_dot2c` instructions and issues another MFMA. Only then does it reach the wait meant for those reads: ```asm LR B 7 \usePLR // eighth ds_read_b128 into vgprValuB_T0_I0+28..31 ... // A conversions and 21-22 MFMAs, no LDS wait v_cvt_pk_bf16_f32 v[vgprValuB_X0_I0+0+0], v[vgprValuB_T0_I0+0+0], v[vgprValuB_T0_I0+0+1] ... // 15 more conversions, 32 v_dot2c, 4 more conversions v_mfma_f32_16x16x32_bf16 ... // mfmaIndex:189 s_waitcnt lgkmcnt(0) // Wait for all LR B0 to finish ``` A conversion can therefore read a temporary whose LDS result has not arrived, and `v_dot2c` can overwrite a temporary that a pending read will still write. The detector reports the first register of each of the eight reads, `v192` through `v220`, on both SIMD-specific loop paths. Its warnings require `lgkmcnt` at most 7 for the first read and at most 4 for the fourth: the number of younger reads that may still be pending. **AFTER — wait once, before the first reader.** The NN and NT kernels place `s_waitcnt lgkmcnt(0) // Wait for LR B0` immediately before their B conversions. This PR moves the TN kernel's existing wait to the same place, ahead of the conversion block. No LDS operation is issued between the old and new positions, so the later conversions that relied on the old position remain covered. `MAINLOOP` is expanded four times: the even- and odd-SIMD main loops, the last full iteration without global loads, and the single-iteration path, where the conversions are compiled out. Each expansion changes by that one moved instruction. Instruction count (26,228), registers, LDS allocation and every other wait are unchanged. The complete hand-written source and kernel disassembly are in the [before](https://gist.github.com/9046d6ae97b74319d395c327418d4724) and [after](https://gist.github.com/77af1eef91145f10c7fddbe8ef4fdbdc) gists. The [diff gist](https://gist.github.com/c77390cdf533af011d691f4052c3affb) holds the source diff and a machine-instruction diff showing exactly the four moved waits. Use **Raw** to download the complete large files. **Performance implications.** After the last B read, the moved wait still follows the same 64–66 instructions, including 21–22 MFMAs. That independent work provides an opportunity to overlap the LDS reads with computation. If the wait stalls, it delays the B conversions only until the data they consume has arrived. The previous placement saved that time only by reading data the hardware did not guarantee to be present. The loop keeps its partial A waits, global-load waits, barriers and instruction order. Physical GPU timing was not measured, and emulator wall times are not performance evidence. This differs from the generator fixes [#14147](#14147) and [#14330](#14330), which change TensileLite's scheduler and cannot reach hand-written custom kernels. ## Test Plan Follow the [reproduction recipe](https://gist.github.com/12f1e2186fdcb357d6095e1ca5c72ac7) to generate baseline and repaired kernels from the same one-solution gfx950 logic entry with `TensileCreateLibrary`, then run them through `tensilelite-client` under rocjitsu's race plugin with full-output numerical validation. The recipe specifies the tool versions, uses portable path variables, selects reduced-library solution 0, and checks the expected kernel dispatch, races and wait warnings as well as numerical validation. Cover a single iteration, short and long K, a K tail, multiple Stream-K tiles, M/N edge tiles and batch 2. Confirm that the regenerated baseline matches the kernel in which the race was first found. ## Test Result The regenerated `develop` kernel is instruction-identical to the CI-library kernel that produced the original report. It reproduces that report's signature: 24 VGPR read races on `v192`–`v220` across waves 0 and 1, with required thresholds of at most 4 to 7. | Shapes (M×N×K, batch) | Before: races / warnings | After: races / warnings | | --- | --- | --- | | 256×256×32 (one iteration) | 0 / 0 | 0 / 0 | | 256×256×64 | 8 / 17 | 0 / 0 | | 256×256×96, ×256, ×300 | 24 / 17 each | 0 / 0 | | 256×256×4096 | 24 / 272 | 0 / 0 | | 512×512×512 | 24 / 136 | 0 / 0 | | 768×512×1024 | 24 / 408 | 0 / 0 | | 300×260×200 | 24 / 68 | 0 / 0 | | 256×256×256, batch 2 | 24 / 34 | 0 / 0 | All 20 runs exit 0, pass numerical validation and dispatch this kernel exactly once. The detector deduplicates races, so the count is not the number of affected lanes. Functional emulation passing numerically does not show whether the race corrupts results on hardware; the tracked pending reads establish the missing dependency. The [race evidence gist](https://gist.github.com/052449eb88bb3555dadb3360e7d8e001) holds all 20 logs, the original CI report and the matrix summary. It also holds the one-solution logic entry, client configuration, detector profile, and source/tool/code-object hashes. Device code was assembled for gfx950 with ROCm 7.1 amdclang (Clang 20), code-object version 5. rocjitsu is at [99b05c7d01c](ROCm/rocm-systems@99b05c7). The defect scan and targeted baseline/repaired replays provide the race evidence for this fix. A dedicated race-check job will follow later. Its behavioral check will select the affected TN XF32 kernel and reject race records and memory-wait warnings, since the hazardous baseline already passes numerical validation. The required Component CI, Math CI, Multi-Arch CI and pre-commit checks pass. The existing gfx950 race-check job also passes, but its general workloads do not select this kernel. The gfx1250 FFM and two aggregate coverage checks remain red; those statuses are separate from the targeted baseline/repaired results above. **Race log, BEFORE.** The first complete record from the original report, solution index 77224 at M=N=K=256. The first arrow marks the pending `ds_read_b128`; the second marks the conversion that reads its destination. The detector printed the elision marker itself. The other 23 records cover `v196` through `v220` and the second loop path. ```text RACE kernel=Custom_Cijk_Alik_Bljk_S_MX_B_BIAS_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x32_MI16x16x1_shortname0_gfx950 symbol=Custom_Cijk_Alik_Bljk_S_MX_B_BIAS_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x32_MI16x16x1_shortname0_gfx950 dispatch=3 type=VGPR access=read reg=192 wave=0 lane=0 wg=0,0,0 conflict=unknown Race on VGPR v192 [workgroup (0, 0, 0), wave 0, lane 0] ==> 0x719a87f16b34 ds_read_b128 v[192:195], v17 ; <-- wave 0 0x719a87f16b3c v_sub_f32_e32 v165, v165, v155 0x719a87f16b40 v_mfma_f32_16x16x32_bf16 acc[236:239], v[142:145], v[42:45], acc[236:239] 0x719a87f16b48 v_cvt_f32_bf16_e32 v156, v21 0x719a87f16b4c v_sub_f32_e32 v166, v166, v156 0x719a87f16b50 v_mfma_f32_16x16x32_bf16 acc[236:239], v[138:141], v[42:45], acc[236:239] 0x719a87f16b58 v_cvt_f32_bf16_sdwa v157, v21 dst_sel:DWORD dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 0x719a87f16b60 v_sub_f32_e32 v167, v167, v157 0x719a87f16b64 v_mfma_f32_16x16x32_bf16 acc[144:147], v[114:117], v[54:57], acc[144:147] 0x719a87f16b6c v_cvt_pk_bf16_f32 v22, v160, v161 0x719a87f16b74 v_cvt_pk_bf16_f32 v23, v162, v163 0x719a87f16b7c v_mfma_f32_16x16x32_bf16 acc[144:147], v[118:121], v[50:53], acc[144:147] 0x719a87f16b84 v_cvt_pk_bf16_f32 v24, v164, v165 0x719a87f16b8c v_cvt_pk_bf16_f32 v25, v166, v167 0x719a87f16b94 v_mfma_f32_16x16x32_bf16 acc[144:147], v[114:117], v[50:53], acc[144:147] 0x719a87f16b9c s_waitcnt vmcnt(63) expcnt(7) lgkmcnt(5) ... 112 instructions elided ... 0x719a87f16e98 v_mfma_f32_16x16x32_bf16 acc[244:247], v[142:145], v[58:61], acc[244:247] 0x719a87f16ea0 v_cvt_f32_bf16_sdwa v155, v44 dst_sel:DWORD dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 0x719a87f16ea8 v_sub_f32_e32 v189, v189, v155 0x719a87f16eac v_mfma_f32_16x16x32_bf16 acc[244:247], v[138:141], v[58:61], acc[244:247] 0x719a87f16eb4 v_cvt_f32_bf16_e32 v156, v45 0x719a87f16eb8 v_sub_f32_e32 v190, v190, v156 0x719a87f16ebc v_mfma_f32_16x16x32_bf16 acc[248:251], v[138:141], v[70:73], acc[248:251] 0x719a87f16ec4 v_cvt_f32_bf16_sdwa v157, v45 dst_sel:DWORD dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 0x719a87f16ecc v_sub_f32_e32 v191, v191, v157 0x719a87f16ed0 v_mfma_f32_16x16x32_bf16 acc[248:251], v[142:145], v[66:69], acc[248:251] 0x719a87f16ed8 v_cvt_pk_bf16_f32 v46, v184, v185 0x719a87f16ee0 v_cvt_pk_bf16_f32 v47, v186, v187 0x719a87f16ee8 v_mfma_f32_16x16x32_bf16 acc[248:251], v[138:141], v[66:69], acc[248:251] 0x719a87f16ef0 s_sub_u32 s8, s8, 1 0x719a87f16ef4 v_cvt_pk_bf16_f32 v49, v190, v191 ==> 0x719a87f16efc v_cvt_pk_bf16_f32 v82, v192, v193 ; <-- wave 0 lane 0 END_RACE ``` **Race log, AFTER.** Every shape above, including M=N=K=256, completes full-output validation with no race records and no memory-wait warnings. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (2/5): one wait moves earlier in a single shipped hand-written gfx950 kernel. That can add a stall only when B data is still in flight, and it cannot remove a dependency. Race-detector and numerical replays cover the affected loop paths and shapes. Physical GPU timing is unmeasured, and a dedicated shared-CI race regression will follow later. [AIHPBLAS-5076]: https://amd-hub.atlassian.net/browse/AIHPBLAS-5076?atlOrigin=eyJpIjoiNWRkNTljNzYxNjVmNDY3MDlhMDU5Y2ZhYzA5YTRkZjUiLCJwIjoiZ2l0aHViLWNvbS1KU1cifQ

… by #14147 and #12676 (#14390) JIRA ID : ROCM-32343, ROCM-32344 ## Motivation The gfx950 custom-schedule `MT256x160x64` kernels used to write wrong values in the TN every-solution tests (ROCM-32343 for bf16, ROCM-32153 for f16). We quarantined the failing cases in `known_bugs.yaml` so the nightlies could stay green. #14147 fixed the cause: the schedule's prologue issued the A loads before the B loads, and the main loop's partial wait counts assumed B first. That fix is on develop, so the quarantines have done their job. This PR removes the four entries they left behind, so the cases run again as real regression gates. The ROCM-32344 entry can go too. It quarantined `matmul_fast_check_short_k` (bf16 TN) because the custom `MT256x256x64` kernels gave wrong values at K = 64. #12676 fixed those kernels on Oct 7, and the quarantine landed a day later in #13022. Davood's retest on 59fb065 found no issue. ## Technical Details - `known_bugs.yaml` loses the two bf16 entries (`matmul_fast_check_every_solution_repeat` and `matmul_fast_check_streamk_every_solution_repeat`, `a_type: bf16_r`) and the two f16 entries (the same two names, `a_type: f16_r`, `M: 4099`, TN) on `gfx950`, along with their comments. - `known_bugs.yaml` also loses the ROCM-32344 entry (`matmul_fast_check_short_k`, `a_type: bf16_r`, `transA: T`) on `gfx950`, along with its comment. - Test data only. No library, kernel, or client code changes. ## Test Plan - Compare the test records `hipblaslt_gentest.py` generates on develop and on this branch. Only the previously quarantined cases should change. - Build develop plus this change for gfx950 and run the formerly quarantined cases on MI350X and MI355X. ## Test Result - `hipblaslt_gentest.py` generates the same 143,631 records on both sides. 24 records change, and all of them just lose `gfx950` from `known_bug_platforms`: the two every-solution names at 9 records each (8 bf16 plus the f16 TN 4,099 case) and `matmul_fast_check_short_k` at 6 (bf16 TN, K = 1, 15, 16, 33, 64 and 129). Nothing else changes. - MI350X (Device 75a0): `*fast_check_every_solution_repeat*` and `*fast_check_streamk_every_solution_repeat*` (without the large-K and xf32 variants) ran 48 tests, 48 passed. - MI355X (Device 75a3): the same filter ran 48 tests, 48 passed. - For comparison, before the fix my A/B on #13276 showed the bf16 TN defaults wrong in 33 of 41 runs on MI350X on a develop that predated #14147, and 0 of 70 with the fix. - I did not run these on MI350P (Device 75a8). The cluster has none. The 75a8 tables carry copies of the same schedule (2 f16 and 19 bf16 solutions), which #14147 should fix the same way. The first check of those is the MI350P nightly after TheRock's rocm-libraries pin moves past #14147 (it is at 565c4c7 today). - Required checks passed before this update. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (1/5): test data only, and these cases are nightly-category, so PR checks are unaffected. If the MI350P nightly still shows failures on the 75a8 copies, we put the quarantine back for that device's cases and look at those kernels again. MI350X and MI355X already pass. ## Risk Acceptance / Waivers None needed. Related: ROCM-32153, #13269, #13793, #13940 (closed), #14147. --------- Co-authored-by: Cursor <cursoragent@cursor.com>

github-actions Bot added project: hipblaslt project: hipsparselt ci:hipsparselt-fast labels Oct 21, 2026

newling marked this pull request as ready for review October 21, 2026 20:26

newling requested a review from a team as a code owner October 21, 2026 20:26

newling requested a review from nakajee October 21, 2026 20:32

assistant-librarian Bot added the organization: ROCm label Oct 21, 2026

nakajee approved these changes Oct 21, 2026

View reviewed changes

newling enabled auto-merge (squash) October 21, 2026 21:03

newling mentioned this pull request Oct 21, 2026

ci(hipblaslt): detect GEMM races in a nightly job #14123

Merged

1 task done

newling merged commit 46a10bc into ROCm:develop Oct 22, 2026
214 of 232 checks passed

tony-davis mentioned this pull request Oct 23, 2026

fix(tensilelite): match DTL custom sched prologue load order to loop #13276

Merged

1 task done

This was referenced Oct 26, 2026

test(hipblaslt): remove the gfx950 TN custom-kernel quarantines fixed by #14147 and #12676 #14390

Merged

fix(hipblaslt): stopgap: remove two wrong f16 custom-schedule solutions from the MI350P (75a8) TN logic #13940

Closed

newling mentioned this pull request Oct 28, 2026

fix(tensilelite): wait for B LDS reads TN XF32 custom kernel #14533

Merged

1 task

Navigation Menu

Repository navigation

fix(tensilelite): match 16-bit TN CMS prefetch order (AIHPBLAS-5228) - #14147

fix(tensilelite): match 16-bit TN CMS prefetch order (AIHPBLAS-5228)#14147
newling merged 1 commit into
ROCm:developfrom
newling:users/newling/aihpblas-5228-race-investigation

newling commented Oct 21, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Oct 21, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Oct 21, 2026

Uh oh!

nakajee left a comment

Uh oh!

codecov Bot commented Oct 21, 2026 •

edited

Loading

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 26, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

newling commented Oct 21, 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

Uh oh!

therock-pr-bot Bot commented Oct 21, 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 Oct 21, 2026

Uh oh!

nakajee left a comment

Choose a reason for hiding this comment

Uh oh!

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

Uh oh!

Codecov Report

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 26, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

newling commented Oct 21, 2026 •

edited

Loading

therock-pr-bot Bot commented Oct 21, 2026 •

edited

Loading

codecov Bot commented Oct 21, 2026 •

edited

Loading