o
    8ήcG                     @   s   d dl Zd dlmZ d dlmZ d dlmZmZmZm	Z	 d dl
mZ d dl
mZ d dlZeeeeeee	e	e	gZejejfZedG dd	 d	eZed
krVe  dS dS )    N)
namedtuple)	vectorize)cudaint32float32float64)skip_on_cudasim)CUDATestCasez&ufunc API unsupported in the simulatorc                   @   s   e Zd Z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"S )#TestCUDAVectorizeiAB c                 C   s:   t tdddd }d}d}|||}| |||  d S )Nr   targetc                 S      | | S N abr   r   M/tmp/pip-target-vg8gfxp4/lib/python/numba/cuda/tests/cudapy/test_vectorize.py
vector_add      z1TestCUDAVectorize.test_scalar.<locals>.vector_addg333333?gffffff@)r   sigassertEqual)selfr   r   r   cr   r   r   test_scalar   s   


zTestCUDAVectorize.test_scalarc                    P   t tdddd }| tj fdd}|tj |tj |tj d S )Nr   r   c                 S   r   r   r   r   r   r   r   r   "   r   z-TestCUDAVectorize.test_1d.<locals>.vector_addc                    sH   t jt jj| d} ||}||}t ||||f d S )Ndtype)nparrayrandomN
assertTrueallclose)tydataresultgold
cuda_ufuncnp_ufuncr   r   r   test,   s   

z'TestCUDAVectorize.test_1d.<locals>.testr   r   r   adddoubler   r   r   r   r+   r   r(   r   test_1d    s   



zTestCUDAVectorize.test_1dc                    r   )Nr   r   c                 S   r   r   r   r   r   r   r   r   9   r   z3TestCUDAVectorize.test_1d_async.<locals>.vector_addc                    sp   t jt jj| d}t }t||} |||d}| }|  ||}	t 
||||f d S )Nr   stream)r   r   r    r!   r   r2   	to_devicecopy_to_hostsynchronizer"   r#   )r$   r%   r2   device_datadresultr&   r'   r(   r   r   r+   C   s   
z-TestCUDAVectorize.test_1d_async.<locals>.testr,   r/   r   r(   r   test_1d_async7   s   



zTestCUDAVectorize.test_1d_asyncc                    s\   t tdddd }| d fdd	}tdd	D ]}tD ]}d
D ]}|||| q!qqd S )Nr   r   c                 S   r   r   r   r   r   r   r   r   V   r   z-TestCUDAVectorize.test_nd.<locals>.vector_add   c                    s   t j|f| | }d|||k< d||tdk< d||tdk< t j|j|d}|| } ||}t ||| |||f d S )Ng333333@gffffff@infgffffffz-inf)order)r   r    astypefloatr   Tr"   r#   )r   r;   ndsizer%   data2r&   
our_resultr)   r   r   r   r+   \   s   

z'TestCUDAVectorize.test_nd.<locals>.test      )CF)r9   )r   r   rangetest_dtypes)r   r   r+   r?   r   r;   r   rC   r   test_ndT   s   

zTestCUDAVectorize.test_ndc                 C   s@   |  d |  d |  d | d | d | d d S )NrE   d   i  )reduce_testreduce_test2)r   r   r   r   test_ufunc_attribm   s   




z#TestCUDAVectorize.test_ufunc_attribc                 C   sf   t tdddd }tjdtjd}tjdtjd}t|}||||d | t|| | d S )Nr   r   c                 S   r   r   r   r   r   r   r   r   v   r   z5TestCUDAVectorize.test_output_arg.<locals>.vector_add
   r   out)r   r   r   aranger   
empty_liker"   r#   )r   r   ABrF   r   r   r   test_output_argu   s   


z!TestCUDAVectorize.test_output_argc                 C   sN   t tdddd }|}tj|tjd}tj|}||}| || d S )Nr   r   c                 S   r   r   r   r   r   r   r   r      r   z1TestCUDAVectorize.reduce_test.<locals>.vector_addr   )r   r   r   rR   r   r-   reducer   )r   nr   r)   xr'   r&   r   r   r   rL      s   


