Repository navigation
clr: allow same-device device-to-device batch swap and indirect copies - #12250
Merged
Merged
Conversation
ihipMemcpyBatch() rejected any extended-op copy whose endpoints had the same memory type, which ruled out same-device device-to-device swap and indirect transfers. The restriction was a conservative guard rather than a hardware limit: the SDMA batch path already routes these correctly, and peer device-to-device is handled by its own case and was never rejected. Narrow the guard to what actually cannot be served: host-to-host, and pageable or unregistered endpoints. Changes: - hip_memory.cpp: reject only host-to-host for extended ops - rocsettings.cpp: gate sdma_indirect_supported_ on gfxipMinor >= 5 rather than == 5, so the capability is not pinned to a single minor version. This selects the same set of architectures today. - memcpyBatchAsync_common.hh: drop the same-device device-to-device rejection from extOpPairingSupported() so the test expectations match, and bring its indirect gate in line with the runtime Without the test change, Unit_hipMemcpyBatchAsync_swap_cp and Unit_hipMemcpyBatchAsync_Broadcast_WithSwap expect hipErrorNotSupported for pairings the runtime now serves. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude <noreply@anthropic.com>
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
Contributor
There was a problem hiding this comment.
Copilot review overview
🟢 Approval recommended
No blocking issues were identified.
Review effort: Lite
Findings: None
What changed in this PR
This PR expands HIP batch extended-operation support for same-device device-to-device swaps and indirect copies while preserving unsupported host/pageable restrictions.
Changes:
- Narrows runtime rejection criteria.
- Enables indirect SDMA copies on gfx12.5+.
- Aligns unit-test expectations with runtime behavior.
| File | Summary |
|---|---|
projects/hip-tests/catch/unit/memory/memcpyBatchAsync_common.hh |
Updates pairing and architecture expectations. |
projects/clr/rocclr/device/rocm/rocsettings.cpp |
Enables indirect SDMA copies on gfx12.5+. |
projects/clr/hipamd/src/hip_memory.cpp |
Allows supported device and pinned-host extended-operation pairings. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
saleelk
approved these changes
Sep 25, 2026
Contributor
TheRock Submodule Bump ActivityNewest first
|
1 task done
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
ihipMemcpyBatch() rejected any extended-op copy whose endpoints had the same memory type, which ruled out same-device device-to-device swap and indirect transfers. The restriction was a conservative guard rather than a hardware limit: the SDMA batch path already routes these correctly, and peer device-to-device is handled by its own case and was never rejected.
Narrow the guard to what actually cannot be served: host-to-host, and pageable or unregistered endpoints.
Technical Details
Changes:
Without the test change, Unit_hipMemcpyBatchAsync_swap_cp and Unit_hipMemcpyBatchAsync_Broadcast_WithSwap expect hipErrorNotSupported for pairings the runtime now serves.
Issue Tracking
JIRA ID : ROCM-31787
Test Plan
Run existing batch tests covering same-device D2D swap, broadcast-with-swap, and indirect src/dst.
Test Result
Before the change these cases skipped the
same-device pairings; after it they execute.
Unit_hipMemcpyBatchAsync_swap_cpUnit_hipMemcpyBatchAsync_Broadcast_WithSwapUnit_hipMemcpyBatchAsync_IndirectSrcUnit_hipMemcpyBatchAsync_IndirectDstUnit_hipMemcpyBatchAsync_Swap(parameterized)Submission Checklist