o
    ۾ik                     @   sn   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r5e
  dS 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 )Nstatue	newspaper)   garden)   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   `/home/ubuntu/.local/lib/python3.10/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      )npfloat64r   r   jitrecarraytestingassert_equal)selfbackyard_type
goose_type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   int32float32r   r   r   r   r    )r!   recordwith2darray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 d S )N)size)outtidinnerr   r   outer7   s   
zDTestCudaSimIssues.test_cuda_module_in_device_function.<locals>.outer
   r   )r      N)numba.cuda.tests.cudasimr/   cuda_module_in_device_functionr   r   r   zerosr*   aranger0   r   r    )r!   r/   r5   arrexpectedr   r3   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 d    n1 s?w   Y  |  d S )Nc                     sR   g } t  D ]}t|tjjsq|d | r  d|  q 	| g  d S )Nr   zBlocked kernel thread: %s)
	threading	enumerate
isinstance	simulatorkernelBlockThreadjoinis_alivefailassertListEqual)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   gridsyncthreads)xyr'   r   r   r   assign_with_syncR   s   
zFTestCudaSimIssues.test_deadlock_on_exception.<locals>.assign_with_syncr   )r   r   )r   r   )	rB   r   r   r;   emptyr   assert_array_equalassertRaises
IndexError)r!   rL   rQ   rO   rP   r   rK   r   test_deadlock_on_exceptionB   s   



z,TestCudaSimIssues.test_deadlock_on_exceptionN)__name__
__module____qualname__r&   r.   r>   r   rV   r   r   r   r   r      s    r   __main__)r?   numpyr   numbar   numba.cuda.testingr   r   numba.cuda.simulatorrB   unittestr   rW   mainr   r   r   r   <module>   s    Z