zTestCUDAVectorize.reduce_testc           	      C   sf   t tdddd }|}tj|tjd}tj|}t }t	||}|j||d}| 
|| d S )Nr   r   c                 S   r   r   r   r   r   r   r   r      r   z2TestCUDAVectorize.reduce_test2.<locals>.vector_addr   r1   )r   r   r   rR   r   r-   rW   r   r2   r3   r   )	r   rX   r   r)   rY   r'   r2   dxr&   r   r   r   rM      s   

zTestCUDAVectorize.reduce_test2c                 C   sZ   t tdddd }|}d}tj|tjd}t|}||| }tj	|||  d S )Nr   r   c                 S   r   r   r   r   r   r   r   r      r   z8TestCUDAVectorize.test_auto_transfer.<locals>.vector_addrO   r   )
r   r   r   rR   r   r   r3   r4   testingassert_equal)r   r   r)   rX   rY   rZ   yr   r   r   test_auto_transfer   s   


z$TestCUDAVectorize.test_auto_transferc                 C   sn   t tdddd }|}d}tj|tjddd}t|}||||d	 | }|| }tj	
|| d S )
Nr   r   c                 S   r   r   r   r   r   r   r   r      r   z=TestCUDAVectorize.test_ufunc_output_ravel.<locals>.vector_addrO   r         rP   )r   r   r   rR   r   reshaper   r3   r4   r[   r\   )r   r   r)   rX   rY   rZ   gotexpectr   r   r   test_ufunc_output_ravel   s   


z)TestCUDAVectorize.test_ufunc_output_ravelc                 C   s@   t tdddd }|||}tjt|t| | d S )Nr   r   c                 S   r   r   r   r   r   r   r   r      r   z5TestCUDAVectorize.check_tuple_arg.<locals>.vector_add)r   r   r   r[   r\   asarray)r   r   r   r   rr   r   r   check_tuple_arg   s   


"z!TestCUDAVectorize.check_tuple_argc                 C   s   d}d}|  || d S )N)      ?       @      @)      @      @      @)rg   )r   r   r   r   r   r   test_tuple_arg   s   z TestCUDAVectorize.test_tuple_argc                 C   s6   t dd}|dddd}|dddd}| || d S )	NPointrY   r]   zrh   ri   rj   rk   rl   rm   r   rg   r   ro   r   r   r   r   r   test_namedtuple_arg   s   
z%TestCUDAVectorize.test_namedtuple_argc                 C   s<   t jdt jd}||d f}|d |d f}| || d S )NrO   r   rD   r_   )r   rR   r   rg   )r   arrr   r   r   r   r   test_tuple_of_array_arg   s   z)TestCUDAVectorize.test_tuple_of_array_argc                 C   sR   t dd}|dddd|ddddf}|d	d
dd|ddddf}| || d S )Nro   rp   rh   ri   rj   g      ?g      @g      @rk   rl   rm   g      @g      @g      @rr   rs   r   r   r   test_tuple_of_namedtuple_arg   s   
z.TestCUDAVectorize.test_tuple_of_namedtuple_argc                 C   sf   t jdt jd}|d }t jdt jdd }|d }tdd}|||d}|||d}| || d S )NrO   r   r_   rD   Points)xsys)r   rR   r   r   rg   )r   xs1ys1xs2ys2rx   r   r   r   r   r   test_namedtuple_of_array_arg   s   
z.TestCUDAVectorize.test_namedtuple_of_array_argN)__name__
__module____qualname__r!   r   r0   r8   rJ   rN   rV   rL   rM   r^   rd   rg   rn   rt   rv   rw   r   r   r   r   r   r
      s$    r
   __main__)numpyr   collectionsr   numbar   r   r   r   r   numba.cuda.testingr   r	   unittestr   rI   r
   r   mainr   r   r   r   <module>   s$     P