fy2dZddlZddlZddlZddlZddlmZddlZddl Z ddl m Z ddl m Z ddl m ZddlmZmZddlmZdd lmZdd lmZdd lmZdd lmZdd lmZ eeddZn#e $rdZYnwxYwdZ!dZ"dZ#Gdde j$Z%Gdde%Z&edZ'Gdde%Z(Gdde)Z*Gdde%ej+Z,Gdde%ej+Z-d*d Z.d*d!Z/d"Z0d#Z1d$Z2d%Z3d+d(Z4d)Z5dS),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)driver)typesconfig)to_fixed_tuple) dummyarray) numpy_support)prepare_shape_strides_dtype)NumbaPerformanceWarning)warn lru_cachec|SN)funcs >lib/python3.11/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.r Array from_descitemsize_dummyr)r$r% functoolsreduceoperatormulr'_drivermemory_size_from_info alloc_sizer get_contextmemallocdevice_memory_sizeUSE_NV_BINDINGbinding CUdeviceptrr MemoryPointergpu_datastream)selfr$r%r&rJrInulls 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|SrrJ)rKrJs 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&rr7r6)rK 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)rIr?rErFrGrrSr`s rrSz'DeviceNDArrayBase.device_ctypes_pointersD = % #221555{{"=6 6rc|jdkrdSt|||}t|t|}}t j|r>t|t ||t j|||j|dStj ||j drdndd|j d }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. rNrkrorprrT WRITEABLE)ordersubokrZ) r'sentry_contiguousrl array_corer?is_device_memorycheck_array_compatibilitydevice_to_devicerAr*arrayrthost_to_device)rKaryrJ self_coreary_cores rcopy_to_devicez DeviceNDArrayBase.copy_to_devices 8q== F$%%f--(.. 38  #C ( ( 2 c " " " %i : : :  $T3 O O O O O Ox&_^<Ecc#! 44 666H &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&rk)r$r&buffer)r$r&r%r)anyr%reformatrArlr*emptybyterr?device_to_hostr'ndarrayr$r&)rKrrJrihostarys 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. r2zonly support 1d arrayrzonly support unit strider&rJrIN)rlr6r.r%r&r9r+mathceilfloatr'rdminrIview DeviceNDArray) rKsectionrJnsectr%r9ibeginendr$rIs 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. )rIr`s r as_cuda_argzDeviceNDArrayBase.as_cuda_arg6s }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) rrBget_ipc_handlerIdictr$r%r&IpcArrayHandle)rKipchdescs 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&rJrI)r:squeezerr$r%r&rlrI)rKrrJ new_dummy_s rrzDeviceNDArrayBase.squeezeFs]&{***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%r9is_c_contiguousr.divmodrrJrI)rKr&r$r%rems rrzDeviceNDArrayBase.viewbsTZ  t|$$ : %. 0 0''))  6 $b DJ//NE"Is axx 6  .GBK;]     rc*|jj|jzSr)r&r9r'r`s rnbyteszDeviceNDArrayBase.nbytess z"TY..rrrrNr)__name__ __module__ __qualname____doc____cuda_memory__rrMpropertyrXr\rar_rlrxrSrrequire_contextrrrrrrrrrrrr0r0>sO1111f  X *  X  # # # #55555X5< 7 7X 7 2222< ,,,,\3333* @ @ @    8# # # J//X///rr0ceZdZdZd fd ZedZedZej dZ ej ddZ dd Z ej d Z ej dd Zdd ZxZS) DeviceRecordz An on-GPU record type rNchd}d}tt||||||dSNr)superrrM)rKr&rJrIr$r% __class__s rrMzDeviceRecord.__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:rtr`s rrtzDeviceRecord.flagsDK%&&&rc4tj|jS)rn)r rur&r`s rrxzDeviceRecord._numba_type_s ' 333rc,||Sr _do_getitemrKitems r __getitem__zDeviceRecord.__getitem__%%%rc.|||Sz0Do `__getitem__(item)` with CUDA stream rrKrrJs rgetitemzDeviceRecord.getitemf---rc||}|jj|\}}|j|}|jdkrT|jt|||Stj d|}tj |||j ||dSt|jd|jdd\}}} t||| ||S) Nrrr2r&dstsrcr'rJrrpr$r%r&rIrJ)rlr&fieldsrIrr$namesrr*rr?rr9r subdtyper) rKrrJr"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_setitemrKkeyrTs r __setitem__zDeviceRecord.__setitem__U+++rc2||||Sz6Do `__setitem__(key, value)` with CUDA stream rkrrKrrTrJs rsetitemzDeviceRecord.setitemU6:::rc||}| }|r'tj}|}|jj|\}}|j|}t||||} t| j||\} } tj | | | jj ||r| dSdS)Nrrk)rlrrBget_default_streamr&rrIrtype auto_devicer?rr9 synchronize) rKrrTrJ synchronousctxr"rrlhsrhsrs rrzDeviceRecord._do_setitems%%f-- !j  .%''C++--Fj', V-$$V,,d4jjs6GDDDSY^^E226BBBQ  c39+=vFFF  !     ! !rrr)rrrrrMrrtrxrrrrrrrr __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)Nr2rcrrr) gridrdr6r$localrrint64r )rrlocation n_elementsridxrr6s rrz_assign_kernel..kernels 99Q<< sx ' 'A #)A, &JJ z ! ! Fjd)+taxR(( & &A 39Q</C1I!CIaL0SYq\A5EFC1I 1 %HH,/s1vt0L0L,MN3q64 ( ()))r)numbarjit)r6rrs` @r_assign_kernelr sx 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_contigr`s ris_f_contiguouszDeviceNDArray.is_f_contiguous&{&&rc4t|jjSrrr`s rrtzDeviceNDArray.flags,rrc|jjS)z; Return true if the array is C-contiguous. )r: is_c_contigr`s rrzDeviceNDArray.is_c_contiguous6rrNc|r'||S|S)zE :return: an `numpy.ndarray`, so copies to the host. )r __array__)rKr&s rrzDeviceNDArray.__array__<sJ  3$$&&0077 7$$&&0022 2rc|jdSr)r$r`s r__len__zDeviceNDArray.__len__Esz!}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') r2r)r$r%r&rIoperation requires copying) r5r r)rrr$r%r&rIr:reshapeextentre)rKnewshapekwsclsnewarrextentss rrzDeviceNDArray.reshapeHs x==A  *Xa[5$-"H"H {H4jj tz ! !3TZ!Z$-AAA A.$+-x?3?? t{)* * *3V\6>!Z$-AAA A&&BCC Crrprc||}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) rlrr:ravelrr$r%r&rIre)rKr|rJr r!r"s rr$zDeviceNDArray.ravel`s %%f--4jj+++%+88 t{)* * *3V\6>!Z$-$&&& & &&BCC Crc,||Srrrs rrzDeviceNDArray.__getitem__rrrc.|||Srrrs rrzDeviceNDArray.getitemvrrc||}|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)Nr2rrrrr)rlr:rriter_contiguous_extentrr5rIris_arrayr&rrr*rr?rr9r$r%r)rKrrJarrr"r rrs rrzDeviceNDArray._do_getitem|sh%%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)rJ user_explicitz Can't assign z -D array to z-D selfrr2zCCan't copy sequence with size %d to array axis %d with dimension %drk) rlrrBrr:rrIrrr r Elementr$r%rr&rr6r.r*onesrr enumeratezipr;r<r=r>r forallr)rKrrTrJrrr*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)rprr)rrrrrrrtrrrrr$rrrrrrrrrrrrr"sV''' ''X'''' 3333DDD0DDDD$ &&& .... 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)rKrrs rrMzIpcArrayHandle.__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. rIr)r:openrrBrr9)rKdptrs rr<zIpcArrayHandle.opens? $$W%8%:%:;;??d?d.>???rc8|jdS)z5 Closes the IPC handle to the array. N)r:closer`s rr?zIpcArrayHandle.closes      rc*|Sr)r<r`s r __enter__zIpcArrayHandle.__enter__syy{{rc.|dSr)r?)rKrrT tracebacks r__exit__zIpcArrayHandle.__exit__s rN) rrrrrMr<r?rArDrrrrrsl$&&&@@@!!! rrceZdZdZddZdS) MappedNDArrayz4 A host array that uses CUDA mapped memory. rc"||_||_dSrrIrJrKrIrJs r device_setupzMappedNDArray.device_setup    rNrrrrrrJrrrrFrF2rrFceZdZdZddZdS)ManagedNDArrayz5 A host array that uses CUDA managed memory. rc"||_||_dSrrHrIs rrJzManagedNDArray.device_setuprKrNrrLrrrrOrOrMrrOcHt|j|j|j||S)z/Create a DeviceNDArray object that is like ary.rJrI)rr$r%r&)rrJrIs rfrom_array_likerSs* CK6"* , , ,,rc0t|j||S)z.Create a DeviceRecord object that is like rec.rR)rr&)recrJrIs rfrom_record_likerV!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. r2rFT)r&r9r2reversedr$r%)rr'r$r[s rrUrU8so 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)Nrorq)rrtr.errmsg_contiguous_buffer)rcores rr~r~NsM c??D :n %3dj.H312223333rTFcJtj|r|dfSt|dr!tj|dfSt |tjrt||}n7tj |dd}t|t||}|rrtjrO|sMt |ts8t |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. FrXrkT)rZr}zGHost array used in CUDA kernel will incur copy overhead to/from device.)r?rrr r as_cuda_arrayr r*voidrVrr~rSrCUDA_WARN_ON_IMPLICIT_COPYrrrr r)rrJrZr.devobjris rrrTs> $$Ez 0 1 1z'',,e33 c27 # # 9%c&999FF(C c " " "$S888F  60 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 rrr{s\\^^T\\^^FF zTZiTZZ122 2 |v|##j***djj233 3 y7V^v~55j,,, 677 77755rr)rTF)6rrr;r=rZctypesrnumpyr*r rnumba.cuda.cudadrvrrr? numba.corerrnumba.np.unsafe.ndarrayr numba.miscr numba.npr numba.cuda.api_utilr numba.core.errorsr warningsrrrrrr,r( DeviceArrayr0rr robjectrrrFrOrSrVrrUr_r~rrrrrrxsn    &&&&&&000000$$$$$$$$222222!!!!!!"""""";;;;;;555555/ ;//55II 333    ;;; N/N/N/N/N/ 0N/N/N/b d!d!d!d!d!$d!d!d!N (( (Vv!v!v!v!v!%v!v!v!r)))))V)))X%rz& ,,,, EEEE """$    3 333 $$$$N 7 7 7 7 7s A66BB