UCT/CUDA_COPY: Skip dmabuf probe for async managed memory - #11922
UCT/CUDA_COPY: Skip dmabuf probe for async managed memory#11922pentschev wants to merge 22 commits into
Conversation
CUDA async managed allocations can report CU_POINTER_ATTRIBUTE_IS_MANAGED with no owning CUDA context. Mark this case as async-managed so mem_flags detection keeps it non-registerable and does not call cuMemGetHandleForAddressRange(). This avoids a CUDA driver crash during UCP memory detection when UCX_CUDA_COPY_DMABUF is enabled. Add a CUDA 13 regression test that queries mem_flags for async managed mempool memory with dmabuf enabled.
|
🤖 Starting review — findings will be posted here when done. |
| * provided address and length as base address and alloc length | ||
| * respectively */ | ||
| mem_info->type = UCS_MEMORY_TYPE_CUDA_MANAGED; | ||
| if (cuda_mem_ctx == NULL) { |
There was a problem hiding this comment.
why not gate on md->config.cuda_async_managed here as well, like the sibling branch below? otherwise pool-managed memory is marked non-registrable even when the user sets CUDA_COPY_ASYNC_MEM_TYPE=cuda.
Also worth confirming that ordinary cuMemAllocManaged memory never reports a NULL context (which would otherwise become a real registration regression), and that the config-mismatch above is intentional.
There was a problem hiding this comment.
Done, added the same md->config.cuda_async_managed gate so CUDA_COPY_ASYNC_MEM_TYPE=cuda keeps this memory registrable.
| cudaMemAllocationTypeManaged)); | ||
| ASSERT_EQ(cudaSuccess, cudaStreamCreate(&stream)); | ||
|
|
||
| cudaError_t cuda_status = cudaMallocFromPoolAsync(&buffer, size, pool, |
There was a problem hiding this comment.
maybe skip instead of asserting here? a device/driver may not support managed mem pools even on CUDART>=13000, and this would turn into a hard test failure.
There was a problem hiding this comment.
Done, the test now skips if the default managed mempool cannot be queried or async managed mempool allocation is unsupported.
|
|
||
| cuda_status = cudaStreamSynchronize(stream); | ||
| if (cuda_status != cudaSuccess) { | ||
| EXPECT_EQ(cudaSuccess, cudaFreeAsync(buffer, stream)); |
There was a problem hiding this comment.
minor: pls cudaStreamSynchronize before cudaStreamDestroy here, to match the error path above and drain the async free.
|
🤖 Starting review — findings will be posted here when done. |
| * provided address and length as base address and alloc length | ||
| * respectively */ | ||
| mem_info->type = UCS_MEMORY_TYPE_CUDA_MANAGED; | ||
| if (cuda_mem_ctx == NULL) { |
There was a problem hiding this comment.
minor: pls add a short comment here explaining why a NULL context indicates an async/stream-ordered (pool) managed allocation; the sibling else-if branch has this reasoning but it's non-obvious in the is_managed path.
There was a problem hiding this comment.
Done, added a short comment explaining that managed pool allocations are stream-ordered and have no owning CUDA context.
| * provided address and length as base address and alloc length | ||
| * respectively */ | ||
| mem_info->type = UCS_MEMORY_TYPE_CUDA_MANAGED; | ||
| if (cuda_mem_ctx == NULL) { |
There was a problem hiding this comment.
why does a NULL context on genuine managed memory imply it's not registrable? regular cudaMallocManaged/cuMemAllocManaged memory is registrable, and with default PREF_LOC=cpu (sys_dev UNKNOWN) it used to report REGISTRABLE. can we confirm the driver always returns a non-NULL context for ordinary managed allocations, including when allocated under the pushed primary context in uct_cuda_copy_mem_alloc?
There was a problem hiding this comment.
Added positive coverage for ordinary cudaMallocManaged memory and uct_mem_alloc(UCS_MEMORY_TYPE_CUDA_MANAGED). Both assert a non-NULL CUDA context and verify the allocation remains registrable.
|
Test coverage gap: can we also add a positive test that ordinary |
|
🤖 Starting review — findings will be posted here when done. |
Respect CUDA_COPY_ASYNC_MEM_TYPE when marking managed allocations without a CUDA context as async-managed. Add coverage for ordinary cudaMallocManaged memory remaining registrable and make the async managed mempool test skip when the device or driver does not support that allocation path.
e502c49 to
1d53937
Compare
Done, added |
| * provided address and length as base address and alloc length | ||
| * respectively */ | ||
| mem_info->type = UCS_MEMORY_TYPE_CUDA_MANAGED; | ||
| if ((cuda_mem_ctx == NULL) && md->config.cuda_async_managed) { |
There was a problem hiding this comment.
why is this gated on md->config.cuda_async_managed? managed pool memory is typed as cuda-managed regardless of ASYNC_MEM_TYPE, so with UCX_CUDA_COPY_ASYNC_MEM_TYPE=cuda we would still report it as registrable.
There was a problem hiding this comment.
Removed the configuration dependency. The regression test now also runs with CUDA_COPY_ASYNC_MEM_TYPE=cuda, and managed-pool memory remains CUDA_MANAGED and non-registrable.
| if ((cuda_mem_ctx == NULL) && md->config.cuda_async_managed) { | ||
| /* Managed pool allocations are stream-ordered and do not have | ||
| * an owning CUDA context. */ | ||
| *is_async_managed = 1; |
There was a problem hiding this comment.
with the default PREF_LOC=cpu this caches sys_dev=UNKNOWN together with mem_flags=0, and ucp_memory_detect_internal() treats that combination as "no info" and falls back to ucp_memory_detect_slowpath() on every operation for these buffers. is the repeated re-detection acceptable?
There was a problem hiding this comment.
This is a valid cache-performance concern. The current cache format cannot distinguish an incomplete UCM allocation event from a completed MD query when both have unknown system device and zero flags. I kept that broader UCS/UCP cache representation change out of this focused bug-fix PR.
| void *buffer = nullptr; | ||
| cudaError_t cuda_status; | ||
|
|
||
| cuda_status = cudaMallocManaged(&buffer, size); |
There was a problem hiding this comment.
maybe use mem_buffer buffer(size, UCS_MEMORY_TYPE_CUDA_MANAGED) here (test/AGENTS.md: reuse mem_buffer for generic allocations); it also releases the buffer if an assertion fails. Currently raw cudaMallocManaged/cudaFree leaks if an assertion aborts the test.
There was a problem hiding this comment.
Replaced the raw cudaMallocManaged/cudaFree pair with mem_buffer, so the allocation is released if an assertion aborts the test.
| UCS_TEST_SKIP_R("failed to allocate async managed mempool memory"); | ||
| } | ||
|
|
||
| cuda_status = cudaMemsetAsync(buffer, 0, size, stream); |
There was a problem hiding this comment.
minor: the three copies of the free/sync/destroy cleanup hide what the test checks; can we wrap the stream and the pool buffer in RAII helpers and keep only the query here?
There was a problem hiding this comment.
Wrapped the stream-ordered managed-pool allocation and stream teardown in cuda_async_managed_mem_buffer. The test body now only performs setup, query, and assertions.
|
🤖 Starting review — findings will be posted here when done. |
| * provided address and length as base address and alloc length | ||
| * respectively */ | ||
| mem_info->type = UCS_MEMORY_TYPE_CUDA_MANAGED; | ||
| if ((cuda_mem_ctx == NULL) && md->config.cuda_async_managed) { |
There was a problem hiding this comment.
Registrability of managed-pool memory is gated on an unrelated config knob. In this branch is_managed is already true, so mem_info->type is UCS_MEMORY_TYPE_CUDA_MANAGED regardless of UCX_CUDA_COPY_ASYNC_MEM_TYPE (md->config.cuda_async_managed only decides the type of non-managed stream-ordered allocations, see the else if at line 758). With UCX_CUDA_COPY_ASYNC_MEM_TYPE=cuda — a configuration already used by existing tests in this file (no_current_context_cuda_registrable) — the same managed-pool buffer is reported as registrable again, i.e. the bug this PR fixes reappears. Registrability is a property of the allocation, not of the reported memory type.
Why gate on md->config.cuda_async_managed here? Suggested alternative:
| if ((cuda_mem_ctx == NULL) && md->config.cuda_async_managed) { | |
| if (cuda_mem_ctx == NULL) { | |
| /* Managed pool allocations are stream-ordered and do not have | |
| * an owning CUDA context. */ | |
| *is_async_managed = 1; | |
| } |
Alternatively, CU_POINTER_ATTRIBUTE_MEMPOOL_HANDLE (already queried in cuda_ipc_md.c) detects pool allocations directly, without the ctx heuristic.
| if ((cuda_mem_ctx == NULL) && md->config.cuda_async_managed) { | ||
| /* Managed pool allocations are stream-ordered and do not have | ||
| * an owning CUDA context. */ | ||
| *is_async_managed = 1; |
There was a problem hiding this comment.
Cached entry becomes indistinguishable from "not detected", forcing a per-operation slow path. When the preferred location resolves to the CPU (PREF_LOC=cpu is the default, and the new positive test uses it), the managed branch sets sys_dev = UCS_SYS_DEVICE_ID_UNKNOWN; with is_async_managed = 1 the cached mem_flags is now 0. ucp_memory_detect_internal() (src/ucp/core/ucp_context.h:746) treats the (sys_dev == UNKNOWN) && (mem_flags == 0) pair as "attributes not detected" and calls ucp_memory_detect_slowpath() on every operation with such a buffer. Before this change these entries were cached with UCS_MEM_FLAG_REGISTRABLE and hit the fast path. Functionally correct, but a hot-path cost for exactly the memory this PR targets — is the per-operation slow path acceptable here?
|
Residual coverage gap: Residual coverage gap: Neither new test covers |
| * provided address and length as base address and alloc length | ||
| * respectively */ | ||
| mem_info->type = UCS_MEMORY_TYPE_CUDA_MANAGED; | ||
| if ((cuda_mem_ctx == NULL) && md->config.cuda_async_managed) { |
There was a problem hiding this comment.
CU_POINTER_ATTRIBUTE_MEMPOOL_HANDLE output might be usable for detection and probably md->config.cuda_async_managed is not exactly related as it should apply to memory that is not is_managed where user can choose memory type to use, so it should be removed.
There was a problem hiding this comment.
Implemented this. The code now queries CU_POINTER_ATTRIBUTE_MEMPOOL_HANDLE and uses a non-NULL pool handle to identify managed-pool allocations. I also removed the unrelated md->config.cuda_async_managed gate from the managed-memory branch. The managed-pool regression test now covers CUDA_COPY_ASYNC_MEM_TYPE=cuda.
|
🤖 Starting review — findings will be posted here when done. |
| * provided address and length as base address and alloc length | ||
| * respectively */ | ||
| mem_info->type = UCS_MEMORY_TYPE_CUDA_MANAGED; | ||
| if (cuda_mempool != NULL) { |
There was a problem hiding this comment.
Hot-path concern: with the default PREF_LOC=cpu this path caches sys_dev=UNKNOWN together with mem_flags=0, which is exactly the condition that makes ucp_memory_detect_internal() call ucp_memory_detect_slowpath() (uct_md_mem_query_v2 + cuPointerGetAttributes + cuMemRangeGetAttribute + memtype-cache update under a spinlock) on every operation for these buffers. The pre-existing is_async_managed branch avoids this because it gets a real sys_dev; this new branch is the first to produce the UNKNOWN+0 combination. Can we avoid the per-op re-detection, e.g. by reporting the buffer's GPU sys_dev for managed-pool allocations?
There was a problem hiding this comment.
Fixed. Managed-pool allocations now use their pointer-reported GPU sys_dev, rather than the CPU preferred-location result. This prevents caching the UNKNOWN plus zero-flags combination and avoids repeated UCP slow-path detection. The regression test now asserts a known system device.
| attr_data[2] = &cuda_device; | ||
| attr_type[3] = CU_POINTER_ATTRIBUTE_CONTEXT; | ||
| attr_data[3] = &cuda_mem_ctx; | ||
| attr_type[4] = CU_POINTER_ATTRIBUTE_MEMPOOL_HANDLE; |
There was a problem hiding this comment.
minor: CU_POINTER_ATTRIBUTE_MEMPOOL_HANDLE and CUmemoryPool are CUDA 11.2+; every other use in the tree is guarded (#if CUDA_VERSION >= 11020 in ucm/cuda/cudamem.c, HAVE_CUDA_FABRIC/HAVE_DECL_SYS_PIDFD_GETFD in cuda_ipc_md.c / cuda_ipc_cache.c). This is the first unguarded use, so building against older CUDA headers would break — do we need a guard here too?
There was a problem hiding this comment.
Added #if CUDA_VERSION >= 11020 guards around the mempool type, fifth pointer attribute, and managed-pool detection path. Older CUDA headers retain the four-attribute query.
| * provided address and length as base address and alloc length | ||
| * respectively */ | ||
| mem_info->type = UCS_MEMORY_TYPE_CUDA_MANAGED; | ||
| if (cuda_mempool != NULL) { |
There was a problem hiding this comment.
minor (compactness): the if can be a direct assignment:
| if (cuda_mempool != NULL) { | |
| /* Managed-pool allocations are stream-ordered and cannot be | |
| * registered. */ | |
| *is_async_managed = (cuda_mempool != NULL); |
There was a problem hiding this comment.
Kept the branch because the managed-pool case now also sets the allocation’s GPU sys_dev, validates it, and skips the preferred-location query. A direct assignment would not express the required locality handling.
|
🤖 Starting review — findings will be posted here when done. |
| /* Managed-pool allocations are stream-ordered and cannot be | ||
| * registered. */ | ||
| *is_async_managed = 1; | ||
| mem_info->sys_dev = uct_cuda_get_sys_dev(cuda_device); |
There was a problem hiding this comment.
does this need the cuda_device == CU_DEVICE_CPU handling that the common path below has? for a host-located managed pool uct_cuda_get_sys_dev() returns UNKNOWN and then the whole query fails with UCS_ERR_NO_DEVICE.
There was a problem hiding this comment.
Fixed. The managed-pool path now mirrors the common CPU-device handling: it reports UCS_SYS_DEVICE_ID_UNKNOWN for CU_DEVICE_CPU instead of failing with UCS_ERR_NO_DEVICE.
| attr_type[3] = CU_POINTER_ATTRIBUTE_CONTEXT; | ||
| attr_data[3] = &cuda_mem_ctx; | ||
| #if CUDA_VERSION >= 11020 | ||
| attr_type[4] = CU_POINTER_ATTRIBUTE_MEMPOOL_HANDLE; |
There was a problem hiding this comment.
maybe reuse cuda_mempool for the (cuda_mem_ctx == NULL) heuristic below as well? it identifies stream-ordered allocations directly, so the context-based guess and its driver-bug comment could go away (fine as a follow-up).
There was a problem hiding this comment.
Leaving this as a follow-up. CU_POINTER_ATTRIBUTE_MEMPOOL_HANDLE is unavailable before CUDA 11.2, while the existing context heuristic remains the compatibility fallback for older CUDA headers. Removing it here would broaden this bug-fix PR and reduce coverage for that configuration.
|
Residual gap (previously reported): CUDART >= 13000 gating plus the concurrent-managed-access / default-managed-pool skips mean the negative tests may never actually execute in CI. Residual gap (previously reported): |
Add positive coverage for ordinary managed memory so CUDA managed allocations with an owning context keep reporting UCS_MEM_FLAG_REGISTRABLE. Cover both user cudaMallocManaged memory and memory allocated through uct_mem_alloc.
Use the CUDA mempool pointer attribute to identify managed stream-ordered pool allocations. Such allocations are not registrable regardless of CUDA_COPY_ASYNC_MEM_TYPE, so skip dmabuf probing for them. Extend the regression test to cover CUDA_COPY_ASYNC_MEM_TYPE=cuda and use RAII for CUDA test allocations.
Report the GPU system device for managed pool allocations so the non-registrable cache entry does not trigger UCP memory detection on every operation. Guard mempool pointer attributes for CUDA versions before 11.2.
Preserve the common CUDA CPU-device handling for managed pool allocations and cover GPU preferred location in the managed pool regression test.
c5fcf7b to
1b27615
Compare
|
🤖 Starting review — findings will be posted here when done. |
| * was not allocated in a context should also allow us to | ||
| * identify virtual/stream-ordered CUDA allocations. Keep this | ||
| * heuristic for non-managed allocations; managed-pool allocations | ||
| * are handled above. */ |
There was a problem hiding this comment.
maybe use cuda_mempool != NULL for this check too, and keep cuda_mem_ctx == NULL only as a fallback when the attribute is not available? the mempool handle is exact for stream-ordered allocations, unlike the driver-bug workaround described above.
| @@ -124,7 +124,7 @@ class mem_buffer { | |||
| /* Allocation mode. */ | |||
| enum class alloc_mode { | |||
| DEFAULT, /* Default allocation mode, using cudaMalloc */ | |||
There was a problem hiding this comment.
minor: DEFAULT is also used for cudaMallocManaged now, so the comment is stale.
| DEFAULT, /* Default allocation mode, using cudaMalloc */ | |
| DEFAULT, /* Default synchronous CUDA allocation mode */ |
| EXPECT_UCS_OK(uct_md_mem_query_v2(md(), buffer.ptr(), size, &mem_attr)); | ||
| EXPECT_EQ(UCS_MEMORY_TYPE_CUDA_MANAGED, mem_attr.mem_type); | ||
| EXPECT_NE(UCS_SYS_DEVICE_ID_UNKNOWN, mem_attr.sys_dev); | ||
| EXPECT_FALSE(mem_attr.mem_flags & UCS_MEM_FLAG_REGISTRABLE); |
There was a problem hiding this comment.
sorry, actually after discussing, should cuda-managed async from mem pool also be marked as registrable without the problematic dmabuf probing removing current mem pool probing changes and simply using something like below as it is real managed memory:
uct_cuda_copy_md_detect_mem_flags(uct_cuda_copy_md_t....)
{
...
if (is_async_managed) {
return 0;
}
/* Managed memory is never dmabuf-exportable, it is registered via ODP */
if (mem_info->type == UCS_MEMORY_TYPE_CUDA_MANAGED) {
return UCS_MEM_FLAG_REGISTRABLE;
}There was a problem hiding this comment.
Fixed. Removed mempool-handle probing and its configure check. CUDA-managed memory now returns UCS_MEM_FLAG_REGISTRABLE through ODP before DMABUF export is considered. The regression now expects async managed-pool memory to be registrable while retaining CUDA_COPY_DMABUF=try and PREF_LOC=gpu coverage of the original crash path.
|
🤖 CI Triage Agent — TL;DR: The Full analysisSummary: Step Root cause: Build, install and installcheck all passed. The last application output is at So the hang is in UCP/UCT setup or the first 128 KB tag transfer (rendezvous over Implicated commit: unknown (no evidence in the log ties the hang to a commit; branch head is File: Suggested fix:
Related: PR #11922 (#11922) — the PR under test; no existing issue found for this perftest hang signature.
|
|
🤖 Starting review — findings will be posted here when done. |
| } | ||
|
|
||
| /* Managed memory is registered through ODP and is not dmabuf-exportable. */ | ||
| if (mem_info->type == UCS_MEMORY_TYPE_CUDA_MANAGED) { |
There was a problem hiding this comment.
does this also flip CUDA device symbols to registrable? ucm/cuda/cudamem.c reports symbols (cudaGetSymbolAddress, cuModuleGetGlobal) as UCS_MEMORY_TYPE_CUDA_MANAGED with mem_flags=0, and uct_cuda_copy_md_mem_query() takes the managed type from the memtype cache, so this branch would report them REGISTRABLE and also overwrite the cached mem_flags via ucs_memtype_cache_update().
Rationale: mem_query() does addr_mem_info.type = cached_mem_info.type when the cached type is CUDA_MANAGED (the "device symbols" workaround), then computes flags from addr_mem_info.type and writes them back to the cache. Before this change the symbol range went through the dmabuf probe and stayed non-registrable when export failed.
There was a problem hiding this comment.
I think this comment is important. @pentschev #11646 made cuda symbols managed memory as a workaround, and the change we suggested made it so they will now be registerable if queried. The fix needs to make sure it doesn't set registerable to device symbols
There was a problem hiding this comment.
I think this is fixed now. mem_query now retains the driver-reported memory type for flag detection before applying the cached CUDA-managed type used to report device symbols. Therefore device symbols continue through the normal non-managed path, remain non-registrable, and their cached zero flags are preserved. I also extended test_cuda_get_symbol_address to verify this after mapping the symbol.
| } | ||
|
|
||
| UCS_TEST_P(test_mem_alloc_device, user_managed_mem_registrable, | ||
| "CUDA_COPY_PREF_LOC=cpu") |
There was a problem hiding this comment.
user_managed_mem_registrable, CUDA_COPY_PREF_LOC=cpu: with PREF_LOC=cpu the managed buffer gets sys_dev=UNKNOWN, so is_registrable() already returned true before this fix and the new branch is not exercised; can we add a PREF_LOC=gpu variant here too? that would also cover the fix on CUDA < 13, where all three async_managed_mem_pool_* tests are compiled out/skipped.
|
|
||
| mem_attr.field_mask = UCT_MD_MEM_ATTR_V2_FIELD_MEM_TYPE | | ||
| UCT_MD_MEM_ATTR_V2_FIELD_MEM_FLAGS; | ||
| EXPECT_UCS_OK(uct_md_mem_query_v2(md(), buffer.ptr(), size, &mem_attr)); |
There was a problem hiding this comment.
minor: these lines duplicate the tail of query_managed_registrable(); can we split the query + expectations into a small helper and call it from both, keeping only the CU_POINTER_ATTRIBUTE_CONTEXT assert in query_managed_registrable()?
| @@ -124,7 +124,7 @@ class mem_buffer { | |||
| /* Allocation mode. */ | |||
| enum class alloc_mode { | |||
| DEFAULT, /* Default allocation mode, using cudaMalloc */ | |||
There was a problem hiding this comment.
minor: the DEFAULT comment is now stale as well — that path also uses cudaMallocManaged.
| DEFAULT, /* Default allocation mode, using cudaMalloc */ | |
| DEFAULT, /* Default CUDA allocation mode */ | |
| ASYNC /* Asynchronous CUDA allocation mode */ |
| * provided address and length as base address and alloc length | ||
| * respectively */ | ||
| mem_info->type = UCS_MEMORY_TYPE_CUDA_MANAGED; | ||
|
|
| * here, but due to a bug in the driver `cudaMalloc` also returns | ||
| * false in that case. Therefore, checking whether the allocation | ||
| * was not allocated in a context should also allows us to | ||
| * was not allocated in a context should also allow us to |
|
🤖 Starting review — findings will be posted here when done. |
| if (status != UCS_OK) { | ||
| printf("memtype cache lookup failed: %s\n", ucs_status_string(status)); | ||
| ++num_errors; | ||
| } else if (cached_mem_info.mem_flags & UCS_MEM_FLAG_REGISTRABLE) { |
There was a problem hiding this comment.
does this hold when dmabuf is not supported? there uct_cuda_copy_md_is_registrable() returns 1 for the symbol (driver type is cuda with a real sys_dev), so the app fails on such a machine. maybe check the flags only when dmabuf is supported?
|
|
||
| /* Managed memory is registered through ODP and is not dmabuf-exportable. */ | ||
| if (mem_info->type == UCS_MEMORY_TYPE_CUDA_MANAGED) { | ||
| return UCS_MEM_FLAG_REGISTRABLE; |
There was a problem hiding this comment.
| return UCS_MEM_FLAG_REGISTRABLE; | |
| return 1; |
minor: the function is a boolean predicate, the other branches return 0/1.
| UCT_MD_MEM_ATTR_V2_FIELD_MEM_FLAGS; | ||
| EXPECT_UCS_OK(uct_md_mem_query_v2(md(), buffer.ptr(), size, &mem_attr)); | ||
| EXPECT_EQ(UCS_MEMORY_TYPE_CUDA_MANAGED, mem_attr.mem_type); | ||
| EXPECT_TRUE(mem_attr.mem_flags & UCS_MEM_FLAG_REGISTRABLE); |
There was a problem hiding this comment.
minor: can we reuse query_managed_registrable() here? the only difference is the context check, which can move to its caller.
| CUdevice avail_cuda_device = CU_DEVICE_INVALID; | ||
| ucs_memory_info_t cached_mem_info; | ||
| ucs_memory_info_t addr_mem_info; | ||
| ucs_memory_info_t detected_mem_info = {}; |
There was a problem hiding this comment.
nit: initialized variables go above unintialized variables.
Issue with symbols looks resolved
|
🤖 Starting review — findings will be posted here when done. |
| if (status != UCS_OK) { | ||
| printf("memtype cache lookup failed: %s\n", ucs_status_string(status)); | ||
| ++num_errors; | ||
| } else if (cached_mem_info.mem_flags & UCS_MEM_FLAG_REGISTRABLE) { |
There was a problem hiding this comment.
this expects dmabuf export to fail for the symbol, but when the build or device has no dmabuf support (CUDA < 11.7, or UCX_CUDA_COPY_DMABUF=n), uct_cuda_copy_md_is_registrable() returns 1 and this check fails. can we make the expectation independent of dmabuf support?
Rationale: uct_cuda_copy_md_is_dmabuf_supported() compiles the attribute query out below CUDA 11.7, so md->config.dmabuf_supported == 0 and is_registrable() returns 1 at the !md->config.dmabuf_supported short-circuit, i.e. the cached symbol flags would contain UCS_MEM_FLAG_REGISTRABLE. buildlib/pr/cuda/test_malloc_hook.sh runs this app unconditionally.
There was a problem hiding this comment.
Removed the zero-flags assertion.
| EXPECT_EQ(CUDA_SUCCESS, cuMemFree(dptr)); | ||
| } | ||
|
|
||
| UCS_TEST_P(test_mem_alloc_device, user_managed_mem_registrable, |
There was a problem hiding this comment.
with PREF_LOC=cpu (which is also the default) the managed branch sets sys_dev to unknown, so is_registrable() returns 1 before reaching the new CUDA_MANAGED check - this test and uct_alloc_managed_mem_registrable pass without the fix too. can we use gpu here?
| UCS_TEST_P(test_mem_alloc_device, user_managed_mem_registrable, | |
| UCS_TEST_P(test_mem_alloc_device, user_managed_mem_registrable, | |
| "CUDA_COPY_PREF_LOC=gpu") |
|
|
||
| void test_async_managed_mem_pool_registrable() | ||
| { | ||
| constexpr size_t size = 192; |
There was a problem hiding this comment.
minor: extra spaces, nothing to align with here.
| constexpr size_t size = 192; | |
| constexpr size_t size = 192; |
|
🤖 CI Triage Agent — TL;DR: The Coverity Full analysisSummary: Azure Pipelines job "Coverity coverity release on coverity_rh7" (build 135586) failed with Root cause: In
Note the sibling variable Implicated commit: File: Suggested fix: Zero-initialize the variable at declaration so the struct copy is well-defined: - ucs_memory_info_t addr_mem_info;
+ ucs_memory_info_t addr_mem_info = {};Alternatives, either of which also silences the checker and is arguably cleaner:
Option 1 is the most robust, since Related: PR #11922 (branch |
|
🤖 Starting review — findings will be posted here when done. |
Fixed. |
| int is_host_located = 0; | ||
| CUdevice cur_cuda_device = CU_DEVICE_INVALID; | ||
| CUdevice avail_cuda_device = CU_DEVICE_INVALID; | ||
| ucs_memory_info_t detected_mem_info = {}; |
There was a problem hiding this comment.
minor: can we drop the = {}? both structs are fully assigned before any field is read, and the initializers break the = alignment of the block above (docs/CodeStyle.md: consecutive assignments align on one column).
| ucs_memory_info_t detected_mem_info = {}; | |
| ucs_memory_info_t detected_mem_info; | |
| ucs_memory_info_t addr_mem_info; | |
| ucs_memory_info_t cached_mem_info; |
|
🤖 CI Triage Agent — TL;DR: The build failed on a single gtest, Full analysisSummary: Root cause: Flaky assertion in the fault-tolerance gtest, not a product regression from this PR. In Implicated commit: db208ee "UCP/FT: probe-gated lane recovery via aux uct_ep_check (#11563)", Evgeny Leksikov (introduced both the probe gating and this test). Timing likely aggravated by fb88400 "UCP/WIREUP: Add a token trailer and an ACK message to the lane address exchange (#11843)", Evgeny Leksikov, merged the same day (2026-09-10), which shortens the reply→probe→connect path. Not caused by File: Suggested fix: Make the probe observation deterministic instead of sampled:
Related: #11563 (introduced probe-gated recovery + test), #11843 (LANES_ADDR ACK/token trailer, changes recovery timing), #11397 (precedent: previously disabled a flaky fault-tolerance test) |
5 participants
What?
Fix CUDA async managed mempool memory detection in
cuda_copyand add a regression test.CUDA async managed allocations can report
CU_POINTER_ATTRIBUTE_IS_MANAGEDwith no owning CUDA context. UCX now marks that case as async-managed, somem_flagsdetection treats the memory as non-registerable and does not probe dmabuf export for it.Why?
With
UCX_CUDA_COPY_DMABUF=try, posting a UCP receive into CUDA 13 async managed mempool memory can crash insidecuMemGetHandleForAddressRange(). SettingUCX_CUDA_COPY_DMABUF=noavoids the crash because UCX skips the dmabuf export probe.The crash happens during UCP memory detection when
cuda_copyqueriesMEM_FLAGS: the allocation is managed memory, but because it has no CUDA context, it must be handled like other async managed allocations and kept off the dmabuf/registration path.How?
The fix is local to
cuda_copymemory attribute detection. WhencuPointerGetAttributes()reports managed memory and the pointer has no CUDA context, mark it asis_async_managed. The existing mem_flags logic then returns non-registerable flags without callingcuMemGetHandleForAddressRange().A CUDA 13 regression test allocates async managed memory from the default managed mempool with
cudaMallocFromPoolAsync(), queriesuct_md_mem_query_v2()withMEM_FLAGSwhileCUDA_COPY_DMABUF=try, and verifies the memory is classified ascuda-managedand not registrable.