DEVICE: Improve error logs in GPU and host device APIs

src/ucp/api/device/ucp_device_impl.h

@@ -88,6 +88,10 @@ UCS_F_DEVICE ucs_status_t ucp_device_prepare_send(
 
     if ((mem_list_h->version != UCP_DEVICE_MEM_LIST_VERSION_V1) ||
         (first_mem_elem_index >= mem_list_h->mem_list_length)) {
+        ucs_device_error("invalid parameters: mem_list version=%u (expected %u), "
+                         "first_mem_elem_index=%u, mem_list_length=%u",
+                         mem_list_h->version, UCP_DEVICE_MEM_LIST_VERSION_V1,
+                         first_mem_elem_index, mem_list_h->mem_list_length);
         return UCS_ERR_INVALID_PARAM;
     }
 
@@ -144,6 +148,8 @@ UCS_F_DEVICE ucs_status_t ucp_device_put_single(
     status = ucp_device_prepare_send(mem_list_h, mem_list_index, req, device_ep,
                                      uct_elem, comp);
     if (status != UCS_OK) {
+        ucs_device_error("send prepare failed with %s, mem_list_index=%u",
+                         ucs_device_status_string(status), mem_list_index);
         return status;
     }
 
@@ -195,6 +201,8 @@ UCS_F_DEVICE ucs_status_t ucp_device_counter_inc(
     status = ucp_device_prepare_send(mem_list_h, mem_list_index, req, device_ep,
                                      uct_elem, comp);
     if (status != UCS_OK) {
+        ucs_device_error("send prepare failed with %s, mem_list_index=%u",
+                         ucs_device_status_string(status), mem_list_index);
         return status;
     }
 
@@ -259,6 +267,8 @@ UCS_F_DEVICE ucs_status_t ucp_device_put_multi(
     status = ucp_device_prepare_send(mem_list_h, 0, req, device_ep,
                                      uct_mem_list, comp);
     if (status != UCS_OK) {
+        ucs_device_error("send prepare failed with %s, mem_list_length=%u",
+                         ucs_device_status_string(status), mem_list_h->mem_list_length);
         return status;
     }
 
@@ -334,6 +344,8 @@ UCS_F_DEVICE ucs_status_t ucp_device_put_multi_partial(
     status = ucp_device_prepare_send(mem_list_h, 0, req, device_ep,
                                      uct_mem_list, comp);
     if (status != UCS_OK) {
+        ucs_device_error("send prepare failed with %s, mem_list_count=%u",
+                         ucs_device_status_string(status), mem_list_count);
         return status;
     }
 

src/ucp/core/ucp_device.c

@@ -380,6 +380,7 @@ ucp_device_mem_list_create(ucp_ep_h ep,
     uct_allocated_memory_t mem;
 
     if (!(ep->flags & UCP_EP_FLAG_REMOTE_CONNECTED)) {
+        ucs_error("ep=%p didn't complete wireup", ep);
         return UCS_ERR_NOT_CONNECTED;
     }
 
@@ -388,6 +389,8 @@ ucp_device_mem_list_create(ucp_ep_h ep,
                                               &local_sys_dev, &local_md_map,
                                               &mem_type);
     if (status != UCS_OK) {
+        ucs_error("ep=%p check parameters failed: %s", ep,
+                  ucs_status_string(status));
         return status;
     }
 
@@ -426,6 +429,7 @@ ucp_device_mem_list_create(ucp_ep_h ep,
     /* Track memory allocator for later release */
     status = ucp_device_mem_handle_hash_insert(&mem);
     if (status != UCS_OK) {
+        ucs_error("failed to insert handle: %s", ucs_status_string(status));
         uct_mem_free(&mem);
     } else {
         *handle_p = mem.address;

src/ucs/sys/device_code.h

@@ -8,6 +8,7 @@
 #define UCS_DEVICE_CODE_H
 
 #include <ucs/sys/compiler_def.h>
+#include <ucs/type/status.h>
 #include <stdint.h>
 
 /*
@@ -35,7 +36,7 @@ typedef enum {
 } ucs_device_level_t;
 
 
-static UCS_F_ALWAYS_INLINE const char*
+UCS_F_DEVICE const char*
 ucs_device_level_name(ucs_device_level_t level)
 {
     switch (level) {
@@ -87,10 +88,10 @@ UCS_F_DEVICE void ucs_device_atomic64_write(uint64_t *ptr, uint64_t value)
 
 
 /* Helper macro to print a message from a device function including the
- * thread and block indices */
+ * thread and block indices, file, line, and function */
 #define ucs_device_printf(_title, _fmt, ...) \
-    printf("(%d:%d) %6s " _fmt "\n", threadIdx.x, blockIdx.x, _title, \
-           ##__VA_ARGS__)
+    printf("(%d:%d) %6s %s:%d %s: " _fmt "\n", threadIdx.x, blockIdx.x, _title, \
+           __FILE__, __LINE__, __func__, ##__VA_ARGS__)
 
 /* Print an error message from a device function */
 #define ucs_device_error(_fmt, ...) \
@@ -101,4 +102,79 @@ UCS_F_DEVICE void ucs_device_atomic64_write(uint64_t *ptr, uint64_t value)
 #define ucs_device_debug(_fmt, ...) \
     ucs_device_printf("DEBUG", _fmt, ##__VA_ARGS__)
 
+
+/**
+ * @brief Device compatible status code to string conversion
+ *
+ * This function provides status code to string conversion that can be called
+ * from device code. Returns a short string representation of the status code.
+ *
+ * @param [in] status  Status code to convert
+ *
+ * @return Short string representation of the status code
+ */
+UCS_F_DEVICE const char* ucs_device_status_string(ucs_status_t status)
+{
+    switch (status) {
+    case UCS_OK:
+        return "Success";
+    case UCS_INPROGRESS:
+        return "Operation in progress";
+    case UCS_ERR_NO_MESSAGE:
+        return "No pending message";
+    case UCS_ERR_NO_RESOURCE:
+        return "No resources are available to initiate the operation";
+    case UCS_ERR_IO_ERROR:
+        return "Input/output error";
+    case UCS_ERR_NO_MEMORY:
+        return "Out of memory";
+    case UCS_ERR_INVALID_PARAM:
+        return "Invalid parameter";
+    case UCS_ERR_UNREACHABLE:
+        return "Destination is unreachable";
+    case UCS_ERR_INVALID_ADDR:
+        return "Address not valid";
+    case UCS_ERR_NOT_IMPLEMENTED:
+        return "Function not implemented";
+    case UCS_ERR_MESSAGE_TRUNCATED:
+        return "Message truncated";
+    case UCS_ERR_NO_PROGRESS:
+        return "No progress";
+    case UCS_ERR_BUFFER_TOO_SMALL:
+        return "Provided buffer is too small";
+    case UCS_ERR_NO_ELEM:
+        return "No such element";
+    case UCS_ERR_SOME_CONNECTS_FAILED:
+        return "Failed to connect some of the requested endpoints";
+    case UCS_ERR_NO_DEVICE:
+        return "No such device";
+    case UCS_ERR_BUSY:
+        return "Device is busy";
+    case UCS_ERR_CANCELED:
+        return "Request canceled";
+    case UCS_ERR_SHMEM_SEGMENT:
+        return "Shared memory error";
+    case UCS_ERR_ALREADY_EXISTS:
+        return "Element already exists";
+    case UCS_ERR_OUT_OF_RANGE:
+        return "Index out of range";
+    case UCS_ERR_TIMED_OUT:
+        return "Operation timed out";
+    case UCS_ERR_EXCEEDS_LIMIT:
+        return "User-defined limit was reached";
+    case UCS_ERR_UNSUPPORTED:
+        return "Unsupported operation";
+    case UCS_ERR_REJECTED:
+        return "Operation rejected by remote peer";
+    case UCS_ERR_NOT_CONNECTED:
+        return "Endpoint is not connected";
+    case UCS_ERR_CONNECTION_RESET:
+        return "Connection reset by remote peer";
+    case UCS_ERR_ENDPOINT_TIMEOUT:
+        return "Endpoint timeout";
+    default:
+        return "Unknown error";
+    };
+}
+
 #endif

src/uct/api/device/uct_device_impl.h

@@ -55,6 +55,8 @@ UCS_F_DEVICE ucs_status_t uct_device_ep_put_single(
                                                  comp);
     }
 
+    ucs_device_error("unsupported device_ep->uct_tl_id=%d",
+                     device_ep->uct_tl_id);
     return UCS_ERR_UNSUPPORTED;
 }
 
@@ -96,6 +98,8 @@ UCS_F_DEVICE ucs_status_t uct_device_ep_atomic_add(
                                                  flags, comp);
     }
 
+    ucs_device_error("unsupported device_ep->uct_tl_id=%d",
+                     device_ep->uct_tl_id);
     return UCS_ERR_UNSUPPORTED;
 }
 
@@ -161,6 +165,8 @@ UCS_F_DEVICE ucs_status_t uct_device_ep_put_multi(
                                                 flags, comp);
     }
 
+    ucs_device_error("unsupported device_ep->uct_tl_id=%d",
+                     device_ep->uct_tl_id);
     return UCS_ERR_UNSUPPORTED;
 }
 
@@ -233,6 +239,9 @@ UCS_F_DEVICE ucs_status_t uct_device_ep_put_multi_partial(
                                                         counter_inc_value, counter_remote_address,
                                                         flags, comp);
     }
+
+    ucs_device_error("unsupported device_ep->uct_tl_id=%d",
+                     device_ep->uct_tl_id);
     return UCS_ERR_UNSUPPORTED;
 }
 
@@ -256,6 +265,8 @@ UCS_F_DEVICE ucs_status_t uct_device_ep_progress(uct_device_ep_h device_ep)
         return UCS_OK;
     }
 
+    ucs_device_error("unsupported device_ep->uct_tl_id=%d",
+                     device_ep->uct_tl_id);
     return UCS_ERR_UNSUPPORTED;
 }
 

src/uct/ib/mlx5/gdaki/gdaki.cuh

@@ -336,6 +336,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_put_multi(
 
     if ((level != UCS_DEVICE_LEVEL_THREAD) &&
         (level != UCS_DEVICE_LEVEL_WARP)) {
+        ucs_device_error("unsupported level: %s", ucs_device_level_name(level));
         return UCS_ERR_UNSUPPORTED;
     }
 
@@ -426,6 +427,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_put_multi_partial(
 
     if ((level != UCS_DEVICE_LEVEL_THREAD) &&
         (level != UCS_DEVICE_LEVEL_WARP)) {
+        ucs_device_error("unsupported level: %s", ucs_device_level_name(level));
         return UCS_ERR_UNSUPPORTED;
     }
 
@@ -600,6 +602,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_progress(uct_device_ep_h tl_ep)
     } else if (level == UCS_DEVICE_LEVEL_THREAD) {
         return uct_rc_mlx5_gda_progress_thread(ep);
     } else {
+        ucs_device_error("unsupported level: %s", ucs_device_level_name(level));
         return UCS_ERR_UNSUPPORTED;
     }
 }