From 8edc7c3763d713affb5f18480f6bf0ae7e7ebe0d Mon Sep 17 00:00:00 2001 From: Jinfeng Date: Thu, 6 Aug 2026 19:45:58 +0000 Subject: [PATCH 1/9] test(cuda.core): verify PDL GraphBuilder capture and overlap Cover GraphBuilder stream-capture for programmatic_stream_serialization: functional launch, programmatic dependency edge mapping, and Hopper+ overlap (xfail if opportunistic), plus 1.2.0 release notes. --- cuda_core/docs/source/release/1.2.0-notes.rst | 12 ++ cuda_core/tests/graph/test_graph_builder.py | 155 +++++++++++++++++- 2 files changed, 166 insertions(+), 1 deletion(-) diff --git a/cuda_core/docs/source/release/1.2.0-notes.rst b/cuda_core/docs/source/release/1.2.0-notes.rst index 4bb81cec759..5e0bbb8f04f 100644 --- a/cuda_core/docs/source/release/1.2.0-notes.rst +++ b/cuda_core/docs/source/release/1.2.0-notes.rst @@ -6,6 +6,18 @@ ``cuda.core`` 1.2.0 Release Notes ================================== +New features +------------ + +- Added the ``programmatic_stream_serialization`` option to + :class:`LaunchConfig`, which sets + ``CU_LAUNCH_ATTRIBUTE_PROGRAMMATIC_STREAM_SERIALIZATION`` so a kernel can + begin executing before the preceding kernel in the same stream has fully + completed (programmatic dependent launch, PDL). Available starting with + devices of compute capability 9.0. + (`#2456 `__, + `#1334 `__) + Fixes and enhancements ---------------------- diff --git a/cuda_core/tests/graph/test_graph_builder.py b/cuda_core/tests/graph/test_graph_builder.py index 6c7c9ef7d64..37a5844d5f7 100644 --- a/cuda_core/tests/graph/test_graph_builder.py +++ b/cuda_core/tests/graph/test_graph_builder.py @@ -7,13 +7,15 @@ import time import weakref +import helpers import numpy as np import pytest from cuda_python_test_helpers.marks import requires_module from helpers.graph_kernels import compile_common_kernels, compile_conditional_kernels from helpers.misc import try_create_condition -from cuda.core import Device, LaunchConfig, LegacyPinnedMemoryResource, launch +from conftest import skipif_need_cuda_headers +from cuda.core import Device, LaunchConfig, LegacyPinnedMemoryResource, Program, ProgramOptions, launch from cuda.core.graph import GraphBuilder, GraphDefinition from cuda.core.graph._graph_builder import ( _capture_callback_with_tail_failure_for_testing, @@ -694,3 +696,154 @@ def test_graph_definition_conditional_body_during_capture_raises(init_cuda): finally: body_gb.end_building() gb.end_building() + + +@requires_module(np, "2.2.5", reason="need numpy 2.2.5+ (numpy GH #28632)") +def test_pdl_launch_graph_capture(init_cuda): + """PDL LaunchConfig is graph-compatible via GraphBuilder stream capture. + + Captures a producer then a secondary launch with + ``programmatic_stream_serialization=True``, instantiates, and launches. + Asserts functional correctness and that capture maps to a programmatic + dependency edge (Programming Guide §4.5.3) — not kernel overlap. + """ + + def _assert_programmatic_dependency_edge(graph_definition): + """Assert capture of ProgrammaticStreamSerialization produced a programmatic edge. + + Programming Guide §4.5.3: stream-capturing a secondary launch with + ``cudaLaunchAttributeProgrammaticStreamSerialization`` maps to + ``CU_GRAPH_DEPENDENCY_TYPE_PROGRAMMATIC`` with + ``CU_GRAPH_KERNEL_NODE_PORT_PROGRAMMATIC``. + """ + from cuda.bindings import driver + + h_graph = graph_definition.handle + err, _, _, _, num_edges = driver.cuGraphGetEdges(h_graph) + assert err == driver.CUresult.CUDA_SUCCESS, err + err, _, _, edge_data, num_edges = driver.cuGraphGetEdges(h_graph, num_edges) + assert err == driver.CUresult.CUDA_SUCCESS, err + assert num_edges == 1, f"expected 1 edge, got {num_edges}" + ed = edge_data[0] + assert ed.type == driver.CUgraphDependencyType.CU_GRAPH_DEPENDENCY_TYPE_PROGRAMMATIC, ed.type + assert ed.from_port == driver.CU_GRAPH_KERNEL_NODE_PORT_PROGRAMMATIC, ed.from_port + + mod = compile_common_kernels() + producer = mod.get_kernel("add_one") + consumer = mod.get_kernel("add_one") + + stream = Device().create_stream() + mr = LegacyPinnedMemoryResource() + buf = mr.allocate(4) + arr = np.from_dlpack(buf).view(np.int32) + arr[0] = 0 + + cfg = LaunchConfig(grid=1, block=1) + pdl = LaunchConfig(grid=1, block=1, programmatic_stream_serialization=True) + + gb = stream.create_graph_builder().begin_building() + launch(gb, cfg, producer, arr.ctypes.data) + launch(gb, pdl, consumer, arr.ctypes.data) + gb.end_building() + _assert_programmatic_dependency_edge(gb.graph_definition) + graph = gb.complete() + + graph.launch(stream) + stream.sync() + assert arr[0] == 2 + + buf.close() + stream.close() + + +@skipif_need_cuda_headers +def test_pdl_primary_secondary_overlap_graph_capture(): + """Primary + secondary PDL via GraphBuilder stream capture can overlap on Hopper+. + + Same kernels / overlap protocol as test_pdl_primary_secondary_overlap_same_stream, + but launches are captured into a CUDA graph (Programming Guide §4.5.3 stream-capture + path). Overlap is opportunistic → miss is xfail. + """ + dev = Device() + if dev.compute_capability < (9, 0): + pytest.skip("Programmatic Dependent Launch requires compute capability >= 9.0") + dev.set_current() + stream = dev.create_stream(options={"nonblocking": True}) + + # clock64 budgets are in GPU cycles; keep the post-trigger window long enough + # for the secondary to boot, but short enough for a unit test. + code = r""" + #include + + extern "C" __global__ void primary_kernel(int* secondary_started, int* overlapped) { + cudaTriggerProgrammaticLaunchCompletion(); + + const long long deadline = clock64() + 100000000LL; // ~50ms @ ~2GHz + if (threadIdx.x == 0 && blockIdx.x == 0) { + while (clock64() < deadline) { + if (atomicAdd(secondary_started, 0) != 0) { + atomicExch(overlapped, 1); + return; + } + __nanosleep(1000); + } + } + } + + extern "C" __global__ void secondary_kernel(int* secondary_started) { + if (threadIdx.x == 0 && blockIdx.x == 0) { + atomicExch(secondary_started, 1); + } + } + """ + + arch = "".join(f"{i}" for i in dev.compute_capability) + pro_opts = ProgramOptions(std="c++17", arch=f"sm_{arch}", include_path=helpers.CUDA_INCLUDE_PATH) + prog = Program(code, code_type="c++", options=pro_opts) + mod = prog.compile("cubin") + primary = mod.get_kernel("primary_kernel") + secondary = mod.get_kernel("secondary_kernel") + + mr = LegacyPinnedMemoryResource() + secondary_started = np.from_dlpack(mr.allocate(4)).view(np.int32) + overlapped = np.from_dlpack(mr.allocate(4)).view(np.int32) + + primary_cfg = LaunchConfig(grid=1, block=1) + secondary_cfg = LaunchConfig(grid=1, block=1, programmatic_stream_serialization=True) + secondary_serial_cfg = LaunchConfig(grid=1, block=1) + + def _run(secondary_launch_cfg: LaunchConfig) -> int: + secondary_started[0] = 0 + overlapped[0] = 0 + gb = stream.create_graph_builder().begin_building() + launch(gb, primary_cfg, primary, secondary_started.ctypes.data, overlapped.ctypes.data) + launch(gb, secondary_launch_cfg, secondary, secondary_started.ctypes.data) + graph = gb.end_building().complete() + try: + graph.launch(stream) + stream.sync() + finally: + graph.close() + gb.close() + return int(overlapped[0]) + + # Without the PDL attribute, same-stream kernels stay serialized even via graph. + assert _run(secondary_serial_cfg) == 0, "Expected no overlap when programmatic_stream_serialization is False" + + # PDL overlap is opportunistic; retry a few times on a quiet GPU. + saw_overlap = False + for _ in range(5): + if _run(secondary_cfg) == 1: + saw_overlap = True + break + + if not saw_overlap: + pytest.xfail( + "PDL (Programmatic Dependent Launch) graph-capture overlap was not observed. " + "If this keeps xfailing in CI, manually re-check on a quiet Hopper+ GPU." + ) + + print( + f"PDL graph-capture overlap verified on {dev.name} compute capability {dev.compute_capability}", + flush=True, + ) From 9d5b2a72d2f071d2928e259a7e66d53a57032968 Mon Sep 17 00:00:00 2001 From: Jinfeng Date: Thu, 6 Aug 2026 19:47:11 +0000 Subject: [PATCH 2/9] docs(cuda.core): use runtime PDL attribute name in 1.2.0 notes Refer to cudaLaunchAttributeProgrammaticStreamSerialization instead of the driver-style CU_LAUNCH_ATTRIBUTE_* spelling. --- cuda_core/docs/source/release/1.2.0-notes.rst | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cuda_core/docs/source/release/1.2.0-notes.rst b/cuda_core/docs/source/release/1.2.0-notes.rst index 5e0bbb8f04f..e64cc2e1a0b 100644 --- a/cuda_core/docs/source/release/1.2.0-notes.rst +++ b/cuda_core/docs/source/release/1.2.0-notes.rst @@ -11,7 +11,7 @@ New features - Added the ``programmatic_stream_serialization`` option to :class:`LaunchConfig`, which sets - ``CU_LAUNCH_ATTRIBUTE_PROGRAMMATIC_STREAM_SERIALIZATION`` so a kernel can + ``cudaLaunchAttributeProgrammaticStreamSerialization`` so a kernel can begin executing before the preceding kernel in the same stream has fully completed (programmatic dependent launch, PDL). Available starting with devices of compute capability 9.0. From 4ca415b2803d61630598c5ecb7d2d88a2e75b0d9 Mon Sep 17 00:00:00 2001 From: Jinfeng Date: Thu, 6 Aug 2026 20:01:20 +0000 Subject: [PATCH 3/9] resolve pre-commit errors --- cuda_core/tests/graph/test_graph_builder.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cuda_core/tests/graph/test_graph_builder.py b/cuda_core/tests/graph/test_graph_builder.py index 37a5844d5f7..5ca44d6fa8b 100644 --- a/cuda_core/tests/graph/test_graph_builder.py +++ b/cuda_core/tests/graph/test_graph_builder.py @@ -10,11 +10,11 @@ import helpers import numpy as np import pytest +from conftest import skipif_need_cuda_headers from cuda_python_test_helpers.marks import requires_module from helpers.graph_kernels import compile_common_kernels, compile_conditional_kernels from helpers.misc import try_create_condition -from conftest import skipif_need_cuda_headers from cuda.core import Device, LaunchConfig, LegacyPinnedMemoryResource, Program, ProgramOptions, launch from cuda.core.graph import GraphBuilder, GraphDefinition from cuda.core.graph._graph_builder import ( From 297d6501b0cd1b850afdccea443477d9ae5b3f0a Mon Sep 17 00:00:00 2001 From: Jinfeng Date: Fri, 7 Aug 2026 18:04:44 +0000 Subject: [PATCH 4/9] test(cuda.core): skip PDL overlap graph capture on NumPy < 2.2.5 The test writes host buffers from np.from_dlpack, which are read-only before NumPy 2.2.5 (GH #28632). --- cuda_core/tests/graph/test_graph_builder.py | 1 + 1 file changed, 1 insertion(+) diff --git a/cuda_core/tests/graph/test_graph_builder.py b/cuda_core/tests/graph/test_graph_builder.py index 5ca44d6fa8b..45f292c0384 100644 --- a/cuda_core/tests/graph/test_graph_builder.py +++ b/cuda_core/tests/graph/test_graph_builder.py @@ -757,6 +757,7 @@ def _assert_programmatic_dependency_edge(graph_definition): @skipif_need_cuda_headers +@requires_module(np, "2.2.5", reason="need numpy 2.2.5+ (numpy GH #28632)") def test_pdl_primary_secondary_overlap_graph_capture(): """Primary + secondary PDL via GraphBuilder stream capture can overlap on Hopper+. From 0c8de79e47119b7b08bba36db0b42c14001183d1 Mon Sep 17 00:00:00 2001 From: Jinfeng Date: Fri, 7 Aug 2026 18:08:51 +0000 Subject: [PATCH 5/9] test(cuda.core): use init_cuda in PDL overlap graph capture test Align with other graph builder tests so context setup and teardown stay consistent. --- cuda_core/tests/graph/test_graph_builder.py | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/cuda_core/tests/graph/test_graph_builder.py b/cuda_core/tests/graph/test_graph_builder.py index 45f292c0384..d03273f5fa1 100644 --- a/cuda_core/tests/graph/test_graph_builder.py +++ b/cuda_core/tests/graph/test_graph_builder.py @@ -758,7 +758,7 @@ def _assert_programmatic_dependency_edge(graph_definition): @skipif_need_cuda_headers @requires_module(np, "2.2.5", reason="need numpy 2.2.5+ (numpy GH #28632)") -def test_pdl_primary_secondary_overlap_graph_capture(): +def test_pdl_primary_secondary_overlap_graph_capture(init_cuda): """Primary + secondary PDL via GraphBuilder stream capture can overlap on Hopper+. Same kernels / overlap protocol as test_pdl_primary_secondary_overlap_same_stream, @@ -768,7 +768,6 @@ def test_pdl_primary_secondary_overlap_graph_capture(): dev = Device() if dev.compute_capability < (9, 0): pytest.skip("Programmatic Dependent Launch requires compute capability >= 9.0") - dev.set_current() stream = dev.create_stream(options={"nonblocking": True}) # clock64 budgets are in GPU cycles; keep the post-trigger window long enough From d7f853d4a481dfc74b61ed4d47bfdae0f4d34845 Mon Sep 17 00:00:00 2001 From: Jinfeng Date: Fri, 7 Aug 2026 18:21:57 +0000 Subject: [PATCH 6/9] test(cuda.core): clarify PDL graph capture doc references Drop fragile Programming Guide section numbers and note Driver vs Runtime enum name equivalence at the edge asserts. --- cuda_core/tests/graph/test_graph_builder.py | 19 +++++++++++-------- 1 file changed, 11 insertions(+), 8 deletions(-) diff --git a/cuda_core/tests/graph/test_graph_builder.py b/cuda_core/tests/graph/test_graph_builder.py index d03273f5fa1..82960750b7a 100644 --- a/cuda_core/tests/graph/test_graph_builder.py +++ b/cuda_core/tests/graph/test_graph_builder.py @@ -702,19 +702,19 @@ def test_graph_definition_conditional_body_during_capture_raises(init_cuda): def test_pdl_launch_graph_capture(init_cuda): """PDL LaunchConfig is graph-compatible via GraphBuilder stream capture. - Captures a producer then a secondary launch with + Captures a first then a secondary launch with ``programmatic_stream_serialization=True``, instantiates, and launches. Asserts functional correctness and that capture maps to a programmatic - dependency edge (Programming Guide §4.5.3) — not kernel overlap. + dependency edge (see Programming Guide, Programmatic Dependent Launch) — + not kernel overlap. """ def _assert_programmatic_dependency_edge(graph_definition): """Assert capture of ProgrammaticStreamSerialization produced a programmatic edge. - Programming Guide §4.5.3: stream-capturing a secondary launch with - ``cudaLaunchAttributeProgrammaticStreamSerialization`` maps to - ``CU_GRAPH_DEPENDENCY_TYPE_PROGRAMMATIC`` with - ``CU_GRAPH_KERNEL_NODE_PORT_PROGRAMMATIC``. + Per Programming Guide (Programmatic Dependent Launch): stream-capturing a + secondary launch with ``cudaLaunchAttributeProgrammaticStreamSerialization`` + maps to a programmatic dependency edge from the programmatic kernel port. """ from cuda.bindings import driver @@ -725,6 +725,9 @@ def _assert_programmatic_dependency_edge(graph_definition): assert err == driver.CUresult.CUDA_SUCCESS, err assert num_edges == 1, f"expected 1 edge, got {num_edges}" ed = edge_data[0] + # Driver (cuda.h) ↔ Runtime / Programming Guide (driver_types.h): + # CU_GRAPH_DEPENDENCY_TYPE_PROGRAMMATIC ↔ cudaGraphDependencyTypeProgrammatic + # CU_GRAPH_KERNEL_NODE_PORT_PROGRAMMATIC ↔ cudaGraphKernelNodePortProgrammatic assert ed.type == driver.CUgraphDependencyType.CU_GRAPH_DEPENDENCY_TYPE_PROGRAMMATIC, ed.type assert ed.from_port == driver.CU_GRAPH_KERNEL_NODE_PORT_PROGRAMMATIC, ed.from_port @@ -762,8 +765,8 @@ def test_pdl_primary_secondary_overlap_graph_capture(init_cuda): """Primary + secondary PDL via GraphBuilder stream capture can overlap on Hopper+. Same kernels / overlap protocol as test_pdl_primary_secondary_overlap_same_stream, - but launches are captured into a CUDA graph (Programming Guide §4.5.3 stream-capture - path). Overlap is opportunistic → miss is xfail. + but launches are captured into a CUDA graph (see CUDA Programming Guide, + Programmatic Dependent Launch). Overlap is opportunistic → miss is xfail. """ dev = Device() if dev.compute_capability < (9, 0): From 5f8405dbf8f456fb95924e8a9f08599af784fec5 Mon Sep 17 00:00:00 2001 From: Jinfeng Date: Fri, 7 Aug 2026 18:39:00 +0000 Subject: [PATCH 7/9] test(cuda.core): share PDL overlap protocol via helper runner Extract kernels and the same-stream / graph-capture overlap check into run_pdl_overlap_check so launcher and GraphBuilder tests stay in sync. --- cuda_core/tests/graph/test_graph_builder.py | 87 +------------- cuda_core/tests/helpers/pdl_kernels.py | 122 ++++++++++++++++++++ cuda_core/tests/test_launcher.py | 83 +------------ 3 files changed, 129 insertions(+), 163 deletions(-) create mode 100644 cuda_core/tests/helpers/pdl_kernels.py diff --git a/cuda_core/tests/graph/test_graph_builder.py b/cuda_core/tests/graph/test_graph_builder.py index 82960750b7a..3bc00b48e56 100644 --- a/cuda_core/tests/graph/test_graph_builder.py +++ b/cuda_core/tests/graph/test_graph_builder.py @@ -7,15 +7,15 @@ import time import weakref -import helpers import numpy as np import pytest from conftest import skipif_need_cuda_headers from cuda_python_test_helpers.marks import requires_module from helpers.graph_kernels import compile_common_kernels, compile_conditional_kernels from helpers.misc import try_create_condition +from helpers.pdl_kernels import run_pdl_overlap_check -from cuda.core import Device, LaunchConfig, LegacyPinnedMemoryResource, Program, ProgramOptions, launch +from cuda.core import Device, LaunchConfig, LegacyPinnedMemoryResource, launch from cuda.core.graph import GraphBuilder, GraphDefinition from cuda.core.graph._graph_builder import ( _capture_callback_with_tail_failure_for_testing, @@ -768,85 +768,4 @@ def test_pdl_primary_secondary_overlap_graph_capture(init_cuda): but launches are captured into a CUDA graph (see CUDA Programming Guide, Programmatic Dependent Launch). Overlap is opportunistic → miss is xfail. """ - dev = Device() - if dev.compute_capability < (9, 0): - pytest.skip("Programmatic Dependent Launch requires compute capability >= 9.0") - stream = dev.create_stream(options={"nonblocking": True}) - - # clock64 budgets are in GPU cycles; keep the post-trigger window long enough - # for the secondary to boot, but short enough for a unit test. - code = r""" - #include - - extern "C" __global__ void primary_kernel(int* secondary_started, int* overlapped) { - cudaTriggerProgrammaticLaunchCompletion(); - - const long long deadline = clock64() + 100000000LL; // ~50ms @ ~2GHz - if (threadIdx.x == 0 && blockIdx.x == 0) { - while (clock64() < deadline) { - if (atomicAdd(secondary_started, 0) != 0) { - atomicExch(overlapped, 1); - return; - } - __nanosleep(1000); - } - } - } - - extern "C" __global__ void secondary_kernel(int* secondary_started) { - if (threadIdx.x == 0 && blockIdx.x == 0) { - atomicExch(secondary_started, 1); - } - } - """ - - arch = "".join(f"{i}" for i in dev.compute_capability) - pro_opts = ProgramOptions(std="c++17", arch=f"sm_{arch}", include_path=helpers.CUDA_INCLUDE_PATH) - prog = Program(code, code_type="c++", options=pro_opts) - mod = prog.compile("cubin") - primary = mod.get_kernel("primary_kernel") - secondary = mod.get_kernel("secondary_kernel") - - mr = LegacyPinnedMemoryResource() - secondary_started = np.from_dlpack(mr.allocate(4)).view(np.int32) - overlapped = np.from_dlpack(mr.allocate(4)).view(np.int32) - - primary_cfg = LaunchConfig(grid=1, block=1) - secondary_cfg = LaunchConfig(grid=1, block=1, programmatic_stream_serialization=True) - secondary_serial_cfg = LaunchConfig(grid=1, block=1) - - def _run(secondary_launch_cfg: LaunchConfig) -> int: - secondary_started[0] = 0 - overlapped[0] = 0 - gb = stream.create_graph_builder().begin_building() - launch(gb, primary_cfg, primary, secondary_started.ctypes.data, overlapped.ctypes.data) - launch(gb, secondary_launch_cfg, secondary, secondary_started.ctypes.data) - graph = gb.end_building().complete() - try: - graph.launch(stream) - stream.sync() - finally: - graph.close() - gb.close() - return int(overlapped[0]) - - # Without the PDL attribute, same-stream kernels stay serialized even via graph. - assert _run(secondary_serial_cfg) == 0, "Expected no overlap when programmatic_stream_serialization is False" - - # PDL overlap is opportunistic; retry a few times on a quiet GPU. - saw_overlap = False - for _ in range(5): - if _run(secondary_cfg) == 1: - saw_overlap = True - break - - if not saw_overlap: - pytest.xfail( - "PDL (Programmatic Dependent Launch) graph-capture overlap was not observed. " - "If this keeps xfailing in CI, manually re-check on a quiet Hopper+ GPU." - ) - - print( - f"PDL graph-capture overlap verified on {dev.name} compute capability {dev.compute_capability}", - flush=True, - ) + run_pdl_overlap_check(Device(), via_graph=True) diff --git a/cuda_core/tests/helpers/pdl_kernels.py b/cuda_core/tests/helpers/pdl_kernels.py new file mode 100644 index 00000000000..0d08d4112a0 --- /dev/null +++ b/cuda_core/tests/helpers/pdl_kernels.py @@ -0,0 +1,122 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +"""Shared helpers for Programmatic Dependent Launch overlap tests.""" + +import helpers +import numpy as np +import pytest + +from cuda.core import LaunchConfig, LegacyPinnedMemoryResource, Program, ProgramOptions, launch + + +def compile_pdl_overlap_kernels(device): + """Compile primary/secondary kernels used to detect PDL same-stream overlap. + + The primary triggers programmatic launch completion then spins briefly looking + for a flag written by the secondary. Seeing that flag proves both grids were + resident at once. clock64 budgets are in GPU cycles: long enough for the + secondary to boot, short enough for a unit test. + + Returns: + (primary_kernel, secondary_kernel) + """ + code = r""" + #include + + extern "C" __global__ void primary_kernel(int* secondary_started, int* overlapped) { + cudaTriggerProgrammaticLaunchCompletion(); + + const long long deadline = clock64() + 100000000LL; // ~50ms @ ~2GHz + if (threadIdx.x == 0 && blockIdx.x == 0) { + while (clock64() < deadline) { + if (atomicAdd(secondary_started, 0) != 0) { + atomicExch(overlapped, 1); + return; + } + __nanosleep(1000); + } + } + } + + extern "C" __global__ void secondary_kernel(int* secondary_started) { + if (threadIdx.x == 0 && blockIdx.x == 0) { + atomicExch(secondary_started, 1); + } + } + """ + arch = "".join(f"{i}" for i in device.compute_capability) + pro_opts = ProgramOptions(std="c++17", arch=f"sm_{arch}", include_path=helpers.CUDA_INCLUDE_PATH) + prog = Program(code, code_type="c++", options=pro_opts) + mod = prog.compile("cubin") + return mod.get_kernel("primary_kernel"), mod.get_kernel("secondary_kernel") + + +def run_pdl_overlap_check(device, *, via_graph: bool = False): + """Run the shared primary/secondary PDL overlap protocol. + + Asserts no overlap without ``programmatic_stream_serialization``, then retries + a few times with it enabled. Overlap is opportunistic → miss is xfail. + + Args: + device: Current CUDA device (compute capability >= 9.0 required). + via_graph: If True, capture launches into a CUDA graph and launch the + graph; otherwise launch kernels directly on the stream. + """ + if device.compute_capability < (9, 0): + pytest.skip("Programmatic Dependent Launch requires compute capability >= 9.0") + + stream = device.create_stream(options={"nonblocking": True}) + primary, secondary = compile_pdl_overlap_kernels(device) + + mr = LegacyPinnedMemoryResource() + secondary_started = np.from_dlpack(mr.allocate(4)).view(np.int32) + overlapped = np.from_dlpack(mr.allocate(4)).view(np.int32) + + primary_cfg = LaunchConfig(grid=1, block=1) + secondary_cfg = LaunchConfig(grid=1, block=1, programmatic_stream_serialization=True) + secondary_serial_cfg = LaunchConfig(grid=1, block=1) + + def _run(secondary_launch_cfg: LaunchConfig) -> int: + secondary_started[0] = 0 + overlapped[0] = 0 + if via_graph: + gb = stream.create_graph_builder().begin_building() + launch(gb, primary_cfg, primary, secondary_started.ctypes.data, overlapped.ctypes.data) + launch(gb, secondary_launch_cfg, secondary, secondary_started.ctypes.data) + graph = gb.end_building().complete() + try: + graph.launch(stream) + stream.sync() + finally: + graph.close() + gb.close() + else: + launch(stream, primary_cfg, primary, secondary_started.ctypes.data, overlapped.ctypes.data) + launch(stream, secondary_launch_cfg, secondary, secondary_started.ctypes.data) + stream.sync() + return int(overlapped[0]) + + path = "graph-capture" if via_graph else "same-stream" + assert _run(secondary_serial_cfg) == 0, ( + f"Expected no overlap when programmatic_stream_serialization is False ({path})" + ) + + saw_overlap = False + for _ in range(5): + if _run(secondary_cfg) == 1: + saw_overlap = True + break + + if not saw_overlap: + # Overlap is never guaranteed by the driver, so a miss is reported as an + # expected failure rather than turning a busy GPU into a red CI run. + pytest.xfail( + f"PDL (Programmatic Dependent Launch) {path} overlap was not observed. " + "If this keeps xfailing in CI, manually re-check on a quiet Hopper+ GPU." + ) + + print( + f"PDL {path} overlap verified on {device.name} compute capability {device.compute_capability}", + flush=True, + ) diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py index e5cf05b435d..8e4b85e8a49 100644 --- a/cuda_core/tests/test_launcher.py +++ b/cuda_core/tests/test_launcher.py @@ -6,6 +6,7 @@ import helpers from cuda_python_test_helpers.marks import requires_module from helpers.misc import StreamWrapper +from helpers.pdl_kernels import run_pdl_overlap_check try: import cupy as cp @@ -203,7 +204,8 @@ def test_to_native_launch_config_pdl(): @skipif_need_cuda_headers -def test_pdl_primary_secondary_overlap_same_stream(): +@requires_module(np, "2.2.5", reason="need numpy 2.2.5+ (numpy GH #28632)") +def test_pdl_primary_secondary_overlap_same_stream(init_cuda): """Primary + secondary PDL launch on one stream can overlap on Hopper+. Secondary is launched with ``programmatic_stream_serialization=True``. After @@ -214,84 +216,7 @@ def test_pdl_primary_secondary_overlap_same_stream(): Note concurrency is opportunistic, so a missing overlap execution is reported as an expected failure. """ - dev = Device() - if dev.compute_capability < (9, 0): - pytest.skip("Programmatic Dependent Launch requires compute capability >= 9.0") - dev.set_current() - stream = dev.create_stream(options={"nonblocking": True}) - - # clock64 budgets are in GPU cycles; keep the post-trigger window long enough - # for the secondary to boot, but short enough for a unit test. - code = r""" - #include - - extern "C" __global__ void primary_kernel(int* secondary_started, int* overlapped) { - cudaTriggerProgrammaticLaunchCompletion(); - - const long long deadline = clock64() + 100000000LL; // ~50ms @ ~2GHz - if (threadIdx.x == 0 && blockIdx.x == 0) { - while (clock64() < deadline) { - if (atomicAdd(secondary_started, 0) != 0) { - atomicExch(overlapped, 1); - return; - } - __nanosleep(1000); - } - } - } - - extern "C" __global__ void secondary_kernel(int* secondary_started) { - if (threadIdx.x == 0 && blockIdx.x == 0) { - atomicExch(secondary_started, 1); - } - } - """ - - arch = "".join(f"{i}" for i in dev.compute_capability) - pro_opts = ProgramOptions(std="c++17", arch=f"sm_{arch}", include_path=helpers.CUDA_INCLUDE_PATH) - prog = Program(code, code_type="c++", options=pro_opts) - mod = prog.compile("cubin") - primary = mod.get_kernel("primary_kernel") - secondary = mod.get_kernel("secondary_kernel") - - mr = LegacyPinnedMemoryResource() - secondary_started = np.from_dlpack(mr.allocate(4)).view(np.int32) - overlapped = np.from_dlpack(mr.allocate(4)).view(np.int32) - - primary_cfg = LaunchConfig(grid=1, block=1) - secondary_cfg = LaunchConfig(grid=1, block=1, programmatic_stream_serialization=True) - secondary_serial_cfg = LaunchConfig(grid=1, block=1) - - def _run(secondary_launch_cfg: LaunchConfig) -> int: - secondary_started[0] = 0 - overlapped[0] = 0 - launch(stream, primary_cfg, primary, secondary_started.ctypes.data, overlapped.ctypes.data) - launch(stream, secondary_launch_cfg, secondary, secondary_started.ctypes.data) - stream.sync() - return int(overlapped[0]) - - # Without the PDL attribute, same-stream kernels stay serialized. - assert _run(secondary_serial_cfg) == 0, "Expected no overlap when programmatic_stream_serialization is False" - - # PDL overlap is opportunistic; retry a few times on a quiet GPU. - saw_overlap = False - for _ in range(5): - if _run(secondary_cfg) == 1: - saw_overlap = True - break - - if not saw_overlap: - # Overlap is never guaranteed by the driver, so a miss is reported as an - # expected failure rather than turning a busy GPU into a red CI run. - pytest.xfail( - "PDL (Programmatic Dependent Launch) overlap was not observed. " - "If this keeps xfailing in CI, manually re-check on a quiet Hopper+ GPU." - ) - - print( - f"PDL (Programmatic Dependent Launch) overlap verified on {dev.name} compute capability {dev.compute_capability}", - flush=True, - ) + run_pdl_overlap_check(Device(), via_graph=False) def test_launch_config_cluster_accepts_hopper_cc(monkeypatch): From 391a66ec9254a316cc35a208944a5844d6482b1d Mon Sep 17 00:00:00 2001 From: Jinfeng Date: Fri, 7 Aug 2026 18:51:41 +0000 Subject: [PATCH 8/9] test(cuda.core): rename PDL overlap tests to emphasize same-stream Both direct and graph-capture paths are same-stream; update names and docs accordingly. --- cuda_core/tests/graph/test_graph_builder.py | 8 ++------ cuda_core/tests/helpers/pdl_kernels.py | 13 +++++++------ cuda_core/tests/test_launcher.py | 6 +++--- 3 files changed, 12 insertions(+), 15 deletions(-) diff --git a/cuda_core/tests/graph/test_graph_builder.py b/cuda_core/tests/graph/test_graph_builder.py index 3bc00b48e56..1591ffca318 100644 --- a/cuda_core/tests/graph/test_graph_builder.py +++ b/cuda_core/tests/graph/test_graph_builder.py @@ -761,11 +761,7 @@ def _assert_programmatic_dependency_edge(graph_definition): @skipif_need_cuda_headers @requires_module(np, "2.2.5", reason="need numpy 2.2.5+ (numpy GH #28632)") -def test_pdl_primary_secondary_overlap_graph_capture(init_cuda): - """Primary + secondary PDL via GraphBuilder stream capture can overlap on Hopper+. - - Same kernels / overlap protocol as test_pdl_primary_secondary_overlap_same_stream, - but launches are captured into a CUDA graph (see CUDA Programming Guide, - Programmatic Dependent Launch). Overlap is opportunistic → miss is xfail. +def test_pdl_same_stream_primary_secondary_overlap_via_graph(init_cuda): + """Same-stream PDL overlap via GraphBuilder stream capture on Hopper+. """ run_pdl_overlap_check(Device(), via_graph=True) diff --git a/cuda_core/tests/helpers/pdl_kernels.py b/cuda_core/tests/helpers/pdl_kernels.py index 0d08d4112a0..e9a2edf8eb2 100644 --- a/cuda_core/tests/helpers/pdl_kernels.py +++ b/cuda_core/tests/helpers/pdl_kernels.py @@ -53,15 +53,16 @@ def compile_pdl_overlap_kernels(device): def run_pdl_overlap_check(device, *, via_graph: bool = False): - """Run the shared primary/secondary PDL overlap protocol. + """Run the shared same-stream primary/secondary PDL overlap protocol. - Asserts no overlap without ``programmatic_stream_serialization``, then retries - a few times with it enabled. Overlap is opportunistic → miss is xfail. + Both paths launch primary then secondary on one stream. Asserts no overlap + without ``programmatic_stream_serialization``, then retries a few times with + it enabled. Overlap is opportunistic → miss is xfail. Args: device: Current CUDA device (compute capability >= 9.0 required). - via_graph: If True, capture launches into a CUDA graph and launch the - graph; otherwise launch kernels directly on the stream. + via_graph: If True, stream-capture the same-stream launches into a CUDA + graph and launch that graph; otherwise launch kernels directly. """ if device.compute_capability < (9, 0): pytest.skip("Programmatic Dependent Launch requires compute capability >= 9.0") @@ -97,7 +98,7 @@ def _run(secondary_launch_cfg: LaunchConfig) -> int: stream.sync() return int(overlapped[0]) - path = "graph-capture" if via_graph else "same-stream" + path = "same-stream (graph)" if via_graph else "same-stream" assert _run(secondary_serial_cfg) == 0, ( f"Expected no overlap when programmatic_stream_serialization is False ({path})" ) diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py index 8e4b85e8a49..f575fcde244 100644 --- a/cuda_core/tests/test_launcher.py +++ b/cuda_core/tests/test_launcher.py @@ -205,13 +205,13 @@ def test_to_native_launch_config_pdl(): @skipif_need_cuda_headers @requires_module(np, "2.2.5", reason="need numpy 2.2.5+ (numpy GH #28632)") -def test_pdl_primary_secondary_overlap_same_stream(init_cuda): - """Primary + secondary PDL launch on one stream can overlap on Hopper+. +def test_pdl_same_stream_primary_secondary_overlap(init_cuda): + """Same-stream primary + secondary PDL launch can overlap on Hopper+. Secondary is launched with ``programmatic_stream_serialization=True``. After the primary triggers completion, it spins until it observes a flag written by the secondary's independent preamble — proving both grids were resident at - once. Without PDL, the secondary cannot start until the primary exits. + once. Without PDL, same-stream kernels stay serialized. Note concurrency is opportunistic, so a missing overlap execution is reported as an expected failure. From fdfdf8b5f9b94f7b208a9577888873c0f7426fe5 Mon Sep 17 00:00:00 2001 From: Jinfeng Date: Fri, 7 Aug 2026 18:53:41 +0000 Subject: [PATCH 9/9] test(cuda.core): tidy PDL overlap helper import order and docstring --- cuda_core/tests/graph/test_graph_builder.py | 3 +-- cuda_core/tests/helpers/pdl_kernels.py | 2 +- 2 files changed, 2 insertions(+), 3 deletions(-) diff --git a/cuda_core/tests/graph/test_graph_builder.py b/cuda_core/tests/graph/test_graph_builder.py index 1591ffca318..825941ab774 100644 --- a/cuda_core/tests/graph/test_graph_builder.py +++ b/cuda_core/tests/graph/test_graph_builder.py @@ -762,6 +762,5 @@ def _assert_programmatic_dependency_edge(graph_definition): @skipif_need_cuda_headers @requires_module(np, "2.2.5", reason="need numpy 2.2.5+ (numpy GH #28632)") def test_pdl_same_stream_primary_secondary_overlap_via_graph(init_cuda): - """Same-stream PDL overlap via GraphBuilder stream capture on Hopper+. - """ + """Same-stream PDL overlap via GraphBuilder stream capture on Hopper+.""" run_pdl_overlap_check(Device(), via_graph=True) diff --git a/cuda_core/tests/helpers/pdl_kernels.py b/cuda_core/tests/helpers/pdl_kernels.py index e9a2edf8eb2..aec15c447af 100644 --- a/cuda_core/tests/helpers/pdl_kernels.py +++ b/cuda_core/tests/helpers/pdl_kernels.py @@ -3,10 +3,10 @@ """Shared helpers for Programmatic Dependent Launch overlap tests.""" -import helpers import numpy as np import pytest +import helpers from cuda.core import LaunchConfig, LegacyPinnedMemoryResource, Program, ProgramOptions, launch