Repository navigation
feat(clr): add new rdma api - #10176
Conversation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
There was a problem hiding this comment.
Pull request overview
Adds a new HIP runtime API to flush inbound GPUDirect RDMA writes and exposes associated capability/ordering attributes, with CLR/hipamd implementation plus contract/unit tests to pin expected behavior.
Changes:
- Introduces
hipDeviceFlushGPUDirectRDMAWritesand new related enums/device attributes in the public runtime header. - Wires the new entry point through hipamd dispatch/trace/profiler plumbing and exports the symbol.
- Adds unit + contract coverage and updates test plans/coverage docs to include the new API.
Reviewed changes
Copilot reviewed 18 out of 18 changed files in this pull request and generated 2 comments.
Show a summary per file
| File | Description |
|---|---|
| projects/hip/include/hip/hip_runtime_api.h | Declares new API, enums, and device attributes for GPUDirect RDMA flush support. |
| projects/hip-tests/CONTRACT_COVERAGE.md | Updates contract coverage snapshot and lists the new API as covered. |
| projects/hip-tests/catch/unit/device/hipDeviceFlushGPUDirectRDMAWrites.cc | Adds unit tests for capability reporting, argument validation, and behavioral invariants. |
| projects/hip-tests/catch/unit/device/CMakeLists.txt | Registers the new unit test source file. |
| projects/hip-tests/catch/TEST_PLAN.md | Adds the new contract cases to the device domain section. |
| projects/hip-tests/catch/contract/README.md | Updates contract domain description to include the new flush contracts. |
| projects/hip-tests/catch/contract/device/test_hip_device_contract.cc | Adds portable contract tests (accepted-or-unsupported + invalid enumerators). |
| projects/hip-tests/catch/config/configs/unit/device.yaml | Adds configuration entries for the new unit test cases. |
| projects/hip-tests/catch/config/configs/contract.yaml | Adds configuration entries for the new contract cases. |
| projects/clr/hipamd/src/hip_table_interface.cpp | Adds table-interface wrapper for the new API. |
| projects/clr/hipamd/src/hip_internal.hpp | Adds helpers for reporting RDMA flush options and write ordering. |
| projects/clr/hipamd/src/hip_hcc.map.in | Exports hipDeviceFlushGPUDirectRDMAWrites in the version script. |
| projects/clr/hipamd/src/hip_device.cpp | Populates new device property fields for RDMA support/options/ordering. |
| projects/clr/hipamd/src/hip_device_runtime.cpp | Implements the new runtime API + exposes new device attributes. |
| projects/clr/hipamd/src/hip_api_trace.cpp | Extends the dispatch table update and ABI enforcement for the new function pointer. |
| projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h | Adds profiling API id/name/args formatting for the new API. |
| projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp | Bumps runtime table step version and adds new dispatch function pointer typedef/slot. |
| projects/clr/CHANGELOG.md | Documents the new GPUDirect RDMA flush API and attributes. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
| deviceProps.gpuDirectRDMASupported = (info.dmabufSupported_ && info.largeBar_) ? 1 : 0; | ||
| // Set option if we have a HDP register | ||
| deviceProps.gpuDirectRDMAFlushWritesOptions = | ||
| hip::ihipRdmaFlushWritesOptions(info.hdpMemFlushCntl); | ||
| deviceProps.gpuDirectRDMAWritesOrdering = hip::ihipRdmaWritesOrdering(); |
| int* flag = nullptr; | ||
| HIP_CHECK(hipHostMalloc(&flag, sizeof(int), hipHostMallocMapped)); | ||
| *flag = 0; |
621fbcb to
daaa33a
Compare
259ffaa to
6099883
Compare
a96fca1 to
9fb30d0
Compare
|
9fb30d0 to
672f501
Compare
672f501 to
71894ae
Compare
71894ae to
8d7e80c
Compare
| deviceProps.gpuDirectRDMASupported = (info.dmabufSupported_ && info.largeBar_) ? 1 : 0; | ||
| // Set option if we have a HDP register | ||
| deviceProps.gpuDirectRDMAFlushWritesOptions = | ||
| hip::ihipRdmaFlushWritesOptions(info.hdpMemFlushCntl); | ||
| deviceProps.gpuDirectRDMAWritesOrdering = hip::ihipRdmaWritesOrdering(); |
There was a problem hiding this comment.
🤖 AI-assisted pre-review (for @mangupta)
[IMPORTANT] gpuDirectRDMASupported can report 1 while the device advertises no flush path and no native ordering
The three RDMA capability fields are now derived from unrelated sources: gpuDirectRDMASupported from (info.dmabufSupported_ && info.largeBar_), gpuDirectRDMAFlushWritesOptions from whether info.hdpMemFlushCntl is non-null, and gpuDirectRDMAWritesOrdering from a helper that always returns hipGPUDirectRDMAWritesOrderingNone (0). A device with dmabuf + large BAR but no HDP flush register therefore advertises supported=1, flushOptions=0, ordering=None — an RDMA capability set that gives a caller no way to make an inbound remote write visible.
Suggested fix: Derive gpuDirectRDMASupported from the same evidence the flush path needs — e.g. (info.dmabufSupported_ && info.largeBar_ && info.hdpMemFlushCntl != nullptr) ? 1 : 0 — or, if a device genuinely can support GPUDirect RDMA without a host flush path, relax the new CapabilityConsistency test and document in the hipDeviceAttributeGPUDirectRDMASupported doxygen that supported=1 does not imply a usable flush/ordering path.
| #if CUDA_VERSION >= CUDA_11000 | ||
| case hipDeviceAttributeGPUDirectRDMASupported: | ||
| cdattr = cudaDevAttrGPUDirectRDMASupported; | ||
| break; | ||
| case hipDeviceAttributeGPUDirectRDMAFlushWritesOptions: | ||
| cdattr = cudaDevAttrGPUDirectRDMAFlushWritesOptions; | ||
| break; | ||
| case hipDeviceAttributeGPUDirectRDMAWritesOrdering: | ||
| cdattr = cudaDevAttrGPUDirectRDMAWritesOrdering; | ||
| break; | ||
| #endif // CUDA_VERSION >= CUDA_11000 |
There was a problem hiding this comment.
🤖 AI-assisted pre-review (for @mangupta)
[IMPORTANT] New hipother GPUDirect RDMA paths are guarded by CUDA_11000, but the CUDA symbols they use appear in CUDA 11.3
cudaDevAttrGPUDirectRDMASupported, cudaDevAttrGPUDirectRDMAFlushWritesOptions, cudaDevAttrGPUDirectRDMAWritesOrdering, cudaDeviceFlushGPUDirectRDMAWrites, cudaFlushGPUDirectRDMAWritesScope and cudaFlushGPUDirectRDMAWritesTargetCurrentDevice are part of the GPUDirect RDMA writes-ordering API introduced in CUDA 11.3, but the new code is compiled whenever CUDA_VERSION >= CUDA_11000. On a CUDA 11.0–11.2 toolkit the header fails to compile with undeclared-identifier errors, breaking every hipother/NVIDIA build on those toolkits.
Suggested fix: Replace CUDA_11000 with the existing CUDA_11030 for both the three attribute cases and the hipDeviceFlushGPUDirectRDMAWrites implementation (confirm the introducing toolkit version first), and drop the newly added CUDA_11000 define if nothing else uses it. Below that version the shim should keep returning hipErrorInvalidValue for a malformed enumerator and hipErrorNotSupported otherwise, as the current #else branch already does.
| SpinUntilFlag<<<1, 1, 0, stream>>>(deviceFlag); | ||
| HIP_CHECK(hipGetLastError()); | ||
|
|
||
| HIP_CHECK(hipDeviceFlushGPUDirectRDMAWrites(hipFlushGPUDirectRDMAWritesTargetCurrentDevice, | ||
| hipFlushGPUDirectRDMAWritesToOwner)); | ||
|
|
||
| // The kernel cannot have retired: nothing has released it yet. | ||
| HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorNotReady); |
There was a problem hiding this comment.
🤖 AI-assisted pre-review (for @mangupta)
[IMPORTANT] DoesNotSynchronize leaves an unbounded spin kernel resident and leaks the stream/pinned flag when an assertion fails
SpinUntilFlag loops forever until the host writes *flag = 1, which only happens after HIP_CHECK(hipDeviceFlushGPUDirectRDMAWrites(...)) and HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorNotReady). If either assertion fails — exactly the regression this test exists to catch — Catch2 aborts the test case, the flag is never set, and the kernel spins on the GPU indefinitely. The stream and the pinned mapped allocation are also never destroyed.
Suggested fix: Release the kernel before asserting: capture the hipStreamQuery result into a local, then set *flag = 1, synchronize, destroy the stream and free the pinned allocation, and only then REQUIRE the captured status. Alternatively wrap the flag/stream/allocation in an RAII guard whose destructor sets the flag and tears everything down, so a failed assertion still releases the kernel.
8d7e80c to
e635b72
Compare
Add the HIP counterpart of cudaDeviceFlushGPUDirectRDMAWrites, a host-ordered visibility barrier on inbound GPUDirect RDMA writes issued by a third-party device such as an RDMA-capable NIC. On AMD the flush is a write to the agent's HDP_MEM_COHERENCY_FLUSH_CNTL register, which ROCr already exposes via HSA_AMD_AGENT_INFO_HDP_FLUSH. Capability is reported by three new device attributes, also mirrored in hipDeviceProp_t, which until now were hardcoded to zero: hipDeviceAttributeGPUDirectRDMASupported hipDeviceAttributeGPUDirectRDMAFlushWritesOptions hipDeviceAttributeGPUDirectRDMAWritesOrdering Argument validation precedes the capability check, so a device with no flush path still reports hipErrorInvalidValue for a malformed enumerator rather than hipErrorNotSupported. Adding a runtime API entry point bumps HIP_RUNTIME_API_TABLE_STEP_VERSION to 35 (ABI slot 545). rocprofiler-sdk mirrors that table and fails a static_assert on any unrecognised step version, so its HIP API id, args union, tracing definition and ABI enforcement are updated in the same change. hipother gains the NVIDIA passthrough to cudaDeviceFlushGPUDirectRDMAWrites and the three attribute mappings. These pieces cannot be split: any commit that bumps the step version without the rocprofiler-sdk mirror, or declares the API without the NVIDIA definition, does not build.
e635b72 to
345a132
Compare
|
Errors are unrelated infra flukes, overriding |
TheRock Submodule Bump ActivityNewest first
|
The checked-in generated files predate the hipDeviceFlushGPUDirectRDMAWrites entry that #10176 added to hip_api_trace.hpp, so a CLR build with CLR_HRR_CODEGEN_CHECK=ON reports them stale. Re-run gen_hrr_api_args.py with --check-hrr-coverage. The new API sorts into slot 545, so the ids of the nine compiler-table entries after it each move up by one and HRR_API_COUNT becomes 555. Signed-off-by: Daniel Olsson <daniel.olsson@amd.com>
…e newer APIs at replay The nightly SDK that HRR CI installs predates hipDeviceFlushGPUDirectRDMAWrites (#10176), so the regenerated playback source no longer compiled against its headers, and would not have linked against its libamdhip64 either. - HRR_HIP_INCLUDE_DIRS puts the in-tree public HIP headers ahead of the SDK's for host C++ (not HIP device code). The unit jobs (Linux and Windows) and the multi-arch build pass it, so an API added to the tree compiles before any nightly declares it. - Headers do not help the link: an SDK library without the symbol fails hrr-playback with an undefined reference. APIs listed in _PLAYBACK_RUNTIME_RESOLVED_APIS are called through hrr_runtime_symbol (dlsym / module enumeration) and skipped with a warning when the runtime lacks them. hipDeviceFlushGPUDirectRDMAWrites is the first entry. - Regenerate hip_playback_generated.cpp; the other generated files are unchanged. Signed-off-by: Daniel Olsson <daniel.olsson@amd.com>
Motivation
Add the runtime API
hipDeviceFlushGPUDirectRDMAWritesand its used enum types.Technical Details
Implement
hipDeviceFlushGPUDirectRDMAWritesAPI. Which allows users to wait for writes to be visible after RDMA.Issue Tracking
JIRA ID: AIRUNTIME-2599
Test Plan
Add new tests for the new API
Test Result
Test passes locally.
Submission Checklist
Co-authored-by: Claude noreply@anthropic.com