o
    ۾i $                     @   s   d dl Zd dlm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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
ej	fZd	Zd
ZedG dd deZedkrje  dS dS )    N)
namedtuple)product)	vectorize)cudaint32float32float64)CudaAPIErrordriver)skip_on_cudasim)CUDATestCase)CF)   d   i  z&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"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   Z/home/ubuntu/.local/lib/python3.10/site-packages/numba/cuda/tests/cudapy/test_vectorize.py
vector_add.      z1TestCUDAVectorize.test_scalar.<locals>.vector_addg333333?gffffff@)r   
signaturesassertEqual)selfr   r   r   cr   r   r   test_scalar,   s   


zTestCUDAVectorize.test_scalarc                 C   sl   t tdddd }tD ]'}tjtj| j|d}t||}|||}tj	|| | 
|j| qd S )Nr   r   c                 S   r   r   r   r   r   r   r   r   9   r   z-TestCUDAVectorize.test_1d.<locals>.vector_adddtype)r   r   dtypesnparrayrandomNaddtestingassert_allcloser   r#   )r   r   tydataexpectedactualr   r   r   test_1d7   s   


zTestCUDAVectorize.test_1dc           	      C   s   t tdddd }t }tD ]3}tjtj| j|d}t	||}||||d}|
 }t||}tj|| | |j| qd S )Nr   r   c                 S   r   r   r   r   r   r   r   r   F   r   z3TestCUDAVectorize.test_1d_async.<locals>.vector_addr"   stream)r   r   r   r2   r$   r%   r&   r'   r(   	to_devicecopy_to_hostr)   r*   r+   r   r#   )	r   r   r2   r,   r-   device_datadresultr/   r.   r   r   r   test_1d_asyncD   s   

zTestCUDAVectorize.test_1d_asyncc           
      C   s   t tdddd }ttddttD ]1\}}}d| }tj||}tj	|j
|d}|| }|||}	tj||	 | |	j| qd S )	Nr   r   c                 S   r   r   r   r   r   r   r   r   Z   r   z-TestCUDAVectorize.test_nd.<locals>.vector_add   r   )   )order)r   r   r   ranger$   ordersr%   r'   astyper&   Tr*   r+   r   r#   )
r   r   ndr#   r:   shaper-   data2r.   r/   r   r   r   test_ndX   s   


zTestCUDAVectorize.test_ndc                 C   sv   t tdddd }tjdtjd}tjdtjd}|| }t|}||||d tj|| | |j	|j	 d S )Nr   r   c                 S   r   r   r   r   r   r   r   r   i   r   z5TestCUDAVectorize.test_output_arg.<locals>.vector_add
   r"   out)
r   r   r%   aranger   
empty_liker*   r+   r   r#   )r   r   ABr.   r/   r   r   r   test_output_argh   s   


z!TestCUDAVectorize.test_output_argc                 C   sh   t tdddd }tj}tD ]"}tj||d}tj|}||}tj	|| | 
||j qd S )Nr   r   c                 S   r   r   r   r   r   r   r   r   x   r   z1TestCUDAVectorize.test_reduce.<locals>.vector_addr"   )r   r   r%   r   input_sizesrF   r)   reducer*   r+   r   r#   )r   r   r#   nxr.   r/   r   r   r   test_reducew   s   


zTestCUDAVectorize.test_reducec           	      C   s   t tdddd }t }tj}tD ]*}tj||d}tj	|}t
||}|j	||d}tj|| | ||j qd S )Nr   r   c                 S   r   r   r   r   r   r   r   r      r   z7TestCUDAVectorize.test_reduce_async.<locals>.vector_addr"   r1   )r   r   r   r2   r%   r   rK   rF   r)   rL   r3   r*   r+   r   r#   )	r   r   r2   r#   rM   rN   r.   dxr/   r   r   r   test_reduce_async   s   

z#TestCUDAVectorize.test_reduce_asyncc                 C   sj   t tdddd }d}tj|tjd}t|}|| }||| }tj	|| | 
|j|j d S )Nr   r   c                 S   r   r   r   r   r   r   r   r      r   z:TestCUDAVectorize.test_manual_transfer.<locals>.vector_addrC   r"   )r   r   r%   rF   r   r   r3   r4   r*   assert_equalr   r#   r   r   rM   rN   rP   r.   r/   r   r   r   test_manual_transfer   s   


z&TestCUDAVectorize.test_manual_transferc                 C   sz   t tdddd }d}tj|tjddd}t|}||||d	 || }| }tj	
|| | |j|j d S )
Nr   r   c                 S   r   r   r   r   r   r   r   r      r   z:TestCUDAVectorize.test_ufunc_output_2d.<locals>.vector_addrC   r"         rD   )r   r   r%   rF   r   reshaper   r3   r4   r*   rR   r   r#   rS   r   r   r   test_ufunc_output_2d   s   


z&TestCUDAVectorize.test_ufunc_output_2dc                 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*   rR   asarray)r   r   r   r   rr   r   r   check_tuple_arg   s   


"z!TestCUDAVectorize.check_tuple_argc                 C   s   d}d}|  || d S )N)      ?       @      @)      @      @      @)r[   )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rN   yzr\   r]   r^   r_   r`   ra   r   r[   r   rc   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 )NrC   r"   r8   rU   )r%   rF   r   r[   )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 )Nrc   rd   r\   r]   r^   g      ?g      @g      @r_   r`   ra   g      @g      @g      @rg   rh   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 )NrC   r"   rU   r8   Points)xsys)r%   rF   r   r   r[   )r   xs1ys1xs2ys2rm   r   r   r   r   r   test_namedtuple_of_array_arg   s   
z.TestCUDAVectorize.test_namedtuple_of_array_argc                 C   s&   t ddddd }| |jd d S )Nzf8(f8)r   r   c                 S   s   | d S )NrU   r   )rN   r   r   r   bar   r   z2TestCUDAVectorize.test_name_attribute.<locals>.barru   )r   r   __name__)r   ru   r   r   r   test_name_attribute   s   

z%TestCUDAVectorize.test_name_attributec              
   C   sZ  t jddddt j}t|}dd }ttdd }ttdd }t	td| t	td| | 
td |  W d    n1 sCw   Y  | 
td tdg W d    n1 s_w   Y  z-td	gd
ddd }|| W |d urt	td| nt`|d urt	td| d S t`d S |d urt	td| nt`|d urt	td| w t`w )Nr8      @   c                  _   s
   t dd)Ni  Transfer not allowed)r	   )argskwargsr   r   r   raising_transfer   s   
zLTestCUDAVectorize.test_no_transfer_for_device_data.<locals>.raising_transfercuMemcpyHtoDcuMemcpyDtoHrz   zfloat32(float32)r   r   c                 S   s   | d S )Nr\   r   )noiser   r   r   func  r   z@TestCUDAVectorize.test_no_transfer_for_device_data.<locals>.func)r%   r'   randnr=   r   r   r3   getattrr
   setattrassertRaisesRegexr	   r4   r   r~   r   )r   r   r}   old_HtoDold_DtoHr   r   r   r    test_no_transfer_for_device_data   s:   



z2TestCUDAVectorize.test_no_transfer_for_device_dataN)rv   
__module____qualname__r(   r!   r0   r7   rB   rJ   rO   rQ   rT   rX   r[   rb   ri   rk   rl   rt   rw   r   r   r   r   r   r   %   s&    
r   __main__)numpyr%   collectionsr   	itertoolsr   numbar   r   r   r   r   numba.cuda.cudadrv.driverr	   r
   numba.cuda.testingr   r   unittestr   r$   r<   rK   r   rv   mainr   r   r   r   <module>   s,    	 u