feat(cuda.core): add synchronization_policy to LaunchConfig

[atiaomar1978-hub]

Summary

• Adds LaunchConfig.synchronization_policy to expose CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY from cuda.core, closing [FEA]: Add support for CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY in LaunchConfig #2628.
• Introduces SynchronizationPolicyType in cuda.core.typing (AUTO, SPIN, YIELD, BLOCKING_SYNC) aligned with cuda.bindings.driver.CUsynchronizationPolicy.
• Wires the attribute through both LaunchConfig._to_native_launch_config() and the standalone _to_native_launch_config() helper used in tests.
• Adds unit and mapping tests in tests/test_launcher.py, plus a GPU launch smoke test (skipped when CUDA 13 bindings run against a CUDA 12 driver).

Motivation

Pool-backed and other LaunchConfig launch attributes already route through cuLaunchKernelEx. Users should be able to set per-launch CPU wait policy without dropping to raw cuda.bindings.driver APIs. This follows the incremental LaunchConfig attribute coverage tracked under #496 (same approach as programmatic_stream_serialization).

Example

from cuda.core import Device, LaunchConfig, Program, ProgramOptions, launch
from cuda.core.typing import ObjectCodeFormatType, SourceCodeType, SynchronizationPolicyType

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)
mod = Program(code, SourceCodeType.CXX, options=ProgramOptions(arch=f"sm_{arch}")).compile(ObjectCodeFormatType.CUBIN)
ker = mod.get_kernel("noop")

config = LaunchConfig(grid=1, block=1, synchronization_policy=SynchronizationPolicyType.SPIN)
launch(stream, config, ker)
stream.sync()

Test plan

• pytest cuda_core/tests/test_launcher.py -k synchronization_policy -v
• Contributor GPU validation on RunPod (RTX A5000 community, $0.16/hr)

RunPod GPU report (contributor-run)

Item	Detail
GPU	NVIDIA RTX A5000 (community)
Branch	feat/launch-config-sync-policy-2628
Build	editable cuda_core, CUDA 13 headers/bindings
Command	pytest tests/test_launcher.py -k synchronization_policy -v
Result	10 passed, 4 skipped, 0 failed

Mapping tests verify CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY for all four policies, invalid-input rejection, default None, and combination with is_cooperative. The four GPU launch smoke tests were skipped on this host because the RunPod image exposes a CUDA 12 driver while the build used CUDA 13 bindings (cubin load is incompatible); upstream CI with matched driver/toolkit should cover the launch path.

Closes #2628.

Made with Cursor

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[atiaomar1978-hub]

GPU validation report (RunPod, contributor-run)

Validated this change on real NVIDIA hardware outside upstream CI.

Environment

• GPU: NVIDIA RTX A5000 (RunPod community cloud, $0.16/hr)
• Image: runpod/pytorch:2.4.0-py3.11-cuda12.4.1-devel-ubuntu22.04
• Branch: feat/launch-config-sync-policy-2628
• Build: editable cuda_core with CUDA 13 headers/bindings (cuda-toolkit==13.*, CUDA_CORE_BUILD_MAJOR=13)

Command

pytest cuda_core/tests/test_launcher.py -k "synchronization_policy" -v

Results: 10 passed, 4 skipped, 0 failed

Test	Result
test_to_native_launch_config_synchronization_policy (AUTO/SPIN/YIELD/BLOCKING_SYNC + driver enum input)	passed (5)
test_launch_config_synchronization_policy_default	passed
test_launch_config_synchronization_policy_invalid (3 cases)	passed (3)
test_to_native_launch_config_synchronization_policy_with_cooperative	passed
test_launch_with_synchronization_policy (4 policies)	skipped

The four launch smoke tests were skipped because this host exposes a CUDA 12 driver while the build used CUDA 13 bindings (CUDA 13 cubin is incompatible with CUDA 12 drivers). Mapping and validation tests fully cover the issue acceptance criteria; upstream CI with matched driver/toolkit should exercise the launch path.

Cost: ~$0.04 total RunPod spend for all validation runs. Pod terminated; no active billing.

[atiaomar1978-hub]

Branch update: history preserved, main merged, no force-push

Force-push fix

The earlier force-push (9f4bb0a → 0008aca) rewrote commit SHAs when re-signing/rebasing. Restored the original lineage without another force-push:

• Merged pre-force-push commits (ce87db27…9f4bb0a) back in → c5d7b86
• Fixed duplicate import from that merge → f3d8d67
• Normal git push only (no --force)

The PR graph now shows both the original and rewritten commits.

Main merge

upstream/main is included (db2801873b — already up to date). Signed confirmation commit: 1e48c4e.

Going forward

• Merge commits from main, not rebase/force-push
• New commits: SSH-signed + DCO Signed-off-by

[atiaomar1978-hub]

Note on commit signing

The two history-restore commits (c5d7b86, f3d8d67) were pushed before SSH commit signing was configured on this machine, so they carry DCO Signed-off-by but not SSH signatures.

Re-signing those commits in place would require rewriting their SHAs and force-pushing the branch — which we're deliberately avoiding to preserve PR history per maintainer guidance.

All new commits going forward on this branch will be SSH-signed with atiaomar1978@gmail.com (plus DCO), starting with 1e48c4e.