Skip to content

Prepare for cuda_bindings 3.15 release - #3368

Merged
rwgk merged 12 commits into
NVIDIA:mainfrom
mdboom:prepare-3.15
Nov 5, 2026
Merged

rwgk merged 12 commits into
NVIDIA:mainfrom
mdboom:prepare-3.15

Conversation

@mdboom

@mdboom mdboom commented Nov 4, 2026 •

Copy link
Copy Markdown
Contributor

Prepare cuda-bindings and cuda-python 13.5.0 for CUDA Toolkit 13.5, including refreshed bindings and API reference documentation.

  • Move the build Toolkit, current-minor CI rows, and cuda-bindings/cuda-core Pixi configurations to 13.5, retaining older-runtime coverage.
  • Keep cuda-core’s published bindings floor unchanged. During the Toolkit minor transition, regular core CI uses bindings built against matching headers instead of the older published floor.
  • Relax the Toolkit wheel minor pin for local tests of the build’s own minor, allowing testing before matching Toolkit wheels are published.
  • Add PTX 9.5 compatibility mapping and recognize Windows CUPTI’s cupti64_13.dll.
  • Add regression coverage for graph-node layouts, embedded batch parameters, aliasing, read-only views, and wrapper lifetimes.
  • Update the 13.5 release notes, documentation version navigation, Cython cache ignores, and file-check exclusions.

Follow-up nightly CI and test-dependency refinements are tracked in #3375.

@copy-pr-bot

copy-pr-bot Bot commented Nov 4, 2026

Copy link
Copy Markdown
Contributor

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

@github-actions github-actions Bot added cuda.bindings Everything related to the cuda.bindings module cuda.pathfinder Everything related to the cuda.pathfinder module labels Nov 4, 2026
@mdboom

mdboom commented Nov 4, 2026

Copy link
Copy Markdown
Contributor Author

/ok to test

@github-actions github-actions Bot added CI/CD CI/CD infrastructure cuda.core Everything related to the cuda.core module labels Nov 4, 2026
@mdboom

mdboom commented Nov 4, 2026

Copy link
Copy Markdown
Contributor Author

/ok to test

@mdboom

mdboom commented Nov 4, 2026

Copy link
Copy Markdown
Contributor Author

/ok to test

@github-actions

github-actions Bot commented Nov 4, 2026

Copy link
Copy Markdown
Contributor

@mdboom

mdboom commented Nov 4, 2026

Copy link
Copy Markdown
Contributor Author

/ok to test

@mdboom

mdboom commented Nov 4, 2026

Copy link
Copy Markdown
Contributor Author

/ok to test

@mdboom
mdboom marked this pull request as ready for review November 5, 2026 01:33
@mdboom mdboom added this to the cuda.bindings 13.5.0 & 12.9.10 milestone Nov 5, 2026
@coderabbitai

coderabbitai Bot commented Nov 5, 2026 •

Copy link
Copy Markdown

Explain Change →

Note

Reviews paused

It looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the reviews.auto_review.auto_pause_after_reviewed_commits setting.

Use the following commands to manage reviews:

  • @coderabbitai resume to resume automatic reviews.
  • @coderabbitai review to trigger a single review.

Use the checkboxes below for quick actions:

  • Resume reviews
  • Trigger review
Summary

Summary by CodeRabbit

  • New Features
    • Updated CUDA bindings for CUDA Toolkit 13.5, including batch memory-copy graph nodes, fabric-copy operations, tensor-map allocation classes, and new NVML capabilities.
    • Added Python 3.15 support and PTX ISA 9.5 version mapping. Python 3.10 is no longer supported.
    • Improved Windows detection of CUDA 13 CUPTI libraries.
  • Bug Fixes
    • Improved compatibility when testing CUDA Toolkit versions and selecting bindings.
    • Improved callback registration safety under free-threaded Python and excluded unpopulated NVML field queries from valid results.
  • Documentation
    • Updated CUDA 13.5 API documentation and release notes, including compatibility details and known issues.

This pull request updates build and test configuration to CUDA 13.5. It adds driver, runtime, and NVML bindings, tests, documentation, and release metadata for the updated toolkit.

Changes: CUDA 13.5 upgrade

