B
     \u                 @   s  d dl mZmZ d dlZd dlmZmZ d dlZd dlZd dl	Z	d dl
Z
d dlZd dlmZ d dlmZmZmZmZ d dlmZmZ d dlmZmZmZmZ d dlmZ d	d
lmZ d	dlm Z  d	dl!m"Z"m#Z#m$Z$ d	dl%m&Z& d	dl'm(Z( d	dl)m*Z* edd Z+edddg dfddZ,G dd de-Z.d/ddZ/d0ddZ0dd Z1G dd de-Z2G d d! d!e-Z3G d"d# d#e-Z4G d$d% d%e-Z5G d&d' d'e-Z6G d(d) d)e-Z7G d*d+ d+e5Z8d,Z9G d-d. d.e5Z:dS )1    )absolute_importprint_functionN)reducewraps)ctypes_support)configcompilertypessigutils)AbstractTemplateConcreteTemplate)funcdesctypingutils	serialize)global_compiler_lock   )	AutoTuner)get_context)nvvmdevicearraydriver)normalize_kernel_dimensions)get_current_device)wrap_argc          	   C   s   ddl m} |j}|j}t }|d |d |rL|d |d |rZ|d tj||| |||i d}	|	j}
|
	  |	S )	Nr   )CUDATargetDescZ
no_compileZno_cpython_wrapperZ
boundcheckZ	debuginfoZforceinline)	typingctx	targetctxfuncargsreturn_typeflagslocals)

descriptorr   r   r   r   ZFlagssetZcompile_extralibraryfinalize)pyfuncr    r   debuginliner   r   r   r!   cresr%    r+   2lib/python3.7/site-packages/numba/cuda/compiler.pycompile_cuda   s*    




r-   Fc             C   sl   t | tj|||d}|jj}	|jj|j|	|jj	|d\}
}t
|
j|j|jj|jj	|j|||j|||d}|S )N)r(   r)   )r(   )llvm_modulenamepretty_nameargtypestype_annotationlinkr(   call_helperfastmath
extensionsmax_registers)r-   r	   ZvoidfndescZllvm_func_nametarget_contextZprepare_cuda_kernelr%   	signaturer   
CUDAKernelZ_final_moduler/   qualnamer2   r4   )r'   r   r3   r(   r)   r5   r6   r7   r*   fnamelibkernelZcukernr+   r+   r,   compile_kernel;   s"    r@   c               @   s4   e Zd ZdZdd Zdd Zedd Zdd	 Zd
S )DeviceFunctionTemplatez#Unmaterialized device function
    c             C   s   || _ || _|| _i | _d S )N)py_funcr(   r)   _compileinfos)selfr'   r(   r)   r+   r+   r,   __init__U   s    zDeviceFunctionTemplate.__init__c             C   s6   t | j}t | j|}| j|| j| jf}t j|fS )N)r   #_get_function_globals_for_reductionrB   _reduce_function	__class__r(   r)   _rebuild_reduction)rD   glblsfunc_reducedr   r+   r+   r,   
__reduce__[   s    z!DeviceFunctionTemplate.__reduce__c             C   s   t j| }t|||dS )N)r(   r)   )r   _rebuild_functioncompile_device_template)clsrK   r(   r)   r   r+   r+   r,   _rebuilda   s    
zDeviceFunctionTemplate._rebuildc             C   sx   || j krht| jd|| j| jd}| j  }|| j |< |jg}|rT|j| |j| qr|j	| |j| n
| j | }|j
S )zCompile the function for the given argument types.

        Each signature is compiled once by caching the compiled function inside
        this object.
        N)r(   r)   )rC   r-   rB   r(   r)   r%   r9   insert_user_functionr8   Zadd_user_functionr:   )rD   r   r*   Zfirst_definitionZlibsr+   r+   r,   compilef   s    



zDeviceFunctionTemplate.compileN)	__name__
__module____qualname____doc__rE   rL   classmethodrP   rR   r+   r+   r+   r,   rA   R   s
   rA   c                sD   ddl m} t| ||d G  fdddt}|j}| |  S )zcCreate a DeviceFunctionTemplate object and register the object to
    the CUDA typing context.
    r   )r   )r(   r)   c                   s   e Zd Z Z fddZdS )z9compile_device_template.<locals>.device_function_templatec                s   |rt  |S )N)AssertionErrorrR   )rD   r   Zkws)dftr+   r,   generic   s    zAcompile_device_template.<locals>.device_function_template.genericN)rS   rT   rU   keyrZ   r+   )rY   r+   r,   device_function_template   s   r\   )r#   r   rA   r   r   rQ   )r'   r(   r)   r   r\   r   r+   )rY   r,   rN      s    rN   Tc             C   s   t | ||dddS )NTF)r)   r(   )DeviceFunction)r'   r    r   r)   r(   r+   r+   r,   compile_device   s    r^   c                st   ddl m} |j}|j}tj|f| t|  G  fdddt}tj	| ||d}|
 | |
 |  S )Nr   )r   c                   s   e Zd Z ZgZdS )z9declare_device_function.<locals>.device_function_templateN)rS   rT   rU   r[   casesr+   )extfnsigr+   r,   r\      s   r\   )r/   restyper1   )r#   r   r   r   r   r:   ExternFunctionr   r   ZExternalFunctionDescriptorrQ   )r/   rb   r1   r   r   r   r\   r8   r+   )r`   ra   r,   declare_device_function   s    
rd   c               @   s0   e Zd Zdd Zdd Zedd Zdd Zd	S )
r]   c                s~   |_ |_|_d_d_tj jjjjd  _G  fdddt} j	|  j
	 j jg d S )NTF)r(   r)   c                   s   e Zd ZZ jgZdS )z9DeviceFunction.__init__.<locals>.device_function_templateN)rS   rT   rU   r[   r:   r_   r+   )r*   rD   r+   r,   r\      s   r\   )rB   r    r   r)   r(   r-   r*   r   Ztyping_contextrQ   r9   r8   r%   )rD   r'   r    r   r)   r(   r\   r+   )r*   rD   r,   rE      s    zDeviceFunction.__init__c             C   s>   t | j}t | j|}| j|| j| j| j| jf}t j	|fS )N)
r   rF   rB   rG   rH   r    r   r)   r(   rI   )rD   ZglobsrK   r   r+   r+   r,   rL      s
    zDeviceFunction.__reduce__c             C   s   | t j| ||||S )N)r   rM   )rO   rK   r    r   r)   r(   r+   r+   r,   rP      s    zDeviceFunction._rebuildc             C   s   d}| | j| jjS )Nz*<DeviceFunction py_func={0} signature={1}>)formatrB   r*   r:   )rD   Zfmtr+   r+   r,   __repr__   s    zDeviceFunction.__repr__N)rS   rT   rU   rE   rL   rW   rP   rf   r+   r+   r+   r,   r]      s   r]   c               @   s   e Zd Zdd ZdS )rc   c             C   s   || _ || _d S )N)r/   ra   )rD   r/   ra   r+   r+   r,   rE      s    zExternFunction.__init__N)rS   rT   rU   rE   r+   r+   r+   r,   rc      s   rc   c               @   s,   e Zd Zdd Zdd Zdd Zdd Zd	S )
ForAllc             C   s"   || _ || _|| _|| _|| _d S )N)r?   ntasksthread_per_blockstream	sharedmem)rD   r?   rh   tpbrj   rk   r+   r+   r,   rE      s
    zForAll.__init__c             G   sZ   t | jtr| jj| }n| j}| |}|d }| j| | }|j||| j| jd| S )Nr   )rj   rk   )	
