B \.@sdZddlmZmZddlZddlZddlmZm Z m Z ddl m Z m Z mZyeWnek rleZYnXe jZe jZe jZed8ddZd d Zd d Zed9ddZeejdddfddZeejddfddZeejdddddfddZejed:ddZddZ ddZ!ddZ"d;d d!Z#ed"d#Z$eejd$d%Z%eejd&d'Z&dz% API that are reported to numba.cuda )print_functionabsolute_importN) devicearraydevicesdriver)InOutInOutc Cs|d}|d}t|d}t|||dd\}}}t|||j}t|dd}tjt |||d}t j ||||d }|S) zCreate 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. shapestridesZtypestrC)orderdatar)sizeowner)r r dtypegpu_data) getnpr_prepare_shape_strides_dtypermemory_size_from_infoitemsizeZget_devptr_for_active_ctxZ MemoryPointercurrent_contextr DeviceNDArray) Zdescrr r rrZdevptrrZdar-lib/python3.7/site-packages/numba/cuda/api.pyfrom_cuda_array_interfaces  rcCs$t|stdnt|j|dSdS)zCreate a DeviceNDArray from any object that implements the 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`. z1*obj* doesn't implement the cuda array interface.)rN) is_cuda_array TypeErrorr__cuda_array_interface__)objrrr as_cuda_array1s r"cCs t|dS)zwTest if the object has defined the `__cuda_array_interface__`. Does not verify the validity of the interface. r )hasattr)r!rrrr?srTcCs6|dkr tj|||d\}}|S|r2|j||d|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) N)streamcopy)r$)r auto_deviceZcopy_to_device)r!r$r%tonewrrr to_deviceGs r)r cCs&t||||\}}}tj||||dS)zdevice_array(shape, dtype=np.float, strides=None, order='C', stream=0) Allocate an empty device ndarray. Similar to :meth:`numpy.empty`. )r r rr$)rrr)r rr rr$rrr device_arrayns  r*cCsDt||||\}}}t|||j}t|}tj|||||dS)zpinned_array(shape, dtype=np.float, strides=None, order='C') Allocate a np.ndarray with a buffer that is pinned (pagelocked). Similar to np.empty(). )r r rrbuffer)rrrrr memhostallocrndarray)r rr rbytesizer+rrr pinned_arrayzs   r/Fc Cslt||||\}}}t|||j}tj|dd}tj|||||d} tjj| t j d} | j ||d| S)a&mapped_array(shape, dtype=np.float, strides=None, order='C', stream=0, portable=False, wc=False) Allocate a mapped ndarray with a buffer that is pinned and mapped on to the device. Similar to np.empty() :param portable: a boolean flag to allow the allocated device memory to be usable in multiple devices. :param wc: a boolean flag to enable writecombined allocation which is faster to write by the host and to read by the device, but slower to write by the host and slower to write by the device. T)mapped)r r rrr+)type)r$) rrrrrr,rr-ZviewrZ MappedNDArrayZ device_setup) r rr rr$ZportableZwcr.r+ZnparyZ mappedviewrrr mapped_arrays  r2ccsZt|}t||j}tjj|}tjd|||d}|jt |||dV| dS)a A context manager that opens a IPC *handle* (*CUipcMemHandle*) that is represented as a sequence of bytes (e.g. *bytes*, tuple of int) and represent it as an array of the given *shape*, *strides* and *dtype*. The *strides* can be omitted. In that case, it is assumed to be a 1D C contiguous array. Yields a device array. The IPC handle is closed automatically when context manager exits. N)offset)r r r) rrZprodrrZdrvapiZcu_ipc_mem_handleZ IpcHandleZ open_arrayrclose)Zhandler rr r3rZ ipchandlerrropen_ipc_arrays    r5cCs tS)z Synchronize the current context.)r synchronizerrrrr6sr6cCsZt|}t|ttfr|f}t|ttfr4|f}n|dkr@d}|pNt|||}|||fS)Nr)r)rr isinstanceintlong_fill_stride_by_order)r r rrrrrrs rcCst|}dg|}|dkr\|j|d<xtt|dD] }||d||d||<q6WnL|dkr|j|d<x8td|D] }||d||d||<qzWntdt|S)Nrr rFzmust be either C/F order)lenrreversedrange ValueErrortuple)r rrZndr drrrr:s  " "r:cCst|j|j|j|dS)z=Call cuda.devicearray() with information from the array. )r rr r$)r*r rr )aryr$rrrdevice_array_likes rDcCs tS)zWstream() Create a CUDA stream that represents a command queue for the device. )rZ create_streamrrrrr$sr$cgsFg}x6|D].}tj|t|t|dd}||q WdVdS)zIA context manager for temporary pinning a sequence of host ndarrays. F)r0N)rmempinr host_pointerhost_memory_sizeappend)arylistpmlistrCpmrrrpinneds rLc os|rd|kstd|dd}g}g}xP|D]H}tj|t|t|dd}||tj |||d}||q.Wz"t |dkr|dVn|VWdx|D] }| qWXdS) zKA context manager for temporarily mapping a sequence of host ndarrays. r$z Only accept 'stream' as keyword.rT)r0)rr$rN) AssertionErrorrrrErrFrGrHrZfrom_array_liker=Zfree)rIZkwsr$rJZ devarylistrCrKZdevaryrrrr0s"       r0cCstj|d}|S)zs Create a CUDA event. Timing data is only recorded by the event if it is created with ``timing=True``. )timing)rZ create_event)rNZevtrrreventsrOcCst|}|jS)z Make the context associated with device *device_id* the current context. Returns a Device instance. Raises exception on error. )r get_contextdevice)Z device_idcontextrrr select_device&s rScCstjS)z5Get current device associated with the current thread)rrQrrrrget_current_device2srTcCstjS)z%Return a list of all detected devices)rgpusrrrr list_devices7srVcCs tdS)z Explicitly clears all contexts in the current thread, and destroys all contexts if the current thread is the main thread. N)rresetrrrrr4<sr4cCstj|||dS)N)r$r%)rr&)rCr$r%rrr _auto_deviceDsrXcCst}tdt|d}x|D]}g}|j}|dd|fg7}|d|jfg7}|d|jfg7}|dkrnd}n d }|d 7}td |j|j|fx |D]\}}td ||fqWq Wtd td|t|f|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 id)rz[NOT SUPPORTED: CC < 2.0]z [SUPPORTED]rzid %d %20s %40sz%40s: %szSummary:z %d/%d devices are supported)rVprintr=Zcompute_capabilityZ PCI_DEVICE_IDZ PCI_BUS_IDidname)ZdevlistZsupported_countZdevZattrsZccZsupportkeyvalrrrdetectHs& r_c cs&tj}| dVWdQRXdS)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)rZ deallocationsdisable)Zdeallocsrrr defer_cleanupfs ra)N)rTN)Nr)r)T)rT)3__doc__Z __future__rr contextlibZnumpyrZcudadrvrrrargsrr r r9 NameErrorr8Zrequire_contextrPrrUrr"rr)floatr*r/r2contextmanagerr5r6rr:rDr$rLr0rOZevent_elapsed_timerSrTrVr4rXr_raZ profilingZ profile_startZ profile_stoprrrrsb   &