o
    iU                     @   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rHe  dS 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 ]}|  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   _/home/ubuntu/veenaModal/venv/lib/python3.10/site-packages/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   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    ) numbar   r	   r
   
numba.corer   ionumpyr1   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   r4   nplogbufhandlercudadrv_loggerN	N_THREADSthreadsthreadexpectedr"   r   )r$   r   r%   r,   r'   r(   r   r)   r*   r1   r   r+   r   
child_test
   sL   



"

rU   c                 C   s6   zt  }d}W n
   t }d}Y | ||f d S )NTF)rU   	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)fnrT   )cuMemcpyHtoD_v2cuLaunchKernelcuMemcpyDtoH_v2F
)mpget_contextQueueProcessr\   rD   rE   getfailsubTestassertInassertNotIn)
selfctxrY   procr[   rZ   ptds_functionsrb   legacy_functions	fn_at_endr   r   r   	test_ptdso   s,   


zTestPTDS.test_ptdsN)__name__
__module____qualname__r   rv   r   r   r   r   r]   k   s    r]   __main__)multiprocessingrg   r8   rV   numba.cuda.testingr   r   r   r   r   numba.tests.supportr   rU   r\   r]   rw   mainr   r   r   r   <module>   s    S&