Batched device memcpy for CUDA/HIP - #791
Conversation
Since CUDA 12.8 we can use batched memory copy to perform multiple operations in a single CUDA call. For tasks with multiple stage in/out flows this can be beneficial. Prominent example are eviction tasks. There is no reason to order pushouts of a single task. In practice, I am observing higher bandwidth for small batched data (9x2GB/s vs 1x12GB/s for 9 8KB chunks). Signed-off-by: Joseph Schuchart <joseph.schuchart@stonybrook.edu>
There was a problem hiding this comment.
Pull request overview
This PR introduces a new multi-item GPU memcpy interface and updates the default GPU stage-in/stage-out paths to batch multiple transfers into a single backend call when possible (notably leveraging CUDA 12.8’s cudaMemcpyBatchAsync) to improve bandwidth for workloads with multiple in/out flows (e.g., evictions).
Changes:
- Adds a new
memcpy_multi_asyncfunction pointer to the GPU device module interface, with a generic fallback that issues multiplememcpy_asynccalls. - Updates default GPU stage-in/stage-out logic to aggregate per-flow copies and issue one combined multi-copy operation.
- Implements a CUDA 12.8+ native backend using
cudaMemcpyBatchAsync, and wires Level Zero to the generic fallback.
Reviewed changes
Copilot reviewed 4 out of 4 changed files in this pull request and generated 1 comment.
| File | Description |
|---|---|
| parsec/mca/device/level_zero/device_level_zero_module.c | Wires Level Zero GPU device to use the generic multi-memcpy fallback. |
| parsec/mca/device/device_gpu.h | Adds the memcpy_multi_async API + documents semantics and declares the generic fallback. |
| parsec/mca/device/device_gpu.c | Implements the generic fallback and updates default stage-in/stage-out and push/pop paths to batch transfers. |
| parsec/mca/device/cuda/device_cuda_module.c | Adds CUDA 12.8+ native batched memcpy implementation and selects it at init time. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
| #if CUDART_VERSION >= 13000 | ||
| cudaStatus = cudaMemcpyBatchAsync(dsts, (const void *const *)srcs, sizes, (size_t)nb_items, | ||
| &attrs, &attrsIdx, 1 /* numAttrs */, | ||
| cuda_stream->cuda_stream); | ||
| #else | ||
| cudaStatus = cudaMemcpyBatchAsync(dsts, srcs, sizes, (size_t)nb_items, | ||
| &attrs, &attrsIdx, 1 /* numAttrs */, | ||
| NULL /* failIdx: unused, we don't need to know which item failed */, | ||
| cuda_stream->cuda_stream); | ||
| #endif |
bosilca
left a comment
bosilca
left a comment
There was a problem hiding this comment.
This is a good idea, but the scope of the current implementation is narrow and only benefit to tasks with small data and only if the copy latency is important, which means for tasks that cannot keep the GPU busy (aka. not compute intensive). It also work per task, so at best it batches MAX_PARAM items.
If we can merge this with the task batching we could have a much stronger support for small compute-trivial tasks.
| { | ||
| int ret, rc = PARSEC_SUCCESS; | ||
|
|
||
| for(int i = 0; i < nb_items; i++) { |
There was a problem hiding this comment.
This loop is too carefree. I know that if any error is returned everything will fail, but if we try to expose the faults at least we should do it properly. Here if anything fails we return the error but we continue to submit all the others copies. However, once we return from this function with an error we call parsec_device_kernel_push_release_readers_on_failure directly, releasing the readers for the entire mask (and that's one problem), and then return without enquing an event (which allow the data use right away on another stream).
The interface needs either atomic submission semantics, a returned submitted mask, or the old per-flow path for backends without native batching.
| parsec_device_module_t *src_dev_mod = parsec_mca_device_get(src->device_index); | ||
| if( (NULL == src_dev_mod) || !PARSEC_DEV_IS_GPU(src_dev_mod->type) ) continue; | ||
| int readers = parsec_gpu_data_copy_release_reader((parsec_device_gpu_module_t*)src_dev_mod, src, 1); | ||
| assert(readers >= 0); |
There was a problem hiding this comment.
releases a GPU reader without locking its original or coordinating with the source GPU manager. When this removes the last reader, parsec_gpu_data_copy_release_reader() can relink the copy into a source-device LRU.
The normal D2D completion path serializes with the source manager, locks the original, and updates data_avail_epoch. The failure path should reuse that protocol.
Since CUDA 12.8 we can use batched memory copy to perform multiple operations in a single CUDA call. For tasks with multiple stage in/out flows this can be beneficial. Prominent example
are eviction tasks. There is no reason to order pushouts of a single task.
In practice, I am observing higher bandwidth for small batched data (9x2GB/s vs 1x12GB/s for 9 8KB chunks).