Repository navigation
UCT/IB/MLX5/GDAKI: Try VMM Fabric Allocation. - #11795
Conversation
5549a4c to
5d499f0
Compare
5d499f0 to
4a5f958
Compare
This reverts commit 4a5f958.
|
🤖 Starting review — findings will be posted here when done. |
|
🤖 Starting review — findings will be posted here when done. |
|
🤖 Starting review — findings will be posted here when done. |
|
Test coverage: This is primarily a refactor plus the GDAKI fabric-path enablement (finding 4). The extracted |
|
🤖 Starting review — findings will be posted here when done. |
|
GDAKI now takes the context-wide sync-memops write-lock (potential deadlock). Before this PR, GDAKI set sync-memops per pointer via cuPointerSetAttribute(SYNC_MEMOPS), which takes no context lock. After the change, uct_rc_gdaki_alloc calls uct_cuda_mem_alloc(), which internally calls uct_cuda_copy_sync_memops() -> uct_cuda_copy_set_ctx_sync_memops(), which may call cuCtxSetFlags(CU_CTX_SYNC_MEMOPS) and take the context write-lock. The comment at that call site documents that this deadlocks a progress thread against an application thread inside cudaLaunchKernel(), and the hazard is only pre-empted for the context current at cuda_copy MD open. GDAKI allocations happen at ep/channel-pool init on the GDAKI primary context (a different context) and can run concurrently with application kernels, so this path can now hit the documented deadlock the old per-pointer call avoided. Blocker-level question: can this deadlock against an app thread in cudaLaunchKernel()?
Return value of fabric release is now silently dropped. uct_cuda_copy_mem_free previously propagated the status from uct_cuda_copy_mem_release_fabric. Now uct_cuda_mem_free is void and uct_cuda_copy_mem_free always returns UCS_OK, so a real fabric-release failure (e.g. cuMemRelease/cuMemUnmap error) is no longer reported to the uct_md_mem_free caller. Minor behavior change; if intentional, fine, but worth confirming since VMM release errors used to be surfaced.
Doc typo: '@PARAM [in,out] granularity_p Allocation granularity,if fabric VMM is used' is missing a space after the comma — should be 'granularity, if fabric VMM is used'. |
|
🤖 Starting review — findings will be posted here when done. |
|
sync-memops context deadlock (blocker-level question): GDAKI's uct_rc_gdaki_alloc now routes through uct_cuda_mem_alloc → uct_cuda_copy_sync_memops → uct_cuda_copy_set_ctx_sync_memops, which can take the context write-lock. The code already contains a cuCtxGetFlags read-first mitigation that only calls cuCtxSetFlags when the flag is not yet set, which reduces (but does not fully eliminate, per its own comment) the documented hazard on non-MD-open contexts. Worth confirming this mitigation is considered sufficient for the GDAKI primary context. Dropped fabric-release status: uct_cuda_mem_free is now void and uct_cuda_copy_mem_free always returns UCS_OK, so VMM release errors are no longer surfaced to the MD caller.
Doc typo: missing space after comma in the granularity_p doc line. Test coverage gap: the newly enabled GDAKI-over-fabric-VMM path (VMM buffer + umem/dmabuf registration, now that enable_fabric is UCS_TRY in uct_rc_gdaki_alloc) has no obvious dedicated CI coverage. |
|
🤖 Starting review — findings will be posted here when done. |
|
🤖 Review complete — no issues found. |
|
🤖 Starting review — findings will be posted here when done. |
|
🤖 Starting review — findings will be posted here when done. |
|
Removed Test coverage gap: The newly enabled GDAKI-over-fabric-VMM path ( |
|
🤖 Starting review — findings will be posted here when done. |
|
🤖 Starting review — findings will be posted here when done. |
|
Residual coverage (already raised in earlier rounds): the GDAKI-over-fabric-VMM path has no dedicated CI coverage. Finding 1 makes that gap more relevant, since the peermem-only configuration is exactly the one no existing test or capability probe exercises. |
|
🤖 Starting review — findings will be posted here when done. |
|
Residual coverage: no dedicated CI coverage for the GDAKI-over-fabric-VMM path (VMM buffer + umem/dmabuf registration). The direct ep_alloc_mode memory footprint with fabric enabled is also untested. |
|
🤖 Starting review — findings will be posted here when done. |
|
Residual coverage (already raised in earlier rounds, not a new comment): the GDAKI-over-fabric-VMM path (VMM buffer + umem/dmabuf registration, |
|
🤖 Starting review — findings will be posted here when done. |
|
🤖 Starting review — findings will be posted here when done. |
|
🤖 Starting review — findings will be posted here when done. |
What?
GDAKI now tries CUDA VMM fabric allocation and falls back to
cuMemAllocif it is not available.Why?
GDAKI previously always allocated with
cuMemAlloc. VMM fabric allocations are created RDMA capability.UCS_TRYkeeps the oldcuMemAllocpath on systems where CUDA VMM fabric allocation is not available.