Repository navigation
fix(rocm): stop leaking alpha/beta pairs in the GEMM pointer wrappers - #2252
Merged
Merged
Conversation
inureyes
added
type:bug
Member
Author
|
|
hipblaslt_gemm_ptrs, hipblaslt_gemm_rowmajor_on_stream and rocblas_gemm_ptrs heap-allocated a never-freed float pair for any scale other than (1, 0) and (1, 1), leaking 8 bytes per call. The justification (the library keeps host pointers past enqueue) does not hold: no pointer mode is set, so the scalars are read during the library call, and launch_kernel runs its functor synchronously. The wrappers now pass the addresses of the lambda's by-value captures (the function parameters in the on-stream variant), as the epilogue and batched paths already did. The use_hip_graphs() comment notes the assumption. Recorded as LOCAL_FIXES 43. Verified on gfx1151: rocm_hipblaslt_concurrency, rocm_rocblas_handle_concurrency and rocm_gather_qmm_expert_batched pass; verify-rocm-overlay passes. No Rust-reachable path calls these wrappers with another scale, so there is no behavior test for it. Closes #2243
inureyes
force-pushed
the
fix/issue-2243-gemm-alpha-beta-leak
branch
from
October 8, 2026 23:26
b54ae04 to
9eb02c9
Compare
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Three ROCm GEMM pointer wrappers heap-allocated a never-freed alpha/beta pair for any scale other than (1, 0) and (1, 1), leaking 8 bytes per call. The comments blamed the library retaining host pointers past enqueue, but no pointer mode is set, so the scalars are read during the call, and
CommandEncoder::launch_kernelruns its functor synchronously.hipblaslt_gemm_ptrs,hipblaslt_gemm_rowmajor_on_streamandrocblas_gemm_ptrsnow pass&alphaand&betaof the lambda's by-value captures (the parameters in the on-stream variant), as the epilogue and batched paths already do. Theuse_hip_graphs()comment records that deferred graph nodes would need a revisit. LOCAL_FIXES 43; stays in mlxcelverse.No Rust-reachable path calls these wrappers with another scale (
addmmgoes throughhipblaslt_gemmandgemm_rocblas; the current callers pass 1/0 or 1/1), so there is no behavior or leak test for it, and no before/after leak measurement was possible.Verification
On gfx1151:
git grep 'new float\['over patches-rocm is empty;make verify-versions verify-kernel-dtype-keys verify-kernel-port-dispatch verify-llama-compat verify-fmt verify-rocm-overlaypass;rocm_hipblaslt_concurrency,rocm_rocblas_handle_concurrencyandrocm_gather_qmm_expert_batchedpass. Fullmake verify-rocm-heldresult is in the comment below.Not verified: Metal and CUDA are not available on this host. Neither is touched (ROCm overlay only).
Closes #2243