isinstancer?   AutoJitCUDAKernel
specialize_compute_thread_per_blockrh   	configurerj   rk   )rD   r   r?   rl   Ztpbm1Zblkctr+   r+   r,   __call__   s    
zForAll.__call__c             C   sr   | j }|dkr|S t }t|j dd | jdd}y|jf |\}}W n  tk
rh   | |} Y nX |S d S )Nr   c             S   s   dS )Nr   r+   )rl   r+   r+   r,   <lambda>   s    z2ForAll._compute_thread_per_block.<locals>.<lambda>i   )r   Zb2d_funcZmemsizeZblocksizelimit)	ri   r   dict_funcgetrk   Zget_max_potential_block_sizeAttributeError_fallback_autotune_best)rD   r?   rl   Zctxkwargs_r+   r+   r,   rp      s    
z ForAll._compute_thread_per_blockc             C   s6   y|j  }W n" tk
r0   td d}Y nX |S )Nz,Could not autotune, using default tpb of 128   )autotuneZbest
ValueErrorwarningswarn)rD   r?   rl   r+   r+   r,   rx     s    

zForAll._fallback_autotune_bestN)rS   rT   rU   rE   rr   rp   rx   r+   r+   r+   r,   rg      s   rg   c               @   sL   e Zd ZdZdd Zdd ZdddZd	d
 ZdddZdd Z	dd Z
dS )CUDAKernelBasez.Define interface for configurable kernels
    c             C   s   d| _ d| _d| _d| _d S )N)r   r   )r   r   r   r   )griddimblockdimrk   rj   )rD   r+   r+   r,   rE     s    zCUDAKernelBase.__init__c             C   s"   | j }||}|j| j |S )z+
        Shallow copy the instance
        )rH   __new____dict__update)rD   rO   newr+   r+   r,   copy  s    
zCUDAKernelBase.copyr   c             C   s:   t ||\}}|  }t||_t||_||_||_|S )N)r   r   tupler   r   rj   rk   )rD   r   r   rj   rk   Zcloner+   r+   r,   rq   $  s    

zCUDAKernelBase.configurec             C   s   t |dkrtd| j| S )N)         z.must specify at least the griddim and blockdim)lenr}   rq   )rD   r   r+   r+   r,   __getitem__.  s    zCUDAKernelBase.__getitem__c             C   s   t | ||||dS )zReturns a configured kernel for 1D kernel of given number of tasks
        ``ntasks``.

        This assumes that:
        - the kernel 1-to-1 maps global thread id ``cuda.grid(1)`` to tasks.
        - the kernel must check if the thread id is valid.)rl   rj   rk   )rg   )rD   rh   rl   rj   rk   r+   r+   r,   forall3  s    zCUDAKernelBase.forallc             C   s   | j | j| jfS )z
        Helper for serializing the grid, block and shared memory configuration.
        CUDA stream config is not serialized.
        )r   r   rk   )rD   r+   r+   r,   _serialize_config=  s    z CUDAKernelBase._serialize_configc             C   s   |\| _ | _| _dS )zc
        Helper for deserializing the grid, block and shared memory
        configuration.
        N)r   r   rk   )rD   r   r+   r+   r,   _deserialize_configD  s    z"CUDAKernelBase._deserialize_configN)r   r   )r   r   r   )rS   rT   rU   rV   rE   r   rq   r   r   r   r   r+   r+   r+   r,   r     s   



r   c               @   s    e Zd ZdZdd Zdd ZdS )	CachedPTXz<A PTX cache that uses compute capability as a cache key
    c             C   s    || _ || _i | _| | _d S )N)r/   llvmircacher   _extra_options)rD   r/   r   optionsr+   r+   r,   rE   O  s    zCachedPTX.__init__c             C   s   t  }|j}|j}| j|}|dkrtj| }tj| jfd|d| j	}|| j|< t
jrtd| j dd t|d td |S )	z9
        Get PTX for the current active context.
        Nr   )ZoptarchzASSEMBLY %sP   -zutf-8zP================================================================================)r   devicecompute_capabilityr   rv   r   Zget_arch_optionZllvm_to_ptxr   r   r   ZDUMP_ASSEMBLYprintr/   centerdecode)rD   cuctxr   ccptxr   r+   r+   r,   rv   U  s    


