ó \ K]c@@sZdZddlmZmZddlZddlZddlmZm Z m Z ddl m Z m Z mZyeWnek reZnXe jZe jZe jZedd„ƒZd„Zd „Zededd „ƒZeejdd dd „ƒZeejdd d „ƒZeejdd deed„ƒZ ej!eddd„ƒƒZ"d„Z#d„Z$d„Z%dd„Z&ed„ƒZ'eej!d„ƒƒZ(eej!d„ƒƒZ)ed„Z*e j+Z+d„Z,d„Z-d„Z.d„Z/ded„Z0d„Z1ej!d„ƒZ2ee j3ƒZ3ee j4ƒZ4ee j5ƒZ5dS(s% API that are reported to numba.cuda i(tprint_functiontabsolute_importNi(t devicearraytdevicestdriver(tIntOuttInOutc C@s |jdƒ}|dkrH|jdƒ}|dk rHtdƒ‚qHn|d}|jdƒ}tj|dƒ}t|||dd ƒ\}}}tj|||jƒ}tj |d d ƒ}tj t ƒ|d |d |ƒ} t j d|d|d|d| ƒ} | S(s½Create 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. tversionitmasksMasked arrays are not supportedtshapetstridesttypestrtordertCtdataitsizetownertdtypetgpu_dataN(tgettNonetNotImplementedErrortnpRt_prepare_shape_strides_dtypeRtmemory_size_from_infotitemsizetget_devptr_for_active_ctxt MemoryPointertcurrent_contextRt DeviceNDArray( tdescRRR R R RRtdevptrRtda((s-lib/python2.7/site-packages/numba/cuda/api.pytfrom_cuda_array_interfaces"   cC@s2t|ƒstdƒ‚nt|jd|ƒSdS(sõCreate 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`. s1*obj* doesn't implement the cuda array interface.RN(t is_cuda_arrayt TypeErrorR"t__cuda_array_interface__(tobj((s-lib/python2.7/site-packages/numba/cuda/api.pyt as_cuda_array8s  cC@s t|dƒS(swTest if the object has defined the `__cuda_array_interface__`. Does not verify the validity of the interface. R%(thasattr(R&((s-lib/python2.7/site-packages/numba/cuda/api.pyR#FscC@sQ|dkr1tj|d|d|ƒ\}}|S|rM|j|d|ƒn|S(sËto_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) tstreamtcopyN(RRt auto_devicetcopy_to_device(R&R)R*ttotnew((s-lib/python2.7/site-packages/numba/cuda/api.pyt to_deviceNs  !Rc C@s@t||||ƒ\}}}tjd|d|d|d|ƒS(s’device_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 R R)((s-lib/python2.7/site-packages/numba/cuda/api.pyt device_arrayus c C@spt||||ƒ\}}}tj|||jƒ}tƒj|ƒ}tjd|d|d|d|d|ƒS(s¢pinned_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 RR tbuffer(RRRRRt memhostallocRtndarray(R RR R tbytesizeR1((s-lib/python2.7/site-packages/numba/cuda/api.pyt pinned_arrays   !c C@sªt||||ƒ\}}}tj|||jƒ}tƒj|dtƒ}tjd|d|d|d|d|ƒ} tjj | dt j ƒ} | j |d|ƒ| S( s&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. tmappedR R RR R1ttypeR)( RRRRRR2tTrueRR3tviewRt MappedNDArrayt device_setup( R RR R R)tportabletwcR4R1tnparyt mappedview((s-lib/python2.7/site-packages/numba/cuda/api.pyt mapped_array‘s ! cc@sƒtj|ƒ}tj|ƒ|j}tjj|Œ}tjd||d|ƒ}|j t ƒd|d|d|ƒV|j ƒdS(s¨ 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. toffsetR R RN( RRtprodRRtdrvapitcu_ipc_mem_handlet IpcHandleRt open_arrayRtclose(thandleR RR RARt ipchandle((s-lib/python2.7/site-packages/numba/cuda/api.pytopen_ipc_arrayªscC@s tƒjƒS(s Synchronize the current context.(Rt synchronize(((s-lib/python2.7/site-packages/numba/cuda/api.pyRKÄscC@s‹tj|ƒ}t|ttfƒr0|f}nt|ttfƒrQ|f}n-|dkrfd}n|p{t|||ƒ}|||fS(Ni((i(RRt isinstancetinttlongt_fill_stride_by_order(R R RR ((s-lib/python2.7/site-packages/numba/cuda/api.pyRÉs    cC@sÞt|ƒ}dg|}|dkrt|j|dscC@stjƒdS(s‡ Explicitly clears all contexts in the current thread, and destroys all contexts if the current thread is the main thread. N(Rtreset(((s-lib/python2.7/site-packages/numba/cuda/api.pyRGCscC@stj|d|d|ƒS(NR)R*(RR+(RXR)R*((s-lib/python2.7/site-packages/numba/cuda/api.pyt _auto_deviceKscC@s!tƒ}tdt|ƒƒd}xÍ|D]Å}g}|j}|dd|fg7}|d|jfg7}|d|jfg7}|dkr—d}nd }|d 7}td |j|j|fƒx(|D] \}}td ||fƒqËWq*Wtd ƒtd|t|ƒfƒ|dkS(s¥ Detect supported CUDA hardware and print a summary of the detected hardware. Returns a boolean indicating whether any supported devices were detected. sFound %d CUDA devicesiscompute capabilitys%d.%ds pci device ids pci bus idis[NOT SUPPORTED: CC < 2.0]s [SUPPORTED]isid %d %20s %40ss%40s: %ssSummary:s %d/%d devices are supported(ii(RutprintRQtcompute_capabilityt PCI_DEVICE_IDt PCI_BUS_IDtidtname(tdevlisttsupported_counttdevtattrstcctsupporttkeytval((s-lib/python2.7/site-packages/numba/cuda/api.pytdetectOs&       cc@s(tƒj}|jƒ dVWdQXdS(sV 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(Rt deallocationstdisable(tdeallocs((s-lib/python2.7/site-packages/numba/cuda/api.pyt defer_cleanupms  (6t__doc__t __future__RRt contextlibtnumpyRtcudadrvRRRtargsRRRRNt NameErrorRMtrequire_contextRnRRtRR"R'R#R8R/tfloatR0R5R^R@tcontextmanagerRJRKRRORYR)RcR6Rmtevent_elapsed_timeRrRsRuRGRwR†RŠt profilingt profile_startt profile_stop(((s-lib/python2.7/site-packages/numba/cuda/api.pytsb         &