a
    
d                      @   s   d dl mZ d dlmZmZ d dlmZmZ d dlm	Z	m
Z
 d dlmZmZmZmZmZmZmZmZmZ d dlmZ d dlZd dlmZmZmZ d dlZeG d	d
 d
ZedG dd deZe dkre!  dS )    )List)	dataclassfield)cudafloat32)compile_ptx_for_current_devicecompile_ptx)	cossintanexploglog10log2powtanh)truedivN)CUDATestCaseskip_on_cudasimskip_unless_cc_75c                   @   sx   e Zd ZU eedZee ed< eedZ	ee ed< eedZ
ee ed< eedZee ed< eeedddZd	S )
FastMathCriterion)default_factoryfast_expectedfast_unexpectedprec_expectedprec_unexpected)testfastprecc                    s|   | t fdd| jD  | t fdd| jD  | tfdd| jD  | tfdd| jD  d S )Nc                 3   s   | ]}| v V  qd S N .0ir   r    n/var/www/html/stable-diffusion-webui/venv/lib/python3.9/site-packages/numba/cuda/tests/cudapy/test_fastmath.py	<genexpr>       z*FastMathCriterion.check.<locals>.<genexpr>c                 3   s   | ]}| vV  qd S r   r    r!   r$   r    r%   r&      r'   c                 3   s   | ]}| v V  qd S r   r    r!   r   r    r%   r&      r'   c                 3   s   | ]}| vV  qd S r   r    r!   r(   r    r%   r&      r'   )
assertTrueallr   r   r   r   )selfr   r   r   r    )r   r   r%   check   s    zFastMathCriterion.checkN)__name__
__module____qualname__r   listr   r   str__annotations__r   r   r   r   r,   r    r    r    r%   r      s
   
r   z4Fastmath and PTX inspection not available on cudasimc                   @   s   e Zd Zdd ZedddZedddZdd	 Zd
d Zdd Z	e
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ejd d! Zd"S )#TestFastMathOptionc           
      C   sx   t j||dd|}t j||d|}|| |||| t|||dd\}}t|||d\}	}|| ||	 d S )NT)devicefastmathr4   )r   jitr,   inspect_asmr   )
r+   Zpyfuncsigr4   	criterionfastverprecverfastptx_precptxr    r    r%   _test_fast_math_common   s    