Layer / File(s) Summary
Build and CI configuration
.github/workflows/*, ci/*, cuda_bindings_13/pixi.toml, cuda_core/pixi.toml, cuda_core/AGENTS.md, .gitignore, .pre-commit-config.yaml, .spdx-ignore
Build settings and CI matrices move to CUDA 13.5. CI exports build-version components and adjusts floor-binding wheel selection and toolkit constraints. Linux and Windows workflows download artifacts for core tests using floor bindings.
Driver declarations and wrappers
cuda_bindings_13/cuda/bindings/_internal/driver*, cuda_bindings_13/cuda/bindings/_v2/driver.*, cuda_bindings_13/cuda/bindings/cy*driver.*, cuda_bindings_13/cuda/bindings/driver.pxd, cuda_bindings_13/docs/source/module/driver.rst
The bindings add batch memcpy and fabric-copy graph-node support, tensor-map allocation-class APIs, and CUDA 13.5 driver enums and attributes. Driver symbol lookup, wrappers, and graph parameter handling are updated.
Runtime API and batch memcpy graph parameters
cuda_bindings_13/cuda/bindings/cyruntime.*, cuda_bindings_13/cuda/bindings/runtime.*, cuda_bindings_13/docs/source/module/runtime.rst, cuda_bindings_13/tests/test_cuda.py, cuda_bindings_13/tests/test_cudart.py
The runtime adds batch memcpy graph-node parameters, locality-domain resource options, and related enums and attributes. Tests cover graph-parameter views, assignment, read-only behavior, and the locality-domain sentinel.
NVML declarations and wrappers
cuda_bindings_13/cuda/bindings/_internal/nvml*, cuda_bindings_13/cuda/bindings/cynvml.*, cuda_bindings_13/cuda/bindings/nvml.*
The bindings add GPU uptime, confidential-compute settings v2, NVLink information v3, remap-table, and NVLE structures and APIs. The NVML documentation describes MPS locality-domain behavior.
Version metadata and release support
cuda_bindings_13/cuda/bindings/*, cuda_bindings_13/docs/*, cuda_python/docs/source/release/*, cuda_pathfinder/cuda/pathfinder/_dynamic_libs/descriptor_catalog.py
Generated-file metadata advances to 13.5 across binding components. The changes add PTX 9.5 mapping, cufile PanFS support, CUDA 13.5 release documentation, and a Windows CUPTI DLL candidate.

Possibly related issues

Suggested reviewers: rwgk

Change: Feature

  • Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 9

Nitpick comments: 1

cuda_bindings_13/cuda/bindings/nvml.pyx: 1

38410-38577: Trivial · Maintainability & Code Quality · Quick win

important: Add targeted tests for the new NVML APIs.

The changed tests cover CUDA driver graph wrappers, not these NVML functions or struct wrappers. Add round-trip coverage through from_buffer, from_data, and from_ptr. Exercise the shared remapTabSize behavior of fla_remap_tab_addr and gpa_remap_tab_addr, and validate the 32-byte nvle_key requirement. Hardware-dependent calls can skip on NotSupportedError.

Sources: Coding guidelines, Learnings

Review info

Run configuration

Commits

Reviewing files that changed from the base of the PR and between f707a1f and f00d958.

Files selected for processing: 91

  • .github/workflows/test-wheel-linux.yml
  • .github/workflows/test-wheel-windows.yml
  • .gitignore
  • .pre-commit-config.yaml
  • .spdx-ignore
  • ci/test-matrix.yml
  • ci/tools/env-vars
  • ci/tools/run-tests
  • ci/versions.yml
  • cuda_bindings_13/cuda/bindings/_internal/cudla.pxd
  • cuda_bindings_13/cuda/bindings/_internal/cudla_linux.pyx
  • cuda_bindings_13/cuda/bindings/_internal/cudla_windows.pyx
  • cuda_bindings_13/cuda/bindings/_internal/cufile.pxd
  • cuda_bindings_13/cuda/bindings/_internal/cufile_linux.pyx
  • cuda_bindings_13/cuda/bindings/_internal/driver.pxd
  • cuda_bindings_13/cuda/bindings/_internal/driver_linux.pyx
  • cuda_bindings_13/cuda/bindings/_internal/driver_windows.pyx
  • cuda_bindings_13/cuda/bindings/_internal/nvfatbin.pxd
  • cuda_bindings_13/cuda/bindings/_internal/nvfatbin_linux.pyx
  • cuda_bindings_13/cuda/bindings/_internal/nvfatbin_windows.pyx
  • cuda_bindings_13/cuda/bindings/_internal/nvjitlink.pxd
  • cuda_bindings_13/cuda/bindings/_internal/nvjitlink_linux.pyx
  • cuda_bindings_13/cuda/bindings/_internal/nvjitlink_windows.pyx
  • cuda_bindings_13/cuda/bindings/_internal/nvml.pxd
  • cuda_bindings_13/cuda/bindings/_internal/nvml_linux.pyx
  • cuda_bindings_13/cuda/bindings/_internal/nvml_windows.pyx
  • cuda_bindings_13/cuda/bindings/_internal/nvrtc.pxd
  • cuda_bindings_13/cuda/bindings/_internal/nvrtc_linux.pyx
  • cuda_bindings_13/cuda/bindings/_internal/nvrtc_windows.pyx
  • cuda_bindings_13/cuda/bindings/_internal/nvvm.pxd
  • cuda_bindings_13/cuda/bindings/_internal/nvvm_linux.pyx
  • cuda_bindings_13/cuda/bindings/_internal/nvvm_windows.pyx
  • cuda_bindings_13/cuda/bindings/_internal/runtime.pxd
  • cuda_bindings_13/cuda/bindings/_internal/runtime_linux.pyx
  • cuda_bindings_13/cuda/bindings/_internal/runtime_ptds.pxd
  • cuda_bindings_13/cuda/bindings/_internal/runtime_ptds_linux.pyx
  • cuda_bindings_13/cuda/bindings/_internal/runtime_ptds_windows.pyx
  • cuda_bindings_13/cuda/bindings/_internal/runtime_windows.pyx
  • cuda_bindings_13/cuda/bindings/_v2/driver.pxd
  • cuda_bindings_13/cuda/bindings/_v2/driver.pyx
  • cuda_bindings_13/cuda/bindings/_v2/nvrtc.pxd
  • cuda_bindings_13/cuda/bindings/_v2/nvrtc.pyx
  • cuda_bindings_13/cuda/bindings/cudla.pxd
  • cuda_bindings_13/cuda/bindings/cudla.pyx
  • cuda_bindings_13/cuda/bindings/cufile.pxd
  • cuda_bindings_13/cuda/bindings/cufile.pyx
  • cuda_bindings_13/cuda/bindings/cycudla.pxd
  • cuda_bindings_13/cuda/bindings/cycudla.pyx
  • cuda_bindings_13/cuda/bindings/cycufile.pxd
  • cuda_bindings_13/cuda/bindings/cycufile.pyx
  • cuda_bindings_13/cuda/bindings/cydriver.pxd
  • cuda_bindings_13/cuda/bindings/cydriver.pyx
  • cuda_bindings_13/cuda/bindings/cynvfatbin.pxd
  • cuda_bindings_13/cuda/bindings/cynvfatbin.pyx
  • cuda_bindings_13/cuda/bindings/cynvjitlink.pxd
  • cuda_bindings_13/cuda/bindings/cynvjitlink.pyx
  • cuda_bindings_13/cuda/bindings/cynvml.pxd
  • cuda_bindings_13/cuda/bindings/cynvml.pyx
  • cuda_bindings_13/cuda/bindings/cynvrtc.pxd
  • cuda_bindings_13/cuda/bindings/cynvrtc.pyx
  • cuda_bindings_13/cuda/bindings/cynvvm.pxd
  • cuda_bindings_13/cuda/bindings/cynvvm.pyx
  • cuda_bindings_13/cuda/bindings/cyruntime.pxd
  • cuda_bindings_13/cuda/bindings/cyruntime.pyx
  • cuda_bindings_13/cuda/bindings/driver.pxd
  • cuda_bindings_13/cuda/bindings/driver.pyx
  • cuda_bindings_13/cuda/bindings/nvfatbin.pxd
  • cuda_bindings_13/cuda/bindings/nvfatbin.pyx
  • cuda_bindings_13/cuda/bindings/nvjitlink.pxd
  • cuda_bindings_13/cuda/bindings/nvjitlink.pyx
  • cuda_bindings_13/cuda/bindings/nvml.pxd
  • cuda_bindings_13/cuda/bindings/nvml.pyx
  • cuda_bindings_13/cuda/bindings/nvrtc.pxd
  • cuda_bindings_13/cuda/bindings/nvrtc.pyx
  • cuda_bindings_13/cuda/bindings/nvvm.pxd
  • cuda_bindings_13/cuda/bindings/nvvm.pyx
  • cuda_bindings_13/cuda/bindings/runtime.pxd
  • cuda_bindings_13/cuda/bindings/runtime.pyx
  • cuda_bindings_13/cuda/bindings/utils/_ptx_utils.py
  • cuda_bindings_13/docs/nv-versions.json
  • cuda_bindings_13/docs/source/module/driver.rst
  • cuda_bindings_13/docs/source/module/nvrtc.rst
  • cuda_bindings_13/docs/source/module/runtime.rst
  • cuda_bindings_13/docs/source/release/13.5.0-notes.rst
  • cuda_bindings_13/pixi.toml
  • cuda_bindings_13/tests/test_cuda.py
  • cuda_bindings_13/tests/test_graph_node_params_lifetime.py
  • cuda_core/AGENTS.md
  • cuda_core/pixi.toml
  • cuda_pathfinder/cuda/pathfinder/_dynamic_libs/descriptor_catalog.py
  • cuda_python/docs/source/release/13.5.0-notes.rst
Files not fully reviewed because their changes were unavailable: 1

  • cuda_bindings_13/cuda/bindings/driver.pyx

Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review.

Comment on lines +25805 to +25806
cdef MemcpyBatchNodeParamsV1 val_ = val
_cyb_memcpy(<void *>&(self._ptr[0].memcpyBatch), <void *>(val_._get_ptr()), sizeof(CUDA_MEMCPY_BATCH_NODE_PARAMS_V1) * 1)

@coderabbitai coderabbitai Bot Nov 5, 2026 •

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Major · Stability & Availability · Quick win

important: Reject None before calling _get_ptr().

GraphNodeParams().memcpy_batch = None passes the typed assignment, then calls a Cython method on None. The fabric_cpy_batch setter has the same failure. This can crash the interpreter instead of raising TypeError; the build does not enable nonecheck. (docs.cython.org)

Add explicit non-null validation in the generator, regenerate both setters, and add regression tests.

Also applies to: 25821-25822

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Probably a real bug, but not a blocker. Can address during the final refinement of the _v2 APIs.

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@mdboom suggestion: Defer this fix to the final _v2 API refinement rather than block this PR. Would you like me to open a follow-up GitHub issue?

You are interacting with an AI system.

Comment on lines +17372 to +17374
out_attrs = [cudaMemcpyAttributes() for _idx in range(self._pvt_ptr[0].numAttrs)]
for _idx in range(len(out_attrs)):
string.memcpy(<cyruntime.cudaMemcpyAttributes*><void_ptr>out_attrs[_idx].getPtr(), &self._pvt_ptr[0].attrs[_idx], sizeof(cyruntime.cudaMemcpyAttributes))

@coderabbitai coderabbitai Bot Nov 5, 2026 •

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Major · Stability & Availability · Quick win

important: Bound attrs reads by its allocation. The public numAttrs setter can set a value larger than the array allocated by attrs. In that state, this getter copies past the array boundary. For example, setting one attribute and then setting numAttrs = 2 makes params.attrs read a second, unallocated entry. Keep the count and allocation consistent, or reject an out-of-range count before reading. Regenerate the binding after correcting its source. (raw.githubusercontent.com)

As per path instructions, “Focus on typed-memoryview safety, GIL and nogil correctness, exception specifications, fused-type behavior, Python reference ownership, and consistency with related declarations and generated type stubs.”

Source: Path instructions

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is a known issue in the legacy (cython_gen) bindings and fixing it would require pervasively breaking existing API (removing the count attributes on structs, in this case numAttrs but there are many examples of this). cybind-generated modules do not have this issue, so will disappear naturally when the _v2 modules ship.

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@mdboom suggestion: I retract the must-fix recommendation for this PR. The bounds risk remains, but your explanation establishes that this is a known cython_gen limitation with an existing migration path through cybind and _v2. No change requested here.


Learnings used

Learnt from: CR
Repo: NVIDIA/cuda-python PR: 0
File: cuda_bindings_13/AGENTS.md:0-0
Timestamp: 2026-10-21T18:37:43.407Z
Learning: Applies to cuda_bindings_13/** : Source excerpt:
## Scope and principles

- **Style**: prioritize correctness and API compatibility over convenience
  wrappers. High-level ergonomics belong in `cuda_core`, not here.
Learnt from: CR
Repo: NVIDIA/cuda-python PR: 0
File: cuda_bindings_13/CLAUDE.md:0-0
Timestamp: 2026-10-31T04:59:49.354Z
Learning: Applies to cuda_bindings_13/** : Source excerpt:
## Scope and principles

- **Style**: prioritize correctness and API compatibility over convenience
  wrappers. High-level ergonomics belong in `cuda_core`, not here.
Learnt from: CR
Repo: NVIDIA/cuda-python PR: 0
File: cuda_bindings_13/CLAUDE.md:0-0
Timestamp: 2026-10-31T04:59:49.354Z
Learning: Source excerpt:
## Generated-source workflow

- **Do not hand-edit generated binding files**: many files under
  `cuda/bindings/` (including `*.pyx` and `*.pxd`) are generated artifacts.
Learnt from: CR
Repo: NVIDIA/cuda-python PR: 0
File: cuda_bindings_13/AGENTS.md:0-0
Timestamp: 2026-10-31T05:00:01.371Z
Learning: Source excerpt:
## Generated-source workflow

- **Do not hand-edit generated binding files**: many files under
  `cuda/bindings/` (including `*.pyx` and `*.pxd`) are generated artifacts.

You are interacting with an AI system.

return self._memcpyBatch
@memcpyBatch.setter
def memcpyBatch(self, memcpyBatch not None : cudaMemcpyBatchNodeParams):
string.memcpy(&self._pvt_ptr[0].memcpyBatch, <void*><void_ptr>memcpyBatch.getPtr(), sizeof(self._pvt_ptr[0].memcpyBatch))

@coderabbitai coderabbitai Bot Nov 5, 2026 •

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Major · Stability & Availability · Quick win

Supported by static analysis

Script executed:

sed -n '17205,17260p' cuda_bindings_13/cuda/bindings/runtime.pyx
sed -n '17960,17985p' cuda_bindings_13/cuda/bindings/runtime.pyx
rg -n '__dealloc__|_dsts|_srcs' cuda_bindings_13/cuda/bindings/runtime.pyx | sed -n '1,40p'

Repository: NVIDIA/cuda-python

Length of output: 4026


Script executed:

set -u
f=cuda_bindings_13/cuda/bindings/runtime.pyx
start=$(rg -n -F 'cdef class cudaMemcpyBatchNodeParams' "$f" | head -1 | cut -d: -f1)
echo "== batch class from line $start =="
sed -n "${start},$((start+190))p" "$f"
echo "== retention pattern in cudaGraphNodeParams setters =="
rg -n -F -e 'self._kernel = ' -e 'self._memcpy = ' -e 'self._memcpyBatch = ' -e 'self._conditional = ' -e 'self._free = ' -e 'self._memset = ' -- "$f" || test "$?" -eq 1
echo "== cudaGraphNodeParams __init__ region =="
gstart=$(rg -n -F 'cdef class cudaGraphNodeParams' "$f" | head -1 | cut -d: -f1)
sed -n "${gstart},$((gstart+60))p" "$f"

Repository: NVIDIA/cuda-python

Length of output: 8335


important: Retain the source wrapper in memcpyBatch.

The memcpyBatch setter copies the native struct but keeps no reference to the source cudaMemcpyBatchNodeParams. The source wrapper frees its attrs allocation in __dealloc__. After the caller drops the source, cudaGraphNodeParams.memcpyBatch.attrs can point to freed memory. A later graph call that reads these parameters can then read freed memory. The source wrapper also owns the dsts, srcs, and sizes helper objects, which can release their backing storage in the same way.

Keep the source wrapper alive for as long as the node parameters hold it. Add a retained attribute in the setter:

Suggested fix

     @memcpyBatch.setter
     def memcpyBatch(self, memcpyBatch not None : cudaMemcpyBatchNodeParams):
+        self._cymemcpyBatch = memcpyBatch
         string.memcpy(&self._pvt_ptr[0].memcpyBatch, <void*><void_ptr>memcpyBatch.getPtr(), sizeof(self._pvt_ptr[0].memcpyBatch))

Declare _cymemcpyBatch with the other cudaGraphNodeParams attributes in runtime.pxd. Apply the change in the binding generator source and regenerate runtime.pyx. Do not edit the generated file directly. Do not assign self._memcpyBatch. The getter depends on it as a view into the node-params struct.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ditto for above.

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Analysis chain

Script executed:

#!/bin/bash
set -euo pipefail
gh api --paginate repos/NVIDIA/cuda-python/pulls/3368/comments --jq '.[] | {id, in_reply_to_id, user: .user.login, path, line, created_at, body}'

Length of output: 36988


@mdboom suggestion: Understood. I will treat this as the same known legacy cython_gen limitation and will not request an ownership change in this PR. The lifetime risk remains. Would you like me to open a follow-up issue?


Learnings used

Learnt from: CR
Repo: NVIDIA/cuda-python PR: 0
File: cuda_bindings_13/CLAUDE.md:0-0
Timestamp: 2026-10-31T04:59:49.354Z
Learning: Source excerpt:
## Generated-source workflow

- **Do not hand-edit generated binding files**: many files under
  `cuda/bindings/` (including `*.pyx` and `*.pxd`) are generated artifacts.
Learnt from: CR
Repo: NVIDIA/cuda-python PR: 0
File: cuda_bindings_13/AGENTS.md:0-0
Timestamp: 2026-10-31T05:00:01.371Z
Learning: Source excerpt:
## Generated-source workflow

- **Do not hand-edit generated binding files**: many files under
  `cuda/bindings/` (including `*.pyx` and `*.pxd`) are generated artifacts.

You are interacting with an AI system.



Treats constraints as a hint, ignoring them if necessary to reach the requested smCount. Lets smCount be a non-multiple of coscheduledSmCount, filling the difference between SM count and already assigned co-scheduled groupings with other SMs. When used with cudaDevSmResourceGroupLocalityDomainId, backfill fills up to the requested smCount using the target locality domain first, then SMs not attributed to any locality domain, then SMs from other locality domains. If no SMs can be found in the requested locality domain, cudaErrorInvalidResourceConfiguration is returned.
Treats constraints as a hint, ignoring them if necessary to reach the requested smCount. Lets smCount be a non-multiple of coscheduledSmCount, filling the difference between SM count and already assigned co-scheduled groupings with other SMs. When used with cudaDevSmResourceGroupLocalityDomainId, backfill fills up to the requested smCount using the target locality domain first, then SMs not attributed to any locality domain, then SMs from other locality domains. If no SMs can be found in the requested locality domain, cudaErrorInvalidResourceConfiguration is returned. When used with ``cudaDevSmResourceGroupLocalityDomainNonPortableId``, backfill follows the same behavior after first consuming all SMs from the target locality domain.

@coderabbitai coderabbitai Bot Nov 5, 2026 •

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Minor · Functional Correctness · Quick win

important: Use the exported flag name in the backfill description.

This text names cudaDevSmResourceGroupLocalityDomainNonPortableId, but the binding exports cudaDevSmResourceGroupLocalityDomainSupplementalSms. The current text directs users to a nonexistent API identifier. Fix the upstream generator input and regenerate this file instead of editing the generated RST directly. (raw.githubusercontent.com)

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is a documentation bug that comes from the upstream headers. We can patch our headers to fix (and report upstream), but I don't think this should hold up the release.

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@mdboom suggestion: This upstream documentation bug should not block the release. My “important” label overstated its impact. The header patch, regeneration, and upstream report can be follow-up work.

Would you like me to open a GitHub issue for that follow-up?

You are interacting with an AI system.

Comment thread cuda_bindings_13/pixi.toml Resolved
params = cuda.CUgraphNodeParams()
batch = cuda.CUDA_MEMCPY_BATCH_NODE_PARAMS_V1()
batch.set_dsts_srcs([10, 20, 30], [40, 50, 60])
params.memcpyBatch = batch

@coderabbitai coderabbitai Bot Nov 5, 2026 •

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Major · Stability & Availability · Heavy lift

Supported by static analysis

Script executed:

sed -n '90,160p' cuda_bindings_13/tests/test_graph_node_params_lifetime.py
rg -n 'memcpyBatch' cuda_bindings_13/cuda/bindings/driver.pyx | head -30

Repository: NVIDIA/cuda-python

Length of output: 3816


Script executed:

rg -n 'class CUDA_MEMCPY_BATCH_NODE_PARAMS_V1|__dealloc__|def set_dsts_srcs|def dsts|def srcs|free\(' cuda_bindings_13/cuda/bindings/driver.pyx | head -120

Repository: NVIDIA/cuda-python

Length of output: 4377


Script executed:

f=cuda_bindings_13/cuda/bindings/driver.pyx
start=$(rg -n '^cdef class CUDA_MEMCPY_BATCH_NODE_PARAMS_V1' "$f" | head -1 | cut -d: -f1 || true)
if [ -z "$start" ]; then start=$(rg -n 'class CUDA_MEMCPY_BATCH_NODE_PARAMS_V1' "$f" | head -1 | cut -d: -f1 || true); fi
echo "class_start=$start"
if [ -n "$start" ]; then
  end=$(awk -v s="$start" 'NR>s && /^(cdef class|class) /{print NR-1; exit}' "$f")
  [ -z "$end" ] && end=$((start+400))
  echo "class_range=$start-$end"
  sed -n "${start},${end}p" "$f" | rg -n -e 'def __dealloc__' -e 'free\(' -e 'def dsts' -e 'def srcs' -e 'def set_dsts_srcs' -e '_dsts' -e '_srcs' -e 'malloc' -e 'PyMem' || test "$?" -eq 1
  sed -n "${start},$((start+40))p" "$f"
else
  echo "class not found"
fi

Repository: NVIDIA/cuda-python

Length of output: 3209


important: Keep memcpyBatch from aliasing arrays owned by batch.

params.memcpyBatch = batch copies the struct bytes with string.memcpy. The copy keeps the dsts and srcs pointers that batch owns. CUDA_MEMCPY_BATCH_NODE_PARAMS_V1_st.__dealloc__ frees both arrays. When batch is deleted, or when batch.set_dsts_srcs(...) replaces the arrays, params.memcpyBatch.dsts and srcs can refer to freed memory. Keeping a reference to batch on params does not cover the set_dsts_srcs case.

Make CUgraphNodeParams own copies of the dsts and srcs arrays. Free those copies on reassignment and when CUgraphNodeParams is deallocated. Update test_memcpy_batch_assignment_is_a_shallow_copy to assert safe reads after batch is deleted. Do not document the aliasing as accepted behavior.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Again, not something we can fix now.

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@mdboom suggestion: Understood. I won’t press for the ownership change in this PR. Would you like me to open a follow-up issue for the aliasing risk?

You are interacting with an AI system.

Comment thread cuda_core/pixi.toml

[feature.cu13.dependencies]
cuda-version = "13.4.*"
cuda-version = "13.5.*"

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Major · Maintainability & Code Quality · Quick win

important: Regenerate cuda_core/pixi.lock. The changed cu13 environment requires CUDA 13.5, but its committed lockfile still contains CUDA 13.4 packages. Locked pixi builds cannot use this manifest and lockfile together. (raw.githubusercontent.com)

As per coding guidelines: “If a PR changes a pixi.toml, update the corresponding pixi.lock in the same PR.”

Source: Coding guidelines

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 1

Review info

Run configuration

Commits

Reviewing files that changed from the base of the PR and between 6a2d754 and e04118b.

Files selected for processing: 3

  • cuda_bindings_13/cuda/bindings/cydriver.pxd
  • cuda_bindings_13/cuda/bindings/cyruntime.pxd
  • cuda_bindings_13/tests/test_cudart.py

Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 10 remain after this review.


cdef enum: RESOURCE_ABI_EXTERNAL_BYTES = 42

cdef const uint32_t CU_DEV_SM_RESOURCE_GROUP_LOCALITY_DOMAIN_ID_ANY = 0xffffffff

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Major · Data Integrity & Integration · Quick win

Supported by static analysis

Script executed:

sed -n '3995,4030p' cuda_bindings_13/cuda/bindings/cydriver.pxd
grep -n 'CUDA_VERSION\|^cdef extern\|^cdef const' cuda_bindings_13/cuda/bindings/cydriver.pxd | tail -20
grep -n 'LocalityDomainIdAny' -B3 cuda_bindings_13/cuda/bindings/cyruntime.pxd
grep -n 'cuda-bindings' cuda_core/pyproject.toml
grep -rln 'cimport.*cydriver\|cimport.*cyruntime' cuda_core/cuda | head

Repository: NVIDIA/cuda-python

Length of output: 2191


Script executed:

# 1) Is cydriver a pxd-only module, or does a runtime .pyx counterpart exist?
fd -HI 'cydriver*' cuda_bindings_13 || test "$?" -eq 1
# 2) Confirm the extern block starting at line 3302 ends before line 4014 (show indentation near the boundary).
awk 'NR>=3300 && NR<=3310 {printf "%5d %s\n", NR, $0}' cuda_bindings_13/cuda/bindings/cydriver.pxd
awk 'NR>=4005 && NR<=4016 {printf "%5d %s\n", NR, $0}' cuda_bindings_13/cuda/bindings/cydriver.pxd
# 3) Find the CUDA_VERSION declaration and any other module-level cdef const in the pxd files.
rg -n -F -- 'CUDA_VERSION' cuda_bindings_13/cuda/bindings/cydriver.pxd || test "$?" -eq 1
rg -n --glob '*.pxd' '^cdef const ' cuda_bindings_13/cuda/bindings || test "$?" -eq 1
# 4) Which cuda_core files cimport cydriver or cyruntime specifically.
rg -n --glob '*.pxd' --glob '*.pyx' -e 'cimport.*cydriver' -e 'cimport.*cyruntime' cuda_core || test "$?" -eq 1
# 5) Inspect the installed Cython export logic, if the sandbox has Cython.
python3 - <<'PY'
import pathlib
try:
    import Cython
except ImportError:
    print("Cython not importable in sandbox")
else:
    root = pathlib.Path(Cython.__file__).parent / "Compiler" / "ModuleNode.py"
    print("Cython version:", Cython.__version__, "file:", root)
    lines = root.read_text().splitlines()
    for number, line in enumerate(lines, 1):
        if "defined_in_pxd" in line or "__pyx_capi__" in line:
            print(f"{number}: {line}")
PY

Repository: NVIDIA/cuda-python

Length of output: 7781


Web query:

Cython 3.2 Compiler/ModuleNode.py generate_c_variable_export_code defined_in_pxd __pyx_capi__ module-level variables

Result:

Sources consulted:
- cython/Cython/Utility/ImportExport.c at 3.2.0 · cython/cython · GitHub — https://github.com/cython/cython/blob/3.2.0/Cython/Utility/ImportExport.c
- cython/Cython/Compiler/ModuleNode.py at 3.2.0 · cython/cython · GitHub — https://github.com/cython/cython/blob/3.2.0/Cython/Compiler/ModuleNode.py
- https://raw.githubusercontent.com/cython/cython/3.2.0/Cython/Compiler/ModuleNode.py
- …and 25 more

Declare the new constant as an enum, not a C variable.

Line 4014 sits at module scope, so Cython exports it as a pxd-defined C variable. An extension built against 13.5 bindings can fail at import when the installed cuda-bindings is 13.4.x, which the cuda-core floor permits. The file already uses anonymous cdef enum constants for values like CUDA_VERSION and CUDA_EGL_INFINITE_TIMEOUT, which do not create an exported variable.

Suggested fix

-cdef const uint32_t CU_DEV_SM_RESOURCE_GROUP_LOCALITY_DOMAIN_ID_ANY = 0xffffffff
+cdef enum: CU_DEV_SM_RESOURCE_GROUP_LOCALITY_DOMAIN_ID_ANY = 0xffffffff
Committable suggestion

IMPORTANT
Review the replacement and indentation, then test before committing.

Suggested change
cdef const uint32_t CU_DEV_SM_RESOURCE_GROUP_LOCALITY_DOMAIN_ID_ANY = 0xffffffff
cdef enum: CU_DEV_SM_RESOURCE_GROUP_LOCALITY_DOMAIN_ID_ANY = 0xffffffff

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Duplicate comments: 1

cuda_bindings_13/cuda/bindings/cydriver.pxd: 1

1847-1847: Major · Potential issue

critical: Emit CUDA_FABRIC_COPY_BATCH_OP_st before this alias. The target struct is declared only at Lines 3244-3250, so Cython resolves this alias before the target type exists and the generated .pxd can fail to compile. Fix the generator ordering for the struct, aliases, and consumers. Do not hand-edit this generated file.

Source: Path instructions

Review info

Run configuration

Commits

Reviewing files that changed from the base of the PR and between e04118b and f2547df.

Files selected for processing: 2

  • cuda_bindings_13/cuda/bindings/cydriver.pxd
  • cuda_bindings_13/cuda/bindings/cyruntime.pxd

Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 10 remain after this review.

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Duplicate comments: 1

cuda_bindings_13/cuda/bindings/cydriver.pxd: 1

1847-1847: Major · Functional Correctness

important: Emit both CUDA_FABRIC_COPY_BATCH_OP_v1 and CUDA_FABRIC_COPY_BATCH_OP after CUDA_FABRIC_COPY_BATCH_OP_st.

CUDA_FABRIC_COPY_BATCH_OP_st is declared at Line 3245, after these aliases. Cython can reject the forward aliases and stop the generated bindings from compiling. Update the generator and regenerate cydriver.pxd; do not hand-edit this generated file. This repeats the unresolved declaration-order issue from the previous review.

Also applies to: 2337-2337

Review info

Run configuration

Commits

Reviewing files that changed from the base of the PR and between f2547df and 91aa3df.

Files selected for processing: 1

  • cuda_bindings_13/cuda/bindings/cydriver.pxd

Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 9 remain after this review.

@rwgk rwgk left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Re-confirming:

codex:

No new release-blocking red flags in MR 615 and PR 3368 at 68e7dcf. Fable’s independent review agrees.

The main fix checks out:

  • The actual CI wheel exports both locality sentinels and the EGL timeout as 4294967295.
  • The unwanted CAPI variables are gone; all six import-cycle checks pass.
  • Completed Linux, ARM, and Windows jobs pass the sentinel and cycle tests.
  • The added NVML documentation names match the exports.

One nonblocking hardening note: the EGL constant declaration can collide with the SDK macro. However, that header combination already failed before these changes, so we haven’t established a newly broken use case.

I’m comfortable proceeding with the agreed plan once the remaining required CI finishes.

@rwgk
rwgk merged commit 3429cba into NVIDIA:main Nov 5, 2026
118 of 119 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

CI/CD CI/CD infrastructure cuda.bindings Everything related to the cuda.bindings module cuda.core Everything related to the cuda.core module cuda.pathfinder Everything related to the cuda.pathfinder module

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants