Skip to content

DEVICE/API: Add NIXL device-side logging infrastructure #886

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 5 commits into
base: main
Choose a base branch
from
Open

DEVICE/API: Add NIXL device-side logging infrastructure #886

Changes from 1 commit
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
18c020f
DEVICE/API: Add NIXL device-side logging infrastructure
michal-shalev Oct 9, 2025
61720cb
DEVICE/API: Align to NIXL log
michal-shalev Oct 20, 2025
a3adf61
Merge branch 'main' into nixl-device-logging
michal-shalev Oct 20, 2025
f44875f
DEVICE/API: call ucs_device_status_string
michal-shalev Oct 21, 2025
7085751
Merge branch 'main' into nixl-device-logging
michal-shalev Oct 21, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
13 changes: 12 additions & 1 deletion src/api/gpu/ucx/nixl_device.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,6 +20,16 @@
#include <nixl_types.h>
#include <ucp/api/device/ucp_device_impl.h>

/* Helper macro to print a message from NIXL device function including the
* thread and block indices, file, line, and function */
#define nixl_device_printf(_title, _fmt, ...) \
Copy link
Contributor

@tvegas1 tvegas1 Oct 9, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

title -> level

Copy link
Contributor Author

@michal-shalev michal-shalev Oct 20, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

renamed to log_level

printf("(%5d:%5d) %5s %s:%d %s: " _fmt "\n", threadIdx.x, blockIdx.x, _title, \
Copy link
Contributor

@rakhmets rakhmets Oct 14, 2025 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Missing a way to disable messages. E.g., using the existing environment variable, or a new one.

Copy link
Contributor

@rakhmets rakhmets Oct 14, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please review this section https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#limitations.
My biggest doubts are related to this:

Note that the buffer is not flushed automatically when the program exits.
The user must call cudaDeviceReset() or cuCtxDestroy() explicitly, as shown in the examples below.

__FILE__, __LINE__, __func__, ##__VA_ARGS__)
Copy link
Contributor

@tvegas1 tvegas1 Oct 9, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

probably __func__ not needed, maybe left pad file and pad lines with %-5d

Copy link
Contributor Author

@michal-shalev michal-shalev Oct 20, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

removed func

I1020 23:39:53.765828 1403216 gpu_xfer_req_h.cpp:84] Created device memory list handle with 1 elements
E T0   :B0                    nixl_device.cuh:88] UCX backend error: Device is busy


/* Print an error message from NIXL device function */
#define nixl_device_error(_fmt, ...) \
nixl_device_printf("ERROR", _fmt, ##__VA_ARGS__)
Comment on lines +41 to +42
Copy link
Contributor

@rakhmets rakhmets Nov 6, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
#define nixl_device_error(_fmt, ...) \
nixl_device_printf("ERROR", _fmt, ##__VA_ARGS__)
#define NIXL_DEVICE_ERROR(_fmt, ...) \
NIXL_DEVICE_PRINTF("ERROR", _fmt, ##__VA_ARGS__)

https://isocpp.github.io/CppCoreGuidelines/CppCoreGuidelines#es32-use-all_caps-for-all-macro-names


struct nixlGpuXferStatusH {
ucp_device_request_t device_request;
};
Expand Down Expand Up @@ -66,7 +76,7 @@ nixlGpuConvertUcsStatus(ucs_status_t status) {
if (!UCS_STATUS_IS_ERR(status)) {
return NIXL_SUCCESS;
}
printf("UCX returned error: %d\n", status);
nixl_device_error("UCX backend error");
return NIXL_ERR_BACKEND;
}

Expand Down Expand Up @@ -242,6 +252,7 @@ nixlGpuGetXferStatus(nixlGpuXferStatusH &xfer_status) {
case UCS_INPROGRESS:
return NIXL_IN_PROG;
default:
nixl_device_error("UCX backend error");
return NIXL_ERR_BACKEND;
}
}
Expand Down