U
    ,d                     @   s,  d dl mZ d dlmZ d dlZd dlZG dd dZG dd deZej	dd	d
d Z
ej	dd	dd Zee
ZeeZdZej	dd	dd Zej	dd	dd Zej	dd	dd ZeeZeeZedejfdejfgZejdejfdejfgddZejdedZeejD ]&Zed ee d< ed ee d< qejeedZej	dd	dd Z ee edZ!ee edZ"dd  Z#e#d!Z$e#d"Z%e#d#Z&e#d$Z'ej	dd	d%d& Z(ee(Z)ej	dd	d'd& Z(ee(Z*ej	dd	d(d) Z+ej	dd	d*d+ Z,ee,Z-G d,d- d-eZ.d.d/ Z/dS )0    )cuda)CUDATestCaseNc                   @   s.   e Zd ZdZd	ddZdd Zedd ZdS )
UseCasea2  
    Provide a way to call a kernel as if it were a function.

    This allows the CUDA cache tests to closely match the CPU cache tests, and
    also to support calling cache use cases as njitted functions. The class
    wraps a function that takes an array for the return value and arguments,
    and provides an interface that accepts arguments, launches the kernel
    appropriately, and returns the stored return value.

    The return type is inferred from the type of the first argument, unless it
    is explicitly overridden by the ``retty`` kwarg.
    Nc                 C   s   || _ || _d S N)_func_retty)selffuncretty r   J/tmp/pip-unpacked-wheel-eu7e0c37/numba/cuda/tests/cudapy/cache_usecases.py__init__   s    zUseCase.__init__c                 G   sL   dd |D }| j r&tjd| j d}nt|d }| j|f|  |d S )Nc                 S   s   g | ]}t |qS r   )npZasarray).0argr   r   r   
<listcomp>   s     z$UseCase.__call__.<locals>.<listcomp>r   dtyper   )r   r   ZndarrayZ
zeros_like_call)r   argsZ
array_argsZarray_returnr   r   r   __call__   s    zUseCase.__call__c                 C   s   | j S r   r   )r   r   r   r   r	   "   s    zUseCase.func)N)__name__
__module____qualname____doc__r   r   propertyr	   r   r   r   r   r      s
   

r   c                   @   s   e Zd Zdd ZdS )CUDAUseCasec                 G   s   | j d |f|  d S )N   r   r   )r   retr   r   r   r   r   (   s    zCUDAUseCase._callN)r   r   r   r   r   r   r   r   r   '   s   r   Tcachec                 C   s   |d |d  t  | d< d S Nr   Zrxyr   r   r   add_usecase_kernel,   s    r*   Fc                 C   s   |d |d  t  | d< d S r#   r$   r&   r   r   r   add_nocache_usecase_kernel1   s    r+   r   c                 C   s   | | t  S r   r$   )r(   r)   r   r   r   inner>   s    r,   c                 C   s   t |d  |d | d< d S r#   r,   r&   r   r   r   outer_kernelC   s    r.   c                 C   s   t |d  |d | d< d S r#   r-   r&   r   r   r   outer_uncached_kernelH   s    r/   ab)Zalign   r   g     @E@c                 C   s   || | d< d S r#   r   )r'   Zaryir   r   r   record_return_   s    r4   )r
   c                    s    t jdd fdd}t|S )NTr!   c                    s    |d  | d< d S r#   r   )r'   r)   r(   r   r   closurek   s    zmake_closure.<locals>.closure)r   jitr   )r(   r6   r   r5   r   make_closurej   s    
r8            	   c                 C   s   |d d | d< d S )Nr   r2   r   r'   r(   r   r   r   ambiguous_functionz   s    r>   c                 C   s   |d d | d< d S )Nr      r   r=   r   r   r   r>      s    c                  C   s  t jdtj} t jdtj}t jdtj}t jdtj}t jdtj}t jdtj}t jdtj}t jdtj}t jdtj}t jdtj}	t jdtj}
t jdtj}t jdtj}t jdtj}t jdtj}t jdtj}t jdtj}t jdtj}t jdtj}t jdtj}t jdtj}t jdtj}t jdtj}t jdtj}d| d d < d|d d < d|d d < d|d d < d|d d < d|d d < d|d d < d|d d < d|d d < d|	d d < d|
d d < d|d d < d|d d < d|d d < d|d d < d|d d < d|d d < d|d d < d|d d < d|d d < d|d d < d|d d < d|d d < d|d d < d S )Nr   r   )r   localarrayr   float64)ZaaabacadZaeafZagZahZaiZajZakalamZanZaoZaparatauavZawZaxZayazr   r   r   many_locals   s`    rN   c                 C   s   |d | d< d S r#   r   r=   r   r   r   simple_usecase_kernel   s    rO   c                   @   s   e Zd ZdZdd ZdS )_TestModulez
    Tests for functionality of this module's functions.
    Note this does not define any "test_*" method, instead check_module()
    should be called by hand.
    c                 C   s   |  |ddd |  |ddd |  |ddd ||jd}|  t|d ||jd}|  t|d |	d d S )Nr2   r9   r?   r   )r2   g     E@)
ZassertPreciseEqualadd_usecaseouter_uncachedouterrecord_return_packed
packed_arrtuplerecord_return_alignedaligned_arrsimple_usecase_caller)r   modZ
packed_recZaligned_recr   r   r   check_module   s    z_TestModule.check_moduleN)r   r   r   r   r[   r   r   r   r   rP      s   rP   c                  C   s   t jt } t |  d S r   )sysmodulesr   rP   r[   )rZ   r   r   r   	self_test   s    
r^   )0Znumbar   Znumba.cuda.testingr   Znumpyr   r\   r   r   r7   r*   r+   rQ   Zadd_nocache_usecaser%   r,   r.   r/   rS   rR   r   Zint8rB   Zpacked_record_typeZaligned_record_typeemptyrU   rangesizer3   rA   rX   r4   rT   rW   r8   Zclosure1Zclosure2Zclosure3Zclosure4r>   Zrenamed_function1Zrenamed_function2rN   rO   rY   rP   r^   r   r   r   r   <module>   sb    

















6

