diff --git a/src/api/gpu/ucx/nixl_device.cuh b/src/api/gpu/ucx/nixl_device.cuh index 08448fad7..6dd7d8187 100644 --- a/src/api/gpu/ucx/nixl_device.cuh +++ b/src/api/gpu/ucx/nixl_device.cuh @@ -20,6 +20,27 @@ #include #include +/* Helper function to extract the base filename */ +__device__ __forceinline__ static const char* nixl_device_basefile(const char* path) { + const char* base = path; + for (const char* p = path; *p; ++p) { + if (*p == '/') { + base = p + 1; + } + } + return base; +} + +/* Helper macro to print a message from NIXL device function including the + * thread and block indices, file and line */ +#define nixl_device_printf(_log_level, _fmt, ...) \ + printf("%c T%-4d:B%-4d%*s%s:%d] " _fmt "\n", _log_level[0], threadIdx.x, blockIdx.x, \ + 17, "", nixl_device_basefile(__FILE__), __LINE__, ##__VA_ARGS__) + +/* Print an error message from NIXL device function */ +#define nixl_device_error(_fmt, ...) \ + nixl_device_printf("ERROR", _fmt, ##__VA_ARGS__) + struct nixlGpuXferStatusH { ucp_device_request_t device_request; }; @@ -67,7 +88,7 @@ nixlGpuConvertUcsStatus(ucs_status_t status) { if (!UCS_STATUS_IS_ERR(status)) { return NIXL_IN_PROG; } - printf("UCX returned error: %d\n", status); + nixl_device_error("UCX backend error: %s", ucs_device_status_string(status)); return NIXL_ERR_BACKEND; } @@ -247,6 +268,7 @@ nixlGpuGetXferStatus(nixlGpuXferStatusH &xfer_status) { case UCS_INPROGRESS: return NIXL_IN_PROG; default: + nixl_device_error("UCX backend error: %s", ucs_device_status_string(status)); return NIXL_ERR_BACKEND; } }