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 #11413 +/- ##
===========================================
+ Coverage 70.78% 70.80% +0.02%
===========================================
Files 2842 2842
Lines 465020 465412 +392
Branches 68499 68572 +73
===========================================
+ Hits 329152 329527 +375
+ Misses 112174 112161 -13
- Partials 23694 23724 +30
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
8b62f1f to
51031ea
Compare
| { | ||
| RocblasltFusedEpilogueInfo fusedInfo; | ||
| if(rocblaslt_resolve_fused_epilogue(prob.fused_epilogue, fusedInfo) | ||
| && fusedInfo.hasA2APrefix) | ||
| { | ||
| tensileProblem.setFusedGemmA2A(true); | ||
| tensileProblem.setFusedA2AExtent(fusedInfo.a2aExtent); | ||
| tensileProblem.setFusedA2AWorld(prob.fused_a2a_world); | ||
| } | ||
| } |
There was a problem hiding this comment.
Nit: Same lines in updateTensileProblem as above? Worth writing a helper.
There was a problem hiding this comment.
These two functions already duplicate the whole translation verbatim, so extracting just the A2A block would leave it the one factored piece in an otherwise duplicated section.
There was a problem hiding this comment.
ConstructTensileProblem (lines 1827-2140) and updateTensileProblem (lines 2142-2427) are already near-copies of each other: 221 of the 314 lines in the first appear verbatim in the second, about 70%. The only structural difference is that one builds a fresh ContractionProblemGemm while the other overwrites an existing one. The fused all-to-all block is copied the same way, which follows the current style of other setters.
51031ea to
c2cc716
Compare
0b4a779 to
a0e4d53
Compare
49595f5 to
2ccdbe9
Compare
2ccdbe9 to
55c7e0b
Compare
55c7e0b to
b458981
Compare
b458981 to
f8129ca
Compare
b538dc7 to
5d2021a
Compare
The mode documented that ordinary stream semantics cover the collective, but its name and its value only spoke for receive completion. A caller that synchronized its stream and then overwrote D raced this card's SDMA engines, which may still be reading it. Rename the value to _FULL and give it the number the kernel's drain mask already uses for both directions, so the name and the value both say what the mode guarantees. The tensilelite client's drain-send default follows: leaving it off contradicted the mode it benchmarks. COMM_CHANNEL had the same gap in the other direction: reusing a channel needs every rank's previous operation on it to have retired, which no rank can establish from its own stream, and the caller's recv buffers carry the same requirement.
This is what is left of the commit that used to add the public API and the device communicator here. #11342 now provides hipblasLtSetDeviceComm, the A2A_PREFIX stage and every attribute, so the header additions, the communicator implementation and its peer-kind bookkeeping are all dropped; what remains is the rocblaslt-level plumbing that PR does not reach. Everything about such a request that is checkable before the heuristic runs goes in a header of its own; handle.h stays limited to the handle definitions and the helpers that resolve an effective value out of handle/desc/pref. JIRA ID : ROCM-27524
buildFusedA2APeerFields turns the resolved epilogue into the per-rank kernarg groups the kernel expects -- flag, recv and the four SDMA queue words per peer, indexed by Tensile's FusedA2APeerSlot so that order stays a single definition -- and the launch path fills rank, drain and the counter, which comes from the Synchronizer the handle already owns rather than a second allocation. FusedA2APeerFields pins that layout: the slot literals in the test are spelled out, so reordering or widening the enum fails there rather than silently downstream. ConstructTensileProblem and updateTensileProblem set fusedGemmA2A from the resolver. Without it the FusedGemmA2A predicate filters every A2A solution out of selection, and FusedA2ATileDivisible rejects a world of zero, so the world has to be carried too. An A2A launch consumes a counter block and a flag block that both have to start at zero, and nothing puts them back, so a second launch would find every flag already READY, wait on nothing and transfer nothing -- silently. A training loop hits that on its second call. The kernel re-arms them itself: it already elects one work-group per launch for the DRAIN barrier, and once that work-group clears a drain poll, the data that poll covers is finished with. The clears go there -- the receive flags after the receive poll, the outbound counter after the send poll, and the SDMA cursors on the way out. counter1/2/3 already wrap back to zero on their last increment and the cursors are the only part of the Synchronizer slice that does not, so the slice still returns to the pool zeroed. Re-arming in the kernel rather than with two hipMemsetAsync calls enqueued after launchKernels keeps two extra dispatches off the caller's stream. The cursor clear needs glc. s_store lands in the per-CU scalar cache, and without writing through, the next launch raised its cursors off a stale value and the SDMA ring deadlocked. No waitcnt is needed: s_endpgm implicitly executes S_WAITCNT 0. A request naming no SDMA queue set is now reported as no match instead of reaching the kernel: every fused A2A solution in the tree emits SDMA packets unconditionally, so null queue handles used to reach the kernarg groups and fault. hipblasLtMatmul rejects it as well, since a caller holding an algo from elsewhere would otherwise still get there. The heuristic reports it as zero algos rather than an error, matching what that function already does when selection yields nothing; both gates restore the dummy bias pointer they may have to return through. hipblaslt_ext::Gemm rejects the stage outright. That path has no reset after its launch and no place to check the per-rank arrays against the world size, so admitting an all-to-all there would answer a second run with stale flags and index past the descriptor's vectors. Leaving fused_a2a_world unset also keeps FusedA2ATileDivisible from matching, should a request reach selection another way. Measured on four gfx950 cards: 50 back-to-back launches verify (250 over five runs). JIRA ID : ROCM-27524
One sample over the public API at the customer shape (18432 features, 2048 tokens, extent 10240), parameterized over the deployment: kProcesses worker processes of kRanksPerProcess ranks each, with kWorld derived from them. That reaches the shapes callers spread over -- torchrun gives a process per card, a single-process launcher gives one process with a thread per card -- and static_asserts hold the pair to a world that divides the extent and fits the eight rank slots the fused-A2A ABI carries. At the default 2 x 2 one communicator carries a self peer, a same-process peer and two cross-process peers, and each kind of recv pointer resolves a different way. The sample checks every rank's recv buffer against every peer's D. The parent only coordinates: it forks one worker per process and holds a socket to each, so a star is the only topology the collectives need. Every collective has the same shape -- a worker contributes its kRanksPerProcess entries behind a byte-stride header, the parent concatenates them in rank order and hands the whole array back -- so the recv IPC handles, the registration allgather and the verdict share one pair of primitives. No A2A solution ships in the logic tree, so the sample needs a device library built for it. Producing one from the gfx950 A2A config would mean giving that config a LibraryLogic section, and the fused path skips the result reporters, so the stage gets an empty benchmark CSV and fails in LogicAnalyzer. Threading the fused path through the listener chain is AIHPBLAS-4534; until then the sample states the prerequisite and skips when the loaded library does not meet it. Each rank releases the barrier it shares on the way out, so one rank failing between the two rendezvous ends the run instead of leaving the others blocked; the descriptors and device buffers are owned rather than leaked, since a sample gets copied into code that loops. A worker that skips closes its socket before the first collective, so the coordinator's own exchange failing is only an error once every worker has claimed to have run; the child exit statuses are consulted first. COMM_CHANNEL selects the flag region. It was validated against nChannels and then dropped -- every channel aliased to the first block, so the guarantee that a channel covers its own flag region did not hold. The channel is a host-side pointer bias of channel * FLAG_BLOCK_BYTES; the kernel is unchanged. The per-rank attributes are passed as the arrays themselves, sized in whole entries, which is what hipblasLtFusedEpilogueSetAttribute expects; a pointer to the array sizes out to less than one entry and is rejected. JIRA ID : ROCM-27524
#11342 closed validate_fused_a2a_launch with an unconditional NOT_SUPPORTED, standing in for the kernels that did not exist yet. Those kernels land on this branch, so the placeholder rejected every request before rocblaslt was reached and both SDMA samples failed at hipblasLtMatmulAlgoGetHeuristic with it. What the function defers is still enforced, in two places rather than one. The shape-independent half -- MacroTile0/1 in {128, 256}, data-parallel only, no split-K, a BF16 D, gfx94x/95x -- is a generation-time solution rejection in Solution.py. The half that depends on the problem is the FusedA2ATileDivisible predicate, which requires extent % (MacroTile0 * world) == 0 and rejects a world of zero, and so runs at selection. Both also filter the size-independent getAllSolutions fallback, which is what keeps a shape with no Equality logic entry from picking up an algo that does not fit it. WellFormedRequestReportsMissingCapability asserted that placeholder. A replacement would have had three different correct answers - SUCCESS with an A2A device library, NOT_SUPPORTED with an ordinary one, and INVALID_VALUE with none - so it is dropped rather than pinned to one environment, and the file's header note is updated to match. Samples 30 and 31 cover that path end to end; the twelve rejection tests around it are untouched. JIRA ID : ROCM-27524
The sample already runs the customer shape, which makes it the natural place to hand users a reproducible latency number they can check on their own machines. The timing pass runs after the correctness check instead of reusing its launch. The closed form that check compares against requires A to be near-zero, which is not representative of real operand traffic, so the timing pass re-fills A and B from the same (x%7)-3 alphabet the tensilelite client draws from. That keeps the two harnesses comparable. IN_KERNEL_FULL completion compares each peer's accumulated arrival count for exact equality, so a rank that ran ahead could push a straggler's slot past tokenTiles and stall its DRAIN poll for good. The loop alternates the communicator's two flag regions, which rules that out structurally: reaching round N+2 requires completing N+1, which requires every peer to have sent its N+1 data, and hence to have finished N -- so the region a run-ahead rank writes is never one a straggler is still polling. Nothing on the host side has to coordinate the iterations, which also keeps socket wake-ups out of the measured window. Latency is reported as the max across ranks per iteration, matching what the tensilelite client reports. It is the one collective that stops at the parent, which reduces it to the single number the parent prints. The kernel re-arms its own A2A state before it retires, so that cost falls inside the measured window, which is correct because every caller pays it. The check now validates the last launch rather than the first. The timing loop leaves its own operands in place and never copies back, so hD and hRecv still held the opening launch and a counter block that went bad mid-loop would have gone unnoticed. The pass now re-fills the near-zero operands the closed form needs and lands a final launch into the buffers checkRank reads. Verified on a four-card hold at 2 x 2, 4 x 1 and 1 x 4, each checking every rank's recv buffer against every peer's D. 8 x 1 compiles but has not been run; it needs a node with eight free cards.
5d2021a to
4e0b054
Compare
FindNUMA.cmake and HsakmtLinkInterface.cmake exist to repair ROCm 7.x's hsakmt package config. Neither is specific to TensileLite, and both now have consumers on the hipBLASLt side. Moving them to the project's cmake/ directory puts them on the module path the root CMakeLists already exports, so find_package(NUMA) and include(HsakmtLinkInterface) resolve from any subdirectory without each caller appending to CMAKE_MODULE_PATH. That also retires the ordering constraint FindNUMA carried: it had to reach the module path before find_package(hsakmt), and nothing enforced it. hipblaslt_sanitize_hsakmt_link_interface now takes no arguments. It reads and rewrites hsakmt::hsakmt's link interface itself and resolves ROCM_PATH internally, so the get/if/call/set block both callers repeated verbatim collapses to one line and the ROCM_PATH fallback lives in one place rather than three. The function returns untouched when the target is absent or carries no link interface. The SDMA sample drops its hand-written libnuma search, which duplicated FindNUMA and looked for numa.h under include/rocm_sysdeps where ROCm does not install it.
`norm.hpp` and `allclose.hpp` define six non-template functions in the header without `inline`: the `void*` plus `hipDataType` dispatch overloads of `norm_check_general` and `allclose_check_general`, `norm_tolerance`, and the two `norm_check` type helpers. Two translation units that include either header emit the same symbols and the link fails. Nothing hit this because each binary pulled the headers into exactly one translation unit - `client.cpp` for the benchmark, `matmul_gtest.cpp` for the tests.
`hipblaslt-bench-a2a` measures the fused all-to-all GEMM path across a rank group started by an external launcher. Rank identity and rendezvous come from the `RANK` / `WORLD_SIZE` / `LOCAL_RANK` / `MASTER_ADDR` / `MASTER_PORT` variables torchrun sets; with none of them set the run is single-rank. The shipped gfx950 logic corpus gains the fused GEMM+A2A entry the benchmark resolves against, so a `--gemm-a2a-fusion` build needs no hand-built device library. It lands in the base `gfx950` tree and declares the default fallback chip ID, which `chipIdFallbacks` maps every gfx950 SKU onto, rather than a `gfx950_id<chip>` directory that would cover one SKU. A build without the fusion gate drops the entry in `generateLogicDataAndSolutions` and reports the filtered count. `hipblaslt-bench-a2a-launch` starts the rank group. It sets all five launcher variables for every rank, so a stale `WORLD_SIZE` in the caller's environment cannot turn a single-rank run into a failed rendezvous that still exits zero. A rank exiting non-zero terminates the survivors, which would otherwise stay blocked in the rendezvous, and each failing rank is reported with its log path instead of the run standing or falling on rank 0's code alone. CMake copies the script next to the benchmark in the build tree and installs it into bin/, so it finds the executable beside itself in either layout. - a TCP rendezvous exchanges IPC handles and recv pointers without MPI - `--adaptive` self-sizes the sample count, and a collective agreement keeps every rank on the same iteration count - latency is reported as the slowest rank, and the derived rate goes through `rate_per_second` so a non-positive time prints `-1` rather than dividing into `inf` - a rank's device resources are released on every path out `--verify` is off by default. When it is on, it runs `--iters` launches as its own pass before the timed batch instead of checking inside the timed callable, so the device-to-host copy and the comparison stay out of the reported latency. Every rank runs the full count even after a mismatch: a short run would leave that rank's launch counter behind the group, the ranks would then pick different flag regions, and the run would deadlock. Operands are HPL-random, which is what `hipblaslt-bench` initializes with by default, so the two clients' numbers are comparable - flip rate drives power draw, and on a power-limited part that reaches the reported latency through clock throttling. The seed is derived from the rank; identical operands would compare equal no matter which source's block landed, so deriving it is what makes a misrouted block visible. `--verify` therefore covers the transport and not the GEMM: the gold is the source rank's own `D`. Each rank reads back the block it exported into every peer's receive buffer and compares it against the rows of its local `D` that block should hold. An all-zero `D` fails outright, since the two sides would otherwise agree trivially. `HIPBLASLT_A2A_COMPLETION_IN_KERNEL_FULL` retires a kernel once that rank's receive buffer is full, which says nothing about this rank's outbound data having landed, so the check takes a rendezvous barrier first. The fixed-count path warms up for `--cold_iters` before the timed batch, so the first launch's lazy code-object load stays out of the sample; a cold load off a network filesystem costs seconds and, divided across the batch, silently multiplies the reported latency. `--adaptive` sizes its own warmup and skips that step, but a `--verify` pass still precedes it. Successive launches alternate over two of the communicator's flag regions, and two is the whole range the benchmark can need. Reaching round N+2 requires completing N+1, which requires every peer to have sent its N+1 data and hence finished N, so ranks stay within one round of each other and the region a run-ahead rank writes is never one a straggler is still polling. Every agreement is one allgather over a caller-supplied reduction, so doubles and flags share one code path instead of one typed callback each. A failed allgather throws rather than answering with the local value, which would leave ranks on different iteration counts. Ranks agree on the default-path launch outcome only after the kernel has been dispatched and completed; a failure before dispatch still hangs, because a peer blocked in `hipStreamSynchronize` never reaches the allgather. `read_launcher_env` clamps `WORLD_SIZE` to at least 1. A zero or non-numeric value used to reach `same_host_group()`, which sizes its result buffer by the world size and then copies a hostname into the empty buffer. `peers_reachable` rejects a multi-rank run whose `LOCAL_RANK` differs from its `RANK`. The peer pre-check indexes devices by `LOCAL_RANK` while the recv-pointer exchange indexes by `RANK`, so a launcher that leaves `LOCAL_RANK` unset puts every rank on device 0 and the two disagree with nothing to catch it unless `--verify` is on. Two places do not reuse an existing client helper: - `parse_a2a_args` handles its ten value options directly rather than pulling in the `boost::program_options` emulation in `bench/include/program_options.hpp` - `RankResources` allocates with `hipMalloc` instead of `device_vector`, because the receive buffer has to export an IPC handle `check_recv` goes through `norm_check_general` and `allclose_check_general`, the same two measures `hipblaslt-bench --verify` turns on, but decides pass and fail itself: `client.cpp` clears `norm_check_assert` and discards the `allclose` result, so consuming them the way it does would leave `--verify` printing diagnostics without gating anything. Requiring a zero norm error is exact rather than approximate, since both sides hold the same bf16 values. `unit_check_general` and `near_check_general` stay out because they expand to nothing unless `GOOGLE_TEST` is defined, which would check nothing here. The A2A SDMA sample leaves timing to the benchmark and picks up the hsakmt, hsa-runtime64 and libnuma targets from the enclosing clients scope, which the benchmark's dependency block resolves before the samples subdirectory is added. gtest coverage: the rendezvous over loopback sockets, including a payload past the socket buffer so the segmented send and receive loops run; the collective agreement's reduction and its failure path; the `Arguments` layout for the added fields; the single-rank benchmark path; and the `LOCAL_RANK` rejection. The suites carry the `smoke` category token so ctest's quick tier and the pull-request gate's `*smoke*` filter select them. The three tests that exercise the fusion build only with `HIPBLASLT_ENABLE_GEMM_A2A_FUSION=ON`, which defaults off and which no lane sets, so those do not run in a shared lane yet. `collective_rendezvous_gtest.cpp` uses POSIX sockets and is excluded on Windows.
| for(uint32_t j = 0; j < arg.a2a_world; ++j) | ||
| { | ||
| res.ownedQueues.push_back(std::make_unique<TensileLite::Client::SdmaQueue>( | ||
| srcNode, TensileLite::Client::sdmaSelectEngine(srcNode, srcNode))); |
There was a problem hiding this comment.
This selects the engine for srcNode -> srcNode for every peer. The public attribute defines entry j as the queue targeting rank j, and the sample resolves a separate destination node before calling sdmaSelectEngine. On a topology whose peer link recommends a different engine, this benchmark can use the wrong copy engine. Please resolve dstNode from j, call sdmaSelectEngine(srcNode, dstNode), and cover the multi-rank case.
| { | ||
| return [&res, &heur, &launchCount, &lastStatus](int64_t) { | ||
| const float alpha = 1.0f, beta = 0.0f; | ||
| const uint32_t channel = launchCount++ % kCommChannels; |
There was a problem hiding this comment.
The benchmark has two channels, but run_measurement can enqueue the default batch of 10 launches without a group-wide barrier. This reuses channel 0 on launch 3 before every rank is known to have retired launch 1. The public header explicitly says that local stream synchronization does not establish safe channel reuse. Please prevent reuse until all ranks agree that the earlier launch retired, or provide a distinct channel for every in-flight launch.
There was a problem hiding this comment.
Two channels are enough here because the alternation is interlocked across ranks, rather than relying on a stream synchronization.
Under HIPBLASLT_A2A_COMPLETION_IN_KERNEL_FULL a rank's kernel does not retire until its own receive buffer is full, and every rank issues its launches on a single stream. So for rank i to start launch k+2, launch k+1 must have retired on i; that requires every peer to have pushed its k+1 shard, which requires every peer to have started k+1, which by stream order requires every peer to have finished k. Launch k and launch k+2 are the pair that share a channel, so by the time i returns to that channel every rank has retired its previous operation on it.
Rebuilding the benchmark with kCommChannels = 1 deadlocks at 8 ranks under --iters 50 --verify 1, with no rank producing output before the run is killed; the two-channel build completes the same run with every launch verified.
| if(rocblaslt_resolve_fused_epilogue(prob.fused_epilogue, fusedInfo) | ||
| && fusedInfo.hasA2APrefix) | ||
| { | ||
| inputs.fusedA2ACounter = prob.Synchronizer; |
There was a problem hiding this comment.
Every fused-A2A launch receives the same handle-wide Synchronizer address. commChannel biases the peer flag blocks, but it does not separate the SDMA cursors or counter1/2/3 storage. Two operations on different streams can therefore race even when they use distinct channels and queue sets, contrary to the documented concurrency rule. Please allocate or derive counter storage per channel or stream, or reject that concurrent use explicitly.
| // kernel in this release. | ||
| log_error(__func__, "no fused all-to-all implementation for the selected device"); | ||
| return HIPBLAS_STATUS_NOT_SUPPORTED; | ||
| return HIPBLAS_STATUS_SUCCESS; |
There was a problem hiding this comment.
Before this returns success, can we revalidate every receive pointer and all four fields of every SDMA queue? The matmul descriptor references the fused descriptor rather than copying it, and SetAttribute can replace these arrays after attachment. The current launch check verifies only their lengths, so a post-attachment null entry reaches kernel argument construction and can fault on the device.
| ) | ||
| endif() | ||
|
|
||
| if(HIPBLASLT_ENABLE_GEMM_A2A_FUSION AND NOT WIN32) |
There was a problem hiding this comment.
These tests are compiled only when the feature is enabled, while the option defaults to OFF and the PR says no shared automated lane enables it. This change adds public launch behavior and device-side state resets, so the enabled path needs a shared gfx950 run; manual results alone do not satisfy the hipBLASLt regression-test requirement.
| comment="some peer incomplete -> spin (poll again)")) | ||
|
|
||
| # Still one lane per slot, and vPollOff still holds lane j's byte offset. | ||
| module.add(VMovB32(dst=vgpr(vPollVal), src=0, comment="fused-A2A: zero for the flag clear")) |
There was a problem hiding this comment.
This changes emitted synchronization behavior by clearing state for the next launch, but the diff adds no focused TensileLite unit test or saved expected-result test for the generated flag, outbound-counter, and cursor clears. Please add a test that fails when these instructions or their ordering are removed, then run it in the shared feature-enabled gfx950 lane.
|
|
||
| if(HIPBLASLT_ENABLE_GEMM_A2A_FUSION AND NOT WIN32) | ||
| rocm_install( | ||
| PROGRAMS "${CMAKE_CURRENT_SOURCE_DIR}/clients/scripts/a2a/hipblaslt-bench-a2a-launch" |
There was a problem hiding this comment.
This installs the launcher, but hipblaslt-bench-a2a is absent from the rocm_install(TARGETS ...) list above. In an installed tree the script therefore cannot find the executable beside itself as the PR description promises. Please install the benchmark target under the same condition and add an installed-tree smoke check.
| return 1; | ||
|
|
||
| hipblasLtMatmulHeuristicResult_t heur{}; | ||
| const uint8_t found = select_algo(res, heur) ? 1 : 0; |
There was a problem hiding this comment.
select_algo returns false both when the heuristic succeeds with zero algorithms and when hipblasLtMatmulAlgoGetHeuristic returns an error. The block below converts both cases into a successful skipped exit. Please preserve the API status separately so an invalid request or internal failure returns non-zero instead of being reported as an unsupported configuration.
| @@ -0,0 +1,497 @@ | |||
| MinimumRequiredVersion: 5.0.0 | |||
There was a problem hiding this comment.
new YAML files under projects/hipblaslt must start with the short copyright and SPDX-License-Identifier: MIT header required by projects/hipblaslt/AGENTS.md. Please add it above MinimumRequiredVersion.
There was a problem hiding this comment.
These files are generated rather than hand-written: none of the 2,797 logic YAML files under Tensile/Logic/ carries the header.
| @@ -0,0 +1,237 @@ | |||
| // Copyright Advanced Micro Devices, Inc., or its affiliates. | |||
There was a problem hiding this comment.
I have the same critique of this one-off script as for the previous PR: #10925 (comment)
There was a problem hiding this comment.
This design did take your earlier comment into account, and the benchmark reuses the common client layer directly: hipblaslt_init_hpl for operand setup (a2a_bench.hpp:202), norm_check_general and allclose_check_general for the host comparison (:264, :269), run_measurement for the timing loop (client_a2a.cpp:157, the same function client_sequence.cpp:658 uses), and the shared Arguments struct throughout. What a2a_bench.hpp adds on top is the part with no counterpart in the common layer: process rendezvous, peer reachability, receive-pointer exchange, and per-rank SDMA queue setup.
A separate executable is also the existing shape here. clients/CMakeLists.txt already builds seven specialized benchmarks from clients/bench/src/ — hipblaslt-api-overhead, hipblaslt-bench-groupedgemm-fixed-mk, the four extop ones, and hipblaslt-sequence — each an add_executable linking hipblaslt-clients-common. hipblaslt-bench-a2a is the eighth. Further fused-A2A benchmarking (e.g. a2a then gemm) is meant to extend this binary rather than add another.
It is not a flag on hipblaslt-bench because this benchmark needs one process per GPU and a launcher to start them, which does not fit that client's single-process problem loop.
3 participants
JIRA ID : ROCM-27524
Depends on #11342.
Motivation
The public host API and the kernels for a fused GEMM + all-to-all epilogue are both in place.
This PR joins them: it carries the A2A attributes from the C-API layer down to the launch path,
builds the per-rank kernarg groups the kernel reads, marks the Tensile problem so that A2A
solutions survive selection, and opens
validate_fused_a2a_launchso a validated requestreaches the heuristic.
Measuring that path takes more than one process.
hipblaslt-benchdrives a single device, andthe collective exists only across a group of ranks. This PR therefore also adds
hipblaslt-bench-a2a, a benchmark that runs the fused path across such a group, and ships thegfx950 kernel-selection entry it resolves against.
Technical Details
Descriptor. The fused-epilogue descriptor is resolved into a POD at the C-API layer, where
its fields are visible, and carried on
RocblasltContractionProblemtogether with the handle'scommunicator view. The checks that can run before the heuristic live in
rocblaslt_fused_a2a_validate.hpp.Kernarg groups.
buildFusedA2APeerFieldsturns the resolved epilogue into the per-rankgroups — flag, recv, and the four SDMA queue words per peer — indexed by Tensile's
FusedA2APeerSlot, so that slot order stays a single definition. The launch path fills rank,drain, and the counter, which comes from the
Synchronizerthe handle already owns.Selection.
ConstructTensileProblemandupdateTensileProblemsetfusedGemmA2Aand theworld size. What the opened placeholder stood for is enforced in two places: the
shape-independent half — MacroTile0/1 in {128, 256}, data-parallel only, no split-K, a BF16 D,
gfx94x/95x — as a generation-time solution rejection in
Solution.py, and theproblem-dependent half as the
FusedA2ATileDivisiblepredicate, which requiresextent % (MacroTile0 * world) == 0and a non-zero world. Both also filter thesize-independent
getAllSolutionsfallback.Repeat launches. The kernel re-arms the counter and flag blocks it consumed before it
retires, so the next launch on the same handle starts from zero without extra dispatches on the
caller's stream. The work-group already elected for the DRAIN barrier does the clears: receive
flags after the receive poll, the outbound counter after the send poll, SDMA cursors on the way
out. The cursor clear writes through (
glc);s_storeotherwise lands in the per-CU scalarcache and the next launch raises its cursors off a stale value.
Completion.
HIPBLASLT_A2A_COMPLETION_IN_KERNEL_FULLcarries the number the kernel's drainmask already uses for both directions, so the kernel does not retire until this rank's receive
buffer is populated and its engines have finished reading
D. Synchronizing the caller's ownstream is therefore enough to reuse
D. It is not enough to reuse a communicator channel(
HIPBLASLT_FUSED_EPILOGUE_COMM_CHANNEL) or the receive buffers: both need every rank'sprevious operation on that channel to have retired, which no rank can establish from its own
stream.
Build gate. All of the above sits behind
HIPBLASLT_ENABLE_GEMM_A2A_FUSION(default OFF),matching the gate introduced in #11342.
Gates. A request naming no SDMA queue set is reported as no match, since every fused A2A
solution in the tree emits SDMA packets unconditionally. The heuristic reports it as zero algos
rather than an error, matching what that function already does when selection yields nothing,
and
hipblasLtMatmulapplies the same gate for a caller holding an algo from elsewhere.Scope.
hipblaslt_ext::Gemmrejects the A2A stage. That path has no launch-time reset andno place to check the per-rank arrays against the world size, so it is left for a follow-up.
Benchmark.
hipblaslt-bench-a2ameasures the fused path across a group of ranks started byan external launcher. Rank identity and the coordinating rank's address come from the
RANK,WORLD_SIZE,LOCAL_RANK,MASTER_ADDRandMASTER_PORTvariables that PyTorch'storchrunsets; with none of them set the run is single-rank. Ranks exchange inter-process memory handles
and receive-buffer addresses over a TCP rendezvous rather than MPI, so the benchmark needs
nothing the library does not already need. Latency is reported as the slowest rank's.
--verifyis off by default. When it is on, it runs its own pass before the timed batch, so thedevice-to-host copy and the comparison stay out of the measurement. It covers the transport
rather than the GEMM: the reference is the source rank's own
D, which each rank reads back outof every peer's receive buffer.
Launcher.
hipblaslt-bench-a2a-launchstarts the group. It sets all five variables forevery rank, so a stale
WORLD_SIZEin the caller's environment cannot turn a single-rank runinto a failed rendezvous that still exits zero. A rank exiting non-zero terminates the others,
and each failing rank is reported with its log path. CMake copies the script next to the
benchmark in the build tree and installs it into
bin/, so it finds the executable besideitself in either layout.
Shipped kernel selection. TensileLite, the kernel generator, picks a kernel per problem size
from selection files shipped with the library. The gfx950 set gains the fused GEMM + A2A entry
the benchmark resolves against, so a build with the fusion gate on needs no hand-built device
library. The entry lands in the base
gfx950tree and declares the default fallback chipidentifier, which every gfx950 part maps onto, rather than a per-part directory that would cover
one. A build without the gate drops the entry and reports the filtered count. Three problem
sizes are covered and any other size still reports no match, and the kernel recorded for each
was chosen from synthetic data rather than from a tuning run on hardware.
Two supporting changes.
norm.hppandallclose.hppdefine six non-template functionswithout
inline, so two translation units that include either header emit the same symbols andthe link fails. Nothing hit this because each binary pulled the headers into exactly one
translation unit, and the benchmark is the first to break that. Both headers sit in
clients/common/include/, shared withhipblaslt-benchandhipblaslt-test. Separately,FindNUMA.cmakeandHsakmtLinkInterface.cmakemove fromtensilelite/cmake/to the project'scmake/, which the rootCMakeLists.txtalready exports, so the samples and the benchmarkresolve them without each caller extending the module path.
Test Plan
FusedA2A*descriptor and attribute rejection matrix and the kernarg slotlayout, plus the benchmark's host-side pieces: the rendezvous, the collective agreement, the
Argumentslayout, and the single-rank path. Runs in CI, and gives the same answer whether ornot an A2A device library is present.
One sample at the customer shape, parameterized over processes and ranks-per-process, so the
same source covers a process per card, one process with a thread per card, and the mixed
deployment that also exercises cross-process peer resolution. It checks every rank's recv
buffer against every peer's D.
hipblaslt-bench-a2a --verifymakes the same comparison fromits own launcher. Run manually, since CI has no 4-GPU runner.
HIPBLASLT_ENABLE_GEMM_A2A_FUSION, since the default-OFF build is what CIcompiles. The library and its gtest are built at every commit with the gate OFF, so the
branch bisects. Three of the benchmark's gtests build only with the gate ON, so no lane runs
them today.
Test Result
with the gate OFF.
over 50 back-to-back launches repeated across five runs.
hipblaslt-bench-a2aon gfx950 against the shipped selection entry, with no device-librarypath set: pass at 1 and 2 ranks,
--verifyreporting no mismatch. Larger rank counts areuntested.
Submission Checklist