U
    9%eU                     @   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re  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   a/var/www/html/Darija-Ai-API/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 qd 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&