zCachedPTX.getN)rS   rT   rU   rV   rE   rv   r+   r+   r+   r,   r   L  s   r   c               @   s<   e Zd ZdZdd Zdd Zdd Zdd	 Zed
d Z	dS )CachedCUFunctionzk
    Get or compile CUDA function for the current active context

    Uses device ID as key for cache.
    c             C   s(   || _ || _|| _i | _i | _|| _d S )N)
entry_namer   linkingr   ccinfosr7   )rD   r   r   r   r7   r+   r+   r,   rE   p  s    zCachedCUFunction.__init__c             C   s   t  }|j}| j|j}|d kr| j }tj| jd}|	| x| j
D ]}|| qLW | \}}|j}	||}
|
| j}|| j|j< |	| j|j< |S )N)r7   )r   r   r   rv   idr   r   ZLinkerr7   Zadd_ptxr   Zadd_file_guess_extZcompleteZinfo_logZcreate_module_imageZget_functionr   r   )rD   r   r   cufuncr   ZlinkerpathZcubinZ_sizeZcompile_infomoduler+   r+   r,   rv   x  s     


zCachedCUFunction.getc             C   s$   |    t }|j}| j|j }|S )N)rv   r   r   r   r   )rD   r   r   Zcir+   r+   r,   get_info  s
    zCachedCUFunction.get_infoc             C   s4   | j rd}t|| j| j| j| j | jf}tj|fS )z
        Reduce the instance for serialization.
        Pre-compiled PTX code string is serialized inside the `ptx` (CachedPTX).
        Loaded CUfunctions are discarded. They are recreated when unserialized.
        zLcannot pickle CUDA kernel function with additional libraries to link against)r   RuntimeErrorrH   r   r   r7   r   rI   )rD   msgr   r+   r+   r,   rL     s
    zCachedCUFunction.__reduce__c             C   s   | ||||S )z&
        Rebuild an instance.
        r+   )rO   r   r   r   r7   r+   r+   r,   rP     s    zCachedCUFunction._rebuildN)
rS   rT   rU   rV   rE   rv   r   rL   rW   rP   r+   r+   r+   r,   r   i  s   r   c                   s   e Zd ZdZddddg df fdd	Ze fddZd	d
 Zdd Zdd Z	e
dd Ze
dd Zdd Zdd Zd"ddZd#ddZdd Ze
dd Ze
d d! Z  ZS )$r;   z
    CUDA Kernel specialized for a given set of argument types. When called, this
    object will validate that the argument types match those for which it is
    specialized, and then launch the kernel on the device.
    r+   FNc                s   t t|   d|i}|r0|tddddd t|t||d}t||||}|| _t	|| _
t	|| _|	| _|| _|| _|| _t|
| _d S )Nr(   TF)ZftzZ	prec_sqrtZprec_divZfma)r   )superr;   rE   r   rt   r   strr   r   r   argument_typesr   _type_annotationru   r(   r4   listr6   )rD   r.   r/   r0   r1   r4   r3   r(   r5   r2   r6   r7   r   r   r   )rH   r+   r,   rE     s"    


zCUDAKernel.__init__c	       
         s^   |  | }	t| |	  ||	_t||	_t||	_d|	_||	_||	_	||	_
||	_|	| |	S )z&
        Rebuild an instance.
        N)r   r   rE   r   r   r   r   r   ru   r(   r4   r6   r   )
rO   r/   r1   r   r3   r(   r4   r6   r   instance)rH   r+   r,   rP     s    



zCUDAKernel._rebuildc          	   C   s8   |   }| j| j| j| j| j| j| j| j|f	}t	j
|fS )a  
        Reduce the instance for serialization.
        Compiled definitions are serialized in PTX form.
        Type annotation are discarded.
        Thread, block and shared memory configuration are serialized.
        Stream information is discarded.
        )r   rH   r   r   ru   r   r(   r4   r6   r   rI   )rD   r   r   r+   r+   r,   rL     s
    
