Skip to content
Merged
Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
22 commits
Select commit Hold shift + click to select a range
d72bcfb
UCT/CUDA_COPY: Skip dmabuf probe for async managed memory
pentschev Sep 8, 2026
65e7d2c
Merge remote-tracking branch 'upstream/master' into cuda-async-skip-d…
pentschev Sep 8, 2026
1d53937
UCT/CUDA_COPY: Address async managed memory review
pentschev Sep 8, 2026
458aeca
GTEST/CUDA: Cover managed memory registration flags
pentschev Sep 8, 2026
31617b6
UCT/CUDA_COPY: Detect managed memory pools
pentschev Sep 8, 2026
0fb53b3
UCT/CUDA_COPY: Preserve managed pool device locality
pentschev Sep 8, 2026
1b27615
UCT/CUDA_COPY: Handle managed pool CPU devices
pentschev Sep 8, 2026
12b3116
UCT/CUDA_COPY: Keep managed-pool detection resilient
pentschev Sep 8, 2026
07409ae
UCT/CUDA_COPY: Use mempool attributes for detection
pentschev Sep 8, 2026
cfc5a8e
UCT/CUDA_COPY: Detect managed memory pools by capability
pentschev Sep 9, 2026
a041773
Merge remote-tracking branch 'upstream/master' into cuda-async-skip-d…
pentschev Sep 9, 2026
ad87c78
GTEST/CUDA: Reuse managed default pool lookup
pentschev Sep 9, 2026
13e872f
UCT/CUDA_COPY: Track queried attribute count
pentschev Sep 9, 2026
2b73283
UCT/CUDA_COPY: Clarify managed pool detection
pentschev Sep 9, 2026
3773519
GTEST/CUDA: Align managed pool test variables
pentschev Sep 9, 2026
15e7781
GTEST/CUDA: Clear managed pool probe errors
pentschev Sep 9, 2026
16a911b
GTEST/CUDA: Share async release path
pentschev Sep 9, 2026
f473f14
UCT/CUDA_COPY: Register managed memory without dmabuf
pentschev Sep 9, 2026
6ece985
Merge remote-tracking branch 'upstream/master' into cuda-async-skip-d…
pentschev Sep 10, 2026
a7cbb96
UCT/CUDA_COPY: Preserve device symbol flags
pentschev Sep 10, 2026
437da17
UCT/CUDA_COPY: Simplify managed memory checks
pentschev Sep 10, 2026
fae7ff1
UCT/CUDA_COPY: Initialize memory query result
pentschev Sep 10, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 3 additions & 0 deletions src/uct/cuda/cuda_copy/cuda_copy_md.c
Original file line number Diff line number Diff line change
Expand Up @@ -730,6 +730,9 @@ uct_cuda_copy_md_query_attributes(const uct_cuda_copy_md_t *md,
* provided address and length as base address and alloc length
* respectively */
mem_info->type = UCS_MEMORY_TYPE_CUDA_MANAGED;
if (cuda_mem_ctx == NULL) {

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

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.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Done, added the same md->config.cuda_async_managed gate so CUDA_COPY_ASYNC_MEM_TYPE=cuda keeps this memory registrable.

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

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.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Done, added a short comment explaining that managed pool allocations are stream-ordered and have no owning CUDA context.

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

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?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

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.

*is_async_managed = 1;

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

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?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

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.

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

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?

}

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.

remove unrelated change

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Fixed.

cu_err = cuMemRangeGetAttribute(
(void*)&pref_loc, sizeof(pref_loc),
Expand Down
63 changes: 63 additions & 0 deletions test/gtest/uct/cuda/test_switch_cuda_device.cc
Original file line number Diff line number Diff line change
Expand Up @@ -441,6 +441,69 @@ UCS_TEST_P(test_mem_alloc_device, no_current_context_user_mem_registrable,
EXPECT_EQ(CUDA_SUCCESS, cuMemFree(dptr));
}

#if CUDART_VERSION >= 13000
UCS_TEST_P(test_mem_alloc_device, async_managed_mem_pool_not_registrable,
"CUDA_COPY_DMABUF=try")
{
constexpr size_t size = 192;
cudaMemLocation location = {};
cudaStream_t stream = nullptr;
cudaMemPool_t pool = nullptr;
void *buffer = nullptr;
int cuda_device;
int concurrent_managed_access;
uct_md_mem_attr_v2_t mem_attr = {};

ASSERT_EQ(cudaSuccess, cudaGetDevice(&cuda_device));
ASSERT_EQ(cudaSuccess,
cudaDeviceGetAttribute(&concurrent_managed_access,
cudaDevAttrConcurrentManagedAccess,
cuda_device));
if (!concurrent_managed_access) {
UCS_TEST_SKIP_R("device does not support concurrent managed access");
}

location.type = cudaMemLocationTypeDevice;
location.id = cuda_device;
ASSERT_EQ(cudaSuccess,
cudaMemGetDefaultMemPool(&pool, &location,
cudaMemAllocationTypeManaged));
ASSERT_EQ(cudaSuccess, cudaStreamCreate(&stream));

cudaError_t cuda_status = cudaMallocFromPoolAsync(&buffer, size, pool,

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

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.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Done, the test now skips if the default managed mempool cannot be queried or async managed mempool allocation is unsupported.

stream);
if (cuda_status != cudaSuccess) {
EXPECT_EQ(cudaSuccess, cudaStreamDestroy(stream));
ASSERT_EQ(cudaSuccess, cuda_status);
}

cuda_status = cudaMemsetAsync(buffer, 0, size, stream);

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

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?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

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.

if (cuda_status != cudaSuccess) {
EXPECT_EQ(cudaSuccess, cudaFreeAsync(buffer, stream));
EXPECT_EQ(cudaSuccess, cudaStreamSynchronize(stream));
EXPECT_EQ(cudaSuccess, cudaStreamDestroy(stream));
ASSERT_EQ(cudaSuccess, cuda_status);
}

cuda_status = cudaStreamSynchronize(stream);
if (cuda_status != cudaSuccess) {
EXPECT_EQ(cudaSuccess, cudaFreeAsync(buffer, stream));

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

minor: pls cudaStreamSynchronize before cudaStreamDestroy here, to match the error path above and drain the async free.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Done.

EXPECT_EQ(cudaSuccess, cudaStreamDestroy(stream));
ASSERT_EQ(cudaSuccess, cuda_status);
}

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, size, &mem_attr));
EXPECT_EQ(UCS_MEMORY_TYPE_CUDA_MANAGED, mem_attr.mem_type);
EXPECT_FALSE(mem_attr.mem_flags & UCS_MEM_FLAG_REGISTRABLE);

EXPECT_EQ(cudaSuccess, cudaFreeAsync(buffer, stream));
EXPECT_EQ(cudaSuccess, cudaStreamSynchronize(stream));
EXPECT_EQ(cudaSuccess, cudaStreamDestroy(stream));
}
#endif

UCS_TEST_P(test_mem_alloc_device, no_current_context_vmm_mem_registrable,
"CUDA_COPY_ASYNC_MEM_TYPE=cuda")
{
Expand Down
Loading