UCM/CUDA/TEST: Install memory hooks for async Cuda allocations

[yosefe]

Why

As discussed in #7194 and #7110 , need to add memory hooks support for cuda async allocations. Without this, applications using these allocations may fail to detect Cuda memory and run into segfault/access error.

[yosefe]

@Akshay-Venkatesh WDYT?

[yosefe]

/azp run

[azure-pipelines]

Azure Pipelines successfully started running 2 pipeline(s).

[Akshay-Venkatesh]

Does this mean that HAVE_CUDA is not set if *Async APIs aren't detected at configure time? That would disallow CUDA for slightly older versions of cuda wouldn't it?

I'm probably missing the commit that defines HAVE_DECL_CUMEMALLOCASYNC/HAVE_DECL_CUMEMFREEASYNC

[yosefe]

it does not affect HAVE_CUDA . it sets a different set of macros, specific for async APIs

[Akshay-Venkatesh]

@yosefe I think we should also intercept cuMemAllocFromPoolAsync

[yosefe]

ok, will add

[Akshay-Venkatesh]

For now this is fine because cuMemAllocAsync can only allocate pinned memory but setting default memory pool to user created pool can alter the behavior in the future when other memory types are supported.

In the future, it would be better to get memory pool associated with current device and examine allocation properties to decide the memory type instead of hard coding to MEMORY_TYPE_CUDA as the same API may be used for other memory types as well. I don't see an API to get MemPool properties from MemPool yet so we'll need to intercept MemPoolCreate/Destroy API for this.

[yosefe]

or we can just set memory type to UKNOWN like we planned to anyway?

[Akshay-Venkatesh]

Technically, the memory should be freed when the stream moves past FreeAsync. When the API itself returns, this may not be true so in that sense it may not be exactly right to change the attributes of the memory range or remove the memory range from pointer cache. But as we don't have a callback per se when free actually occurs, this should be ok for now as users would be very unlikely to issue ucp transfer operations after freeasync knowing that it may not be actually freed yet.

[yosefe]

yes, i guess once this is submitted it's no longer legal to issue data transfer from CPU.
do you know at which exact point the GPU can map a new physical memory to same virtual address?

[Akshay-Venkatesh]

It would have to be at the next cu*Alloc* call. Since we intercept all of those, I guess we don't have to worry about stream semantics on FreeAsync.

[Akshay-Venkatesh]

@yosefe forgot to bring up the issue of lack of sync memops support on MallocAsync memory that may come up because of this PR. Adding this PR would likely result in IB or cuda-ipc UCTs to be used to move memory allocated through MallocAsync but the following sequence could lead to stale data being transferred:

cudaMallocAsync(&x, length1, stream1);
cudaStreamSynchromize(stream1);
...
cudaMemcpy(x, y, length2, cudaMemcpyHostToDevice); // potentially non-blocking wrt CPU and copy to destination x may still be in flight
ucp_tag_send_nbx(x, ...); // region pointed by x is not valid yet because previous memcpy is still in flight

Setting sync memops attribute on x would synchronize all outstanding memory operations on it but it's not supported on MallocAsync memory so this could lead to data validation issues irrespective of zcopy operations through ib/cuda_ipc or through pipeline protocols.

[simonbyrne]

Any update on this?

[Akshay-Venkatesh]

Any update on this?

@simonbyrne SYNC_MEMOPS is still yet to be supported with Malloc Async API. We plan to support such memory once it becomes available.

[yosefe]

replaced by #8623