fix(rocBLAS): repair d_vector guard logic and HMM accounting - #11398
Open
bobbyphilip-amd wants to merge 5 commits into
Open
bobbyphilip-amd wants to merge 5 commits into
bobbyphilip-amd wants to merge 5 commits into
Conversation
4 tasks
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Contributor
Author
|
@TorreZuk, take a look at this. If it makes sense I'll change from draft PR status. It's overly verbose at this point, but I can fix that later. Note there's a follow on that addresses a comment you made on the syrk PR |
bobbyphilip-amd
force-pushed
the
fix/dvector-guard-and-hmm-accounting
branch
from
August 28, 2026 05:45
857d4a1 to
b1d004f
Compare
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project check has failed because the head coverage (48.25%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage. Additional details and impacted files@@ Coverage Diff @@
## develop #11398 +/- ##
===========================================
- Coverage 70.71% 70.70% -0.01%
===========================================
Files 2847 2847
Lines 466776 466775 -1
Branches 68835 68835
===========================================
- Hits 330060 330026 -34
- Misses 112771 112812 +41
+ Partials 23945 23937 -8
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
TorreZuk
reviewed
Aug 28, 2026
Contributor
|
@bobbyphilip-amd can you get this one moving again or need help? |
TorreZuk
reviewed
Sep 16, 2026
bobbyphilip-amd
force-pushed
the
fix/dvector-guard-and-hmm-accounting
branch
2 times, most recently
from
September 16, 2026 21:44
8816664 to
151ee89
Compare
bobbyphilip-amd
force-pushed
the
fix/dvector-guard-and-hmm-accounting
branch
from
September 18, 2026 18:13
151ee89 to
8eeb747
Compare
bobbyphilip-amd
marked this pull request as ready for review
September 18, 2026 18:30
bobbyphilip-amd
force-pushed
the
fix/dvector-guard-and-hmm-accounting
branch
from
September 18, 2026 18:31
4e4d546 to
423c9ff
Compare
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
One or more issues must be addressed before approval.
Get a fresh assessment by requesting another Copilot review.
Pull request overview
Repairs rocBLAS client device-guard handling and HMM memory accounting, while adding regression tests and making host allocation initialization thread-safe.
Changes:
- Fixes guard setup/check/teardown behavior and pointer accounting.
- Adds guard-corruption and HMM lifecycle tests.
- Simplifies host-memory environment parsing and initialization.
File summaries
| File | Description |
|---|---|
| projects/rocblas/clients/include/d_vector.hpp | Updated as part of this pull request. |
| projects/rocblas/clients/gtest/host_alloc_gtest.cpp | Updated as part of this pull request. |
| projects/rocblas/clients/gtest/CMakeLists.txt | Updated as part of this pull request. |
| projects/rocblas/clients/common/host_alloc.cpp | Updated as part of this pull request. |
Review details
Suppressed comments (1)
projects/rocblas/clients/common/host_alloc.cpp:198
- The rocBLAS contribution guide explicitly prefers function-local static initialization over
std::call_oncefor one-time setup (projects/rocblas/.github/CONTRIBUTING.rst:375-392). Since this value is immutable after reading the environment, astatic const int value = [] { ... }();is thread-safe, removes the extra flag, and matches the repository's documented pattern.
static std::once_flag once_flag;
std::call_once(once_flag, [] {
- Files reviewed: 4/4 changed files
- Comments generated: 3
- Review effort level: Lite
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
bobbyphilip-amd
force-pushed
the
fix/dvector-guard-and-hmm-accounting
branch
3 times, most recently
from
September 18, 2026 22:17
f12440f to
373280f
Compare
Contributor
Author
|
@TorreZuk, FYI, this should have passed but of course I got a GPU hang. Rerunning. |
Contributor
|
I don't want to push you too hard, but as you are changing HMM next time you would want to add the ci:extended label as those are the ones that run additional HMM test with HSA_XNACK=1. It is also too late on Friday for me to push a new instruction file that hopefully will suggest this to you when getting co-pilot reviews. If you didn't have a stack I would say you could set auto-merge if you are happy with it. Have a good weekend! |
- device_vector_teardown released managed memory count keyed on the offset (post-guard) pointer rather than the base allocation pointer, leaving the RAM ceiling inflated for the process lifetime; fixed by releasing before the pointer advances - device_vector_check read both guard regions into one buffer; a failed hipMemcpy left stale bytes that caused a missed corruption or a false one; each guard now reads into its own buffer, conditioned on its copy - a failed guard write in device_vector_setup now zeros m_guard_len so device_vector_check does not compare uninitialized device memory against the guard pattern and report corruption that never happened; applies to both the pre- and post-guard write; the pointer still advances by m_pad so teardown arithmetic stays consistent; device_vector_check now gates on m_guard_len > 0 - m_guard_len(0 * sizeof(T)) in the non-test constructor simplified to m_guard_len(0); stale m_pad(0) comment removed Adds three tests: guard_detects_post_overwrite and guard_detects_pre_overwrite verify that device_vector_check fires when either guard region is overwritten; hmm_count_returns_to_its_baseline verifies that the HMM accounting count returns to its baseline after a managed device_vector lifecycle with a nonzero guard pad. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Co-authored-by: bobbyphilip-amd <267730874+bobbyphilip-amd@users.noreply.github.com>
Record dv.memcheck() under the fake reporter so a failed 1024-element allocation cannot greenwash an empty failure list.
bobbyphilip-amd
force-pushed
the
fix/dvector-guard-and-hmm-accounting
branch
from
September 19, 2026 03:43
373280f to
db41589
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
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.
JIRA ID: AIROCBLAS-1381
AIROCBLAS-1381
Summary
Core bug fixes (
d_vector.hpp)device_vector_teardownreleased the HMM memory count keyed on the base allocation pointer instead of the offset (post-guard) pointer thatalloc_ptr_useregistered; the map miss was silent so the ceiling rose permanently for the life of the process — fixed by callingfree_ptr_usebefored -= m_paddevice_vector_checkread both guard regions into a single shared buffer; a failedhipMemcpyleft stale bytes from the previous guard that caused a missed or false corruption report — each guard now reads into its own stack bufferdevice_vector_setupnow zerosm_guard_lensodevice_vector_checkdoes not compare uninitialized device memory against the guard pattern; the pointer still advances bym_padso teardown arithmetic stays consistent;device_vector_checkgates onm_guard_len > 0hipMemcpyDefault(preferred for HMM / managed memory; explicit HostToDevice/DeviceToHost was tried and reverted)m_guardinitialized viastd::call_once+std::once_flagreplacing a racystatic boolpatterndevice_vector_checkinstead of heap-allocating containers (safe to use from a destructor;std::vectorconstruction can throwstd::bad_alloc→std::terminate)#include <mutex>andm_init_flagdeclaration and out-of-line definition wrapped in#ifdef GOOGLE_TEST(only used in test builds)Pre-existing bug fixes
d_vector.hpp:#if __GLIBC__ < 3 && __GLIBC_MINOR__ < 39→#if defined(__GLIBC__) && __GLIBC__ < 3 && __GLIBC_MINOR__ < 39— on non-glibc platforms__GLIBC__is undefined and evaluated as 0, silently suppressing_GLIBCXX_USE_C99_INTTYPES_TR1on every non-glibc buildhost_alloc.cpp(host_malloc): racystatic auto once = false/once = truepattern replaced withstd::call_once+std::once_flag; unnecessary[&]capture removed (static locals are accessible without capture)host_alloc.cpp(host_bytes_available):mem_tokenselection andstrlencollapsed into two statics resolved once at first call — removes per-callstrlenand the mutable-static-pointer patternhost_alloc.cpp: comment// B to GBcorrected to// GB to BTest plan
guard_detects_post_overwrite— deliberately writes zeros into the post-guard; assertsdevice_vector_checkfires with"post-guard"viaEXPECT_NONFATAL_FAILURE; usesASSERT(notEXPECT) onhipMemsetso a memset failure produceskFatalFailurerather than a nonfatal failure that would satisfySingleFailureCheckerand mask the real problemguard_detects_pre_overwrite— mirror of above for the pre-guard regionguard_no_false_positive_on_clean_alloc— allocates and destroys adevice_vectorwithout touching the guards; usesScopedFakeTestPartResultReporterto assert zero nonfatal failures are emittedhmm_count_returns_to_its_baseline— verifies the HMM accounting count returns to its pre-allocation baseline after a full manageddevice_vectorlifecycle with a nonzero guard pad