Merge pull request #6452 from G-071/fix_nvcc_crashes

Fix nvcc crashes in transform_stream.cu and synchronize.cu
STEllAR-GROUP · Apr 3, 2024 · c516c88 · c516c88
2 parents fe66c1b + 3c88a73
commit c516c88
Show file tree

Hide file tree

Showing 3 changed files with 13 additions and 6 deletions.
diff --git a/libs/core/async_cuda/tests/performance/synchronize.cu b/libs/core/async_cuda/tests/performance/synchronize.cu
@@ -6,8 +6,9 @@
 
 #include <hpx/config.hpp>
 
-// NVCC fails unceremoniously with this test at least until V12.1
-#if !defined(HPX_CUDA_VERSION) || (HPX_CUDA_VERSION > 1201)
+// NVCC fails unceremoniously with this test until V12.2
+// Fixed in CUDA 12.3
+#if !defined(HPX_CUDA_VERSION) || (HPX_CUDA_VERSION > 1202)
 
 #include <hpx/chrono.hpp>
 #include <hpx/execution.hpp>

diff --git a/libs/core/async_cuda/tests/unit/transform_stream.cu b/libs/core/async_cuda/tests/unit/transform_stream.cu
@@ -6,8 +6,9 @@
 
 #include <hpx/config.hpp>
 
-// NVCC fails unceremoniously with this test at least until V12.1
-#if !defined(HPX_CUDA_VERSION) || (HPX_CUDA_VERSION > 1201)
+// NVCC fails unceremoniously with this test until 12.2
+// Fixed in CUDA 12.3
+#if !defined(HPX_CUDA_VERSION) || (HPX_CUDA_VERSION > 1202)
 
 #include <hpx/execution.hpp>
 #include <hpx/init.hpp>
@@ -16,6 +17,7 @@
 
 #include <atomic>
 #include <cstddef>
+#include <iostream>
 #include <utility>
 
 __global__ void dummy_kernel() {}
@@ -292,6 +294,7 @@ int hpx_main()
         cu::check_cuda_error(cudaFree(p));
     }
 
+    std::cout << "Test finished!" << std::endl;
     return hpx::local::finalize();
 }
 

diff --git a/libs/core/execution/include/hpx/execution/algorithms/just.hpp b/libs/core/execution/include/hpx/execution/algorithms/just.hpp
@@ -102,8 +102,11 @@ namespace hpx::execution::experimental {
 
             template <typename Receiver>
             friend auto tag_invoke(
-                connect_t, just_sender&& s, Receiver&& receiver) noexcept(util::
-                    all_of_v<std::is_nothrow_move_constructible<Ts>...>)
+                connect_t, just_sender&& s, Receiver&& receiver)
+#if !defined(HPX_COMPUTE_DEVICE_CODE)
+                noexcept(
+                    util::all_of_v<std::is_nothrow_move_constructible<Ts>...>)
+#endif
             {
                 return operation_state<Receiver>{
                     HPX_FORWARD(Receiver, receiver), HPX_MOVE(s.ts)};