[Issue] Linux gfx950-dcgpu: hipblaslt `S_B_Bias_HA_S_SAV_UserArgs` standard suite hits HSA memory fault / GPU hang

[jayhawk-commits]

Context

• Workflow: Multi-Arch CI / TheRock rocm-libraries submodule bump validation
• Bump PRs:
  • Bump rocm-libraries from d316419 to c1555e0 TheRock#8166 — Bump rocm-libraries from d316419 to c1555e0
    • Run: https://github.com/ROCm/TheRock/actions/runs/34631952799
  • Bump rocm-libraries from d316419 to a3eb50b TheRock#8172 — Bump rocm-libraries from d316419 to a3eb50b
    • Run: https://github.com/ROCm/TheRock/actions/runs/34662050610
• Previous passing bump: Bump rocm-libraries from 8265858 to d316419 TheRock#8152 — Bump rocm-libraries from 8265858 to d316419
  • Run: https://github.com/ROCm/TheRock/actions/runs/34560816923
• Affected job: Linux::release / Test gfx950-dcgpu / Test hipblaslt / Test hipblaslt (shard N/6)
• Affected arch: gfx950-dcgpu
• Platform: Linux
• Component: hipBLASLt

Summary

Both ROCm/TheRock#8166 and ROCm/TheRock#8172 fail all six gfx950 hipBLASLt shards with the same HSA_STATUS_ERROR_MEMORY_FAULT / GPU hang signature in hipblaslt-test_standard_suite.

This appears deterministic rather than infrastructure flakiness:

• 6/6 gfx950 hipBLASLt shards fail in Bump rocm-libraries from d316419 to c1555e0 TheRock#8166.
• 6/6 gfx950 hipBLASLt shards fail in Bump rocm-libraries from d316419 to a3eb50b TheRock#8172.
• The previous rocm-libraries bump to d316419 passed all gfx950 hipBLASLt shards.
• hipblasltprovider passes in the failing runs; the failure is specific to the full hipBLASLt standard-suite test job.

Failure Signature

The failing CTest target is consistently:

hipblaslt-test_standard_suite

The immediate failure is consistently:

Warning: Queue error - HSA_STATUS_ERROR_MEMORY_FAULT
GPU HANG ANALYSIS
hipblaslt-test_standard_suite (Subprocess aborted)
0% tests passed, 1 tests failed out of 1

The faulting dispatch shape is consistent across the inspected shards:

grid=[4096, 1, 1]
workgroup=[256, 1, 1]
group_seg_size=32768

The faulting kernel family is consistently a gfx950 S_B_Bias_HA_S_SAV_UserArgs hipBLASLt kernel, for example:

Cijk_Ailk_Bjlk_S_B_Bias_HA_S_SAV_UserArgs_MT32x32x32_MI16x16x1_..._ISA950_..._SIA3_..._SK3_..._WS64_WG32_8_1_WGMXCC1

The final [ RUN ] test differs by shard, but ROCm/TheRock#8166 and ROCm/TheRock#8172 match shard-for-shard. Examples include:

pre_checkin_matmul_128_f32_rf32_rf32_rf32_rf32_r_NT_128_128_128_5_128_128_0_128_128_1
pre_checkin_matmul_bias_relu_f32_rf32_rf32_rf32_rf32_r_relu_BIASd_f32_r_NT_128_128_128_1_128_128_0_128_128_1
pre_checkin_matmul_bias_gelu_f32_rf32_rf32_rf32_rf32_r_gelu_NT_128_128_128_1_128_128_0_128_128_1
pre_checkin_matmul_bias_relu_streamk_f32_rf32_rf32_rf32_rf32_r_relu_NT_128_128_128_1_128_128_2_128_128_1

Suspected Root Cause

The leading suspect is:

8d9e03f56d7 User/mtahvapp/hipblaslt/gfx950 bbs 9102026 (#11925)

Reasoning:

• The failing bump range starts from the previous passing rocm-libraries commit d316419.
• Only three commits in the failing range touched projects/hipblaslt.
• 8d9e03f56d7 made the large gfx950-specific logic change: seven *_HAS_SAV_UserArgs.yaml files under projects/hipblaslt/.../Tensile/Logic/asm_full/gfx950/....
• The faulting kernel name is in the same S_B_Bias_HA_S_SAV_UserArgs / gfx950 area changed by that commit.

Triage Runs

Created four rocm-libraries branches and launched TheRock Multi-Arch CI with the narrow reproduction scope:

linux_amdgpu_families: gfx950
linux_test_labels: test:hipblaslt,test_filter:standard
windows_amdgpu_families: none

Runs:

Purpose	Branch / commit	Run
Before suspected commit	triage/gfx950-hipblaslt-before-11925 / 9637ceab04154cd7e82dc3560b2245a5c58566ed	https://github.com/ROCm/rocm-libraries/actions/runs/34715288859
Suspected commit	triage/gfx950-hipblaslt-at-11925 / 8d9e03f56d7bb7c998ed05c9444c8015aa538216	https://github.com/ROCm/rocm-libraries/actions/runs/34715289535
Latest failing ROCm/TheRock#8172 target	triage/gfx950-hipblaslt-pr8172-tip / a3eb50b9ea8678ea80a1a4b2ac5a26a909252f58	https://github.com/ROCm/rocm-libraries/actions/runs/34715289482
Latest failing target with suspect gfx950 YAML restored	triage/gfx950-hipblaslt-pr8172-revert-11925 / 7b612d4c5667a48390fc429797b1eb411d7a3782	https://github.com/ROCm/rocm-libraries/actions/runs/34715289527

Expected result if 8d9e03f56d7 is causal:

before suspect: pass
at suspect: fail
latest failing tip: fail
latest failing tip with suspect gfx950 YAML restored: pass

Note: the last branch is not a clean git revert because the reverse patch no longer applies at a3eb50b; it restores only the seven gfx950 logic YAML files touched by 8d9e03f56d7 to their parent-commit blobs.

Impact

This blocks promotion of the current rocm-libraries submodule in ROCm/TheRock because every gfx950 hipBLASLt shard fails on the bump PRs.

Suggested Next Steps

• Check the four rocm-libraries triage runs above when they complete.
• If the expected pass/fail pattern holds, inspect or revert the gfx950 *_HAS_SAV_UserArgs.yaml changes from 8d9e03f56d7.
• Consider adding a targeted gfx950 hipBLASLt test or quarantine for the offending S_B_Bias_HA_S_SAV_UserArgs kernel shape while the kernel/logic issue is fixed.

Notes

The logs also contain component-checker noise about missing GCC tools, but the jobs continue past that. The actionable failure is the gfx950 hipBLASLt HSA_STATUS_ERROR_MEMORY_FAULT / GPU hang in hipblaslt-test_standard_suite.

[jayhawk-commits]

Triage workflow results

The narrow TheRock Multi-Arch CI runs completed with linux_amdgpu_families=gfx950, linux_test_labels=test:hipblaslt,test_filter:standard, and Windows disabled.

Branch / commit	Run	Result	hipBLASLt gfx950 shards
triage/gfx950-hipblaslt-before-11925 / 9637ceab04154cd7e82dc3560b2245a5c58566ed	https://github.com/ROCm/rocm-libraries/actions/runs/34715288859	PASS	6/6 passed
triage/gfx950-hipblaslt-at-11925 / 8d9e03f56d7bb7c998ed05c9444c8015aa538216	https://github.com/ROCm/rocm-libraries/actions/runs/34715289535	FAIL	6/6 failed
triage/gfx950-hipblaslt-pr8172-tip / a3eb50b9ea8678ea80a1a4b2ac5a26a909252f58	https://github.com/ROCm/rocm-libraries/actions/runs/34715289482	FAIL	6/6 failed
triage/gfx950-hipblaslt-pr8172-revert-11925 / 7b612d4c5667a48390fc429797b1eb411d7a3782	https://github.com/ROCm/rocm-libraries/actions/runs/34715289527	PASS	6/6 passed

The failure signature observed in the failed shard logs is a gfx950 GPU memory access fault during hipblaslt-test_standard_suite, followed by Subprocess aborted. The failing test context is in generated ISA950 matmul cases from the HAS_SAV_UserArgs area.

This strongly implicates 8d9e03f56d7bb7c998ed05c9444c8015aa538216 / #11925. Restoring only the seven gfx950 *_HAS_SAV_UserArgs.yaml files changed by that commit on top of the current failing tip is sufficient to make the narrow gfx950 hipBLASLt standard suite pass again.

[jayhawk-commits]

Revert: #12018

Test Build: https://github.com/ROCm/rocm-libraries/actions/runs/34731255214

This PR restores the seven gfx950 HAS_SAV_UserArgs YAML files changed by #11925 to their pre-#11925 state. Prior triage showed:

• pre-User/mtahvapp/hipblaslt/gfx950 bbs 9102026 #11925: 6/6 gfx950 hipBLASLt shards passed
• User/mtahvapp/hipblaslt/gfx950 bbs 9102026 #11925: 6/6 failed
• PR8172 tip: 6/6 failed
• PR8172 tip with these seven files restored: 6/6 passed