U
    ,‰di	  ã                   @   sh   d dl Zd dlmZmZ d dlmZmZmZ dd„ Z	dd„ Z
edƒG d	d
„ d
eƒƒZedkrde ¡  dS )é    N)ÚcudaÚfloat64)ÚunittestÚCUDATestCaseÚskip_on_cudasimc                 C   s8   t  d¡}|t|ƒkrd S tt| | || ƒƒ||< d S ©Né   )r   ÚgridÚlenr   Úmax©ÚAÚBÚCÚi© r   úG/tmp/pip-unpacked-wheel-eu7e0c37/numba/cuda/tests/cudapy/test_minmax.pyÚbuiltin_max   s    
r   c                 C   s8   t  d¡}|t|ƒkrd S tt| | || ƒƒ||< d S r   )r   r	   r
   r   Úminr   r   r   r   Úbuiltin_min   s    
r   zTests PTX emissionc                   @   sV   e Zd Zddd„Zdd„ Zdd„ Zdd	„ Zd
d„ Zdd„ Zdd„ Z	dd„ Z
dd„ ZdS )ÚTestCudaMinMaxé   c                 C   s   t  |¡}tj|tjd}tj||dd }tj|d|d}	|d|jf ||	|ƒ tj 	||||	ƒ¡ t
dd„ | ¡  ¡ D ƒƒ}
|  ||
¡ d S )N)Údtypeg      à?é   )Z
fill_valuer   r   c                 s   s   | ]
}|V  qd S )Nr   )Ú.0Úpr   r   r   Ú	<genexpr>,   s     z&TestCudaMinMax._run.<locals>.<genexpr>)r   ZjitÚnpÚzerosr   ZarangeÚfullÚshapeZtestingZassert_allcloseÚnextZinspect_asmÚvaluesZassertIn)ÚselfÚkernelZnumpy_equivalentZptx_instructionZ
dtype_leftZdtype_rightÚnÚcÚaÚbZptxr   r   r   Ú_run   s    
zTestCudaMinMax._runc                 C   s   |   ttjdtjtj¡ d S ©Nzmax.f64)r)   r   r   Úmaximumr   ©r#   r   r   r   Útest_max_f8f8/   s    ûzTestCudaMinMax.test_max_f8f8c                 C   s   |   ttjdtjtj¡ d S r*   )r)   r   r   r+   Úfloat32r   r,   r   r   r   Útest_max_f4f87   s    ûzTestCudaMinMax.test_max_f4f8c                 C   s   |   ttjdtjtj¡ d S r*   )r)   r   r   r+   r   r.   r,   r   r   r   Útest_max_f8f4?   s    ûzTestCudaMinMax.test_max_f8f4c                 C   s   |   ttjdtjtj¡ d S )Nzmax.f32)r)   r   r   r+   r.   r,   r   r   r   Útest_max_f4f4G   s    ûzTestCudaMinMax.test_max_f4f4c                 C   s   |   ttjdtjtj¡ d S ©Nzmin.f64)r)   r   r   Úminimumr   r,   r   r   r   Útest_min_f8f8O   s    ûzTestCudaMinMax.test_min_f8f8c                 C   s   |   ttjdtjtj¡ d S r2   )r)   r   r   r3   r.   r   r,   r   r   r   Útest_min_f4f8W   s    ûzTestCudaMinMax.test_min_f4f8c                 C   s   |   ttjdtjtj¡ d S r2   )r)   r   r   r3   r   r.   r,   r   r   r   Útest_min_f8f4_   s    ûzTestCudaMinMax.test_min_f8f4c                 C   s   |   ttjdtjtj¡ d S )Nzmin.f32)r)   r   r   r3   r.   r,   r   r   r   Útest_min_f4f4g   s    ûzTestCudaMinMax.test_min_f4f4N)r   )Ú__name__Ú
__module__Ú__qualname__r)   r-   r/   r0   r1   r4   r5   r6   r7   r   r   r   r   r      s   	 ù
r   Ú__main__)Znumpyr   Znumbar   r   Znumba.cuda.testingr   r   r   r   r   r   r8   Úmainr   r   r   r   Ú<module>   s   		V