zCUDAKernel.__reduce__c             O   s(   |rt | j|| j| j| j| jd d S )N)r   r   r   rj   rk   )rX   _kernel_callr   r   rj   rk   )rD   r   ry   r+   r+   r,   rr     s    zCUDAKernel.__call__c             C   s   | j   dS )z7
        Force binding to current CUDA context
        N)ru   rv   )rD   r+   r+   r,   bind  s    zCUDAKernel.bindc             C   s   | j j dS )z+
        PTX code for this kernel.
        utf8)ru   r   rv   r   )rD   r+   r+   r,   r     s    zCUDAKernel.ptxc             C   s   t  S )z,
        Get current active context
        )r   )rD   r+   r+   r,   r      s    zCUDAKernel.devicec             C   s   t | jjjS )z6
        Returns the LLVM IR for this kernel.
        )r   ru   r   r   )rD   r+   r+   r,   inspect_llvm  s    zCUDAKernel.inspect_llvmc             C   s   | j j dS )z7
        Returns the PTX code for this kernel.
        ascii)ru   r   rv   r   )rD   r+   r+   r,   inspect_asm  s    zCUDAKernel.inspect_asmc             C   sb   | j dkrtd|dkr tj}td| j| jf |d td|d t| j |d td|d dS )z
        Produce a dump of the Python source of this function annotated with the
        corresponding Numba IR and type information. The dump is written to
        *file*, or *sys.stdout* if *file* is *None*.
        Nz Type annotation is not availablez%s %s)filezP--------------------------------------------------------------------------------zP================================================================================)r   r}   sysstdoutr   r   r   )rD   r   r+   r+   r,   inspect_types  s    
zCUDAKernel.inspect_typesr   c                s  | j   | jrT jd } j|\}}|ttjks>t	t }	|j
d|d g }
g }x*t| j|D ]\}}| ||||
| qjW  j||||d}||  | jrtt|	|| |	jdkr fddfddd	D }fd
dd	D }|	j}| j|\}}}|d kr"d}n$|\}}}tj|}d|||f }d|||f }|rzd||d f f|dd   }n|f}|| x|
D ]}|  qW d S )NZ__errcode__r   )rj   )rj   rk   c                s<    j d j| f \}}t }tt||| |jS )Nz%s__%s__)	r   get_global_symbolr/   ctypesc_intr   device_to_host	addressofvalue)r/   ZmemZszval)r   r+   r,   load_symbolA  s    z,CUDAKernel._kernel_call.<locals>.load_symbolc                s   g | ]} d | qS )tidr+   ).0i)r   r+   r,   
<listcomp>I  s    z+CUDAKernel._kernel_call.<locals>.<listcomp>Zzyxc                s   g | ]} d | qS )ctaidr+   )r   r   )r   r+   r,   r   J  s     z"In function %r, file %s, line %s, z%stid=%s ctaid=%sz%s: %sr   )ru   rv   r(   r/   r   r   r   Zsizeofr   rX   Zmemsetzipr   _prepare_argsrq   r   r   r   r   r4   Zget_exceptionosr   relpath)rD   r   r   r   rj   rk   ZexcnameZexcmemZexcszZexcvalretr
kernelargstvZcu_funcr   r   codeZexcclsZexc_argsZlocZlocinfoZsymfilepathlinenoprefixwbr+   )r   r   r,   r   $  sH    



 
zCUDAKernel._kernel_callc                s  x(t | jD ]}|j| ||d\} qW t|tjrPt|tjrv d}| fdd t	
d}|| nt ||}t	j}	t	
d}
t	
d}|	|j}|	|jj}t	
t|}||
 || || || || x(t|jD ]}||	|j|  qW x(t|jD ]}||	|j|  q.W n`t|tjr~tt	d|  }|| n2|tjkrt	 }|| n|tjkrt	 }|| n|tjkrt	t }|| n|tj kr|t	 j! |t	 j" n|tj#krN|t	 j! |t	 j" nbt|tj$tj%fr||t	& 't(j) n4t|tj*rt ||}|| n
t+| dS )zF
        Convert arguments to ctypes and append to kernelargs
        )rj   r   gpuc                  s
     dS )Nr   )Zmark_changedr+   )r   r+   r,   rs   r  s    z*CUDAKernel._prepare_args.<locals>.<lambda>r   zc_%sN),reversedr6   Zprepare_argsrm   r	   ZArrayZSmartArrayTyperv   appendr   Zc_void_pr   Z	to_deviceZ	c_ssize_tsizeZdtypeitemsizer   Zdevice_pointerrangendimshapestridesZIntegergetattrZfloat64Zc_doubleZfloat32Zc_floatZbooleanZc_uint8intZ	complex64realimagZ
complex128Z
NPDatetimeZNPTimedeltaZc_int64ZviewnpZint64ZRecordNotImplementedError)rD   Ztyr   rj   r   r   	extensionZdevaryZouter_parentZc_intpZmeminfoparentZnitemsr   dataZaxZcvalZdevrecr+   )r   r,   r   b  sh    











zCUDAKernel._prepare_argsc             C   s`   t tdt t| d}|r4| jj| jkr4| jS | j	
 }t|j|jjd}|| _| jS dS )z8Return the autotuner object associated with this kernel.r|   	_autotune)infor   N)r~   r   _deprec_warn_msgre   DeprecationWarninghasattrr   Zdynsmemrk   ru   rv   r   Zattrsr   r   )rD   Zhas_autotuner   Zatr+   r+   r,   r|     s    

