-
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
NVCC bug with __device__ on defaulted function #1470
Labels
Bug
Broken / incorrect code; it could be Kokkos' responsibility, or others’ (e.g., Trilinos)
Compiler Issue
An issue that Kokkos cannot / should not fix; Kokkos must communicate to relevant vendor
Milestone
Comments
ibaned
added
Bug
Broken / incorrect code; it could be Kokkos' responsibility, or others’ (e.g., Trilinos)
Compiler Issue
An issue that Kokkos cannot / should not fix; Kokkos must communicate to relevant vendor
labels
Mar 12, 2018
ibaned
added a commit
that referenced
this issue
Mar 12, 2018
#1471 is the PR, and we should consider patching Trilinos as well (currently Tacho just segfaults). |
Filed NVIDIA bug 2083242 for this contradictory behavior. |
Based on discussion in #1473 , we've found that this warning can be suppressed:
|
ndellingwood
added a commit
to ndellingwood/kokkos-kernels
that referenced
this issue
Mar 16, 2018
Required following changes from issue kokkos/kokkos#1470 Addresses issue kokkos#186
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Labels
Bug
Broken / incorrect code; it could be Kokkos' responsibility, or others’ (e.g., Trilinos)
Compiler Issue
An issue that Kokkos cannot / should not fix; Kokkos must communicate to relevant vendor
So it turns out that having
KOKKOS_INLINE_FUNCTION_DEFAULTED=inline
for GCC+NVCC isn't right. Even though the NVCC compiler spits out tons of warnings like this:It turns out that removing
__host__ __device__
gives incorrect behavior. In particular, NVCC will choose to call the__host__
version of a defaulted copy constructor even if the call is inside device code. It emits no warnings about the declaration but does spit out this warning at the call site:None of our tests were actually testing this case, but Tacho was (by having a
MemoryPool
be a member of a task).By adding a test that did have a
MemoryPool
as a member of a task, I replicated the failure (the test segfaults).MemoryPool
's default copy constructor usesKOKKOS_INLINE_FUNCTION_DEFAULTED
. If one changes that toKOKKOS_INLINE_FUNCTION
, many warnings of the first kind ensue, but the code runs without segfaulting.I consider this a very frustrating bug on the part of NVIDIA, and can see no better option than to simply stop using defaulted methods altogether.
I'm putting together a PR with the reproducer test as one commit and another commit that removes defaulted functions.
@kyungjoo-kim @crtrott @ndellingwood
The text was updated successfully, but these errors were encountered: