From caf42e7bef6d4cdc99ad67dd77d22432ad0c9539 Mon Sep 17 00:00:00 2001 From: Andrei Alexandrescu Date: Fri, 17 Jul 2026 23:06:23 -0400 Subject: [PATCH 1/5] =?UTF-8?q?cudax/stf:=20cuda=5Ftry=20migration=20?= =?UTF-8?q?=E2=80=94=20stream=5Ftask=20(PR9)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Migrate throwing-capable stream task setup to cuda_try while retaining cuda_safe_call in noexcept scope guards. Make timing-event ownership exception-safe and use the capture-aware stream-device helper so partial failures restore device state and release events. --- .../experimental/__stf/stream/stream_task.cuh | 120 ++++++++++-------- 1 file changed, 67 insertions(+), 53 deletions(-) diff --git a/cudax/include/cuda/experimental/__stf/stream/stream_task.cuh b/cudax/include/cuda/experimental/__stf/stream/stream_task.cuh index 0b6623f7636..88f9af5e326 100644 --- a/cudax/include/cuda/experimental/__stf/stream/stream_task.cuh +++ b/cudax/include/cuda/experimental/__stf/stream/stream_task.cuh @@ -291,23 +291,29 @@ public: auto& dot = ctx.get_dot(); - bool record_time = reserved::dot::instance().is_timing(); + const bool record_time = reserved::dot::instance().is_timing(); - cudaEvent_t start_event, end_event; + cudaEvent_t start_event = nullptr; + cudaEvent_t end_event = nullptr; + bool timing_active = false; - if (record_time) + SCOPE(exit) { - // Events must be created here to avoid issues with multi-gpu - cuda_safe_call(cudaEventCreate(&start_event)); - cuda_safe_call(cudaEventCreate(&end_event)); - cuda_safe_call(cudaEventRecord(start_event, get_stream())); - } + if (start_event) + { + cuda_safe_call(cudaEventDestroy(start_event)); + } + if (end_event) + { + cuda_safe_call(cudaEventDestroy(end_event)); + } + }; SCOPE(exit) { end_uncleared(); - if (record_time) + if (timing_active) { cuda_safe_call(cudaEventRecord(end_event, get_stream())); cuda_safe_call(cudaEventSynchronize(end_event)); @@ -324,6 +330,17 @@ public: clear(); }; + if (record_time) + { + // Events must be created here to avoid issues with multi-gpu. + // cudaEventCreate is an overload set, so use the non-overloaded + // cudaEventCreateWithFlags with the default flags. + start_event = cuda_try(cudaEventDefault); + end_event = cuda_try(cudaEventDefault); + cuda_try(start_event, get_stream()); + timing_active = true; + } + // Default for the first argument is a `cudaStream_t`. if constexpr (::std::is_invocable_v) { @@ -399,45 +416,25 @@ private: // record the event, and restore the current device to its original // value. - // TODO leverage dev_id if known ? - - // Find the stream structure in the driver API - CUcontext ctx; - cuda_safe_call(cuStreamGetCtx(CUstream(streams[0].stream), &ctx)); - - // Query the context associated with a stream by using the underlying driver API - cuda_safe_call(cuCtxPushCurrent(ctx)); - const CUdevice s0_dev = cuda_try(); - cuda_safe_call(cuCtxPopCurrent(&ctx)); - - const int current_dev = cuda_try(); - - if (current_dev != s0_dev) - { - cuda_safe_call(cudaSetDevice(s0_dev)); - } - - // Create a dependency between the last stream and the current stream - cudaEvent_t sync_event; - // Disable timing to avoid implicit barriers - cuda_safe_call(cudaEventCreateWithFlags(&sync_event, cudaEventDisableTiming)); - - cuda_safe_call(cudaEventRecord(sync_event, streams[0].stream)); + const int s0_dev = streams[0].dev_id == -1 ? get_device_from_stream(streams[0].stream) : streams[0].dev_id; - // According to documentation "event may be from a different device than stream." - for (size_t i = 0; i < streams.size(); i++) - { - cuda_safe_call(cudaStreamWaitEvent(streams[i].stream, sync_event, 0)); - } + exec_place::device(s0_dev)->*[&] { + // Disable timing to avoid implicit barriers. + const cudaEvent_t sync_event = cuda_try(cudaEventDisableTiming); + SCOPE(exit) + { + // Asynchronously destroy the event to avoid a memory leak. + cuda_safe_call(cudaEventDestroy(sync_event)); + }; - // Asynchronously destroy event to avoid a memleak - cuda_safe_call(cudaEventDestroy(sync_event)); + cuda_try(sync_event, streams[0].stream); - if (current_dev != s0_dev) - { - // Restore current device - cuda_safe_call(cudaSetDevice(current_dev)); - } + // According to documentation, the event may be from a different device than the stream. + for (const auto& stream : streams) + { + cuda_try(stream.stream, sync_event, 0); + } + }; } bool automatic_stream = true; // `true` if the stream is automatically fetched from the internal pool @@ -539,7 +536,9 @@ public: auto& dot = ctx.get_dot(); auto& statistics = reserved::task_statistics::instance(); - cudaEvent_t start_event, end_event; + cudaEvent_t start_event = nullptr; + cudaEvent_t end_event = nullptr; + bool timing_active = false; bool record_time = schedule_task(); @@ -551,19 +550,23 @@ public: nvtx_range nr(get_symbol().c_str()); start(); - if (record_time) + SCOPE(exit) { - // Events must be created here to avoid issues with multi-gpu - cuda_safe_call(cudaEventCreate(&start_event)); - cuda_safe_call(cudaEventCreate(&end_event)); - cuda_safe_call(cudaEventRecord(start_event, get_stream())); - } + if (start_event) + { + cuda_safe_call(cudaEventDestroy(start_event)); + } + if (end_event) + { + cuda_safe_call(cudaEventDestroy(end_event)); + } + }; SCOPE(exit) { end_uncleared(); - if (record_time) + if (timing_active) { cuda_safe_call(cudaEventRecord(end_event, get_stream())); cuda_safe_call(cudaEventSynchronize(end_event)); @@ -585,6 +588,17 @@ public: clear(); }; + if (record_time) + { + // Events must be created here to avoid issues with multi-gpu. + // cudaEventCreate is an overload set, so use the non-overloaded + // cudaEventCreateWithFlags with the default flags. + start_event = cuda_try(cudaEventDefault); + end_event = cuda_try(cudaEventDefault); + cuda_try(start_event, get_stream()); + timing_active = true; + } + if constexpr (::std::is_invocable_v) { // Invoke passing this task's stream as the first argument, followed by the slices From 19bb001304301a226989351906e2f6c7aa0a6ddd Mon Sep 17 00:00:00 2001 From: Andrei Alexandrescu Date: Fri, 17 Jul 2026 23:46:49 -0400 Subject: [PATCH 2/5] cudax/stf: make task timing cleanup exception-safe Destroy timing events reliably and separate successful timing collection from failure cleanup across stream and launch tasks. --- .../experimental/__stf/graph/graph_task.cuh | 14 +--- .../__stf/internal/host_launch_scope.cuh | 31 +++++--- .../experimental/__stf/internal/launch.cuh | 62 ++++++++++------ .../experimental/__stf/stream/stream_task.cuh | 70 ++++++++++--------- 4 files changed, 101 insertions(+), 76 deletions(-) diff --git a/cudax/include/cuda/experimental/__stf/graph/graph_task.cuh b/cudax/include/cuda/experimental/__stf/graph/graph_task.cuh index fe4b2c15735..085911b7e7c 100644 --- a/cudax/include/cuda/experimental/__stf/graph/graph_task.cuh +++ b/cudax/include/cuda/experimental/__stf/graph/graph_task.cuh @@ -355,12 +355,7 @@ public: // cudaEvent_t start_event, end_event; - bool record_time = schedule_task(); - - if (statistics.is_calibrating_to_file()) - { - record_time = true; - } + const bool record_time = schedule_task() || statistics.is_calibrating_to_file(); start(); @@ -702,12 +697,7 @@ public: // cudaEvent_t start_event, end_event; - bool record_time = schedule_task(); - - if (statistics.is_calibrating_to_file()) - { - record_time = true; - } + const bool record_time = schedule_task() || statistics.is_calibrating_to_file(); start(); diff --git a/cudax/include/cuda/experimental/__stf/internal/host_launch_scope.cuh b/cudax/include/cuda/experimental/__stf/internal/host_launch_scope.cuh index dd81d675202..6885305164f 100644 --- a/cudax/include/cuda/experimental/__stf/internal/host_launch_scope.cuh +++ b/cudax/include/cuda/experimental/__stf/internal/host_launch_scope.cuh @@ -240,22 +240,32 @@ public: } cudaEvent_t start_event = nullptr, end_event = nullptr; + + SCOPE(exit) + { + if (start_event) + { + cuda_safe_call(cudaEventDestroy(start_event)); + } + if (end_event) + { + cuda_safe_call(cudaEventDestroy(end_event)); + } + }; + const bool record_time = t.schedule_task() || statistics.is_calibrating_to_file(); - // Set only once both timing events exist and the start event has been recorded. - // The timing setup is done below, after the SCOPE(exit) guard is installed, so a - // throw from those cuda_try calls cannot skip t.end_uncleared()/t.clear(). - bool timing_active = false; t.start(); - SCOPE(exit) + // If things go well, end the task with time measurements. + SCOPE(success) { t.end_uncleared(); if constexpr (::std::is_same_v) { - if (timing_active) + if (start_event && end_event) { - // Inside the noexcept SCOPE(exit) body; keep cuda_safe_call so a CUDA + // Inside the noexcept SCOPE body; keep cuda_safe_call so a CUDA // error aborts rather than throwing through the guard. cuda_safe_call(cudaEventRecord(end_event, t.get_stream())); cuda_safe_call(cudaEventSynchronize(end_event)); @@ -277,6 +287,12 @@ public: t.clear(); }; + // And if they don't, just end the task. + SCOPE(fail) + { + t.end(); + }; + if constexpr (::std::is_same_v) { if (record_time) @@ -287,7 +303,6 @@ public: start_event = cuda_try(cudaEventDefault); end_event = cuda_try(cudaEventDefault); cuda_try(start_event, t.get_stream()); - timing_active = true; } } diff --git a/cudax/include/cuda/experimental/__stf/internal/launch.cuh b/cudax/include/cuda/experimental/__stf/internal/launch.cuh index b4b88714ad9..97c97a6dc16 100644 --- a/cudax/include/cuda/experimental/__stf/internal/launch.cuh +++ b/cudax/include/cuda/experimental/__stf/internal/launch.cuh @@ -351,24 +351,29 @@ public: t.set_symbol(symbol); } - bool record_time = t.schedule_task(); + const bool scheduled = t.schedule_task(); // Execution place may have changed during scheduling task - e_place = t.get_exec_place(); - - if (statistics.is_calibrating_to_file()) - { - record_time = true; - } + e_place = t.get_exec_place(); + const bool record_time = scheduled || statistics.is_calibrating_to_file(); nvtx_range nr(t.get_symbol().c_str()); - t.start(); int device = -1; cudaEvent_t start_event = nullptr, end_event = nullptr; - // Set only once both timing events exist and the start event has been recorded. - // The timing setup is done below, after the SCOPE(exit) guard is installed, so a - // throw from those cuda_try calls cannot skip t.end_uncleared()/t.clear(). - bool timing_active = false; + + SCOPE(exit) + { + if (start_event) + { + cuda_safe_call(cudaEventDestroy(start_event)); + } + if (end_event) + { + cuda_safe_call(cudaEventDestroy(end_event)); + } + }; + + t.start(); const size_t grid_size = e_place.size(); @@ -380,13 +385,10 @@ public: auto interpreted_policy = interpreted_execution_policy(spec, e_place, reserved::launch_kernel); - SCOPE(exit) - { - t.end_uncleared(); - + // Free launch-scoped managed temps between end_uncleared and clear on both paths. + auto free_launch_temps = [&] { if constexpr (::std::is_same_v) { - /* If there was managed memory allocated we need to deallocate it */ void* sys_mem = interpreted_policy.get_system_mem(); if (sys_mem) { @@ -399,12 +401,21 @@ public: { deallocateManagedMemory(hostMemoryArrivedList, grid_size, t.get_stream()); } + } + }; + + // If things go well, end the task with time measurements. + SCOPE(success) + { + t.end_uncleared(); + free_launch_temps(); - if (timing_active) + if constexpr (::std::is_same_v) + { + if (start_event && end_event) { - // These run inside the enclosing SCOPE(exit) body, which is noexcept; - // keep cuda_safe_call so a CUDA error aborts rather than throwing - // through the guard (which would call std::terminate). + // Inside the SCOPE body; keep cuda_safe_call so a CUDA error aborts + // rather than throwing through a fail/exit guard (std::terminate). cuda_safe_call(cudaEventRecord(end_event, t.get_stream())); cuda_safe_call(cudaEventSynchronize(end_event)); @@ -426,6 +437,14 @@ public: t.clear(); }; + // And if they don't, free temps and just end the task. + SCOPE(fail) + { + t.end_uncleared(); + free_launch_temps(); + t.clear(); + }; + if constexpr (::std::is_same_v) { if (record_time) @@ -438,7 +457,6 @@ public: start_event = cuda_try(cudaEventDefault); end_event = cuda_try(cudaEventDefault); cuda_try(start_event, t.get_stream()); - timing_active = true; } } diff --git a/cudax/include/cuda/experimental/__stf/stream/stream_task.cuh b/cudax/include/cuda/experimental/__stf/stream/stream_task.cuh index 88f9af5e326..d5500c75802 100644 --- a/cudax/include/cuda/experimental/__stf/stream/stream_task.cuh +++ b/cudax/include/cuda/experimental/__stf/stream/stream_task.cuh @@ -285,17 +285,7 @@ public: template void operator->*(Fun&& fun) { - // Apply function to the stream (in the first position) and the data tuple - nvtx_range nr(get_symbol().c_str()); - start(); - - auto& dot = ctx.get_dot(); - - const bool record_time = reserved::dot::instance().is_timing(); - - cudaEvent_t start_event = nullptr; - cudaEvent_t end_event = nullptr; - bool timing_active = false; + cudaEvent_t start_event = nullptr, end_event = nullptr; SCOPE(exit) { @@ -309,11 +299,19 @@ public: } }; - SCOPE(exit) + // Apply function to the stream (in the first position) and the data tuple + nvtx_range nr(get_symbol().c_str()); + auto& dot = ctx.get_dot(); + const bool record_time = reserved::dot::instance().is_timing(); + + start(); + + // If things go well, end the task with time measuremments, + SCOPE(success) { end_uncleared(); - if (timing_active) + if (start_event && end_event) { cuda_safe_call(cudaEventRecord(end_event, get_stream())); cuda_safe_call(cudaEventSynchronize(end_event)); @@ -330,6 +328,12 @@ public: clear(); }; + // And if they don't, just end the task. + SCOPE(fail) + { + end(); + }; + if (record_time) { // Events must be created here to avoid issues with multi-gpu. @@ -338,7 +342,6 @@ public: start_event = cuda_try(cudaEventDefault); end_event = cuda_try(cudaEventDefault); cuda_try(start_event, get_stream()); - timing_active = true; } // Default for the first argument is a `cudaStream_t`. @@ -532,23 +535,7 @@ public: template auto operator->*(Fun&& fun) { - // Apply function to the stream (in the first position) and the data tuple - auto& dot = ctx.get_dot(); - auto& statistics = reserved::task_statistics::instance(); - - cudaEvent_t start_event = nullptr; - cudaEvent_t end_event = nullptr; - bool timing_active = false; - - bool record_time = schedule_task(); - - if (statistics.is_calibrating_to_file()) - { - record_time = true; - } - - nvtx_range nr(get_symbol().c_str()); - start(); + cudaEvent_t start_event = nullptr, end_event = nullptr; SCOPE(exit) { @@ -562,11 +549,21 @@ public: } }; - SCOPE(exit) + // Apply function to the stream (in the first position) and the data tuple + auto& dot = ctx.get_dot(); + auto& statistics = reserved::task_statistics::instance(); + + const bool record_time = schedule_task() || statistics.is_calibrating_to_file(); + + nvtx_range nr(get_symbol().c_str()); + start(); + + // If things go well, end the task with time measurements, + SCOPE(success) { end_uncleared(); - if (timing_active) + if (start_event && end_event) { cuda_safe_call(cudaEventRecord(end_event, get_stream())); cuda_safe_call(cudaEventSynchronize(end_event)); @@ -588,6 +585,12 @@ public: clear(); }; + // And if they don't, just end the task. + SCOPE(fail) + { + end(); + }; + if (record_time) { // Events must be created here to avoid issues with multi-gpu. @@ -596,7 +599,6 @@ public: start_event = cuda_try(cudaEventDefault); end_event = cuda_try(cudaEventDefault); cuda_try(start_event, get_stream()); - timing_active = true; } if constexpr (::std::is_invocable_v) From 7a17bd915bb8fd05a7d99ba9c3ad4ae71a045418 Mon Sep 17 00:00:00 2001 From: Andrei Alexandrescu Date: Sat, 18 Jul 2026 00:32:22 -0400 Subject: [PATCH 3/5] cudax/stf: close launch exception-safety gaps Build the execution policy before starting the task so teardown guards cover occupancy failures, adopt arrived-list ownership atomically, and install the device-temp free guard before set_device_tmp. --- .../experimental/__stf/internal/launch.cuh | 73 +++++++++++++------ 1 file changed, 51 insertions(+), 22 deletions(-) diff --git a/cudax/include/cuda/experimental/__stf/internal/launch.cuh b/cudax/include/cuda/experimental/__stf/internal/launch.cuh index 97c97a6dc16..def819e0b60 100644 --- a/cudax/include/cuda/experimental/__stf/internal/launch.cuh +++ b/cudax/include/cuda/experimental/__stf/internal/launch.cuh @@ -118,14 +118,10 @@ void launch_impl(interpreted_spec interpreted_policy, exec_place& p, Fun f, Arg th.set_system_tmp(sys_mem); } - if (th_mem_config[1] > 0) - { - cuda_try(cudaMallocAsync(&th_dev_tmp_ptr, th_mem_config[1], stream)); - th.set_device_tmp(th_dev_tmp_ptr); - } - - // Free the temporary device memory on the way out, even if the launch throws. - // cuda_safe_call (not cuda_try) because SCOPE(exit) is noexcept. + // Free the temporary device memory on the way out, even if set_device_tmp + // or the launch throws. Installed before the malloc so a throw from + // set_device_tmp cannot skip the free. cuda_safe_call (not cuda_try) + // because SCOPE(exit) is noexcept. SCOPE(exit) { if (th_dev_tmp_ptr) @@ -134,6 +130,12 @@ void launch_impl(interpreted_spec interpreted_policy, exec_place& p, Fun f, Arg } }; + if (th_mem_config[1] > 0) + { + cuda_try(cudaMallocAsync(&th_dev_tmp_ptr, th_mem_config[1], stream)); + th.set_device_tmp(th_dev_tmp_ptr); + } + auto kernel_args = tuple_prepend(mv(th), mv(arg)); using args_type = decltype(kernel_args); void* all_args[] = {&f, &kernel_args}; @@ -222,7 +224,7 @@ public: unsigned char* hostMemoryArrivedList = interpreted_policy.cg_system.get_arrived_list(); if (hostMemoryArrivedList) { - deallocateManagedMemory(hostMemoryArrivedList, grid_size, streams[0]); + deallocateManagedMemory(hostMemoryArrivedList, grid_size - 1, streams[0]); } }; @@ -232,10 +234,22 @@ public: { if (interpreted_policy.last_level_scope() == hw_scope::device) { - auto hostMemoryArrivedList = (unsigned char*) allocateManagedMemory(grid_size - 1); - // printf("About to allocate hostmemarrivedlist : %lu bytes\n", grid_size - 1); - memset(hostMemoryArrivedList, 0, grid_size - 1); + const size_t arrived_bytes = grid_size - 1; + auto* hostMemoryArrivedList = static_cast(allocateManagedMemory(arrived_bytes)); + // Own the buffer until cg_system adopts it; if the assignment below + // throws, free immediately so we do not leak. cuda_safe_call because + // SCOPE(fail) is noexcept (pool insert can throw). + bool arrived_list_adopted = false; + SCOPE(fail) + { + if (!arrived_list_adopted) + { + cuda_safe_call(cudaFree(hostMemoryArrivedList)); + } + }; + memset(hostMemoryArrivedList, 0, arrived_bytes); interpreted_policy.cg_system = reserved::cooperative_group_system(hostMemoryArrivedList); + arrived_list_adopted = true; } } @@ -373,15 +387,13 @@ public: } }; - t.start(); - const size_t grid_size = e_place.size(); - // Put all data instances in a tuple - auto args = data2inst(t); - + // Build the policy before start(): occupancy queries can throw, and we must + // not leave a started task without teardown guards. args_type only needs the + // unevaluated return type of data2inst, so get() is not invoked here. using th_t = typename thread_hierarchy_spec_t::thread_hierarchy_t; - using args_type = decltype(tuple_prepend(th_t(), args)); + using args_type = decltype(tuple_prepend(th_t(), data2inst(t))); auto interpreted_policy = interpreted_execution_policy(spec, e_place, reserved::launch_kernel); @@ -399,11 +411,13 @@ public: unsigned char* hostMemoryArrivedList = interpreted_policy.cg_system.get_arrived_list(); if (hostMemoryArrivedList) { - deallocateManagedMemory(hostMemoryArrivedList, grid_size, t.get_stream()); + deallocateManagedMemory(hostMemoryArrivedList, grid_size - 1, t.get_stream()); } } }; + t.start(); + // If things go well, end the task with time measurements. SCOPE(success) { @@ -445,6 +459,9 @@ public: t.clear(); }; + // Put all data instances in a tuple (requires a started task). + auto args = data2inst(t); + if constexpr (::std::is_same_v) { if (record_time) @@ -466,10 +483,22 @@ public: { if (interpreted_policy.last_level_scope() == hw_scope::device) { - unsigned char* hostMemoryArrivedList; - hostMemoryArrivedList = (unsigned char*) allocateManagedMemory(grid_size - 1); - memset(hostMemoryArrivedList, 0, grid_size - 1); + const size_t arrived_bytes = grid_size - 1; + auto* hostMemoryArrivedList = static_cast(allocateManagedMemory(arrived_bytes)); + // Own the buffer until cg_system adopts it; if the assignment below + // throws, free immediately so we do not leak. cuda_safe_call because + // SCOPE(fail) is noexcept (pool insert can throw). + bool arrived_list_adopted = false; + SCOPE(fail) + { + if (!arrived_list_adopted) + { + cuda_safe_call(cudaFree(hostMemoryArrivedList)); + } + }; + memset(hostMemoryArrivedList, 0, arrived_bytes); interpreted_policy.cg_system = reserved::cooperative_group_system(hostMemoryArrivedList); + arrived_list_adopted = true; } } From f89181bec3c068fa6f41bc7111f798285418306b Mon Sep 17 00:00:00 2001 From: Andrei Alexandrescu Date: Sat, 18 Jul 2026 00:32:42 -0400 Subject: [PATCH 4/5] cudax/stf: fix swapped streams/arg in launch delegating ctor The 3-arg constructor forwarded arg and streams in the wrong order to the primary constructor. --- cudax/include/cuda/experimental/__stf/internal/launch.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cudax/include/cuda/experimental/__stf/internal/launch.cuh b/cudax/include/cuda/experimental/__stf/internal/launch.cuh index def819e0b60..a1433ac75ff 100644 --- a/cudax/include/cuda/experimental/__stf/internal/launch.cuh +++ b/cudax/include/cuda/experimental/__stf/internal/launch.cuh @@ -182,7 +182,7 @@ public: {} launch(exec_place e_place, ::std::vector streams, Arg arg) - : launch(spec_t(), mv(e_place), mv(arg), mv(streams)) + : launch(spec_t(), mv(e_place), mv(streams), mv(arg)) {} template From 01833c576a1ee5ab9b326b52b69af3d9234a8ba1 Mon Sep 17 00:00:00 2001 From: Andrei Alexandrescu Date: Sat, 18 Jul 2026 00:57:14 -0400 Subject: [PATCH 5/5] cudax/stf: simplify launch failure guards A SCOPE(fail) guard only runs during unwinding, so it can directly free an unadopted arrived list without tracking redundant state. --- .../experimental/__stf/internal/launch.cuh | 38 +++++++------------ 1 file changed, 14 insertions(+), 24 deletions(-) diff --git a/cudax/include/cuda/experimental/__stf/internal/launch.cuh b/cudax/include/cuda/experimental/__stf/internal/launch.cuh index a1433ac75ff..d7b9d6ef1e6 100644 --- a/cudax/include/cuda/experimental/__stf/internal/launch.cuh +++ b/cudax/include/cuda/experimental/__stf/internal/launch.cuh @@ -102,6 +102,18 @@ void launch_impl(interpreted_spec interpreted_policy, exec_place& p, Fun f, Arg void* th_dev_tmp_ptr = nullptr; + // Free the temporary device memory on the way out, even if set_device_tmp + // or the launch throws. Installed before the malloc so a throw from + // set_device_tmp cannot skip the free. cuda_safe_call (not cuda_try) + // because SCOPE(exit) is noexcept. + SCOPE(exit) + { + if (th_dev_tmp_ptr) + { + cuda_safe_call(cudaFreeAsync(th_dev_tmp_ptr, stream)); + } + }; + /* Allocate temporary device memory */ auto th_mem_config = interpreted_policy.get_mem_config(); if (th_mem_config[0] > 0) @@ -118,18 +130,6 @@ void launch_impl(interpreted_spec interpreted_policy, exec_place& p, Fun f, Arg th.set_system_tmp(sys_mem); } - // Free the temporary device memory on the way out, even if set_device_tmp - // or the launch throws. Installed before the malloc so a throw from - // set_device_tmp cannot skip the free. cuda_safe_call (not cuda_try) - // because SCOPE(exit) is noexcept. - SCOPE(exit) - { - if (th_dev_tmp_ptr) - { - cuda_safe_call(cudaFreeAsync(th_dev_tmp_ptr, stream)); - } - }; - if (th_mem_config[1] > 0) { cuda_try(cudaMallocAsync(&th_dev_tmp_ptr, th_mem_config[1], stream)); @@ -239,17 +239,12 @@ public: // Own the buffer until cg_system adopts it; if the assignment below // throws, free immediately so we do not leak. cuda_safe_call because // SCOPE(fail) is noexcept (pool insert can throw). - bool arrived_list_adopted = false; SCOPE(fail) { - if (!arrived_list_adopted) - { - cuda_safe_call(cudaFree(hostMemoryArrivedList)); - } + cuda_safe_call(cudaFree(hostMemoryArrivedList)); }; memset(hostMemoryArrivedList, 0, arrived_bytes); interpreted_policy.cg_system = reserved::cooperative_group_system(hostMemoryArrivedList); - arrived_list_adopted = true; } } @@ -488,17 +483,12 @@ public: // Own the buffer until cg_system adopts it; if the assignment below // throws, free immediately so we do not leak. cuda_safe_call because // SCOPE(fail) is noexcept (pool insert can throw). - bool arrived_list_adopted = false; SCOPE(fail) { - if (!arrived_list_adopted) - { - cuda_safe_call(cudaFree(hostMemoryArrivedList)); - } + cuda_safe_call(cudaFree(hostMemoryArrivedList)); }; memset(hostMemoryArrivedList, 0, arrived_bytes); interpreted_policy.cg_system = reserved::cooperative_group_system(hostMemoryArrivedList); - arrived_list_adopted = true; } }