U
    hâËdÊ  ã                   @   sH  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-ej	dd	�d,d-„ ƒZ.ee.ƒZ/G d.d/„ d/eƒZ0d0d1„ Z1dS )2é    )Ú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   ú_/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/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)Úaligné   r   g     @E@c                 C   s   || | d< d S r#   r   )r'   ZaryÚir   r   r   Úrecord_return_   s    r5   )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(   r7   r   r6   r   Úmake_closurej   s    
r9   é   é   é   é	   c                 C   s   |d d | d< d S )Nr   r3   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ÚafÚagZahZaiZajZakÚalÚamZanZaoZapÚarÚatÚauÚavZawZaxZayÚazr   r   r   Úmany_localsŠ   s`    rP   c                 C   s   |d | d< d S r#   r   r>   r   r   r   Úsimple_usecase_kernelÁ   s    rQ   c                 C   s   t j ¡ }| ¡  d S r   )r   ZcgZ	this_gridÚsync)r'   r(   Úgridr   r   r   Úcg_usecase_kernelË   s    
rT   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 )Nr3   r:   r@   r   )r3   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   rU   Ô   s   rU   c                  C   s   t jt } tƒ  | ¡ d S r   )ÚsysÚmodulesr   rU   r`   )r_   r   r   r   Ú	self_testè   s    
rc   )2Znumbar   Znumba.cuda.testingr   Únumpyr   ra   r   r   r8   r*   r+   rV   Zadd_nocache_usecaser%   r,   r.   r/   rX   rW   r   Zint8rC   Zpacked_record_typeZaligned_record_typeÚemptyrZ   ÚrangeÚsizer4   rB   r]   r5   rY   r\   r9   Zclosure1Zclosure2Zclosure3Zclosure4r?   Zrenamed_function1Zrenamed_function2rP   rQ   r^   rT   Z
cg_usecaserU   rc   r   r   r   r   Ú<module>   sh    

















6