z)TestFastMathOption._test_fast_math_common)r:   c                    sP    fdd} fdd}| j |td d d tfd|d | j |tfd|d d S )	Nc                    s    || d< d S Nr   r    rxopr    r%   kernel2   s    z8TestFastMathOption._test_fast_math_unary.<locals>.kernelc                    s    | S r   r    )rD   rE   r    r%   device_function5   s    zATestFastMathOption._test_fast_math_unary.<locals>.device_function   Fr4   r:   Tr@   r   )r+   rF   r:   rG   rH   r    rE   r%   _test_fast_math_unary1   s    
z(TestFastMathOption._test_fast_math_unaryc                    sT    fdd} fdd}| j |td d d ttfd|d | j |ttfd|d d S )	Nc                    s    ||| d< d S rA   r    rC   rD   yrE   r    r%   rG   @   s    z9TestFastMathOption._test_fast_math_binary.<locals>.kernelc                    s
    | |S r   r    )rD   rN   rE   r    r%   r4   C   s    z9TestFastMathOption._test_fast_math_binary.<locals>.devicerI   FrJ   TrK   )r+   rF   r:   rG   r4   r    rE   r%   _test_fast_math_binary?   s    z)TestFastMathOption._test_fast_math_binaryc                 C   s   |  ttdgdgd d S )Ncos.approx.ftz.f32 r   r   )rL   r	   r   r+   r    r    r%   	test_cosfN   s    zTestFastMathOption.test_cosfc                 C   s   |  ttdgdgd d S )Nsin.approx.ftz.f32 rQ   )rL   r
   r   rR   r    r    r%   	test_sinfW   s    zTestFastMathOption.test_sinfc                 C   s   |  ttg ddgd d S )N)rT   rP   div.approx.ftz.f32 rT   rQ   )rL   r   r   rR   r    r    r%   	test_tanf`   s    zTestFastMathOption.test_tanfc                 C   s   |  ttdgdgd d S )Ntanh.approx.f32 rQ   )rL   r   r   rR   r    r    r%   
test_tanhfj   s    zTestFastMathOption.test_tanhfc                    sL   dd  fdd}|dt dgdgdd |d	t d
dgdgdd d S )Nc                 S   s   t || d< d S rA   )r   rB   r    r    r%   tanh_kernelv   s    z>TestFastMathOption.test_tanhf_compile_ptx.<locals>.tanh_kernelc                    sT   t td d d tfd| d\}}t td d d tf| d\}}| || d S )NrI   T)r5   cc)r[   )r   r   r,   )r[   r:   r=   r>   r?   r+   rZ   r    r%   tanh_common_testy   s    

zCTestFastMathOption.test_tanhf_compile_ptx.<locals>.tanh_common_test)      rX   rQ   )r[   r:   )r^   r   zex2.approx.ftz.f32 zrcp.approx.ftz.f32 )r   )r+   r]   r    r\   r%   test_tanhf_compile_ptxu   s    
z)TestFastMathOption.test_tanhf_compile_ptxc                 C   s   |  ttdgdgd d S )Nzfma.rn.f32 )r   r   )rL   r   r   rR   r    r    r%   	test_expf   s    zTestFastMathOption.test_expfc                 C   s   |  ttddgdgd d S )Nlg2.approx.ftz.f32 Z
0f3F317218rQ   )rL   r   r   rR   r    r    r%   	test_logf   s    zTestFastMathOption.test_logfc                 C   s   |  ttddgdgd d S )Nrb   Z
0f3E9A209BrQ   )rL   r   r   rR   r    r    r%   test_log10f   s    zTestFastMathOption.test_log10fc                 C   s   |  ttdgdgd d S Nrb   rQ   )rL   r   r   rR   r    r    r%   
test_log2f   s    zTestFastMathOption.test_log2fc                 C   s   |  ttdgdgd d S re   )rO   r   r   rR   r    r    r%   	test_powf   s    zTestFastMathOption.test_powfc              	   C   s$   |  ttdgdgdgdgd d S )NrV   
div.rn.f32)r   r   r   r   )rO   r   r   rR   r    r    r%   	test_divf   s    zTestFastMathOption.test_divfc                 C   s   dd }t d d d t t f}tj|ddd|}tj|dd|}d}tj|tj d}| t$ |d|f |d	d
 W d    n1 s0    Y  z|d|f |d	d
 W n ty   | d Y n0 d S )Nc                 S   s   || | d< d S rA   r    rM   r    r    r%   f10   s    z3TestFastMathOption.test_divf_exception.<locals>.f10rI   T)r5   debug)rk   
   )Zdtypeg      $@g        z5Divide in fastmath should not throw ZeroDivisionError)r   r   r7   npemptyassertRaisesZeroDivisionErrorZfail)r+   rj   r9   r;   r<   ZnelemZaryr    r    r%   test_divf_exception   s    2z&TestFastMathOption.test_divf_exceptionc                    s   t jddddd   fdd}td d d tf}t j|dd	|}t ||}| d
|| | d|| | d
|| | d|| d S )Nzfloat32(float32, float32)Tr6   c                 S   s   | | S r   r    )abr    r    r%   foo   s    z@TestFastMathOption.test_device_fastmath_propagation.<locals>.fooc                    s&   t d}|| jk r" ||| |< d S )NrI   )r   Zgridsize)Zarrvalr#   rt   r    r%   bar   s    

z@TestFastMathOption.test_device_fastmath_propagation.<locals>.barrI   )r5   zdiv.approx.f32rh   zdiv.full.f32)r   r7   r   ZassertInr8   ZassertNotIn)r+   rx   r9   r;   r<   r    rw   r%    test_device_fastmath_propagation   s    
z3TestFastMathOption.test_device_fastmath_propagationN)r-   r.   r/   r@   r   rL   rO   rS   rU   rW   r   rY   r`   ra   rc   rd   rf   rg   ri   rq   unittestZexpectedFailurery   r    r    r    r%   r3      s$   		


			
r3   __main__)"typingr   Zdataclassesr   r   Znumbar   r   Znumba.cuda.compilerr   r   mathr	   r
   r   r   r   r   r   r   r   operatorr   numpyrm   Znumba.cuda.testingr   r   r   rz   r   r3   r-   mainr    r    r    r%   <module>   s   , T