From 9965f6c518956cd88be18de6633b28e2bbaf0ea4 Mon Sep 17 00:00:00 2001
From: Edgar Gabriel <Edgar.Gabriel@amd.com>
Date: Fri, 12 Jan 2024 13:12:47 -0800
Subject: [PATCH] opal/accelerator: allow 0 size copies

it looks like zero size messages with valid buffer pointers can occur
for send-to-self operations and with persistent request. The current
accelerator components return however an error for size 0 making those
tests fail in our testsuite.

This commit allows size 0 messages, and returns immediatly.

Signed-off-by: Edgar Gabriel <Edgar.Gabriel@amd.com>
---
 opal/mca/accelerator/cuda/accelerator_cuda.c        | 10 ++++++++--
 opal/mca/accelerator/rocm/accelerator_rocm_module.c | 10 ++++++++--
 opal/mca/accelerator/ze/accelerator_ze_module.c     | 12 +++++++++---
 3 files changed, 25 insertions(+), 7 deletions(-)

diff --git a/opal/mca/accelerator/cuda/accelerator_cuda.c b/opal/mca/accelerator/cuda/accelerator_cuda.c
index 5c926700a25..e5ec16294a8 100644
--- a/opal/mca/accelerator/cuda/accelerator_cuda.c
+++ b/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
diff --git a/opal/mca/accelerator/rocm/accelerator_rocm_module.c b/opal/mca/accelerator/rocm/accelerator_rocm_module.c
index abb6449ba13..3bee233b004 100644
--- a/opal/mca/accelerator/rocm/accelerator_rocm_module.c
+++ b/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);
diff --git a/opal/mca/accelerator/ze/accelerator_ze_module.c b/opal/mca/accelerator/ze/accelerator_ze_module.c
index 01bc4a00a43..40c13eb8d56 100644
--- a/opal/mca/accelerator/ze/accelerator_ze_module.c
+++ b/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;