Skip to content

feat(hip): add hipExtMemcpyBatch graph node type - #13512

Open
saleelk wants to merge 22 commits into
developfrom
users/saleelk/graphMemcpyFanOut_Batch_asymswap
Open

saleelk wants to merge 22 commits into
developfrom
users/saleelk/graphMemcpyFanOut_Batch_asymswap

Conversation

@saleelk

@saleelk saleelk commented Oct 14, 2026

Copy link
Copy Markdown
Contributor

hipExtMemcpyBatchAsync has no graph support. Add a hipGraphNodeTypeExtMemcpyBatch
node whose parameters mirror hipExtMemcpyBatchAsync: sizesDst for asymmetric
swap, the reserved waits/signals arrays, and hipExtMemcpyAttributes.

The node is reachable three ways: explicitly via hipGraphAddExtMemcpyBatchNode,
or by capturing hipExtMemcpyBatchAsync or hipMemcpyBatchAsync, both of which go
through STREAM_CAPTURE. hipExtMemcpyBatchAsync previously rejected capture with
hipErrorStreamCaptureUnsupported.

Argument validation is shared with the stream path (ihipValidateMemcpyBatch), so
node add/set and the stream APIs reject the same inputs. Classification is
shared too, so asymmetric swaps decompose identically on both paths.

Parameter updates work through hipGraphExtMemcpyBatchNodeSetParams,
hipGraphExecExtMemcpyBatchNodeSetParams and hipGraphExecUpdate.

Tests: catch/unit/graph/hipGraphAddExtMemcpyBatchNode.cc covers explicit
construction and capture of both batch APIs, parameter get/set, exec update, a
batch sandwiched between two kernel nodes, a peer-to-peer broadcast and an
asymmetric swap. Unit_hipExtMemcpyBatchAsync_StreamCapture_Negative becomes
_Positive.

Stacked on #9098 (hipExtMemcpyBatchAsync); targets its branch so only this commit shows. Retarget to develop once #9098 merges.

iassiour and others added 22 commits October 14, 2026 21:20
Remove the size_t* failIdx output parameter from hipExtMemcpyBatchAsync,
mirroring CUDA's C base batch-memcpy signature. The legacy
hipMemcpyBatchAsync and the 3D batch API keep failIdx; the shared
ihipMemcpyBatch still computes it for the legacy path, and the Ext
wrapper now passes nullptr.

Changes:
- Public header + doc: drop failIdx parameter
- hip_memory.cpp: Ext wrapper passes nullptr to ihipMemcpyBatch
- Dispatch-table typedef, table thunk, forward decl updated
- hip_prof_str.h: drop failIdx from arg struct, INIT macro, roctracer
  deep-copy, and ostream printer
- rocprofiler-sdk: drop failIdx from api_args.h and hip.def.cpp
- Tests: strip failIdx from Ext call sites; delete FailIdx_SwapRejected
  (its sole purpose was the Ext failIdx) + its yaml entry; keep the
  hipErrorInvalidValue checks in UnknownOpsBits/SizesB negatives

The dispatch-table slot (offset 542) is unchanged (still one function
pointer), so HIP_ENFORCE_ABI / ROCP_SDK_ENFORCE_ABI and the
STEP_VERSION bump remain valid.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
The previous commit removed the failIdx argument too broadly in the test
file, stripping it from legacy hipMemcpyBatchAsync calls as well (the
Negative / NegativeTsts / swap_cp / HD_Swap cases), which take failIdx and
broke the build ("too few arguments", caught by the NVIDIA hip-tests build).

