U
    ,d:                      @   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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_53c                 C   s
   t | S N)npint8x r   H/tmp/pip-unpacked-wheel-eu7e0c37/numba/cuda/tests/cudapy/test_casting.py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   float16r   r   r   r   
to_float16-   s    r)   c                 C   s   t | t d S N   )r   float32r   r   r   r   
to_float323   s    r-   c                 C   s   t | t d S r*   )r   float64r   r   r   r   
to_float647   s    r/   c                 C   s
   t | S r   )r   	complex64r   r   r   r   to_complex64;   s    r1   c                 C   s
   t | S r   )r   
complex128r   r   r   r   to_complex128?   s    r3   c                 C   s   t jt| dS r*   r   Zfp16Zhmulr   r(   r   r   r   r   cuda_int_literal_to_float16J   s    r5   c                 C   s   t | t d S r*   r'   r   r   r   r    reference_int_literal_to_float16P   s    r6   c                 C   s   t jt| dS Ng      @r4   r   r   r   r   cuda_float_literal_to_float16T   s    r8   c                 C   s   t | t d S r7   r'   r   r   r   r   "reference_float_literal_to_float16Z   s    r9   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Z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   )Zdtyper   )r@   r@   )r   zeros)r<   ZargarrayZresarray)r?   intypeouttyper   r   
wrapper_fnf   s
    z/TestCasting._create_wrapped.<locals>.wrapper_fn)r   Zjit)selfpyfuncrB   rC   rD   r   )r?   rB   rC   r>   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 ]\}}|D ]|}| j||dd | |||}| |d|d | |dtd | |d|d | |dtd W 5 Q R X qFq:d S )Nfromtytoty皙(@皙()r   r   r   r   r   r   r   r   r   r(   r,   r.   zipsubTestrG   assertEqualintrE   pyfuncstotysfromtysrF   rJ   rI   cfuncr   r   r   test_float_to_into   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gdd\}}| d| d| qd S )N          @   Tr;   z	cvt.rni.s.f16)r   r   r   r   rM   r   r   assertInrE   rR   sizesrF   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 ]b\}}|D ]T}| j||d< | |||}| |d|d | |dtd W 5 Q R X qFq:d S )NrH   rK   )r   r   r   r   r   r   r!   r#   r%   r(   r,   r.   rM   rN   rG   rO   rP   rQ   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gdd\}}| d| d| qd S )NrW   Tr;   z	cvt.rni.ur\   )r    r"   r$   r&   rM   r   r   r]   r^   r   r   r   test_float16_to_uint_ptx   s
    z$TestCasting.test_float16_to_uint_ptxc              
   C   sn   t ttf}tjtjtjf}t||D ]D\}}| j|d* | 	|tj
|}| |d|d W 5 Q R X q$d S )N)rJ   A  )r)   r-   r/   r   r(   r,   r.   rM   rN   rG   r   rO   )rE   rR   rS   rF   rJ   rU   r   r   r   test_int_to_float   s    
zTestCasting.test_int_to_floatc              
   C   sf   t tf}ttf}t||D ]F\}}| j|d, | |tjtj}| 	|d|d W 5 Q R X qd S )N)funcrf   )
r5   r8   r6   r9   rM   rN   rG   r   r(   rO   )rE   Z	cudafuncsZ	hostfuncsZcudafuncZhostfuncrU   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|gdd\}}| d| | qd S )NrW   Tr;   zcvt.rn.f16.s)r   r   r   r   rM   r   r)   r]   rE   rT   r_   tyr`   ra   rb   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|gdd\}}| d| | qd S )NrW   Tr;   zcvt.rn.f16.u)r   r	   r
   r   rM   r   r)   r]   rj   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 ]x\\}}}| j
||dX | |||}tjj|d|d|d dd tjj|d|d|d dd W 5 Q R X q,d S )NrH   rK   r+   ga2U0*3?)ZrtolrL   )r)   r-   r/   r   r(   r,   r.   	itertoolsproductrM   rN   rG   testingassert_allclose)rE   rR   ZtysrF   rI   rJ   rU   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gdd\}}| d| d| qd S )N)Zf32Zf64Tr;   zcvt.r\   )r-   r/   rM   r   r   r]   )rE   rR   Z	postfixesrF   Zpostfixra   rb   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 ]r\}}|D ]d}| j	||dL | 
|||}tj|d||d tj|d||dd  W 5 Q R X q:q.d S )NrH   gGz	@gGz	y                )r1   r3   r   r0   r2   r(   r,   r.   rM   rN   rG   rp   rq   rQ   r   r   r   test_float_to_complex   s    
z!TestCasting.test_float_to_complexN)__name__
__module____qualname__rG   rV   r   rc   rd   re   rg   r   ri   rl   rm   rr   rs   rt   r   r   r   r   r:   ^   s$   

	



r:   __main__),Z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   rn   Zunittestr   r   r   r   r    r"   r$   r&   r)   r-   r/   r1   r3   r5   r6   r8   r9   r:   ru   mainr   r   r   r   <module>   s8   , 
