ó \ K]c@`südZddlmZmZmZddlZddlZddlZddlZddl m Z ddl m Z ddl Zddl Z ddlmZddlmZdd l mZmZmZdd lmZyeed ƒeƒZWnek r d „ZnXd „Zd„Zd„Zdefd„ƒYZ de fd„ƒYZ!ed„ƒZ"de fd„ƒYZ#defd„ƒYZ$de ej%fd„ƒYZ&ded„Z'ded„Z(d„Z)dZ*d„Z+de,d „Z-d!„Z.dS("sÑ 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. i(tprint_functiontabsolute_importtdivisionN(tsix(tc_void_pi(tdriver(tdevices(t dummyarrayttypest numpy_support(tto_fixed_tuplet lru_cachecC`s|S(N((tfunc((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyR scC`st|dtƒS(s$Check if an object is a CUDA ndarrayt__cuda_ndarray__(tgetattrtFalse(tobj((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pytis_cuda_ndarraysc`sWtˆƒ‡fd†}|dtƒ|dtƒ|dtjƒ|dtjƒdS(s,Verify the CUDA ndarray interface for an objc`sStˆ|ƒst|ƒ‚nttˆ|ƒ|ƒsOtd||fƒ‚ndS(Ns%s must be of type %s(thasattrtAttributeErrort isinstanceR(tattrttyp(R(s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyt requires_attr(stshapetstridestdtypetsizeN(trequire_cuda_ndarrayttupletnpRRt integer_types(RR((Rs=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pytverify_cuda_ndarray_interface$s    cC`st|ƒstdƒ‚ndS(s9Raises ValueError is is_cuda_ndarray(obj) evaluates Falsesrequire an cuda ndarray objectN(Rt ValueError(R((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyR4s tDeviceNDArrayBasecB`seZdZeZeZdddd„Zed„ƒZ dd„Z ed„ƒZ dd„Z d„Z ed„ƒZed „ƒZejdd „ƒZejddd „ƒZdd „Zdd „Zd„Zd„Zddd„Zd„Zed„ƒZRS(s$A on GPU NDArray representation icC`sœt|tjƒr|f}nt|tjƒr<|f}nt|ƒ|_t|ƒ|jkrotdƒ‚ntjjd|||j ƒ|_ t |ƒ|_ t |ƒ|_ tj|ƒ|_ttj|j ƒƒ|_|jdkrJ|dkr5tj|j |j |jj ƒ|_tjƒj|jƒ}q}tj|ƒ|_n3tjdtjƒdtdƒddƒ}d|_||_||_||_dS(sN 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 sstrides not match ndimitcontexttpointerRN( RRRtlentndimR!RtArrayt from_desctitemsizet_dummyRRRRRtinttprodRtNonet_drivertmemory_size_from_infot alloc_sizeRt get_contexttmemalloctdevice_memory_sizet MemoryPointerRtgpu_datat_DeviceNDArrayBase__writebacktstream(tselfRRRR7t writebackR5((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyt__init__@s4       cC`sKit|jƒd6t|jƒd6|jjtfd6|jjd6dd6S(NRRtdatattypestritversion(RRRtdevice_ctypes_pointertvalueRRtstr(R8((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyt__cuda_array_interface__ts  cC`stj|ƒ}||_|S(sBind a CUDA stream to this object so that all subsequent operation on this array defaults to the given stream. (tcopyR7(R8R7tclone((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pytbind~s cC`s |jƒS(N(t transpose(R8((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pytT†scC`sª|r+t|ƒtt|jƒƒkr+|S|jdkrItdƒ‚n]|dk rŒt|ƒtt|jƒƒkrŒtd|fƒ‚nddlm}||ƒSdS(Nis2transposing a non-2D DeviceNDArray isn't supportedsinvalid axes list %ri(RE( RtrangeR&tNotImplementedErrorR-tsetR!tnumba.cuda.kernels.transposeRE(R8taxesRE((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyREŠs'-cC`s|s |jS|S(N(R7(R8R7((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyt_default_stream•scC`s(tj|jƒ}tj||jdƒS(sn Magic attribute expected by Numba to get the numba type that represents this object. tA(R t from_dtypeRRR'R&(R8R((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyt _numba_type_˜scC`s'|jdkrtdƒS|jjSdS(s:Returns the ctypes pointer to the GPU data buffer iN(R5R-RR>(R8((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyR>¡s c C`só|jdkrdSt|ƒ|j|ƒ}t|ƒt|ƒ}}tj|ƒrŠt|ƒt||ƒtj|||jd|ƒnet j |d|j dr©dnddt d |j d  ƒ}t||ƒtj |||jd|ƒdS( sŸCopy `ary` to `self`. If `ary` is a CUDA memory, perform a device-to-device transfer. Otherwise, perform a a host-to-device transfer. iNR7tordert C_CONTIGUOUStCtFtsubokRBt WRITEABLE(Rtsentry_contiguousRLt array_coreR.tis_device_memorytcheck_array_compatibilitytdevice_to_deviceR0RtarraytflagstTruethost_to_device(R8taryR7t self_coretary_core((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pytcopy_to_deviceªs     c C`sHtd„|jDƒƒr:d}t|j|jƒƒ‚n|jdksUtdƒ‚|j|ƒ}|d kr‘tj d|jdtj ƒ}nt ||ƒ|}|jdkrÒt j |||jd|ƒn|d krD|jdkrtjd|jd|jd|ƒ}qDtjd|jd|jd |jd|ƒ}n|S( s^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() cs`s|]}|dkVqdS(iN((t.0ts((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pys ßss2D->H copy not implemented for negative strides: {}isNegative memory sizeRRR7tbufferRN(tanyRRHtformatR0tAssertionErrorRLR-RtemptytbyteRYR.tdevice_to_hostRtndarrayRR(R8R_R7tmsgthostary((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyt copy_to_hostÈs$ !   cC`sW|j|ƒ}tjdtƒ|jdkr=tdƒ‚n|j|jd|ƒdS(Ns+to_host() is deprecated and will be removedsno associated writeback arrayR7(RLtwarningstwarntDeprecationWarningR6R-R!Ro(R8R7((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pytto_hostös  c c`s|j|ƒ}|jdkr-tdƒ‚n|jd|jjkrUtdƒ‚nttjt |j ƒ|ƒƒ}|j}|jj}xt |ƒD]s}||}t |||j ƒ}||f} |j j||||ƒ} t| |d|jd|d| ƒVq™WdS( sžSplit the array into equal partition of the `section` size. If the array cannot be equally divided, the last section will be smaller. isonly support 1d arrayisonly support unit strideRR7R5N(RLR&R!RRR)R+tmathtceiltfloatRRGtminR5tviewt DeviceNDArray( R8tsectionR7tnsectRR)titbegintendRR5((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pytsplitþs"    cC`s|jS(sEReturns a device memory object that is used as the argument. (R5(R8((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyt as_cuda_argscC`sOtjƒj|jƒ}td|jd|jd|jƒ}td|d|ƒS(sÌ 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. RRRt ipc_handlet array_desc( RR1tget_ipc_handleR5tdictRRRtIpcArrayHandle(R8tipchtdesc((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyRƒs$c C`sU|jjd|ƒ\}}td|jd|jd|jd|j|ƒd|jƒS(s( 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. taxisRRRR7R5(R*tsqueezeRyRRRRLR5(R8RˆR7t new_dummyt_((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyR‰#s   c C`sdtj|ƒ}|j|jjkr3tdƒ‚ntd|jd|jd|d|jd|jƒS(seReturns a new object by reinterpretting the dtype without making a copy of the data. s new dtype itemsize doesn't matchRRRR7R5( RRR)t TypeErrorRyRRR7R5(R8R((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyRx?s   cC`s|jj|jS(N(RR)R(R8((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pytnbytesNsN(t__name__t __module__t__doc__R]t__cuda_memory__R R-R:tpropertyRARDRFRERLROR>Rtrequire_contextRbRoRsRR€RƒR‰RxR(((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyR":s. 3    -     t DeviceRecordcB`s;eZdZddd„Zed„ƒZed„ƒZRS(s An on-GPU record type icC`s2d}d}tt|ƒj|||||ƒdS(N(((tsuperR”R:(R8RR7R5RR((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyR:ZscC`st|jjƒS(sÿ 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). (R„R*R\(R8((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyR\`scC`stj|jƒS(sn Magic attribute expected by Numba to get the numba type that represents this object. (R RNR(R8((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyROjsN(RŽRRR-R:R’R\RO(((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyR”Vs c`s/ddlm‰ˆj‡‡fd†ƒ}|S(sÙ 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 i(tcudac`sˆjdƒ}d}x(t|jƒD]}||j|9}q%W||krPdSˆjjddˆfdtjƒ}xvtˆdddƒD]^}||j||d|f<||j||j|dk|d|f<||j|}q‹W|t|dˆƒ|t|dˆƒtjj|ƒtfSt|tj ƒret |d|ƒ}n7tj |dtdt ƒ}t |ƒt|d|ƒ}|r¸|j|d|ƒn|t fSdS(sº 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. RAR7RBRTN(R.RXRRR R–t as_cuda_arrayRRtvoidR×R[R]RVRÕRb(RR7RBtdevobj((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyRÃs   cC`s°|jƒ|jƒ}}|j|jkrJtd|j|jfƒ‚n|j|jkr{td|j|jfƒ‚n|j|jkr¬td|j|jfƒ‚ndS(Nsincompatible dtype: %s vs. %ssincompatible shape: %s vs. %ssincompatible strides: %s vs. %s(R‰RRŒRR!R(tary1tary2tary1sqtary2sq((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pyRY»s(/Rt __future__RRRRpRtt functoolsRBR RtctypesRtnumpyRtRR.RRRR tnumba.unsafe.ndarrayR RR-R RRR RtobjectR"R”R¢RyR…RlRÓRÕR×RWRÝRVR]RÃRY(((s=lib/python2.7/site-packages/numba/cuda/cudadrv/devicearray.pytsD           ÿ%¥,