o
    8ήc:                      @   s,  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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(G d)d* d*eZ)e*d+kre+  dS dS ),    N)compile_ptx)	f2i1i2i4i8u1u2u4u8)cuda)types)CUDATestCaseskip_on_cudasimskip_unless_cc_53c                 C   
   t | S N)npint8x r   K/tmp/pip-target-vg8gfxp4/lib/python/numba/cuda/tests/cudapy/test_casting.pyto_int8      
r   c                 C   r   r   )r   int16r   r   r   r   to_int16   r   r   c                 C   r   r   )r   int32r   r   r   r   to_int32   r   r   c                 C   r   r   )r   int64r   r   r   r   to_int64   r   r    c                 C   r   r   )r   uint8r   r   r   r   to_uint8   r   r"   c                 C   r   r   )r   uint16r   r   r   r   	to_uint16!   r   r$   c                 C   r   r   )r   uint32r   r   r   r   	to_uint32%   r   r&   c                 C   r   r   )r   uint64r   r   r   r   	to_uint64)   r   r(   c                 C      t | t d S )Ng      ?r   float16r   r   r   r   
to_float16-      r,   c                 C      t | t d S N   )r   float32r   r   r   r   
to_float323      r2   c                 C   r.   r/   )r   float64r   r   r   r   
to_float647   r3   r5   c                 C   r   r   )r   	complex64r   r   r   r   to_complex64;   r   r7   c                 C   r   r   )r   
complex128r   r   r   r   to_complex128?   r   r9   c                 C      t jt| dS r/   r   fp16hmulr   r+   r   r   r   r   cuda_int_literal_to_float16J   r-   r>   c                 C   r)   r/   r*   r   r   r   r    reference_int_literal_to_float16P   r3   r?   c                 C   r:   Ng      @r;   r   r   r   r   cuda_float_literal_to_float16T   r-   rA   c                 C   r)   r@   r*   r   r   r   r   "reference_float_literal_to_float16Z   r3   rB   c                   @   s   e Zd Zdd Zdd Zeddd Zdd	 Zedd
d Zdd Z	e
dd Zeddd Zeddd Zdd Zeddd Zdd ZdS )TestCastingc                    s6   t jdd|t jfdd  fdd}|S )NTdevicec                    s    | d |d< d S )Nr   r   )argres)wrapped_funcr   r   cuda_wrapper_fnb   s   z4TestCasting._create_wrapped.<locals>.cuda_wrapper_fnc                    s:   t jdd}| |d< t jdd} d || |d S )N   )dtyper   )rJ   rJ   )r   zeros)rF   argarrayresarray)rI   intypeouttyper   r   
wrapper_fnf   s
   z/TestCasting._create_wrapped.<locals>.wrapper_fn)r   jit)selfpyfuncrO   rP   rQ   r   )rI   rO   rP   rH   r   _create_wrapped_   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 ]O\}}|D ]H}| j||d7 | |||}| |d|d | |dtd | |d|d | |dtd W d    n1 sfw   Y  q#qd S )Nfromtytoty皙(@皙()r   r   r   r    r   r   r   r   r   r+   r1   r4   zipsubTestrU   assertEqualintrS   pyfuncstotysfromtysrT   rX   rW   cfuncr   r   r   test_float_to_into   s   zTestCasting.test_float_to_intz(Compilation unsupported in the simulatorc                 C   P   t tttf}d}t||D ]\}}t|tgdd\}}| d| d| qd S )N          @   TrD   z	cvt.rni.s.f16)r   r   r   r    r[   r   r   assertInrS   r`   sizesrT   sizeptx_r   r   r   test_float16_to_int_ptx}      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 ];\}}|D ]4}| j||d# | |||}| |d|d | |dtd W d    n1 sRw   Y  q#qd S )NrV   rY   )r   r   r   r    r   r!   r#   r%   r'   r+   r1   r4   r[   r\   rU   r]   r^   r_   r   r   r   test_float_to_uint   s   zTestCasting.test_float_to_uintc                 C   re   )Nrf   TrD   z	cvt.rni.urk   )r"   r$   r&   r(   r[   r   r   rl   rm   r   r   r   test_float16_to_uint_ptx   rs   z$TestCasting.test_float16_to_uint_ptxc              	   C   s   t ttf}tjtjtjf}t||D ],\}}| j|d | 	|tj
|}| |d|d W d    n1 s9w   Y  qd S )N)rX   A  )r,   r2   r5   r   r+   r1   r4   r[   r\   rU   r   r]   )rS   r`   ra   rT   rX   rc   r   r   r   test_int_to_float   s   
zTestCasting.test_int_to_floatc              	   C   sz   t tf}ttf}t||D ]-\}}| j|d | |tjtj}| 	|d|d W d    n1 s5w   Y  qd S )N)funcrv   )
r>   rA   r?   rB   r[   r\   rU   r   r+   r]   )rS   	cudafuncs	hostfuncscudafunchostfuncrc   r   r   r   test_literal_to_float16   s   z#TestCasting.test_literal_to_float16c                 C   N   t tttf}d}t||D ]\}}tt|gdd\}}| d| | qd S )Nrf   TrD   zcvt.rn.f16.s)r   r   r   r   r[   r   r,   rl   rS   rb   rn   tyro   rp   rq   r   r   r   test_int_to_float16_ptx      z#TestCasting.test_int_to_float16_ptxc                 C   r~   )Nrf   TrD   zcvt.rn.f16.u)r   r	   r
   r   r[   r   r,   rl   r   r   r   r   test_uint_to_float16_ptx   r   z$TestCasting.test_uint_to_float16_ptxc              	   C   s   t ttf}tjtjtjf}tt	|||D ]F\\}}}| j
||d1 | |||}tjj|d|d|d dd tjj|d|d|d dd W d    n1 sWw   Y  qd S )NrV   rY   r0   ga2U0*3?)rtolrZ   )r,   r2   r5   r   r+   r1   r4   	itertoolsproductr[   r\   rU   testingassert_allclose)rS   r`   tysrT   rW   rX   rc   r   r   r   test_float_to_float   s   
zTestCasting.test_float_to_floatc                 C   sL   t tf}d}t||D ]\}}t|tgdd\}}| d| d| qd S )N)f32f64TrD   zcvt.rk   )r2   r5   r[   r   r   rl   )rS   r`   	postfixesrT   postfixrp   rq   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 ]C\}}|D ]<}| j	||d+ | 
|||}tj|d||d tj|d||dd  W d    n1 sTw   Y  qqd S )NrV   gGz	@gGz	y                )r7   r9   r   r6   r8   r+   r1   r4   r[   r\   rU   r   r   r_   r   r   r   test_float_to_complex   s"   
z!TestCasting.test_float_to_complexN)__name__
__module____qualname__rU   rd   r   rr   rt   ru   rw   r   r}   r   r   r   r   r   r   r   r   r   rC   ^   s&    

	



rC   __main__),numpyr   
numba.cudar   numba.core.typesr   r   r   r   r   r   r	   r
   r   numbar   
numba.corer   numba.cuda.testingr   r   r   r   unittestr   r   r   r    r"   r$   r&   r(   r,   r2   r5   r7   r9   r>   r?   rA   rB   rC   r   mainr   r   r   r   <module>   s<    , 