Restrict the removal to hipExtMemcpyBatchAsync call sites only; legacy calls
keep failIdx. All 17 Ext calls now pass 12 args; legacy calls are unchanged.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
The file-header doxygen snippet documents the legacy hipMemcpyBatchAsync
signature, which keeps failIdx. The earlier failIdx removal wrongly stripped
it from this comment too. Restore it to match hip_runtime_api.h.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
The waits/signals descriptors are reserved (must be NULL) today. Fix their
layout to 64 bytes now, while nothing consumes them, so future fields can be
added without an ABI break: each carries an explicit uint32 reserved slot
(fills the alignment gap after compareOp/signalOp) plus a uint64 reserved
tail. Callers must zero-initialize; reserved fields must be 0.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Indirect copy was gated to heterogeneous src/dst memory types (H<->D), which
rejected device-to-device indirect - the primary use case. The SDMA indirect
packet has no such restriction (all system-scope, no placement logic), so the
gate was an over-conservative CLR-layer check, not a HW limit.

Relax the indirect validation to require only SDMA-accessible operands (both
have an amd::Memory); device<->device (same device or peer) and host<->device
are all accepted, pageable host is still rejected. Add a D2D indirect-source
test (skips where indirect is unsupported).

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
The contract-coverage gate flagged hipExtMemcpyBatchAsync as declared but
uncovered. Add a device-only contract test (two-op D2D round-trip with all
extended parameters null, i.e. passthrough behavior) alongside the legacy
hipMemcpyBatchAsync contract tests. Skips where the extended batch path is
unsupported. Verified with check_contract_coverage.py --check (RESULT: OK).

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
- Drop ops[]; op selection via attrs[i].flags (new hipExtMemcpyAttributes)
- Rename sizesA/B -> sizes/sizesDst; legacy wrapper converts attrs by value
- Add hipMemcpyFlagExtPreferCU (shader copy path); reject with PreferCE/swap/indirect
- Keep develop's swap/indirect scope: H<->D and peer P2P only
- Attributes carry a reserved tail (reserved[8]) for forward-compat

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
- Reject non-zero sizesDst[i] on non-swap entries (hipErrorInvalidValue) + test
- Validate attrs/attrsIdxs/numAttrs in the legacy wrapper before the conversion loop
- Add HIP_PUBLIC_API to the hipExtMemcpyBatchAsync table thunk
- static_assert sizeof(hipExtMemcpyWait) == 64 and sizeof(hipExtMemcpySignal) == 64
- Document the disjoint-range invariant of the asymmetric-swap decomposition and add
  a mixed-batch test (asymmetric swap + linear copy in one submission)
- Fix dispatch-table versioning after rebase (versioning 546, step 35)

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
"CU" reads as a hardware Compute Unit, which does not map cleanly onto
every architecture (e.g. WGP-based designs). Rename the new flag to
hipMemcpyFlagExtPreferComputeEngine so it pairs with the existing
hipMemcpyFlagExtPreferCE ("copy engine over compute engine") and does
not name a hardware unit. The value (0x1000) and behavior are unchanged.

Changes:
- Rename hipMemcpyFlagExtPreferCU -> hipMemcpyFlagExtPreferComputeEngine
  in the public header, runtime validation/mapping, and tests.
- Rename Unit_hipExtMemcpyBatchAsync_PreferCU_D2D to
  Unit_hipExtMemcpyBatchAsync_PreferComputeEngine_D2D.
- Keep flag docs engine-neutral: restore the PreferCE doc text to match
  develop and drop the implementation detail from the new flag's doc.

Tested on gfx950: hipMemcpyBatchAsync suite 28 passed / 8 skipped,
contract test 4/4.

Co-Authored-By: Claude <noreply@anthropic.com>
Add two grouping hints to hipMemcpyFlags for copies that share a
source: hipMemcpyFlagExtPreferLinear (issue each copy individually)
and hipMemcpyFlagExtPreferBroadcast (combine copies with the same src
and size into one broadcast). Like the other Prefer* flags they are
hints; the runtime accepts them now and does not yet act on them, so
applications can set them ahead of the grouping logic honoring them.

Changes:
- driver_types.h: add hipMemcpyFlagExtPreferLinear = 0x2000 and
  hipMemcpyFlagExtPreferBroadcast = 0x4000.
- hip_memory.cpp: accept both bits; reject PreferLinear together with
  PreferBroadcast, and PreferBroadcast with swap or indirect
  (hipErrorInvalidValue).
