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..e64cc2e1a0b 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 + ``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. + (`#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..825941ab774 100644 --- a/cuda_core/tests/graph/test_graph_builder.py +++ b/cuda_core/tests/graph/test_graph_builder.py @@ -9,9 +9,11 @@ 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, launch from cuda.core.graph import GraphBuilder, GraphDefinition @@ -694,3 +696,71 @@ 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 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 (see Programming Guide, Programmatic Dependent Launch) — + not kernel overlap. + """ + + def _assert_programmatic_dependency_edge(graph_definition): + """Assert capture of ProgrammaticStreamSerialization produced a programmatic edge. + + 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 + + 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] + # 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 + + 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 +@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+.""" + 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..aec15c447af --- /dev/null +++ b/cuda_core/tests/helpers/pdl_kernels.py @@ -0,0 +1,123 @@ +# 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 numpy as np +import pytest + +import helpers +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 same-stream primary/secondary PDL overlap protocol. + + 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, 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") + + 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 = "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})" + ) + + 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..f575fcde244 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,95 +204,19 @@ def test_to_native_launch_config_pdl(): @skipif_need_cuda_headers -def test_pdl_primary_secondary_overlap_same_stream(): - """Primary + secondary PDL launch on one stream can overlap on Hopper+. +@requires_module(np, "2.2.5", reason="need numpy 2.2.5+ (numpy GH #28632)") +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. """ - 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):