Repository navigation
fix(tensilelite): copy the source tensor in CopyTensorVoid - #13302
Merged
Merged
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
✅ 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. |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #13302 +/- ##
===========================================
+ Coverage 75.79% 75.82% +0.04%
===========================================
Files 2880 2882 +2
Lines 478823 480202 +1379
Branches 71949 72197 +248
===========================================
+ Hits 362888 364110 +1222
- Misses 91072 91105 +33
- Partials 24863 24987 +124
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
1 task done
asleepzzz
approved these changes
Oct 12, 2026
CurtisFu1002
approved these changes
Oct 12, 2026
CopyTensorVoid took its source address from dst, so every hipMemcpyAsync copied the destination onto itself. Its only caller is copyBadInputBuffers, which fills a NaN bounds-check buffer with sentinels and then relies on CopyTensorVoid to write the real tensor into the middle. The tensor never landed, so with --bounds-check=NaN every input and output kept the inf sentinel: - the kernel and the CPU reference both computed D from inf inputs, so validation compared inf against inf and passed for any result; - the Synchronizer buffer started as inf instead of zero, so the Synchronizer clean-on-exit check reported StreamK solutions as dirty. Measured with the tensilelite client on gfx950, BF16 GEMM, 3 sizes x 8 solutions, --bounds-check=NaN --print-valids: - before: every compared element printed inf==inf; 12/24 FAILED, all from the Synchronizer check - after: compared elements print ordinary values; 24/24 PASSED - Disable mode: 24/24 PASSED before and after
newling
force-pushed
the
users/yiding12/tensilelite-nan-boundscheck-copy
branch
from
October 23, 2026 20:13
a701c39 to
0640152
Compare
1 task done
Contributor
Author
Contributor
TheRock Submodule Bump ActivityNewest first
|
DDEle
added a commit
that referenced
this pull request
Oct 27, 2026
…th through the shared benchmark pipeline (#12472) ## Motivation The TensileLite client (`tensilelite-client`) benchmarks and validates GPU kernels, the functions that run on the GPU. For every problem and every candidate solution, the loop in `client/main.cpp` calls a chain of listeners that initialize data, time the launches, compare the output with a CPU reference, and write result rows. The fused GEMM + all-to-all (A2A) path, a GEMM whose final stage also sends slices of its output D to the other GPUs, bypassed that loop: `FusedA2AClient.cpp` (about 1,100 lines) had its own solution loop, data initialization, CPU reference, timing and exit codes. As a result, options such as `--num-warmups`, `--use-gpu-timer`, `--check-synchronizer` and rotating buffers (extra copies of the inputs used in turn so the caches stay cold) had no effect on it, no result rows were written, and every fix to the shared loop had to be repeated there. JIRA ID : AIHPBLAS-4534 ## Technical Details `FusedA2AClient.cpp` is deleted. Both modes now run through the loop in `main.cpp`, and the number of GPUs is a dimension of that loop rather than a second code path. - `DeviceContext` holds one GPU's state: stream, kernel loader, `DataInitialization`, inputs and their rotating copies. On the fused path it also holds the flag block the cards use to signal each other, and at startup it enables P2P access to the other GPUs (reading and writing their memory directly) and opens one SDMA (copy engine) queue per peer. The single-device path is a list of one. Each launch is enqueued on every GPU before the next is enqueued anywhere, because the fused kernel's last work-group (group of GPU threads) waits for its peers. - `FusedA2ARunner`, what remains of the fused driver, no longer launches, times or validates. At startup it checks the options and that every pair of GPUs has P2P access; per solution it zeroes each card's flag block and Synchronizers, allocates the receive buffers and writes the fused kernel arguments. - The listener hooks take the device list and the rotating slots that were launched. With one device they compute what they did before. With more than one, `BenchmarkTimer` reports the slowest card, `ReferenceValidator` keeps one CPU reference per card and also checks each card's receive buffer against the sending card's reference, and `SynchronizerValidator` checks every card's Synchronizer (the counter buffer StreamK and the fused kernel use) and flag block. - Fused launches now run back to back instead of clearing the counter, flag and receive buffers and synchronizing every GPU before each launch. The counter is each card's Synchronizer, which the kernel leaves at zero; consecutive launches alternate between two flag channels, so a peer one launch ahead never writes the channel a card is still reading; and each rotating slot has its own receive buffer, zeroed once per solution. User-visible changes on the fused path: - Results are standard rows with `time-us` (slowest card), `gflops` (one card's GEMM work; A2A traffic is not counted) and `Validation`. The exit code is the number of listener errors, as on the single-device path, instead of the fused-only codes 1, 2, 3 and -1. - With `--use-gpu-timer=False`, the time is wall-clock as on the single-device path and includes launching on and waiting for every card; the old client always used per-card GPU events. - A, B, C, alpha and beta come from `DataInitialization`, so `--init-c`, `--init-alpha` and `--init-beta` now apply. `--init-beta` defaults to `Two`, so a hand-written config without it now runs with beta = 2; `fused_a2a.yaml` sets `DataInitTypeBeta: 0`. - A solution the runner cannot use for this problem (tile shape, predicate, copy geometry) reports `INVALID` and the run moves on. As on the single-device path, this is not counted as an error, so a run in which every solution is rejected exits 0. - Startup now rejects settings the fused path cannot honor: validation with `--print-max` at or below 0 (a mismatch would otherwise be reported as `PASSED`; the default is -1), `--init-a`/`--init-b` other than `Random`, `--icache-flush-args`, `--use-user-args`, and `--fused-a2a-drain-recv=0` or `--fused-a2a-drain-send=0`. - With several cards, the hardware-monitor columns (temperature, clocks, fan) sample card 0 only. - If a launch fails on one card after its peers have launched, the peers wait without a timeout in the kernel's drain step, where each card waits for its sends and receives to finish, so the client hangs. Changes on every path: - With rotating buffers, the first single-device enqueue of a pass continues after the warmups' slots instead of restarting at slot 0. - A failed lazy-loading initialization now stops the client; the exception used to be constructed but never thrown. The kernel generator changes only in Python comments that named the deleted files; the generated assembly is unchanged. ## Test Plan All runs are on 4× MI355X (gfx950), comparing this branch with develop. The develop build is from `fbbc964569c`; the client, host library and their tests have not changed on develop since. Both builds include the fix from #13302, without which the NaN bounds-check mode copies no input data and its validation is meaningless. - Fused path, W = 4, on the fixture built from `Tensile/Tests/common/comm/gfx950/fused_a2a.yaml` (BF16, 2 problems × 2 solutions, one problem with N = 511 to hit a partial tile), full validation, Synchronizer check on (the default): - bounds check off, NaN, and NaN with rotating buffers, each with 1 launch per sync (20 syncs, 2 warmups) and back to back (10 launches per sync, 5 syncs, 3 warmups); - 3 rotating slots, 5 launches per sync, 3 syncs, so the kernel table cycles through all 6 slot and channel pairs; - beta = 2 with a random C; - `--icache-rotate-copies=3`, which loads 3 extra copies of the kernels and rotates launches across the 4 copies, once with bounds check off and once with 3 rotating slots and NaN. - Single-device path on a BF16 GEMM fixture (3 problems × 8 solutions), rotating buffers, 10 warmups, 50 launches per sync, full validation, bounds check off and NaN. - `tensilelite-tests`, the client's C++ unit tests, with 4 GPUs visible. The multi-card cases need at least 2 GPUs and skip otherwise. ## Test Result - Fused path, this branch: all 10 configurations report 4 of 4 rows `PASSED`, exit 0, with no receive-buffer mismatch and no Synchronizer or flag-block residue. Develop, run on the first 8, prints `overall PASSED` for all 4 in each; it forces beta to 0 and ignores rotating buffers and `--icache-rotate-copies`, so those variants only exercise this branch. - Single-device path: 24 of 24 `PASSED` in both modes, on develop and on this branch. - `tensilelite-tests`: develop 637 passed, 10 skipped; this branch 647 passed, 10 skipped. The 10 extra tests are the new ones; the same 10 tests skip in both builds. - Not covered: a GPU pair without P2P access, and a launch failing on one card. In an earlier run, forcing every launch onto flag channel 0 still passed 8,000 back-to-back launches, so these tests do not show that the second channel is needed. - After the review fixes (`db950d06ff3`): the fused comm test and `tensilelite-tests` pass on 4× MI355X, and with `--check-synchronizer=0` a solution that leaves the flag block or Synchronizer dirty no longer hangs the next one. - Automated GitHub checks are pending. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
archana-ramalingam
added a commit
that referenced
this pull request
Oct 29, 2026
…TensorVoid on CPU - WorkgroupMappingXCCCheck: the lazily queried CU count and processor were mutable std::optional, filled on first use from whichever thread evaluates the predicate first. Use atomics instead (processor stored before the count, released/acquired), so concurrent solution selection is safe. - Drop the MessagePack catch narrowing: with the device query out of the constructor, deserialization no longer throws device errors, and the narrowed catch let objectToMap's std::runtime_error escape the lazy loader. - Component CI gtest lane: exclude CopyTensorVoid (added by #13302), which needs a HIP device for hipMemcpyAsync. Co-Authored-By: Claude <noreply@anthropic.com>
newling
added a commit
that referenced
this pull request
Nov 5, 2026
…ling (#14315) JIRA ID : AIHPBLAS-4195 ## Motivation The hipBLASLt and TensileLite clients prepare matrix multiplication inputs and check GPU results. An isolated data-initialization test exposed a crash when optional settings were omitted: uninitialized values could enable features the caller had not requested. Buffer helpers could also accept invalid transfers or miss writes outside a buffer. On Windows, the client could continue after YAML conversion failed and use incomplete generated input. ## Technical Details Each change has regression coverage: - Initialize five optional feature settings to disabled in `ClientProblemFactory`, and default the bias source to output matrix D. Remove the unused `m_useSynchronizer` setting. - Add contiguous and padded-copy tests for `CopyTensorVoid`, including checks that padding remains untouched. The source-pointer correction has already landed in [#13302](#13302); this PR adds regression coverage. Both tests fail when the source address is computed from the destination, as it was before #13302, and pass with the current code. - Return `hipErrorInvalidValue` for zero transfer counts and padded copies whose destination column spacing is smaller than the row count. - Fill all 65536 bytes of the guard reference array with nonzero values for buffers whose data type is selected at runtime. Guards detect writes outside the buffer; the previous zero-filled tail hid zero writes beyond the first 8192 bytes. - Stop the Windows client when the Python YAML converter returns a nonzero status. Tests cover malformed input and a valid empty test list, checking only the exit status so that they do not depend on PyYAML's error wording. `TESTING.md` lists the new host unit test. The Linux YAML converter uses Python from the active environment, so it can find PyYAML installed by the test jobs. The standalone ASAN lane now prepares that environment. ## Test Plan Build the native clients, tests, and CPU matrix multiplication driver. Run regressions, adjacent tests, and AddressSanitizer/UndefinedBehaviorSanitizer checks. Verify that both YAML cases use the prepared Python environment. Existing shared test lanes select the regressions. ## Test Result After rebasing onto develop, including #13630's MX free-block options, the affected native targets build on a Radeon AI PRO R9700 (`gfx1201`). Focused regressions and adjacent factory, copy, tensor-descriptor, and MX data-initialization tests pass; two existing architecture-specific cases skip. The factory regression needs no MX free-block arguments: an omitted option reads as zero, and `setMXScaleA`/`setMXScaleB` raise it to the default of one. Both YAML cases pass with PyYAML 6.0.3. Formatting, root hooks, and whitespace checks pass. On the previous revision, nine focused sanitizer tests passed with the copies, buffer helpers, and production YAML parser instrumented. Linked libraries remained native; local suppressions cover separately reproduced GPU-runtime leaks and a monitoring-library startup error. [Native Windows CI](https://github.com/ROCm/rocm-libraries/actions/runs/40023459817/job/120734677969) passed the YAML, buffer, and guard regressions on an earlier revision. [Automated checks for this update](https://github.com/ROCm/rocm-libraries/pull/14315/checks) are pending. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (2/5): narrow client fixes and test-environment changes with regressions and valid-input controls.
cycheng
pushed a commit
that referenced
this pull request
Nov 6, 2026
…ling (#14315) JIRA ID : AIHPBLAS-4195 ## Motivation The hipBLASLt and TensileLite clients prepare matrix multiplication inputs and check GPU results. An isolated data-initialization test exposed a crash when optional settings were omitted: uninitialized values could enable features the caller had not requested. Buffer helpers could also accept invalid transfers or miss writes outside a buffer. On Windows, the client could continue after YAML conversion failed and use incomplete generated input. ## Technical Details Each change has regression coverage: - Initialize five optional feature settings to disabled in `ClientProblemFactory`, and default the bias source to output matrix D. Remove the unused `m_useSynchronizer` setting. - Add contiguous and padded-copy tests for `CopyTensorVoid`, including checks that padding remains untouched. The source-pointer correction has already landed in [#13302](#13302); this PR adds regression coverage. Both tests fail when the source address is computed from the destination, as it was before #13302, and pass with the current code. - Return `hipErrorInvalidValue` for zero transfer counts and padded copies whose destination column spacing is smaller than the row count. - Fill all 65536 bytes of the guard reference array with nonzero values for buffers whose data type is selected at runtime. Guards detect writes outside the buffer; the previous zero-filled tail hid zero writes beyond the first 8192 bytes. - Stop the Windows client when the Python YAML converter returns a nonzero status. Tests cover malformed input and a valid empty test list, checking only the exit status so that they do not depend on PyYAML's error wording. `TESTING.md` lists the new host unit test. The Linux YAML converter uses Python from the active environment, so it can find PyYAML installed by the test jobs. The standalone ASAN lane now prepares that environment. ## Test Plan Build the native clients, tests, and CPU matrix multiplication driver. Run regressions, adjacent tests, and AddressSanitizer/UndefinedBehaviorSanitizer checks. Verify that both YAML cases use the prepared Python environment. Existing shared test lanes select the regressions. ## Test Result After rebasing onto develop, including #13630's MX free-block options, the affected native targets build on a Radeon AI PRO R9700 (`gfx1201`). Focused regressions and adjacent factory, copy, tensor-descriptor, and MX data-initialization tests pass; two existing architecture-specific cases skip. The factory regression needs no MX free-block arguments: an omitted option reads as zero, and `setMXScaleA`/`setMXScaleB` raise it to the default of one. Both YAML cases pass with PyYAML 6.0.3. Formatting, root hooks, and whitespace checks pass. On the previous revision, nine focused sanitizer tests passed with the copies, buffer helpers, and production YAML parser instrumented. Linked libraries remained native; local suppressions cover separately reproduced GPU-runtime leaks and a monitoring-library startup error. [Native Windows CI](https://github.com/ROCm/rocm-libraries/actions/runs/40023459817/job/120734677969) passed the YAML, buffer, and guard regressions on an earlier revision. [Automated checks for this update](https://github.com/ROCm/rocm-libraries/pull/14315/checks) are pending. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (2/5): narrow client fixes and test-environment changes with regressions and valid-input controls.
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.
Motivation
With
--bounds-check=NaN, the TensileLite client fills every input and output buffer with aninfsentinel and then copies the real tensor into the middle, so that reads or writes outside the tensor show up. The copy never happened: every tensor kept the sentinel. The kernel and the CPU reference both computed D frominfinputs, so validation comparedinfagainstinfand passed whatever the kernel produced. The Synchronizer buffer also started asinfinstead of zero, so the Synchronizer clean-on-exit check reported StreamK solutions as dirty.JIRA ID : AIHPBLAS-4534
Technical Details
CopyTensorVoid(include/Tensile/hip/HipUtils.hpp) took its source address fromdst, so eachhipMemcpyAsynccopied the destination onto itself. It now reads fromsrc. Its only caller iscopyBadInputBuffers, the NaN bounds-check path; the Disable and GuardPage modes do not use it.A new
CopyTensorVoid.CopiesSourcecase intests/DataInit_test.cppcopies a small host tensor and checks that the destination matches the source.Test Plan
On one MI355X (gfx950), run the TensileLite client on a BF16 GEMM fixture (3 sizes x 8 solutions) with
--bounds-check=NaNand--bounds-check=Disable, with--print-valids, before and after the change. Runtensilelite-testswith and without the fix. Run the two common tests that use NaN bounds checking,nan_bounds_check_odd_padding.yamlanduse_beta_false.yaml, with the fixed client.Test Result
inf==inf; 12 of 24 FAILED, all from the Synchronizer check on StreamK solutions.tensilelite-tests: 635 passed, 10 skipped, 0 failed. Without the fix,CopyTensorVoid.CopiesSourcefails with the destination still all zeros.nan_bounds_check_odd_padding.yamlanduse_beta_false.yamlpass with the fixed client.Submission Checklist