- hip_runtime_api.h: list the new hints in the API description.
- Tests: negative sections for the invalid combinations, and
  Unit_hipExtMemcpyBatchAsync_PreferLinearBroadcast_D2D (one source to
  three destinations, per hint) registered in memory.yaml.

Tested on gfx950: hipMemcpyBatchAsync suite 29 passed / 8 skipped,
contract test 4/4.

Co-Authored-By: Claude <noreply@anthropic.com>
Replace the hand-rolled loop over the parallel src/dst size vectors
with a direct vector comparison. Behavior is unchanged.

Tested on gfx950: hipMemcpyBatchAsync suite 29 passed / 8 skipped,
contract test 4/4.

Co-Authored-By: Claude <noreply@anthropic.com>
The scalar swap path no longer rejects asymmetric sizes itself; it
forwards both sizes and the SDMA blit rejects A != B unless native
asymmetric swap is supported. Update the comment to match.

Co-Authored-By: Claude <noreply@anthropic.com>
…ep 37

Co-Authored-By: Claude <noreply@anthropic.com>
The asymmetric LINEAR_SWAP rejection lived only in BlitSdma::SubmitBodies,
so entries routed to the coordinator engine on the fused WaitSignal path
went through SubmitFusedCoordinator unchecked. There the chunk count is
derived from max(size_a, size_b) while each swap packet's count comes from
the A side, so src_size != dst_size could advance past the smaller region
or underflow the count. It was also evaluated per engine group, after the
prologue and other groups could already have been submitted.

Changes:
- Check dst_size_list once in GpuAgent::DmaCopyFanOutOp, before any engine
  is submitted, covering both the fused coordinator and body paths
- Expose NativeAsymmetricSwapSupported() on BlitSdmaBase (still false)
- Drop the per-group check from SubmitBodies and update the related docs

HIP is unaffected (CLR always decomposes asymmetric swaps); the scalar
swap op of hsa_amd_memory_async_batch_copy reaches this path directly.

Tested on gfx950: hipMemcpyBatchAsync suite 29 passed / 8 skipped, batch
copy contract 4/4; a direct scalar swap with src_size != dst_size returns
HSA_STATUS_ERROR_INVALID_ARGUMENT with the completion signal untouched,
and a symmetric swap still swaps.

Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
The CLR decomposition treated sizes[i] as the dsts[i] length and
sizesDst[i] as the srcs[i] length, opposite to the documented contract,
the parameter name, and the native pass-through (src_size = sizes[i] on
srcs[i], dst_size = sizesDst[i] on dsts[i]).

Changes:
- hip_memory.cpp: the tail copy now goes srcs[i][sizesDst..sizes) ->
  dsts[i][sizesDst..sizes), matching the swap's src->dst direction
- hip_runtime_api.h: state which operand each size describes
- hipMemcpyBatchAsync.cc: swap the pointer roles in the asymmetric tests
  so the larger (A) side is srcs[i]

Tested on MI350X (gfx950): batch suite 37 cases, 29 passed, 8 skipped,
0 failed; H<->D asymmetric and mixed asymmetric+linear cases pass;
contract tests 4/4.

Co-Authored-By: Claude <noreply@anthropic.com>
hipExtMemcpyBatchAsync executed eagerly on a capturing stream, returned
success, and left an empty graph, so graph replay silently skipped the
copy. Graph capture is not supported for the Ext API yet: a capturing
stream now has its capture invalidated and the call returns
hipErrorStreamCaptureUnsupported.

In ROCr, the swap packet's addr_a is the dst pointer but was paired with
the src size (and addr_b with the dst size). This only matters for
asymmetric swaps, which are rejected up front today, but would corrupt
the operand lengths once a native asymmetric path is enabled.

Changes:
- hip_memory.cpp: reject/invalidate stream capture in
  hipExtMemcpyBatchAsync
