You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
{{ message }}
Repository navigation
test(hipblaslt): hold the testing_matmul_with_bias stream in RAII so ~PlacedRegion drains it #14310
Follow-up from #14260 (and its earlier attempt, #13517, reverted in #14021).
Problem
testing_matmul_with_bias creates its stream as a raw hipStream_t near the top of the function and destroys it explicitly at the end of the function body. The device buffers (dA, dB, ..., placedWorkspace) are declared after the stream, so their destructors run after hipStreamDestroy.
~PlacedRegion already calls hipDeviceSynchronize() before it unmaps and frees its address reservation. The thing is, by then the stream has already been destroyed without a sync. Before ROCm 7.1 that leaks a reference to the stream's last command (SWDEV-533220), which keeps the placed window alive and breaks a later hipMemAddressFree at the same address. #14260 works around this with a hipStreamSynchronize gated on arg.placement.
There's also a second, smaller bug: more than 20 early returns in testing_matmul_with_bias (and the gtest FAIL() paths) happen after the stream is created and never destroy it.
Proposal
Hold the stream in an RAII wrapper (a unique_ptr with a hipStreamDestroy deleter, or a small HipStream class in clients/common) declared where the raw stream is today. Locals are destroyed in reverse order, so the buffers would go first, ~PlacedRegion's existing sync would drain the stream, and then the stream would be destroyed with nothing pending.
Non-placement tests and hipblaslt-bench pay nothing extra.
Any future PlacedRegion user (grouped GEMM, bench) is covered automatically.
No stream leak on early return.
Test plan
ROCm 7.0.2: hipblaslt-test --gtest_filter='*fast_check_placement*:*fast_check_mx_fp8_placement*-*known_bug*' passes 30/30, with no failed hsa_amd_vmem_address_free calls under AMD_LOG_LEVEL=1.
Activity
test(hipblaslt): hold the testing_matmul_with_bias stream in RAII so …