Early testing of new driver API

.pre-commit-config.yaml

@@ -108,7 +108,7 @@ repos:
     rev: "3e8a8703264a2f4a69428a0aa4dcb512790b2c8c"  # frozen: v6.0.0
     hooks:
     - id: check-added-large-files
-      exclude: cuda_bindings/cuda/bindings/nvml.pyx
+      exclude: cuda_bindings/cuda/bindings/nvml.pyx|cuda_bindings/cuda/bindings/_v2/driver.pyx
     - id: check-case-conflict
     - id: check-docstring-first
     - id: check-merge-conflict

cuda_bindings/cuda/bindings/_example_helpers/common.py

@@ -8,8 +8,8 @@
 import numpy as np
 
 from cuda import pathfinder
-from cuda.bindings import driver as cuda
 from cuda.bindings import runtime as cudart
+from cuda.bindings._v2 import driver as cuda
 from cuda.bindings._v2 import nvrtc
 
 from .helper_cuda import check_cuda_errors
@@ -84,7 +84,9 @@ def __init__(self, code, dev_id):
         else:
             data = nvrtc.get_ptx(prog)
 
-        self.module = check_cuda_errors(cuda.cuModuleLoadData(np.char.array(data)))
+        self.module = cuda.module_load_data(np.char.array(data))
 
     def get_function(self, name):
-        return check_cuda_errors(cuda.cuModuleGetFunction(self.module, name))
+        if isinstance(name, bytes):
+            name = name.decode()
+        return cuda.module_get_function(self.module, name)

cuda_bindings/cuda/bindings/_example_helpers/helper_cuda.py

@@ -4,6 +4,7 @@
 from cuda.bindings import driver as cuda
 from cuda.bindings import nvrtc
 from cuda.bindings import runtime as cudart
+from cuda.bindings._v2 import driver as cuda_v2
 
 from .helper_string import check_cmd_line_flag, get_cmd_line_argument_int
 
@@ -43,6 +44,6 @@ def find_cuda_device_drv():
     dev_id = 0
     if check_cmd_line_flag("device="):
         dev_id = get_cmd_line_argument_int("device=")
-    check_cuda_errors(cuda.cuInit(0))
-    cu_device = check_cuda_errors(cuda.cuDeviceGet(dev_id))
+    cuda_v2.init(0)
+    cu_device = cuda_v2.device_get(dev_id)
     return cu_device

cuda_bindings/cuda/bindings/_v2/nvrtc.pyx

@@ -3,14 +3,13 @@
 # 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=a36c7e54cf29166832dd9aebc1fa71cc3649498794a2846e707396419caebe10
+# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=f9b340180d3296dbf8b6145df99c5475fb6fabd158d985d1c5f6e080a1852886
 
 
 # <<<< PREAMBLE CONTENT >>>>
 
 cimport cpython as _cyb_cpython
 cimport cpython.buffer as _cyb_cpython_buffer
