U
    hâËdU  ã                   @   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l
mZ dd„ Zdd„ Ze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_memcheck)Ú
linux_onlyc                     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   ú[/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/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   )ÚgridÚlenÚrange)r   r   ÚiÚj)ÚN_ADDITIONSr   r   r   Úf3   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   D   s   ÿ) Znumbar   r	   r
   Z
numba.corer   ÚioÚ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synchronizeÚ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    



"ÿrE   c                 C   s:   zt ƒ }d}W n   t ¡ }d}Y nX |  ||f¡ d S )NTF)rE   Ú	tracebackÚ
format_excÚput)Úresult_queueÚoutputÚsuccessr   r   r   Úchild_test_wrapper]   s    
rL   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 qˆd S )	NÚspawnr)   )ZcuMemcpyHtoD_v2_ptdsZcuLaunchKernel_ptszZcuMemcpyDtoH_v2_ptdsT)ÚfnrD   )ZcuMemcpyHtoD_v2ZcuLaunchKernelZcuMemcpyDtoH_v2FÚ
)ÚmpZget_contextÚQueueÚProcessrL   r:   r;   ÚgetZfailZsubTestZassertInZassertNotIn)
ÚselfÚctxrI   ÚprocrK   rJ   Zptds_functionsrO   Zlegacy_functionsZ	fn_at_endr   r   r   Ú	test_ptdso   s"    


zTestPTDS.test_ptdsN)Ú__name__Ú
__module__Ú__qualname__r   rX   r   r   r   r   rM   k   s   rM   Ú__main__)ÚmultiprocessingrQ   r1   rF   Znumba.cuda.testingr   r   r   r   r   Znumba.tests.supportr   rE   rL   rM   rY   Úmainr   r   r   r   Ú<module>   s   S&