U
    ,‰d  ã                   @   sŒ   d dl Z d dlZd dlZd dlZd dlmZ d dlmZm	Z	m
Z
 dd„ Ze
dƒG dd„ de	ƒƒZe
dƒG d	d
„ d
e	ƒƒZedkrˆe ¡  dS )é    N)Úcuda)ÚunittestÚCUDATestCaseÚskip_on_cudasimc                    s   t  ˆ ¡‡ fdd„ƒ}|S )Nc                     s6   t  ¡ }| d¡ z| ˆ | |Ž¡W ¢S | ¡  X d S )NT)ÚasyncioZnew_event_loopZ	set_debugÚcloseZrun_until_complete)ÚargsÚkwdsZloop©Úf© úI/tmp/pip-unpacked-wheel-eu7e0c37/numba/cuda/tests/cudadrv/test_streams.pyÚrunner
   s
    
z!with_asyncio_loop.<locals>.runner)Ú	functoolsÚwraps)r   r   r   r
   r   Úwith_asyncio_loop	   s    r   z,CUDA Driver API unsupported in the simulatorc                   @   sP   e Zd Zdd„ Zedd„ ƒZedd„ ƒZedd„ ƒZed	d
„ ƒZedd„ ƒZ	dS )ÚTestCudaStreamc                 C   s8   dd„ }t  ¡ }t ¡ }| ||¡ |  | d¡¡ d S )Nc                 S   s   |  ¡  d S ©N)Úset)ÚstreamÚstatusÚeventr   r   r   Úcallback   s    z2TestCudaStream.test_add_callback.<locals>.callbackg      ð?)r   r   Ú	threadingÚEventZadd_callbackÚ
assertTrueÚwait)Úselfr   r   Zcallback_eventr   r   r   Útest_add_callback   s
    z TestCudaStream.test_add_callbackc                 Ã   s   t  ¡ }| ¡ I d H  d S r   )r   r   Ú
async_done)r   r   r   r   r   Útest_async_done    s    zTestCudaStream.test_async_donec                 ƒ   sX   t t dœ‡fdd„‰ ddddg}‡ fdd	„|D ƒ}tj|Ž I d H }ˆ t ||¡¡ d S )
N)Úvalue_inÚreturnc                 “   sh   t  ¡ }t  d¡t  d¡ }}| |d d …< t j||d}|j||d | ¡ I d H }ˆ  ||¡ | ¡ S )Né   ©r   )r   r   Zpinned_arrayZ	to_deviceZcopy_to_hostr   ÚassertEqualZmean)r!   r   Zh_srcZh_dstZd_aryZdone_result)r   r   r   Úasync_cuda_fn'   s    z9TestCudaStream.test_parallel_tasks.<locals>.async_cuda_fné   é   é   é   c                    s   g | ]}t  ˆ |ƒ¡‘qS r   )r   Zcreate_task)Ú.0Úv)r&   r   r   Ú
<listcomp>2   s     z6TestCudaStream.test_parallel_tasks.<locals>.<listcomp>)Úfloatr   Úgatherr   ÚnpZallclose)r   Z	values_inZtasksZ
values_outr   )r&   r   r   Útest_parallel_tasks%   s
    
z"TestCudaStream.test_parallel_tasksc                 ƒ   sH   t  ¡ ‰ ‡ fdd„tdƒD ƒ}tj|Ž I d H }|D ]}|  |ˆ ¡ q2d S )Nc                    s   g | ]}ˆ   ¡ ‘qS r   ©r   ©r+   Ú_r$   r   r   r-   9   s     z;TestCudaStream.test_multiple_async_done.<locals>.<listcomp>r*   )r   r   Úranger   r/   r%   )r   Údone_awsÚdoneÚdr   r$   r   Útest_multiple_async_done6   s
    z'TestCudaStream.test_multiple_async_donec                 Ã   sH   dd„ t dƒD ƒ}dd„ |D ƒ}tj|Ž I d H }|  t|ƒt|ƒ¡ d S )Nc                 S   s   g | ]}t  ¡ ‘qS r   )r   r   r3   r   r   r   r-   @   s     zLTestCudaStream.test_multiple_async_done_multiple_streams.<locals>.<listcomp>r*   c                 S   s   g | ]}|  ¡ ‘qS r   r2   )r+   r   r   r   r   r-   A   s     )r5   r   r/   ZassertSetEqualr   )r   Zstreamsr6   r7   r   r   r   Ú)test_multiple_async_done_multiple_streams>   s    z8TestCudaStream.test_multiple_async_done_multiple_streamsc                 Ã   sL   t  ¡ }| ¡ | ¡  }}| ¡  |I d H  |  | ¡ ¡ |  | ¡ ¡ d S r   )r   r   r   Úcancelr   Z	cancelledr7   )r   r   Zdone1Zdone2r   r   r   Útest_cancelled_futureG   s    
z$TestCudaStream.test_cancelled_futureN)
Ú__name__Ú
__module__Ú__qualname__r   r   r    r1   r9   r:   r<   r   r   r   r   r      s   	



r   c                   @   s   e Zd Zejedd„ ƒƒZdS )ÚTestFailingStreamc              	   Ã   sr   t  ¡ }| d¡}| d¡}t  ¡ }|jdd|d ¡  | ¡ }|  t	¡ |I d H  W 5 Q R X |  
| ¡ ¡ d S )Nz
            .version 6.5
            .target sm_30
            .address_size 64
            .visible .entry failing_kernel() { trap; }
        Úfailing_kernel)r'   r$   )r   Zcurrent_contextZcreate_module_ptxZget_functionr   Ú	configureÚ__call__r   ÚassertRaisesÚ	ExceptionZassertIsNotNoneÚ	exception)r   ÚctxÚmodulerA   r   r7   r   r   r   Útest_failed_streamZ   s    

z$TestFailingStream.test_failed_streamN)r=   r>   r?   r   Úskipr   rI   r   r   r   r   r@   Q   s   	r@   Ú__main__)r   r   r   Znumpyr0   Znumbar   Znumba.cuda.testingr   r   r   r   r   r@   r=   Úmainr   r   r   r   Ú<module>   s   ;