Merge pull request #12229 from edgargabriel/topic/accelerator-size-0-fix

opal/accelerator: allow 0 size copies

Author: edgargabriel

opal/mca/accelerator/cuda/accelerator_cuda.c

@@ -399,9 +399,12 @@ static int accelerator_cuda_memcpy_async(int dest_dev_id, int src_dev_id, void *
     }
 
     if ((MCA_ACCELERATOR_STREAM_DEFAULT != stream && NULL == stream) ||
-        NULL == dest || NULL == src || size <= 0) {
+        NULL == dest || NULL == src || size < 0) {
         return OPAL_ERR_BAD_PARAM;
     }
+    if (0 == size) {
+        return OPAL_SUCCESS;
+    }
 
     result = cuMemcpyAsync((CUdeviceptr) dest, (CUdeviceptr) src, size, GET_STREAM(stream));
     if (OPAL_UNLIKELY(CUDA_SUCCESS != result)) {
@@ -422,9 +425,12 @@ static int accelerator_cuda_memcpy(int dest_dev_id, int src_dev_id, void *dest,
         return delayed_init;
     }
 
-    if (NULL == dest || NULL == src || size <= 0) {
+    if (NULL == dest || NULL == src || size < 0) {
         return OPAL_ERR_BAD_PARAM;
     }
+    if (0 == size) {
+        return OPAL_SUCCESS;
+    }
 
     /* Async copy then synchronize is the default behavior as some applications
      * cannot utilize synchronous copies. In addition, host memory does not need

opal/mca/accelerator/rocm/accelerator_rocm_module.c

@@ -303,9 +303,12 @@ static int mca_accelerator_rocm_memcpy_async(int dest_dev_id, int src_dev_id, vo
 {
     if ((MCA_ACCELERATOR_STREAM_DEFAULT != stream &&
         (NULL == stream || NULL == stream->stream)) ||
-        NULL == src || NULL == dest || size <= 0) {
+        NULL == src || NULL == dest || size < 0) {
         return OPAL_ERR_BAD_PARAM;
     }
+    if (0 == size) {
+        return OPAL_SUCCESS;
+    }
 
     hipError_t err = hipMemcpyAsync(dest, src, size, hipMemcpyDefault,
 				    GET_STREAM(stream));
@@ -324,9 +327,12 @@ static int mca_accelerator_rocm_memcpy(int dest_dev_id, int src_dev_id, void *de
 {
     hipError_t err;
 
-    if (NULL == src || NULL == dest || size <=0) {
+    if (NULL == src || NULL == dest || size < 0) {
         return OPAL_ERR_BAD_PARAM;
     }
+    if (0 == size) {
+        return OPAL_SUCCESS;
+    }
 
     if (type == MCA_ACCELERATOR_TRANSFER_DTOH && size <= opal_accelerator_rocm_memcpyD2H_limit) {
         memcpy(dest, src, size);

opal/mca/accelerator/ze/accelerator_ze_module.c

@@ -403,11 +403,14 @@ static int mca_accelerator_ze_memcpy_async(int dest_dev_id, int src_dev_id, void
    opal_accelerator_ze_stream_t *ze_stream = NULL;
 
    if (NULL == stream || NULL == src ||
-        NULL == dest   || size <= 0) {
+        NULL == dest   || size < 0) {
         return OPAL_ERR_BAD_PARAM;
     }
+   if (0 == size) {
+       return OPAL_SUCCESS;
+   }
 
-    ze_stream = (opal_accelerator_ze_stream_t  *)stream->stream;
+   ze_stream = (opal_accelerator_ze_stream_t  *)stream->stream;
     assert(NULL != ze_stream);
 
     zret = zeCommandListAppendMemoryCopy(ze_stream->hCommandList,
@@ -435,9 +438,12 @@ static int mca_accelerator_ze_memcpy(int dest_dev_id, int src_dev_id, void *dest
 
     opal_accelerator_ze_stream_t *ze_stream = NULL;
 
-    if (NULL == src || NULL == dest || size <=0) {
+    if (NULL == src || NULL == dest || size <0) {
         return OPAL_ERR_BAD_PARAM;
     }                           
+    if (0 == size) {
+        return OPAL_SUCCESS;
+    }
 
     if (MCA_ACCELERATOR_NO_DEVICE_ID == src_dev_id) {
         dev_id = 0;