Skip to content

UCM: Propagate memory attributes in allocation events - v1.22.x #11655

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
gleon99 merged 2 commits into from
Jul 15, 2026
Merged

UCM: Propagate memory attributes in allocation events - v1.22.x #11655

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
0ce9ebc
UCM: Propagate memory attributes in allocation events - v1.22.x
gleon99 Jul 15, 2026
844a16f
UCM: CR fix
gleon99 Jul 15, 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
5 changes: 5 additions & 0 deletions src/ucm/api/ucm.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,6 +15,7 @@ BEGIN_C_DECLS

#include <ucs/config/types.h>
#include <ucs/memory/memory_type.h>
#include <ucs/sys/topo/base/topo.h>
#include <ucs/type/status.h>

#include <sys/types.h>
Expand Down Expand Up @@ -192,11 +193,15 @@ typedef union ucm_event {
* Memory type allocation and deallocation event.
* If mem_type is @ref UCS_MEMORY_TYPE_LAST, the memory type is unknown, and
* further memory type detection is required.
* If sys_dev is @ref UCS_SYS_DEVICE_ID_UNKNOWN and mem_flags is 0, the
* remaining memory attributes are unknown and require detection.
*/
struct {
void *address;
size_t size;
ucs_memory_type_t mem_type;
ucs_sys_device_t sys_dev;
uint8_t mem_flags;
} mem_type;

} ucm_event_t;
Expand Down
59 changes: 43 additions & 16 deletions src/ucm/cuda/cudamem.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -78,7 +78,9 @@
if (ret == CUDA_SUCCESS) { \
ucm_trace("%s(size_ptr=%p, obj=%p, name=%s) returned dptr=%p", \
__func__, bytes, obj, name, (void*)(*dptr)); \
ucm_cuda_dispatch_mem_alloc(*dptr, *size_ptr); \
ucm_cuda_dispatch_mem_type_alloc(*dptr, *size_ptr, \
UCS_MEMORY_TYPE_CUDA_MANAGED, \
0); \
} \
ucm_event_leave(); \
return ret; \
Expand Down Expand Up @@ -153,18 +155,27 @@ UCM_DEFINE_REPLACE_DLSYM_PTR_FUNC(cudaMallocPitch, cudaError_t, -1, void**,
UCM_DEFINE_REPLACE_DLSYM_PTR_FUNC(cudaGetSymbolAddress, cudaError_t, -1, void**,
const void*)

static void ucm_cuda_dispatch_mem_alloc(CUdeviceptr ptr, size_t length)
static void ucm_cuda_dispatch_mem_type_alloc(CUdeviceptr ptr, size_t length,
ucs_memory_type_t mem_type,
uint8_t mem_flags)
{
ucm_event_t event;

event.mem_type.address = (void*)ptr;
event.mem_type.size = length;
event.mem_type.mem_type = UCS_MEMORY_TYPE_LAST; /* indicate unknown type
and let cuda_md detect
attributes */
event.mem_type.address = (void*)ptr;
event.mem_type.size = length;
event.mem_type.mem_type = mem_type;
event.mem_type.sys_dev = UCS_SYS_DEVICE_ID_UNKNOWN;
event.mem_type.mem_flags = mem_flags;
ucm_event_dispatch(UCM_EVENT_MEM_TYPE_ALLOC, &event);
}

static void ucm_cuda_dispatch_mem_alloc(CUdeviceptr ptr, size_t length)
{
/* Indicate unknown type and let cuda_md detect attributes. */
ucm_cuda_dispatch_mem_type_alloc(ptr, length, UCS_MEMORY_TYPE_LAST,
UCS_MEM_FLAG_REGISTRABLE);
}

static void ucm_cuda_dispatch_mem_free(CUdeviceptr ptr, size_t length,
ucs_memory_type_t mem_type,
const char *func_name)
Expand All @@ -191,9 +202,11 @@ static void ucm_cuda_dispatch_mem_free(CUdeviceptr ptr, size_t length,
}
}

event.mem_type.address = (void*)ptr;
event.mem_type.size = length;
event.mem_type.mem_type = mem_type;
event.mem_type.address = (void*)ptr;
event.mem_type.size = length;
event.mem_type.mem_type = mem_type;
event.mem_type.sys_dev = UCS_SYS_DEVICE_ID_UNKNOWN;
event.mem_type.mem_flags = UCS_MEM_FLAG_REGISTRABLE;
ucm_event_dispatch(UCM_EVENT_MEM_TYPE_FREE, &event);
}

Expand Down Expand Up @@ -294,9 +307,21 @@ UCM_CUDA_ALLOC_FUNC(cudaMallocFromPoolAsync, cudaError_t, cudaSuccess, arg0,
void*, *, "size=%zu pool=%p stream=%p", size_t,
cudaMemPool_t, cudaStream_t)
#endif
UCM_CUDA_ALLOC_FUNC(cudaGetSymbolAddress, cudaError_t, cudaSuccess,
ucm_cuda_get_symbol_size(arg0), void*, *, "symbol=%p",
const void*)
cudaError_t ucm_cudaGetSymbolAddress(void **devPtr, const void *symbol)
{
cudaError_t ret;

ucm_event_enter();
ret = ucm_orig_cudaGetSymbolAddress(devPtr, symbol);
if (ret == cudaSuccess) {
ucm_trace("%s(symbol=%p) allocated %p", __func__, symbol, *devPtr);
ucm_cuda_dispatch_mem_type_alloc((CUdeviceptr)*devPtr,
ucm_cuda_get_symbol_size(symbol),
UCS_MEMORY_TYPE_CUDA_MANAGED, 0);
}
ucm_event_leave();
return ret;
}
UCM_CUDA_FREE_FUNC(cudaFree, UCS_MEMORY_TYPE_CUDA, cudaError_t, arg0, 0,
"devPtr=%p", void*)
UCM_CUDA_FREE_FUNC(cudaFreeHost, UCS_MEMORY_TYPE_HOST, cudaError_t, arg0, 0,
Expand Down Expand Up @@ -438,9 +463,11 @@ static int ucm_cudamem_scan_regions_cb(void *arg, void *addr, size_t length,
ucm_debug("dispatching initial memtype allocation for %p..%p %s", addr,
UCS_PTR_BYTE_OFFSET(addr, length), path);

event.mem_type.address = addr;
event.mem_type.size = length;
event.mem_type.mem_type = UCS_MEMORY_TYPE_LAST; /* unknown memory type */
event.mem_type.address = addr;
event.mem_type.size = length;
event.mem_type.mem_type = UCS_MEMORY_TYPE_LAST; /* unknown memory type */
event.mem_type.sys_dev = UCS_SYS_DEVICE_ID_UNKNOWN;
event.mem_type.mem_flags = UCS_MEM_FLAG_REGISTRABLE;

ucm_event_enter();
handler->cb(UCM_EVENT_MEM_TYPE_ALLOC, &event, handler->arg);
Expand Down
24 changes: 15 additions & 9 deletions src/ucm/rocm/rocmmem.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,9 +36,11 @@ ucm_dispatch_mem_type_alloc(void *addr, size_t length, ucs_memory_type_t mem_typ
{
ucm_event_t event;

event.mem_type.address = addr;
event.mem_type.size = length;
event.mem_type.mem_type = mem_type;
event.mem_type.address = addr;
event.mem_type.size = length;
event.mem_type.mem_type = mem_type;
event.mem_type.sys_dev = UCS_SYS_DEVICE_ID_UNKNOWN;
event.mem_type.mem_flags = UCS_MEM_FLAG_REGISTRABLE;
ucm_event_dispatch(UCM_EVENT_MEM_TYPE_ALLOC, &event);
}

Expand All @@ -47,9 +49,11 @@ ucm_dispatch_mem_type_free(void *addr, size_t length, ucs_memory_type_t mem_type
{
ucm_event_t event;

event.mem_type.address = addr;
event.mem_type.size = length;
event.mem_type.mem_type = mem_type;
event.mem_type.address = addr;
event.mem_type.size = length;
event.mem_type.mem_type = mem_type;
event.mem_type.sys_dev = UCS_SYS_DEVICE_ID_UNKNOWN;
event.mem_type.mem_flags = UCS_MEM_FLAG_REGISTRABLE;
ucm_event_dispatch(UCM_EVENT_MEM_TYPE_FREE, &event);
}

Expand Down Expand Up @@ -183,9 +187,11 @@ static int ucm_rocm_scan_regions_cb(void *arg, void *addr, size_t length,
ucm_debug("dispatching initial memtype allocation for %p..%p %s", addr,
UCS_PTR_BYTE_OFFSET(addr, length), path);

event.mem_type.address = addr;
event.mem_type.size = length;
event.mem_type.mem_type = UCS_MEMORY_TYPE_LAST; /* unknown memory type */
event.mem_type.address = addr;
event.mem_type.size = length;
event.mem_type.mem_type = UCS_MEMORY_TYPE_LAST; /* unknown memory type */
event.mem_type.sys_dev = UCS_SYS_DEVICE_ID_UNKNOWN;
event.mem_type.mem_flags = UCS_MEM_FLAG_REGISTRABLE;

ucm_event_enter();
handler->cb(UCM_EVENT_MEM_TYPE_ALLOC, &event, handler->arg);
Expand Down
24 changes: 15 additions & 9 deletions src/ucm/ze/zemem.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,9 +45,11 @@ ucm_dispatch_mem_type_alloc(void *addr, size_t length,
{
ucm_event_t event;

event.mem_type.address = addr;
event.mem_type.size = length;
event.mem_type.mem_type = mem_type;
event.mem_type.address = addr;
event.mem_type.size = length;
event.mem_type.mem_type = mem_type;
event.mem_type.sys_dev = UCS_SYS_DEVICE_ID_UNKNOWN;
event.mem_type.mem_flags = UCS_MEM_FLAG_REGISTRABLE;
ucm_event_dispatch(UCM_EVENT_MEM_TYPE_ALLOC, &event);
}

Expand All @@ -57,9 +59,11 @@ ucm_dispatch_mem_type_free(void *addr, size_t length,
{
ucm_event_t event;

event.mem_type.address = addr;
event.mem_type.size = length;
event.mem_type.mem_type = mem_type;
event.mem_type.address = addr;
event.mem_type.size = length;
event.mem_type.mem_type = mem_type;
event.mem_type.sys_dev = UCS_SYS_DEVICE_ID_UNKNOWN;
event.mem_type.mem_flags = UCS_MEM_FLAG_REGISTRABLE;
ucm_event_dispatch(UCM_EVENT_MEM_TYPE_FREE, &event);
}

Expand Down Expand Up @@ -234,9 +238,11 @@ static int ucm_zemem_scan_regions_cb(void *arg, void *addr, size_t length,
ucm_debug("dispatching initial memtype allocation for %p..%p %s", addr,
UCS_PTR_BYTE_OFFSET(addr, length), path);

event.mem_type.address = addr;
event.mem_type.size = length;
event.mem_type.mem_type = UCS_MEMORY_TYPE_LAST; /* unknown memory type */
event.mem_type.address = addr;
event.mem_type.size = length;
event.mem_type.mem_type = UCS_MEMORY_TYPE_LAST; /* unknown memory type */
event.mem_type.sys_dev = UCS_SYS_DEVICE_ID_UNKNOWN;
event.mem_type.mem_flags = UCS_MEM_FLAG_REGISTRABLE;

ucm_event_enter();
handler->cb(UCM_EVENT_MEM_TYPE_ALLOC, &event, handler->arg);
Expand Down
8 changes: 5 additions & 3 deletions src/ucp/core/ucp_context.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -736,9 +736,11 @@ ucp_memory_detect_internal(ucp_context_h context, const void *address,
address, length);
goto out_host_mem;
} else if (ucs_likely(status == UCS_OK)) {
if (ucs_unlikely(mem_info->type == UCS_MEMORY_TYPE_UNKNOWN)) {
ucs_trace_req(
"address %p length %zu: memtype cache returned 'unknown'",
if (ucs_unlikely(
(mem_info->type == UCS_MEMORY_TYPE_UNKNOWN) ||
((mem_info->sys_dev == UCS_SYS_DEVICE_ID_UNKNOWN) &&
(mem_info->mem_flags == 0)))) {
ucs_trace_req("address %p length %zu: querying memory attributes",
address, length);
ucp_memory_detect_slowpath(context, address, length, mem_info);
} else {
Expand Down
14 changes: 6 additions & 8 deletions src/ucs/memory/memtype_cache.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -310,19 +310,17 @@ static void ucs_memtype_cache_event_callback(ucm_event_type_t event_type,
return;
}

ucs_trace("dispatching mem event %d address %p length %zu mem_type %s",
ucs_trace("dispatching mem event %d address %p length %zu mem_type %s "
"sys_dev %u flags 0x%x",
event_type, event->mem_type.address, event->mem_type.size,
ucs_memory_type_names[event->mem_type.mem_type]);
ucs_memory_type_names[event->mem_type.mem_type],
event->mem_type.sys_dev, event->mem_type.mem_flags);

/* UCM allocation events carry only the memory type, not registrability.
* Default to "registrable": for a definite type this matches all current
* providers (e.g. ZE), and for UCS_MEMORY_TYPE_LAST/UNKNOWN the entry is
* re-detected via the MD, which overwrites this default with real flags. */
ucs_memtype_cache_update_internal(arg, event->mem_type.address,
event->mem_type.size,
event->mem_type.mem_type,
UCS_SYS_DEVICE_ID_UNKNOWN,
UCS_MEM_FLAG_REGISTRABLE, action);
event->mem_type.sys_dev,
event->mem_type.mem_flags, action);
}

static void ucs_memtype_cache_purge(ucs_memtype_cache_t *memtype_cache)
Expand Down
6 changes: 4 additions & 2 deletions src/uct/cuda/cuda_copy/cuda_copy_ep.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -86,7 +86,9 @@ uct_cuda_copy_get_mem_type(uct_md_h md, const void *address, size_t length,
}

if (ucs_unlikely((status == UCS_ERR_UNSUPPORTED) ||
(mem_info.type == UCS_MEMORY_TYPE_UNKNOWN))) {
(mem_info.type == UCS_MEMORY_TYPE_UNKNOWN) ||
((mem_info.sys_dev == UCS_SYS_DEVICE_ID_UNKNOWN) &&
(mem_info.mem_flags == 0)))) {
mem_attr.field_mask = UCT_MD_MEM_ATTR_V2_FIELD_MEM_TYPE |
UCT_MD_MEM_ATTR_V2_FIELD_SYS_DEV;

Expand Down Expand Up @@ -252,6 +254,7 @@ static UCS_F_ALWAYS_INLINE ucs_status_t uct_cuda_copy_ctx_rsc_get(
return status;
}


static UCS_F_ALWAYS_INLINE ucs_status_t uct_cuda_copy_ep_get_ctx(
uct_cuda_copy_iface_t *iface, const void *src, const void *dst,
size_t length, uct_cuda_copy_ep_ctx_t *ctx_p)
Expand Down Expand Up @@ -467,4 +470,3 @@ UCS_PROFILE_FUNC(ucs_status_t, uct_cuda_copy_ep_get_short,
length, (void *)remote_addr, buffer);
return status;
}

46 changes: 38 additions & 8 deletions test/apps/test_cuda_get_symbol_address.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,9 +16,15 @@
__device__ int device_int;


typedef struct test_ctx {
int num_events;
int num_errors;
} test_ctx_t;


static void event_cb(ucm_event_type_t event_type, ucm_event_t *event, void *arg)
{
int *count_p = (int *)arg;
test_ctx_t *ctx = (test_ctx_t *)arg;
const char *title;

if (event_type == UCM_EVENT_MEM_TYPE_ALLOC) {
Expand All @@ -33,7 +39,29 @@ static void event_cb(ucm_event_type_t event_type, ucm_event_t *event, void *arg)
printf("%s %s address %p size %zu\n", title,
ucs_memory_type_names[event->mem_type.mem_type],
event->mem_type.address, event->mem_type.size);
++(*count_p);
++ctx->num_events;

if ((event_type == UCM_EVENT_MEM_TYPE_ALLOC) &&
(event->mem_type.mem_type != UCS_MEMORY_TYPE_CUDA_MANAGED)) {
printf("unexpected symbol memory type %s, expected %s\n",
ucs_memory_type_names[event->mem_type.mem_type],
ucs_memory_type_names[UCS_MEMORY_TYPE_CUDA_MANAGED]);
++ctx->num_errors;
}

if ((event_type == UCM_EVENT_MEM_TYPE_ALLOC) &&
(event->mem_type.sys_dev != UCS_SYS_DEVICE_ID_UNKNOWN)) {
printf("unexpected symbol system device %u, expected unknown\n",
event->mem_type.sys_dev);
++ctx->num_errors;
}

if ((event_type == UCM_EVENT_MEM_TYPE_ALLOC) &&
(event->mem_type.mem_flags != 0)) {
printf("unexpected symbol memory flags 0x%x, expected 0\n",
event->mem_type.mem_flags);
++ctx->num_errors;
}
}

int main(int argc, char **argv)
Expand All @@ -44,7 +72,7 @@ int main(int argc, char **argv)
ucp_context_h context;
ucs_status_t status;
ucp_params_t params;
int num_events;
test_ctx_t ctx;
void *dptr;
cudaError_t res;

Expand All @@ -56,15 +84,17 @@ int main(int argc, char **argv)
return -1;
}

num_events = 0;
ucm_set_event_handler(memtype_events, 1000, event_cb, &num_events);
ctx.num_events = 0;
ctx.num_errors = 0;
ucm_set_event_handler(memtype_events, 1000, event_cb, &ctx);

res = cudaGetSymbolAddress(&dptr, device_int);
printf("cudaGetSymbolAddress() returned %p result %d\n", dptr, res);

ucm_unset_event_handler(memtype_events, event_cb, &num_events);
printf("got %d/%d memory events\n", num_events, num_expected_events);
ucm_unset_event_handler(memtype_events, event_cb, &ctx);
printf("got %d/%d memory events\n", ctx.num_events, num_expected_events);

ucp_cleanup(context);
return (num_events == num_expected_events) ? 0 : -1;
return ((ctx.num_events == num_expected_events) &&
(ctx.num_errors == 0)) ? 0 : -1;
}
24 changes: 24 additions & 0 deletions test/gtest/ucs/test_memtype_cache.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -479,6 +479,30 @@ UCS_TEST_P(test_memtype_cache, diff_mem_types_diff_bufs_keep_mem) {
test_memtype_cache_alloc_diff_mem_types(true, false);
}

UCS_TEST_P(test_memtype_cache, event_attributes) {
const size_t size = ucs_get_page_size();
void *address = reinterpret_cast<void*>(0xdead0000ul);
const ucs_sys_device_t dev = 0;
ucm_event_t event;
ucs_memory_info_t mem_info;

event.mem_type.address = address;
event.mem_type.size = size;
event.mem_type.mem_type = UCS_MEMORY_TYPE_CUDA_MANAGED;
event.mem_type.sys_dev = dev;
event.mem_type.mem_flags = 0;
ucm_event_dispatch(UCM_EVENT_MEM_TYPE_ALLOC, &event);

ASSERT_UCS_OK(ucs_memtype_cache_lookup(address, size, &mem_info));
EXPECT_EQ(event.mem_type.mem_type, mem_info.type);
EXPECT_EQ(dev, mem_info.sys_dev);
EXPECT_EQ(0, mem_info.mem_flags);

ucm_event_dispatch(UCM_EVENT_MEM_TYPE_FREE, &event);
EXPECT_EQ(UCS_ERR_NO_ELEM,
ucs_memtype_cache_lookup(address, size, &mem_info));
}

INSTANTIATE_TEST_SUITE_P(mem_type, test_memtype_cache,
::testing::ValuesIn(mem_buffer::supported_mem_types()));

Expand Down
Loading