zCUDAKernel.autotunec             C   s.   t tdt ttj| jd}| j	
|S )a  Occupancy is the ratio of the number of active warps per multiprocessor to the maximum
        number of warps that can be active on the multiprocessor at once.
        Calculate the theoretical occupancy of the kernel given the
        current configuration.	occupancyr   )r~   r   r   re   r   r   operatormulr   r|   Zclosest)rD   ri   r+   r+   r,   r     s    zCUDAKernel.occupancy)N)r   r   )rS   rT   rU   rV   rE   rW   rP   rL   rr   r   propertyr   r   r   r   r   r   r   r|   r   __classcell__r+   r+   )rH   r,   r;     s    

>Ir;   zJThe .{} attribute is is deprecated and will be removed in a future releasec                   sv   e Zd ZdZ fddZedd Zdd Zdd	 Zd
d Z	dddZ
dddZdddZedd Zdd Z  ZS )rn   a  
    CUDA Kernel object. When called, the kernel object will specialize itself
    for the given arguments (if no suitable specialized version already exists)
    & compute capability, and launch on the device associated with the current
    context.

    Kernel objects are not to be constructed by the user, but instead are
    created using the :func:`numba.cuda.jit` decorator.
    c                sV   t t|   || _|| _i | _|| _t| jdg | jd< ddl	m
} |j| _d S )Nr6   r   )r   )r   rn   rE   rB   r   definitionstargetoptionsr   rv   r#   r   r   )rD   r   r   r   r   )rH   r+   r,   rE     s    zAutoJitCUDAKernel.__init__c             C   s
   | j d S )aS  
        A list of objects that must have a `prepare_args` function. When a
        specialized kernel is called, each argument will be passed through
        to the `prepare_args` (from the last object in this list to the
        first). The arguments to `prepare_args` are:

        - `ty` the numba type of the argument
        - `val` the argument value itself
        - `stream` the CUDA stream used for the current call to the kernel
        - `retr` a list of zero-arg functions that you may want to append
          post-call cleanup work to.

        The `prepare_args` function must return a tuple `(ty, val)`, which
        will be passed in turn to the next right-most `extension`. After all
        the extensions have been called, the resulting `(ty, val)` will be
        passed into Numba's default argument marshalling logic.
        r6   )r   )rD   r+   r+   r,   r6     s    zAutoJitCUDAKernel.extensionsc             G   s.   | j | }|| j| j| j| jf }||  dS )z@
        Specialize and invoke this kernel with *args*.
        N)ro   r   r   rj   rk   )rD   r   r?   Zcfgr+   r+   r,   rr     s    
