-
Notifications
You must be signed in to change notification settings - Fork 407
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
Fix compilation error when using HIP with RDC #4553
Conversation
core/src/HIP/Kokkos_HIP_Abort.hpp
Outdated
#ifndef KOKKOS_ENABLE_HIP_RELOCATABLE_DEVICE_CODE | ||
__attribute__((noinline)) | ||
#else | ||
inline | ||
#endif | ||
void | ||
hip_abort(char const *msg) { |
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.
Hmmm... can't you just do
#ifndef KOKKOS_ENABLE_HIP_RELOCATABLE_DEVICE_CODE | |
__attribute__((noinline)) | |
#else | |
inline | |
#endif | |
void | |
hip_abort(char const *msg) { | |
[[noreturn]] __device__ __attribute__((noinline)) inline void hip_abort( | |
char const *msg) { |
Is there a problem declaring the function always inline
?
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 decided not to inline because the compiler was generating poor code if you inline the printf
. I don't know if the problem has been fixed or not.
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.
Sure, I didn't argue about __attribute__((noinline))
but was asking about inline
.
9a854a9
to
f553fb5
Compare
Do you want to add a RDC build on the Jenkins CI? |
Retest this please |
Ideally yes but we need ROCm 4.5 and linking with RDC is still very slow :/ |
With this PR and ROCm 4.5, RDC works.