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

Author: mdboom

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=370c109c32bc68fce1daa6d02b41abf3f021d9ca99f7686976a9a338430c3649
 
 
 # <<<< 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,10 +684,10 @@ 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)
+    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]
     with nogil:
         __status__ = nvrtcGetSupportedArchs(supported_archs_ptr)
     check_status(__status__)
@@ -881,6 +884,9 @@ cpdef bytes get_optix_ir(intptr_t prog):
 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.
 
@@ -925,6 +931,9 @@ cpdef bytes get_program_log(intptr_t prog):
 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 +970,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 +987,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.
 

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):

cuda_bindings/cuda/bindings/cufile.pyx

@@ -3,15 +3,15 @@
 # SPDX-License-Identifier: Apache-2.0
 #
 # This code was automatically generated across versions from 12.9.1 to 13.3.0. Do not modify it directly.
-# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=df46a6921d93f83249134c7705b2809f57145b6fb72f6f40c4657ecd1b443b81
+# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=b46f5a16585b36d52e8bf6a7ca3009d8c2350427cff03f68ab4b558b93c905fc
 
 
 # <<<< PREAMBLE CONTENT >>>>
 
 cimport cpython as _cyb_cpython
 cimport cpython.buffer as _cyb_cpython_buffer
 cimport cpython.memoryview as _cyb_cpython_memoryview
-from cython cimport view as _cyb_view
+from cpython.memoryview cimport PyMemoryView_FromMemory as _cyb_PyMemoryView_FromMemory
 from libc.stdint cimport (
     intptr_t,
     uint64_t,
@@ -2532,36 +2532,40 @@ cdef class StatsLevel2:
     @property
     def read_size_kb_hist(self):
         """~_numpy.uint64: (array of length 32)."""
-        cdef _cyb_view.array arr = _cyb_view.array(shape=(32,), itemsize=sizeof(uint64_t), format="Q", mode="c", allocate_buffer=False)
-        arr.data = <char *>(&(self._ptr[0].read_size_kb_hist))
-        return _numpy.asarray(arr)
+        cdef object _mv_ = _cyb_PyMemoryView_FromMemory(
+            <char *>(&(self._ptr[0].read_size_kb_hist)),
+            <Py_ssize_t>(sizeof(uint64_t) * (32)),
+            _cyb_cpython_buffer.PyBUF_WRITE if not self._readonly else _cyb_cpython_buffer.PyBUF_READ,
+        )
+        return _numpy.frombuffer(_mv_, dtype=_numpy.uint64)
 
     @read_size_kb_hist.setter
     def read_size_kb_hist(self, val):
         if self._readonly:
             raise ValueError("This StatsLevel2 instance is read-only")
         if len(val) != 32:
             raise ValueError(f"Expected length { 32 } for field read_size_kb_hist, got {len(val)}")
-        cdef _cyb_view.array arr = _cyb_view.array(shape=(32,), itemsize=sizeof(uint64_t), format="Q", mode="c")
-        arr[:] = _numpy.asarray(val, dtype=_numpy.uint64)
-        _cyb_memcpy(<void *>(&(self._ptr[0].read_size_kb_hist)), <void *>(arr.data), sizeof(uint64_t) * len(val))
+        _val_ = _numpy.ascontiguousarray(_numpy.asarray(val, dtype=_numpy.uint64))
+        _cyb_memcpy(<void *>(&(self._ptr[0].read_size_kb_hist)), <void *><intptr_t>(_val_.ctypes.data), sizeof(uint64_t) * (32))
 
     @property
     def write_size_kb_hist(self):
         """~_numpy.uint64: (array of length 32)."""
-        cdef _cyb_view.array arr = _cyb_view.array(shape=(32,), itemsize=sizeof(uint64_t), format="Q", mode="c", allocate_buffer=False)
-        arr.data = <char *>(&(self._ptr[0].write_size_kb_hist))
-        return _numpy.asarray(arr)
+        cdef object _mv_ = _cyb_PyMemoryView_FromMemory(
+            <char *>(&(self._ptr[0].write_size_kb_hist)),
+            <Py_ssize_t>(sizeof(uint64_t) * (32)),
+            _cyb_cpython_buffer.PyBUF_WRITE if not self._readonly else _cyb_cpython_buffer.PyBUF_READ,
+        )
+        return _numpy.frombuffer(_mv_, dtype=_numpy.uint64)
 
     @write_size_kb_hist.setter
     def write_size_kb_hist(self, val):
         if self._readonly:
             raise ValueError("This StatsLevel2 instance is read-only")
         if len(val) != 32:
             raise ValueError(f"Expected length { 32 } for field write_size_kb_hist, got {len(val)}")
-        cdef _cyb_view.array arr = _cyb_view.array(shape=(32,), itemsize=sizeof(uint64_t), format="Q", mode="c")
-        arr[:] = _numpy.asarray(val, dtype=_numpy.uint64)
-        _cyb_memcpy(<void *>(&(self._ptr[0].write_size_kb_hist)), <void *>(arr.data), sizeof(uint64_t) * len(val))
+        _val_ = _numpy.ascontiguousarray(_numpy.asarray(val, dtype=_numpy.uint64))
+        _cyb_memcpy(<void *>(&(self._ptr[0].write_size_kb_hist)), <void *><intptr_t>(_val_.ctypes.data), sizeof(uint64_t) * (32))
 
     @staticmethod
     def from_buffer(buffer):