
    shk                         S SK r S SKrS SKJr  S SKJrJr  S SKJs  J	r	  S SK
r
 " S S\5      r\S:X  a  \
R                  " 5         gg)    N)cuda)CUDATestCaseskip_unless_cudasimc                   B    \ rS rSrS rS rS r\" S5      S 5       rSr	g)	TestCudaSimIssues   c                    S[         R                  4S[         R                  S4/nS[         R                  S4S[         R                  S4SU4/n[         R                  " US	S
9n[        R                  S 5       n[         R
                  " SUS9nUS   " US   5        [         R                  R                  US   S   S   S5        [         R                  R                  US   S   S   S   S5        g )Nstatue	newspaper)   garden)   town)*   backyardT)alignc                     SU R                   S'   SU R                  R                  S'   U R                  R                  S   S-   U R                  R                  S'   g )Ng     F@r   g       @   g      @)r   r   r   fs    ڀ/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/tests/cudasim/test_cudasim_issues.pysimple_kernel;TestCudaSimIssues.test_record_access.<locals>.simple_kernel   sK    AHHQK&)AJJ  #&'jj&:&:1&=&CAJJ  #       dtyper   r   r   -   r      )npfloat64r   r   jitrecarraytestingassert_equal)selfbackyard_type
goose_typegoose_np_typer   items         r   test_record_access$TestCudaSimIssues.test_record_access   s    "BJJ/%rzz48:  U3rzz51!=13
 48		D 
	D
 {{1M2dDG$


Q 1! 4b9


Q
 3K @ CQGr   c                 J   [         R                  " S[         R                  4S[         R                  S4/5      n[         R                  " SUS9nSUS   S'   [
        R                  S 5       nUS	   " U5        [         R                  R                  US   S   US
   S   5        g )Nij)r      r1   r   r   r   c                     U S   U S'   g )Nr   r    r   s    r   r   >TestCudaSimIssues.test_recarray_setting.<locals>.simple_kernel'   s    Q4AaDr   r   r   )	r!   r   int32float32r$   r   r#   r%   r&   )r'   recordwith2darrayrecr   s       r   test_recarray_setting'TestCudaSimIssues.test_recarray_setting!   s    HHsBHHo'*BJJ&?&A Bkk!#45As		 
	dC 


AsSVC[9r   c                 P  ^ SSK Jn  UR                  m[        R                  U4S j5       n[
        R                  " S[
        R                  S9nUS   " U5        [
        R                  " UR                  [
        R                  S9n[
        R                  R                  XC5        g)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                 <   > T" 5       nXR                   :  a  XU'   g g )N)size)outtidinners     r   outerDTestCudaSimIssues.test_cuda_module_in_device_function.<locals>.outer7   s    'CXX~C r   
   r   )r   r   N)numba.cuda.tests.cudasimr<   cuda_module_in_device_functionr   r#   r!   zerosr5   aranger>   r%   r&   )r'   r<   rB   arrexpectedrA   s        @r   #test_cuda_module_in_device_function5TestCudaSimIssues.test_cuda_module_in_device_function-   su     	566		 
	
 hhr*eS99SXXRXX6


.r   zOnly works on CUDASIMc                 z  ^  U 4S jn[         R                  S 5       n[        R                  " S5      n[        R                  " S5      nUS   " X45        [        R
                  R                  X45        U" 5         T R                  [        5         US   " X45        S S S 5        U" 5         g ! , (       d  f       N= f)Nc                  4  > / n [         R                  " 5        Hj  n[        U[        R                  R
                  5      (       d  M.  UR                  S5        UR                  5       (       d  MV  TR                  SU-  5        Ml     TR                  U / 5        g )Nr   zBlocked kernel thread: %s)
	threading	enumerate
isinstance	simulatorkernelBlockThreadjoinis_alivefailassertListEqual)blockthreadstr'   s     r   assert_no_blockthreadsLTestCudaSimIssues.test_deadlock_on_exception.<locals>.assert_no_blockthreadsD   ss    L((*!!Y%5%5%A%ABB q	::<<II9A=> +   r2r   c                     [         R                  " S5      nX   X'   [         R                  " 5         [         R                  " 5         g )Nr   )r   gridsyncthreads)xyr/   s      r   assign_with_syncFTestCudaSimIssues.test_deadlock_on_exception.<locals>.assign_with_syncR   s0    		!A4ADr   r   )r   r   )r   r   )	rR   r#   r!   rH   emptyr%   assert_array_equalassertRaises
IndexError)r'   r[   rb   r`   ra   s   `    r   test_deadlock_on_exception,TestCudaSimIssues.test_deadlock_on_exceptionB   s    	3 
	 
	 IIaLHHQKq$


%%a+ z*T"1( +  +*s   B,,
B:r3   N)
__name__
__module____qualname____firstlineno__r,   r9   rK   r   rh   __static_attributes__r3   r   r   r   r      s,    H*
:/* 01! 2!r   r   __main__)rO   numpyr!   numbar   numba.cuda.testingr   r   numba.cuda.simulatorrR   unittestr   rj   mainr3   r   r   <module>rv      sA       @ ( ( W! W!t zMMO r   