B \\@s~dZddlmZmZmZddlZddlZddlZddlZddl m Z ddl m Z ddl Zddl Z ddlmZddlmZdd l mZmZmZdd lmZyeed dZWnek rd d ZYnXd dZddZddZGdddeZGdddeZ eddZ!GdddeZ"GdddeZ#Gdddeej$Z%d-dd Z&d.d!d"Z'd#d$Z(d%Z)d&d'Z*d/d)d*Z+d+d,Z,dS)0z A CUDA ND Array is recognized by checking the __cuda_memory__ attribute on the object. If it exists and evaluate to True, it must define shape, strides, dtype and size attributes similar to a NumPy ndarray. )print_functionabsolute_importdivisionN)six)c_void_p)driver)devices) dummyarraytypes numpy_support)to_fixed_tuple lru_cachecCs|S)N)funcrr=lib/python3.7/site-packages/numba/cuda/cudadrv/devicearray.pyrscCs t|ddS)z$Check if an object is a CUDA ndarray__cuda_ndarray__F)getattr)objrrris_cuda_ndarraysrcsDtfdd}|dt|dt|dtj|dtjdS)z,Verify the CUDA ndarray interface for an objcs6t|st|tt||s2td||fdS)Nz%s must be of type %s)hasattrAttributeError isinstancer)attrtyp)rrr requires_attr(s z4verify_cuda_ndarray_interface..requires_attrshapestridesdtypesizeN)require_cuda_ndarraytuplenprr integer_types)rrr)rrverify_cuda_ndarray_interface$s     r$cCst|stddS)z9Raises ValueError is is_cuda_ndarray(obj) evaluates Falsezrequire an cuda ndarray objectN)r ValueError)rrrrr 4sr c@seZdZdZdZdZd%ddZeddZd&d d Z ed d Z d'd dZ ddZ eddZ eddZejd(ddZejd)ddZd*ddZd+ddZddZdd Zd,d!d"Zd#d$ZdS)-DeviceNDArrayBasez$A on GPU NDArray representation TrNcCs t|tjr|f}t|tjr$|f}t||_t||jkrDtdtjd|||j |_ t ||_ t ||_ t||_tt|j |_|jdkr|dkrt|j |j |jj |_t|j}qt||_ntjttddd}d|_||_||_||_dS)aN Args ---- shape array shape. strides array strides. dtype data type as np.dtype. stream cuda stream. writeback Deprecated. gpu_data user provided device memory for the ndarray data buffer zstrides not match ndimrN)contextZpointerr)rrr#lenndimr%r ArrayZ from_descitemsize_dummyr!rrr"rintprodr_driverZmemory_size_from_info alloc_sizer get_contextZmemallocZdevice_memory_sizeZ MemoryPointerrgpu_data_DeviceNDArrayBase__writebackstream)selfrrrr4 writebackr2rrr__init__@s4           zDeviceNDArrayBase.__init__cCs(t|jt|j|jjdf|jjddS)NFr)rrdataZtypestrversion)r!rrdevice_ctypes_pointervaluerstr)r5rrr__cuda_array_interface__ts  z*DeviceNDArrayBase.__cuda_array_interface__cCst|}||_|S)zBind a CUDA stream to this object so that all subsequent operation on this array defaults to the given stream. )copyr4)r5r4Zclonerrrbind~s zDeviceNDArrayBase.bindcCs|S)N) transpose)r5rrrTszDeviceNDArrayBase.TcCsx|rt|tt|jkr|S|jdkr2tdnB|dk r`t|tt|jkr`td|fnddlm}||SdS)Nz2transposing a non-2D DeviceNDArray isn't supportedzinvalid axes list %rr)r@)r!ranger)NotImplementedErrorsetr%Znumba.cuda.kernels.transposer@)r5Zaxesr@rrrr@s   zDeviceNDArrayBase.transposecCs|s |jS|S)N)r4)r5r4rrr_default_streamsz!DeviceNDArrayBase._default_streamcCst|j}t||jdS)zn Magic attribute expected by Numba to get the numba type that represents this object. A)r from_dtyperr r*r))r5rrrr _numba_type_s zDeviceNDArrayBase._numba_type_cCs|jdkrtdS|jjSdS)z:Returns the ctypes pointer to the GPU data buffer Nr)r2rr:)r5rrrr:s z'DeviceNDArrayBase.device_ctypes_pointercCs|jdkrdSt|||}t|t|}}t|rdt|t||tj|||j|dnFt j ||j drxdndd|j d d }t||tj |||j|ddS) zCopy `ary` to `self`. If `ary` is a CUDA memory, perform a device-to-device transfer. Otherwise, perform a a host-to-device transfer. rN)r4 C_CONTIGUOUSCFTZ WRITEABLE)ordersubokr>) rsentry_contiguousrF array_corer/is_device_memorycheck_array_compatibilityZdevice_to_devicer0r"arrayflagsZhost_to_device)r5aryr4Z self_coreZary_corerrrcopy_to_devices      z DeviceNDArrayBase.copy_to_devicecCstdd|jDr(d}t||j|jdks:td||}|dkr`tj|jtj d}nt |||}|jdkrt j |||j|d|dkr|j dkrtj|j|j|d }ntj|j|j|j|d }|S) a^Copy ``self`` to ``ary`` or create a new Numpy ndarray if ``ary`` is ``None``. If a CUDA ``stream`` is given, then the transfer will be made asynchronously as part as the given stream. Otherwise, the transfer is synchronous: the function returns after the copy is finished. Always returns the host array. Example:: import numpy as np from numba import cuda arr = np.arange(1000) d_arr = cuda.to_device(arr) my_kernel[100, 100](d_arr) result_array = d_arr.copy_to_host() css|]}|dkVqdS)rNr).0srrr sz1DeviceNDArrayBase.copy_to_host..z2D->H copy not implemented for negative strides: {}rzNegative memory sizeN)rr)r4)rrbuffer)rrrrZ)anyrrDformatr0AssertionErrorrFr"emptyZbyterRr/device_to_hostrndarrayrr)r5rUr4msghostaryrrr copy_to_hosts$        zDeviceNDArrayBase.copy_to_hostcCs<||}tdt|jdkr(td|j|j|ddS)Nz+to_host() is deprecated and will be removedzno associated writeback array)r4)rFwarningswarnDeprecationWarningr3r%rc)r5r4rrrto_hosts   zDeviceNDArrayBase.to_hostc cs||}|jdkrtd|jd|jjkr6tdttt |j |}|j}|jj}x^t |D]R}||}t |||j }||f} |j ||||} t| ||j|| dVqfWdS)zSplit the array into equal partition of the `section` size. If the array cannot be equally divided, the last section will be smaller. rzonly support 1d arrayrzonly support unit stride)rr4r2N)rFr)r%rrr+r-mathZceilfloatrrCminr2view DeviceNDArray) r5Zsectionr4Znsectrr+iZbeginendrr2rrrsplits    zDeviceNDArrayBase.splitcCs|jS)zEReturns a device memory object that is used as the argument. )r2)r5rrr as_cuda_argszDeviceNDArrayBase.as_cuda_argcCs0t|j}t|j|j|jd}t||dS)z Returns a *IpcArrayHandle* object that is safe to serialize and transfer to another process to share the local allocation. Note: this feature is only available on Linux. )rrr) ipc_handle array_desc) r r1get_ipc_handler2dictrrrIpcArrayHandle)r5ZipchZdescrrrrssz DeviceNDArrayBase.get_ipc_handlecCs2|jj|d\}}t|j|j|j|||jdS)a( Remove axes of size one from the array shape. Parameters ---------- axis : None or int or tuple of ints, optional Subset of dimensions to remove. A `ValueError` is raised if an axis with size greater than one is selected. If `None`, all axes with size one are removed. stream : cuda stream or 0, optional Default stream for the returned view of the array. Returns ------- DeviceNDArray Squeezed view into the array. )axis)rrrr4r2)r,squeezerlrrrrFr2)r5rvr4Z new_dummy_rrrrw#szDeviceNDArrayBase.squeezecCs:t|}|j|jjkr tdt|j|j||j|jdS)zeReturns a new object by reinterpretting the dtype without making a copy of the data. z new dtype itemsize doesn't match)rrrr4r2) r"rr+ TypeErrorrlrrr4r2)r5rrrrrk?s zDeviceNDArrayBase.view)rNN)r)N)r)Nr)r)r)Nr)__name__ __module__ __qualname____doc__Z__cuda_memory__rr7propertyr=r?rAr@rFrIr:r require_contextrVrcrgrorprsrwrkrrrrr&:s, 3      -   r&cs:eZdZdZd fdd ZeddZedd ZZS) DeviceRecordz An on-GPU record type rNcs$d}d}tt||||||dS)Nr)superrr7)r5rr4r2rr) __class__rrr7SszDeviceRecord.__init__cCs t|jjS)z For `numpy.ndarray` compatibility. Ideally this would return a `np.core.multiarray.flagsobj`, but that needs to be constructed with an existing `numpy.ndarray` (as the C- and F- contiguous flags aren't writeable). )rtr,rT)r5rrrrTYszDeviceRecord.flagscCs t|jS)zn Magic attribute expected by Numba to get the numba type that represents this object. )r rHr)r5rrrrIcszDeviceRecord._numba_type_)rN) rzr{r|r}r7r~rTrI __classcell__rr)rrrOs rcs$ddlmjfdd}|S)z A separate method so we don't need to compile code every assignment (!). :param ndim: We need to have static array sizes for cuda.local.array, so bake in the number of dimensions into the kernel r)cudacsd}d}x t|jD]}||j|9}qW||kr.kernel)numbarZjit)r)rr)rr)r_assign_kernells rc@seZdZdZddZeddZddZdd d Zd d Z d dZ d ddZ e j ddZd!ddZd"ddZe j ddZd#ddZd$ddZdS)%rlz An on-GPU array type cCs|jjS)zA Return true if the array is Fortran-contiguous. )r,Z is_f_contig)r5rrris_f_contiguousszDeviceNDArray.is_f_contiguouscCs t|jjS)z For `numpy.ndarray` compatibility. Ideally this would return a `np.core.multiarray.flagsobj`, but that needs to be constructed with an existing `numpy.ndarray` (as the C- and F- contiguous flags aren't writeable). )rtr,rT)r5rrrrTszDeviceNDArray.flagscCs|jjS)z; Return true if the array is C-contiguous. )r,Z is_c_contig)r5rrris_c_contiguousszDeviceNDArray.is_c_contiguousNcCs||S)zE :return: an `numpy.ndarray`, so copies to the host. )rc __array__)r5rrrrrszDeviceNDArray.__array__cCs |jdS)Nr)r)r5rrr__len__szDeviceNDArray.__len__cOst|dkr&t|dttfr&|d}t|}||jkrP||j|j|j|jdS|j j ||\}}||j j gkr||j|j|j|jdSt ddS)z Reshape the array without changing its contents, similarly to :meth:`numpy.ndarray.reshape`. Example:: d_arr = d_arr.reshape(20, 50, order='F') rr)rrrr2zoperation requires copyingN) r(rr!listtyperrrr2r,reshapeextentrD)r5ZnewshapeZkwsclsnewarrextentsrrrrs   zDeviceNDArray.reshaperKrcCsX||}t|}|jj|d\}}||jjgkrL||j|j|j|j|dSt ddS)zr Flatten the array without changing its contents, similar to :meth:`numpy.ndarray.ravel`. )rM)rrrr2r4zoperation requires copyingN) rFrr,ravelrrrrr2rD)r5rMr4rrrrrrrs  zDeviceNDArray.ravelcCs ||S)N) _do_getitem)r5itemrrr __getitem__szDeviceNDArray.__getitem__cCs |||S)z0Do `__getitem__(item)` with CUDA stream )r)r5rr4rrrgetitemszDeviceNDArray.getitemcCs||}|j|}t|}t|}t|dkr|jj|d}|j szt j d|j d}t j|||jj|d|dS||j|j|j ||dSn&|jj|j}||j|j|j ||dSdS)Nrr)r)dstsrcrr4)rrrr2r4)rFr,rrZiter_contiguous_extentrr(r2rkZis_arrayr"r^rr/r_r+rrr)r5rr4arrrrnewdatarbrrrrs"      zDeviceNDArray._do_getitemcCs |||S)N) _do_setitem)r5keyr;rrr __setitem__szDeviceNDArray.__setitem__cCs||||S)z6Do `__setitem__(key, value)` with CUDA stream )Z _so_getitem)r5rr;r4rrrsetitemszDeviceNDArray.setitemcCs2||}|j|}|jj|j}t|tjr@d}|j j f}n |j }|j }t ||||j ||d}t||d\} } | j|jkrtd| j|jftj|jtjd} | j | | j d<| j| } xDtt|j | j D].\} \} }|dkr| |krtd|| | fqWt|j }t|jj||d|| dS)N)r)rrrr2r4)r4z$Can't assign %s-D array to %s-D self)rrzCCan't copy sequence with size %d to array axis %d with dimension %d)rFr,rr2rkrrr ZElementrr+rrr auto_devicer)r%r"Zonesrr enumeratezipr.rZforall)r5rr;r4rrrrrrrxZ rhs_shapermlrrrrrrs<         zDeviceNDArray._do_setitem)N)rKr)r)r)r)r)rzr{r|r}rr~rTrrrrrr rrrrrrrrrrrrls      rlc@s8eZdZdZddZddZddZdd Zd d Zd S) rua" An IPC array handle that can be serialized and transfer to another process in the same machine for share a GPU allocation. On the destination process, use the *.open()* method to creates a new *DeviceNDArray* object that shares the allocation from the original process. To release the resources, call the *.close()* method. After that, the destination can no longer use the shared array object. (Note: the underlying weakref to the resource is now dead.) This object implements the context-manager interface that calls the *.open()* and *.close()* method automatically:: with the_ipc_array_handle as ipc_array: # use ipc_array here as a normal gpu array object some_code(ipc_array) # ipc_array is dead at this point cCs||_||_dS)N) _array_desc _ipc_handle)r5rqrrrrrr7IszIpcArrayHandle.__init__cCs$|jt}tfd|i|jS)z Returns a new *DeviceNDArray* that shares the allocation from the original process. Must not be used on the original process. r2)ropenr r1rlr)r5ZdptrrrrrMszIpcArrayHandle.opencCs|jdS)z5 Closes the IPC handle to the array. N)rclose)r5rrrrUszIpcArrayHandle.closecCs|S)N)r)r5rrr __enter__[szIpcArrayHandle.__enter__cCs |dS)N)r)r5rr; tracebackrrr__exit__^szIpcArrayHandle.__exit__N) rzr{r|r}r7rrrrrrrrru6s ruc@seZdZdZdddZdS) MappedNDArrayz4 A host array that uses CUDA mapped memory. rcCs ||_dS)N)r2)r5r2r4rrr device_setupgszMappedNDArray.device_setupN)r)rzr{r|r}rrrrrrbsrcCs.|jdkr|d}t|j|j|j|||dS)z/Create a DeviceNDArray object that is like ary.rr)r6r4r2)r)rrlrrr)rUr4r2rrrfrom_array_likeks  rcCst|j||dS)z.Create a DeviceRecord object that is like rec.)r4r2)rr)Zrecr4r2rrrfrom_record_likessrcCsD|js |Sg}x(|jD]}||dkr*dntdqW|t|S)aG Extract the repeated core of a broadcast array. Broadcast arrays are by definition non-contiguous due to repeated dimensions, i.e., dimensions with stride 0. In order to ascertain memory contiguity and copy the underlying data from such arrays, we must create a view without the repeated dimensions. rN)rappendslicer!)rUZ core_indexZstriderrrrPxs  rPzArray contains non-contiguous buffer and cannot be transferred as a single memory region. Please ensure contiguous buffer with numpy .ascontiguousarray()cCs(t|}|jds$|jds$ttdS)NrJZ F_CONTIGUOUS)rPrTr%errmsg_contiguous_buffer)rUZcorerrrrOsrOTcCst|r|dfSt|dr,tj|dfSt|tjrFt ||d}n$tj |ddd}t |t ||d}|r||j ||d|dfSdS)z Create a DeviceRecord or DeviceArray like obj and optionally copy data from host to device. If obj already represents device memory, it is returned and no copy is made. Fr=)r4T)r>rNN)r/rQrrrZ as_cuda_arrayrr"ZvoidrrSrOrrV)rr4r>Zdevobjrrrrs    rcCsv||}}|j|jkr2td|j|jf|j|jkrRtd|j|jf|j|jkrrtd|j|jfdS)Nzincompatible dtype: %s vs. %szincompatible shape: %s vs. %szincompatible strides: %s vs. %s)rwrryrr%r)Zary1Zary2Zary1sqZary2sqrrrrRs   rR)rN)rN)rT)-r}Z __future__rrrrdrh functoolsr>rrZctypesrZnumpyr"rr/r r r r Znumba.unsafe.ndarrayr rrrrr$r objectr&rrrlrur`rrrrPrrOrrRrrrrsH      %&,