From e03fd9c4f53468eb555f9bbd2218646ddcd1ab00 Mon Sep 17 00:00:00 2001 From: rtmadduri Date: Thu, 9 Jul 2026 17:09:23 -0500 Subject: [PATCH 1/4] fix hip graph --- warp/_src/context.py | 38 ++++++---- warp/_src/optim/linear.py | 24 ++++-- warp/examples/fem/example_kelvin_helmholtz.py | 11 ++- warp/examples/fem/example_shallow_water.py | 11 ++- warp/examples/fem/example_taylor_green.py | 9 ++- warp/native/warp.cu | 74 ++++++++---------- warp/tests/cuda/test_async.py | 11 ++- warp/tests/cuda/test_streams.py | 19 ++++- warp/tests/geometry/test_bvh.py | 70 +++++++++++++++++ warp/tests/geometry/test_grouped_bvh.py | 75 +++++++++++++++++++ warp/tests/test_apic.py | 12 +-- warp/tests/test_apic_mesh.py | 8 +- warp/tests/test_fixedarray.py | 9 +-- warp/tests/unittest_utils.py | 3 +- 14 files changed, 281 insertions(+), 93 deletions(-) diff --git a/warp/_src/context.py b/warp/_src/context.py index 29a790209d..e3b89face2 100644 --- a/warp/_src/context.py +++ b/warp/_src/context.py @@ -3970,15 +3970,18 @@ def is_capturing(self) -> bool: def supports_graph_capture(self) -> bool: """A boolean indicating whether native graph capture is supported on this device. - Returns ``True`` for CPU (always recordable via APIC) and CUDA devices. - Returns ``False`` for HIP/ROCm devices: hipGraph is not yet mature in - Warp's runtime (mempool pointer remapping bugs, event timing inside - graphs, etc.), so :func:`capture_begin` / :class:`ScopedCapture` are - no-ops on HIP and :func:`capture_save` is unavailable for HIP-targeted - graphs. + Returns ``True`` for CPU (always recordable via APIC), CUDA, and HIP/ROCm + devices. HIP graph capture uses the same native capture/replay path as + CUDA (all ``hipGraph*`` entry points are aliased in ``hip_util.h``). + + Note: two graph features remain gated separately on HIP: + + - Conditional graph nodes (if/else/while subgraphs), tracked by + :func:`is_conditional_graph_supported` (still ``False`` on HIP). + - APIC ``.wrp`` serialization via :func:`capture_save`, which encodes a + CUDA ``sm`` architecture integer and does not yet support HIP ``gfx`` + string architectures. """ - if self.is_hip: - return False return True @property @@ -9453,9 +9456,7 @@ def capture_begin( if stream is None: stream = device.stream - # HIP/ROCm graph capture is not yet mature (mempool pointer remapping bugs, - # event timing inside graphs, etc.). Disable until hipGraph stabilizes. - # See ``Device.supports_graph_capture`` -- callers (notably + # Devices that cannot record a native graph return early. Callers (notably # :class:`ScopedCapture`) treat a ``False`` return as "capture disabled" and # skip the matching :func:`capture_end`. if not device.supports_graph_capture: @@ -9946,11 +9947,9 @@ def capture_launch(graph: Graph, stream: Stream | None = None): if graph is None: # ScopedCapture leaves ``graph=None`` when ``capture_begin`` returned - # False (currently only on HIP, where graph capture is unsupported; - # see ``Device.supports_graph_capture``). + # False on a device where ``Device.supports_graph_capture`` is False. raise RuntimeError( - "capture_launch() received graph=None: capture was not started. " - "On HIP/ROCm, native graph capture is unsupported " + "capture_launch() received graph=None: capture was not started " "(Device.supports_graph_capture is False); guard call sites or " "filter test devices with get_graph_capture_test_devices()." ) @@ -10026,6 +10025,15 @@ def capture_save(graph: Graph, path: str, inputs: dict | None = None, outputs: d import os # noqa: PLC0415 import shutil # noqa: PLC0415 + if graph.device.is_hip: + # The APIC .wrp format encodes a CUDA sm architecture integer and derives + # the device type from it; HIP gfx string architectures are not yet + # representable. In-memory HIP graph capture/replay is unaffected. + raise RuntimeError( + "capture_save() is not supported on HIP/ROCm: the APIC .wrp format does not " + "yet support HIP gfx architectures. In-memory graph capture/replay is supported." + ) + if not graph.apic: raise RuntimeError( "Graph was not captured with apic=True. Pass apic=True to capture_begin() or ScopedCapture()." diff --git a/warp/_src/optim/linear.py b/warp/_src/optim/linear.py index f7431c1141..927e4afc5f 100644 --- a/warp/_src/optim/linear.py +++ b/warp/_src/optim/linear.py @@ -1112,8 +1112,22 @@ def do_cycle_with_condition(): if callback_launch is not None: callback_launch.launch() + def fixed_count_loop(): + # Fallback that avoids conditional graph nodes: run a fixed number of + # cycles. The CG/CR update kernels zero out alpha/beta once the residual + # drops below tolerance (see _cg_kernel_1/_cg_kernel_2), so iterating past + # convergence is a no-op and the result is unchanged. When invoked inside an + # outer capture, these launches are recorded into that graph. + for _ in range(0, maxiter, cycle_size): + do_cycle_with_condition() + if use_cuda_graph and device.is_cuda: - if device.is_capturing: + if not wp.is_conditional_graph_supported(): + # Conditional graph nodes are unavailable (e.g. HIP/ROCm). Use a + # fixed-count loop instead of wp.capture_while so the solver still runs + # (and stays capturable) without requiring if/while subgraph support. + fixed_count_loop() + elif device.is_capturing: wp.capture_while(condition, do_cycle_with_condition) else: with wp.ScopedCapture() as capture: @@ -1121,12 +1135,10 @@ def do_cycle_with_condition(): if capture.graph is not None: wp.capture_launch(capture.graph) else: - # Graph capture not supported (e.g. HIP), fall back to eager loop - for _ in range(0, maxiter, cycle_size): - do_cycle_with_condition() + # Graph capture not supported, fall back to eager loop + fixed_count_loop() else: - for _ in range(0, maxiter, cycle_size): - do_cycle_with_condition() + fixed_count_loop() return cur_iter, r_norm_sq, atol_sq diff --git a/warp/examples/fem/example_kelvin_helmholtz.py b/warp/examples/fem/example_kelvin_helmholtz.py index cc38399038..661becdf29 100644 --- a/warp/examples/fem/example_kelvin_helmholtz.py +++ b/warp/examples/fem/example_kelvin_helmholtz.py @@ -510,14 +510,19 @@ def __init__( self.renderer = fem_example_utils.Plot() self.renderer.add_field("rho", self.rho_field) - # Capture CUDA graph for the simulation step - self.use_cuda_graph = wp.get_device().is_cuda + # Capture CUDA graph for the simulation step. + # NOTE: on HIP/ROCm 7.2, instantiating the large, mempool-allocation-heavy + # graph this DG step produces crashes inside hipGraphInstantiate (heap + # corruption in libamdhip64). Capture itself succeeds, so the graph-is-None + # check below does not catch it; skip capture up front and run eagerly. + device = wp.get_device() + self.use_cuda_graph = device.is_cuda and not device.is_hip if self.use_cuda_graph: with wp.ScopedCapture() as capture: self.simulate() self.graph = capture.graph if self.graph is None: - self.use_cuda_graph = False # Graph capture is disabled on HIP. + self.use_cuda_graph = False def _state_delta(self, trial_state): """Evaluate the DG spatial operator: M^{-1} [volume_rhs + side_rhs].""" diff --git a/warp/examples/fem/example_shallow_water.py b/warp/examples/fem/example_shallow_water.py index 9a7433a6aa..18305fe53a 100644 --- a/warp/examples/fem/example_shallow_water.py +++ b/warp/examples/fem/example_shallow_water.py @@ -433,14 +433,19 @@ def __init__(self, quiet=False, resolution=50, degree=1, domain_size=1.0): self.renderer = fem_example_utils.Plot() self.renderer.add_field("height", self.height_field) - # Capture CUDA graph for the simulation step - self.use_cuda_graph = wp.get_device().is_cuda + # Capture CUDA graph for the simulation step. + # NOTE: on HIP/ROCm 7.2, instantiating the large, mempool-allocation-heavy + # graph this DG step produces crashes inside hipGraphInstantiate (heap + # corruption in libamdhip64). Capture itself succeeds, so the graph-is-None + # check below does not catch it; skip capture up front and run eagerly. + device = wp.get_device() + self.use_cuda_graph = device.is_cuda and not device.is_hip if self.use_cuda_graph: with wp.ScopedCapture() as capture: self.simulate() self.graph = capture.graph if self.graph is None: - self.use_cuda_graph = False # Graph capture is disabled on HIP. + self.use_cuda_graph = False def _state_delta(self, trial_state): """Evaluate the DG spatial operator: M^{-1} [volume_rhs + side_rhs]. diff --git a/warp/examples/fem/example_taylor_green.py b/warp/examples/fem/example_taylor_green.py index 197fde5dae..0a59b56f3c 100644 --- a/warp/examples/fem/example_taylor_green.py +++ b/warp/examples/fem/example_taylor_green.py @@ -557,7 +557,12 @@ def __init__(self, quiet=False, degree=2, resolution=25, Re=100.0, mesh: str = " # reduced launch overhead. Before each replay, step() computes # the time-dependent BC contributions into the working buffers # (_u_bd_rhs_buf, _p_rhs) that the graph references. - self.use_cuda_graph = wp.get_device().is_cuda + # NOTE: on HIP/ROCm 7.2, instantiating the large, mempool-allocation-heavy + # graph this step produces crashes inside hipGraphInstantiate (heap + # corruption in libamdhip64). Capture itself succeeds, so the graph-is-None + # check below does not catch it; skip capture up front and run eagerly. + device = wp.get_device() + self.use_cuda_graph = device.is_cuda and not device.is_hip if self.use_cuda_graph: import gc # noqa: PLC0415 @@ -569,7 +574,7 @@ def __init__(self, quiet=False, degree=2, resolution=25, Re=100.0, mesh: str = " finally: gc.enable() if self.graph is None: - self.use_cuda_graph = False # Graph capture is disabled on HIP. + self.use_cuda_graph = False def _compute_bc(self, t): """Compute time-dependent Dirichlet BC contributions at time ``t``. diff --git a/warp/native/warp.cu b/warp/native/warp.cu index e36053eaee..b401a2e47d 100644 --- a/warp/native/warp.cu +++ b/warp/native/warp.cu @@ -3240,6 +3240,39 @@ bool wp_cuda_graph_update_memcpy_batch( return true; } +// Pause/resume of an in-progress stream capture. These use only standard +// capture APIs (cudaStreamEndCapture / cudaStreamBeginCaptureToGraph, both +// aliased for HIP in hip_util.h) and are independent of conditional graph +// nodes, so they are compiled for all backends including HIP/ROCm. +bool wp_cuda_graph_pause_capture(void* context, void* stream, void** graph_ret) +{ + ContextGuard guard(context); + + CUstream cuda_stream = static_cast(stream); + if (!check_cuda(cudaStreamEndCapture(cuda_stream, (cudaGraph_t*)graph_ret))) + return false; + return true; +} + +bool wp_cuda_graph_resume_capture(void* context, void* stream, void* graph) +{ + ContextGuard guard(context); + + CUstream cuda_stream = static_cast(stream); + cudaGraph_t cuda_graph = static_cast(graph); + + std::vector leaf_nodes; + if (!get_graph_leaf_nodes(cuda_graph, leaf_nodes)) + return false; + + if (!check_cuda(cudaStreamBeginCaptureToGraph( + cuda_stream, cuda_graph, leaf_nodes.data(), nullptr, leaf_nodes.size(), cudaStreamCaptureModeThreadLocal + ))) + return false; + + return true; +} + // Support for conditional graph nodes available with CUDA 12.4+. #if !defined(__HIP_PLATFORM_AMD__) && CUDA_VERSION >= 12040 @@ -3387,35 +3420,6 @@ static CUfunction get_conditional_kernel(void* context, int arch, bool use_ptx, return kernel; } -bool wp_cuda_graph_pause_capture(void* context, void* stream, void** graph_ret) -{ - ContextGuard guard(context); - - CUstream cuda_stream = static_cast(stream); - if (!check_cuda(cudaStreamEndCapture(cuda_stream, (cudaGraph_t*)graph_ret))) - return false; - return true; -} - -bool wp_cuda_graph_resume_capture(void* context, void* stream, void* graph) -{ - ContextGuard guard(context); - - CUstream cuda_stream = static_cast(stream); - cudaGraph_t cuda_graph = static_cast(graph); - - std::vector leaf_nodes; - if (!get_graph_leaf_nodes(cuda_graph, leaf_nodes)) - return false; - - if (!check_cuda(cudaStreamBeginCaptureToGraph( - cuda_stream, cuda_graph, leaf_nodes.data(), nullptr, leaf_nodes.size(), cudaStreamCaptureModeThreadLocal - ))) - return false; - - return true; -} - // https://developer.nvidia.com/blog/constructing-cuda-graphs-with-dynamic-parameters/#combined_approach // https://developer.nvidia.com/blog/dynamic-control-flow-in-cuda-graphs-with-conditional-nodes/ // condition is a gpu pointer @@ -3807,18 +3811,6 @@ bool wp_cuda_graph_set_condition(void* context, void* stream, int arch, bool use #else // stubs for conditional graph node API if CUDA toolkit is too old. -bool wp_cuda_graph_pause_capture(void* context, void* stream, void** graph_ret) -{ - wp::set_error_string("Warp error: Warp must be built with CUDA Toolkit 12.4+ to enable conditional graph nodes"); - return false; -} - -bool wp_cuda_graph_resume_capture(void* context, void* stream, void* graph) -{ - wp::set_error_string("Warp error: Warp must be built with CUDA Toolkit 12.4+ to enable conditional graph nodes"); - return false; -} - bool wp_cuda_graph_insert_if_else( void* context, void* stream, int arch, bool use_ptx, int* condition, void** if_graph_ret, void** else_graph_ret ) diff --git a/warp/tests/cuda/test_async.py b/warp/tests/cuda/test_async.py index e87558a1dc..1103da7d55 100644 --- a/warp/tests/cuda/test_async.py +++ b/warp/tests/cuda/test_async.py @@ -583,9 +583,14 @@ def test_func( else: grad_flags = [False] - # graph capture options (CUDA-only, not supported on HIP) - graph_supported = ( - (src_device.is_cuda and not src_device.is_hip) or (dst_device.is_cuda and not dst_device.is_hip) + # graph capture options. Graph capture of async copies is not yet + # enabled on HIP/ROCm: copies that allocate a staging buffer during + # capture and the expected-error paths do not clean up the stream + # capture reliably on HIP, poisoning subsequent tests. Basic capture + # and replay is validated separately (see test_graph). Deferred until + # the async-copy-capture path is hardened on HIP. + graph_supported = (src_device.is_cuda and not src_device.is_hip) or ( + dst_device.is_cuda and not dst_device.is_hip ) if graph_supported: graph_flags = [False, True] diff --git a/warp/tests/cuda/test_streams.py b/warp/tests/cuda/test_streams.py index 00922fa449..a3623f6e96 100644 --- a/warp/tests/cuda/test_streams.py +++ b/warp/tests/cuda/test_streams.py @@ -331,6 +331,11 @@ def test_event_elapsed_time(test, device): def test_event_elapsed_time_graph(test, device): + if wp.get_device(device).is_hip: + # ROCm limitation: hipEventElapsedTime() on timing events recorded inside + # a captured graph fails with "invalid resource handle" (HIP error 400). + test.skipTest("Event timing inside a captured graph is not supported on HIP/ROCm") + stream = wp.get_stream(device) e1 = wp.Event(device, enable_timing=True) e2 = wp.Event(device, enable_timing=True) @@ -354,6 +359,12 @@ def test_event_elapsed_time_graph(test, device): def test_event_external(test, device): + if wp.get_device(device).is_hip: + # ROCm limitation: an external event recorded in one captured graph does + # not reliably synchronize a wait_event() in a separately-captured graph + # launched on another stream, producing a race (non-deterministic result). + test.skipTest("External event synchronization across captured graphs is not supported on HIP/ROCm") + with wp.ScopedDevice(device): # event used to synchronize two graphs (external event) event = wp.Event() @@ -572,8 +583,8 @@ def test_stream_exceptions(self): @unittest.skipUnless(len(wp.get_cuda_devices()) > 1, "Requires at least two CUDA devices") @unittest.skipUnless(check_p2p(), "Peer-to-Peer transfers not supported") def test_stream_arg_graph_mgpu(self): - if any(d.is_hip for d in wp.get_cuda_devices()): - self.skipTest("CUDA graph capture is not supported on HIP") + if not all(d.supports_graph_capture for d in wp.get_cuda_devices()): + self.skipTest("Graph capture is not supported on all CUDA devices") wp.load_module(device="cuda:0") wp.load_module(device="cuda:1") @@ -624,8 +635,8 @@ def test_stream_arg_graph_mgpu(self): @unittest.skipUnless(len(wp.get_cuda_devices()) > 1, "Requires at least two CUDA devices") @unittest.skipUnless(check_p2p(), "Peer-to-Peer transfers not supported") def test_stream_scope_graph_mgpu(self): - if any(d.is_hip for d in wp.get_cuda_devices()): - self.skipTest("CUDA graph capture is not supported on HIP") + if not all(d.supports_graph_capture for d in wp.get_cuda_devices()): + self.skipTest("Graph capture is not supported on all CUDA devices") wp.load_module(device="cuda:0") wp.load_module(device="cuda:1") diff --git a/warp/tests/geometry/test_bvh.py b/warp/tests/geometry/test_bvh.py index 53ea4a801a..37360cfdf4 100644 --- a/warp/tests/geometry/test_bvh.py +++ b/warp/tests/geometry/test_bvh.py @@ -202,6 +202,14 @@ def compute_num_contact_with_checksums( def test_capture_bvh_rebuild(test, device): + if wp.get_device(device).is_hip: + # The LBVH builder's radix sort (rocPRIM/hipcub DeviceRadixSort) is not + # graph-capturable on HIP/ROCm: it triggers an error inside the stream + # capture (Bvh.rebuild() works fine outside capture). Basic graph + # capture/replay is validated by test_graph; this in-capture BVH rebuild + # path is deferred pending a graph-capturable HIP sort. + test.skipTest("BVH rebuild inside a captured graph is not supported on HIP/ROCm") + with wp.ScopedDevice(device): rng = np.random.default_rng(123) @@ -283,6 +291,67 @@ def test_capture_bvh_rebuild(test, device): assert_array_equal(checksums_1, checksums_2) +def test_capture_bvh_query(test, device): + # Graph capture of BVH *queries* (the BVH is built outside the capture). + # Unlike test_capture_bvh_rebuild, this does not build/rebuild inside the + # capture, so it is supported on HIP/ROCm as well as CUDA. + with wp.ScopedDevice(device): + rng = np.random.default_rng(123) + + num_item_bounds = 100000 + item_bound_size = 0.01 + relative_shift = 2 + + num_test_bounds = 10000 + test_bound_relative_size = 0.05 + + center = np.array([0.0, 0.0, 0.0]) + + item_lowers_np, item_uppers_np = get_random_aabbs(num_item_bounds, center, relative_shift, item_bound_size, rng) + item_lowers = wp.array(item_lowers_np, dtype=wp.vec3) + item_uppers = wp.array(item_uppers_np, dtype=wp.vec3) + + # Build the BVH outside of the capture. + bvh = wp.Bvh(item_lowers, item_uppers) + + test_lowers_np, test_uppers_np = get_random_aabbs( + num_test_bounds, center, relative_shift, test_bound_relative_size, rng + ) + test_lowers = wp.array(test_lowers_np, dtype=wp.vec3) + test_uppers = wp.array(test_uppers_np, dtype=wp.vec3) + + counts_ref = wp.empty(n=num_test_bounds, dtype=int) + checksums_ref = wp.empty(n=num_test_bounds, dtype=int) + counts_graph = wp.empty(n=num_test_bounds, dtype=int) + checksums_graph = wp.empty(n=num_test_bounds, dtype=int) + + # Reference result computed with an eager launch (no capture). + wp.launch( + compute_num_contact_with_checksums, + dim=num_test_bounds, + inputs=[test_lowers, test_uppers, bvh.id], + outputs=[counts_ref, checksums_ref], + ) + + # Capture only the query launch, then replay it several times. + wp.load_module(device=device) + with wp.ScopedCapture(force_module_load=False) as capture: + wp.launch( + compute_num_contact_with_checksums, + dim=num_test_bounds, + inputs=[test_lowers, test_uppers, bvh.id], + outputs=[counts_graph, checksums_graph], + ) + + test.assertIsNotNone(capture.graph) + + for _ in range(10): + wp.capture_launch(capture.graph) + + assert_array_equal(counts_graph, counts_ref) + assert_array_equal(checksums_graph, checksums_ref) + + @wp.kernel def tile_bvh_query_aabb_kernel( bvh_id: wp.uint64, @@ -750,6 +819,7 @@ def test_bvh_new_del(self): add_function_test(TestBvh, "test_bvh_query_ray_tiled", test_bvh_query_ray_tiled, devices=cuda_devices) add_function_test(TestBvh, "test_capture_bvh_rebuild", test_capture_bvh_rebuild, devices=cuda_graph_devices) +add_function_test(TestBvh, "test_capture_bvh_query", test_capture_bvh_query, devices=cuda_graph_devices) if __name__ == "__main__": unittest.main(verbosity=2) diff --git a/warp/tests/geometry/test_grouped_bvh.py b/warp/tests/geometry/test_grouped_bvh.py index 3e6a73b96a..99ef5a717f 100644 --- a/warp/tests/geometry/test_grouped_bvh.py +++ b/warp/tests/geometry/test_grouped_bvh.py @@ -420,6 +420,14 @@ def compute_num_contact_with_checksums( def test_capture_bvh_rebuild_grouped(test, device): + if wp.get_device(device).is_hip: + # The LBVH builder's radix sort (rocPRIM/hipcub DeviceRadixSort) is not + # graph-capturable on HIP/ROCm: it triggers an error inside the stream + # capture (Bvh.rebuild() works fine outside capture). Basic graph + # capture/replay is validated by test_graph; this in-capture BVH rebuild + # path is deferred pending a graph-capturable HIP sort. + test.skipTest("BVH rebuild inside a captured graph is not supported on HIP/ROCm") + with wp.ScopedDevice(device): rng = np.random.default_rng(123) @@ -517,6 +525,70 @@ def test_capture_bvh_rebuild_grouped(test, device): assert_array_equal(checksums_1, checksums_brutal) +def test_capture_bvh_query_grouped(test, device): + # Graph capture of grouped BVH *queries* (the BVH is built outside the + # capture). Unlike test_capture_bvh_rebuild_grouped, nothing is built or + # rebuilt inside the capture, so it is supported on HIP/ROCm as well as CUDA. + with wp.ScopedDevice(device): + rng = np.random.default_rng(123) + + num_item_bounds = 100000 + num_groups = 100 + item_bound_size = 0.01 + relative_shift = 2 + + num_test_bounds = 10000 + test_bound_relative_size = 0.05 + + center = np.array([0.0, 0.0, 0.0]) + + item_lowers_np, item_uppers_np = get_random_aabbs(num_item_bounds, center, relative_shift, item_bound_size, rng) + item_lowers = wp.array(item_lowers_np, dtype=wp.vec3) + item_uppers = wp.array(item_uppers_np, dtype=wp.vec3) + item_groups = wp.array(np.arange(num_item_bounds, dtype=np.int32) % num_groups, dtype=int) + + # Build the grouped BVH outside of the capture. + bvh = wp.Bvh(item_lowers, item_uppers, groups=item_groups) + + test_lowers_np, test_uppers_np = get_random_aabbs( + num_test_bounds, center, relative_shift, test_bound_relative_size, rng + ) + test_lowers = wp.array(test_lowers_np, dtype=wp.vec3) + test_uppers = wp.array(test_uppers_np, dtype=wp.vec3) + test_groups = wp.array(np.arange(num_item_bounds, dtype=np.int32) % num_groups, dtype=int) + + counts_ref = wp.empty(n=num_test_bounds, dtype=int) + checksums_ref = wp.empty(n=num_test_bounds, dtype=int) + counts_graph = wp.empty(n=num_test_bounds, dtype=int) + checksums_graph = wp.empty(n=num_test_bounds, dtype=int) + + # Reference result computed with an eager launch (no capture). + wp.launch( + kernel=compute_num_contact_with_checksums, + dim=num_test_bounds, + inputs=[test_lowers, test_uppers, test_groups, bvh.id], + outputs=[counts_ref, checksums_ref], + ) + + # Capture only the query launch, then replay it several times. + wp.load_module(device=device) + with wp.ScopedCapture(force_module_load=False) as capture: + wp.launch( + kernel=compute_num_contact_with_checksums, + dim=num_test_bounds, + inputs=[test_lowers, test_uppers, test_groups, bvh.id], + outputs=[counts_graph, checksums_graph], + ) + + test.assertIsNotNone(capture.graph) + + for _ in range(10): + wp.capture_launch(capture.graph) + + assert_array_equal(counts_graph, counts_ref) + assert_array_equal(checksums_graph, checksums_ref) + + devices = get_test_devices() cuda_devices = get_cuda_test_devices() cuda_graph_devices = [d for d in cuda_devices if d.supports_graph_capture] @@ -559,6 +631,9 @@ def test_bvh_new_del(self): add_function_test( TestGroupedBvh, "test_grouped_capture_bvh_rebuild", test_capture_bvh_rebuild_grouped, devices=cuda_graph_devices ) +add_function_test( + TestGroupedBvh, "test_grouped_capture_bvh_query", test_capture_bvh_query_grouped, devices=cuda_graph_devices +) if __name__ == "__main__": unittest.main(verbosity=2) diff --git a/warp/tests/test_apic.py b/warp/tests/test_apic.py index 4061e34e6a..422aaa2814 100644 --- a/warp/tests/test_apic.py +++ b/warp/tests/test_apic.py @@ -543,9 +543,11 @@ def test_get_param_ptr(test, device): plain_capture.graph.get_param_ptr("a") -# APIC capture/save/load requires native graph capture, which is not yet -# supported on HIP/ROCm (see Device.supports_graph_capture). -devices = get_graph_capture_test_devices() +# APIC capture/save/load requires native graph capture (see +# Device.supports_graph_capture). The .wrp serialization format does not yet +# support HIP gfx architectures, so HIP/ROCm devices are excluded here (see +# capture_save()). +devices = [d for d in get_graph_capture_test_devices() if not d.is_hip] add_function_test(TestApic, "test_save_apic_false_error", test_save_apic_false_error, devices=devices) add_function_test(TestApic, "test_save_single_kernel", test_save_single_kernel, devices=devices) @@ -565,8 +567,8 @@ def test_get_param_ptr(test, device): TestApic, "test_save_load_fill", test_save_load_fill, - devices=[d for d in get_cuda_test_devices() if d.supports_graph_capture], -) # CPU: wp_memtile_host not recorded; HIP: no native graph capture + devices=[d for d in get_cuda_test_devices() if d.supports_graph_capture and not d.is_hip], +) # CPU: wp_memtile_host not recorded; HIP: .wrp format lacks gfx arch support (see capture_save) add_function_test(TestApic, "test_save_load_alloc_only", test_save_load_alloc_only, devices=devices) add_function_test(TestApic, "test_get_param_ptr", test_get_param_ptr, devices=devices) diff --git a/warp/tests/test_apic_mesh.py b/warp/tests/test_apic_mesh.py index 96a3053be5..5f14b260d5 100644 --- a/warp/tests/test_apic_mesh.py +++ b/warp/tests/test_apic_mesh.py @@ -563,9 +563,11 @@ class TestApicMesh(unittest.TestCase): pass -# APIC capture/save/load requires native graph capture, which is not yet -# supported on HIP/ROCm (see Device.supports_graph_capture). -devices = get_graph_capture_test_devices() +# APIC capture/save/load requires native graph capture (see +# Device.supports_graph_capture). The .wrp serialization format does not yet +# support HIP gfx architectures, so HIP/ROCm devices are excluded here (see +# capture_save()). +devices = [d for d in get_graph_capture_test_devices() if not d.is_hip] add_function_test(TestApicMesh, "test_apic_mesh_query_point", test_apic_mesh_query_point, devices=devices) add_function_test(TestApicMesh, "test_apic_mesh_query_ray", test_apic_mesh_query_ray, devices=devices) diff --git a/warp/tests/test_fixedarray.py b/warp/tests/test_fixedarray.py index 72e9f24d5e..0e9022e5f6 100644 --- a/warp/tests/test_fixedarray.py +++ b/warp/tests/test_fixedarray.py @@ -157,12 +157,9 @@ def test_capture_if_kernel(): def test_capture_if(test, device): - if ( - not wp.get_device(device).is_cuda - or not wp.get_device(device).supports_graph_capture - or wp._src.context.runtime.toolkit_version < (12, 4) - or wp._src.context.runtime.driver_version < (12, 4) - ): + # Conditional graph nodes require CUDA 12.4+ toolkit/driver and are not + # supported on HIP/ROCm; is_conditional_graph_supported() covers all cases. + if not wp.get_device(device).is_cuda or not wp.is_conditional_graph_supported(): return def foo(): diff --git a/warp/tests/unittest_utils.py b/warp/tests/unittest_utils.py index 4622a05914..63612f657b 100644 --- a/warp/tests/unittest_utils.py +++ b/warp/tests/unittest_utils.py @@ -128,8 +128,7 @@ def get_cuda_test_devices(mode=None): def get_graph_capture_test_devices(mode: str | None = None): """Subset of :func:`get_test_devices` filtered to devices that support - native graph capture (CPU and non-HIP CUDA). HIP/ROCm devices are excluded - because :func:`warp.capture_begin` is a no-op on them; see + native graph capture (CPU, CUDA, and HIP/ROCm); see ``Device.supports_graph_capture``. """ return [d for d in get_test_devices(mode=mode) if d.supports_graph_capture] From 1181bee56902635603d32fe1d94f3213f4194386 Mon Sep 17 00:00:00 2001 From: rtmadduri Date: Thu, 16 Jul 2026 22:58:50 -0400 Subject: [PATCH 2/4] fix sprase.cu for test_sparse, test_fem_integrate --- warp/_src/context.py | 14 +++- warp/native/sparse.cu | 111 ++++++++++++++++++--------- warp/tests/cuda/test_async.py | 7 +- warp/tests/fem/test_fem_integrate.py | 9 ++- warp/tests/test_apic.py | 15 ++-- warp/tests/test_sparse.py | 11 ++- 6 files changed, 117 insertions(+), 50 deletions(-) diff --git a/warp/_src/context.py b/warp/_src/context.py index e3b89face2..97afe7cd2c 100644 --- a/warp/_src/context.py +++ b/warp/_src/context.py @@ -8746,13 +8746,21 @@ def synchronize(): # save the original context to avoid side effects saved_context = runtime.core.wp_cuda_context_get_current() + # Refuse to synchronize while a graph capture is active on ANY device + # before performing any partial synchronization. On HIP/ROCm graph + # capture is process-global, so synchronizing a non-capturing device's + # context while another device is capturing fails and invalidates the + # in-progress capture. Checking all devices up front keeps the error + # clean and leaves the capture intact on both CUDA and HIP. + for device in runtime.cuda_devices: + # avoid creating primary context if the device has not been used yet + if device.has_context and device.is_capturing: + raise RuntimeError(f"Cannot synchronize device {device} while graph capture is active") + # TODO: only synchronize devices that have outstanding work for device in runtime.cuda_devices: # avoid creating primary context if the device has not been used yet if device.has_context: - if device.is_capturing: - raise RuntimeError(f"Cannot synchronize device {device} while graph capture is active") - runtime.core.wp_cuda_context_synchronize(device.context) # restore the original context to avoid side effects diff --git a/warp/native/sparse.cu b/warp/native/sparse.cu index 80fe3dd263..bd7c1c114b 100644 --- a/warp/native/sparse.cu +++ b/warp/native/sparse.cu @@ -7,6 +7,7 @@ #include "temp_buffer.h" #include +#include #define THRUST_IGNORE_CUB_VERSION_CHECK @@ -24,6 +25,41 @@ extern CUcontext get_current_context(); namespace { +// Persistent per-stream scratch for CUB device algorithms (sort/RLE/scan) used +// while building BSR matrices. Allocating CUB temp storage on the fly through a +// graph-captured mempool allocation causes GPU memory faults on replay on +// non-default HIP devices (peer-mapped pool memory reallocated inside the graph +// is not reliably accessible during replay). Reusing a persistent buffer that +// was sized during warmup avoids any allocation inside the capture, mirroring +// the approach already used for radix sort (see g_radix_sort_temp_map). +struct BsrCubTemp { + void* mem = nullptr; + size_t size = 0; +}; + +// Keyed by (stream, slot). Distinct slots are needed for buffers that must be +// live simultaneously (e.g. CUB temp storage and the sort key/value buffers). +static std::unordered_map> g_bsr_scratch_map; + +static void* bsr_reserve_scratch(void* context, int slot, size_t size) +{ + void* stream = wp_cuda_stream_get_current(); + BsrCubTemp& temp = g_bsr_scratch_map[stream][slot]; + if (size > temp.size) { + // Only grows outside of capture (warmup); during capture the buffer is + // already large enough and is reused without a new allocation. + wp_free_device(context, temp.mem); + temp.mem = wp_alloc_device(context, size, "(native:sparse)"); + temp.size = size; + } + return temp.mem; +} + +static void* bsr_reserve_cub_temp(void* context, size_t size) +{ + return bsr_reserve_scratch(context, 0, size); +} + // Combined row+column value that can be radix-sorted with CUB using BsrRowCol = uint64_t; @@ -273,19 +309,26 @@ WP_API void wp_bsr_matrix_from_triplets_device( cudaStream_t stream = static_cast(wp_cuda_stream_get_current()); - ScopedTemporary combined_row_col(context, 2 * size_t(nnz)); - ScopedTemporary unique_triplet_count(context, 1); + BsrRowCol* combined_row_col_buf = + static_cast(bsr_reserve_scratch(context, 1, 2 * size_t(nnz) * sizeof(BsrRowCol))); + int* unique_triplet_count_buf = static_cast(bsr_reserve_scratch(context, 2, sizeof(int))); bool return_summed_blocks = tpl_block_offsets != nullptr && tpl_block_indices != nullptr; - if (!return_summed_blocks) { - // if not provided, allocate temporary offset and indices buffers - tpl_block_offsets = static_cast(wp_alloc_device(context, size_t(nnz) * sizeof(int), "(native:sparse)")); - tpl_block_indices = static_cast(wp_alloc_device(context, size_t(nnz) * sizeof(int), "(native:sparse)")); - } - - cub::DoubleBuffer d_keys(tpl_block_indices, tpl_block_offsets); - cub::DoubleBuffer d_values(combined_row_col.buffer(), combined_row_col.buffer() + nnz); + // All buffers read/written by CUB device algorithms (radix sort, run-length + // encode, inclusive scan) must be persistent rather than graph-captured + // mempool allocations. On non-default HIP devices, CUB operating on memory + // that was allocated inside a graph capture faults during replay (the + // caller-provided summed_triplet_offsets/indices arrays are allocated inside + // the capture). We therefore compute the topology entirely in persistent + // per-stream scratch and copy the results into the caller-provided output + // buffers at the end. + int* block_indices_buf = static_cast(bsr_reserve_scratch(context, 3, size_t(nnz) * sizeof(int))); + int* block_offsets_buf = static_cast(bsr_reserve_scratch(context, 4, size_t(nnz) * sizeof(int))); + int* run_offsets_buf = static_cast(bsr_reserve_scratch(context, 5, size_t(nnz) * sizeof(int))); + + cub::DoubleBuffer d_keys(block_indices_buf, block_offsets_buf); + cub::DoubleBuffer d_values(combined_row_col_buf, combined_row_col_buf + nnz); // Combine rows and columns so we can sort on them both, // ensuring that blocks that should be pruned are moved to the end @@ -331,16 +374,8 @@ WP_API void wp_bsr_matrix_from_triplets_device( { size_t buff_size = 0; check_cuda(cub::DeviceRadixSort::SortPairs(nullptr, buff_size, d_values, d_keys, nnz, 0, 64, stream)); - ScopedTemporary<> temp(context, buff_size); - check_cuda(cub::DeviceRadixSort::SortPairs(temp.buffer(), buff_size, d_values, d_keys, nnz, 0, 64, stream)); - - // Depending on data size and GPU architecture buffers may have been swapped or not - // Ensures the sorted keys are available in summed_block_indices if needed - if (return_summed_blocks && d_keys.Current() != tpl_block_indices) { - check_cuda(cudaMemcpyAsync( - tpl_block_indices, d_keys.Current(), nnz * sizeof(int), cudaMemcpyDeviceToDevice, stream - )); - } + void* temp = bsr_reserve_cub_temp(context, buff_size); + check_cuda(cub::DeviceRadixSort::SortPairs(temp, buff_size, d_values, d_keys, nnz, 0, 64, stream)); } // Runlength encode row-col sequences @@ -348,15 +383,15 @@ WP_API void wp_bsr_matrix_from_triplets_device( size_t buff_size = 0; check_cuda( cub::DeviceRunLengthEncode::Encode( - nullptr, buff_size, d_values.Current(), d_values.Alternate(), tpl_block_offsets, - unique_triplet_count.buffer(), nnz, stream + nullptr, buff_size, d_values.Current(), d_values.Alternate(), run_offsets_buf, + unique_triplet_count_buf, nnz, stream ) ); - ScopedTemporary<> temp(context, buff_size); + void* temp = bsr_reserve_cub_temp(context, buff_size); check_cuda( cub::DeviceRunLengthEncode::Encode( - temp.buffer(), buff_size, d_values.Current(), d_values.Alternate(), tpl_block_offsets, - unique_triplet_count.buffer(), nnz, stream + temp, buff_size, d_values.Current(), d_values.Alternate(), run_offsets_buf, + unique_triplet_count_buf, nnz, stream ) ); } @@ -364,7 +399,7 @@ WP_API void wp_bsr_matrix_from_triplets_device( // Compute row offsets from sorted unique blocks wp_launch_device( WP_CURRENT_CONTEXT, bsr_find_row_offsets, row_count + 1, - (row_count, unique_triplet_count.buffer(), d_values.Alternate(), bsr_offsets) + (row_count, unique_triplet_count_buf, d_values.Alternate(), bsr_offsets) ); if (bsr_nnz) { @@ -383,20 +418,24 @@ WP_API void wp_bsr_matrix_from_triplets_device( WP_CURRENT_CONTEXT, bsr_set_column, nnz, (bsr_offsets + row_count, d_values.Alternate(), bsr_columns) ); - // Scan repeated block counts + // Scan repeated block counts and copy the summed topology into the + // caller-provided output buffers (which may be graph-captured allocations). if (return_summed_blocks) { size_t buff_size = 0; check_cuda( - cub::DeviceScan::InclusiveSum(nullptr, buff_size, tpl_block_offsets, tpl_block_offsets, nnz, stream) + cub::DeviceScan::InclusiveSum(nullptr, buff_size, run_offsets_buf, run_offsets_buf, nnz, stream) ); - ScopedTemporary<> temp(context, buff_size); + void* temp = bsr_reserve_cub_temp(context, buff_size); check_cuda( - cub::DeviceScan::InclusiveSum(temp.buffer(), buff_size, tpl_block_offsets, tpl_block_offsets, nnz, stream) + cub::DeviceScan::InclusiveSum(temp, buff_size, run_offsets_buf, run_offsets_buf, nnz, stream) ); - } else { - // free our temporary buffers - wp_free_device(context, tpl_block_offsets); - wp_free_device(context, tpl_block_indices); + + check_cuda(cudaMemcpyAsync( + tpl_block_offsets, run_offsets_buf, size_t(nnz) * sizeof(int), cudaMemcpyDeviceToDevice, stream + )); + check_cuda(cudaMemcpyAsync( + tpl_block_indices, d_keys.Current(), size_t(nnz) * sizeof(int), cudaMemcpyDeviceToDevice, stream + )); } } @@ -431,8 +470,8 @@ WP_API void wp_bsr_transpose_device( { size_t buff_size = 0; check_cuda(cub::DeviceRadixSort::SortPairs(nullptr, buff_size, d_values, d_keys, nnz, 0, 64, stream)); - ScopedTemporary<> temp(context, buff_size); - check_cuda(cub::DeviceRadixSort::SortPairs(temp.buffer(), buff_size, d_values, d_keys, nnz, 0, 64, stream)); + void* temp = bsr_reserve_cub_temp(context, buff_size); + check_cuda(cub::DeviceRadixSort::SortPairs(temp, buff_size, d_values, d_keys, nnz, 0, 64, stream)); // Depending on data size and GPU architecture buffers may have been swapped or not // Ensures the sorted keys are available in summed_block_indices if needed diff --git a/warp/tests/cuda/test_async.py b/warp/tests/cuda/test_async.py index 1103da7d55..b5d9f25088 100644 --- a/warp/tests/cuda/test_async.py +++ b/warp/tests/cuda/test_async.py @@ -586,9 +586,10 @@ def test_func( # graph capture options. Graph capture of async copies is not yet # enabled on HIP/ROCm: copies that allocate a staging buffer during # capture and the expected-error paths do not clean up the stream - # capture reliably on HIP, poisoning subsequent tests. Basic capture - # and replay is validated separately (see test_graph). Deferred until - # the async-copy-capture path is hardened on HIP. + # capture reliably on HIP, poisoning subsequent tests (and cross-stream + # d2d capture can crash the process). Basic capture and replay is + # validated separately (see test_graph). Deferred until the async-copy + # capture path is hardened on HIP. graph_supported = (src_device.is_cuda and not src_device.is_hip) or ( dst_device.is_cuda and not dst_device.is_hip ) diff --git a/warp/tests/fem/test_fem_integrate.py b/warp/tests/fem/test_fem_integrate.py index 4b6476cee2..71bd2bbbbe 100644 --- a/warp/tests/fem/test_fem_integrate.py +++ b/warp/tests/fem/test_fem_integrate.py @@ -574,7 +574,14 @@ class TestFemIntegrate(unittest.TestCase): add_function_test(TestFemIntegrate, "test_grad_decomposition", test_grad_decomposition, devices=devices) add_function_test(TestFemIntegrate, "test_integrate_high_order", test_integrate_high_order, devices=cuda_graph_devices) add_function_test(TestFemIntegrate, "test_interpolate_reduction", test_interpolate_reduction, devices=devices) -add_function_test(TestFemIntegrate, "test_capturability", test_capturability, devices=cuda_graph_devices) +# On HIP this test faults on non-primary devices: `fem.integrate` builds BSR +# topology by chaining data-dependent matrix creations inside the capture, which +# a ROCm graph-mempool limitation cannot replay on a non-primary device (the +# in-capture allocations are backed by device-0 memory and fault on replay). +# Since a GPU memory-access fault aborts the whole process, restrict to the +# primary device on HIP. See HIP_GRAPH_CAPTURE_TODO.md section 2.1a. +capturability_devices = [d for d in cuda_graph_devices if not (d.is_hip and d.ordinal != 0)] +add_function_test(TestFemIntegrate, "test_capturability", test_capturability, devices=capturability_devices) if __name__ == "__main__": unittest.main(verbosity=2, failfast=True) diff --git a/warp/tests/test_apic.py b/warp/tests/test_apic.py index 422aaa2814..3f33f9ed14 100644 --- a/warp/tests/test_apic.py +++ b/warp/tests/test_apic.py @@ -543,12 +543,15 @@ def test_get_param_ptr(test, device): plain_capture.graph.get_param_ptr("a") -# APIC capture/save/load requires native graph capture (see -# Device.supports_graph_capture). The .wrp serialization format does not yet -# support HIP gfx architectures, so HIP/ROCm devices are excluded here (see -# capture_save()). +# APIC save/load requires .wrp serialization, which does not yet support HIP gfx +# architectures, so HIP/ROCm devices are excluded from tests that call +# capture_save()/capture_load() (see capture_save()). devices = [d for d in get_graph_capture_test_devices() if not d.is_hip] +# Tests that only exercise in-memory APIC capture/replay (no .wrp serialization) +# run on every graph-capture device, including HIP/ROCm. +graph_devices = get_graph_capture_test_devices() + add_function_test(TestApic, "test_save_apic_false_error", test_save_apic_false_error, devices=devices) add_function_test(TestApic, "test_save_single_kernel", test_save_single_kernel, devices=devices) add_function_test(TestApic, "test_save_load_round_trip", test_save_load_round_trip, devices=devices) @@ -556,11 +559,11 @@ def test_get_param_ptr(test, device): add_function_test(TestApic, "test_save_load_memcpy", test_save_load_memcpy, devices=devices) add_function_test(TestApic, "test_save_load_memset", test_save_load_memset, devices=devices) add_function_test(TestApic, "test_bindings_param_update", test_bindings_param_update, devices=devices) -add_function_test(TestApic, "test_array_slicing", test_array_slicing, devices=devices) +add_function_test(TestApic, "test_array_slicing", test_array_slicing, devices=graph_devices) add_function_test(TestApic, "test_complex_pipeline", test_complex_pipeline, devices=devices) add_function_test(TestApic, "test_internal_allocation", test_internal_allocation, devices=devices) add_function_test(TestApic, "test_multiple_internal_allocations", test_multiple_internal_allocations, devices=devices) -add_function_test(TestApic, "test_graph_execution_unchanged", test_graph_execution_unchanged, devices=devices) +add_function_test(TestApic, "test_graph_execution_unchanged", test_graph_execution_unchanged, devices=graph_devices) add_function_test(TestApic, "test_save_load_with_param_update", test_save_load_with_param_update, devices=devices) add_function_test(TestApic, "test_save_load_memcpy_and_kernel", test_save_load_memcpy_and_kernel, devices=devices) add_function_test( diff --git a/warp/tests/test_sparse.py b/warp/tests/test_sparse.py index 92ef00959d..5ad8328969 100644 --- a/warp/tests/test_sparse.py +++ b/warp/tests/test_sparse.py @@ -864,7 +864,16 @@ def test_bsr_copy_scale(self): add_function_test(TestSparse, "test_bsr_mv_1_3", make_test_bsr_mv((1, 3), wp.float32), devices=devices) add_function_test(TestSparse, "test_bsr_mv_3_3", make_test_bsr_mv((3, 3), wp.float64), devices=devices) -add_function_test(TestSparse, "test_capturability", test_capturability, devices=cuda_graph_devices) +# test_capturability chains data-dependent BSR matrix *creations* inside a single +# graph capture (e.g. `A = bsr_from_triplets(...)` then `A + bsr_copy(A*2.0)`). +# On HIP this faults on non-primary devices due to a ROCm graph-mempool +# limitation: allocations made during capture on a non-primary device are backed +# by device-0 memory and fault under graph replay. Because a GPU memory-access +# fault aborts the whole process (poisoning the rest of the suite), restrict the +# test to the primary device on HIP. CUDA devices are unaffected. +# See HIP_GRAPH_CAPTURE_TODO.md section 2.1a. +capturability_devices = [d for d in cuda_graph_devices if not (d.is_hip and d.ordinal != 0)] +add_function_test(TestSparse, "test_capturability", test_capturability, devices=capturability_devices) add_function_test(TestSparse, "test_bsr_mm_max_new_nnz", test_bsr_mm_max_new_nnz, devices=devices, check_output=False) add_function_test(TestSparse, "test_bsr_alloc", test_bsr_alloc, devices=devices) From 5460235237e86082ff5cbcc32e147a8f6412d6e2 Mon Sep 17 00:00:00 2001 From: Zhihui Du Date: Sun, 19 Jul 2026 18:24:42 -0700 Subject: [PATCH 3/4] fix: enable mempool by default on HIP when graph capture is supported (#16) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit PR #15 enables graph capture on HIP (supports_graph_capture=True). However, the device initialization still disables mempool by default (enable_mempools_at_init=False), which causes hipGraph allocation to fail: Failed to allocate memory during graph capture because memory pools are not enabled on device. Try calling wp.set_mempool_enabled(True). Fix: when a HIP device supports both mempool and graph capture, enable mempool by default. This is safe on ROCm 7.2+ where graph capture has been validated. Users can still disable mempool explicitly via wp.set_mempool_enabled(device, False) or by unsetting warp.config.enable_mempools_at_init before warp.init(). Without this fix, users must call wp.config.enable_mempools_at_init=True before wp.init() as a workaround — which is not discoverable. --- warp/_src/context.py | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/warp/_src/context.py b/warp/_src/context.py index 97afe7cd2c..15159846ee 100644 --- a/warp/_src/context.py +++ b/warp/_src/context.py @@ -3858,6 +3858,12 @@ def __init__(self, runtime, alias, ordinal=-1, is_primary=False, context=None): if warp.config.enable_mempools_at_init: # enable if supported self.is_mempool_enabled = self.is_mempool_supported + elif self.is_hip and self.is_mempool_supported: + # HIP/ROCm: enable mempool by default when graph capture is supported + # (hipGraph requires mempool for in-capture allocations). + # Previously disabled due to hipMemsetAsync unreliability, but + # graph capture is now validated on ROCm 7.2 (Warp PR #15). + self.is_mempool_enabled = True else: # disable by default self.is_mempool_enabled = False From 8ca65dd5f8a444785408ecaa956bac0d2c427d6f Mon Sep 17 00:00:00 2001 From: rtmadduri Date: Mon, 20 Jul 2026 00:11:59 -0400 Subject: [PATCH 4/4] fix hip multi process error --- warp/native/cuda_util.cpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/warp/native/cuda_util.cpp b/warp/native/cuda_util.cpp index 0bcc67532a..4c8e6db312 100644 --- a/warp/native/cuda_util.cpp +++ b/warp/native/cuda_util.cpp @@ -798,7 +798,12 @@ CUresult cuCtxEnablePeerAccess_f(CUcontext peer_ctx, unsigned int flags) #if defined(__HIP_PLATFORM_AMD__) if (!peer_ctx) return CUDA_ERROR_NOT_INITIALIZED; - return hipDeviceEnablePeerAccess(peer_ctx->device, flags); + hipError_t status = hipDeviceEnablePeerAccess(peer_ctx->device, flags); + // hipDeviceEnablePeerAccess is a runtime API that sets the sticky last-error state (e.g. + // hipErrorPeerAccessAlreadyEnabled/704). Clear it so a benign status does not leak into an + // unrelated runtime call (such as a subsequent hipcub call) and get reported there. + ignore_cuda_error(hipGetLastError()); + return status; #else return pfn_cuCtxEnablePeerAccess ? pfn_cuCtxEnablePeerAccess(peer_ctx, flags) : DRIVER_ENTRY_POINT_ERROR; #endif // defined(__HIP_PLATFORM_AMD__)