a
    
db"                     @   sB  d dl Zd dl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mZ d dlmZmZmZ d dlmZmZ d dlZd dl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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*d+ Z,G d,d- d-eZ-e.d.kr>e/  dS )/    N)compile_ptx)	f2i1i2i4i8u1u2u4u8)cuda)types)CUDATestCaseskip_on_cudasimskip_unless_cc_53)float16float32c                 C   s   t | S N)floatx r   m/var/www/html/stable-diffusion-webui/venv/lib/python3.9/site-packages/numba/cuda/tests/cudapy/test_casting.pynative_cast   s    r   c                 C   s
   t | S r   )npint8r   r   r   r   to_int8   s    r   c                 C   s
   t | S r   )r   int16r   r   r   r   to_int16   s    r   c                 C   s
   t | S r   )r   int32r   r   r   r   to_int32   s    r    c                 C   s
   t | S r   )r   int64r   r   r   r   to_int64   s    r"   c                 C   s
   t | S r   )r   uint8r   r   r   r   to_uint8"   s    r$   c                 C   s
   t | S r   )r   uint16r   r   r   r   	to_uint16&   s    r&   c                 C   s
   t | S r   )r   uint32r   r   r   r   	to_uint32*   s    r(   c                 C   s
   t | S r   )r   uint64r   r   r   r   	to_uint64.   s    r*   c                 C   s   t | t d S )Ng      ?r   r   r   r   r   r   
to_float162   s    r,   c                 C   s   t | t d S N   )r   r   r   r   r   r   
to_float328   s    r/   c                 C   s   t | t d S r-   )r   float64r   r   r   r   
to_float64<   s    r1   c                 C   s
   t | S r   )r   	complex64r   r   r   r   to_complex64@   s    r3   c                 C   s
   t | S r   )r   
complex128r   r   r   r   to_complex128D   s    r5   c                 C   s   t jt| dS r-   r   Zfp16Zhmulr   r   r   r   r   r   cuda_int_literal_to_float16O   s    r7   c                 C   s   t | t d S r-   r+   r   r   r   r    reference_int_literal_to_float16U   s    r8   c                 C   s   t jt| dS Ng      @r6   r   r   r   r   cuda_float_literal_to_float16Y   s    r:   c                 C   s   t | t d S r9   r+   r   r   r   r   "reference_float_literal_to_float16_   s    r;   c                   @   s   e Zd Zdd Zedd Zeddd Zedd	 Zedd
d Z	edd Z
edd Zeddd Zeddd Zedd Zeddd Zedd Zeddd ZdS )TestCastingc                    s6   t jdd|t jfdd  fdd}|S )NTZdevicec                    s    | d |d< d S )Nr   r   )argres)wrapped_funcr   r   cuda_wrapper_fng   s    z4TestCasting._create_wrapped.<locals>.cuda_wrapper_fnc                    s:   t jdd}| |d< t jdd} d || |d S )N   )Zdtyper   )rB   rB   )r   zeros)r>   ZargarrayZresarray)rA   intypeouttyper   r   
wrapper_fnk   s
    z/TestCasting._create_wrapped.<locals>.wrapper_fn)r   Zjit)selfpyfuncrD   rE   rF   r   )rA   rD   rE   r@   r   _create_wrappedd   s
    zTestCasting._create_wrappedc              
   C   s   t tttf}tjtjtjtjf}tj	tj
tjf}t||D ]\}}|D ]}| j||dn | |||}| |d|d | |dtd | |d|d | |dtd W d    qF1 s0    Y  qFq:d S )Nfromtytoty皙(@皙()r   r   r    r"   r   r   r   r   r!   r   r   r0   zipsubTestrI   assertEqualintrG   pyfuncstotysfromtysrH   rL   rK   cfuncr   r   r   test_float_to_intt   s    zTestCasting.test_float_to_intz(Compilation unsupported in the simulatorc                 C   sP   t tttf}d}t||D ]0\}}t|tfdd\}}| d| d| qd S )N          @   Tr=   z	cvt.rni.s.f16)r   r   r    r"   rO   r   r   assertInrG   rT   sizesrH   sizeptx_r   r   r   test_float16_to_int_ptx   s
    z#TestCasting.test_float16_to_int_ptxc              
   C   s   t tttf}tjtjtjtjf}tj	tj
