From ef629cef0671dd42b9d9b5766e127b95b65d068c Mon Sep 17 00:00:00 2001 From: Mickael Ide Date: Wed, 11 Sep 2024 14:24:00 +0200 Subject: [PATCH 1/3] Use runtime check of cudart version for eig --- cpp/include/raft/linalg/detail/eig.cuh | 33 ++++++++++++++------------ 1 file changed, 18 insertions(+), 15 deletions(-) diff --git a/cpp/include/raft/linalg/detail/eig.cuh b/cpp/include/raft/linalg/detail/eig.cuh index ba7ed3dcdf..c342d5f4e4 100644 --- a/cpp/include/raft/linalg/detail/eig.cuh +++ b/cpp/include/raft/linalg/detail/eig.cuh @@ -95,16 +95,19 @@ void eigDC(raft::resources const& handle, return; #endif -#if CUDART_VERSION <= 12040 - // Use a new stream instead of `cudaStreamPerThread` to avoid cusolver bug # 4580093. - rmm::cuda_stream stream_new_wrapper; - cudaStream_t stream_new = stream_new_wrapper.value(); - cudaEvent_t sync_event = resource::detail::get_cuda_stream_sync_event(handle); - RAFT_CUDA_TRY(cudaEventRecord(sync_event, stream)); - RAFT_CUDA_TRY(cudaStreamWaitEvent(stream_new, sync_event)); -#else - cudaStream_t stream_new = stream; -#endif + int cudart_version = 0; + RAFT_CUDA_TRY(cudaRuntimeGetVersion(&cudart_version)); + cudaStream_t stream_new; + if (cudart_version < 12050) { + // Use a new stream instead of `cudaStreamPerThread` to avoid cusolver bug # 4580093. + rmm::cuda_stream stream_new_wrapper; + stream_new = stream_new_wrapper.value(); + cudaEvent_t sync_event = resource::detail::get_cuda_stream_sync_event(handle); + RAFT_CUDA_TRY(cudaEventRecord(sync_event, stream)); + RAFT_CUDA_TRY(cudaStreamWaitEvent(stream_new, sync_event)); + } else { + stream_new = stream; + } cusolverDnHandle_t cusolverH = resource::get_cusolver_dn_handle(handle); cusolverDnParams_t dn_params = nullptr; @@ -152,11 +155,11 @@ void eigDC(raft::resources const& handle, "eig.cuh: eigensolver couldn't converge to a solution. " "This usually occurs when some of the features do not vary enough."); -#if CUDART_VERSION <= 12040 - // Synchronize the created stream with the original stream before return - RAFT_CUDA_TRY(cudaEventRecord(sync_event, stream_new)); - RAFT_CUDA_TRY(cudaStreamWaitEvent(stream, sync_event)); -#endif + if (cudart_version < 12050) { + // Synchronize the created stream with the original stream before return + RAFT_CUDA_TRY(cudaEventRecord(sync_event, stream_new)); + RAFT_CUDA_TRY(cudaStreamWaitEvent(stream, sync_event)); + } } enum EigVecMemUsage { OVERWRITE_INPUT, COPY_INPUT }; From e76eb18b631d49ea2763e7fca58a4108eb30ed9a Mon Sep 17 00:00:00 2001 From: Mickael Ide Date: Wed, 11 Sep 2024 18:10:33 +0200 Subject: [PATCH 2/3] Fix compilation --- cpp/include/raft/linalg/detail/eig.cuh | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/cpp/include/raft/linalg/detail/eig.cuh b/cpp/include/raft/linalg/detail/eig.cuh index c342d5f4e4..0da548a5ac 100644 --- a/cpp/include/raft/linalg/detail/eig.cuh +++ b/cpp/include/raft/linalg/detail/eig.cuh @@ -98,11 +98,12 @@ void eigDC(raft::resources const& handle, int cudart_version = 0; RAFT_CUDA_TRY(cudaRuntimeGetVersion(&cudart_version)); cudaStream_t stream_new; + cudaEvent_t sync_event; + rmm::cuda_stream stream_new_wrapper; if (cudart_version < 12050) { // Use a new stream instead of `cudaStreamPerThread` to avoid cusolver bug # 4580093. - rmm::cuda_stream stream_new_wrapper; - stream_new = stream_new_wrapper.value(); - cudaEvent_t sync_event = resource::detail::get_cuda_stream_sync_event(handle); + stream_new = stream_new_wrapper.value(); + sync_event = resource::detail::get_cuda_stream_sync_event(handle); RAFT_CUDA_TRY(cudaEventRecord(sync_event, stream)); RAFT_CUDA_TRY(cudaStreamWaitEvent(stream_new, sync_event)); } else { From 64a9e2ef5206b919bf3532d37be238fd0d6b7ea6 Mon Sep 17 00:00:00 2001 From: Mickael Ide Date: Wed, 11 Sep 2024 15:17:54 -0700 Subject: [PATCH 3/3] Fix uninitialized cuda event --- cpp/include/raft/linalg/detail/eig.cuh | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/cpp/include/raft/linalg/detail/eig.cuh b/cpp/include/raft/linalg/detail/eig.cuh index 0da548a5ac..561187178c 100644 --- a/cpp/include/raft/linalg/detail/eig.cuh +++ b/cpp/include/raft/linalg/detail/eig.cuh @@ -98,12 +98,11 @@ void eigDC(raft::resources const& handle, int cudart_version = 0; RAFT_CUDA_TRY(cudaRuntimeGetVersion(&cudart_version)); cudaStream_t stream_new; - cudaEvent_t sync_event; + cudaEvent_t sync_event = resource::detail::get_cuda_stream_sync_event(handle); rmm::cuda_stream stream_new_wrapper; if (cudart_version < 12050) { // Use a new stream instead of `cudaStreamPerThread` to avoid cusolver bug # 4580093. stream_new = stream_new_wrapper.value(); - sync_event = resource::detail::get_cuda_stream_sync_event(handle); RAFT_CUDA_TRY(cudaEventRecord(sync_event, stream)); RAFT_CUDA_TRY(cudaStreamWaitEvent(stream_new, sync_event)); } else {