o
    8ήc                     @   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 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 ]}|  q|
D ]}|  q  
  }t|	D ]}|j|  | q|  | S )Nr   )cudaint32void)config   znumba.cuda.cudadrv.driveri   
   i   i  )lowhighsizedtypec                       g | ]}  qS  	to_device.0_)r   xr   I/tmp/pip-target-vg8gfxp4/lib/python/numba/cuda/tests/cudadrv/test_ptds.py
<listcomp>(       zchild_test.<locals>.<listcomp>c                    r   r   r   r   )r   rr   r   r   )   r      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   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    ) numbar   r   r	   
numba.corer
   ionumpyr0   CUDA_PER_THREAD_DEFAULT_STREAMStringIOloggingStreamHandler	getLogger
addHandlersetLevelDEBUGrandomseedrandint
zeros_liker    default_streamjitstartjoinsynchronizetestingassert_equalcopy_to_hostflushgetvalue)r   r	   r
   r3   nplogbufhandlercudadrv_loggerN	N_THREADSthreadsthreadexpectedr!   r   )r#   r   r$   r+   r&   r'   r   r(   r)   r0   r   r*   r   
child_test	   sL   



"

rT   c                 C   s6   zt  }d}W n
   t }d}Y | ||f d S )NTF)rT   	traceback
format_excput)result_queueoutputsuccessr   r   r   child_test_wrapper\   s   r[   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}|  |  | \}}|s&| | d}|D ]}| j	|dd | 
|| W d    n1 sDw   Y  q*d}|D ]$}| j	|dd | d}	| |	| W d    n1 smw   Y  qNd S )	Nspawnr,   )cuMemcpyHtoD_v2_ptdscuLaunchKernel_ptszcuMemcpyDtoH_v2_ptdsT)fnrS   )cuMemcpyHtoD_v2cuLaunchKernelcuMemcpyDtoH_v2F
)mpget_contextQueueProcessr[   rC   rD   getfailsubTestassertInassertNotIn)
selfctxrX   procrZ   rY   ptds_functionsra   legacy_functions	fn_at_endr   r   r   	test_ptdsk   s,   


zTestPTDS.test_ptdsN)__name__
__module____qualname__r   ru   r   r   r   r   r\   h   s    r\   __main__)multiprocessingrf   r7   rU   numba.cuda.testingr   r   r   r   r   rT   r[   r\   rv   mainr   r   r   r   <module>   s    S&