-from cython cimport view as _cyb_view
 from libc.stdint cimport intptr_t
 from libc.stdlib cimport (
     calloc as _cyb_calloc,
@@ -630,7 +629,7 @@ cpdef str get_error_string(int result):
     cdef const char *_output_cstr_
     cdef bytes _output_
     with nogil:
-        _output_cstr_ = nvrtcGetErrorString(<_Result>result)
+        _output_cstr_ = nvrtcGetErrorString(<nvrtcResult>result)
     _output_ = _output_cstr_
     return _output_.decode()
 
@@ -657,6 +656,8 @@ cpdef tuple version():
 cpdef int get_num_supported_archs() except? -1:
     """nvrtcGetNumSupportedArchs sets the output parameter ``num_archs`` with the number of architectures supported by NVRTC. This can then be used to pass an array to ``nvrtcGetSupportedArchs`` to get the supported architectures.
 
+    see ``nvrtcGetSupportedArchs``.
+
     Returns:
         int: number of supported architectures.
 
@@ -672,6 +673,8 @@ cpdef int get_num_supported_archs() except? -1:
 cpdef object get_supported_archs():
     """nvrtcGetSupportedArchs populates the array passed via the output parameter ``supported_archs`` with the architectures supported by NVRTC. The array is sorted in the ascending order. The size of the array to be passed can be determined using ``nvrtcGetNumSupportedArchs``.
 
+    see ``nvrtcGetNumSupportedArchs``.
+
     Returns:
         int: sorted array of supported architectures.
 
@@ -681,13 +684,14 @@ cpdef object get_supported_archs():
     with nogil:
         __status__ = nvrtcGetNumSupportedArchs(&numArchs)
     check_status(__status__)
-    if numArchs == 0:
-        return _cyb_view.array(shape=(1,), itemsize=sizeof(int), format="i", mode="c")[:0]
-    cdef _cyb_view.array supported_archs = _cyb_view.array(shape=(numArchs,), itemsize=sizeof(int), format="i", mode="c")
-    cdef int *supported_archs_ptr = <int *>(supported_archs.data)
-    with nogil:
-        __status__ = nvrtcGetSupportedArchs(supported_archs_ptr)
-    check_status(__status__)
+    cdef object _supported_archs_alloc_ = _numpy.empty(max(numArchs, 1), dtype=_numpy.int32)
+    cdef intptr_t _supported_archs_data_ = <intptr_t>_supported_archs_alloc_.ctypes.data
+    cdef int *supported_archs_ptr = <int *>_supported_archs_data_
+    cdef object supported_archs = _supported_archs_alloc_[:numArchs]
+    if numArchs != 0:
+        with nogil:
+            __status__ = nvrtcGetSupportedArchs(supported_archs_ptr)
+        check_status(__status__)
     return supported_archs
 
 
@@ -699,9 +703,9 @@ cpdef destroy_program(intptr_t prog):
 
     .. seealso:: `nvrtcDestroyProgram`
     """
-    cdef Program _prog_ = <Program>prog
+    cdef nvrtcProgram _prog_ = <nvrtcProgram>prog
     with nogil:
-        __status__ = nvrtcDestroyProgram(&_prog_)
+        __status__ = nvrtcDestroyProgram(<nvrtcProgram*>prog)
     check_status(__status__)
 
 
@@ -719,7 +723,7 @@ cpdef size_t get_ptx_size(intptr_t prog) except? 0:
     """
     cdef size_t ptx_size_ret
     with nogil:
-        __status__ = nvrtcGetPTXSize(<Program>prog, &ptx_size_ret)
+        __status__ = nvrtcGetPTXSize(<nvrtcProgram>prog, &ptx_size_ret)
     check_status(__status__)
     return ptx_size_ret
 
@@ -737,15 +741,14 @@ cpdef bytes get_ptx(intptr_t prog):
     """
     cdef size_t ptxSizeRet
     with nogil:
-        __status__ = nvrtcGetPTXSize(<Program>prog, &ptxSizeRet)
+        __status__ = nvrtcGetPTXSize(<nvrtcProgram>prog, &ptxSizeRet)
     check_status(__status__)
-    if ptxSizeRet == 0:
-        return b""
     cdef bytes _ptx_ = bytes(ptxSizeRet)
     cdef char* ptx = _ptx_
-    with nogil:
-        __status__ = nvrtcGetPTX(<Program>prog, ptx)
-    check_status(__status__)
+    if ptxSizeRet != 0:
+        with nogil:
+            __status__ = nvrtcGetPTX(<nvrtcProgram>prog, ptx)
+        check_status(__status__)
     return _ptx_
 
 
@@ -762,7 +765,7 @@ cpdef size_t get_cubin_size(intptr_t prog) except? 0:
     """
     cdef size_t cubin_size_ret
     with nogil:
-        __status__ = nvrtcGetCUBINSize(<Program>prog, &cubin_size_ret)
+        __status__ = nvrtcGetCUBINSize(<nvrtcProgram>prog, &cubin_size_ret)
     check_status(__status__)
     return cubin_size_ret
 
@@ -780,15 +783,14 @@ cpdef bytes get_cubin(intptr_t prog):
     """
     cdef size_t cubinSizeRet
     with nogil:
-        __status__ = nvrtcGetCUBINSize(<Program>prog, &cubinSizeRet)
+        __status__ = nvrtcGetCUBINSize(<nvrtcProgram>prog, &cubinSizeRet)
     check_status(__status__)
-    if cubinSizeRet == 0:
-        return b""
     cdef bytes _cubin_ = bytes(cubinSizeRet)
     cdef char* cubin = _cubin_
-    with nogil:
-        __status__ = nvrtcGetCUBIN(<Program>prog, cubin)
-    check_status(__status__)
+    if cubinSizeRet != 0:
+        with nogil:
+            __status__ = nvrtcGetCUBIN(<nvrtcProgram>prog, cubin)
+        check_status(__status__)
     return _cubin_
 
 
@@ -805,7 +807,7 @@ cpdef size_t get_ltoir_size(intptr_t prog) except? 0:
     """
     cdef size_t ltoir_size_ret
     with nogil:
-        __status__ = nvrtcGetLTOIRSize(<Program>prog, &ltoir_size_ret)
+        __status__ = nvrtcGetLTOIRSize(<nvrtcProgram>prog, &ltoir_size_ret)
     check_status(__status__)
     return ltoir_size_ret
 
@@ -823,15 +825,14 @@ cpdef bytes get_ltoir(intptr_t prog):
     """
     cdef size_t LTOIRSizeRet
     with nogil:
-        __status__ = nvrtcGetLTOIRSize(<Program>prog, &LTOIRSizeRet)
+        __status__ = nvrtcGetLTOIRSize(<nvrtcProgram>prog, &LTOIRSizeRet)
     check_status(__status__)
-    if LTOIRSizeRet == 0:
-        return b""
     cdef bytes _ltoir_ = bytes(LTOIRSizeRet)
     cdef char* ltoir = _ltoir_
-    with nogil:
-        __status__ = nvrtcGetLTOIR(<Program>prog, ltoir)
-    check_status(__status__)
+    if LTOIRSizeRet != 0:
+        with nogil:
+            __status__ = nvrtcGetLTOIR(<nvrtcProgram>prog, ltoir)
+        check_status(__status__)
     return _ltoir_
 
 
@@ -848,7 +849,7 @@ cpdef size_t get_optix_ir_size(intptr_t prog) except? 0:
     """
     cdef size_t optixir_size_ret
     with nogil:
-        __status__ = nvrtcGetOptiXIRSize(<Program>prog, &optixir_size_ret)
+        __status__ = nvrtcGetOptiXIRSize(<nvrtcProgram>prog, &optixir_size_ret)
     check_status(__status__)
     return optixir_size_ret
 
@@ -866,21 +867,23 @@ cpdef bytes get_optix_ir(intptr_t prog):
     """
     cdef size_t optixirSizeRet
     with nogil:
-        __status__ = nvrtcGetOptiXIRSize(<Program>prog, &optixirSizeRet)
+        __status__ = nvrtcGetOptiXIRSize(<nvrtcProgram>prog, &optixirSizeRet)
     check_status(__status__)
-    if optixirSizeRet == 0:
-        return b""
     cdef bytes _optixir_ = bytes(optixirSizeRet)
     cdef char* optixir = _optixir_
-    with nogil:
-        __status__ = nvrtcGetOptiXIR(<Program>prog, optixir)
-    check_status(__status__)
+    if optixirSizeRet != 0:
+        with nogil:
+            __status__ = nvrtcGetOptiXIR(<nvrtcProgram>prog, optixir)
+        check_status(__status__)
     return _optixir_
 
 
 cpdef size_t get_program_log_size(intptr_t prog) except? 0:
     """nvrtcGetProgramLogSize sets ``log_size_ret`` with the size of the log generated by the previous compilation of ``prog`` (including the trailing ``NULL``).
 
+    Note that compilation log may be generated with warnings and informative
+    messages, even when the compilation of ``prog`` succeeds.
+
     Args:
         prog (intptr_t): CUDA Runtime Compilation program.
 
@@ -892,7 +895,7 @@ cpdef size_t get_program_log_size(intptr_t prog) except? 0:
     """
     cdef size_t log_size_ret
     with nogil:
-        __status__ = nvrtcGetProgramLogSize(<Program>prog, &log_size_ret)
+        __status__ = nvrtcGetProgramLogSize(<nvrtcProgram>prog, &log_size_ret)
     check_status(__status__)
     return log_size_ret
 
@@ -910,21 +913,23 @@ cpdef bytes get_program_log(intptr_t prog):
     """
     cdef size_t logSizeRet
     with nogil:
-        __status__ = nvrtcGetProgramLogSize(<Program>prog, &logSizeRet)
+        __status__ = nvrtcGetProgramLogSize(<nvrtcProgram>prog, &logSizeRet)
     check_status(__status__)
-    if logSizeRet == 0:
-        return b""
     cdef bytes _log_ = bytes(logSizeRet)
     cdef char* log = _log_
-    with nogil:
-        __status__ = nvrtcGetProgramLog(<Program>prog, log)
-    check_status(__status__)
+    if logSizeRet != 0:
+        with nogil:
+            __status__ = nvrtcGetProgramLog(<nvrtcProgram>prog, log)
+        check_status(__status__)
     return _log_
 
 
 cpdef add_name_expression(intptr_t prog, name_expression):
     """nvrtcAddNameExpression notes the given name expression denoting the address of a global function or device/__constant__ variable.
 
+    The identical name expression string must be provided on a subsequent call
+    to nvrtcGetLoweredName to extract the lowered name.
+
     Args:
         prog (intptr_t): CUDA Runtime Compilation program.
         name_expression (str): constant expression denoting the
@@ -938,7 +943,7 @@ cpdef add_name_expression(intptr_t prog, name_expression):
     cdef bytes _temp_name_expression_ = (<str>name_expression).encode()
     cdef char* _name_expression_ = _temp_name_expression_
     with nogil:
-        __status__ = nvrtcAddNameExpression(<Program>prog, <const char* const>_name_expression_)
+        __status__ = nvrtcAddNameExpression(<nvrtcProgram>prog, <const char*>_name_expression_)
     check_status(__status__)
 
 
@@ -961,6 +966,10 @@ cpdef size_t get_pch_heap_size() except? 0:
 cpdef set_pch_heap_size(size_t size):
     """set the size of the PCH Heap.
 
+    The requested size may be rounded up to a platform dependent alignment
+    (e.g. page size). If the PCH Heap has already been allocated, the heap
+    memory will be freed and a new PCH Heap will be allocated.
+
     Args:
         size (size_t): requested size of the PCH Heap, in bytes.
 
@@ -974,14 +983,28 @@ cpdef set_pch_heap_size(size_t size):
 cpdef int get_pch_create_status(intptr_t prog) except? -1:
     """returns the PCH creation status.
 
+    NVRTC_SUCCESS indicates that the PCH was successfully created.
+    NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED indicates that no PCH creation was
+    attempted, either because PCH functionality was not requested during the
+    preceding nvrtcCompileProgram call, or automatic PCH processing was
+    requested, and compiler chose not to create a PCH file.
+    NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED indicates that a PCH file could
+    potentially have been created, but the compiler ran out space in the PCH
+    heap. In this scenario, the :func:`get_pch_heap_size_required` can be used
+    to query the required heap size, the heap can be reallocated for this size
+    with :func:`set_pch_heap_size` and PCH creation may be reattempted again
+    invoking :func:`compile_program` with a new NVRTC program instance.
+    NVRTC_ERROR_PCH_CREATE indicates that an error condition prevented the PCH
+    file from being created.
+
     Args:
         prog (intptr_t): CUDA Runtime Compilation program.
 
     .. seealso:: `nvrtcGetPCHCreateStatus`
     """
     cdef int ret
     with nogil:
-        ret = <int>nvrtcGetPCHCreateStatus(<Program>prog)
+        ret = <int>nvrtcGetPCHCreateStatus(<nvrtcProgram>prog)
     return ret
 
 
@@ -999,7 +1022,7 @@ cpdef size_t get_pch_heap_size_required(intptr_t prog) except? 0:
     """
     cdef size_t size
     with nogil:
-        __status__ = nvrtcGetPCHHeapSizeRequired(<Program>prog, &size)
+        __status__ = nvrtcGetPCHHeapSizeRequired(<nvrtcProgram>prog, &size)
     check_status(__status__)
     return size
 
@@ -1017,7 +1040,7 @@ cpdef size_t get_tile_ir_size(intptr_t prog) except? 0:
     """
     cdef size_t tile_ir_size_ret
     with nogil:
-        __status__ = nvrtcGetTileIRSize(<Program>prog, &tile_ir_size_ret)
+        __status__ = nvrtcGetTileIRSize(<nvrtcProgram>prog, &tile_ir_size_ret)
     check_status(__status__)
     return tile_ir_size_ret
 
@@ -1035,15 +1058,14 @@ cpdef bytes get_tile_ir(intptr_t prog):
     """
     cdef size_t TileIRSizeRet
     with nogil:
-        __status__ = nvrtcGetTileIRSize(<Program>prog, &TileIRSizeRet)
+        __status__ = nvrtcGetTileIRSize(<nvrtcProgram>prog, &TileIRSizeRet)
     check_status(__status__)
-    if TileIRSizeRet == 0:
-        return b""
     cdef bytes _tile_ir_ = bytes(TileIRSizeRet)
     cdef char* tile_ir = _tile_ir_
-    with nogil:
-        __status__ = nvrtcGetTileIR(<Program>prog, tile_ir)
-    check_status(__status__)
+    if TileIRSizeRet != 0:
+        with nogil:
+            __status__ = nvrtcGetTileIR(<nvrtcProgram>prog, tile_ir)
+        check_status(__status__)
     return _tile_ir_
 
 

cuda_bindings/cuda/bindings/cudla.pxd

@@ -2,7 +2,7 @@
 # SPDX-License-Identifier: Apache-2.0
 
 # This code was automatically generated across versions from 1.5.0 to 13.3.0. Do not modify it directly.
-# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=b82680ec867e23638b173760105c35030e0cba5c9a8b3bb536ce5bb3381ec1fb
+# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=3b52467aafbb78d33c99d8be6d3acdd552c8aa0851758d01932080aebaec5213
 
 
 # <<<< PREAMBLE CONTENT >>>>
@@ -55,7 +55,7 @@ cpdef intptr_t create_device(uint64_t device, uint32_t flags) except *
 cpdef intptr_t mem_register(intptr_t dev_handle, intptr_t ptr, size_t size, uint32_t flags) except *
 cpdef intptr_t module_load_from_memory(intptr_t dev_handle, p_module, size_t module_size, uint32_t flags) except *
 cpdef module_unload(intptr_t h_module, uint32_t flags)
-cpdef submit_task(intptr_t dev_handle, intptr_t ptr_to_tasks, uint32_t num_tasks, intptr_t stream, uint32_t flags)
+cpdef submit_task(intptr_t dev_handle, ptr_to_tasks, uint32_t num_tasks, intptr_t stream, uint32_t flags)
 cpdef object device_get_attribute(intptr_t dev_handle, int attrib)
 cpdef mem_unregister(intptr_t dev_handle, intptr_t dev_ptr)
 cpdef int get_last_error(intptr_t dev_handle) except? 0