New lowpp layer for nvrtc

cuda_bindings/AGENTS.md

@@ -42,6 +42,11 @@ subpackage in the `cuda-python` monorepo.
 - **Examples**: example coverage is pytest-based under `examples/`.
 - **Benchmarks**: run with `pytest --benchmark-only benchmarks/` when needed.
 
+The `legacy_tests` subdirectory tests the old pre-v2 APIs of `driver`, `runtime`
+and `nvrtc`.  These test files should not be added to, only updated when
+necessary to fix test failures.  The canonical set of tests are those outside of
+the `legacy_tests` subdirectory.
+
 ## Build and environment notes
 
 - `CUDA_HOME` or `CUDA_PATH` must point to a valid CUDA Toolkit for source

cuda_bindings/build_hooks.py

@@ -196,7 +196,7 @@ def _cleanup_dst_files():
 
     # Build extension list
     extensions = []
-    cuda_bindings_files = glob.glob("cuda/bindings/*.pyx")
+    cuda_bindings_files = glob.glob("cuda/bindings/*.pyx") + glob.glob("cuda/bindings/_v2/*.pyx")
     if sys.platform == "win32":
         cuda_bindings_files = [f for f in cuda_bindings_files if "cufile" not in f]
 

cuda_bindings/cuda/bindings/_example_helpers/common.py

@@ -9,8 +9,8 @@
 
 from cuda import pathfinder
 from cuda.bindings import driver as cuda
-from cuda.bindings import nvrtc
 from cuda.bindings import runtime as cudart
+from cuda.bindings._v2 import nvrtc
 
 from .helper_cuda import check_cuda_errors
 
@@ -44,7 +44,7 @@ def __init__(self, code, dev_id):
                 requirement_not_met(f'pathfinder.find_nvidia_header_directory("{libname}") returned None')
             include_dirs.append(hdr_dir)
 
-        prog = check_cuda_errors(nvrtc.nvrtcCreateProgram(str.encode(code), b"sourceCode.cu", 0, None, None))
+        prog = nvrtc.create_program(str.encode(code), b"sourceCode.cu")
 
         # Initialize CUDA
         check_cuda_errors(cudart.cudaFree(0))
@@ -55,7 +55,7 @@ def __init__(self, code, dev_id):
         minor = check_cuda_errors(
             cudart.cudaDeviceGetAttribute(cudart.cudaDeviceAttr.cudaDevAttrComputeCapabilityMinor, dev_id)
         )
-        _, nvrtc_minor = check_cuda_errors(nvrtc.nvrtcVersion())
+        _, nvrtc_minor = nvrtc.version()
         use_cubin = nvrtc_minor >= 1
         prefix = "sm" if use_cubin else "compute"
         arch_arg = bytes(f"--gpu-architecture={prefix}_{major}{minor}", "ascii")
@@ -70,25 +70,19 @@ def __init__(self, code, dev_id):
             opts.append(f"--include-path={inc_dir}".encode())
 
         try:
-            check_cuda_errors(nvrtc.nvrtcCompileProgram(prog, len(opts), opts))
-        except RuntimeError as err:
-            log_size = check_cuda_errors(nvrtc.nvrtcGetProgramLogSize(prog))
-            log = b" " * log_size
-            check_cuda_errors(nvrtc.nvrtcGetProgramLog(prog, log))
+            nvrtc.compile_program(prog, opts)
+        except nvrtc.NvrtcError as err:
+            log = nvrtc.get_program_log(prog)
             import sys
 
             print(log.decode(), file=sys.stderr)  # noqa: T201
             print(err, file=sys.stderr)  # noqa: T201
             sys.exit(1)
 
         if use_cubin:
-            data_size = check_cuda_errors(nvrtc.nvrtcGetCUBINSize(prog))
-            data = b" " * data_size
-            check_cuda_errors(nvrtc.nvrtcGetCUBIN(prog, data))
+            data = nvrtc.get_cubin(prog)
         else:
-            data_size = check_cuda_errors(nvrtc.nvrtcGetPTXSize(prog))
-            data = b" " * data_size
-            check_cuda_errors(nvrtc.nvrtcGetPTX(prog, data))
+            data = nvrtc.get_ptx(prog)
 
         self.module = check_cuda_errors(cuda.cuModuleLoadData(np.char.array(data)))
 

cuda_bindings/cuda/bindings/_v2/nvrtc.pxd

@@ -0,0 +1,60 @@
+# SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+#
+# SPDX-License-Identifier: Apache-2.0
+#
+# This code was automatically generated across versions from 12.9.0 to 13.3.0. Do not modify it directly.
+
+# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=223716a3d6961dfe7231f2c88e76a9ab4c0a8d888cc4bf3e889bfbcbf8580346
+from libc.stdint cimport intptr_t
+
+from ..cynvrtc cimport *
+
+
+###############################################################################
+# Types
+###############################################################################
+
+ctypedef nvrtcProgram Program
+
+
+###############################################################################
+# Enum
+###############################################################################
+
+ctypedef nvrtcResult _Result
+
+
+###############################################################################
+# Functions
+###############################################################################
+
+cpdef intptr_t create_program(bytes src, name, headers=*, include_names=*) except? 0
+cpdef compile_program(intptr_t prog, options=*)
+cpdef set_flow_callback(intptr_t prog, intptr_t callback, intptr_t payload)
+cpdef bytes get_lowered_name(intptr_t prog, bytes name_expression)
+cpdef install_bundled_headers(bytes install_path, unsigned int flags)
+cpdef tuple get_bundled_headers_info()
+cpdef remove_bundled_headers(bytes install_path)
+
+cpdef str get_error_string(int result)
+cpdef tuple version()
+cpdef int get_num_supported_archs() except? -1
+cpdef object get_supported_archs()
+cpdef destroy_program(intptr_t prog)
+cpdef size_t get_ptx_size(intptr_t prog) except? 0
+cpdef bytes get_ptx(intptr_t prog)
+cpdef size_t get_cubin_size(intptr_t prog) except? 0
+cpdef bytes get_cubin(intptr_t prog)
+cpdef size_t get_ltoir_size(intptr_t prog) except? 0
+cpdef bytes get_ltoir(intptr_t prog)
+cpdef size_t get_optix_ir_size(intptr_t prog) except? 0
+cpdef bytes get_optix_ir(intptr_t prog)
+cpdef size_t get_program_log_size(intptr_t prog) except? 0
+cpdef bytes get_program_log(intptr_t prog)
+cpdef add_name_expression(intptr_t prog, name_expression)
+cpdef size_t get_pch_heap_size() except? 0
+cpdef set_pch_heap_size(size_t size)
+cpdef int get_pch_create_status(intptr_t prog) except? -1
+cpdef size_t get_pch_heap_size_required(intptr_t prog) except? 0
+cpdef size_t get_tile_ir_size(intptr_t prog) except? 0
+cpdef bytes get_tile_ir(intptr_t prog)