U
    ?hi	                     @   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   U/var/www/html/venv/lib/python3.8/site-packages/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npZzerosr   ZarangefullshapeZtestingZassert_allclosenextZinspect_asmvaluesZassertIn)selfZkernelZnumpy_equivalentZptx_instructionZ
dtype_leftZdtype_rightncab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 r1   )r(   r   r   r2   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 r1   )r(   r   r   r2   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   r2   r-   r+   r   r   r   test_min_f4f4g   s    zTestCudaMinMax.test_min_f4f4N)r   )__name__
__module____qualname__r(   r,   r.   r/   r0   r3   r4   r5   r6   r   r   r   r   r      s   	 
r   __main__)numpyr   Znumbar   r   Znumba.cuda.testingr   r   r   r   r   r   r7   mainr   r   r   r   <module>   s   		V