U
    ,d,                     @   sX   d dl mZ d dlZd dlmZmZ d dlZd dlZG dd deZ	e
dkrTe  dS )    )cudaN)skip_on_cudasimCUDATestCasec                   @   sv   e Zd Zeeejdk ddd Ze	ddd Z
eeejdk ddd	 Zeeejdk dd
d ZdS )TestMultiGPUContext   zneed more than 1 gpusc           
   
   C   sp  t ddd }dd }d}tj|tjd}tj|tjd}t jd  |d	|f || W 5 Q R X ||| |d	|f || ||| t jd  tj|tjd}tj|tjd}|d	|f || t jd	 8 tj|tjd}tj|tjd}	|d	|f ||	 W 5 Q R X W 5 Q R X ||| |||	 tj|tjd}tj|tjd}|d	|f || ||| d S )
Nzvoid(float64[:], float64[:])c                 S   s(   t d}||jk r$| | d ||< d S N   r   Zgridsize)inpouti r   I/tmp/pip-unpacked-wheel-eu7e0c37/numba/cuda/tests/cudapy/test_multigpu.pycopy_plus_1   s    

z>TestMultiGPUContext.test_multigpu_context.<locals>.copy_plus_1c                 S   s   t j| d | d S r   )nptestingassert_equal)r   r   r   r   r   check   s    z8TestMultiGPUContext.test_multigpu_context.<locals>.check    Zdtyper   r   )r   jitr   arangeZfloat64gpus)
selfr   r   NABZA0B0A1ZB1r   r   r   test_multigpu_context	   s2    


&

z)TestMultiGPUContext.test_multigpu_contextz+Simulator does not support multiple threadsc                    s   dd t td d}d g|  fddt|D }|D ]}|  qD|D ]}|  qVD ]}t|tr||qh| 	| qhd S )Nc              
   S   sd   z|  |  }W 5 Q R X W n* tk
rF } z|||< W 5 d }~X Y nX t|tdk||< d S )N
   )copy_to_host	Exceptionr   allr   )ZgpudAresultsZridxarrer   r   r   work4   s    z4TestMultiGPUContext.test_multithreaded.<locals>.workr!   c                    s&   g | ]}t jtjj |fd qS ))targetargs)	threadingThreadr   r   current).0r   r%   r&   r)   r   r   
<listcomp>C   s     z:TestMultiGPUContext.test_multithreaded.<locals>.<listcomp>)
r   	to_devicer   r   rangestartjoin
isinstanceBaseException
assertTrue)r   Znthreadsthreadsthrr   r0   r   test_multithreaded2   s    



z&TestMultiGPUContext.test_multithreadedc              	   C   s  t jdd }tjdtjd}t jd  t |}W 5 Q R X t jd  t |}W 5 Q R X t jd  |d |d W 5 Q R X t jd  |d |d W 5 Q R X t jd  tj|	 |d  W 5 Q R X t jd  tj|	 |d  W 5 Q R X d S )	Nc                 S   s(   t d}|| jk r$| |  |7  < d S r   r	   )r'   valr   r   r   r   vector_add_scalarU   s    

z@TestMultiGPUContext.test_with_context.<locals>.vector_add_scalarr!   r   r   r   )r   r!   r   )
r   r   r   r   float32r   r2   r   r   r"   )r   r>   hostarrarr1arr2r   r   r   test_with_contextR   s    
 z%TestMultiGPUContext.test_with_contextc              	   C   s   t jd " t  }|ds(| d W 5 Q R X tjdtjd}t jd  t |}W 5 Q R X t jd  t t	|}W 5 Q R X t jd " |
| tj| | W 5 Q R X d S )Nr   r   z!Peer access between GPUs disabledr!   r   )r   r   Zcurrent_contextZcan_access_peerZskipTestr   r   r?   r2   Z
zeros_likeZcopy_to_devicer   r   r"   )r   ctxr@   rA   rB   r   r   r   test_with_context_peer_copyn   s    

z/TestMultiGPUContext.test_with_context_peer_copyN)__name__
__module____qualname__unittestZskipIflenr   r   r    r   r<   rC   rE   r   r   r   r   r      s   
(

r   __main__)Znumbar   Znumpyr   Znumba.cuda.testingr   r   r,   rI   r   rF   mainr   r   r   r   <module>   s    