o
    8ήc                     @   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rFe  dS dS )    N)cuda)unittestCUDATestCaseskip_on_cudasimc                    s   t   fdd}|S )Nc                     s<   t  }|d z| | i |W |  S |  w )NT)asyncionew_event_loop	set_debugrun_until_completeclose)argskwdsloopf L/tmp/pip-target-vg8gfxp4/lib/python/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Eventadd_callback
assertTruewait)selfr   r   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   dt dt ffdd g d} fdd|D }tj| I d H }t|| d S )Nvalue_inreturnc                    sj   t  }t dt d}}| |d d < t j||d}|j||d | I d H } || | S )N   r   )r   r   pinned_array	to_devicecopy_to_hostr%   assertEqualmean)r'   r   h_srch_dstd_ary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   create_task).0v)r4   r   r   
<listcomp>2   s    z6TestCudaStream.test_parallel_tasks.<locals>.<listcomp>)floatr   gatherr    npallclose)r"   	values_intasks
values_outr   )r4   r"   r   test_parallel_tasks%   s   
z"TestCudaStream.test_parallel_tasksc                    sJ   t    fddtdD }tj| I d H }|D ]}| |  qd S )Nc                    s   g | ]}   qS r   r%   r:   _r*   r   r   r<   9       z;TestCudaStream.test_multiple_async_done.<locals>.<listcomp>r8   )r   r   ranger   r>   r.   )r"   done_awsdonedr   r*   r   test_multiple_async_done6   s   z'TestCudaStream.test_multiple_async_donec                    sJ   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   rF   r   r   r   r<   @   rH   zLTestCudaStream.test_multiple_async_done_multiple_streams.<locals>.<listcomp>r8   c                 S   s   g | ]}|  qS r   rE   )r:   r   r   r   r   r<   A   rH   )rI   r   r>   assertSetEqualr   )r"   streamsrJ   rK   r   r   r   )test_multiple_async_done_multiple_streams>   s
   z8TestCudaStream.test_multiple_async_done_multiple_streamsc                    sN   t  }| | }}|  |I d H  | |  | |  d S r   )r   r   r%   cancelr    	cancelledrK   )r"   r   done1done2r   r   r   test_cancelled_futureG   s   
z$TestCudaStream.test_cancelled_futureN)
__name__
__module____qualname__r$   r   r&   rD   rM   rP   rU   r   r   r   r   r      s    	



r   c                   @   s   e Zd Zejedd ZdS )TestFailingStreamc                    s   t  }|d}|d}t  }|jdd|d  | }| t	 |I d H  W d    n1 s6w   Y  | 
|  d S )Nz
            .version 6.5
            .target sm_30
            .address_size 64
            .visible .entry failing_kernel() { trap; }
        failing_kernel)r5   r*   )r   current_contextcreate_module_ptxget_functionr   	configure__call__r%   assertRaises	ExceptionassertIsNotNone	exception)r"   ctxmodulerZ   r   rK   r   r   r   test_failed_streamZ   s   

z$TestFailingStream.test_failed_streamN)rV   rW   rX   r   skipr   rf   r   r   r   r   rY   Q   s    	rY   __main__)r   r   r   numpyr?   numbar   numba.cuda.testingr   r   r   r   r   rY   rV   mainr   r   r   r   <module>   s    ;