DEVICE: Improve error logs in GPU and host device APIs #10921
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
michal-shalev
force-pushed
the
add-error-logs
branch
9 times, most recently
from
0deab63 to
341b9c3
Compare
michal-shalev
force-pushed
the
add-error-logs
branch
from
341b9c3 to
325142a
Compare
brminich
reviewed
Oct 1, 2025
michal-shalev
changed the title
tvegas1
previously approved these changes
Oct 6, 2025
michal-shalev
force-pushed
the
add-error-logs
branch
from
59e1cb7 to
e00e62a
Compare
iyastreb
reviewed
Oct 9, 2025
| UCS_F_DEVICE const char* ucs_device_status_string(ucs_status_t status) | ||
| { | ||
| switch (status) { | ||
| UCS_STATUS_STRING_CASES |
There was a problem hiding this comment.
IMO X macro is an ideal candidate here, because it's much more flexible.
Also we keep all related things in one place, so it's less error prone
#define UCS_FOREACH_STATUS(_macro) \
_macro(UCS_OK, 0, "Success") \
_macro(UCS_INPROGRESS, 1, "Operation in progress") \
...
#define UCS_STATUS_ENUMIFY(ID, VALUE, _) ID = VALUE,
typedef enum {
UCS_FOREACH_STATUS(UCS_STATUS_ENUMIFY)
} ucs_status_t;
There was a problem hiding this comment.
Thank you, it's a great idea, I pushed a commit that changes the implementation to use X macro.
iyastreb
reviewed
Oct 9, 2025
src/ucs/sys/device_code.h
Outdated
| #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, \ |
There was a problem hiding this comment.
Just some ideas for improvements:
- Now messages are not aligned, because thread id vary from 0 to 1023. Maybe we can use a fixed width for thread:block part?
- I think we don't need a function name - it would make the output too long. Also not sure if it really helps in case of inline functions - did you check?
- Maybe we could also fix max width for file:line, so than the output looks like:
(0:0) DEBUG filename_long.cu:1234 Message 1
(255:0) ERROR short.cu:9 Message 2
IMO it's important that message starts at the same length, otherwise it's hard to read
There was a problem hiding this comment.
- Fixed with Thomas' suggestion.
- It works in the case of the inline functions, I tested it. IMO function name is very helpful for debug.
- Yes, I changed filename max width to 40 character, line number to max 4 digits and function name max to 30 characters.
This is how it looks like now:
( 0: 0) ERROR /hpc/newhome/mshalev/workspace/gdaki/ucx:255 uct_rc_mlx5_gda_ep_single : This is an example error message from lane 0
( 0: 0) DEBUG /hpc/newhome/mshalev/workspace/gdaki/ucx:253 uct_rc_mlx5_gda_ep_single : This is an example debug message addr=0x7fdeacc00200 remote=0x7fdeacc44400 len=8
There was a problem hiding this comment.
Thanks, much better now
- But still I'm against the function name, because:
- it makes output too long
- it's redundant since file:line is present
- it's absent in UCX logs (I think it would be nice to have a similar layout)
UCX logs
[1760011677.835218] [rock14:3182762:0] debug.c:1157 UCX DEBUG using signal stack 0x7f1a76826000 size 141824
[1760011677.848311] [rock14:3182762:0] cpu.c:338 UCX DEBUG measured tsc frequency 2595.125 MHz after 0.08 ms
[1760011677.848337] [rock14:3182762:0] init.c:120 UCX DEBUG /labhome/iyastrebov/ws/ucx2/bld-devel/lib/libucs.so.0 loaded at 0x7f1a75d50000
It looks weird if we have different log formats within the same library
- Can we print just the basename of the file? Long log lines are hard to read
- Maybe align
(0:1)on the right or left?
iyastreb
reviewed
Oct 9, 2025
| _macro(UCS_ERR_CONNECTION_RESET, -25, "Connection reset by remote peer") \ | ||
| _macro(UCS_ERR_FIRST_LINK_FAILURE, -40, "First link failure") \ | ||
| _macro(UCS_ERR_LAST_LINK_FAILURE, -59, "Last link failure") \ | ||
| _macro(UCS_ERR_FIRST_ENDPOINT_FAILURE,-60, "First endpoint failure") \ |
tvegas1
reviewed
Oct 13, 2025
| uct_allocated_memory_t mem; | ||
|
|
||
| if (!(ep->flags & UCP_EP_FLAG_REMOTE_CONNECTED)) { | ||
| ucs_error("ep=%p didn't complete wireup", ep); |
There was a problem hiding this comment.
do we need to remove that one eventually?
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
What?
Improve error logging to UCP and UCT device APIs to improve debugging.
Why?
The device API functions lacked proper error reporting, making it difficult to debug failures in GPU kernels.
Additionally, some UCP host functions were returning error codes without logging what went wrong,
How?
ucs_device_printfmacro to automatically include file, line, and function informationucs_device_status_string()function for device-compatible status code to string conversionucp_device_impl.h,uct_device_impl.h,ucp_device.candgdaki.cuh