1@j}6dZddlZddlZddlZddlZddlmZddlZddl Z ddl m Z ddl m Z m Z ddl mZddlmZmZddlmZdd lmZdd lmZdd lmZdd lmZdd lmZ eeddZ n#e!$rdZ YnwxYwdZ"dZ#dZ$Gdde j%Z&Gdde&Z'e dZ(Gdde&Z)Gdde*Z+Gdde&ej,Z-Gdde&ej,Z.d*d Z/d*d!Z0d"Z1d#Z2d$Z3d%Z4d+d(Z5d)Z6dS),z 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. N)c_void_p) _devicearray)devices dummyarray)driver)typesconfig)to_fixed_tuple) numpy_version) numpy_support)prepare_shape_strides_dtype)NumbaPerformanceWarning)warn lru_cachec|SN)funcs OC:\Projects\Historical\venv\Lib\site-packages\numba/cuda/cudadrv/devicearray.pyrrs c$t|ddS)z$Check if an object is a CUDA ndarray__cuda_ndarray__F)getattrobjs ris_cuda_ndarrayr#s 3*E 2 22rctfd}|dt|dt|dtj|dtdS)z,Verify the CUDA ndarray interface for an objct|st|tt||st|d|dS)Nz must be of type )hasattrAttributeError isinstancer)attrtyprs r requires_attrz4verify_cuda_ndarray_interface..requires_attr,sesD!! ' && &'#t,,c22 H DDD##!FGG G H HrshapestridesdtypesizeN)require_cuda_ndarraytuplenpr'int)rr$s` rverify_cuda_ndarray_interfacer-(sHHHHH M'5!!!M)U###M'28$$$M&#rcBt|stddS)z9Raises ValueError is is_cuda_ndarray(obj) evaluates Falsezrequire an cuda ndarray objectN)r ValueErrorrs rr)r)8s+ 3  ;9:::;;rc eZdZdZdZdZddZedZddZ edZ dd Z d Z ed Z ed Zejdd ZejddZddZdZdZddZdZedZdS)DeviceNDArrayBasez$A on GPU NDArray representation TrNct|tr|f}t|tr|f}tj|}t ||_t ||jkrt dtj d|||j |_ t||_ t||_||_ttjt"j|j d|_|jdkrw|[t)j|j |j|jj |_t/j|j}nt)j||_njt(jr t(jd}nt=d}t)jt/j|d}d|_||_ ||_!dS)a5 Args ---- shape array shape. strides array strides. dtype data type as np.dtype coercible object. stream cuda stream. gpu_data user provided device memory for the ndarray data buffer zstrides not match ndimrN)contextpointerr()"r!r,r+r'lenndimr/rArray from_descitemsize_dummyr*r%r& functoolsreduceoperatormulr(_drivermemory_size_from_info alloc_sizer get_contextmemallocdevice_memory_sizeUSE_NV_BINDINGbinding CUdeviceptrr MemoryPointergpu_datastream)selfr%r&r'rKrJnulls r__init__zDeviceNDArrayBase.__init__Ds eS ! ! HE gs # # !jGJJ w<<49 $ $566 6 &00E716AA 5\\ W~~   (tz1EEFF 9q==")"?J dj.A#C#C".0099$/JJ")"> > > > > > > >9T?? "rc|s|jn|SrrK)rLrKs r_default_streamz!DeviceNDArrayBase._default_streams"(4t{{f4rcd|jv}|jdr|sd}n|jdr|sd}nd}tj|j}t j||j|S)n Magic attribute expected by Numba to get the numba type that represents this object. r C_CONTIGUOUSC F_CONTIGUOUSFA)r&flagsr from_dtyper'rr8r7)rL broadcastlayoutr's r _numba_type_zDeviceNDArrayBase._numba_type_sw(% :n % i FF Z '   FFF(44{5$)V444rc|j:tjrtjdSt dS|jjS)z:Returns the ctypes pointer to the GPU data buffer Nr)rJr@rFrGrHrrTras rrTz'DeviceNDArrayBase.device_ctypes_pointersD = % #221555{{"=6 6rc"|jdkrdSt|||}t|t|}}t j|r>t|t ||t j|||j|dStj ||j drdnddtdkr|j d  nd }t ||t j |||j|dS) zCopy `ary` to `self`. If `ary` is a CUDA memory, perform a device-to-device transfer. Otherwise, perform a a host-to-device transfer. rNrlrprqrsTrdr WRITEABLE)ordersubokr[)r(sentry_contiguousrm array_corer@is_device_memorycheck_array_compatibilitydevice_to_devicerBr+arrayrur host_to_device)rLaryrK self_coreary_cores rcopy_to_devicez DeviceNDArrayBase.copy_to_devices1 8q== F$%%f--(.. 38  #C ( ( 2 c " " " %i : : :  $T3 O O O O O Ox&_^<Ecc# 6))#.555/3 555H &i : : :  "44?*0 2 2 2 2 2 2rcXtd|jDr)d}t||j|jdks Jd||}|&t j|jt j}nt|||}|jdkrtj |||j||T|j dkr"t j |j|j|}n't j |j|j|j| }|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() c3"K|] }|dkV dSrNr).0ss r z1DeviceNDArrayBase.copy_to_host.. s&++q1u++++++rz2D->H copy not implemented for negative strides: {}rzNegative memory sizeNr%r'rl)r%r'buffer)r%r'r&r)anyr&rfformatrBrmr+emptybyterr@device_to_hostr(ndarrayr%r')rLrrKrjhostarys r copy_to_hostzDeviceNDArrayBase.copy_to_hosts@. ++dl+++ + + @FC%cjj&>&>?? ?!###%;###%%f-- ;hT_BGDDDGG %dC 0 0 0G ?a    "7D$/*0 2 2 2 2 ;yA~~*4:TZ,3555*4:TZ-1\'KKKrc#0K||}|jdkrtd|jd|jjkrtdt tjt|j |z }|j}|jj}t|D]a}||z}t||z|j }||z f} |j ||z||z} t| ||j|| VbdS)zSplit the array into equal partition of the `section` size. If the array cannot be equally divided, the last section will be smaller. r3zonly support 1d arrayrzonly support unit strider'rKrJN)rmr7r/r&r'r:r,mathceilfloatr(reminrJview DeviceNDArray) rLsectionrKnsectr&r:ibeginendr%rJs rsplitzDeviceNDArrayBase.split"s, %%f-- 9>>455 5 <?dj1 1 1788 8DIeDI..899::,:&u 3 3AKEegoty11C5[NE}))%(*:C(NKKHwdj)1333 3 3 3 3  3 3rc|jS)zEReturns a device memory object that is used as the argument. )rJras r as_cuda_argzDeviceNDArrayBase.as_cuda_arg7s }rctj|j}t |j|j|j}t||S)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. )r%r&r') ipc_handle array_desc) rrCget_ipc_handlerJdictr%r&r'IpcArrayHandle)rLipchdescs rrz DeviceNDArrayBase.get_ipc_handle<sO"$$33DMBB$*dl$*MMM$????rc|j|\}}t|j|j|j|||jS)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. )axisr%r&r'rKrJ)r;squeezerr%r&r'rmrJ)rLrrK new_dummy_s rrzDeviceNDArrayBase.squeezeGs]&{***55 1/%*''//]     rctj|}t|j}t|j}|jj|jkrp|stdt|d|jjz|j\|d<}|dkrtd|j|d<t||||j |j S)zeReturns a new object by reinterpretting the dtype without making a copy of the data. zHTo change to a dtype of a different size, the array must be C-contiguousrzuWhen changing to a larger dtype, its size must be a divisor of the total size in bytes of the last axis of the array.r) r+r'listr%r&r:is_c_contiguousr/divmodrrKrJ)rLr'r%r&rems rrzDeviceNDArrayBase.viewcsTZ  t|$$ : %. 0 0''))  6 $b DJ//NE"Is axx 6  .GBK;]     rc*|jj|jzSr)r'r:r(ras rnbyteszDeviceNDArrayBase.nbytess z"TY..rrrrNr)__name__ __module__ __qualname____doc____cuda_memory__rrNpropertyrYr]rbr`rmryrTrrequire_contextrrrrrrrrrrrr1r1>sO1111f  X *  X  # # # #55555X5< 7 7X 7 2222> ,,,,\3333* @ @ @    8# # # J//X///rr1ceZdZdZd fd ZedZedZej dZ ej ddZ dd Z ej d Z ej dd Zdd ZxZS) DeviceRecordz An on-GPU record type rNchd}d}tt||||||dSNr)superrrN)rLr'rKrJr%r& __class__s rrNzDeviceRecord.__init__sD lD!!**5'5&+3 5 5 5 5 5rc4t|jjSz 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). rr;ruras rruzDeviceRecord.flagsDK%&&&rc4tj|jS)ro)r rvr'ras rryzDeviceRecord._numba_type_s ' 333rc,||Sr _do_getitemrLitems r __getitem__zDeviceRecord.__getitem__%%%rc.|||Sz0Do `__getitem__(item)` with CUDA stream rrLrrKs rgetitemzDeviceRecord.getitemf---rc||}|jj|\}}|j|}|jdkrT|jt|||Stj d|}tj |||j ||dSt|jd|jdd\}}} t||| ||S) Nrrr3r'dstsrcr(rKrrqr%r&r'rJrK)rmr'fieldsrJrr%namesrr+rr@rr:r subdtyper) rLrrKr#offsetnewdatarr%r&r's rrzDeviceRecord._do_getitems %%f--j'- V-$$V,, 9??y$##f-46666(1C000&7,/L.466661: ,CI,0,/LOSBB "E7E!ug',w(.000 0rc.|||Sr _do_setitemrLkeyrUs r __setitem__zDeviceRecord.__setitem__U+++rc2||||Sz6Do `__setitem__(key, value)` with CUDA stream rlrrLrrUrKs rsetitemzDeviceRecord.setitemU6:::rc||}| }|r'tj}|}|jj|\}}|j|}t||||} t| j||\} } tj | | | jj ||r| dSdS)Nrrl)rmrrCget_default_streamr'rrJrtype auto_devicer@rr: synchronize) rLrrUrK synchronousctxr#rrlhsrhsrs rrzDeviceRecord._do_setitems%%f-- !j  .%''C++--Fj', V-$$V,,d4jjs6GDDDSY^^E226BBBQ  c39+=vFFF  !     ! !rrr)rrrrrNrruryrrrrrrrr __classcell__)rs@rrrs+555555 ''X'44X4 &&& .... 00000 ,,, ;;;; !!!!!!!!rrclddlmdkrjd}|Sjfd}|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)cudac|d|d<dSrr)rrs rkernelz_assign_kernel..kernels"gCGGGrcd}d}t|jD]}||j|z}||krdSjdft j}tdz ddD]N}||j|z|d|f<||j|z|j|dkz|d|f<||j|z}O|t|d|t|d<dS)Nr3rdrrr) gridrer7r%localrrint64r )rrlocation n_elementsridxrr7s rrz_assign_kernel..kernels 99Q<< sx ' 'A #)A, &JJ z ! ! Fjd)+taxR(( & &A 39Q</C1I!CIaL0SYq\A5EFC1I 1 %HH,/s1vt0L0L,MN3q64 ( ()))r)numbarjit)r7rrs` @r_assign_kernelrsx qyy       XNNNNNXN. MrceZdZdZdZedZdZddZdZ dZ dd Z e j d Ze j dd ZddZe j dZe j ddZddZdS)rz An on-GPU array type c|jjS)zA Return true if the array is Fortran-contiguous. )r; is_f_contigras ris_f_contiguouszDeviceNDArray.is_f_contiguous'{&&rc4t|jjSrrras rruzDeviceNDArray.flags-rrc|jjS)z; Return true if the array is C-contiguous. )r; is_c_contigras rrzDeviceNDArray.is_c_contiguous7rrNc|r'||S|S)zE :return: an `numpy.ndarray`, so copies to the host. )r __array__)rLr's rrzDeviceNDArray.__array__=sJ  3$$&&0077 7$$&&0022 2rc|jdSr)r%ras r__len__zDeviceNDArray.__len__Fsz!}rct|dkr*t|dttfr|d}t |}||jkr#||j|j|j|jS|j j |i|\}}||j j gkr#||j|j|j|jStd)z Reshape the array without changing its contents, similarly to :meth:`numpy.ndarray.reshape`. Example:: d_arr = d_arr.reshape(20, 50, order='F') r3r)r%r&r'rJoperation requires copying) r6r!r*rrr%r&r'rJr;reshapeextentrf)rLnewshapekwsclsnewarrextentss rrzDeviceNDArray.reshapeIs x==A  *Xa[5$-"H"H {H4jj tz ! !3TZ!Z$-AAA A.$+-x?3?? t{)* * *3V\6>!Z$-AAA A&&BCC Crrqrc||}t|}|j|\}}||jjgkr$||j|j|j|j|Std)z Flattens a contiguous array without changing its contents, similar to :meth:`numpy.ndarray.ravel`. If the array is not contiguous, raises an exception. )r~rr) rmrr;ravelrr%r&r'rJrf)rLr~rKr"r#r$s rr&zDeviceNDArray.ravelcs %%f--4jj+++%+88 t{)* * *3V\6>!Z$-$&&& & &&BCC Crc,||Srrrs rrzDeviceNDArray.__getitem__urrc.|||Srrrs rrzDeviceNDArray.getitemyrrc||}|j|}t|}t |}t |dkr|jj|d}|j sh|j j t|j ||Stjd|j }tj|||jj||dS||j|j|j ||S|jj|j}||j|j|j ||S)Nr3rrrrr)rmr;rriter_contiguous_extentrr6rJris_arrayr'rrr+rr@rr:r%r&r)rLrrKarrr$r"rrs rrzDeviceNDArray._do_getitemsh%%f--k%%d++s1133444jj w<<1  (dm('!*5G< N:#/'dj18::::!hq ;;;G*wG04 0D28::::qz!sCK!%gfNNNN)dm(#*5G3SY !Z'&JJJ Jrc.|||Srrrs rrzDeviceNDArray.__setitem__rrc2||||Srrrs rrzDeviceNDArray.setitemrrc||}| }|r'tj}|}|j|}|jj|j}t|tj rd}d} n|j }|j } t||| |j||} t!||d\} } | j| jkr t%d| jd| jdt'j| jt&j} | j | | j| jz d<| j| } t/t1| j | j D])\}\}}|d kr||krt%d |||fz*t3jt6j| j d }t;| j|| | | |r|dSdS) NrrT)rK user_explicitz Can't assign z -D array to z-D selfrr3zCCan't copy sequence with size %d to array axis %d with dimension %drl) rmrrCrr;rrJrrr!rElementr%r&rr'rr7r/r+onesr r enumeratezipr<r=r>r?rforallr)rLrrUrKrrr,rr%r&rrr rhs_shaperlrr s rrzDeviceNDArray._do_setitems?%%f-- !j  .%''C++--Fk%%c**$$-$cj1 c:- . . "EGGIEkGd4jj* U6FFFQ 8ch  * GCHBH555 *-) #(SX%&&'ck9%"3sy#)#<#<== K KIAv1Avv!q&& "=ABAqz"JKKK %hlCIqAA Bsx  '' 6'BB3LLL  !     ! !rr)rqrr)rrrrrrrurrrrr&rrrrrrrrrrrrr#sV''' ''X'''' 3333DDD4DDDD$ &&& .... JJJJ: ,,, ;;;; 5!5!5!5!5!5!rrc0eZdZdZdZdZdZdZdZdS)ra" 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 c"||_||_dSr) _array_desc _ipc_handle)rLrrs rrNzIpcArrayHandle.__init__s%%rc~|jtj}t dd|i|jS)z Returns a new *DeviceNDArray* that shares the allocation from the original process. Must not be used on the original process. rJr)r<openrrCrr;)rLdptrs rr>zIpcArrayHandle.opens? $$W%8%:%:;;??d?d.>???rc8|jdS)z5 Closes the IPC handle to the array. N)r<closeras rrAzIpcArrayHandle.closes      rc*|Sr)r>ras r __enter__zIpcArrayHandle.__enter__syy{{rc.|dSr)rA)rLrrU tracebacks r__exit__zIpcArrayHandle.__exit__s rN) rrrrrNr>rArCrFrrrrrsl$&&&@@@!!! rrceZdZdZddZdS) MappedNDArrayz4 A host array that uses CUDA mapped memory. rc"||_||_dSrrJrKrLrJrKs r device_setupzMappedNDArray.device_setup   rNrrrrrrLrrrrHrH 2rrHceZdZdZddZdS)ManagedNDArrayz5 A host array that uses CUDA managed memory. rc"||_||_dSrrJrKs rrLzManagedNDArray.device_setuprMrNrrNrrrrQrQrOrrQcHt|j|j|j||S)z/Create a DeviceNDArray object that is like ary.rKrJ)rr%r&r')rrKrJs rfrom_array_likerUs* CK6"* , , ,,rc0t|j||S)z.Create a DeviceRecord object that is like rec.rT)rr')recrKrJs rfrom_record_likerX$s  &8 D D DDrc|jr|js|Sg}|jD],}||dkrdntd-|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)r&r(appendslicer*)r core_indexstrides rrr)sn ;ch J+==v{{!!d <<<< uZ  !!rc|jj}tt|jt|jD]\}}|dkr|dkr||krdS||z} dS)z Returns True iff `ary` is C-style contiguous while ignoring broadcasted and 1-sized dimensions. As opposed to array_core(), it does not call require_context(), which can be quite expensive. r3rFT)r'r:r4reversedr%r&)rr(r%r]s rrVrV;so 9 DXci00(3;2G2GHH v 1991v~~uu EMD 4rzArray contains non-contiguous buffer and cannot be transferred as a single memory region. Please ensure contiguous buffer with numpy .ascontiguousarray()ct|}|jds!|jdsttdSdS)Nrprr)rrur/errmsg_contiguous_buffer)rcores rrrQsM c??D :n %3dj.H312223333rTFc|tj|r|dfSt|dr!tj|dfSt |tjrt||}nDtj |tdkrdndd}t|t||}|r~tjr[tjsO|sMt |t"s8t |tjrd}t't)|||||dfS) 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. FrYrlr|NT)r[rzGHost array used in CUDA kernel will incur copy overhead to/from device.)r@rrr r as_cuda_arrayr!r+voidrXrr rrUr CUDA_WARN_ON_IMPLICIT_COPYDISABLE_PERFORMANCE_WARNINGSrrrrr)rrKr[r0devobjrjs rrrWsY $$!Ez 0 1 1z'',,e33 c27 # # 9%c&999FF(+f44UU$C c " " "$S888F  61 7; 7 &7#C777$C447 ;C055666  ! !#f ! 5 5 5t|rc||}}|j|jkrtd|jd|j|j|jkrt d|jd|j|jr/|j|jkr!t d|jd|jdSdS)Nzincompatible dtype: z vs. zincompatible shape: zincompatible strides: )rr' TypeErrorr%r/r(r&)ary1ary2ary1sqary2sqs rrrs\\^^T\\^^FF zTZiTZZ122 2 |v|##j***djj233 3 y7V^v~55j,,, 677 77755rr)rTF)7rrr<r>r[ctypesrnumpyr+r rnumba.cuda.cudadrvrrrr@ numba.corerr numba.np.unsafe.ndarrayr numba.np.numpy_supportr numba.npr numba.cuda.api_utilr numba.core.errorsrwarningsrrrr rr-r) DeviceArrayr1rrrobjectrrrHrQrUrXrrVrarrrrrrr{st    22222222000000$$$$$$$$222222000000"""""";;;;;;555555/ ;//55II 333    ;;; O/O/O/O/O/ 0O/O/O/d d!d!d!d!d!$d!d!d!N (( (Vx!x!x!x!x!%x!x!x!v)))))V)))X%rz& ,,,, EEEE """$    3 333 ''''T 7 7 7 7 7s"A88BB