- amd_blit_sdma.cpp: pair dst size with addr_a and src size with addr_b
  in the coordinator and SubmitBodies; src/dst wording in comments
- hipMemcpyBatchAsync.cc / memory.yaml: add
  Unit_hipExtMemcpyBatchAsync_StreamCapture_Negative

Tested on gfx950: batch suite 38 cases (30 passed, 8 skipped), contract
tests 4/4.

🤖 Generated with [Claude Code](https://claude.com/claude-code)

Co-Authored-By: Claude <noreply@anthropic.com>
Shorten the comments added for the extended batch-copy API so they state
only what the code does not already say.

Changes:
- hip_memory.cpp: drop comments that restate the validation checks;
  condense the asymmetric swap decomposition and sizesDst notes
- driver_types.h / hip_runtime_api.h: one-line flag and struct docs;
  document that the API is not supported during stream capture and
  returns hipErrorStreamCaptureUnsupported
- rocclr / ROCr: use src/dst side wording and shorten dst_size_list docs
- Contract test: remove a comment duplicating the @asserts tag

Comment-only apart from the doc additions. Batch unit tests (30 passed,
8 skipped) and mem_batch_copy contract tests (4/4) pass on gfx950.

Co-Authored-By: Claude <noreply@anthropic.com>
"Always" read as "wait forever" at the HIP level. The zero value means no
wait, so name it to match hipExtMemcpySignalNone. The value is unchanged
(0) and still maps to the HSA ALWAYS compare.

Changes:
- driver_types.h: hipExtMemcpyWaitAlways -> hipExtMemcpyWaitNone

Header-only rename; no other users of the enumerator in this branch.

Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
hipExtMemcpyBatchAsync has no graph support. Add a hipGraphNodeTypeExtMemcpyBatch
node whose parameters mirror hipExtMemcpyBatchAsync: sizesDst for asymmetric
swap, the reserved waits/signals arrays, and hipExtMemcpyAttributes.

The node is reachable three ways: explicitly via hipGraphAddExtMemcpyBatchNode,
or by capturing hipExtMemcpyBatchAsync or hipMemcpyBatchAsync, both of which go
through STREAM_CAPTURE. hipExtMemcpyBatchAsync previously rejected capture with
hipErrorStreamCaptureUnsupported.

Argument validation is shared with the stream path (ihipValidateMemcpyBatch), so
node add/set and the stream APIs reject the same inputs. Classification is
shared too, so asymmetric swaps decompose identically on both paths.

Parameter updates work through hipGraphExtMemcpyBatchNodeSetParams,
hipGraphExecExtMemcpyBatchNodeSetParams and hipGraphExecUpdate.

Tests: catch/unit/graph/hipGraphAddExtMemcpyBatchNode.cc covers explicit
construction and capture of both batch APIs, parameter get/set, exec update, a
batch sandwiched between two kernel nodes, a peer-to-peer broadcast and an
asymmetric swap. Unit_hipExtMemcpyBatchAsync_StreamCapture_Negative becomes
_Positive.

Depends on #9098 (hipExtMemcpyBatchAsync).

@mangupta mangupta left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🤖 AI-assisted pre-review (for @mangupta)

[IMPORTANT] No contract test or uncovered_apis.txt rationale for the four new public graph APIs

The PR adds four public HIP APIs to hip_runtime_api.h: hipGraphAddExtMemcpyBatchNode, hipGraphExtMemcpyBatchNodeGetParams, hipGraphExtMemcpyBatchNodeSetParams and hipGraphExecExtMemcpyBatchNodeSetParams. It also adds the public enum value hipGraphNodeTypeExtMemcpyBatch. It does not add a projects/hip-tests/catch/contract/ test with an @asserts: API - invariant annotation, does not add an entry to catch/contract/uncovered_apis.txt, and does not update catch/TEST_PLAN.md.

Suggested fix: For each new API, add contract tests under catch/contract/ following AUTHORING.md, for example: null pGraphNode → hipErrorInvalidValue; wrong node type for GetParams → hipErrorInvalidValue; a count change through ExecSetParams → hipErrorInvalidValue. Alternatively, add a concrete rationale to uncovered_apis.txt. Then regenerate catch/TEST_PLAN.md.

[IMPORTANT] rocprofiler-sdk gets ABI asserts for the new dispatch entries but no API IDs, argument plumbing, or hip.def.cpp registration

The PR adds four HipDispatchTable entries (step version 39) and HIP_API_IDs 492–495 in hip_prof_str.h. On the rocprofiler-sdk side it only adds ROCP_SDK_ENFORCE_ABI lines in abi.cpp and a format.hpp case for the new node type. It does not add the API to runtime_api_id.h, api_args.h, or hip.def.cpp. Without those, rocprofiler-sdk cannot assign callback/buffer-tracing operations to the new APIs, so calls to them would not show up in HIP API traces.

Suggested fix: Following the pattern used for earlier step versions, add ROCPROFILER_HIP_RUNTIME_API_ID_hipGraphAddExtMemcpyBatchNode and the other three IDs, guarded by HIP_RUNTIME_API_TABLE_STEP_VERSION >= 39, to runtime_api_id.h. Add matching argument structs to api_args.h and HIP_API_INFO_DEFINITION entries to hip.def.cpp. If they are intentionally deferred, say so in the PR and add a follow-up.

[SUGGESTION] No CHANGELOG entry for the new graph APIs and capture behavior; hipMemcpyBatchAsync doc doesn't mention capture

The PR adds four public APIs, a new public hipGraphNodeType value, and stream-capture support for hipExtMemcpyBatchAsync, which previously returned hipErrorStreamCaptureUnsupported. It also adds capture support for hipMemcpyBatchAsync. projects/clr/CHANGELOG.md is not updated. Only the hipExtMemcpyBatchAsync doxygen was updated for capture; the hipMemcpyBatchAsync doxygen was not changed to say that capture now records a hipGraphNodeTypeExtMemcpyBatch node.

Suggested fix: Add a CLR CHANGELOG entry listing the new APIs, the node type, and the capture support. Update the hipMemcpyBatchAsync doxygen to describe its capture behavior and the restrictions (no host-to-host, single device).

Comment on lines +3449 to +3463
// outside that model - a CPU memcpy after draining the stream, and submissions to another
// device's null stream joined by a marker - have no equivalent here.
if (!classification.host_to_host_indices.empty()) {
return hipErrorNotSupported;
}
const size_t stream_device = static_cast<size_t>(stream.DeviceId());
for (size_t device_id = 0; device_id < classification.copy_ops_by_device.size(); ++device_id) {
if (device_id != stream_device &&
(!classification.copy_ops_by_device[device_id].empty() ||
!classification.write_ops_by_device[device_id].empty() ||
!classification.read_ops_by_device[device_id].empty())) {
return hipErrorNotSupported;
}
}

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🤖 AI-assisted pre-review (for @mangupta)

[IMPORTANT] Graph batch node accepts host-to-host and cross-device batches at add/capture/SetParams time and only fails later at instantiate/launch

The graph path cannot run host-to-host entries or entries that classify to a device other than the stream's device. ihipMemcpyBatchCreateCommands returns hipErrorNotSupported for both. That check only runs from GraphExtMemcpyBatchNode::CreateCommand, which executes when the graph is instantiated or launched. hipGraphAddExtMemcpyBatchNode, hipGraphExtMemcpyBatchNodeSetParams, hipGraphExecExtMemcpyBatchNodeSetParams and the two capture paths (capturehipMemcpyBatchAsync / capturehipExtMemcpyBatchAsync) call only ValidateExtMemcpyBatchNodeParams, which calls ihipValidateMemcpyBatch. That function checks argument shape and never classifies the copies. So these calls return hipSuccess for a batch the node can never execute.

Suggested fix: Run ClassifyMemcpyBatch at node add, capture, and SetParams time, using the target device: the capturing stream's device for capture, and the current device for explicit add, recorded on the node. Return hipErrorNotSupported, with failIdx where one is available, for host-to-host or cross-device entries. When capture rejects a batch, invalidate the capture the same way other unsupported captured operations do. Then align the doxygen with the error-reporting point, and add negative tests that add and capture a host-to-host batch and expect hipErrorNotSupported from the add/capture call.

Comment on lines +3523 to 3525
failIdx);
if (failIdx != nullptr) *failIdx = SIZE_MAX;
if (!hip::isValid(stream)) {

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🤖 AI-assisted pre-review (for @mangupta)

[SUGGESTION] Capture path of hipMemcpyBatchAsync returns before failIdx is set to SIZE_MAX

STREAM_CAPTURE is placed before if (failIdx != nullptr) *failIdx = SIZE_MAX;. On a capturing stream the API returns from inside STREAM_CAPTURE, so failIdx is never reset. For success, and for errors not tied to one entry (for example count == 0 or invalid attrs), the caller gets whatever stale value was in failIdx instead of SIZE_MAX. The non-capture path does reset it.

Suggested fix: Move the *failIdx = SIZE_MAX assignment above STREAM_CAPTURE, or set it at the start of capturehipMemcpyBatchAsync, so both paths report failIdx the same way.

@andraghetti andraghetti left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

HRR side only: the hipExtMemcpyBatchAsync capture and replay, with pointer translation for the arrays and the 16-entry limit documented, look right.

HRR_API_HIPREGISTERTEXTURE = 555,
HRR_API_HIPREGISTERVAR = 556,
HRR_API_HIPUNREGISTERFATBINARY = 557,
HRR_API_COUNT = 558

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Running gen_hrr_api_args.py at a039890 gives HRR_API_COUNT = 562, not 558: the four hipGraph*ExtMemcpyBatchNode* slots this PR adds to HipDispatchTable are missing from the checked-in HRR files. So a graph built with hipGraphAddExtMemcpyBatchNode is captured without that node and replays without its copies. Regenerating the three files fixes it.

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Regenerating adds the slots but does not replay the node. hipExtMemcpyBatchNodeParams holds the same pointer arrays as hipExtMemcpyBatchAsync, so hipGraphAddExtMemcpyBatchNode and both SetParams calls need matching Deref entries in gen_hrr_api_args.py.

* reader decodes those events as the wrong API.
* v9: hipKernelSetAttributeForDevice and hipExtMemcpyBatchAsync were
* appended to HipDispatchTable, shifting the compiler API IDs up by two. */
#define HRR_VERSION ((uint16_t)9u)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

develop is already at v9, with hipKernelSetAttributeForDevice at 547 (#13472). The APIs added here move the compiler IDs again, so once this targets develop it needs v10; otherwise a v9 archive from develop decodes __hipPopCallConfiguration (548) as hipExtMemcpyBatchAsync.

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

#9098 landed on develop as v12 with HRR_API_COUNT 559, so after the rebase this needs v13.

@andraghetti
andraghetti requested a review from a team October 19, 2026 08:08

@cfreeamd cfreeamd left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🤖 AI-assisted pre-review (for @cfreeamd)

The ROCr changes thread an optional dst_size_list (per-entry dst-side swap sizes) through DmaCopyFanOutOp → SubmitBodies → BuildWaitSignalSwapCommand, plumbing asymmetric-swap sizing while keeping the feature gated behind the new NativeAsymmetricSwapSupported() (false for all current BlitSdma instantiations). Symmetric copy/swap throughput is unchanged — b_size(d) collapses to size_list[d] whenever the op is non-swap or dst_size_list == nullptr. Completion-signal ordering is untouched; only per-chunk byte counts and the addr↔size pairing changed.

BuildWaitSignalSwapCommand(command_addr, chunks[i],
dsts[i], const_cast<void*>(srcs[i]),
sizes_a[i], sizes_b[i],
sizes_b[i], sizes_a[i], // addr_a is dst: pair it with the dst size

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🤖 AI-assisted pre-review (for @cfreeamd)

[Question] The swap here flips sizes_a[i], sizes_b[i] → sizes_b[i], sizes_a[i] so addr_a (dst) is paired with the dst size — correct pairing, inert for symmetric entries where sizes_a==sizes_b. But is the fused coordinator actually reachable for asymmetric swaps in this PR, or only in the follow-up asym-swap branch? If asymmetric swaps can't reach it yet, consider adding an assert here too (sizes_a[i]==sizes_b[i]) to mirror the classic and wait/signal-swap builders and make all three paths fail consistently until native support lands.


if (is_swap) {
// Single-count packet; DmaCopyFanOutOp rejects asymmetric swaps.
assert(b_size(d) == size_list[d] && "asymmetric swap reached classic swap builder");

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🤖 AI-assisted pre-review (for @cfreeamd)

[Optional] Good defensive assertion. Note that it is debug-only: in a release build, if a future NativeAsymmetricSwapSupported()==true blit somehow routes an asymmetric entry through this classic path, BuildSwapCopyCommand(... size_list[d]) would silently use the src size and emit a wrong copy with no error. The real guarantee is that native-asymmetric blits must never use this builder. Consider a hard return HSA_STATUS_ERROR_INVALID_ARGUMENT instead of (or in addition to) the assert, or a comment pointing at the follow-up branch so reviewers know this assert is load-bearing until that work merges.

}
}

// Reject asymmetric swaps before this op submits; the caller decomposes them.

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🤖 AI-assisted pre-review (for @cfreeamd)

[Optional] engines[d].blit->NativeAsymmetricSwapSupported() is dereferenced for each entry before the main dispatch. Worth confirming that engines[d].blit is guaranteed non-null for all d < num_entries at this point (the downstream fan-out code clearly assumes valid engines, so this is likely fine — just worth a quick check in the context of the broader function).

@iassiour
iassiour force-pushed the users/iassiour/asym-swap branch 6 times, most recently from 238c9d5 to 9caa902 Compare October 26, 2026 11:38
@iassiour
iassiour force-pushed the users/iassiour/asym-swap branch from d31d0e0 to efe7ba0 Compare October 27, 2026 11:20
Base automatically changed from users/iassiour/asym-swap to develop October 27, 2026 21:13
@therock-pr-bot

Copy link
Copy Markdown

✅ All Policy Checks Passed

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —

🎉 All policy checks passed!

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

🙋 Wish to Override Policy?

@andraghetti andraghetti left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Checked on MI350 (gfx950) at a039890 with a standalone program, develop b09c3ab as control.

return extAttrs;
}
// Convert by value; the struct strides differ.
extAttrs.resize(numAttrs);

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

hipMemcpyBatchAsync lost the numAttrs > count check that ran before this conversion. count=1, numAttrs=SIZE_MAX now throws in resize and returns hipErrorUnknown, both immediately and under capture; develop returns hipErrorInvalidValue. Restore the check ahead of ihipToExtMemcpyAttributes in hipMemcpyBatchAsync and capturehipMemcpyBatchAsync.

}

hip::GraphNode* node = new hip::GraphExtMemcpyBatchNode(nodeParams);
status = ihipGraphAddNode(node, reinterpret_cast<hip::Graph*>(graph),

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This call leaves capture at its default true, so a batch node added to the graph from hipStreamGetCaptureInfo_v2 is not registered as a manual node. hipStreamEndCapture then returns hipErrorStreamCaptureUnjoined; the same flow with hipGraphAddEmptyNode returns hipSuccess. Pass false, as hipGraphAddEmptyNode does.

HIP_CHECK(hipStreamSynchronize(0));
buffers.VerifyCopied();

// The node owns a copy of the parameter arrays, so a relaunch must repeat the same work.

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

BatchBuffers keeps dsts_, srcs_ and sizes_ alive and unchanged, so a node that aliased the caller's arrays would also pass this relaunch. Overwrite the arrays after BuildBatchGraph so the test checks the copy.

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants