Replies: 12 comments 13 replies
-
|
If the application knows this, it makes sense to pass that data through. But does middleware know what call was used for the allocation? Can it get this through the buffer attributes? |
Beta Was this translation helpful? Give feedback.
-
|
Both CUDA and L0 have calls to find this out so it's possible for the middleware to do this discovery if they don't have the information already to pass through. The same call that the hooking provider will need to do to register the memory will have the type of memory as well to use in the registration. |
Beta Was this translation helpful? Give feedback.
-
|
Does this data come 'for free' as part of querying if a buffer is a GPU or host buffer? Or is it an extra call? |
Beta Was this translation helpful? Give feedback.
-
|
Fo free! |
Beta Was this translation helpful? Give feedback.
-
|
An 'easy' solution is to extend enum fi_hmem_iface with device allocation details. For example: This may add a bunch of enum values. However, I'm not sure if additional details would be necessary or not, based on which call was used and how it was invoked. As an alternative, rather than indicate which call was used, we could pass in the necessary information more directly. Using a flag, rather than extending the enum seems to make more sense here: These flags would be passed into fi_mr_regattr. I don't like have 2 flags, but I think we need both in order to indicate that the IPC'ish is unknown to the caller. |
Beta Was this translation helpful? Give feedback.
-
|
In addition to whether a buffer supports IPC, I think we may need to distinguish between managed (cudaMallocManaged) and non-managed (cudaMalloc). With MOFED, RDMA to CUDA managed buffers cannot be supported due to the following: CUDA Unified Memory is not explicitly supported in combination with GPUDirect RDMA. While the page table returned by nvidia_p2p_get_pages() is valid for managed memory buffers and provides a mapping of GPU memory at any given moment in time, the GPU device copy of that memory may be incoherent with the writable copy of the page which is not on the GPU. Using the page table in this circumstance may result in accessing stale data, or data loss, because of a DMA write access to device memory that is subsequently overwritten by the Unified Memory run-time. cuPointerGetAttribute() may be used to determine if an address is being managed by the Unified Memory runtime. |
Beta Was this translation helpful? Give feedback.
-
|
Is libfabric support for AMD GPU card ? Is it will work for windows. |
Beta Was this translation helpful? Give feedback.
-
|
AMD GPUs are supported. I don't know about windows support. I doubt it, and I'm not sure what it would take to support that. The code paths for Linux and Windows are mostly shared, but I don't know that Windows has a similar lower-level functionality that's needed. |
Beta Was this translation helpful? Give feedback.
-
|
@iziemba comment suggests that the enum extension suggested above would work for both the IPC and GPUDirect RDMA cases. Or we need to define more flags. Enum is looking like a better approach, and might be easier for the user. |
Beta Was this translation helpful? Give feedback.
-
|
As we add interfaces and types of memory, there are going to be a lot of combinations and I think it might make it hard to code for specific protocols. For example in shm to figure out if I can do IPC then I have to do that can apply the same to all interfaces. |
Beta Was this translation helpful? Give feedback.
-
This is exactly how I view the various GPU allocation functions. WRT to AMD, the HIP API (ROCM), which can either layer over CUDA or AMD HSA (ROCR), also has these three types of calls: hipMalloc(), hipMallocHost(), and hipMallocManaged(). But since libfabric supports ROCR and not explicitly ROCM, I need to spend some time understanding how hipMallocManaged() translates to ROCR. |
Beta Was this translation helpful? Give feedback.
-
|
I have verified that it is possible to identify these three types using ROCR. |
Beta Was this translation helpful? Give feedback.
-
|
Why would anyone use a device API to allocate host memory? I don't understand why those calls exist at all. Looking at the CUDA documentation, there's a ridiculous number of allocation calls. Sigh... What is 'managed' memory actually mean? That the allocation uses a shared virtual address space? And are we getting into an area where a restriction today may not be a restriction tomorrow? |
Beta Was this translation helpful? Give feedback.
-
Using a device API to allocate host memory can lead to improved host <-> GPU memcpy latency. It has been a while since I have looked into it, but one optimization this device API has is the pining of registering of host pages with the GPU device. But, the same thing can be achieved using standard host memory allocators and registering the host buffer with the GPU. This is what RXM does today with its bounce buffers.
Managed memory, or unified memory in CUDA terms, is an allocation where the pages can be faulted and migrated between the CPU and GPU. https://developer.nvidia.com/blog/unified-memory-cuda-beginners/ |
Beta Was this translation helpful? Give feedback.
-
|
Thanks. This is helpful. With the kernel hmem driver, malloc supports page migration between the host and GPU. If we can get this down to 3 flags, that's not too bad. But the flags appear to be exclusive. Is it enough to indicate that the memory has FI_HMEM_HOST_ACCESS and/or FI_HMEM_DEVICE_ACCESS? The absence of any flag would indicate the caller has no idea, I guess. |
Beta Was this translation helpful? Give feedback.
-
This is a good point. Using FI_HMEM_HOST_ACCESS and FI_HMEM_DEVICE_ACCESS, how would a user distinguish between a malloc allocation which migrates between host and GPU and a GPU managed allocation? |
Beta Was this translation helpful? Give feedback.
-
|
If it can migrate between them then it can be accessed by both, right? So |
Beta Was this translation helpful? Give feedback.
-
|
The hmem driver allows migration without the app even being aware that it's a thing. |
Beta Was this translation helpful? Give feedback.
-
|
Ahh shoot. Is there a way to detect if this is enabled/happening? Does it always migrate when necessary or is there no guarantee? |
Beta Was this translation helpful? Give feedback.
-
|
Trying to decipher back through the thread... IPC is supported on device-only memory, correct? And GPUDirect is also only supported on device-only memory, correct? Host-memory buffers may or may not have the ability to migrate to/from the device. This can occur without the application even being aware that migration is possible. DMAbuf support enables peer-to-peer transfers, similar to GPUDirect, but using an upstream mechanism. DMAbuf requires on-demand paging support from the NIC, and would work with page migration. Even if MOFED today has a limitation on peer-to-peer, I don't think we can assume that restriction will hold. And I'm not confident that the restrictions in place on CUDA, Ze, RoCR will hold either. |
Beta Was this translation helpful? Give feedback.
-
|
Basically, I'm wondering if we can get away with only having this flag |
Beta Was this translation helpful? Give feedback.
-
|
For my purposes, yes, but based on what @iziemba I don't think it would cover his case since his restriction is just not managed memory. So it sounds like explicitly host memory would still be ok. |
Beta Was this translation helpful? Give feedback.
-
|
PCI peer-to-peer transfers between the GPU and NIC either need non-migrating device regions or the NIC capable of handling the migration. The latter today is handled by the hmem driver + RDMA on-demand paging, though I have no idea if Nvidia SW supports this yet. The former is indicated by setting the FI_HMEM_DEVICE_ONLY flag. AFAIK, no special handling is needed for accessing host regions, even if allocated using a proprietary device API. I'm assuming that the device API still hooks into the standard kernel page table mechanisms. This should be a normal memory registration call to the NIC. There is a problem that migrating ('managed') regions may not work, depending on what lower layer software is used, and possibly what version or drivers are used. |
Beta Was this translation helpful? Give feedback.
-
|
Ignoring hmem driver + RDMA on-demand paging, I think what we have found is host memory allocated with a GPU allocator supports RDMA, IPC/GPUDirect only supports device memory, and IPC/GPUDirect cannot support managed memory. Could the man pages be updated to say FI_HMEM_CUDA, and others, should only be used for device memory, use FI_HMEM_SYSTEM for host memory allocated with GPU allocators, and managed memory is not supported?
I agree. I'm not sure if defining FI_HMEM_DEVICE_ONLY or updating documentation adds restrictions. It seems like if a device/provider supports hmem driver + RDMA on-demand paging, this would be a new capability bit. |
Beta Was this translation helpful? Give feedback.
-
|
Based on OFIWG discussions, the FI_HMEM_DEVICE_ONLY flag passed into the memory registration call should work. When given, the provider can optimize transfers. The shm provider can use IPC for transfers, and verbs can setup P2P RDMA. If not given, the provider would need to query the hmem APIs to discover this data. |
Beta Was this translation helpful? Give feedback.
-
To extend for use with FI_MR_HMEM, fi_mr_attr was extended to include the interface used to allocate and manage the memory (ie FI_HMEM_CUDA, FI_HMEM_ZE) as well as interface-specific fields.
Each of these interface types is defined to include memory allocated in a variety of different ways. For example, for CUDA:
"Uses Nvidia CUDA interfaces such as cuMemAlloc, cuMemAllocHost, cuMemAllocManaged, cuMemFree, cudaMalloc, cudaFree."
However for both CUDA and L0 (maybe others), IPC is not supported for all types of allocations. See here for CUDA and here for ZE.
fi_mr_attr needs to include what type of memory/allocation the buffer is as well so that the provider can know what protocols, like IPC, can or cannot be used.
Possible solution - add allocation types to fi_hmem_iface
Beta Was this translation helpful? Give feedback.
All reactions