feat(hipblaslt): add host API for the fused GEMM + all-to-all epilogue by yzhang93 · Pull Request #11342 · ROCm/rocm-libraries

yzhang93 · 2026-08-27T00:29:17Z

Motivation

The fused GEMM + all-to-all epilogue redistributes the leading run of D's feature dimension across the ranks of a device communicator from the GEMM's own store path, removing a separate collective pass. This lands the host side ahead of the TensileLite kernel so callers can build and validate a request, and so the family's enum and attribute ranges are reserved before the other fused-epilogue families arrive.

JIRA ID : ROCM-27524

Technical Details

Public API: HIPBLASLT_FUSEABLE_EPILOGUE_A2A_PREFIX (7, leaving 0-6 to the chainable families) and its five attributes (12-16), plus hipblasLtSdmaQueue_t, hipblasLtA2ACompletionMode_t, and the shared hipblasLtFusedEpilogueCreate / Add / SetAttribute / Destroy builder. The builder documents family-neutral rules; per-family specifics stay on the enums.
hipblasLtSetDeviceComm registers a per-handle view of a device communicator, allocating the fine-grained flag regions and resolving peer addresses through a caller-supplied allgather, with IPC mapping for out-of-process peers. The flag-block layout is twinned with FUSED_A2A_FLAG_BLOCK_BYTES in the kernel's FusedA2AKernArg.hpp.
The descriptor attaches through HIPBLASLT_MATMUL_DESC_FUSED_EPILOGUE (106, after develop's UNIFORM_SUMMATION_ORDER_EXT) and is carried by _rocblaslt_matmul_desc as a non-owning pointer.
Validation is split three ways: stage composition at Add, required attributes at attach, and communicator, shape, and layout constraints before the heuristic and again at launch. No kernel is wired, so a well-formed request reports HIPBLAS_STATUS_NOT_SUPPORTED.

Test Plan

Build hipblaslt-test and run --gtest_filter='FusedA2A*' on gfx950, covering descriptor lifecycle, builder composition rules, attribute range and width checks, communicator registration, attach-time completeness, and each dispatch-time rejection.

Test Result

39 tests across 6 suites pass on gfx950. Nothing dispatches a GEMM: a valid request reports NOT_SUPPORTED, confirming the validation chain runs to completion.

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

therock-pr-bot · 2026-08-27T00:34:33Z

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

therock-pr-bot · 2026-08-27T00:34:34Z

🎉 All checks passed! This PR is ready for review.

codecov · 2026-08-27T02:41:42Z

Codecov Report

✅ All modified and coverable lines are covered by tests.

Additional details and impacted files

@@           Coverage Diff            @@
##           develop   #11342   +/-   ##
========================================
  Coverage    70.29%   70.29%           
========================================
  Files         2812     2812           
  Lines       462769   462779   +10     
  Branches     68122    68124    +2     
========================================
+ Hits        325271   325281   +10     
  Misses      113943   113943           
  Partials     23555    23555

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	`46.40% <ø> (ø)`
TensileLite-Unit	`76.10% <100.00%> (+<0.01%)`	⬆️
hipBLAS	`90.62% <ø> (ø)`		Carriedforward from 8011ff3
hipBLASLt	`35.24% <ø> (ø)`		Carriedforward from 8011ff3
hipCUB	`82.68% <ø> (ø)`		Carriedforward from 8011ff3
hipDNN	`87.01% <ø> (ø)`		Carriedforward from 8011ff3
hipFFT	`43.07% <ø> (ø)`		Carriedforward from 8011ff3
hipRAND	`76.12% <ø> (ø)`		Carriedforward from 8011ff3
hipSOLVER	`68.92% <ø> (ø)`		Carriedforward from 8011ff3
hipSPARSE	`86.99% <ø> (ø)`		Carriedforward from 8011ff3
rocBLAS	`48.31% <ø> (ø)`		Carriedforward from 8011ff3
rocFFT	`47.88% <ø> (ø)`		Carriedforward from 8011ff3
rocRAND	`57.42% <ø> (ø)`		Carriedforward from 8011ff3
rocSOLVER	`76.83% <ø> (ø)`		Carriedforward from 8011ff3
rocSPARSE	`74.61% <ø> (ø)`		Carriedforward from 8011ff3
rocThrust	`91.60% <ø> (ø)`		Carriedforward from 8011ff3

*This pull request uses carry forward flags. Click here to find out more.

Files with missing lines	Coverage Δ
...ts/hipblaslt/library/include/hipblaslt/hipblaslt.h	`25.00% <ø> (ø)`
...cts/hipblaslt/library/src/amd_detail/hipblaslt.cpp	`63.50% <ø> (ø)`
...rary/src/amd_detail/rocblaslt/src/include/handle.h	`72.83% <ø> (ø)`
...c/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp	`68.40% <ø> (ø)`
...rary/src/amd_detail/rocblaslt/src/tensile_host.cpp	`41.84% <ø> (ø)`
...t/library/src/amd_detail/rocblaslt/src/utility.cpp	`28.64% <ø> (ø)`
...slt/tensilelite/Tensile/Common/GlobalParameters.py	`95.79% <ø> (ø)`
...ite/Tensile/TensileCreateLibrary/ParseArguments.py	`100.00% <100.00%> (ø)`
...lt/tensilelite/Tensile/TensileCreateLibrary/Run.py	`88.29% <100.00%> (+0.15%)`	⬆️

🚀 New features to boost your workflow:

❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

A peer in another process gets its flag region through hipIpcOpenMemHandle with hipIpcMemLazyEnablePeerAccess, but a peer in this process had its raw device pointer stored and nothing else, leaving another device's memory unreachable when the copy engine writes to it. The ordinal that hipDeviceEnablePeerAccess needs was not in the exchange payload, so it travels there now. Already-enabled is a success: in a single process the W handles enable each other pairwise.

DDEle

This reads well and the API surface matches what we agreed on. Two comments below. LGTM for other parts.

#11342 defines hipblasLtFusedEpilogueDescriptor inside amd_detail/hipblaslt.cpp and hands rocblaslt only an opaque pointer, so the launch path cannot read the A2A attributes it needs. Resolve the descriptor into a POD at the C-API layer, where its fields are visible, and carry that plus the handle's communicator view on RocblasltContractionProblem. 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. 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. The macro-tile requirements the function defers to are enforced by the kernel generator as solution rejections, which is what its own comment already describes. 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

#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

#11342 defines hipblasLtFusedEpilogueDescriptor inside amd_detail/hipblaslt.cpp and hands rocblaslt only an opaque pointer, so the launch path cannot read the A2A attributes it needs. Resolve the descriptor into a POD at the C-API layer, where its fields are visible, and carry that plus the handle's communicator view on RocblasltContractionProblem. 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

#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

## Motivation Review of #11342 found the two peer branches of hipblasLtSetDeviceComm asymmetric: a peer in another process is mapped with hipIpcMemLazyEnablePeerAccess, while a peer in this process had only its raw device pointer recorded. A pointer from this process is not addressable on another device until peer access is enabled, so the copy engine would fault on a region it was handed. The same review found the communicator state on the handle holding two of its values twice. ## Technical Details - DeviceCommExchange carries the contributing rank's device ordinal, which is what hipDeviceEnablePeerAccess names. A same-process peer on another device gets peer access enabled, already-enabled counting as success since the W handles enable each other pairwise. A peer on this device needs nothing, and HIP rejects a request to reach a device from itself. - device_comm_flags is dropped: the handle's own allocation is entry device_comm_rank of device_comm_peer_flags. release_device_comm tells the three kinds of entry apart by index and by mapped, so it frees its own region, closes its IPC mappings, and leaves a same-process peer's pointer to the handle that allocated it. - device_comm_registered is dropped for device_comm_world != 0. Registration rejects a world below one and publishes world last, so nonzero already means what the bool did. ## Test Plan FusedA2ACommPeers.TwoRanksInOneProcessRegister: two ranks on two handles rendezvousing through a real allgather, which is the only way more than one rank can register in a single process. Where a second device exists the ranks take one each, so the peer branch has to make the region reachable rather than merely record its address; destroying both handles then exercises the ownership split. ## Test Result 40 tests across 7 suites pass on gfx950, with aux_handle_test, aux_attr_test, and aux_ext_test unaffected. The test machine has one device, so the cross-device hipDeviceEnablePeerAccess call itself is reached only where a second is present. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: Ding, Yi <yi.ding@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com>

…used-A2A API base

#11342 defines hipblasLtFusedEpilogueDescriptor inside amd_detail/hipblaslt.cpp and hands rocblaslt only an opaque pointer, so the launch path cannot read the A2A attributes it needs. Resolve the descriptor into a POD at the C-API layer, where its fields are visible, and carry that plus the handle's communicator view on RocblasltContractionProblem. 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

#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

Review of #11342 found the two peer branches of hipblasLtSetDeviceComm asymmetric: a peer in another process is mapped with hipIpcMemLazyEnablePeerAccess, while a peer in this process had only its raw device pointer recorded. A pointer from this process is not addressable on another device until peer access is enabled, so the copy engine would fault on a region it was handed. The same review found the communicator state on the handle holding two of its values twice. - DeviceCommExchange carries the contributing rank's device ordinal, which is what hipDeviceEnablePeerAccess names. A same-process peer on another device gets peer access enabled, already-enabled counting as success since the W handles enable each other pairwise. A peer on this device needs nothing, and HIP rejects a request to reach a device from itself. - device_comm_flags is dropped: the handle's own allocation is entry device_comm_rank of device_comm_peer_flags. release_device_comm tells the three kinds of entry apart by index and by mapped, so it frees its own region, closes its IPC mappings, and leaves a same-process peer's pointer to the handle that allocated it. - device_comm_registered is dropped for device_comm_world != 0. Registration rejects a world below one and publishes world last, so nonzero already means what the bool did. FusedA2ACommPeers.TwoRanksInOneProcessRegister: two ranks on two handles rendezvousing through a real allgather, which is the only way more than one rank can register in a single process. Where a second device exists the ranks take one each, so the peer branch has to make the region reachable rather than merely record its address; destroying both handles then exercises the ownership split. 40 tests across 7 suites pass on gfx950, with aux_handle_test, aux_attr_test, and aux_ext_test unaffected. The test machine has one device, so the cross-device hipDeviceEnablePeerAccess call itself is reached only where a second is present. - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: Ding, Yi <yi.ding@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com>

…ed-A2A API base

#11342 defines hipblasLtFusedEpilogueDescriptor inside amd_detail/hipblaslt.cpp and hands rocblaslt only an opaque pointer, so the launch path cannot read the A2A attributes it needs. Resolve the descriptor into a POD at the C-API layer, where its fields are visible, and carry that plus the handle's communicator view on RocblasltContractionProblem. 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

davidd-amd · 2026-09-11T16:00:18Z

HIPBLASLT_HAS_GEMM_A2A_FUSION is used to conditionally compile this feature in the installed public header, descriptor layout, attribute dispatch, and implementation files. Consequently, enabled and disabled builds expose different symbols and enum values, while every future change must preserve both configurations.

Are we required to exclude this feature at compile time? If not, compile it unconditionally and return HIPBLAS_STATUS_NOT_SUPPORTED when device support is unavailable. Keep the build option only for device-code generation.

If the implementation must remain optional, isolate it behind a stable interface and select enabled or stub translation units at build time. The disabled implementation should return NOT_SUPPORTED.

davidd-amd · 2026-09-14T18:09:51Z

@talumbau

This change makes hipBLASLt own communication-transport work in addition to the GEMM-specific flag layout. That code becomes a long-term maintenance responsibility of the hipBLASLt team. rocSHMEM exposes the corresponding primitives today for one processing element per process: bootstrap, symmetric allocation, peer-address lookup, synchronization, and release. RCCL develop also exposes symmetric-window peer-pointer APIs.

I implemented the rocSHMEM path behind a separate C ABI adapter against the current PR head (ec9fb614ef2) and reran the evidence on 8 MI210 GPUs. The run passed peer-pointer all-to-all at 2, 4, and 8 ranks; four communicator-registration cases; seven ported communicator cases from the PR plus a repeated-registration case; descriptor and heuristic validation at 2, 4, and 8 ranks; and a peer hipMemcpyAsync visibility probe. The hipblaslt.cpp compile command contained no -fgpu-rdc; libhipblaslt.so contained no rocSHMEM symbols and recorded liba2a_transport.so as a runtime dependency. This demonstrates that the transport can live outside hipBLASLt; it does not establish production equivalence. No fused kernel exists yet, the hipBLASLt build does not build or install the adapter, the test that asserted hipDeviceEnablePeerAccess was replaced with a two-device registration test, and the tested rocSHMEM implementation supports one processing element per process plus collective allocation/free and process-global lifetime.

Current PR:

// Error handling and logging omitted.
if(j == rank || peer.pid == mine.pid)
{
    if(j != rank && peer.device != mine.device)
        hipDeviceEnablePeerAccess(peer.device, 0);

    h->device_comm_peer_flags[j]        = peer.flags;
    h->device_comm_peer_flags_mapped[j] = false;
}
else
{
    void* mapped = nullptr;
    hipIpcOpenMemHandle(&mapped, peer.ipc, hipIpcMemLazyEnablePeerAccess);
    h->device_comm_peer_flags[j]        = mapped;
    h->device_comm_peer_flags_mapped[j] = true;
}

External rocSHMEM adapter:

// Error handling omitted.
rocshmem_uniqueid_t uid{};
rocshmem_get_uniqueid(&uid); // rank 0 obtains it, then distributes it
rocshmem_init_attr_t attr{};
rocshmem_set_attr_uniqueid_args(rank, world, &uid, &attr);
rocshmem_init_attr(ROCSHMEM_INIT_WITH_UNIQUEID, &attr);

void* localFlags = rocshmem_malloc(bytes); // collective symmetric allocation
hipMemset(localFlags, 0, bytes);
rocshmem_barrier_all();

for(int peer = 0; peer < world; ++peer)
    peerFlags[peer] = rocshmem_ptr(localFlags, peer);

// On coordinated teardown:
rocshmem_free(localFlags); // also collective

Review of #11342 found the two peer branches of hipblasLtSetDeviceComm asymmetric: a peer in another process is mapped with hipIpcMemLazyEnablePeerAccess, while a peer in this process had only its raw device pointer recorded. A pointer from this process is not addressable on another device until peer access is enabled, so the copy engine would fault on a region it was handed. The same review found the communicator state on the handle holding two of its values twice. - DeviceCommExchange carries the contributing rank's device ordinal, which is what hipDeviceEnablePeerAccess names. A same-process peer on another device gets peer access enabled, already-enabled counting as success since the W handles enable each other pairwise. A peer on this device needs nothing, and HIP rejects a request to reach a device from itself. - device_comm_flags is dropped: the handle's own allocation is entry device_comm_rank of device_comm_peer_flags. release_device_comm tells the three kinds of entry apart by index and by mapped, so it frees its own region, closes its IPC mappings, and leaves a same-process peer's pointer to the handle that allocated it. - device_comm_registered is dropped for device_comm_world != 0. Registration rejects a world below one and publishes world last, so nonzero already means what the bool did. FusedA2ACommPeers.TwoRanksInOneProcessRegister: two ranks on two handles rendezvousing through a real allgather, which is the only way more than one rank can register in a single process. Where a second device exists the ranks take one each, so the peer branch has to make the region reachable rather than merely record its address; destroying both handles then exercises the ownership split. 40 tests across 7 suites pass on gfx950, with aux_handle_test, aux_attr_test, and aux_ext_test unaffected. The test machine has one device, so the cross-device hipDeviceEnablePeerAccess call itself is reached only where a second is present. - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: Ding, Yi <yi.ding@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com>

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

#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 fused GEMM + all-to-all epilogue redistributes the leading run of D's feature dimension across the ranks of a device communicator from the GEMM's own store path, removing a separate collective pass. This lands the host side ahead of the TensileLite kernel so callers can build and validate a request, and so the family's enum and attribute ranges are reserved before the other fused-epilogue families arrive. - Public API: `HIPBLASLT_FUSEABLE_EPILOGUE_A2A_PREFIX` (7, leaving 0-6 to the chainable families) and its five attributes (12-16), plus `hipblasLtSdmaQueue_t`, `hipblasLtA2ACompletionMode_t`, and the shared `hipblasLtFusedEpilogueCreate` / `Add` / `SetAttribute` / `Destroy` builder. The builder documents family-neutral rules; per-family specifics stay on the enums. - `hipblasLtSetDeviceComm` registers a per-handle view of a device communicator, allocating the fine-grained flag regions and resolving peer addresses through a caller-supplied allgather, with IPC mapping for out-of-process peers. The flag-block layout is twinned with `FUSED_A2A_FLAG_BLOCK_BYTES` in the kernel's `FusedA2AKernArg.hpp`. - The descriptor attaches through `HIPBLASLT_MATMUL_DESC_FUSED_EPILOGUE` (106, after develop's `UNIFORM_SUMMATION_ORDER_EXT`) and is carried by `_rocblaslt_matmul_desc` as a non-owning pointer. - Validation is split three ways: stage composition at `Add`, required attributes at attach, and communicator, shape, and layout constraints before the heuristic and again at launch. No kernel is wired, so a well-formed request reports `HIPBLAS_STATUS_NOT_SUPPORTED`. Build `hipblaslt-test` and run `--gtest_filter='FusedA2A*'` on gfx950, covering descriptor lifecycle, builder composition rules, attribute range and width checks, communicator registration, attach-time completeness, and each dispatch-time rejection. 39 tests across 6 suites pass on gfx950. Nothing dispatches a GEMM: a valid request reports `NOT_SUPPORTED`, confirming the validation chain runs to completion. - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: Cursor <cursoragent@cursor.com>

Review of #11342 found the two peer branches of hipblasLtSetDeviceComm asymmetric: a peer in another process is mapped with hipIpcMemLazyEnablePeerAccess, while a peer in this process had only its raw device pointer recorded. A pointer from this process is not addressable on another device until peer access is enabled, so the copy engine would fault on a region it was handed. The same review found the communicator state on the handle holding two of its values twice. - DeviceCommExchange carries the contributing rank's device ordinal, which is what hipDeviceEnablePeerAccess names. A same-process peer on another device gets peer access enabled, already-enabled counting as success since the W handles enable each other pairwise. A peer on this device needs nothing, and HIP rejects a request to reach a device from itself. - device_comm_flags is dropped: the handle's own allocation is entry device_comm_rank of device_comm_peer_flags. release_device_comm tells the three kinds of entry apart by index and by mapped, so it frees its own region, closes its IPC mappings, and leaves a same-process peer's pointer to the handle that allocated it. - device_comm_registered is dropped for device_comm_world != 0. Registration rejects a world below one and publishes world last, so nonzero already means what the bool did. FusedA2ACommPeers.TwoRanksInOneProcessRegister: two ranks on two handles rendezvousing through a real allgather, which is the only way more than one rank can register in a single process. Where a second device exists the ranks take one each, so the peer branch has to make the region reachable rather than merely record its address; destroying both handles then exercises the ownership split. 40 tests across 7 suites pass on gfx950, with aux_handle_test, aux_attr_test, and aux_ext_test unaffected. The test machine has one device, so the cross-device hipDeviceEnablePeerAccess call itself is reached only where a second is present. - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: Ding, Yi <yi.ding@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com>

- Guard the getpid() include so Windows finds <process.h>, which was breaking the Windows CI build. - Stop re-deriving the flag region's geometry on the host. Its size now comes from the kernel's FUSED_A2A_FLAG_BLOCK_BYTES via rocblaslt_device_comm_flag_block_bytes(), declared in rocblaslt-auxiliary.h so hipblaslt.cpp gains no Tensile dependency. Only HIPBLASLT_DEVICE_COMM_MAX_WORLD still needs a static_assert, being public API that sizes the handle's per-rank arrays. - Split the two-rank registration test. The same-device case stays in pre_checkin; the cross-device case becomes a multi_gpu suite that skips below two GPUs instead of placing both ranks on device 0, where it passed without running the branch it was meant to cover. - Have that test assert registration's effect: each rank disables peer access first, so a later request reporting already-enabled can only be registration's work. - Add a multi_gpu ctest category so a two-device runner can select the suite. - Scope the allgather doc's "memcpy" to world == 1 and note that returning before every rank contributes fails registration rather than hanging. Co-authored-by: Cursor <cursoragent@cursor.com>

- register EnableGemmA2AFusion in _GLOBAL_PARAMETER_IGNORE_KEYS, without which the strict global-parameter gate aborted every device-library codegen run with ConfigTypeError, failing every Linux, Windows, ASAN and internal build on all architectures - admit a logic entry whose problem type is absent rather than subscripting it, which raised TypeError in three generateLogicDataAndSolutions characterization tests Co-authored-by: Cursor <cursoragent@cursor.com>

…A gate A MagicMock problem type is truthy under both subscript and .get, so the fusion filter read the synthetic logic files in test_gfx1250_asic_revision.py as fused and dropped them, leaving masterLibraries empty and failing three tests. Exclude only on an explicit True, so a problem type that cannot answer is admitted rather than silently emptying the library. Co-authored-by: Cursor <cursoragent@cursor.com>

The gate's only unit coverage called the predicate directly, so the branch that actually drops a logic file and reports the count went untested. That is the path a previously over-matching filter took when it emptied the library behind a successful build, which is what makes the reported count worth asserting alongside the drop. Drives the real merge loop over one fused logic file with the gate off. Co-authored-by: Cursor <cursoragent@cursor.com>

Review feedback on the fused GEMM + all-to-all suite. The fake allgathers silenced unused-parameter warnings with (void) no-op statements; leaving the names in comments does the same thing without the statements, and keeps the send and receive buffers distinguishable in failingAllgather, which takes two void* in a row. This matches the /*overwrite*/ style already used in testing_aux_check_numerics_gemm.hpp and tensile_host.cpp. Also tighten the file header and the helper comments. What stays is what the code cannot show: that suite names must keep their category token or the ctest presets silently stop selecting them, and why the multi_gpu test has to assert on peer access as a side effect. Co-authored-by: Cursor <cursoragent@cursor.com>

hipblasLtSetDeviceComm establishes a device communicator itself: rank exchange, peer access, IPC mapping of a peer's flag region, and ownership of the flag allocations. RCCL already owns communicator topology and peer transport, so leaving this here risks a second communication ABI. Point at the tracking issue for moving the setup to RCCL or a shared transport layer. Packet emission stays in TensileLite regardless, since it depends on GEMM tile completion and matrix layout. Co-authored-by: Cursor <cursoragent@cursor.com>

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

#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

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

#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

assistant-librarian · 2026-09-17T19:13:25Z

TheRock Submodule Bump Activity

Newest first

2026-09-17 — Included in TheRock via ROCm/TheRock#8289.

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

#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

yzhang93 requested review from a team as code owners August 27, 2026 00:29

github-actions Bot added the project: hipblaslt label Aug 27, 2026

therock-pr-bot Bot added the Not ready to Review label Aug 27, 2026

assistant-librarian Bot added the organization: ROCm label Aug 27, 2026

DDEle reviewed Aug 28, 2026

View reviewed changes

Comment thread projects/hipblaslt/library/src/amd_detail/hipblaslt.cpp Resolved

Comment thread projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/handle.h Resolved

DDEle mentioned this pull request Aug 28, 2026

feat(hipblaslt): wire the fused GEMM + all-to-all epilogue into the launch path #11413

Merged

1 task done

therock-pr-bot Bot removed the Not ready to Review label Aug 28, 2026

yzhang93 requested a review from DDEle August 28, 2026 18:42

DDEle added a commit that referenced this pull request Aug 31, 2026

Merge hipBLASLt PR #11342 (1901f36) into PR #10925 (f92f577) as the f…

7163d3e

…used-A2A API base

yzhang93 force-pushed the users/vivizhan/fused-epilogue-a2a branch from 1901f36 to f332481 Compare September 1, 2026 17:36

yzhang93 requested a review from Alex-Vasile September 1, 2026 18:10

DDEle added a commit that referenced this pull request Sep 2, 2026

Merge develop (be47443) into hipBLASLt PR #11342 (f332481) as the fus…

2c8910a

…ed-A2A API base

yzhang93 requested a review from davidd-amd September 10, 2026 20:58

davidd-amd approved these changes Sep 14, 2026

View reviewed changes

yzhang93 force-pushed the users/vivizhan/fused-epilogue-a2a branch from ec9fb61 to ebc274d Compare September 14, 2026 18:36

yzhang93 force-pushed the users/vivizhan/fused-epilogue-a2a branch from ebc274d to a8c370e Compare September 14, 2026 18:38

yzhang93 and others added 10 commits September 15, 2026 22:51

build(hipblaslt): gate GEMM A2A fusion

4e83349

fix(hipblaslt): name GEMM A2A build flag

88e7c22

yzhang93 force-pushed the users/vivizhan/fused-epilogue-a2a branch from a8c370e to 7e28eb9 Compare September 16, 2026 04:44

yzhang93 merged commit 3cb75b3 into develop Sep 16, 2026
328 of 337 checks passed

yzhang93 deleted the users/vivizhan/fused-epilogue-a2a branch September 16, 2026 16:08

Navigation Menu

Repository navigation

feat(hipblaslt): add host API for the fused GEMM + all-to-all epilogue - #11342

feat(hipblaslt): add host API for the fused GEMM + all-to-all epilogue#11342
yzhang93 merged 10 commits into
developfrom
users/vivizhan/fused-epilogue-a2a

yzhang93 commented Aug 27, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Aug 27, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Aug 27, 2026 •

edited

Loading

Uh oh!

codecov Bot commented Aug 27, 2026 •

edited

Loading

Uh oh!

DDEle left a comment •

edited

Loading

Uh oh!

Uh oh!

Uh oh!

davidd-amd commented Sep 11, 2026

Uh oh!

davidd-amd commented Sep 14, 2026

Uh oh!

Uh oh!

assistant-librarian Bot commented Sep 17, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

Conversation

yzhang93 commented Aug 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Uh oh!

therock-pr-bot Bot commented Aug 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

therock-pr-bot Bot commented Aug 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

codecov Bot commented Aug 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

DDEle left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

davidd-amd commented Sep 11, 2026

Uh oh!

davidd-amd commented Sep 14, 2026

Uh oh!

Uh oh!

assistant-librarian Bot commented Sep 17, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

yzhang93 commented Aug 27, 2026 •

edited

Loading

therock-pr-bot Bot commented Aug 27, 2026 •

edited

Loading

therock-pr-bot Bot commented Aug 27, 2026 •

edited

Loading

codecov Bot commented Aug 27, 2026 •

edited

Loading

DDEle left a comment •

edited

Loading