Skip to content

test(hipblaslt): hold the testing_matmul_with_bias stream in RAII so ~PlacedRegion drains it #14310

Open

Description

@tony-davis

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.

I think this gets us:

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.
  • Same filter on a current ROCm build.
  • hipblaslt-bench timing unchanged.

Activity

  1. self-assigned this
    on Oct 23, 2026
  2. added a commit that references this issue on Oct 26, 2026
    9850374
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

Labels

No labels
No labels

Type

No type

Projects

No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions