diff --git a/cuda_bindings/cuda/bindings/cyruntime.pxd b/cuda_bindings/cuda/bindings/cyruntime.pxd index 58e0a14ca60..7eb76712237 100644 --- a/cuda_bindings/cuda/bindings/cyruntime.pxd +++ b/cuda_bindings/cuda/bindings/cyruntime.pxd @@ -4,7 +4,7 @@ # # 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=2b62a3b56226f5cf6953f578057af792369a100b7775328c5aa66d81d780d2ac +# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=a7205ec1f1749acd8f9e32e85660d77d8ccb60a0d038bb63fb4b015496f23d07 from libc.stdint cimport uint32_t, uint64_t @@ -1254,6 +1254,20 @@ cdef extern from 'library_types.h': CUDA_EMULATION_STRATEGY_EAGER ctypedef cudaEmulationStrategy_t cudaEmulationStrategy +cdef extern from 'library_types.h': + cdef enum cudaEmulationMantissaControl_t: + CUDA_EMULATION_MANTISSA_CONTROL_DYNAMIC + CUDA_EMULATION_MANTISSA_CONTROL_FIXED + ctypedef cudaEmulationMantissaControl_t cudaEmulationMantissaControl + +cdef extern from 'library_types.h': + cdef enum cudaEmulationSpecialValuesSupport_t: + CUDA_EMULATION_SPECIAL_VALUES_SUPPORT_DEFAULT + CUDA_EMULATION_SPECIAL_VALUES_SUPPORT_NONE + CUDA_EMULATION_SPECIAL_VALUES_SUPPORT_INFINITY + CUDA_EMULATION_SPECIAL_VALUES_SUPPORT_NAN + ctypedef cudaEmulationSpecialValuesSupport_t cudaEmulationSpecialValuesSupport + cdef extern from 'driver_types.h': cdef enum cudaDevSmResourceGroup_flags: cudaDevSmResourceGroupDefault @@ -1276,20 +1290,6 @@ cdef extern from 'driver_types.h': cudaDevWorkqueueConfigScopeDeviceCtx cudaDevWorkqueueConfigScopeGreenCtxBalanced -cdef extern from 'library_types.h': - cdef enum cudaEmulationMantissaControl_t: - CUDA_EMULATION_MANTISSA_CONTROL_DYNAMIC - CUDA_EMULATION_MANTISSA_CONTROL_FIXED - ctypedef cudaEmulationMantissaControl_t cudaEmulationMantissaControl - -cdef extern from 'library_types.h': - cdef enum cudaEmulationSpecialValuesSupport_t: - CUDA_EMULATION_SPECIAL_VALUES_SUPPORT_DEFAULT - CUDA_EMULATION_SPECIAL_VALUES_SUPPORT_NONE - CUDA_EMULATION_SPECIAL_VALUES_SUPPORT_INFINITY - CUDA_EMULATION_SPECIAL_VALUES_SUPPORT_NAN - ctypedef cudaEmulationSpecialValuesSupport_t cudaEmulationSpecialValuesSupport - cdef extern from 'driver_types.h': cdef enum cudaHostTaskSyncMode: cudaHostTaskBlocking diff --git a/cuda_bindings/cuda/bindings/driver.pyx b/cuda_bindings/cuda/bindings/driver.pyx index 44b7c4c567c..5ab7f55e487 100644 --- a/cuda_bindings/cuda/bindings/driver.pyx +++ b/cuda_bindings/cuda/bindings/driver.pyx @@ -2,7 +2,7 @@ # SPDX-License-Identifier: Apache-2.0 # This code was automatically generated with version 13.3.0. Do not modify it directly. -# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=4a70be46627269cff03a2b89504463158d6e50d738cfde91e8c3ed1bf88fd9e0 +# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=6cc0baa1c711f87e25f49626d1d5ab66c770b0f81b21ae45da420e2632024565 from typing import Any, Optional import cython import ctypes @@ -10810,6 +10810,7 @@ cdef class CUDA_BATCH_MEM_OP_NODE_PARAMS_v1_st: return [CUstreamBatchMemOpParams(_ptr=arr) for arr in arrs] @paramArray.setter def paramArray(self, val): + cdef cydriver.CUstreamBatchMemOpParams* _paramArray_new if len(val) == 0: free(self._paramArray) self._paramArray = NULL @@ -10817,14 +10818,22 @@ cdef class CUDA_BATCH_MEM_OP_NODE_PARAMS_v1_st: self._pvt_ptr[0].paramArray = NULL else: if self._paramArray_length != len(val): - free(self._paramArray) - self._paramArray = calloc(len(val), sizeof(cydriver.CUstreamBatchMemOpParams)) - if self._paramArray is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _paramArray_new = calloc(len(val), sizeof(cydriver.CUstreamBatchMemOpParams)) + if _paramArray_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cydriver.CUstreamBatchMemOpParams))) + for idx in range(len(val)): + string.memcpy(&_paramArray_new[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUstreamBatchMemOpParams)) + free(self._paramArray) + self._paramArray = _paramArray_new self._paramArray_length = len(val) - self._pvt_ptr[0].paramArray = self._paramArray - for idx in range(len(val)): - string.memcpy(&self._paramArray[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUstreamBatchMemOpParams)) + self._pvt_ptr[0].paramArray = _paramArray_new + else: + for idx in range(len(val)): + string.memcpy(&self._paramArray[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUstreamBatchMemOpParams)) @@ -10945,6 +10954,7 @@ cdef class CUDA_BATCH_MEM_OP_NODE_PARAMS_v2_st: return [CUstreamBatchMemOpParams(_ptr=arr) for arr in arrs] @paramArray.setter def paramArray(self, val): + cdef cydriver.CUstreamBatchMemOpParams* _paramArray_new if len(val) == 0: free(self._paramArray) self._paramArray = NULL @@ -10952,14 +10962,22 @@ cdef class CUDA_BATCH_MEM_OP_NODE_PARAMS_v2_st: self._pvt_ptr[0].paramArray = NULL else: if self._paramArray_length != len(val): - free(self._paramArray) - self._paramArray = calloc(len(val), sizeof(cydriver.CUstreamBatchMemOpParams)) - if self._paramArray is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _paramArray_new = calloc(len(val), sizeof(cydriver.CUstreamBatchMemOpParams)) + if _paramArray_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cydriver.CUstreamBatchMemOpParams))) + for idx in range(len(val)): + string.memcpy(&_paramArray_new[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUstreamBatchMemOpParams)) + free(self._paramArray) + self._paramArray = _paramArray_new self._paramArray_length = len(val) - self._pvt_ptr[0].paramArray = self._paramArray - for idx in range(len(val)): - string.memcpy(&self._paramArray[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUstreamBatchMemOpParams)) + self._pvt_ptr[0].paramArray = _paramArray_new + else: + for idx in range(len(val)): + string.memcpy(&self._paramArray[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUstreamBatchMemOpParams)) @@ -14390,6 +14408,7 @@ cdef class CUlaunchConfig_st: return [CUlaunchAttribute(_ptr=arr) for arr in arrs] @attrs.setter def attrs(self, val): + cdef cydriver.CUlaunchAttribute* _attrs_new if len(val) == 0: free(self._attrs) self._attrs = NULL @@ -14397,14 +14416,22 @@ cdef class CUlaunchConfig_st: self._pvt_ptr[0].attrs = NULL else: if self._attrs_length != len(val): - free(self._attrs) - self._attrs = calloc(len(val), sizeof(cydriver.CUlaunchAttribute)) - if self._attrs is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _attrs_new = calloc(len(val), sizeof(cydriver.CUlaunchAttribute)) + if _attrs_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cydriver.CUlaunchAttribute))) + for idx in range(len(val)): + string.memcpy(&_attrs_new[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUlaunchAttribute)) + free(self._attrs) + self._attrs = _attrs_new self._attrs_length = len(val) - self._pvt_ptr[0].attrs = self._attrs - for idx in range(len(val)): - string.memcpy(&self._attrs[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUlaunchAttribute)) + self._pvt_ptr[0].attrs = _attrs_new + else: + for idx in range(len(val)): + string.memcpy(&self._attrs[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUlaunchAttribute)) @@ -14733,6 +14760,7 @@ cdef class CUctxCreateParams_st: return [CUexecAffinityParam(_ptr=arr) for arr in arrs] @execAffinityParams.setter def execAffinityParams(self, val): + cdef cydriver.CUexecAffinityParam* _execAffinityParams_new if len(val) == 0: free(self._execAffinityParams) self._execAffinityParams = NULL @@ -14740,14 +14768,22 @@ cdef class CUctxCreateParams_st: self._pvt_ptr[0].execAffinityParams = NULL else: if self._execAffinityParams_length != len(val): - free(self._execAffinityParams) - self._execAffinityParams = calloc(len(val), sizeof(cydriver.CUexecAffinityParam)) - if self._execAffinityParams is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _execAffinityParams_new = calloc(len(val), sizeof(cydriver.CUexecAffinityParam)) + if _execAffinityParams_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cydriver.CUexecAffinityParam))) + for idx in range(len(val)): + string.memcpy(&_execAffinityParams_new[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUexecAffinityParam)) + free(self._execAffinityParams) + self._execAffinityParams = _execAffinityParams_new self._execAffinityParams_length = len(val) - self._pvt_ptr[0].execAffinityParams = self._execAffinityParams - for idx in range(len(val)): - string.memcpy(&self._execAffinityParams[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUexecAffinityParam)) + self._pvt_ptr[0].execAffinityParams = _execAffinityParams_new + else: + for idx in range(len(val)): + string.memcpy(&self._execAffinityParams[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUexecAffinityParam)) @@ -14765,6 +14801,7 @@ cdef class CUctxCreateParams_st: return [CUctxCigParam(_ptr=arr) for arr in arrs] @cigParams.setter def cigParams(self, val): + cdef cydriver.CUctxCigParam* _cigParams_new if len(val) == 0: free(self._cigParams) self._cigParams = NULL @@ -14772,14 +14809,22 @@ cdef class CUctxCreateParams_st: self._pvt_ptr[0].cigParams = NULL else: if self._cigParams_length != len(val): - free(self._cigParams) - self._cigParams = calloc(len(val), sizeof(cydriver.CUctxCigParam)) - if self._cigParams is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _cigParams_new = calloc(len(val), sizeof(cydriver.CUctxCigParam)) + if _cigParams_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cydriver.CUctxCigParam))) + for idx in range(len(val)): + string.memcpy(&_cigParams_new[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUctxCigParam)) + free(self._cigParams) + self._cigParams = _cigParams_new self._cigParams_length = len(val) - self._pvt_ptr[0].cigParams = self._cigParams - for idx in range(len(val)): - string.memcpy(&self._cigParams[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUctxCigParam)) + self._pvt_ptr[0].cigParams = _cigParams_new + else: + for idx in range(len(val)): + string.memcpy(&self._cigParams[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUctxCigParam)) @@ -14904,6 +14949,7 @@ cdef class CUstreamCigCaptureParams_st: return [CUstreamCigParam(_ptr=arr) for arr in arrs] @streamCigParams.setter def streamCigParams(self, val): + cdef cydriver.CUstreamCigParam* _streamCigParams_new if len(val) == 0: free(self._streamCigParams) self._streamCigParams = NULL @@ -14911,14 +14957,22 @@ cdef class CUstreamCigCaptureParams_st: self._pvt_ptr[0].streamCigParams = NULL else: if self._streamCigParams_length != len(val): - free(self._streamCigParams) - self._streamCigParams = calloc(len(val), sizeof(cydriver.CUstreamCigParam)) - if self._streamCigParams is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _streamCigParams_new = calloc(len(val), sizeof(cydriver.CUstreamCigParam)) + if _streamCigParams_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cydriver.CUstreamCigParam))) + for idx in range(len(val)): + string.memcpy(&_streamCigParams_new[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUstreamCigParam)) + free(self._streamCigParams) + self._streamCigParams = _streamCigParams_new self._streamCigParams_length = len(val) - self._pvt_ptr[0].streamCigParams = self._streamCigParams - for idx in range(len(val)): - string.memcpy(&self._streamCigParams[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUstreamCigParam)) + self._pvt_ptr[0].streamCigParams = _streamCigParams_new + else: + for idx in range(len(val)): + string.memcpy(&self._streamCigParams[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUstreamCigParam)) @@ -20078,6 +20132,7 @@ cdef class CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_st: return [CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS(_ptr=arr) for arr in arrs] @paramsArray.setter def paramsArray(self, val): + cdef cydriver.CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS* _paramsArray_new if len(val) == 0: free(self._paramsArray) self._paramsArray = NULL @@ -20085,14 +20140,22 @@ cdef class CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_st: self._pvt_ptr[0].paramsArray = NULL else: if self._paramsArray_length != len(val): - free(self._paramsArray) - self._paramsArray = calloc(len(val), sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS)) - if self._paramsArray is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _paramsArray_new = calloc(len(val), sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS)) + if _paramsArray_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS))) + for idx in range(len(val)): + string.memcpy(&_paramsArray_new[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS)) + free(self._paramsArray) + self._paramsArray = _paramsArray_new self._paramsArray_length = len(val) - self._pvt_ptr[0].paramsArray = self._paramsArray - for idx in range(len(val)): - string.memcpy(&self._paramsArray[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS)) + self._pvt_ptr[0].paramsArray = _paramsArray_new + else: + for idx in range(len(val)): + string.memcpy(&self._paramsArray[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS)) @@ -20205,6 +20268,7 @@ cdef class CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_v2_st: return [CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS(_ptr=arr) for arr in arrs] @paramsArray.setter def paramsArray(self, val): + cdef cydriver.CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS* _paramsArray_new if len(val) == 0: free(self._paramsArray) self._paramsArray = NULL @@ -20212,14 +20276,22 @@ cdef class CUDA_EXT_SEM_SIGNAL_NODE_PARAMS_v2_st: self._pvt_ptr[0].paramsArray = NULL else: if self._paramsArray_length != len(val): - free(self._paramsArray) - self._paramsArray = calloc(len(val), sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS)) - if self._paramsArray is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _paramsArray_new = calloc(len(val), sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS)) + if _paramsArray_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS))) + for idx in range(len(val)): + string.memcpy(&_paramsArray_new[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS)) + free(self._paramsArray) + self._paramsArray = _paramsArray_new self._paramsArray_length = len(val) - self._pvt_ptr[0].paramsArray = self._paramsArray - for idx in range(len(val)): - string.memcpy(&self._paramsArray[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS)) + self._pvt_ptr[0].paramsArray = _paramsArray_new + else: + for idx in range(len(val)): + string.memcpy(&self._paramsArray[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS)) @@ -20332,6 +20404,7 @@ cdef class CUDA_EXT_SEM_WAIT_NODE_PARAMS_st: return [CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS(_ptr=arr) for arr in arrs] @paramsArray.setter def paramsArray(self, val): + cdef cydriver.CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS* _paramsArray_new if len(val) == 0: free(self._paramsArray) self._paramsArray = NULL @@ -20339,14 +20412,22 @@ cdef class CUDA_EXT_SEM_WAIT_NODE_PARAMS_st: self._pvt_ptr[0].paramsArray = NULL else: if self._paramsArray_length != len(val): - free(self._paramsArray) - self._paramsArray = calloc(len(val), sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS)) - if self._paramsArray is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _paramsArray_new = calloc(len(val), sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS)) + if _paramsArray_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS))) + for idx in range(len(val)): + string.memcpy(&_paramsArray_new[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS)) + free(self._paramsArray) + self._paramsArray = _paramsArray_new self._paramsArray_length = len(val) - self._pvt_ptr[0].paramsArray = self._paramsArray - for idx in range(len(val)): - string.memcpy(&self._paramsArray[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS)) + self._pvt_ptr[0].paramsArray = _paramsArray_new + else: + for idx in range(len(val)): + string.memcpy(&self._paramsArray[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS)) @@ -20459,6 +20540,7 @@ cdef class CUDA_EXT_SEM_WAIT_NODE_PARAMS_v2_st: return [CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS(_ptr=arr) for arr in arrs] @paramsArray.setter def paramsArray(self, val): + cdef cydriver.CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS* _paramsArray_new if len(val) == 0: free(self._paramsArray) self._paramsArray = NULL @@ -20466,14 +20548,22 @@ cdef class CUDA_EXT_SEM_WAIT_NODE_PARAMS_v2_st: self._pvt_ptr[0].paramsArray = NULL else: if self._paramsArray_length != len(val): - free(self._paramsArray) - self._paramsArray = calloc(len(val), sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS)) - if self._paramsArray is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _paramsArray_new = calloc(len(val), sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS)) + if _paramsArray_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS))) + for idx in range(len(val)): + string.memcpy(&_paramsArray_new[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS)) + free(self._paramsArray) + self._paramsArray = _paramsArray_new self._paramsArray_length = len(val) - self._pvt_ptr[0].paramsArray = self._paramsArray - for idx in range(len(val)): - string.memcpy(&self._paramsArray[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS)) + self._pvt_ptr[0].paramsArray = _paramsArray_new + else: + for idx in range(len(val)): + string.memcpy(&self._paramsArray[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS)) @@ -22840,6 +22930,7 @@ cdef class CUDA_MEM_ALLOC_NODE_PARAMS_v1_st: return [CUmemAccessDesc(_ptr=arr) for arr in arrs] @accessDescs.setter def accessDescs(self, val): + cdef cydriver.CUmemAccessDesc* _accessDescs_new if len(val) == 0: free(self._accessDescs) self._accessDescs = NULL @@ -22847,14 +22938,22 @@ cdef class CUDA_MEM_ALLOC_NODE_PARAMS_v1_st: self._pvt_ptr[0].accessDescs = NULL else: if self._accessDescs_length != len(val): - free(self._accessDescs) - self._accessDescs = calloc(len(val), sizeof(cydriver.CUmemAccessDesc)) - if self._accessDescs is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _accessDescs_new = calloc(len(val), sizeof(cydriver.CUmemAccessDesc)) + if _accessDescs_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cydriver.CUmemAccessDesc))) + for idx in range(len(val)): + string.memcpy(&_accessDescs_new[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUmemAccessDesc)) + free(self._accessDescs) + self._accessDescs = _accessDescs_new self._accessDescs_length = len(val) - self._pvt_ptr[0].accessDescs = self._accessDescs - for idx in range(len(val)): - string.memcpy(&self._accessDescs[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUmemAccessDesc)) + self._pvt_ptr[0].accessDescs = _accessDescs_new + else: + for idx in range(len(val)): + string.memcpy(&self._accessDescs[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUmemAccessDesc)) @@ -23001,6 +23100,7 @@ cdef class CUDA_MEM_ALLOC_NODE_PARAMS_v2_st: return [CUmemAccessDesc(_ptr=arr) for arr in arrs] @accessDescs.setter def accessDescs(self, val): + cdef cydriver.CUmemAccessDesc* _accessDescs_new if len(val) == 0: free(self._accessDescs) self._accessDescs = NULL @@ -23008,14 +23108,22 @@ cdef class CUDA_MEM_ALLOC_NODE_PARAMS_v2_st: self._pvt_ptr[0].accessDescs = NULL else: if self._accessDescs_length != len(val): - free(self._accessDescs) - self._accessDescs = calloc(len(val), sizeof(cydriver.CUmemAccessDesc)) - if self._accessDescs is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _accessDescs_new = calloc(len(val), sizeof(cydriver.CUmemAccessDesc)) + if _accessDescs_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cydriver.CUmemAccessDesc))) + for idx in range(len(val)): + string.memcpy(&_accessDescs_new[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUmemAccessDesc)) + free(self._accessDescs) + self._accessDescs = _accessDescs_new self._accessDescs_length = len(val) - self._pvt_ptr[0].accessDescs = self._accessDescs - for idx in range(len(val)): - string.memcpy(&self._accessDescs[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUmemAccessDesc)) + self._pvt_ptr[0].accessDescs = _accessDescs_new + else: + for idx in range(len(val)): + string.memcpy(&self._accessDescs[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUmemAccessDesc)) @@ -24005,6 +24113,7 @@ cdef class CUcheckpointRestoreArgs_st: return [CUcheckpointGpuPair(_ptr=arr) for arr in arrs] @gpuPairs.setter def gpuPairs(self, val): + cdef cydriver.CUcheckpointGpuPair* _gpuPairs_new if len(val) == 0: free(self._gpuPairs) self._gpuPairs = NULL @@ -24012,14 +24121,22 @@ cdef class CUcheckpointRestoreArgs_st: self._pvt_ptr[0].gpuPairs = NULL else: if self._gpuPairs_length != len(val): - free(self._gpuPairs) - self._gpuPairs = calloc(len(val), sizeof(cydriver.CUcheckpointGpuPair)) - if self._gpuPairs is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _gpuPairs_new = calloc(len(val), sizeof(cydriver.CUcheckpointGpuPair)) + if _gpuPairs_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cydriver.CUcheckpointGpuPair))) + for idx in range(len(val)): + string.memcpy(&_gpuPairs_new[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUcheckpointGpuPair)) + free(self._gpuPairs) + self._gpuPairs = _gpuPairs_new self._gpuPairs_length = len(val) - self._pvt_ptr[0].gpuPairs = self._gpuPairs - for idx in range(len(val)): - string.memcpy(&self._gpuPairs[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUcheckpointGpuPair)) + self._pvt_ptr[0].gpuPairs = _gpuPairs_new + else: + for idx in range(len(val)): + string.memcpy(&self._gpuPairs[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUcheckpointGpuPair)) @@ -25121,6 +25238,7 @@ cdef class CUdevResource_st: return [CUdevResource_st(_ptr=arr) for arr in arrs] @nextResource.setter def nextResource(self, val): + cdef cydriver.CUdevResource_st* _nextResource_new if len(val) == 0: free(self._nextResource) self._nextResource = NULL @@ -25128,14 +25246,22 @@ cdef class CUdevResource_st: self._pvt_ptr[0].nextResource = NULL else: if self._nextResource_length != len(val): - free(self._nextResource) - self._nextResource = calloc(len(val), sizeof(cydriver.CUdevResource_st)) - if self._nextResource is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _nextResource_new = calloc(len(val), sizeof(cydriver.CUdevResource_st)) + if _nextResource_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cydriver.CUdevResource_st))) + for idx in range(len(val)): + string.memcpy(&_nextResource_new[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUdevResource_st)) + free(self._nextResource) + self._nextResource = _nextResource_new self._nextResource_length = len(val) - self._pvt_ptr[0].nextResource = self._nextResource - for idx in range(len(val)): - string.memcpy(&self._nextResource[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUdevResource_st)) + self._pvt_ptr[0].nextResource = _nextResource_new + else: + for idx in range(len(val)): + string.memcpy(&self._nextResource[idx], (val[idx])._pvt_ptr, sizeof(cydriver.CUdevResource_st)) @@ -35463,7 +35589,7 @@ def cuMemSetAccess(ptr, size_t size, desc : Optional[tuple[CUmemAccessDesc] | li See Also -------- - :py:obj:`~.cuMemSetAccess`, :py:obj:`~.cuMemCreate`, :py:obj:`~.py`:obj:`~.cuMemMap` + :py:obj:`~.cuMemSetAccess`, :py:obj:`~.cuMemCreate`, :py:obj:`~.cuMemMap` """ desc = [] if desc is None else desc if not all(isinstance(_x, (CUmemAccessDesc,)) for _x in desc): @@ -36718,11 +36844,11 @@ def cuMulticastCreate(prop : Optional[CUmulticastObjectProp]): :py:obj:`~.cuMulticastBindAddr`, or :py:obj:`~.cuMulticastBindAddr_v2`. and can be unbound via :py:obj:`~.cuMulticastUnbind`. The total amount of memory that can be bound per device is specified by - :py:obj:`~.py`:obj:`~.CUmulticastObjectProp.size`. This size must be a - multiple of the value returned by :py:obj:`~.cuMulticastGetGranularity` - with the flag :py:obj:`~.CU_MULTICAST_GRANULARITY_MINIMUM`. For best - performance however, the size should be aligned to the value returned - by :py:obj:`~.cuMulticastGetGranularity` with the flag + :py:obj:`~.CUmulticastObjectProp.size`. This size must be a multiple of + the value returned by :py:obj:`~.cuMulticastGetGranularity` with the + flag :py:obj:`~.CU_MULTICAST_GRANULARITY_MINIMUM`. For best performance + however, the size should be aligned to the value returned by + :py:obj:`~.cuMulticastGetGranularity` with the flag :py:obj:`~.CU_MULTICAST_GRANULARITY_RECOMMENDED`. After all participating devices have been added, multicast objects can @@ -38081,13 +38207,12 @@ def cuPointerGetAttribute(attribute not None : CUpointer_attribute, ptr): :py:obj:`~.CUDA_POINTER_ATTRIBUTE_P2P_TOKENS`. - `ptr` must be a pointer to memory obtained from - :py:obj:`~.py`:obj:`~.cuMemAlloc()`. Note that p2pToken and - vaSpaceToken are only valid for the lifetime of the source - allocation. A subsequent allocation at the same address may return - completely different tokens. Querying this attribute has a side - effect of setting the attribute - :py:obj:`~.CU_POINTER_ATTRIBUTE_SYNC_MEMOPS` for the region of memory - that `ptr` points to. + :py:obj:`~.cuMemAlloc()`. Note that p2pToken and vaSpaceToken are + only valid for the lifetime of the source allocation. A subsequent + allocation at the same address may return completely different + tokens. Querying this attribute has a side effect of setting the + attribute :py:obj:`~.CU_POINTER_ATTRIBUTE_SYNC_MEMOPS` for the region + of memory that `ptr` points to. - :py:obj:`~.CU_POINTER_ATTRIBUTE_SYNC_MEMOPS`: diff --git a/cuda_bindings/cuda/bindings/nvml.pyx b/cuda_bindings/cuda/bindings/nvml.pyx index a51d4264363..fe750b4c04a 100644 --- a/cuda_bindings/cuda/bindings/nvml.pyx +++ b/cuda_bindings/cuda/bindings/nvml.pyx @@ -3,7 +3,7 @@ # 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=e47a16bd9956de14a991ded5d1aef667cdd26a141e27db5b2015b91be6918d3c +# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=6dc4cfe2b034290ebc642bf79eea9ec031190b3a02615d5a2c8c0e6c26100bbb # <<<< PREAMBLE CONTENT >>>> @@ -840,7 +840,7 @@ class GpmMetricId(_cyb_FastEnum): GPM_METRIC_HMMA_TENSOR_UTIL = (NVML_GPM_METRIC_HMMA_TENSOR_UTIL, "Percentage of time the GPU's SMs were doing HMMA tensor operations. 0.0 - 100.0.") GPM_METRIC_DMMA_TENSOR_UTIL = (NVML_GPM_METRIC_DMMA_TENSOR_UTIL, "Percentage of time the GPU's SMs were doing DMMA tensor operations. 0.0 - 100.0.") GPM_METRIC_IMMA_TENSOR_UTIL = (NVML_GPM_METRIC_IMMA_TENSOR_UTIL, "Percentage of time the GPU's SMs were doing IMMA tensor operations. 0.0 - 100.0.") - GPM_METRIC_DRAM_BW_UTIL = (NVML_GPM_METRIC_DRAM_BW_UTIL, 'Percentage of DRAM bw used vs theoretical maximum. 0.0 - 100.0 *\u200d/.') + GPM_METRIC_DRAM_BW_UTIL = (NVML_GPM_METRIC_DRAM_BW_UTIL, 'Percentage of DRAM bw used vs theoretical maximum. `0.0 - 100.0 */`.') GPM_METRIC_FP64_UTIL = (NVML_GPM_METRIC_FP64_UTIL, "Percentage of time the GPU's SMs were doing non-tensor FP64 math. 0.0 - 100.0.") GPM_METRIC_FP32_UTIL = (NVML_GPM_METRIC_FP32_UTIL, "Percentage of time the GPU's SMs were doing non-tensor FP32 math. 0.0 - 100.0.") GPM_METRIC_FP16_UTIL = (NVML_GPM_METRIC_FP16_UTIL, "Percentage of time the GPU's SMs were doing non-tensor FP16 math. 0.0 - 100.0.") @@ -25091,7 +25091,7 @@ cpdef int device_get_virtualization_mode(intptr_t device) except? -1: device (intptr_t): Identifier of the target device. Returns: - int: Reference to virtualization mode. One of NVML_GPU_VIRTUALIZATION_?. + int: Reference to virtualization mode. One of ``NVML_GPU_VIRTUALIZATION_?``. .. seealso:: `nvmlDeviceGetVirtualizationMode` """ @@ -25125,7 +25125,7 @@ cpdef device_set_virtualization_mode(intptr_t device, int virtual_mode): Args: device (intptr_t): Identifier of the target device. - virtual_mode (GpuVirtualizationMode): virtualization mode. One of NVML_GPU_VIRTUALIZATION_?. + virtual_mode (GpuVirtualizationMode): virtualization mode. One of ``NVML_GPU_VIRTUALIZATION_?``. .. seealso:: `nvmlDeviceSetVirtualizationMode` """ @@ -28464,10 +28464,4 @@ cpdef str vgpu_type_get_name(unsigned int vgpu_type_id): __status__ = nvmlVgpuTypeGetName(vgpu_type_id, vgpu_type_name, size) check_status(__status__) return cpython.PyUnicode_FromStringAndSize(vgpu_type_name, size[0]) - - -# Cleanup some docstrings that don't parse as rst. -device_get_virtualization_mode.__doc__ = device_get_virtualization_mode.__doc__.replace("NVML_GPU_VIRTUALIZATION_?", "``NVML_GPU_VIRTUALIZATION_?``") -device_set_virtualization_mode.__doc__ = device_set_virtualization_mode.__doc__.replace("NVML_GPU_VIRTUALIZATION_?", "``NVML_GPU_VIRTUALIZATION_?``") -GpmMetricId.GPM_METRIC_DRAM_BW_UTIL.__doc__ = "Percentage of DRAM bw used vs theoretical maximum. ``0.0 - 100.0 *\u200d/``." del _cyb_FastEnum diff --git a/cuda_bindings/cuda/bindings/runtime.pxd b/cuda_bindings/cuda/bindings/runtime.pxd index 7cb680d1743..d704250e5cd 100644 --- a/cuda_bindings/cuda/bindings/runtime.pxd +++ b/cuda_bindings/cuda/bindings/runtime.pxd @@ -2,7 +2,7 @@ # SPDX-License-Identifier: Apache-2.0 # This code was automatically generated with version 13.3.0. Do not modify it directly. -# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=0f2431380680008795336b7acb5ccd83dba7a6e05d0c81c2b5f825bc95576ccd +# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=ea44514ba11796759754ddfb1ae93c9b7e1d510a09f7349a82a2201dd0c0c159 cimport cuda.bindings.cyruntime as cyruntime include "_lib/utils.pxd" @@ -4023,12 +4023,12 @@ cdef class cudaLaunchAttributeValue: with the following fields: - `x` - The X dimension of the preferred cluster, in blocks. Must be a divisor of the grid X dimension, and must be a multiple of the `x` field of - ::cudaLaunchAttributeValue::clusterDim. - `y` - The Y dimension - of the preferred cluster, in blocks. Must be a divisor of the grid - Y dimension, and must be a multiple of the `y` field of - ::cudaLaunchAttributeValue::clusterDim. - `z` - The Z dimension - of the preferred cluster, in blocks. Must be equal to the `z` field - of ::cudaLaunchAttributeValue::clusterDim. + cudaLaunchAttributeValue::clusterDim. - `y` - The Y dimension of + the preferred cluster, in blocks. Must be a divisor of the grid Y + dimension, and must be a multiple of the `y` field of + cudaLaunchAttributeValue::clusterDim. - `z` - The Z dimension of + the preferred cluster, in blocks. Must be equal to the `z` field of + cudaLaunchAttributeValue::clusterDim. launchCompletionEvent : anon_struct20 @@ -4830,12 +4830,12 @@ cdef class cudaStreamAttrValue(cudaLaunchAttributeValue): with the following fields: - `x` - The X dimension of the preferred cluster, in blocks. Must be a divisor of the grid X dimension, and must be a multiple of the `x` field of - ::cudaLaunchAttributeValue::clusterDim. - `y` - The Y dimension - of the preferred cluster, in blocks. Must be a divisor of the grid - Y dimension, and must be a multiple of the `y` field of - ::cudaLaunchAttributeValue::clusterDim. - `z` - The Z dimension - of the preferred cluster, in blocks. Must be equal to the `z` field - of ::cudaLaunchAttributeValue::clusterDim. + cudaLaunchAttributeValue::clusterDim. - `y` - The Y dimension of + the preferred cluster, in blocks. Must be a divisor of the grid Y + dimension, and must be a multiple of the `y` field of + cudaLaunchAttributeValue::clusterDim. - `z` - The Z dimension of + the preferred cluster, in blocks. Must be equal to the `z` field of + cudaLaunchAttributeValue::clusterDim. launchCompletionEvent : anon_struct20 @@ -4959,12 +4959,12 @@ cdef class cudaKernelNodeAttrValue(cudaLaunchAttributeValue): with the following fields: - `x` - The X dimension of the preferred cluster, in blocks. Must be a divisor of the grid X dimension, and must be a multiple of the `x` field of - ::cudaLaunchAttributeValue::clusterDim. - `y` - The Y dimension - of the preferred cluster, in blocks. Must be a divisor of the grid - Y dimension, and must be a multiple of the `y` field of - ::cudaLaunchAttributeValue::clusterDim. - `z` - The Z dimension - of the preferred cluster, in blocks. Must be equal to the `z` field - of ::cudaLaunchAttributeValue::clusterDim. + cudaLaunchAttributeValue::clusterDim. - `y` - The Y dimension of + the preferred cluster, in blocks. Must be a divisor of the grid Y + dimension, and must be a multiple of the `y` field of + cudaLaunchAttributeValue::clusterDim. - `z` - The Z dimension of + the preferred cluster, in blocks. Must be equal to the `z` field of + cudaLaunchAttributeValue::clusterDim. launchCompletionEvent : anon_struct20 diff --git a/cuda_bindings/cuda/bindings/runtime.pyx b/cuda_bindings/cuda/bindings/runtime.pyx index 6e11bdfc932..01157267892 100644 --- a/cuda_bindings/cuda/bindings/runtime.pyx +++ b/cuda_bindings/cuda/bindings/runtime.pyx @@ -2,7 +2,7 @@ # SPDX-License-Identifier: Apache-2.0 # This code was automatically generated with version 13.3.0. Do not modify it directly. -# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=783d27cfa66fa4f5818edcf88141aeb96dcc0e27e73a77a36571f82f30f3bb47 +# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=c251080d776c07c133fb4c06e4aa5cb7567642467a5fd4260405524c18edf1e8 from typing import Any, Optional import cython import ctypes @@ -1748,8 +1748,8 @@ class cudaLaunchAttributeID(_FastEnum): cyruntime.cudaLaunchAttributeID.cudaLaunchAttributeDeviceUpdatableKernelNode, 'Valid for graph nodes, launches. This attribute is graphs-only, and passing\n' 'it to a launch in a non-capturing stream will result in an error.\n' - ' :cudaLaunchAttributeValue::deviceUpdatableKernelNode::deviceUpdatable can\n' - 'only be set to 0 or 1. Setting the field to 1 indicates that the\n' + ':py:obj:`~.cudaLaunchAttributeValue.deviceUpdatableKernelNode.deviceUpdatable`\n' + 'can only be set to 0 or 1. Setting the field to 1 indicates that the\n' 'corresponding kernel node should be device-updatable. On success, a handle\n' 'will be returned via\n' ':py:obj:`~.cudaLaunchAttributeValue.deviceUpdatableKernelNode.devNode`\n' @@ -6590,8 +6590,8 @@ class cudaStreamAttrID(_FastEnum): cyruntime.cudaLaunchAttributeID.cudaLaunchAttributeDeviceUpdatableKernelNode, 'Valid for graph nodes, launches. This attribute is graphs-only, and passing\n' 'it to a launch in a non-capturing stream will result in an error.\n' - ' :cudaLaunchAttributeValue::deviceUpdatableKernelNode::deviceUpdatable can\n' - 'only be set to 0 or 1. Setting the field to 1 indicates that the\n' + ':py:obj:`~.cudaLaunchAttributeValue.deviceUpdatableKernelNode.deviceUpdatable`\n' + 'can only be set to 0 or 1. Setting the field to 1 indicates that the\n' 'corresponding kernel node should be device-updatable. On success, a handle\n' 'will be returned via\n' ':py:obj:`~.cudaLaunchAttributeValue.deviceUpdatableKernelNode.devNode`\n' @@ -6835,8 +6835,8 @@ class cudaKernelNodeAttrID(_FastEnum): cyruntime.cudaLaunchAttributeID.cudaLaunchAttributeDeviceUpdatableKernelNode, 'Valid for graph nodes, launches. This attribute is graphs-only, and passing\n' 'it to a launch in a non-capturing stream will result in an error.\n' - ' :cudaLaunchAttributeValue::deviceUpdatableKernelNode::deviceUpdatable can\n' - 'only be set to 0 or 1. Setting the field to 1 indicates that the\n' + ':py:obj:`~.cudaLaunchAttributeValue.deviceUpdatableKernelNode.deviceUpdatable`\n' + 'can only be set to 0 or 1. Setting the field to 1 indicates that the\n' 'corresponding kernel node should be device-updatable. On success, a handle\n' 'will be returned via\n' ':py:obj:`~.cudaLaunchAttributeValue.deviceUpdatableKernelNode.devNode`\n' @@ -11338,6 +11338,7 @@ cdef class cudaMemAllocNodeParams: return [cudaMemAccessDesc(_ptr=arr) for arr in arrs] @accessDescs.setter def accessDescs(self, val): + cdef cyruntime.cudaMemAccessDesc* _accessDescs_new if len(val) == 0: free(self._accessDescs) self._accessDescs = NULL @@ -11345,14 +11346,22 @@ cdef class cudaMemAllocNodeParams: self._pvt_ptr[0].accessDescs = NULL else: if self._accessDescs_length != len(val): - free(self._accessDescs) - self._accessDescs = calloc(len(val), sizeof(cyruntime.cudaMemAccessDesc)) - if self._accessDescs is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _accessDescs_new = calloc(len(val), sizeof(cyruntime.cudaMemAccessDesc)) + if _accessDescs_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cyruntime.cudaMemAccessDesc))) + for idx in range(len(val)): + string.memcpy(&_accessDescs_new[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaMemAccessDesc)) + free(self._accessDescs) + self._accessDescs = _accessDescs_new self._accessDescs_length = len(val) - self._pvt_ptr[0].accessDescs = self._accessDescs - for idx in range(len(val)): - string.memcpy(&self._accessDescs[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaMemAccessDesc)) + self._pvt_ptr[0].accessDescs = _accessDescs_new + else: + for idx in range(len(val)): + string.memcpy(&self._accessDescs[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaMemAccessDesc)) @@ -11487,6 +11496,7 @@ cdef class cudaMemAllocNodeParamsV2: return [cudaMemAccessDesc(_ptr=arr) for arr in arrs] @accessDescs.setter def accessDescs(self, val): + cdef cyruntime.cudaMemAccessDesc* _accessDescs_new if len(val) == 0: free(self._accessDescs) self._accessDescs = NULL @@ -11494,14 +11504,22 @@ cdef class cudaMemAllocNodeParamsV2: self._pvt_ptr[0].accessDescs = NULL else: if self._accessDescs_length != len(val): - free(self._accessDescs) - self._accessDescs = calloc(len(val), sizeof(cyruntime.cudaMemAccessDesc)) - if self._accessDescs is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _accessDescs_new = calloc(len(val), sizeof(cyruntime.cudaMemAccessDesc)) + if _accessDescs_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cyruntime.cudaMemAccessDesc))) + for idx in range(len(val)): + string.memcpy(&_accessDescs_new[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaMemAccessDesc)) + free(self._accessDescs) + self._accessDescs = _accessDescs_new self._accessDescs_length = len(val) - self._pvt_ptr[0].accessDescs = self._accessDescs - for idx in range(len(val)): - string.memcpy(&self._accessDescs[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaMemAccessDesc)) + self._pvt_ptr[0].accessDescs = _accessDescs_new + else: + for idx in range(len(val)): + string.memcpy(&self._accessDescs[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaMemAccessDesc)) @@ -16237,6 +16255,7 @@ cdef class cudaDevResource_st: return [cudaDevResource_st(_ptr=arr) for arr in arrs] @nextResource.setter def nextResource(self, val): + cdef cyruntime.cudaDevResource_st* _nextResource_new if len(val) == 0: free(self._nextResource) self._nextResource = NULL @@ -16244,14 +16263,22 @@ cdef class cudaDevResource_st: self._pvt_ptr[0].nextResource = NULL else: if self._nextResource_length != len(val): - free(self._nextResource) - self._nextResource = calloc(len(val), sizeof(cyruntime.cudaDevResource_st)) - if self._nextResource is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _nextResource_new = calloc(len(val), sizeof(cyruntime.cudaDevResource_st)) + if _nextResource_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cyruntime.cudaDevResource_st))) + for idx in range(len(val)): + string.memcpy(&_nextResource_new[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaDevResource_st)) + free(self._nextResource) + self._nextResource = _nextResource_new self._nextResource_length = len(val) - self._pvt_ptr[0].nextResource = self._nextResource - for idx in range(len(val)): - string.memcpy(&self._nextResource[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaDevResource_st)) + self._pvt_ptr[0].nextResource = _nextResource_new + else: + for idx in range(len(val)): + string.memcpy(&self._nextResource[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaDevResource_st)) @@ -16861,6 +16888,7 @@ cdef class cudaExternalSemaphoreSignalNodeParams: return [cudaExternalSemaphoreSignalParams(_ptr=arr) for arr in arrs] @paramsArray.setter def paramsArray(self, val): + cdef cyruntime.cudaExternalSemaphoreSignalParams* _paramsArray_new if len(val) == 0: free(self._paramsArray) self._paramsArray = NULL @@ -16868,14 +16896,22 @@ cdef class cudaExternalSemaphoreSignalNodeParams: self._pvt_ptr[0].paramsArray = NULL else: if self._paramsArray_length != len(val): - free(self._paramsArray) - self._paramsArray = calloc(len(val), sizeof(cyruntime.cudaExternalSemaphoreSignalParams)) - if self._paramsArray is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _paramsArray_new = calloc(len(val), sizeof(cyruntime.cudaExternalSemaphoreSignalParams)) + if _paramsArray_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cyruntime.cudaExternalSemaphoreSignalParams))) + for idx in range(len(val)): + string.memcpy(&_paramsArray_new[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaExternalSemaphoreSignalParams)) + free(self._paramsArray) + self._paramsArray = _paramsArray_new self._paramsArray_length = len(val) - self._pvt_ptr[0].paramsArray = self._paramsArray - for idx in range(len(val)): - string.memcpy(&self._paramsArray[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaExternalSemaphoreSignalParams)) + self._pvt_ptr[0].paramsArray = _paramsArray_new + else: + for idx in range(len(val)): + string.memcpy(&self._paramsArray[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaExternalSemaphoreSignalParams)) @@ -16988,6 +17024,7 @@ cdef class cudaExternalSemaphoreSignalNodeParamsV2: return [cudaExternalSemaphoreSignalParams(_ptr=arr) for arr in arrs] @paramsArray.setter def paramsArray(self, val): + cdef cyruntime.cudaExternalSemaphoreSignalParams* _paramsArray_new if len(val) == 0: free(self._paramsArray) self._paramsArray = NULL @@ -16995,14 +17032,22 @@ cdef class cudaExternalSemaphoreSignalNodeParamsV2: self._pvt_ptr[0].paramsArray = NULL else: if self._paramsArray_length != len(val): - free(self._paramsArray) - self._paramsArray = calloc(len(val), sizeof(cyruntime.cudaExternalSemaphoreSignalParams)) - if self._paramsArray is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _paramsArray_new = calloc(len(val), sizeof(cyruntime.cudaExternalSemaphoreSignalParams)) + if _paramsArray_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cyruntime.cudaExternalSemaphoreSignalParams))) + for idx in range(len(val)): + string.memcpy(&_paramsArray_new[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaExternalSemaphoreSignalParams)) + free(self._paramsArray) + self._paramsArray = _paramsArray_new self._paramsArray_length = len(val) - self._pvt_ptr[0].paramsArray = self._paramsArray - for idx in range(len(val)): - string.memcpy(&self._paramsArray[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaExternalSemaphoreSignalParams)) + self._pvt_ptr[0].paramsArray = _paramsArray_new + else: + for idx in range(len(val)): + string.memcpy(&self._paramsArray[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaExternalSemaphoreSignalParams)) @@ -17115,6 +17160,7 @@ cdef class cudaExternalSemaphoreWaitNodeParams: return [cudaExternalSemaphoreWaitParams(_ptr=arr) for arr in arrs] @paramsArray.setter def paramsArray(self, val): + cdef cyruntime.cudaExternalSemaphoreWaitParams* _paramsArray_new if len(val) == 0: free(self._paramsArray) self._paramsArray = NULL @@ -17122,14 +17168,22 @@ cdef class cudaExternalSemaphoreWaitNodeParams: self._pvt_ptr[0].paramsArray = NULL else: if self._paramsArray_length != len(val): - free(self._paramsArray) - self._paramsArray = calloc(len(val), sizeof(cyruntime.cudaExternalSemaphoreWaitParams)) - if self._paramsArray is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _paramsArray_new = calloc(len(val), sizeof(cyruntime.cudaExternalSemaphoreWaitParams)) + if _paramsArray_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cyruntime.cudaExternalSemaphoreWaitParams))) + for idx in range(len(val)): + string.memcpy(&_paramsArray_new[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaExternalSemaphoreWaitParams)) + free(self._paramsArray) + self._paramsArray = _paramsArray_new self._paramsArray_length = len(val) - self._pvt_ptr[0].paramsArray = self._paramsArray - for idx in range(len(val)): - string.memcpy(&self._paramsArray[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaExternalSemaphoreWaitParams)) + self._pvt_ptr[0].paramsArray = _paramsArray_new + else: + for idx in range(len(val)): + string.memcpy(&self._paramsArray[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaExternalSemaphoreWaitParams)) @@ -17242,6 +17296,7 @@ cdef class cudaExternalSemaphoreWaitNodeParamsV2: return [cudaExternalSemaphoreWaitParams(_ptr=arr) for arr in arrs] @paramsArray.setter def paramsArray(self, val): + cdef cyruntime.cudaExternalSemaphoreWaitParams* _paramsArray_new if len(val) == 0: free(self._paramsArray) self._paramsArray = NULL @@ -17249,14 +17304,22 @@ cdef class cudaExternalSemaphoreWaitNodeParamsV2: self._pvt_ptr[0].paramsArray = NULL else: if self._paramsArray_length != len(val): - free(self._paramsArray) - self._paramsArray = calloc(len(val), sizeof(cyruntime.cudaExternalSemaphoreWaitParams)) - if self._paramsArray is NULL: + # Allocate and fill a new buffer before touching the + # live state so a failure leaves this object unchanged + # (strong exception guarantee); the old buffer is only + # freed once the resize is known to succeed. + _paramsArray_new = calloc(len(val), sizeof(cyruntime.cudaExternalSemaphoreWaitParams)) + if _paramsArray_new is NULL: raise MemoryError('Failed to allocate length x size memory: ' + str(len(val)) + 'x' + str(sizeof(cyruntime.cudaExternalSemaphoreWaitParams))) + for idx in range(len(val)): + string.memcpy(&_paramsArray_new[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaExternalSemaphoreWaitParams)) + free(self._paramsArray) + self._paramsArray = _paramsArray_new self._paramsArray_length = len(val) - self._pvt_ptr[0].paramsArray = self._paramsArray - for idx in range(len(val)): - string.memcpy(&self._paramsArray[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaExternalSemaphoreWaitParams)) + self._pvt_ptr[0].paramsArray = _paramsArray_new + else: + for idx in range(len(val)): + string.memcpy(&self._paramsArray[idx], (val[idx])._pvt_ptr, sizeof(cyruntime.cudaExternalSemaphoreWaitParams)) @@ -19167,12 +19230,12 @@ cdef class cudaLaunchAttributeValue: with the following fields: - `x` - The X dimension of the preferred cluster, in blocks. Must be a divisor of the grid X dimension, and must be a multiple of the `x` field of - ::cudaLaunchAttributeValue::clusterDim. - `y` - The Y dimension - of the preferred cluster, in blocks. Must be a divisor of the grid - Y dimension, and must be a multiple of the `y` field of - ::cudaLaunchAttributeValue::clusterDim. - `z` - The Z dimension - of the preferred cluster, in blocks. Must be equal to the `z` field - of ::cudaLaunchAttributeValue::clusterDim. + cudaLaunchAttributeValue::clusterDim. - `y` - The Y dimension of + the preferred cluster, in blocks. Must be a divisor of the grid Y + dimension, and must be a multiple of the `y` field of + cudaLaunchAttributeValue::clusterDim. - `z` - The Z dimension of + the preferred cluster, in blocks. Must be equal to the `z` field of + cudaLaunchAttributeValue::clusterDim. launchCompletionEvent : anon_struct20 diff --git a/cuda_bindings/docs/source/module/runtime.rst b/cuda_bindings/docs/source/module/runtime.rst index 80b11f16644..000c5fa188f 100644 --- a/cuda_bindings/docs/source/module/runtime.rst +++ b/cuda_bindings/docs/source/module/runtime.rst @@ -1,7 +1,7 @@ .. SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. .. SPDX-License-Identifier: Apache-2.0 -.. CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=0eccd5db44f7406acb693f5dd95ad2fa17c787c11659578cf54e116eb0b06e01 +.. CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=86b1e47dc2cb6ed343cc979b646bdfb99db65d352605464573a8e1d5b248fce2 ------- runtime ------- @@ -4185,7 +4185,7 @@ Data types used by CUDA Runtime Valid for graph nodes, launches. This attribute is graphs-only, and passing it to a launch in a non-capturing stream will result in an error. - :cudaLaunchAttributeValue::deviceUpdatableKernelNode::deviceUpdatable can only be set to 0 or 1. Setting the field to 1 indicates that the corresponding kernel node should be device-updatable. On success, a handle will be returned via :py:obj:`~.cudaLaunchAttributeValue.deviceUpdatableKernelNode.devNode` which can be passed to the various device-side update functions to update the node's kernel parameters from within another kernel. For more information on the types of device updates that can be made, as well as the relevant limitations thereof, see :py:obj:`~.cudaGraphKernelNodeUpdatesApply`. + :py:obj:`~.cudaLaunchAttributeValue.deviceUpdatableKernelNode.deviceUpdatable` can only be set to 0 or 1. Setting the field to 1 indicates that the corresponding kernel node should be device-updatable. On success, a handle will be returned via :py:obj:`~.cudaLaunchAttributeValue.deviceUpdatableKernelNode.devNode` which can be passed to the various device-side update functions to update the node's kernel parameters from within another kernel. For more information on the types of device updates that can be made, as well as the relevant limitations thereof, see :py:obj:`~.cudaGraphKernelNodeUpdatesApply`. Nodes which are device-updatable have additional restrictions compared to regular kernel nodes. Firstly, device-updatable nodes cannot be removed from their graph via :py:obj:`~.cudaGraphDestroyNode`. Additionally, once opted-in to this functionality, a node cannot opt out, and any attempt to set the deviceUpdatable attribute to 0 will result in an error. Device-updatable kernel nodes also cannot have their attributes copied to/from another kernel node via :py:obj:`~.cudaGraphKernelNodeCopyAttributes`. Graphs containing one or more device-updatable nodes also do not allow multiple instantiation, and neither the graph nor its instantiated version can be passed to :py:obj:`~.cudaGraphExecUpdate`. diff --git a/cuda_bindings/tests/test_cuda.py b/cuda_bindings/tests/test_cuda.py index 192ad0f72fe..e201d569216 100644 --- a/cuda_bindings/tests/test_cuda.py +++ b/cuda_bindings/tests/test_cuda.py @@ -1355,3 +1355,56 @@ def test_dealloc_clears_array_field_in_external_struct(): f"external struct still holds a dangling pointer ({attrs_after:#x}) " "where attrs was, after the aliasing wrapper was destroyed" ) + + +def test_array_setter_preserves_state_on_failed_resize(): + # Regression test for an OOM state-desync in the generated array-member + # setters (Glasswing V1.1 / NVBUG 6268871). On a length-changing assignment + # the setter freed the old buffer, then allocated the new one; if the + # allocation failed it raised *before* updating the cached length and the + # C-struct pointer, leaving the wrapper with a freed buffer but a stale + # length. A later same-length assignment then took the "no resize" path and + # memcpy'd through the NULL/freed buffer -> near-NULL segfault. + # + # The fix builds and fills the new buffer first and only frees/swaps the old + # one once the resize is known to succeed, so a failed (over)allocation + # leaves the object completely unchanged (strong exception guarantee). + # + # CUlaunchConfig.attrs is exercised as one representative of the many + # affected setters. calloc failure is forced without real host OOM by + # assigning an object whose __len__ reports a huge length, so + # calloc(len, sizeof) overflows size_t and returns NULL. The reproducer runs + # in a subprocess so a segfault surfaces as a non-zero return code instead + # of tearing down the pytest process. + code = textwrap.dedent( + """ + import cuda.bindings.driver as cuda + + class HugeLen: + # __len__ only: the setter raises on the overflowing calloc before + # it ever indexes this object. + def __len__(self): + return 2**62 + + params = cuda.CUlaunchConfig() + # Establish a live 4-element buffer and cached length. + params.attrs = [cuda.CUlaunchAttribute() for _ in range(4)] + + # Force a failed resize (calloc overflow -> NULL -> MemoryError). + try: + params.attrs = HugeLen() + except MemoryError: + pass + else: + raise SystemExit("expected MemoryError on overflowing resize") + + # Same-length reassignment. Pre-fix this took the no-resize path and + # wrote through the freed/NULL buffer -> segfault. Post-fix the object + # was left intact, so this succeeds. + params.attrs = [cuda.CUlaunchAttribute() for _ in range(4)] + """ + ) + proc = subprocess.run([sys.executable, "-c", code], capture_output=True, cwd=os.path.dirname(__file__)) # noqa: S603 + assert proc.returncode == 0, ( + f"reproducer subprocess exited with code {proc.returncode}; stderr: {proc.stderr.decode(errors='replace')}" + )