o
    Y۷ii	                     @   sl   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r4e  dS dS )    N)cudafloat64)unittestCUDATestCaseskip_on_cudasimc                 C   8   t d}|t|krd S tt| | || ||< d S N   )r   gridlenr   maxABCi r   Y/home/ubuntu/vllm_env/lib/python3.10/site-packages/numba/cuda/tests/cudapy/test_minmax.pybuiltin_max      
r   c                 C   r   r   )r   r
   r   r   minr   r   r   r   builtin_min   r   r   zTests PTX emissionc                   @   sX   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      ?   )
fill_valuer   r	   c                 s   s    | ]}|V  qd S )Nr   ).0pr   r   r   	<genexpr>,   s    z&TestCudaMinMax._run.<locals>.<genexpr>)r   jitnpzerosr   arangefullshapetestingassert_allclosenextinspect_asmvaluesassertIn)selfkernelnumpy_equivalentptx_instruction
dtype_leftdtype_rightncabptxr   r   r   _run   s   
zTestCudaMinMax._runc                 C      |  ttjdtjtj d S Nzmax.f64)r7   r   r!   maximumr   r,   r   r   r   test_max_f8f8/      zTestCudaMinMax.test_max_f8f8c                 C      |  ttjdtjtj d S r9   )r7   r   r!   r:   float32r   r;   r   r   r   test_max_f4f87   r=   zTestCudaMinMax.test_max_f4f8c                 C   r>   r9   )r7   r   r!   r:   r   r?   r;   r   r   r   test_max_f8f4?   r=   zTestCudaMinMax.test_max_f8f4c                 C   r8   )Nzmax.f32)r7   r   r!   r:   r?   r;   r   r   r   test_max_f4f4G   r=   zTestCudaMinMax.test_max_f4f4c                 C   r8   Nzmin.f64)r7   r   r!   minimumr   r;   r   r   r   test_min_f8f8O   r=   zTestCudaMinMax.test_min_f8f8c                 C   r>   rC   )r7   r   r!   rD   r?   r   r;   r   r   r   test_min_f4f8W   r=   zTestCudaMinMax.test_min_f4f8c                 C   r>   rC   )r7   r   r!   rD   r   r?   r;   r   r   r   test_min_f8f4_   r=   zTestCudaMinMax.test_min_f8f4c                 C   r8   )Nzmin.f32)r7   r   r!   rD   r?   r;   r   r   r   test_min_f4f4g   r=   zTestCudaMinMax.test_min_f4f4N)r   )__name__
__module____qualname__r7   r<   r@   rA   rB   rE   rF   rG   rH   r   r   r   r   r      s    	
r   __main__)numpyr!   numbar   r   numba.cuda.testingr   r   r   r   r   r   rI   mainr   r   r   r   <module>   s    		V