From ce87db27d6f8b4bc1103ea61343c38809a0e63e9 Mon Sep 17 00:00:00 2001 From: Omar Atie Date: Fri, 14 Aug 2026 16:20:17 -0700 Subject: [PATCH 01/17] feat(cuda.core): add synchronization_policy to LaunchConfig Expose CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY through LaunchConfig so cuda.core users can set per-launch CPU wait policies without dropping to cuda.bindings.driver. Adds SynchronizationPolicyType and tests for native attribute mapping and real kernel launches. Closes #2628. Co-authored-by: Cursor --- cuda_core/cuda/core/_launch_config.pxd | 1 + cuda_core/cuda/core/_launch_config.pyi | 11 ++- cuda_core/cuda/core/_launch_config.pyx | 42 +++++++++ cuda_core/cuda/core/typing.py | 19 +++++ cuda_core/tests/test_launcher.py | 104 ++++++++++++++++++++++- cuda_core/tests/test_object_protocols.py | 2 +- 6 files changed, 175 insertions(+), 4 deletions(-) diff --git a/cuda_core/cuda/core/_launch_config.pxd b/cuda_core/cuda/core/_launch_config.pxd index 892a73f8efc..a4dc4b00c53 100644 --- a/cuda_core/cuda/core/_launch_config.pxd +++ b/cuda_core/cuda/core/_launch_config.pxd @@ -16,6 +16,7 @@ cdef class LaunchConfig: public int shmem_size public bint is_cooperative public bint programmatic_stream_serialization + public object synchronization_policy vector[cydriver.CUlaunchAttribute] _attrs object __weakref__ diff --git a/cuda_core/cuda/core/_launch_config.pyi b/cuda_core/cuda/core/_launch_config.pyi index 47187fb03d6..1082ffa1825 100644 --- a/cuda_core/cuda/core/_launch_config.pyi +++ b/cuda_core/cuda/core/_launch_config.pyi @@ -4,6 +4,8 @@ from __future__ import annotations from typing import Any +from cuda.core.typing import SynchronizationPolicyType + class LaunchConfig: """Customizable launch options. @@ -39,9 +41,12 @@ class LaunchConfig: Whether to allow programmatic stream serialization (PDL). When True, the kernel may overlap with a previous kernel in the same stream that signals completion via programmatic means. + synchronization_policy : SynchronizationPolicyType | None, optional + CPU wait policy applied when synchronizing the launch stream after this + kernel. Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY``. """ - def __init__(self, grid: int | tuple[int, ...] | None=None, cluster: int | tuple[int, ...] | None=None, block: int | tuple[int, ...] | None=None, shmem_size: int | None=None, is_cooperative: bool=False, programmatic_stream_serialization: bool=False) -> None: + def __init__(self, grid: int | tuple[int, ...] | None=None, cluster: int | tuple[int, ...] | None=None, block: int | tuple[int, ...] | None=None, shmem_size: int | None=None, is_cooperative: bool=False, programmatic_stream_serialization: bool=False, synchronization_policy: SynchronizationPolicyType | int | None=None) -> None: """Initialize LaunchConfig with validation. Parameters @@ -58,6 +63,8 @@ class LaunchConfig: Whether to launch as cooperative kernel (default: False) programmatic_stream_serialization : bool, optional Whether to allow programmatic stream serialization / PDL (default: False) + synchronization_policy : SynchronizationPolicyType | None, optional + CPU wait policy for synchronizing the launch stream (default: None) """ def _identity(self) -> tuple[Any, ...]: @@ -71,7 +78,7 @@ class LaunchConfig: def __hash__(self) -> int: ... -_LAUNCH_CONFIG_ATTRS = ('grid', 'cluster', 'block', 'shmem_size', 'is_cooperative', 'programmatic_stream_serialization') +_LAUNCH_CONFIG_ATTRS = ('grid', 'cluster', 'block', 'shmem_size', 'is_cooperative', 'programmatic_stream_serialization', 'synchronization_policy') __all__ = ['LaunchConfig'] def _to_native_launch_config(config: LaunchConfig) -> object: diff --git a/cuda_core/cuda/core/_launch_config.pyx b/cuda_core/cuda/core/_launch_config.pyx index adbf9a16c5d..991bd3bbab7 100644 --- a/cuda_core/cuda/core/_launch_config.pyx +++ b/cuda_core/cuda/core/_launch_config.pyx @@ -20,8 +20,32 @@ _LAUNCH_CONFIG_ATTRS = ( 'shmem_size', 'is_cooperative', 'programmatic_stream_serialization', + 'synchronization_policy', ) + +cdef object _validate_synchronization_policy(object policy): + from cuda.core.typing import SynchronizationPolicyType + + if policy is None: + return None + if isinstance(policy, SynchronizationPolicyType): + return policy + try: + value = int(policy) + except (TypeError, ValueError) as exc: + raise TypeError( + "LaunchConfig.synchronization_policy must be a SynchronizationPolicyType, " + f"cuda.bindings.driver.CUsynchronizationPolicy, or int; got {type(policy).__name__}" + ) from exc + try: + return SynchronizationPolicyType(value) + except ValueError as exc: + raise ValueError( + f"LaunchConfig.synchronization_policy must be one of " + f"{[member.name for member in SynchronizationPolicyType]}; got {policy!r}" + ) from exc + __all__ = ['LaunchConfig'] @@ -59,6 +83,9 @@ cdef class LaunchConfig: Whether to allow programmatic stream serialization (PDL). When True, the kernel may overlap with a previous kernel in the same stream that signals completion via programmatic means. + synchronization_policy : SynchronizationPolicyType | None, optional + CPU wait policy applied when synchronizing the launch stream after this + kernel. Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY``. """ # TODO: expand LaunchConfig to include other attributes @@ -72,6 +99,7 @@ cdef class LaunchConfig: shmem_size: int | None = None, is_cooperative: bool = False, programmatic_stream_serialization: bool = False, + synchronization_policy: object = None, ) -> None: """Initialize LaunchConfig with validation. @@ -89,6 +117,8 @@ cdef class LaunchConfig: Whether to launch as cooperative kernel (default: False) programmatic_stream_serialization : bool, optional Whether to allow programmatic stream serialization / PDL (default: False) + synchronization_policy : SynchronizationPolicyType | None, optional + CPU wait policy for synchronizing the launch stream (default: None) """ # Convert and validate grid and block dimensions self.grid = cast_to_3_tuple("LaunchConfig.grid", grid) @@ -116,6 +146,7 @@ cdef class LaunchConfig: self.is_cooperative = is_cooperative self.programmatic_stream_serialization = programmatic_stream_serialization + self.synchronization_policy = _validate_synchronization_policy(synchronization_policy) if self.is_cooperative and not Device().properties.cooperative_launch: raise CUDAError("cooperative kernels are not supported on this device") @@ -169,6 +200,11 @@ cdef class LaunchConfig: attr.value.programmaticStreamSerializationAllowed = 1 self._attrs.push_back(attr) + if self.synchronization_policy is not None: + attr.id = cydriver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY + attr.value.syncPolicy = int(self.synchronization_policy) + self._attrs.push_back(attr) + drv_cfg.numAttrs = self._attrs.size() drv_cfg.attrs = self._attrs.data() @@ -230,6 +266,12 @@ cpdef object _to_native_launch_config(LaunchConfig config): attr.value.programmaticStreamSerializationAllowed = 1 attrs.append(attr) + if config.synchronization_policy is not None: + attr = driver.CUlaunchAttribute() + attr.id = driver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY + attr.value.syncPolicy = int(config.synchronization_policy) + attrs.append(attr) + drv_cfg.numAttrs = len(attrs) drv_cfg.attrs = attrs diff --git a/cuda_core/cuda/core/typing.py b/cuda_core/cuda/core/typing.py index 1bf9bb7c0d2..e6a4acaba0e 100644 --- a/cuda_core/cuda/core/typing.py +++ b/cuda_core/cuda/core/typing.py @@ -5,6 +5,7 @@ """Public type aliases, protocols, and enumerations used in cuda.core API signatures.""" import sys +from enum import IntEnum from typing import TYPE_CHECKING from typing import Literal as _Literal from typing import TypeAlias as _TypeAlias @@ -47,6 +48,7 @@ class StrEnum(str, Enum): "ProcessStateType", "ReadModeType", "SourceCodeType", + "SynchronizationPolicyType", "VirtualMemoryAccessType", "VirtualMemoryAllocationType", "VirtualMemoryGranularityType", @@ -124,6 +126,23 @@ class PCHStatusType(StrEnum): FAILED = "failed" +class SynchronizationPolicyType(IntEnum): + """CPU wait policy for host-side stream synchronization after a launch. + + Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY``. + + * ``AUTO`` — inherit the stream's synchronization policy. + * ``SPIN`` — busy-wait on the CPU (lowest latency). + * ``YIELD`` — yield the CPU while waiting. + * ``BLOCKING_SYNC`` — block in the OS scheduler while waiting. + """ + + AUTO = 0 + SPIN = 1 + YIELD = 2 + BLOCKING_SYNC = 3 + + class GraphConditionalType(StrEnum): """Conditional node flavor for :class:`~cuda.core.graph.GraphBuilder`. diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py index e5cf05b435d..e7ac4d62d0c 100644 --- a/cuda_core/tests/test_launcher.py +++ b/cuda_core/tests/test_launcher.py @@ -13,6 +13,7 @@ cp = None import numpy as np import pytest +from cuda.bindings import driver from conftest import skipif_need_cuda_headers from cuda.core import ( @@ -26,7 +27,7 @@ ) from cuda.core._memory._legacy import _SynchronousMemoryResource from cuda.core._utils.cuda_utils import CUDAError -from cuda.core.typing import ObjectCodeFormatType, SourceCodeType +from cuda.core.typing import ObjectCodeFormatType, SourceCodeType, SynchronizationPolicyType def test_launch_config_init(init_cuda): @@ -202,6 +203,107 @@ def test_to_native_launch_config_pdl(): ) +@pytest.mark.parametrize( + ("policy", "expected_value"), + [ + (SynchronizationPolicyType.AUTO, driver.CUsynchronizationPolicy.CU_SYNC_POLICY_AUTO), + (SynchronizationPolicyType.SPIN, driver.CUsynchronizationPolicy.CU_SYNC_POLICY_SPIN), + (SynchronizationPolicyType.YIELD, driver.CUsynchronizationPolicy.CU_SYNC_POLICY_YIELD), + ( + SynchronizationPolicyType.BLOCKING_SYNC, + driver.CUsynchronizationPolicy.CU_SYNC_POLICY_BLOCKING_SYNC, + ), + (driver.CUsynchronizationPolicy.CU_SYNC_POLICY_SPIN, driver.CUsynchronizationPolicy.CU_SYNC_POLICY_SPIN), + ], +) +def test_to_native_launch_config_synchronization_policy(policy, expected_value): + """LaunchConfig.synchronization_policy maps to CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY.""" + from cuda.core._launch_config import _to_native_launch_config + + config = LaunchConfig(grid=1, block=1, synchronization_policy=policy) + assert config.synchronization_policy is SynchronizationPolicyType(int(expected_value)) + + native = _to_native_launch_config(config) + assert native.numAttrs == 1 + attr = native.attrs[0] + assert attr.id == driver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY + assert attr.value.syncPolicy == expected_value + + +def test_launch_config_synchronization_policy_default(): + config = LaunchConfig(grid=1, block=1) + assert config.synchronization_policy is None + + from cuda.core._launch_config import _to_native_launch_config + + native = _to_native_launch_config(config) + assert native.numAttrs == 0 + + +@pytest.mark.parametrize("invalid_policy", ["spin", -1, 99]) +def test_launch_config_synchronization_policy_invalid(invalid_policy): + with pytest.raises((TypeError, ValueError)): + LaunchConfig(grid=1, block=1, synchronization_policy=invalid_policy) + + +def test_to_native_launch_config_synchronization_policy_with_cooperative(monkeypatch): + """synchronization_policy can be combined with other launch attributes.""" + from cuda.core import _launch_config as _lc_mod + from cuda.core._launch_config import _to_native_launch_config + + class _FakeProps: + cooperative_launch = True + + class _FakeDev: + properties = _FakeProps() + + monkeypatch.setattr(_lc_mod, "Device", lambda: _FakeDev()) + + config = LaunchConfig( + grid=1, + block=1, + is_cooperative=True, + synchronization_policy=SynchronizationPolicyType.SPIN, + ) + native = _to_native_launch_config(config) + assert native.numAttrs == 2 + attr_ids = {attr.id for attr in native.attrs} + assert driver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_COOPERATIVE in attr_ids + assert driver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY in attr_ids + sync_attrs = [ + attr + for attr in native.attrs + if attr.id == driver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY + ] + assert len(sync_attrs) == 1 + assert sync_attrs[0].value.syncPolicy == driver.CUsynchronizationPolicy.CU_SYNC_POLICY_SPIN + + +@pytest.mark.parametrize( + "policy", + [ + SynchronizationPolicyType.AUTO, + SynchronizationPolicyType.SPIN, + SynchronizationPolicyType.YIELD, + SynchronizationPolicyType.BLOCKING_SYNC, + ], +) +def test_launch_with_synchronization_policy(init_cuda, policy): + """Driver accepts per-launch synchronization policies on a real kernel launch.""" + dev = Device() + dev.set_current() + stream = dev.create_stream() + + code = 'extern "C" __global__ void noop() {}' + program = Program(code, SourceCodeType.CXX) + mod = program.compile(ObjectCodeFormatType.CUBIN) + ker = mod.get_kernel("noop") + + config = LaunchConfig(grid=1, block=1, synchronization_policy=policy) + launch(stream, config, ker) + stream.sync() + + @skipif_need_cuda_headers def test_pdl_primary_secondary_overlap_same_stream(): """Primary + secondary PDL launch on one stream can overlap on Hopper+. diff --git a/cuda_core/tests/test_object_protocols.py b/cuda_core/tests/test_object_protocols.py index baf790abea8..9a54c38cfd0 100644 --- a/cuda_core/tests/test_object_protocols.py +++ b/cuda_core/tests/test_object_protocols.py @@ -685,7 +685,7 @@ def sample_switch_node_alt(sample_graphdef): "sample_launch_config", r"LaunchConfig\(grid=\(\d+, \d+, \d+\), cluster=.+, block=\(\d+, \d+, \d+\), " r"shmem_size=\d+, is_cooperative=(?:True|False), " - r"programmatic_stream_serialization=(?:True|False)\)", + r"programmatic_stream_serialization=(?:True|False), synchronization_policy=(?:None|SynchronizationPolicyType\.\w+)\)", ), ("sample_kernel", r""), # ObjectCode variations (by code_type) From bbff8d224e56270eeb8f19c2ffccf72a710c9346 Mon Sep 17 00:00:00 2001 From: Omar Atie Date: Fri, 14 Aug 2026 16:22:24 -0700 Subject: [PATCH 02/17] fix(cuda.core): use cdef int for sync policy Cython cast Declare sync_policy_value at function scope so Cython can cast to CUsynchronizationPolicy when building the launch attribute. Co-authored-by: Cursor --- cuda_core/cuda/core/_launch_config.pyx | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/cuda_core/cuda/core/_launch_config.pyx b/cuda_core/cuda/core/_launch_config.pyx index 991bd3bbab7..afc9f130109 100644 --- a/cuda_core/cuda/core/_launch_config.pyx +++ b/cuda_core/cuda/core/_launch_config.pyx @@ -170,6 +170,7 @@ cdef class LaunchConfig: cdef cydriver.CUlaunchConfig _to_native_launch_config(self): cdef cydriver.CUlaunchConfig drv_cfg cdef cydriver.CUlaunchAttribute attr + cdef int sync_policy_value memset(&drv_cfg, 0, sizeof(drv_cfg)) self._attrs.resize(0) @@ -201,8 +202,9 @@ cdef class LaunchConfig: self._attrs.push_back(attr) if self.synchronization_policy is not None: + sync_policy_value = int(self.synchronization_policy) attr.id = cydriver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY - attr.value.syncPolicy = int(self.synchronization_policy) + attr.value.syncPolicy = sync_policy_value self._attrs.push_back(attr) drv_cfg.numAttrs = self._attrs.size() From b4da62b21a7ea22b8dc2ffc3c3a71aba89c172d6 Mon Sep 17 00:00:00 2001 From: Omar Atie Date: Fri, 14 Aug 2026 16:25:01 -0700 Subject: [PATCH 03/17] fix(cuda.core): align SynchronizationPolicyType with driver enum values Use cuda.bindings.driver.CUsynchronizationPolicy constants for the public IntEnum and skip GPU launch smoke tests when CUDA 13 bindings run against a CUDA 12 driver. Co-authored-by: Cursor --- cuda_core/cuda/core/typing.py | 11 ++++++----- cuda_core/tests/test_launcher.py | 11 +++++++++-- 2 files changed, 15 insertions(+), 7 deletions(-) diff --git a/cuda_core/cuda/core/typing.py b/cuda_core/cuda/core/typing.py index e6a4acaba0e..bc05a76f5d5 100644 --- a/cuda_core/cuda/core/typing.py +++ b/cuda_core/cuda/core/typing.py @@ -129,7 +129,8 @@ class PCHStatusType(StrEnum): class SynchronizationPolicyType(IntEnum): """CPU wait policy for host-side stream synchronization after a launch. - Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY``. + Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY`` and + ``cuda.bindings.driver.CUsynchronizationPolicy``. * ``AUTO`` — inherit the stream's synchronization policy. * ``SPIN`` — busy-wait on the CPU (lowest latency). @@ -137,10 +138,10 @@ class SynchronizationPolicyType(IntEnum): * ``BLOCKING_SYNC`` — block in the OS scheduler while waiting. """ - AUTO = 0 - SPIN = 1 - YIELD = 2 - BLOCKING_SYNC = 3 + AUTO = driver.CUsynchronizationPolicy.CU_SYNC_POLICY_AUTO + SPIN = driver.CUsynchronizationPolicy.CU_SYNC_POLICY_SPIN + YIELD = driver.CUsynchronizationPolicy.CU_SYNC_POLICY_YIELD + BLOCKING_SYNC = driver.CUsynchronizationPolicy.CU_SYNC_POLICY_BLOCKING_SYNC class GraphConditionalType(StrEnum): diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py index e7ac4d62d0c..e53d97c5c2a 100644 --- a/cuda_core/tests/test_launcher.py +++ b/cuda_core/tests/test_launcher.py @@ -221,7 +221,7 @@ def test_to_native_launch_config_synchronization_policy(policy, expected_value): from cuda.core._launch_config import _to_native_launch_config config = LaunchConfig(grid=1, block=1, synchronization_policy=policy) - assert config.synchronization_policy is SynchronizationPolicyType(int(expected_value)) + assert config.synchronization_policy == SynchronizationPolicyType(int(expected_value)) native = _to_native_launch_config(config) assert native.numAttrs == 1 @@ -290,12 +290,19 @@ class _FakeDev: ) def test_launch_with_synchronization_policy(init_cuda, policy): """Driver accepts per-launch synchronization policies on a real kernel launch.""" + import cuda.bindings + from cuda.core._utils.version import driver_version + + if int(cuda.bindings.__version__.split(".")[0]) >= 13 and driver_version()[0] < 13: + pytest.skip("CUDA 13 bindings produce modules incompatible with CUDA 12 drivers") + dev = Device() dev.set_current() stream = dev.create_stream() code = 'extern "C" __global__ void noop() {}' - program = Program(code, SourceCodeType.CXX) + arch = "".join(f"{i}" for i in dev.compute_capability) + program = Program(code, SourceCodeType.CXX, options=ProgramOptions(arch=f"sm_{arch}")) mod = program.compile(ObjectCodeFormatType.CUBIN) ker = mod.get_kernel("noop") From 9f4bb0a8f62aa71aebe21096b1b1f99e94bf6995 Mon Sep 17 00:00:00 2001 From: Omar Atie Date: Fri, 14 Aug 2026 16:31:09 -0700 Subject: [PATCH 04/17] style: fix pre-commit failures for PR #2637 Sort imports in test_launcher.py (ruff I001) and regenerate _launch_config.pyi via stubgen-pyx after LaunchConfig changes. Co-authored-by: Cursor --- cuda_core/cuda/core/_launch_config.pyi | 4 +--- cuda_core/tests/test_launcher.py | 2 +- 2 files changed, 2 insertions(+), 4 deletions(-) diff --git a/cuda_core/cuda/core/_launch_config.pyi b/cuda_core/cuda/core/_launch_config.pyi index 1082ffa1825..f7427eec51e 100644 --- a/cuda_core/cuda/core/_launch_config.pyi +++ b/cuda_core/cuda/core/_launch_config.pyi @@ -4,8 +4,6 @@ from __future__ import annotations from typing import Any -from cuda.core.typing import SynchronizationPolicyType - class LaunchConfig: """Customizable launch options. @@ -46,7 +44,7 @@ class LaunchConfig: kernel. Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY``. """ - def __init__(self, grid: int | tuple[int, ...] | None=None, cluster: int | tuple[int, ...] | None=None, block: int | tuple[int, ...] | None=None, shmem_size: int | None=None, is_cooperative: bool=False, programmatic_stream_serialization: bool=False, synchronization_policy: SynchronizationPolicyType | int | None=None) -> None: + def __init__(self, grid: int | tuple[int, ...] | None=None, cluster: int | tuple[int, ...] | None=None, block: int | tuple[int, ...] | None=None, shmem_size: int | None=None, is_cooperative: bool=False, programmatic_stream_serialization: bool=False, synchronization_policy: object=None) -> None: """Initialize LaunchConfig with validation. Parameters diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py index e53d97c5c2a..ca3a053a1e7 100644 --- a/cuda_core/tests/test_launcher.py +++ b/cuda_core/tests/test_launcher.py @@ -13,8 +13,8 @@ cp = None import numpy as np import pytest -from cuda.bindings import driver from conftest import skipif_need_cuda_headers +from cuda.bindings import driver from cuda.core import ( Device, From df550c9b7540c8505b98fa0f7db1ceff9d461016 Mon Sep 17 00:00:00 2001 From: Omar Atie Date: Fri, 14 Aug 2026 16:35:24 -0700 Subject: [PATCH 05/17] feat(cuda.core): add synchronization_policy to LaunchConfig Expose CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY through LaunchConfig so cuda.core users can set per-launch CPU wait policies without dropping to cuda.bindings.driver. Adds SynchronizationPolicyType and tests for native attribute mapping and real kernel launches. Closes #2628. Co-authored-by: Cursor Signed-off-by: Omar Atie --- cuda_core/cuda/core/_launch_config.pxd | 1 + cuda_core/cuda/core/_launch_config.pyi | 11 ++- cuda_core/cuda/core/_launch_config.pyx | 42 +++++++++ cuda_core/cuda/core/typing.py | 19 +++++ cuda_core/tests/test_launcher.py | 104 ++++++++++++++++++++++- cuda_core/tests/test_object_protocols.py | 2 +- 6 files changed, 175 insertions(+), 4 deletions(-) diff --git a/cuda_core/cuda/core/_launch_config.pxd b/cuda_core/cuda/core/_launch_config.pxd index 892a73f8efc..a4dc4b00c53 100644 --- a/cuda_core/cuda/core/_launch_config.pxd +++ b/cuda_core/cuda/core/_launch_config.pxd @@ -16,6 +16,7 @@ cdef class LaunchConfig: public int shmem_size public bint is_cooperative public bint programmatic_stream_serialization + public object synchronization_policy vector[cydriver.CUlaunchAttribute] _attrs object __weakref__ diff --git a/cuda_core/cuda/core/_launch_config.pyi b/cuda_core/cuda/core/_launch_config.pyi index 47187fb03d6..1082ffa1825 100644 --- a/cuda_core/cuda/core/_launch_config.pyi +++ b/cuda_core/cuda/core/_launch_config.pyi @@ -4,6 +4,8 @@ from __future__ import annotations from typing import Any +from cuda.core.typing import SynchronizationPolicyType + class LaunchConfig: """Customizable launch options. @@ -39,9 +41,12 @@ class LaunchConfig: Whether to allow programmatic stream serialization (PDL). When True, the kernel may overlap with a previous kernel in the same stream that signals completion via programmatic means. + synchronization_policy : SynchronizationPolicyType | None, optional + CPU wait policy applied when synchronizing the launch stream after this + kernel. Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY``. """ - def __init__(self, grid: int | tuple[int, ...] | None=None, cluster: int | tuple[int, ...] | None=None, block: int | tuple[int, ...] | None=None, shmem_size: int | None=None, is_cooperative: bool=False, programmatic_stream_serialization: bool=False) -> None: + def __init__(self, grid: int | tuple[int, ...] | None=None, cluster: int | tuple[int, ...] | None=None, block: int | tuple[int, ...] | None=None, shmem_size: int | None=None, is_cooperative: bool=False, programmatic_stream_serialization: bool=False, synchronization_policy: SynchronizationPolicyType | int | None=None) -> None: """Initialize LaunchConfig with validation. Parameters @@ -58,6 +63,8 @@ class LaunchConfig: Whether to launch as cooperative kernel (default: False) programmatic_stream_serialization : bool, optional Whether to allow programmatic stream serialization / PDL (default: False) + synchronization_policy : SynchronizationPolicyType | None, optional + CPU wait policy for synchronizing the launch stream (default: None) """ def _identity(self) -> tuple[Any, ...]: @@ -71,7 +78,7 @@ class LaunchConfig: def __hash__(self) -> int: ... -_LAUNCH_CONFIG_ATTRS = ('grid', 'cluster', 'block', 'shmem_size', 'is_cooperative', 'programmatic_stream_serialization') +_LAUNCH_CONFIG_ATTRS = ('grid', 'cluster', 'block', 'shmem_size', 'is_cooperative', 'programmatic_stream_serialization', 'synchronization_policy') __all__ = ['LaunchConfig'] def _to_native_launch_config(config: LaunchConfig) -> object: diff --git a/cuda_core/cuda/core/_launch_config.pyx b/cuda_core/cuda/core/_launch_config.pyx index adbf9a16c5d..991bd3bbab7 100644 --- a/cuda_core/cuda/core/_launch_config.pyx +++ b/cuda_core/cuda/core/_launch_config.pyx @@ -20,8 +20,32 @@ _LAUNCH_CONFIG_ATTRS = ( 'shmem_size', 'is_cooperative', 'programmatic_stream_serialization', + 'synchronization_policy', ) + +cdef object _validate_synchronization_policy(object policy): + from cuda.core.typing import SynchronizationPolicyType + + if policy is None: + return None + if isinstance(policy, SynchronizationPolicyType): + return policy + try: + value = int(policy) + except (TypeError, ValueError) as exc: + raise TypeError( + "LaunchConfig.synchronization_policy must be a SynchronizationPolicyType, " + f"cuda.bindings.driver.CUsynchronizationPolicy, or int; got {type(policy).__name__}" + ) from exc + try: + return SynchronizationPolicyType(value) + except ValueError as exc: + raise ValueError( + f"LaunchConfig.synchronization_policy must be one of " + f"{[member.name for member in SynchronizationPolicyType]}; got {policy!r}" + ) from exc + __all__ = ['LaunchConfig'] @@ -59,6 +83,9 @@ cdef class LaunchConfig: Whether to allow programmatic stream serialization (PDL). When True, the kernel may overlap with a previous kernel in the same stream that signals completion via programmatic means. + synchronization_policy : SynchronizationPolicyType | None, optional + CPU wait policy applied when synchronizing the launch stream after this + kernel. Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY``. """ # TODO: expand LaunchConfig to include other attributes @@ -72,6 +99,7 @@ cdef class LaunchConfig: shmem_size: int | None = None, is_cooperative: bool = False, programmatic_stream_serialization: bool = False, + synchronization_policy: object = None, ) -> None: """Initialize LaunchConfig with validation. @@ -89,6 +117,8 @@ cdef class LaunchConfig: Whether to launch as cooperative kernel (default: False) programmatic_stream_serialization : bool, optional Whether to allow programmatic stream serialization / PDL (default: False) + synchronization_policy : SynchronizationPolicyType | None, optional + CPU wait policy for synchronizing the launch stream (default: None) """ # Convert and validate grid and block dimensions self.grid = cast_to_3_tuple("LaunchConfig.grid", grid) @@ -116,6 +146,7 @@ cdef class LaunchConfig: self.is_cooperative = is_cooperative self.programmatic_stream_serialization = programmatic_stream_serialization + self.synchronization_policy = _validate_synchronization_policy(synchronization_policy) if self.is_cooperative and not Device().properties.cooperative_launch: raise CUDAError("cooperative kernels are not supported on this device") @@ -169,6 +200,11 @@ cdef class LaunchConfig: attr.value.programmaticStreamSerializationAllowed = 1 self._attrs.push_back(attr) + if self.synchronization_policy is not None: + attr.id = cydriver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY + attr.value.syncPolicy = int(self.synchronization_policy) + self._attrs.push_back(attr) + drv_cfg.numAttrs = self._attrs.size() drv_cfg.attrs = self._attrs.data() @@ -230,6 +266,12 @@ cpdef object _to_native_launch_config(LaunchConfig config): attr.value.programmaticStreamSerializationAllowed = 1 attrs.append(attr) + if config.synchronization_policy is not None: + attr = driver.CUlaunchAttribute() + attr.id = driver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY + attr.value.syncPolicy = int(config.synchronization_policy) + attrs.append(attr) + drv_cfg.numAttrs = len(attrs) drv_cfg.attrs = attrs diff --git a/cuda_core/cuda/core/typing.py b/cuda_core/cuda/core/typing.py index 1bf9bb7c0d2..e6a4acaba0e 100644 --- a/cuda_core/cuda/core/typing.py +++ b/cuda_core/cuda/core/typing.py @@ -5,6 +5,7 @@ """Public type aliases, protocols, and enumerations used in cuda.core API signatures.""" import sys +from enum import IntEnum from typing import TYPE_CHECKING from typing import Literal as _Literal from typing import TypeAlias as _TypeAlias @@ -47,6 +48,7 @@ class StrEnum(str, Enum): "ProcessStateType", "ReadModeType", "SourceCodeType", + "SynchronizationPolicyType", "VirtualMemoryAccessType", "VirtualMemoryAllocationType", "VirtualMemoryGranularityType", @@ -124,6 +126,23 @@ class PCHStatusType(StrEnum): FAILED = "failed" +class SynchronizationPolicyType(IntEnum): + """CPU wait policy for host-side stream synchronization after a launch. + + Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY``. + + * ``AUTO`` — inherit the stream's synchronization policy. + * ``SPIN`` — busy-wait on the CPU (lowest latency). + * ``YIELD`` — yield the CPU while waiting. + * ``BLOCKING_SYNC`` — block in the OS scheduler while waiting. + """ + + AUTO = 0 + SPIN = 1 + YIELD = 2 + BLOCKING_SYNC = 3 + + class GraphConditionalType(StrEnum): """Conditional node flavor for :class:`~cuda.core.graph.GraphBuilder`. diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py index e5cf05b435d..e7ac4d62d0c 100644 --- a/cuda_core/tests/test_launcher.py +++ b/cuda_core/tests/test_launcher.py @@ -13,6 +13,7 @@ cp = None import numpy as np import pytest +from cuda.bindings import driver from conftest import skipif_need_cuda_headers from cuda.core import ( @@ -26,7 +27,7 @@ ) from cuda.core._memory._legacy import _SynchronousMemoryResource from cuda.core._utils.cuda_utils import CUDAError -from cuda.core.typing import ObjectCodeFormatType, SourceCodeType +from cuda.core.typing import ObjectCodeFormatType, SourceCodeType, SynchronizationPolicyType def test_launch_config_init(init_cuda): @@ -202,6 +203,107 @@ def test_to_native_launch_config_pdl(): ) +@pytest.mark.parametrize( + ("policy", "expected_value"), + [ + (SynchronizationPolicyType.AUTO, driver.CUsynchronizationPolicy.CU_SYNC_POLICY_AUTO), + (SynchronizationPolicyType.SPIN, driver.CUsynchronizationPolicy.CU_SYNC_POLICY_SPIN), + (SynchronizationPolicyType.YIELD, driver.CUsynchronizationPolicy.CU_SYNC_POLICY_YIELD), + ( + SynchronizationPolicyType.BLOCKING_SYNC, + driver.CUsynchronizationPolicy.CU_SYNC_POLICY_BLOCKING_SYNC, + ), + (driver.CUsynchronizationPolicy.CU_SYNC_POLICY_SPIN, driver.CUsynchronizationPolicy.CU_SYNC_POLICY_SPIN), + ], +) +def test_to_native_launch_config_synchronization_policy(policy, expected_value): + """LaunchConfig.synchronization_policy maps to CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY.""" + from cuda.core._launch_config import _to_native_launch_config + + config = LaunchConfig(grid=1, block=1, synchronization_policy=policy) + assert config.synchronization_policy is SynchronizationPolicyType(int(expected_value)) + + native = _to_native_launch_config(config) + assert native.numAttrs == 1 + attr = native.attrs[0] + assert attr.id == driver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY + assert attr.value.syncPolicy == expected_value + + +def test_launch_config_synchronization_policy_default(): + config = LaunchConfig(grid=1, block=1) + assert config.synchronization_policy is None + + from cuda.core._launch_config import _to_native_launch_config + + native = _to_native_launch_config(config) + assert native.numAttrs == 0 + + +@pytest.mark.parametrize("invalid_policy", ["spin", -1, 99]) +def test_launch_config_synchronization_policy_invalid(invalid_policy): + with pytest.raises((TypeError, ValueError)): + LaunchConfig(grid=1, block=1, synchronization_policy=invalid_policy) + + +def test_to_native_launch_config_synchronization_policy_with_cooperative(monkeypatch): + """synchronization_policy can be combined with other launch attributes.""" + from cuda.core import _launch_config as _lc_mod + from cuda.core._launch_config import _to_native_launch_config + + class _FakeProps: + cooperative_launch = True + + class _FakeDev: + properties = _FakeProps() + + monkeypatch.setattr(_lc_mod, "Device", lambda: _FakeDev()) + + config = LaunchConfig( + grid=1, + block=1, + is_cooperative=True, + synchronization_policy=SynchronizationPolicyType.SPIN, + ) + native = _to_native_launch_config(config) + assert native.numAttrs == 2 + attr_ids = {attr.id for attr in native.attrs} + assert driver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_COOPERATIVE in attr_ids + assert driver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY in attr_ids + sync_attrs = [ + attr + for attr in native.attrs + if attr.id == driver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY + ] + assert len(sync_attrs) == 1 + assert sync_attrs[0].value.syncPolicy == driver.CUsynchronizationPolicy.CU_SYNC_POLICY_SPIN + + +@pytest.mark.parametrize( + "policy", + [ + SynchronizationPolicyType.AUTO, + SynchronizationPolicyType.SPIN, + SynchronizationPolicyType.YIELD, + SynchronizationPolicyType.BLOCKING_SYNC, + ], +) +def test_launch_with_synchronization_policy(init_cuda, policy): + """Driver accepts per-launch synchronization policies on a real kernel launch.""" + dev = Device() + dev.set_current() + stream = dev.create_stream() + + code = 'extern "C" __global__ void noop() {}' + program = Program(code, SourceCodeType.CXX) + mod = program.compile(ObjectCodeFormatType.CUBIN) + ker = mod.get_kernel("noop") + + config = LaunchConfig(grid=1, block=1, synchronization_policy=policy) + launch(stream, config, ker) + stream.sync() + + @skipif_need_cuda_headers def test_pdl_primary_secondary_overlap_same_stream(): """Primary + secondary PDL launch on one stream can overlap on Hopper+. diff --git a/cuda_core/tests/test_object_protocols.py b/cuda_core/tests/test_object_protocols.py index baf790abea8..9a54c38cfd0 100644 --- a/cuda_core/tests/test_object_protocols.py +++ b/cuda_core/tests/test_object_protocols.py @@ -685,7 +685,7 @@ def sample_switch_node_alt(sample_graphdef): "sample_launch_config", r"LaunchConfig\(grid=\(\d+, \d+, \d+\), cluster=.+, block=\(\d+, \d+, \d+\), " r"shmem_size=\d+, is_cooperative=(?:True|False), " - r"programmatic_stream_serialization=(?:True|False)\)", + r"programmatic_stream_serialization=(?:True|False), synchronization_policy=(?:None|SynchronizationPolicyType\.\w+)\)", ), ("sample_kernel", r""), # ObjectCode variations (by code_type) From c1c0e85855e10a744972faacf9e72ee9e144a72f Mon Sep 17 00:00:00 2001 From: Omar Atie Date: Fri, 14 Aug 2026 16:35:25 -0700 Subject: [PATCH 06/17] fix(cuda.core): use cdef int for sync policy Cython cast Declare sync_policy_value at function scope so Cython can cast to CUsynchronizationPolicy when building the launch attribute. Co-authored-by: Cursor Signed-off-by: Omar Atie --- cuda_core/cuda/core/_launch_config.pyx | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/cuda_core/cuda/core/_launch_config.pyx b/cuda_core/cuda/core/_launch_config.pyx index 991bd3bbab7..afc9f130109 100644 --- a/cuda_core/cuda/core/_launch_config.pyx +++ b/cuda_core/cuda/core/_launch_config.pyx @@ -170,6 +170,7 @@ cdef class LaunchConfig: cdef cydriver.CUlaunchConfig _to_native_launch_config(self): cdef cydriver.CUlaunchConfig drv_cfg cdef cydriver.CUlaunchAttribute attr + cdef int sync_policy_value memset(&drv_cfg, 0, sizeof(drv_cfg)) self._attrs.resize(0) @@ -201,8 +202,9 @@ cdef class LaunchConfig: self._attrs.push_back(attr) if self.synchronization_policy is not None: + sync_policy_value = int(self.synchronization_policy) attr.id = cydriver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY - attr.value.syncPolicy = int(self.synchronization_policy) + attr.value.syncPolicy = sync_policy_value self._attrs.push_back(attr) drv_cfg.numAttrs = self._attrs.size() From b1abc19c7d0d4d5a89e2119c654aa0b807c7eb27 Mon Sep 17 00:00:00 2001 From: Omar Atie Date: Fri, 14 Aug 2026 16:35:25 -0700 Subject: [PATCH 07/17] fix(cuda.core): align SynchronizationPolicyType with driver enum values Use cuda.bindings.driver.CUsynchronizationPolicy constants for the public IntEnum and skip GPU launch smoke tests when CUDA 13 bindings run against a CUDA 12 driver. Co-authored-by: Cursor Signed-off-by: Omar Atie --- cuda_core/cuda/core/typing.py | 11 ++++++----- cuda_core/tests/test_launcher.py | 11 +++++++++-- 2 files changed, 15 insertions(+), 7 deletions(-) diff --git a/cuda_core/cuda/core/typing.py b/cuda_core/cuda/core/typing.py index e6a4acaba0e..bc05a76f5d5 100644 --- a/cuda_core/cuda/core/typing.py +++ b/cuda_core/cuda/core/typing.py @@ -129,7 +129,8 @@ class PCHStatusType(StrEnum): class SynchronizationPolicyType(IntEnum): """CPU wait policy for host-side stream synchronization after a launch. - Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY``. + Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY`` and + ``cuda.bindings.driver.CUsynchronizationPolicy``. * ``AUTO`` — inherit the stream's synchronization policy. * ``SPIN`` — busy-wait on the CPU (lowest latency). @@ -137,10 +138,10 @@ class SynchronizationPolicyType(IntEnum): * ``BLOCKING_SYNC`` — block in the OS scheduler while waiting. """ - AUTO = 0 - SPIN = 1 - YIELD = 2 - BLOCKING_SYNC = 3 + AUTO = driver.CUsynchronizationPolicy.CU_SYNC_POLICY_AUTO + SPIN = driver.CUsynchronizationPolicy.CU_SYNC_POLICY_SPIN + YIELD = driver.CUsynchronizationPolicy.CU_SYNC_POLICY_YIELD + BLOCKING_SYNC = driver.CUsynchronizationPolicy.CU_SYNC_POLICY_BLOCKING_SYNC class GraphConditionalType(StrEnum): diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py index e7ac4d62d0c..e53d97c5c2a 100644 --- a/cuda_core/tests/test_launcher.py +++ b/cuda_core/tests/test_launcher.py @@ -221,7 +221,7 @@ def test_to_native_launch_config_synchronization_policy(policy, expected_value): from cuda.core._launch_config import _to_native_launch_config config = LaunchConfig(grid=1, block=1, synchronization_policy=policy) - assert config.synchronization_policy is SynchronizationPolicyType(int(expected_value)) + assert config.synchronization_policy == SynchronizationPolicyType(int(expected_value)) native = _to_native_launch_config(config) assert native.numAttrs == 1 @@ -290,12 +290,19 @@ class _FakeDev: ) def test_launch_with_synchronization_policy(init_cuda, policy): """Driver accepts per-launch synchronization policies on a real kernel launch.""" + import cuda.bindings + from cuda.core._utils.version import driver_version + + if int(cuda.bindings.__version__.split(".")[0]) >= 13 and driver_version()[0] < 13: + pytest.skip("CUDA 13 bindings produce modules incompatible with CUDA 12 drivers") + dev = Device() dev.set_current() stream = dev.create_stream() code = 'extern "C" __global__ void noop() {}' - program = Program(code, SourceCodeType.CXX) + arch = "".join(f"{i}" for i in dev.compute_capability) + program = Program(code, SourceCodeType.CXX, options=ProgramOptions(arch=f"sm_{arch}")) mod = program.compile(ObjectCodeFormatType.CUBIN) ker = mod.get_kernel("noop") From 0008acaf03243512d6bf5441da5aa002dc67a2c6 Mon Sep 17 00:00:00 2001 From: Omar Atie Date: Fri, 14 Aug 2026 16:35:25 -0700 Subject: [PATCH 08/17] style: fix pre-commit failures for PR #2637 Sort imports in test_launcher.py (ruff I001) and regenerate _launch_config.pyi via stubgen-pyx after LaunchConfig changes. Co-authored-by: Cursor Signed-off-by: Omar Atie --- cuda_core/cuda/core/_launch_config.pyi | 4 +--- cuda_core/tests/test_launcher.py | 2 +- 2 files changed, 2 insertions(+), 4 deletions(-) diff --git a/cuda_core/cuda/core/_launch_config.pyi b/cuda_core/cuda/core/_launch_config.pyi index 1082ffa1825..f7427eec51e 100644 --- a/cuda_core/cuda/core/_launch_config.pyi +++ b/cuda_core/cuda/core/_launch_config.pyi @@ -4,8 +4,6 @@ from __future__ import annotations from typing import Any -from cuda.core.typing import SynchronizationPolicyType - class LaunchConfig: """Customizable launch options. @@ -46,7 +44,7 @@ class LaunchConfig: kernel. Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY``. """ - def __init__(self, grid: int | tuple[int, ...] | None=None, cluster: int | tuple[int, ...] | None=None, block: int | tuple[int, ...] | None=None, shmem_size: int | None=None, is_cooperative: bool=False, programmatic_stream_serialization: bool=False, synchronization_policy: SynchronizationPolicyType | int | None=None) -> None: + def __init__(self, grid: int | tuple[int, ...] | None=None, cluster: int | tuple[int, ...] | None=None, block: int | tuple[int, ...] | None=None, shmem_size: int | None=None, is_cooperative: bool=False, programmatic_stream_serialization: bool=False, synchronization_policy: object=None) -> None: """Initialize LaunchConfig with validation. Parameters diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py index e53d97c5c2a..ca3a053a1e7 100644 --- a/cuda_core/tests/test_launcher.py +++ b/cuda_core/tests/test_launcher.py @@ -13,8 +13,8 @@ cp = None import numpy as np import pytest -from cuda.bindings import driver from conftest import skipif_need_cuda_headers +from cuda.bindings import driver from cuda.core import ( Device, From dd692198c46a730800545c925b4e118dda01a641 Mon Sep 17 00:00:00 2001 From: Omar Atie Date: Fri, 14 Aug 2026 17:00:34 -0700 Subject: [PATCH 09/17] style: sort imports in test_launcher.py (#2628) Fix ruff I001 unsorted-imports for pre-commit.ci on PR #2637. Signed-off-by: Omar Atie --- cuda_core/tests/test_launcher.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py index ca3a053a1e7..3296386543b 100644 --- a/cuda_core/tests/test_launcher.py +++ b/cuda_core/tests/test_launcher.py @@ -14,8 +14,8 @@ import numpy as np import pytest from conftest import skipif_need_cuda_headers -from cuda.bindings import driver +from cuda.bindings import driver from cuda.core import ( Device, DeviceMemoryResource, From f3d8d67e072ca630a3101e08682f279b61bae724 Mon Sep 17 00:00:00 2001 From: Omar Atie Date: Sat, 15 Aug 2026 10:12:48 -0700 Subject: [PATCH 10/17] fix: resolve duplicate import after history-restore merge (#2637) Clean up test_launcher.py after merging the pre-force-push commit lineage back into the branch. Signed-off-by: Omar Atie Co-authored-by: Cursor --- cuda_core/tests/test_launcher.py | 1 - 1 file changed, 1 deletion(-) diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py index cc2b81a06ad..3296386543b 100644 --- a/cuda_core/tests/test_launcher.py +++ b/cuda_core/tests/test_launcher.py @@ -14,7 +14,6 @@ import numpy as np import pytest from conftest import skipif_need_cuda_headers -from cuda.bindings import driver from cuda.bindings import driver from cuda.core import ( From 1e48c4eddc4cc64a443fcc4d6efbf2c0dfa84a18 Mon Sep 17 00:00:00 2001 From: Omar Atie Date: Sat, 15 Aug 2026 10:15:25 -0700 Subject: [PATCH 11/17] chore: confirm branch synced with upstream main (#2637) Record that feat/launch-config-sync-policy-2628 includes upstream main at db2801873b. Future updates will use merge commits, not force-push. Signed-off-by: Omar Atie From 21c56b041a499555ddfdee28ed9e62ee9cef6466 Mon Sep 17 00:00:00 2001 From: Omar Atie Date: Mon, 17 Aug 2026 14:47:39 -0700 Subject: [PATCH 12/17] fix(cuda.core): correct sync policy mapping and document enum Use CUsynchronizationPolicy when setting syncPolicy in the Python _to_native_launch_config helper, and add SynchronizationPolicyType to api_private.rst for docs consistency checks. Co-authored-by: Cursor --- cuda_core/cuda/core/_launch_config.pyx | 4 +++- cuda_core/docs/source/api_private.rst | 1 + 2 files changed, 4 insertions(+), 1 deletion(-) diff --git a/cuda_core/cuda/core/_launch_config.pyx b/cuda_core/cuda/core/_launch_config.pyx index afc9f130109..16cb950aebe 100644 --- a/cuda_core/cuda/core/_launch_config.pyx +++ b/cuda_core/cuda/core/_launch_config.pyx @@ -271,7 +271,9 @@ cpdef object _to_native_launch_config(LaunchConfig config): if config.synchronization_policy is not None: attr = driver.CUlaunchAttribute() attr.id = driver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY - attr.value.syncPolicy = int(config.synchronization_policy) + attr.value.syncPolicy = driver.CUsynchronizationPolicy( + int(config.synchronization_policy) + ) attrs.append(attr) drv_cfg.numAttrs = len(attrs) diff --git a/cuda_core/docs/source/api_private.rst b/cuda_core/docs/source/api_private.rst index 80675799c07..de161c96ae0 100644 --- a/cuda_core/docs/source/api_private.rst +++ b/cuda_core/docs/source/api_private.rst @@ -35,6 +35,7 @@ CUDA runtime typing.ProcessStateType typing.ReadModeType typing.SourceCodeType + typing.SynchronizationPolicyType typing.VirtualMemoryAccessType typing.VirtualMemoryAllocationType typing.VirtualMemoryGranularityType From 71763c5fc677c56038ea16f3b756e43e3aded7a0 Mon Sep 17 00:00:00 2001 From: Omar Atie Date: Mon, 17 Aug 2026 18:02:16 -0700 Subject: [PATCH 13/17] revert(cuda.core): drop LaunchConfig.synchronization_policy (#2637) CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY is not accepted by cuLaunchKernelEx, so the per-launch API requested in #2628 cannot be implemented. Every GPU test job rejected the launch with CUDA_ERROR_INVALID_VALUE on CUDA 12.9, 13.0 and 13.3 across linux-64, linux-aarch64 and win-64. CUDA documents this attribute as "Valid for streams", unlike CU_LAUNCH_ATTRIBUTE_COOPERATIVE which is "Valid for graph nodes, launches". A driver-level probe on an RTX A4000 (driver 550.144.03, cuda-bindings 12.9.7) confirms the distinction is real: the same no-op kernel launches successfully with no attributes and with COOPERATIVE, but returns CUDA_ERROR_INVALID_VALUE for all four synchronization policies, while cuStreamSetAttribute accepts the same values. Revert the feature so this branch matches main, and defer to maintainers on whether #2628 should be re-scoped to a stream-level API. Co-authored-by: Cursor --- cuda_core/cuda/core/_launch_config.pxd | 1 - cuda_core/cuda/core/_launch_config.pyi | 9 +- cuda_core/cuda/core/_launch_config.pyx | 46 ---------- cuda_core/cuda/core/typing.py | 20 ---- cuda_core/docs/source/api_private.rst | 1 - cuda_core/tests/test_launcher.py | 111 +---------------------- cuda_core/tests/test_object_protocols.py | 2 +- 7 files changed, 4 insertions(+), 186 deletions(-) diff --git a/cuda_core/cuda/core/_launch_config.pxd b/cuda_core/cuda/core/_launch_config.pxd index a4dc4b00c53..892a73f8efc 100644 --- a/cuda_core/cuda/core/_launch_config.pxd +++ b/cuda_core/cuda/core/_launch_config.pxd @@ -16,7 +16,6 @@ cdef class LaunchConfig: public int shmem_size public bint is_cooperative public bint programmatic_stream_serialization - public object synchronization_policy vector[cydriver.CUlaunchAttribute] _attrs object __weakref__ diff --git a/cuda_core/cuda/core/_launch_config.pyi b/cuda_core/cuda/core/_launch_config.pyi index f7427eec51e..47187fb03d6 100644 --- a/cuda_core/cuda/core/_launch_config.pyi +++ b/cuda_core/cuda/core/_launch_config.pyi @@ -39,12 +39,9 @@ class LaunchConfig: Whether to allow programmatic stream serialization (PDL). When True, the kernel may overlap with a previous kernel in the same stream that signals completion via programmatic means. - synchronization_policy : SynchronizationPolicyType | None, optional - CPU wait policy applied when synchronizing the launch stream after this - kernel. Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY``. """ - def __init__(self, grid: int | tuple[int, ...] | None=None, cluster: int | tuple[int, ...] | None=None, block: int | tuple[int, ...] | None=None, shmem_size: int | None=None, is_cooperative: bool=False, programmatic_stream_serialization: bool=False, synchronization_policy: object=None) -> None: + def __init__(self, grid: int | tuple[int, ...] | None=None, cluster: int | tuple[int, ...] | None=None, block: int | tuple[int, ...] | None=None, shmem_size: int | None=None, is_cooperative: bool=False, programmatic_stream_serialization: bool=False) -> None: """Initialize LaunchConfig with validation. Parameters @@ -61,8 +58,6 @@ class LaunchConfig: Whether to launch as cooperative kernel (default: False) programmatic_stream_serialization : bool, optional Whether to allow programmatic stream serialization / PDL (default: False) - synchronization_policy : SynchronizationPolicyType | None, optional - CPU wait policy for synchronizing the launch stream (default: None) """ def _identity(self) -> tuple[Any, ...]: @@ -76,7 +71,7 @@ class LaunchConfig: def __hash__(self) -> int: ... -_LAUNCH_CONFIG_ATTRS = ('grid', 'cluster', 'block', 'shmem_size', 'is_cooperative', 'programmatic_stream_serialization', 'synchronization_policy') +_LAUNCH_CONFIG_ATTRS = ('grid', 'cluster', 'block', 'shmem_size', 'is_cooperative', 'programmatic_stream_serialization') __all__ = ['LaunchConfig'] def _to_native_launch_config(config: LaunchConfig) -> object: diff --git a/cuda_core/cuda/core/_launch_config.pyx b/cuda_core/cuda/core/_launch_config.pyx index 16cb950aebe..adbf9a16c5d 100644 --- a/cuda_core/cuda/core/_launch_config.pyx +++ b/cuda_core/cuda/core/_launch_config.pyx @@ -20,32 +20,8 @@ _LAUNCH_CONFIG_ATTRS = ( 'shmem_size', 'is_cooperative', 'programmatic_stream_serialization', - 'synchronization_policy', ) - -cdef object _validate_synchronization_policy(object policy): - from cuda.core.typing import SynchronizationPolicyType - - if policy is None: - return None - if isinstance(policy, SynchronizationPolicyType): - return policy - try: - value = int(policy) - except (TypeError, ValueError) as exc: - raise TypeError( - "LaunchConfig.synchronization_policy must be a SynchronizationPolicyType, " - f"cuda.bindings.driver.CUsynchronizationPolicy, or int; got {type(policy).__name__}" - ) from exc - try: - return SynchronizationPolicyType(value) - except ValueError as exc: - raise ValueError( - f"LaunchConfig.synchronization_policy must be one of " - f"{[member.name for member in SynchronizationPolicyType]}; got {policy!r}" - ) from exc - __all__ = ['LaunchConfig'] @@ -83,9 +59,6 @@ cdef class LaunchConfig: Whether to allow programmatic stream serialization (PDL). When True, the kernel may overlap with a previous kernel in the same stream that signals completion via programmatic means. - synchronization_policy : SynchronizationPolicyType | None, optional - CPU wait policy applied when synchronizing the launch stream after this - kernel. Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY``. """ # TODO: expand LaunchConfig to include other attributes @@ -99,7 +72,6 @@ cdef class LaunchConfig: shmem_size: int | None = None, is_cooperative: bool = False, programmatic_stream_serialization: bool = False, - synchronization_policy: object = None, ) -> None: """Initialize LaunchConfig with validation. @@ -117,8 +89,6 @@ cdef class LaunchConfig: Whether to launch as cooperative kernel (default: False) programmatic_stream_serialization : bool, optional Whether to allow programmatic stream serialization / PDL (default: False) - synchronization_policy : SynchronizationPolicyType | None, optional - CPU wait policy for synchronizing the launch stream (default: None) """ # Convert and validate grid and block dimensions self.grid = cast_to_3_tuple("LaunchConfig.grid", grid) @@ -146,7 +116,6 @@ cdef class LaunchConfig: self.is_cooperative = is_cooperative self.programmatic_stream_serialization = programmatic_stream_serialization - self.synchronization_policy = _validate_synchronization_policy(synchronization_policy) if self.is_cooperative and not Device().properties.cooperative_launch: raise CUDAError("cooperative kernels are not supported on this device") @@ -170,7 +139,6 @@ cdef class LaunchConfig: cdef cydriver.CUlaunchConfig _to_native_launch_config(self): cdef cydriver.CUlaunchConfig drv_cfg cdef cydriver.CUlaunchAttribute attr - cdef int sync_policy_value memset(&drv_cfg, 0, sizeof(drv_cfg)) self._attrs.resize(0) @@ -201,12 +169,6 @@ cdef class LaunchConfig: attr.value.programmaticStreamSerializationAllowed = 1 self._attrs.push_back(attr) - if self.synchronization_policy is not None: - sync_policy_value = int(self.synchronization_policy) - attr.id = cydriver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY - attr.value.syncPolicy = sync_policy_value - self._attrs.push_back(attr) - drv_cfg.numAttrs = self._attrs.size() drv_cfg.attrs = self._attrs.data() @@ -268,14 +230,6 @@ cpdef object _to_native_launch_config(LaunchConfig config): attr.value.programmaticStreamSerializationAllowed = 1 attrs.append(attr) - if config.synchronization_policy is not None: - attr = driver.CUlaunchAttribute() - attr.id = driver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY - attr.value.syncPolicy = driver.CUsynchronizationPolicy( - int(config.synchronization_policy) - ) - attrs.append(attr) - drv_cfg.numAttrs = len(attrs) drv_cfg.attrs = attrs diff --git a/cuda_core/cuda/core/typing.py b/cuda_core/cuda/core/typing.py index bc05a76f5d5..1bf9bb7c0d2 100644 --- a/cuda_core/cuda/core/typing.py +++ b/cuda_core/cuda/core/typing.py @@ -5,7 +5,6 @@ """Public type aliases, protocols, and enumerations used in cuda.core API signatures.""" import sys -from enum import IntEnum from typing import TYPE_CHECKING from typing import Literal as _Literal from typing import TypeAlias as _TypeAlias @@ -48,7 +47,6 @@ class StrEnum(str, Enum): "ProcessStateType", "ReadModeType", "SourceCodeType", - "SynchronizationPolicyType", "VirtualMemoryAccessType", "VirtualMemoryAllocationType", "VirtualMemoryGranularityType", @@ -126,24 +124,6 @@ class PCHStatusType(StrEnum): FAILED = "failed" -class SynchronizationPolicyType(IntEnum): - """CPU wait policy for host-side stream synchronization after a launch. - - Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY`` and - ``cuda.bindings.driver.CUsynchronizationPolicy``. - - * ``AUTO`` — inherit the stream's synchronization policy. - * ``SPIN`` — busy-wait on the CPU (lowest latency). - * ``YIELD`` — yield the CPU while waiting. - * ``BLOCKING_SYNC`` — block in the OS scheduler while waiting. - """ - - AUTO = driver.CUsynchronizationPolicy.CU_SYNC_POLICY_AUTO - SPIN = driver.CUsynchronizationPolicy.CU_SYNC_POLICY_SPIN - YIELD = driver.CUsynchronizationPolicy.CU_SYNC_POLICY_YIELD - BLOCKING_SYNC = driver.CUsynchronizationPolicy.CU_SYNC_POLICY_BLOCKING_SYNC - - class GraphConditionalType(StrEnum): """Conditional node flavor for :class:`~cuda.core.graph.GraphBuilder`. diff --git a/cuda_core/docs/source/api_private.rst b/cuda_core/docs/source/api_private.rst index de161c96ae0..80675799c07 100644 --- a/cuda_core/docs/source/api_private.rst +++ b/cuda_core/docs/source/api_private.rst @@ -35,7 +35,6 @@ CUDA runtime typing.ProcessStateType typing.ReadModeType typing.SourceCodeType - typing.SynchronizationPolicyType typing.VirtualMemoryAccessType typing.VirtualMemoryAllocationType typing.VirtualMemoryGranularityType diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py index 3296386543b..e5cf05b435d 100644 --- a/cuda_core/tests/test_launcher.py +++ b/cuda_core/tests/test_launcher.py @@ -15,7 +15,6 @@ import pytest from conftest import skipif_need_cuda_headers -from cuda.bindings import driver from cuda.core import ( Device, DeviceMemoryResource, @@ -27,7 +26,7 @@ ) from cuda.core._memory._legacy import _SynchronousMemoryResource from cuda.core._utils.cuda_utils import CUDAError -from cuda.core.typing import ObjectCodeFormatType, SourceCodeType, SynchronizationPolicyType +from cuda.core.typing import ObjectCodeFormatType, SourceCodeType def test_launch_config_init(init_cuda): @@ -203,114 +202,6 @@ def test_to_native_launch_config_pdl(): ) -@pytest.mark.parametrize( - ("policy", "expected_value"), - [ - (SynchronizationPolicyType.AUTO, driver.CUsynchronizationPolicy.CU_SYNC_POLICY_AUTO), - (SynchronizationPolicyType.SPIN, driver.CUsynchronizationPolicy.CU_SYNC_POLICY_SPIN), - (SynchronizationPolicyType.YIELD, driver.CUsynchronizationPolicy.CU_SYNC_POLICY_YIELD), - ( - SynchronizationPolicyType.BLOCKING_SYNC, - driver.CUsynchronizationPolicy.CU_SYNC_POLICY_BLOCKING_SYNC, - ), - (driver.CUsynchronizationPolicy.CU_SYNC_POLICY_SPIN, driver.CUsynchronizationPolicy.CU_SYNC_POLICY_SPIN), - ], -) -def test_to_native_launch_config_synchronization_policy(policy, expected_value): - """LaunchConfig.synchronization_policy maps to CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY.""" - from cuda.core._launch_config import _to_native_launch_config - - config = LaunchConfig(grid=1, block=1, synchronization_policy=policy) - assert config.synchronization_policy == SynchronizationPolicyType(int(expected_value)) - - native = _to_native_launch_config(config) - assert native.numAttrs == 1 - attr = native.attrs[0] - assert attr.id == driver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY - assert attr.value.syncPolicy == expected_value - - -def test_launch_config_synchronization_policy_default(): - config = LaunchConfig(grid=1, block=1) - assert config.synchronization_policy is None - - from cuda.core._launch_config import _to_native_launch_config - - native = _to_native_launch_config(config) - assert native.numAttrs == 0 - - -@pytest.mark.parametrize("invalid_policy", ["spin", -1, 99]) -def test_launch_config_synchronization_policy_invalid(invalid_policy): - with pytest.raises((TypeError, ValueError)): - LaunchConfig(grid=1, block=1, synchronization_policy=invalid_policy) - - -def test_to_native_launch_config_synchronization_policy_with_cooperative(monkeypatch): - """synchronization_policy can be combined with other launch attributes.""" - from cuda.core import _launch_config as _lc_mod - from cuda.core._launch_config import _to_native_launch_config - - class _FakeProps: - cooperative_launch = True - - class _FakeDev: - properties = _FakeProps() - - monkeypatch.setattr(_lc_mod, "Device", lambda: _FakeDev()) - - config = LaunchConfig( - grid=1, - block=1, - is_cooperative=True, - synchronization_policy=SynchronizationPolicyType.SPIN, - ) - native = _to_native_launch_config(config) - assert native.numAttrs == 2 - attr_ids = {attr.id for attr in native.attrs} - assert driver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_COOPERATIVE in attr_ids - assert driver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY in attr_ids - sync_attrs = [ - attr - for attr in native.attrs - if attr.id == driver.CUlaunchAttributeID.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY - ] - assert len(sync_attrs) == 1 - assert sync_attrs[0].value.syncPolicy == driver.CUsynchronizationPolicy.CU_SYNC_POLICY_SPIN - - -@pytest.mark.parametrize( - "policy", - [ - SynchronizationPolicyType.AUTO, - SynchronizationPolicyType.SPIN, - SynchronizationPolicyType.YIELD, - SynchronizationPolicyType.BLOCKING_SYNC, - ], -) -def test_launch_with_synchronization_policy(init_cuda, policy): - """Driver accepts per-launch synchronization policies on a real kernel launch.""" - import cuda.bindings - from cuda.core._utils.version import driver_version - - if int(cuda.bindings.__version__.split(".")[0]) >= 13 and driver_version()[0] < 13: - pytest.skip("CUDA 13 bindings produce modules incompatible with CUDA 12 drivers") - - dev = Device() - dev.set_current() - stream = dev.create_stream() - - code = 'extern "C" __global__ void noop() {}' - arch = "".join(f"{i}" for i in dev.compute_capability) - program = Program(code, SourceCodeType.CXX, options=ProgramOptions(arch=f"sm_{arch}")) - mod = program.compile(ObjectCodeFormatType.CUBIN) - ker = mod.get_kernel("noop") - - config = LaunchConfig(grid=1, block=1, synchronization_policy=policy) - launch(stream, config, ker) - stream.sync() - - @skipif_need_cuda_headers def test_pdl_primary_secondary_overlap_same_stream(): """Primary + secondary PDL launch on one stream can overlap on Hopper+. diff --git a/cuda_core/tests/test_object_protocols.py b/cuda_core/tests/test_object_protocols.py index 9a54c38cfd0..baf790abea8 100644 --- a/cuda_core/tests/test_object_protocols.py +++ b/cuda_core/tests/test_object_protocols.py @@ -685,7 +685,7 @@ def sample_switch_node_alt(sample_graphdef): "sample_launch_config", r"LaunchConfig\(grid=\(\d+, \d+, \d+\), cluster=.+, block=\(\d+, \d+, \d+\), " r"shmem_size=\d+, is_cooperative=(?:True|False), " - r"programmatic_stream_serialization=(?:True|False), synchronization_policy=(?:None|SynchronizationPolicyType\.\w+)\)", + r"programmatic_stream_serialization=(?:True|False)\)", ), ("sample_kernel", r""), # ObjectCode variations (by code_type) From 64b318aa6322d2eca64086dd81d40668cb333cbc Mon Sep 17 00:00:00 2001 From: Omar Atie Date: Mon, 17 Aug 2026 18:13:32 -0700 Subject: [PATCH 14/17] feat(cuda.core): add Stream.synchronization_policy via cuStreamSetAttribute Re-scope #2628 from the invalid per-launch attribute to the documented stream attribute path. CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY is rejected by cuLaunchKernelEx but accepted by cuStreamSetAttribute on real hardware (RunPod RTX A4000, driver 550 / CUDA 12.4). Expose SynchronizationPolicyType, Stream.synchronization_policy get/set, and an optional StreamOptions.synchronization_policy for creation-time configuration. Add GPU tests covering set/get, defaults, options, and launch+sync for all four policies. Co-authored-by: Cursor --- cuda_core/cuda/core/_stream.pxd | 1 + cuda_core/cuda/core/_stream.pyi | 11 ++++ cuda_core/cuda/core/_stream.pyx | 89 +++++++++++++++++++++++++++ cuda_core/cuda/core/typing.py | 20 ++++++ cuda_core/docs/source/api_private.rst | 1 + cuda_core/tests/test_stream.py | 87 ++++++++++++++++++++++++++ 6 files changed, 209 insertions(+) diff --git a/cuda_core/cuda/core/_stream.pxd b/cuda_core/cuda/core/_stream.pxd index dc9a2da826c..abc4e396fd0 100644 --- a/cuda_core/cuda/core/_stream.pxd +++ b/cuda_core/cuda/core/_stream.pxd @@ -13,6 +13,7 @@ cdef class Stream: int _device_id int _nonblocking int _priority + int _synchronization_policy object __weakref__ @staticmethod diff --git a/cuda_core/cuda/core/_stream.pyi b/cuda_core/cuda/core/_stream.pyi index 99af5f9b15b..b094cf67354 100644 --- a/cuda_core/cuda/core/_stream.pyi +++ b/cuda_core/cuda/core/_stream.pyi @@ -25,10 +25,13 @@ class StreamOptions: priority : int, optional Stream priority where lower number represents a higher priority. (Default to lowest priority) + synchronization_policy : SynchronizationPolicyType | None, optional + CPU wait policy applied when synchronizing this stream from the host. """ nonblocking: cython.bint = True priority: int | None = None + synchronization_policy: object = None class IsStreamType(Protocol): @@ -110,6 +113,14 @@ class Stream: def priority(self) -> int: """Return the stream priority.""" + @property + def synchronization_policy(self): + """Return the stream's CPU wait policy for host-side synchronization.""" + + @synchronization_policy.setter + def synchronization_policy(self, policy) -> None: + """Set the stream's CPU wait policy for subsequent work on this stream.""" + def sync(self) -> None: """Synchronize the stream.""" diff --git a/cuda_core/cuda/core/_stream.pyx b/cuda_core/cuda/core/_stream.pyx index c8c5faf74bc..1171d536c13 100644 --- a/cuda_core/cuda/core/_stream.pyx +++ b/cuda_core/cuda/core/_stream.pyx @@ -49,6 +49,53 @@ if TYPE_CHECKING: __all__ = ['LEGACY_DEFAULT_STREAM', 'PER_THREAD_DEFAULT_STREAM', 'Stream', 'StreamOptions'] +cdef int _SYNC_POLICY_ATTR_ID = ( + cydriver.CUlaunchAttributeID_enum.CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY +) + + +cdef object _validate_synchronization_policy(object policy): + from cuda.core.typing import SynchronizationPolicyType + + if policy is None: + raise TypeError("synchronization_policy must not be None") + if isinstance(policy, SynchronizationPolicyType): + return policy + try: + value = int(policy) + except (TypeError, ValueError) as exc: + raise TypeError( + "Stream.synchronization_policy must be a SynchronizationPolicyType, " + f"cuda.bindings.driver.CUsynchronizationPolicy, or int; got {type(policy).__name__}" + ) from exc + try: + return SynchronizationPolicyType(value) + except ValueError as exc: + raise ValueError( + f"Stream.synchronization_policy must be one of " + f"{[member.name for member in SynchronizationPolicyType]}; got {policy!r}" + ) from exc + + +cdef int Stream_query_synchronization_policy(Stream self) except?-1 nogil: + cdef cydriver.CUstreamAttrValue value + HANDLE_RETURN(cydriver.cuStreamGetAttribute( + as_cu(self._h_stream), + _SYNC_POLICY_ATTR_ID, + &value, + )) + return value.syncPolicy + + +cdef void Stream_apply_synchronization_policy(Stream self, int policy) except * nogil: + cdef cydriver.CUstreamAttrValue value + value.syncPolicy = policy + HANDLE_RETURN(cydriver.cuStreamSetAttribute( + as_cu(self._h_stream), + _SYNC_POLICY_ATTR_ID, + &value, + )) + @dataclass cdef class StreamOptions: @@ -61,11 +108,16 @@ cdef class StreamOptions: priority : int, optional Stream priority where lower number represents a higher priority. (Default to lowest priority) + synchronization_policy : SynchronizationPolicyType | None, optional + CPU wait policy applied when synchronizing this stream from the host. + Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY`` via + ``cuStreamSetAttribute``. (Default to None, leaving the driver default) """ nonblocking : cython.bint = True priority: int | None = None + synchronization_policy: object = None class IsStreamType(Protocol): @@ -110,6 +162,7 @@ cdef class Stream: s._device_id = -1 # lazy init'd (invalid sentinel) s._nonblocking = -1 # lazy init'd s._priority = INT32_MIN # lazy init'd + s._synchronization_policy = INT32_MIN # lazy init'd return s @classmethod @@ -146,6 +199,7 @@ cdef class Stream: cdef StreamOptions opts = check_or_create_options(StreamOptions, options, "Stream options") nonblocking = opts.nonblocking priority = opts.priority + sync_policy_opt = opts.synchronization_policy cdef unsigned int flags = (cydriver.CUstream_flags.CU_STREAM_NON_BLOCKING if nonblocking else cydriver.CUstream_flags.CU_STREAM_DEFAULT) @@ -190,6 +244,12 @@ cdef class Stream: cdef Stream self = Stream._from_handle(cls, h_stream) self._nonblocking = int(nonblocking) self._priority = prio + if sync_policy_opt is not None: + validated = _validate_synchronization_policy(sync_policy_opt) + cdef int sync_policy = int(validated) + with nogil: + Stream_apply_synchronization_policy(self, sync_policy) + self._synchronization_policy = sync_policy if device_id is not None: self._device_id = device_id return self @@ -251,6 +311,35 @@ cdef class Stream: self._priority = prio return self._priority + @property + def synchronization_policy(self): + """Return the stream's CPU wait policy for host-side synchronization. + + Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY``, queried via + ``cuStreamGetAttribute``. + """ + from cuda.core.typing import SynchronizationPolicyType + + cdef int policy + if self._synchronization_policy == INT32_MIN or Stream_is_default_token(self): + with nogil: + policy = Stream_query_synchronization_policy(self) + if not Stream_is_default_token(self): + self._synchronization_policy = policy + else: + policy = self._synchronization_policy + return SynchronizationPolicyType(policy) + + @synchronization_policy.setter + def synchronization_policy(self, object policy): + """Set the stream's CPU wait policy for subsequent work on this stream.""" + cdef object validated = _validate_synchronization_policy(policy) + cdef int value = int(validated) + with nogil: + Stream_apply_synchronization_policy(self, value) + if not Stream_is_default_token(self): + self._synchronization_policy = value + def sync(self) -> None: """Synchronize the stream.""" with nogil: diff --git a/cuda_core/cuda/core/typing.py b/cuda_core/cuda/core/typing.py index 1bf9bb7c0d2..0d40d11d9a3 100644 --- a/cuda_core/cuda/core/typing.py +++ b/cuda_core/cuda/core/typing.py @@ -5,6 +5,7 @@ """Public type aliases, protocols, and enumerations used in cuda.core API signatures.""" import sys +from enum import IntEnum from typing import TYPE_CHECKING from typing import Literal as _Literal from typing import TypeAlias as _TypeAlias @@ -47,6 +48,7 @@ class StrEnum(str, Enum): "ProcessStateType", "ReadModeType", "SourceCodeType", + "SynchronizationPolicyType", "VirtualMemoryAccessType", "VirtualMemoryAllocationType", "VirtualMemoryGranularityType", @@ -124,6 +126,24 @@ class PCHStatusType(StrEnum): FAILED = "failed" +class SynchronizationPolicyType(IntEnum): + """CPU wait policy when synchronizing a stream from the host. + + Maps to ``cuda.bindings.driver.CUsynchronizationPolicy`` and is applied via + ``cuStreamSetAttribute`` with ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY``. + + * ``AUTO`` — driver selects the policy (default for new streams). + * ``SPIN`` — busy-wait on the CPU (lowest latency). + * ``YIELD`` — yield the CPU while waiting. + * ``BLOCKING_SYNC`` — block in the OS scheduler while waiting. + """ + + AUTO = driver.CUsynchronizationPolicy.CU_SYNC_POLICY_AUTO + SPIN = driver.CUsynchronizationPolicy.CU_SYNC_POLICY_SPIN + YIELD = driver.CUsynchronizationPolicy.CU_SYNC_POLICY_YIELD + BLOCKING_SYNC = driver.CUsynchronizationPolicy.CU_SYNC_POLICY_BLOCKING_SYNC + + class GraphConditionalType(StrEnum): """Conditional node flavor for :class:`~cuda.core.graph.GraphBuilder`. diff --git a/cuda_core/docs/source/api_private.rst b/cuda_core/docs/source/api_private.rst index 80675799c07..de161c96ae0 100644 --- a/cuda_core/docs/source/api_private.rst +++ b/cuda_core/docs/source/api_private.rst @@ -35,6 +35,7 @@ CUDA runtime typing.ProcessStateType typing.ReadModeType typing.SourceCodeType + typing.SynchronizationPolicyType typing.VirtualMemoryAccessType typing.VirtualMemoryAllocationType typing.VirtualMemoryGranularityType diff --git a/cuda_core/tests/test_stream.py b/cuda_core/tests/test_stream.py index 55f34bbc9ec..28af8167868 100644 --- a/cuda_core/tests/test_stream.py +++ b/cuda_core/tests/test_stream.py @@ -469,3 +469,90 @@ def test_created_stream_keeps_its_own_context(): finally: dev0.set_current() stream.close() + + +# ============================================================================ +# Stream synchronization policy tests +# ============================================================================ + + +@pytest.mark.parametrize( + "policy", + [ + "AUTO", + "SPIN", + "YIELD", + "BLOCKING_SYNC", + ], +) +def test_stream_synchronization_policy_set_get(init_cuda, policy): + from cuda.core.typing import SynchronizationPolicyType + + stream = Device().create_stream() + enum_policy = SynchronizationPolicyType[policy] + stream.synchronization_policy = enum_policy + assert stream.synchronization_policy == enum_policy + stream.close() + + +@pytest.mark.parametrize("invalid_policy", ["spin", -1, 99]) +def test_stream_synchronization_policy_invalid(init_cuda, invalid_policy): + stream = Device().create_stream() + with pytest.raises((TypeError, ValueError)): + stream.synchronization_policy = invalid_policy + stream.close() + + +def test_stream_synchronization_policy_default(init_cuda): + from cuda.core.typing import SynchronizationPolicyType + + stream = Device().create_stream() + assert stream.synchronization_policy == SynchronizationPolicyType.AUTO + stream.close() + + +def test_stream_options_synchronization_policy(init_cuda): + from cuda.core.typing import SynchronizationPolicyType + + stream = Device().create_stream( + options=StreamOptions(synchronization_policy=SynchronizationPolicyType.YIELD) + ) + assert stream.synchronization_policy == SynchronizationPolicyType.YIELD + stream.close() + + +@pytest.mark.parametrize( + "policy", + [ + "AUTO", + "SPIN", + "YIELD", + "BLOCKING_SYNC", + ], +) +@pytest.mark.agent_authored(model="composer-2.5") +def test_stream_synchronization_policy_launch_and_sync(init_cuda, policy): + """Stream-level sync policy is applied via cuStreamSetAttribute and survives launch+sync.""" + import cuda.bindings + from cuda.core import Device, LaunchConfig, Program, ProgramOptions, launch + from cuda.core._utils.version import driver_version + from cuda.core.typing import ObjectCodeFormatType, SourceCodeType, SynchronizationPolicyType + + if int(cuda.bindings.__version__.split(".")[0]) >= 13 and driver_version()[0] < 13: + pytest.skip("CUDA 13 bindings produce modules incompatible with CUDA 12 drivers") + + dev = Device() + dev.set_current() + stream = dev.create_stream() + stream.synchronization_policy = SynchronizationPolicyType[policy] + + code = 'extern "C" __global__ void noop() {}' + arch = "".join(f"{i}" for i in dev.compute_capability) + program = Program(code, SourceCodeType.CXX, options=ProgramOptions(arch=f"sm_{arch}")) + mod = program.compile(ObjectCodeFormatType.CUBIN) + ker = mod.get_kernel("noop") + + launch(stream, LaunchConfig(grid=1, block=1), ker) + stream.sync() + assert stream.synchronization_policy == SynchronizationPolicyType[policy] + stream.close() From e18dfb2f7aecbe1b870d3838cba7f59f2f18b3e3 Mon Sep 17 00:00:00 2001 From: Omar Atie Date: Mon, 17 Aug 2026 18:20:10 -0700 Subject: [PATCH 15/17] fix(cuda.core): declare sync_policy cdef at function scope in Stream._init Cython rejects cdef declarations inside conditional blocks. Co-authored-by: Cursor --- cuda_core/cuda/core/_stream.pyx | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/cuda_core/cuda/core/_stream.pyx b/cuda_core/cuda/core/_stream.pyx index 1171d536c13..cd224d147c1 100644 --- a/cuda_core/cuda/core/_stream.pyx +++ b/cuda_core/cuda/core/_stream.pyx @@ -206,6 +206,7 @@ cdef class Stream: # TODO: we might want to consider memoizing high/low per CUDA context and avoid this call cdef int high, low cdef cydriver.CUresult res_code + cdef int sync_policy with nogil: res_code = cydriver.cuCtxGetStreamPriorityRange(&high, &low) if res_code != cydriver.CUresult.CUDA_SUCCESS: @@ -246,7 +247,7 @@ cdef class Stream: self._priority = prio if sync_policy_opt is not None: validated = _validate_synchronization_policy(sync_policy_opt) - cdef int sync_policy = int(validated) + sync_policy = int(validated) with nogil: Stream_apply_synchronization_policy(self, sync_policy) self._synchronization_policy = sync_policy From d221efb71ad4da613d76fba0b80d0b5c6684ca5d Mon Sep 17 00:00:00 2001 From: Omar Atie Date: Mon, 17 Aug 2026 18:54:12 -0700 Subject: [PATCH 16/17] fix(cuda.core): treat unset stream sync policy query as AUTO cuStreamGetAttribute can return -1 when synchronization_policy was never explicitly set on a new stream. Map that sentinel to CU_SYNC_POLICY_AUTO so the documented default is returned instead of raising ValueError. Co-authored-by: Cursor --- cuda_core/cuda/core/_stream.pyx | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/cuda_core/cuda/core/_stream.pyx b/cuda_core/cuda/core/_stream.pyx index cd224d147c1..066dbdfe6f5 100644 --- a/cuda_core/cuda/core/_stream.pyx +++ b/cuda_core/cuda/core/_stream.pyx @@ -87,6 +87,13 @@ cdef int Stream_query_synchronization_policy(Stream self) except?-1 nogil: return value.syncPolicy +cdef int Stream_normalize_queried_sync_policy(int policy) nogil: + # cuStreamGetAttribute may return -1 when the attribute was never set. + if policy < 0: + return cydriver.CUsynchronizationPolicy_enum.CU_SYNC_POLICY_AUTO + return policy + + cdef void Stream_apply_synchronization_policy(Stream self, int policy) except * nogil: cdef cydriver.CUstreamAttrValue value value.syncPolicy = policy @@ -325,6 +332,7 @@ cdef class Stream: if self._synchronization_policy == INT32_MIN or Stream_is_default_token(self): with nogil: policy = Stream_query_synchronization_policy(self) + policy = Stream_normalize_queried_sync_policy(policy) if not Stream_is_default_token(self): self._synchronization_policy = policy else: From a4fc6c11fb60b99b1519a96140dc731a353bddf3 Mon Sep 17 00:00:00 2001 From: Omar Atie Date: Mon, 17 Aug 2026 18:56:44 -0700 Subject: [PATCH 17/17] style(cuda.core): apply pre-commit ruff format and regenerate _stream.pyi Fix pre-commit.ci failures on ruff-format and stubgen-pyx-cuda-core. Co-authored-by: Cursor --- cuda_core/cuda/core/_stream.pyi | 10 ++++++++-- cuda_core/tests/test_stream.py | 4 +--- 2 files changed, 9 insertions(+), 5 deletions(-) diff --git a/cuda_core/cuda/core/_stream.pyi b/cuda_core/cuda/core/_stream.pyi index b094cf67354..52a87299fbb 100644 --- a/cuda_core/cuda/core/_stream.pyi +++ b/cuda_core/cuda/core/_stream.pyi @@ -27,6 +27,8 @@ class StreamOptions: higher priority. (Default to lowest priority) synchronization_policy : SynchronizationPolicyType | None, optional CPU wait policy applied when synchronizing this stream from the host. + Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY`` via + ``cuStreamSetAttribute``. (Default to None, leaving the driver default) """ nonblocking: cython.bint = True @@ -115,10 +117,14 @@ class Stream: @property def synchronization_policy(self): - """Return the stream's CPU wait policy for host-side synchronization.""" + """Return the stream's CPU wait policy for host-side synchronization. + + Maps to ``CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY``, queried via + ``cuStreamGetAttribute``. + """ @synchronization_policy.setter - def synchronization_policy(self, policy) -> None: + def synchronization_policy(self, policy: object): """Set the stream's CPU wait policy for subsequent work on this stream.""" def sync(self) -> None: diff --git a/cuda_core/tests/test_stream.py b/cuda_core/tests/test_stream.py index 28af8167868..a89d8c963e1 100644 --- a/cuda_core/tests/test_stream.py +++ b/cuda_core/tests/test_stream.py @@ -514,9 +514,7 @@ def test_stream_synchronization_policy_default(init_cuda): def test_stream_options_synchronization_policy(init_cuda): from cuda.core.typing import SynchronizationPolicyType - stream = Device().create_stream( - options=StreamOptions(synchronization_policy=SynchronizationPolicyType.YIELD) - ) + stream = Device().create_stream(options=StreamOptions(synchronization_policy=SynchronizationPolicyType.YIELD)) assert stream.synchronization_policy == SynchronizationPolicyType.YIELD stream.close()