Use numpy for array allocation instead of cython.view.array

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=b73f313203f01d91dab664fb485858edcaab834bc1b5a181db05c2d53d660117
 
 
 # <<<< 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,
@@ -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,17 @@ 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
+    cdef intptr_t _supported_archs_data_
+    cdef int *supported_archs_ptr
+    supported_archs_alloc = _numpy.empty(max(numArchs, 1), dtype=_numpy.int32)
+    supported_archs = supported_archs_alloc[:numArchs]
+    _supported_archs_data_ = <intptr_t>supported_archs_alloc.ctypes.data
+    supported_archs_ptr = <int *>_supported_archs_data_
+    if numArchs != 0:
+        with nogil:
+            __status__ = nvrtcGetSupportedArchs(supported_archs_ptr)
+        check_status(__status__)
     return supported_archs
 
 
@@ -739,13 +746,12 @@ cpdef bytes get_ptx(intptr_t prog):
     with nogil:
         __status__ = nvrtcGetPTXSize(<Program>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(<Program>prog, ptx)
+        check_status(__status__)
     return _ptx_
 
 
@@ -782,13 +788,12 @@ cpdef bytes get_cubin(intptr_t prog):
     with nogil:
         __status__ = nvrtcGetCUBINSize(<Program>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(<Program>prog, cubin)
+        check_status(__status__)
     return _cubin_
 
 
@@ -825,13 +830,12 @@ cpdef bytes get_ltoir(intptr_t prog):
     with nogil:
         __status__ = nvrtcGetLTOIRSize(<Program>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(<Program>prog, ltoir)
+        check_status(__status__)
     return _ltoir_
 
 
@@ -868,19 +872,21 @@ cpdef bytes get_optix_ir(intptr_t prog):
     with nogil:
         __status__ = nvrtcGetOptiXIRSize(<Program>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(<Program>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.
 
@@ -912,19 +918,21 @@ cpdef bytes get_program_log(intptr_t prog):
     with nogil:
         __status__ = nvrtcGetProgramLogSize(<Program>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(<Program>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
@@ -961,6 +969,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,6 +986,20 @@ 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.
 
@@ -1037,13 +1063,12 @@ cpdef bytes get_tile_ir(intptr_t prog):
     with nogil:
         __status__ = nvrtcGetTileIRSize(<Program>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(<Program>prog, tile_ir)
+        check_status(__status__)
     return _tile_ir_
 
 

cuda_bindings/cuda/bindings/cudla.pyx

@@ -2,14 +2,14 @@
 # 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=3c177b7a0328c0f6f16067c8c9f4e5a002bd019e8c17c017ba9f77af21da8d75
+# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=5c0d6715f105108fd83cc471fd0978ecab1f1cdfdd78f41ac3bcf7dc1a470ab5
 
 
 # <<<< PREAMBLE CONTENT >>>>
 
 cimport cpython as _cyb_cpython
 cimport cpython.buffer as _cyb_cpython_buffer
-from cython cimport view as _cyb_view
+from cpython.memoryview cimport PyMemoryView_FromMemory as _cyb_PyMemoryView_FromMemory
 from libc.stdint cimport (
     intptr_t,
     uint32_t,
@@ -604,19 +604,21 @@ cdef class ModuleTensorDescriptor:
     @property
     def stride(self):
         """~_numpy.uint32: (array of length 8)."""
-        cdef _cyb_view.array arr = _cyb_view.array(shape=(8,), itemsize=sizeof(uint32_t), format="I", mode="c", allocate_buffer=False)
-        arr.data = <char *>(&(self._ptr[0].stride))
-        return _numpy.asarray(arr)
+        cdef object _mv_ = _cyb_PyMemoryView_FromMemory(
+            <char *>(&(self._ptr[0].stride)),
+            <Py_ssize_t>(sizeof(uint32_t) * (8)),
+            _cyb_cpython_buffer.PyBUF_WRITE if not self._readonly else _cyb_cpython_buffer.PyBUF_READ,
+        )
+        return _numpy.frombuffer(_mv_, dtype=_numpy.uint32)
 
     @stride.setter
     def stride(self, val):
         if self._readonly:
             raise ValueError("This ModuleTensorDescriptor instance is read-only")
         if len(val) != 8:
             raise ValueError(f"Expected length { 8 } for field stride, got {len(val)}")
-        cdef _cyb_view.array arr = _cyb_view.array(shape=(8,), itemsize=sizeof(uint32_t), format="I", mode="c")
-        arr[:] = _numpy.asarray(val, dtype=_numpy.uint32)
-        _cyb_memcpy(<void *>(&(self._ptr[0].stride)), <void *>(arr.data), sizeof(uint32_t) * len(val))
+        _val_ = _numpy.ascontiguousarray(_numpy.asarray(val, dtype=_numpy.uint32))
+        _cyb_memcpy(<void *>(&(self._ptr[0].stride)), <void *><intptr_t>(_val_.ctypes.data), sizeof(uint32_t) * (8))
 
     @staticmethod
     def from_buffer(buffer):
@@ -1344,10 +1346,13 @@ cdef class SignalEvents:
     def dev_ptrs(self):
         """int: """
         if self._ptr[0].devPtrs == NULL or self._ptr[0].numEvents == 0:
-            return _cyb_view.array(shape=(1,), itemsize=sizeof(intptr_t), format="q", mode="c")[:0]
-        cdef _cyb_view.array arr = _cyb_view.array(shape=(self._ptr[0].numEvents,), itemsize=sizeof(intptr_t), format="q", mode="c", allocate_buffer=False)
-        arr.data = <char *>(self._ptr[0].devPtrs)
-        return arr
+            return _numpy.empty(0, dtype=_numpy.intp)
+        cdef object _mv_ = _cyb_PyMemoryView_FromMemory(
+            <char *>(self._ptr[0].devPtrs),
+            <Py_ssize_t>(self._ptr[0].numEvents * sizeof(intptr_t)),
+            _cyb_cpython_buffer.PyBUF_WRITE,
+        )
+        return _numpy.frombuffer(_mv_, dtype=_numpy.intp)
 
     @dev_ptrs.setter
     def dev_ptrs(self, val):
@@ -1357,13 +1362,9 @@ cdef class SignalEvents:
         self._ptr[0].numEvents = _n
         if _n == 0:
             return
-        cdef _cyb_view.array arr = _cyb_view.array(shape=(_n,), itemsize=sizeof(intptr_t), format="q", mode="c")
-        cdef intptr_t[:] mv = arr
-        cdef Py_ssize_t i
-        for i in range(_n):
-            mv[i] = val[i]
-        self._ptr[0].devPtrs = <uint64_t**><intptr_t>(arr.data)
-        self._refs["dev_ptrs"] = arr
+        _arr_ = _numpy.ascontiguousarray(_numpy.asarray(val, dtype=_numpy.intp))
+        self._ptr[0].devPtrs = <uint64_t**><intptr_t>_arr_.ctypes.data
+        self._refs["dev_ptrs"] = _arr_
 
     @property
     def eof_fences(self):
@@ -1530,10 +1531,13 @@ cdef class Task:
     def output_tensor(self):
         """int: """
         if self._ptr[0].outputTensor == NULL or self._ptr[0].numOutputTensors == 0:
-            return _cyb_view.array(shape=(1,), itemsize=sizeof(intptr_t), format="q", mode="c")[:0]
-        cdef _cyb_view.array arr = _cyb_view.array(shape=(self._ptr[0].numOutputTensors,), itemsize=sizeof(intptr_t), format="q", mode="c", allocate_buffer=False)
-        arr.data = <char *>(self._ptr[0].outputTensor)
-        return arr
+            return _numpy.empty(0, dtype=_numpy.intp)
+        cdef object _mv_ = _cyb_PyMemoryView_FromMemory(
+            <char *>(self._ptr[0].outputTensor),
+            <Py_ssize_t>(self._ptr[0].numOutputTensors * sizeof(intptr_t)),
+            _cyb_cpython_buffer.PyBUF_WRITE,
+        )
+        return _numpy.frombuffer(_mv_, dtype=_numpy.intp)
 
     @output_tensor.setter
     def output_tensor(self, val):
@@ -1543,22 +1547,21 @@ cdef class Task:
         self._ptr[0].numOutputTensors = _n
         if _n == 0:
             return
-        cdef _cyb_view.array arr = _cyb_view.array(shape=(_n,), itemsize=sizeof(intptr_t), format="q", mode="c")
-        cdef intptr_t[:] mv = arr
-        cdef Py_ssize_t i
-        for i in range(_n):
-            mv[i] = val[i]
-        self._ptr[0].outputTensor = <uint64_t**><intptr_t>(arr.data)
-        self._refs["output_tensor"] = arr
+        _arr_ = _numpy.ascontiguousarray(_numpy.asarray(val, dtype=_numpy.intp))
+        self._ptr[0].outputTensor = <uint64_t**><intptr_t>_arr_.ctypes.data
+        self._refs["output_tensor"] = _arr_
 
     @property
     def input_tensor(self):
         """int: """
         if self._ptr[0].inputTensor == NULL or self._ptr[0].numInputTensors == 0:
-            return _cyb_view.array(shape=(1,), itemsize=sizeof(intptr_t), format="q", mode="c")[:0]
-        cdef _cyb_view.array arr = _cyb_view.array(shape=(self._ptr[0].numInputTensors,), itemsize=sizeof(intptr_t), format="q", mode="c", allocate_buffer=False)
-        arr.data = <char *>(self._ptr[0].inputTensor)
-        return arr
+            return _numpy.empty(0, dtype=_numpy.intp)
+        cdef object _mv_ = _cyb_PyMemoryView_FromMemory(
+            <char *>(self._ptr[0].inputTensor),
+            <Py_ssize_t>(self._ptr[0].numInputTensors * sizeof(intptr_t)),
+            _cyb_cpython_buffer.PyBUF_WRITE,
+        )
+        return _numpy.frombuffer(_mv_, dtype=_numpy.intp)
 
     @input_tensor.setter
     def input_tensor(self, val):
@@ -1568,13 +1571,9 @@ cdef class Task:
         self._ptr[0].numInputTensors = _n
         if _n == 0:
             return
-        cdef _cyb_view.array arr = _cyb_view.array(shape=(_n,), itemsize=sizeof(intptr_t), format="q", mode="c")
-        cdef intptr_t[:] mv = arr
-        cdef Py_ssize_t i
-        for i in range(_n):
-            mv[i] = val[i]
-        self._ptr[0].inputTensor = <uint64_t**><intptr_t>(arr.data)
-        self._refs["input_tensor"] = arr
+        _arr_ = _numpy.ascontiguousarray(_numpy.asarray(val, dtype=_numpy.intp))
+        self._ptr[0].inputTensor = <uint64_t**><intptr_t>_arr_.ctypes.data
+        self._refs["input_tensor"] = _arr_
 
     @property
     def wait_events(self):