tjf}t||D ]v\}}|D ]h}| j||dF | |||}| |d|d | |dtd W d    qF1 s0    Y  qFq:d S )NrJ   rM   )r   r   r    r"   r   r#   r%   r'   r)   r   r   r0   rO   rP   rI   rQ   rR   rS   r   r   r   test_float_to_uint   s    zTestCasting.test_float_to_uintc                 C   sP   t tttf}d}t||D ]0\}}t|tfdd\}}| d| d| qd S )NrY   Tr=   z	cvt.rni.ur^   )r$   r&   r(   r*   rO   r   r   r_   r`   r   r   r   test_float16_to_uint_ptx   s
    z$TestCasting.test_float16_to_uint_ptxc              	   C   s   t ttf}tjtjtjf}t||D ]X\}}| j|d4 | 	|tj
|}| |d|d W d    q$1 sr0    Y  q$d S )N)rL   A  )r,   r/   r1   r   r   r   r0   rO   rP   rI   r!   rQ   )rG   rT   rU   rH   rL   rW   r   r   r   test_int_to_float   s    
zTestCasting.test_int_to_floatc              	   C   sz   t tf}ttf}t||D ]Z\}}| j|d6 | |tjtj}| 	|d|d W d    q1 sj0    Y  qd S )N)funcrh   )
r7   r:   r8   r;   rO   rP   rI   r   r   rQ   )rG   Z	cudafuncsZ	hostfuncsZcudafuncZhostfuncrW   r   r   r   test_literal_to_float16   s    z#TestCasting.test_literal_to_float16c                 C   sN   t tttf}d}t||D ].\}}tt|fdd\}}| d| | qd S )NrY   Tr=   zcvt.rn.f16.s)r   r   r   r   rO   r   r,   r_   rG   rV   ra   tyrb   rc   rd   r   r   r   test_int_to_float16_ptx   s
    z#TestCasting.test_int_to_float16_ptxc                 C   sN   t tttf}d}t||D ].\}}tt|fdd\}}| d| | qd S )NrY   Tr=   zcvt.rn.f16.u)r   r	   r
   r   rO   r   r,   r_   rl   r   r   r   test_uint_to_float16_ptx   s
    z$TestCasting.test_uint_to_float16_ptxc              	   C   s   t ttf}tjtjtjf}tt	|||D ]\\}}}| j
||db | |||}tjj|d|d|d dd tjj|d|d|d dd W d    q,1 s0    Y  q,d S )NrJ   rM   r.   ga2U0*3?)ZrtolrN   )r,   r/   r1   r   r   r   r0   	itertoolsproductrO   rP   rI   testingassert_allclose)rG   rT   ZtysrH   rK   rL   rW   r   r   r   test_float_to_float   s    
zTestCasting.test_float_to_floatc                 C   sL   t tf}d}t||D ]0\}}t|tfdd\}}| d| d| qd S )N)Zf32Zf64Tr=   zcvt.r^   )r/   r1   rO   r   r   r_   )rG   rT   Z	postfixesrH   Zpostfixrc   rd   r   r   r   test_float16_to_float_ptx   s
    z%TestCasting.test_float16_to_float_ptxc              
   C   s   t tf}tjtjf}tjtjtjf}t||D ]\}}|D ]x}| j	||dV | 
|||}tj|d||d tj|d||dd  W d    q:1 s0    Y  q:q.d S )NrJ   gGz	@gGz	y                )r3   r5   r   r2   r4   r   r   r0   rO   rP   rI   rr   rs   rS   r   r   r   test_float_to_complex   s    
z!TestCasting.test_float_to_complexc                 C   sH   t jttfdd\}}| d| t jttfdd\}}| d| d S )NTr=   zst.f32zst.u16)r   r   r   r   r_   r   )rG   Zfloat32_ptxrd   Zfloat16_ptxr   r   r   test_native_cast   s    zTestCasting.test_native_castN)__name__
__module____qualname__rI   r   rX   r   re   rf   rg   ri   rk   rn   ro   rt   ru   rv   rw   r   r   r   r   r<   c   s2   




	





r<   __main__)0numpyr   Z
numba.cudar   Znumba.core.typesr   r   r   r   r   r   r	   r
   r   Znumbar   Z
numba.corer   Znumba.cuda.testingr   r   r   Znumba.typesr   r   rp   Zunittestr   r   r   r    r"   r$   r&   r(   r*   r,   r/   r1   r3   r5   r7   r8   r:   r;   r<   rx   mainr   r   r   r   <module>   s<   , 
