U
    ?hk                     @   sj   d dl Z d dlZd dlmZ d dlmZmZ d dlm  m	Z	 d dl
Z
G dd deZedkrfe
  dS )    N)cuda)CUDATestCaseskip_unless_cudasimc                   @   s4   e Zd Zdd Zdd Zdd Zeddd	 Zd
S )TestCudaSimIssuesc                 C   s   dt jfdt jdfg}dt jdfdt jdfd|fg}t j|d	d
}tjdd }t jd|d}|d |d  t j|d d d d t j|d d d d d d S )NZstatue	newspaper)   garden)   Ztown)*   backyardT)alignc                 S   s2   d| j d< d| jjd< | jjd d | jjd< d S )Ng     F@r   g       @   g      @)r   r   r   f r   ^/var/www/html/venv/lib/python3.8/site-packages/numba/cuda/tests/cudasim/test_cudasim_issues.pysimple_kernel   s    
z;TestCudaSimIssues.test_record_access.<locals>.simple_kernel   dtyper   r   r   -   r      )npZfloat64r   r   jitrecarraytestingassert_equal)selfZbackyard_typeZ
goose_typeZgoose_np_typer   itemr   r   r   test_record_access   s    



z$TestCudaSimIssues.test_record_accessc                 C   sr   t dt jfdt jdfg}t jd|d}d|d d< tjdd	 }|d
 | t j|d d |d d  d S )Nij)r      r#   r   r   r   c                 S   s   | d | d< d S )Nr   r   r   r   r   r   r   r   '   s    z>TestCudaSimIssues.test_recarray_setting.<locals>.simple_kernelr   r   )	r   r   int32Zfloat32r   r   r   r   r   )r   Zrecordwith2darrayZrecr   r   r   r   test_recarray_setting!   s    

z'TestCudaSimIssues.test_recarray_settingc                    sd   ddl m} |j tj fdd}tjdtjd}|d | tj|j	tjd}tj
|| dS )	z
        Discovered in https://github.com/numba/numba/issues/1837.
        When the `cuda` module is referenced in a device function,
        it does not have the kernel API (e.g. cuda.threadIdx, cuda.shared)
        r   )supportc                    s     }|| j k r|| |< d S )N)size)outtidinnerr   r   outer7   s    
zDTestCudaSimIssues.test_cuda_module_in_device_function.<locals>.outer
   r   )r      N)Znumba.cuda.tests.cudasimr&   Zcuda_module_in_device_functionr   r   r   Zzerosr$   aranger'   r   r   )r   r&   r,   Zarrexpectedr   r*   r   #test_cuda_module_in_device_function-   s    z5TestCudaSimIssues.test_cuda_module_in_device_functionzOnly works on CUDASIMc              	      s~    fdd}t jdd }td}td}|d || tj|| |   t |d || W 5 Q R X |  d S )Nc                     sR   g } t  D ]4}t|tjjs q|d | r d|  q 	| g  d S )Nr   zBlocked kernel thread: %s)
	threading	enumerate
isinstance	simulatorZkernelZBlockThreadjoinis_aliveZfailZassertListEqual)Zblockthreadstr   r   r   assert_no_blockthreadsD   s    
zLTestCudaSimIssues.test_deadlock_on_exception.<locals>.assert_no_blockthreadsc                 S   s*   t d}| | ||< t   t   d S )Nr   )r   gridZsyncthreads)xyr!   r   r   r   assign_with_syncR   s    
zFTestCudaSimIssues.test_deadlock_on_exception.<locals>.assign_with_syncr   )r   r   )r   r   )	r5   r   r   r/   emptyr   Zassert_array_equalZassertRaises
IndexError)r   r:   r>   r<   r=   r   r9   r   test_deadlock_on_exceptionB   s    


z,TestCudaSimIssues.test_deadlock_on_exceptionN)__name__
__module____qualname__r    r%   r1   r   rA   r   r   r   r   r      s
   r   __main__)r2   numpyr   Znumbar   Znumba.cuda.testingr   r   Znumba.cuda.simulatorr5   Zunittestr   rB   mainr   r   r   r   <module>   s   Z