J/PhDLdZddlZddlZddlZddlmZmZmZddl m Z ddl m Z ej Z ejZejZe d)dZd*d Zd Zd Ze d+d Ze ejdd dfdZe ejdd ddfdZe ejdd fdZe ejdd dddfdZeje d,dZdZdZdZd-dZ d.dZ!dZ"e dZ#e dZ$e dZ%e dZ&e dZ'e ejdZ(e ejd Z)d*d!Z*ej+Z+d"Z,d#Z-d$Z.d%Z/d/d&Z0d'Z1ejd(Z2e ej3Z3e ej4Z4e ej5Z5dS)0z% API that are reported to numba.cuda N) devicearraydevicesdriver)config)prepare_shape_strides_dtypeTc|d}d|kr&|d}|td|d}|d}tj|d}t |||d \}}}t j|||j}t j|d d } t j t| || } |dd} | 2t| } |r tj r| nd } tj|||| | } | S)aCreate a DeviceNDArray from a cuda-array-interface description. The ``owner`` is the owner of the underlying memory. The resulting DeviceNDArray will acquire a reference from it. If ``sync`` is ``True``, then the imported stream (if present) will be synchronized. versionrmaskNzMasked arrays are not supportedshapestridestypestrC)orderdatar)sizeownerstream)r r dtypegpu_datar)getNotImplementedErrornprrrmemory_size_from_infoitemsizeget_devptr_for_active_ctx MemoryPointercurrent_contextexternal_streamrCUDA_ARRAY_INTERFACE_SYNC synchronizer DeviceNDArray)descrsyncr r r r rrdevptrr stream_ptrrdas N/var/www/html/test/jupyter/venv/lib/python3.11/site-packages/numba/cuda/api.pyfrom_cuda_array_interfacer)schhy!!GG||xx  %&GHH H MEhhy!!G HT)_ % %E7 wS***E7E  'w G GD  -d6l1o > >F  6E ; ; ;D(D))J ,,  !F4 !      ").*0 2 2 2B Iclt|stdt|j||S)aqCreate a DeviceNDArray from any object that implements the :ref:`cuda array interface `. A view of the underlying GPU buffer is created. No copying of the data is done. The resulting DeviceNDArray will acquire a reference from `obj`. If ``sync`` is ``True``, then the imported stream (if present) will be synchronized. z1*obj* doesn't implement the cuda array interface.)rr$) is_cuda_array TypeErrorr)__cuda_array_interface__)objr$s r( as_cuda_arrayr0?sD   ?KLLL()E/2??? ?r*c"t|dS)zTest if the object has defined the `__cuda_array_interface__` attribute. Does not verify the validity of the interface. r.)hasattr)r/s r(r,r,Ps 32 3 33r*cdS)zwWhether 16-bit floats are supported. float16 is always supported in current versions of Numba - returns True. Tr4r*r(is_float16_supportedr5Xs 4r*cv|tj|||d\}}|S|r||||S)ato_device(obj, stream=0, copy=True, to=None) Allocate and transfer a numpy ndarray or structured scalar to the device. To copy host->device a numpy array:: ary = np.arange(10) d_ary = cuda.to_device(ary) To enqueue the transfer to a stream:: stream = cuda.stream() d_ary = cuda.to_device(ary, stream=stream) The resulting ``d_ary`` is a ``DeviceNDArray``. To copy device->host:: hary = d_ary.copy_to_host() To copy device->host to an existing array:: ary = np.empty(shape=d_ary.shape, dtype=d_ary.dtype) d_ary.copy_to_host(ary) To enqueue the transfer to a stream:: hary = d_ary.copy_to_host(stream=stream) NT)rcopy user_explicitr)r auto_devicecopy_to_device)r/rr7tonews r( to_devicer>`sX> z)#f48<>>>C  . #f--- Ir*rc^t||||\}}}tj||||S)zdevice_array(shape, dtype=np.float64, strides=None, order='C', stream=0) Allocate an empty device ndarray. Similar to :meth:`numpy.empty`. )r r rr)rrr"r rr rrs r( device_arrayrAsF 8w8=??E7E  $5',2 4 4 44r*cdt||||\}}}tj|||j}t ||}t j|||||}t j|tj } | ||| S)amanaged_array(shape, dtype=np.float64, strides=None, order='C', stream=0, attach_global=True) Allocate a np.ndarray with a buffer that is managed. Similar to np.empty(). Managed memory is supported on Linux / x86 and PowerPC, and is considered experimental on Windows and Linux / AArch64. :param attach_global: A flag indicating whether to attach globally. Global attachment implies that the memory is accessible from any stream on any device. If ``False``, attachment is *host*, and memory is only accessible by devices with Compute Capability 6.0 and later. ) attach_globalr r rrbuffertyper9) rrrrrmemallocmanagedrndarrayviewrManagedNDArray device_setup) r rr rrrCbytesizerEnpary managedviews r( managed_arrayrPs$8w8=??E7E+E7ENKKH    . .x=J/LLF JUG5$ & & &E*//%k.H/IIKVF333 r*ct||||\}}}tj|||j}t |}t j|||||S)zpinned_array(shape, dtype=np.float64, strides=None, order='C') Allocate an :class:`ndarray ` with a buffer that is pinned (pagelocked). Similar to :func:`np.empty() `. rD)rrrrr memhostallocrrI)r rr rrMrEs r( pinned_arrayrSs{8w8=??E7E+E7,1N<>EN *D A5577 !'  7@  }d6JJJI   00'.e  = ==== OOr*cBtS)z Synchronize the current context.)rr!r4r*r(r!r!s    ( ( * **r*cZ|jds|jds |jdkrdSdt|jD}|ddgt |jz}|jj}|D]\}}|||<||j|z}t|S) z[ Given an array, compute strides for a new contiguous array of the same shape. C_CONTIGUOUS F_CONTIGUOUSrNcg|]}|Sr4r4).0xs r( z2_contiguous_strides_like_array..s6661666r*c|dS)Nrr4)rrs r(z0_contiguous_strides_like_array..s !A$r*)keyr) flagsndim enumerater sortlenrrr tuple)ary stridepermr stridei_perm_s r(_contiguous_strides_like_arrayrs y CIn$=Qt76i 44666JOOO'''cC $$$G Y F$$  #)F## >>r*c>|jdr|jdsdSdS)NrornFr)rw)r}s r(_order_like_arrayrs) y >)Bssr*czt|}t|}t|j|j|||S)zd Call :func:`device_array() ` with information from the array. r@)rrrAr r)r}rr rs r(device_array_liker#sD -S11G c " "E cisy'#F 4 4 44r*c ~t|}t|}t|j|j|||||S)zh Call :func:`mapped_array() ` with the information from the array. )r rr rrrXrY)rrr[r r)r}rrXrYr rs r(mapped_array_liker.sL -S11G c " "E cisy'#FX" N N NNr*cxt|}t|}t|j|j||S)zh Call :func:`pinned_array() ` with the information from the array. )r rr r)rrrSr r)r}r rs r(pinned_array_liker9sB -S11G c " "E cisy'# % % %%r*cBtS)zN Create a CUDA stream that represents a command queue for the device. )r create_streamr4r*r(rrEs    * * , ,,r*cBtS)a~ Get the default CUDA stream. CUDA semantics in general are that the default stream is either the legacy default stream or the per-thread default stream depending on which CUDA APIs are in use. In Numba, the APIs for the legacy default stream are always the ones in use, but an option to use APIs for the per-thread default stream may be provided in future. )rget_default_streamr4r*r(default_streamrMs    / / 1 11r*cBtS)z- Get the legacy default CUDA stream. )rget_legacy_default_streamr4r*r(legacy_default_streamrYs    6 6 8 88r*cBtS)z1 Get the per-thread default CUDA stream. )rget_per_thread_default_streamr4r*r(per_thread_default_streamras    : : < <s+C+C&,&=c&B&B-2&44  b EEEEEr*c/0K|rd|vs Jd|dd}g}g}|D]}t|tj|tj|d}||tj|||}|| t|dkr |dVn|V|D]}| dS#|D]}| wxYw) zKA context manager for temporarily mapping a sequence of host ndarrays. rz Only accept 'stream' as keyword.rTrU)rrrN) rrrrrrrrfrom_array_liker{free)rkwsrr devarylistr}rdevarys r(rVrVsT Ih#ooo'Iooo WWXq ! !F FJ""    % %c6+>s+C+C&,&=c&B&B-1&33  b,S2fMMM&!!!!  z??a  Q-          B GGIIII  &  B GGIIII s ="C::DcJt|}|S)zs Create a CUDA event. Timing data is only recorded by the event if it is created with ``timing=True``. )timing)r create_event)revts r(eventrs%    ( ( ( 7 7C Jr*c8tj|}|jS)z Make the context associated with device *device_id* the current context. Returns a Device instance. Raises exception on error. )r get_contextdevice) device_idcontexts r( select_devicers!),,G >r*c(tjS)z5Get current device associated with the current thread)rrr4r*r(get_current_devicers    ##r*ctjS)z%Return a list of all detected devices)rgpusr4r*r( list_devicesrs <r*c,tjdS)z Explicitly clears all contexts in the current thread, and destroys all contexts if the current thread is the main thread. N)rresetr4r*r(rgrgs  MOOOOOr*c0tj|||S)N)rr7)rr:)r}rr7s r( _auto_devicers  "3vD A A AAr*ct}tdt|zd}|D]}g}|j}|j}|j}|j}|dd|zfgz }|d|jfgz }|d|jfgz }|d|j fgz }|d|rd nd fgz }tj d kr |d |rd ndfgz }|d|fgz }|dkrd}n|dkrd}|dz }nd}|dz }td|j |j |fz|D]\} } t| dd| tdtd|t|fz|dkS)z Detect supported CUDA hardware and print a summary of the detected hardware. Returns a boolean indicating whether any supported devices were detected. zFound %d CUDA devicesrzCompute Capabilityz%d.%dz PCI Device IDz PCI Bus IDUUIDWatchdogEnabledDisabledntz Compute ModeTCCWDDMzFP32/FP64 Performance Ratio)z[NOT SUPPORTED: CC < 3.5])rrz[SUPPORTED (DEPRECATED)]rz [SUPPORTED]zid %d %20s %40sz>40z: zSummary:z %d/%d devices are supported) rprintr{compute_capabilityKERNEL_EXEC_TIMEOUT TCC_DRIVER%SINGLE_TO_DOUBLE_PRECISION_PERF_RATIO PCI_DEVICE_ID PCI_BUS_IDuuidosnameid) devlistsupported_countdevattrscckernel_timeouttccfp32_to_fp64_ratiosupportrvvals r(detectrs nnG !CLL 0111O++  #0n F '2677 ?C$5677 <011 638$%% :NJyy KLL 7d?? ~'?uu@A AE 02DEFF ;;1GG &[[0G q OO#G q O "cfch%@@AAA + +HC SS) * * * * + * )_c'll,K KLLL Q r*c#Kt5dVddddS#1swxYwYdS)aV Temporarily disable memory deallocation. Use this to prevent resource deallocation breaking asynchronous execution. For example:: with defer_cleanup(): # all cleanup is deferred in here do_speed_critical_code() # cleanup can occur here Note: this context manager can be nested. N)r defer_cleanupr4r*r(rrs    ( ( * * s 599)NT)T)rTN)Nr)r)rFF)rT)6__doc__ contextlibrnumpyrcudadrvrrr numba.corernumba.cuda.api_utilrrequire_contextrrrr)r0r,r5r>float64rArPrSr[contextmanagerrkr!rrrrrrrrrrrrVrevent_elapsed_timerrrrgrrr profiling profile_start profile_stopr4r*r(rs  1111111111;;;;;;)%|%%%%P????"444$$$$N j$c!4444!z4s1 $: j$c % % % % j$c!E2 8+++ :4444NNNN%%%---222999===999       6.   $$$  BBBB&&&R $ OF, - -  455 v233 r*