-
Notifications
You must be signed in to change notification settings - Fork 443
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.
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
Apply [[noreturn]] to Kokkos::abort when applicable #3106
Conversation
…ntaiton is [[noreturn]]
Can one of the admins verify this patch? |
OK to test |
I would say it does at least for CUDA versions we support: https://www.nvidia.com/en-us/drivers/cuda/418_163/macosx-cuda-418_163-driver/ and |
OK, if people are still using CUDA on Mac, then I'll put it back. |
b037b3a
to
b0d8a00
Compare
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
we use the KOKKOS_IMPL_
prefix in macros where we're not going to guarantee stability for users, which is probably the case here, at least at first (we can always add #define KOKKOS_ABORT_NORETURN KOKKOS_IMPL_ABORT_NORETURN
later if we change our minds)
Otherwise, thanks! Looks good.
Co-authored-by: D. S. Hollman <[email protected]>
Retest this please. |
Retest This Please |
#endif | ||
|
||
#elif defined(KOKKOS_ENABLE_HIP) && defined(__HIP_DEVICE_COMPILE__) | ||
// HIP does not abort |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
This changed recently, see
kokkos/core/src/HIP/Kokkos_HIP_Abort.hpp
Lines 56 to 63 in 89e9d26
__device__ __attribute__((noinline)) void hip_abort(char const *msg) { | |
#ifndef NDEBUG | |
// disable printf on release builds, as it has a non-trivial performance | |
// impact | |
printf("Aborting with message `%s'.\n", msg); | |
#endif | |
abort(); | |
} |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
oh man. I was about to merge it :-(
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I would be fine with addressing that in a follow-up pull request.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I am currently seeing
/tmp/kokkos/core/src/HIP/Kokkos_HIP_Abort.hpp:64:1: warning: function declared 'noreturn' should not return [-Winvalid-noreturn]
if trying to mark the function with [[noreturn]]
so we might just postpone fixing it.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
LGTM
581ba70
to
70b7dcc
Compare
comment that guards added to workaround issues with cuda random num gen with KOKKOS_ENABLE_DEBUG_BOUNDS_CHECK following changes from PR kokkos#3106
I haven't run any tests beyond just building. Pushing to run CI.
I removed the
__APPLE__
check forcuda_abort
since I don't think macOS supports CUDA in recent versions anyway. I can put it back if this configuration is still supported, though.