Make sure that CUDA instances are destroyed prior to destroying streams #1050
Conversation
aprokop
added
bug
benchmarks/execution_space_instances/execution_space_instances_driver.cpp
Outdated
Show resolved
Hide resolved
aprokop
force-pushed
the
fix_stream_instance_destroy_order
branch
from
47759bf
to
52f352d
Compare
aprokop
force-pushed
the
fix_stream_instance_destroy_order
branch
from
52f352d
to
b19cb1c
Compare
| // Place the code using CUDA instance in a block to make sure the instance | ||
| // is destroyed prior to destroying the stream. |
There was a problem hiding this comment.
What about
Kokkos::push_finalize_hook([&stream](){ cudaStreamDestroy(stream); });instead?
There was a problem hiding this comment.
Would it guarantee the destroy order with the instance, though?
There was a problem hiding this comment.
The instance would get out of scope and destroyed first, then the kokkos execution environment scope guard destructor kicks in, calls Kokkos::finalize which invokes all registered callbacks before releasing resources.
There was a problem hiding this comment.
That is "yes it would do the right thing".
I prefer that solution but I am not blocking.
There was a problem hiding this comment.
Alternatively, you could just move Kokkos::ScopeGuard guard(argc, argv); inside this block so that the stream is created/destroyed before/after Kokkos is initialized/finalized.
There was a problem hiding this comment.
I don't think it matters whether the stream is destroyed before/after Kokkos is finalized. As long as the instance is destroyed before, that's the only thing that matters, if I understand it right. So, I would prefer to keep the current code.
There was a problem hiding this comment.
The change I suggested is more explicit about what needs to happen, hence more readable IMO.
Co-authored-by: Damien L-G <dalg24@gmail.com>
aprokop
force-pushed
the
fix_stream_instance_destroy_order
branch
from
2cd4715
to
d236452
Compare
|
HIP failed, but this PR does not include #1048. |
|
This indeed fixed CUDA nightlies. |
Fix #1049. Should fix nightly failures.