U
    ,d                     @   s   d dl Zd dlZd dlZd dlmZmZ d dlmZmZm	Z	 dd Z
dd Ze	ded	G d
d deZedkr|e  dS )    N)unittestCUDATestCase)skip_on_cudasimskip_with_cuda_pythonskip_under_cuda_memcheckc                     s  ddl mm} m} ddlm} dd l}dd l}dd l	d|_	|
 }t|}td}|| |tj d}d}	d |jd |jjdd	||jd

|

fddt|	D fddt|	D d|  || d d d | d d d  fddfdd	fddt|	D }
|
D ]}|  qP|
D ]}|  qd  
  }t|	D ]}|j|  | q|  | S )Nr   )cudaint32void)config   znumba.cuda.cudadrv.driveri   
   i   i  )lowhighsizeZdtypec                    s   g | ]}  qS  Z	to_device.0_)r   xr   F/tmp/pip-unpacked-wheel-eu7e0c37/numba/cuda/tests/cudadrv/test_ptds.py
<listcomp>(   s     zchild_test.<locals>.<listcomp>c                    s   g | ]}  qS r   r   r   )r   rr   r   r   )   s        c                    s@    d}|t| krd S t D ]}| |  || 7  < q"d S )Nr   )Zgridlenrange)r   r   ij)N_ADDITIONSr   r   r   f2   s
    
zchild_test.<locals>.fc                    s     f |  |   d S )Nr   )n)r   n_blocks	n_threadsrsstreamxsr   r   kernel_thread?   s    z!child_test.<locals>.kernel_threadc                    s   g | ]}j  |fd qS )targetargs)Thread)r   r   )r&   	threadingr   r   r   C   s   ) Znumbar   r   r	   Z
numba.corer
   ioZnumpyr+   ZCUDA_PER_THREAD_DEFAULT_STREAMStringIOloggingStreamHandler	getLogger
addHandlersetLevelDEBUGrandomseedrandintZ
zeros_liker   Zdefault_streamZjitstartjoinZsynchronizeZtestingZassert_equalZcopy_to_hostflushgetvalue)r   r	   r
   r,   npZlogbufhandlerZcudadrv_loggerNZ	N_THREADSthreadsthreadexpectedr   r   )r   r   r   r&   r!   r"   r   r#   r$   r+   r   r%   r   
child_test	   sL    



"rA   c                 C   s:   zt  }d}W n   t }d}Y nX | ||f d S )NTF)rA   	traceback
format_excput)result_queueoutputsuccessr   r   r   child_test_wrapper\   s    
rH   zHangs cuda-memcheckz&Streams not supported on the simulatorc                   @   s   e Zd Zeddd ZdS )TestPTDSz1Function names unchanged for PTDS with NV Bindingc           
   
   C   s   t d}| }|jt|fd}|  |  | \}}|sL| | d}|D ]*}| j	|dd | 
|| W 5 Q R X qTd}|D ]4}| j	|dd | d}	| |	| W 5 Q R X qd S )	Nspawnr'   )ZcuMemcpyHtoD_v2_ptdsZcuLaunchKernel_ptszZcuMemcpyDtoH_v2_ptdsT)fnr@   )ZcuMemcpyHtoD_v2ZcuLaunchKernelZcuMemcpyDtoH_v2F
)mpZget_contextQueueProcessrH   r7   r8   getZfailZsubTestZassertInZassertNotIn)
selfctxrE   procrG   rF   Zptds_functionsrK   Zlegacy_functionsZ	fn_at_endr   r   r   	test_ptdsk   s"    


zTestPTDS.test_ptdsN)__name__
__module____qualname__r   rT   r   r   r   r   rI   h   s   rI   __main__)multiprocessingrM   r.   rB   Znumba.cuda.testingr   r   r   r   r   rA   rH   rI   rU   mainr   r   r   r   <module>   s   S&