zAutoJitCUDAKernel.__call__c                s$   t  fdd|D } |}|S )z|
        Compile and bind to the current context a version of this kernel
        specialized for the given *args*.
        c                s   g | ]} j |qS r+   )r   Zresolve_argument_type)r   a)rD   r+   r,   r     s    z0AutoJitCUDAKernel.specialize.<locals>.<listcomp>)r   rR   )rD   r   r1   r?   r+   )rD   r,   ro     s    
zAutoJitCUDAKernel.specializec             C   s   t |\}}|dkstt j}| j||f}|dkr|d| jkrNd| jd< t| j	|f| j}|| j||f< | j
r||
  |S )z
        Compile and bind to the current context a version of this kernel
        specialized for the given signature.
        Nr3   r+   )r
   Znormalize_signaturerX   r   r   r   rv   r   r@   rB   r   )rD   ra   r1   r    r   r?   r+   r+   r,   rR     s    


zAutoJitCUDAKernel.compileNc             C   sB   |p
t  j}|dk	r&| j||f  S tdd | j D S dS )z
        Return the LLVM IR for all signatures encountered thus far, or the LLVM
        IR for a specific signature and compute_capability if given.
        Nc             s   s   | ]\}}||  fV  qd S )N)r   )r   ra   defnr+   r+   r,   	<genexpr>'  s   z1AutoJitCUDAKernel.inspect_llvm.<locals>.<genexpr>)r   r   r   r   rt   items)rD   r:   r   r   r+   r+   r,   r     s
    zAutoJitCUDAKernel.inspect_llvmc             C   sB   |p
t  j}|dk	r&| j||f  S tdd | j D S dS )z
        Return the generated assembly code for all signatures encountered thus
        far, or the LLVM IR for a specific signature and compute_capability
        if given.
        Nc             s   s   | ]\}}||  fV  qd S )N)r   )r   ra   r   r+   r+   r,   r   4  s   z0AutoJitCUDAKernel.inspect_asm.<locals>.<genexpr>)r   r   r   r   rt   r   )rD   r:   r   r   r+   r+   r,   r   *  s
    zAutoJitCUDAKernel.inspect_asmc             C   s8   |dkrt j}x$t| jD ]\}}|j|d qW dS )z
        Produce a dump of the Python source of this function annotated with the
        corresponding Numba IR and type information. The dump is written to
        *file*, or *sys.stdout* if *file* is *None*.
        N)r   )r   r   r   Z	iteritemsr   r   )rD   r   rz   r   r+   r+   r,   r   7  s    zAutoJitCUDAKernel.inspect_typesc             C   s$   t j| }| |||}|| |S )z&
        Rebuild an instance.
        )r   rM   r   )rO   rK   r   r   r   r   r   r+   r+   r,   rP   C  s    

zAutoJitCUDAKernel._rebuildc             C   s@   t | j}t | j|}|  }| j|| j| j|f}t j|fS )zd
        Reduce the instance for serialization.
        Compiled definitions are discarded.
        )	r   rF   rB   rG   r   rH   r   r   rI   )rD   rJ   rK   r   r   r+   r+   r,   rL   M  s    zAutoJitCUDAKernel.__reduce__)NN)NN)N)rS   rT   rU   rV   rE   r   r6   rr   ro   rR   r   r   r   rW   rP   rL   r   r+   r+   )rH   r,   rn     s   	




rn   )FF)TF);Z
__future__r   r   r   	functoolsr   r   r   r   Z	threadingr~   Znumpyr   Znumbar   r   r   r   r	   r
   Znumba.typing.templatesr   r   r   r   r   r   Znumba.compiler_lockr   Zcudadrv.autotuner   Zcudadrv.devicesr   Zcudadrvr   r   r   errorsr   Zapir   r   r   r-   r@   objectrA   rN   r^   rd   r]   rc   rg   r   r   r   r;   r   rn   r+   r+   r+   r,   <module>   sJ    .

&6>A  