
    wgk                         d dl Z d dlZd dlmZ d dlmZmZ d dlmc m	Z	 d dl
Z
 G d de      Zedk(  r e
j                          yy)    N)cuda)CUDATestCaseskip_unless_cudasimc                   :    e Zd Zd Zd Zd Z ed      d        Zy)TestCudaSimIssuesc                    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	
      }t        j                  d        }t        j
                  d|      } |d   |d          t         j                  j                  |d   d   d   d       t         j                  j                  |d   d   d   d   d       y )Nstatue	newspaper)   garden)   town)*   backyardT)alignc                     d| j                   d<   d| j                  j                  d<   | j                  j                  d   dz   | j                  j                  d<   y )Ng     F@r   g       @   g      @)r   r   r
   fs    q/home/mcse/projects/flask/flask-venv/lib/python3.12/site-packages/numba/cuda/tests/cudasim/test_cudasim_issues.pysimple_kernelz;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_accessz$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                 L   t        j                  dt         j                  fdt         j                  dfg      }t        j                  d|      }d|d   d<   t
        j                  d        } |d	   |       t         j                  j                  |d   d   |d
   d          y )Nij)r      r.   r   r   r   c                     | d   | d<   y )Nr   r    r   s    r   r   z>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_settingz'TestCudaSimIssues.test_recarray_setting!   s    HHsBHHo'*BJJ&?&A Bkk!#45As		 
	dC 


AsSVC[9r   c                 X   ddl m} |j                  t        j                  fd       }t        j                  dt
        j                        } |d   |       t        j                  |j                  t
        j                        }t
        j                  j                  ||       y)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                 >            }|| j                   k  r|| |<   y y )N)size)outtidinners     r   outerzDTestCudaSimIssues.test_cuda_module_in_device_function.<locals>.outer7   s#    'CSXX~C r   
   r   )r      N)numba.cuda.tests.cudasimr7   cuda_module_in_device_functionr   r!   r   zerosr1   aranger9   r#   r$   )r%   r7   r=   arrexpectedr<   s        @r   #test_cuda_module_in_device_functionz5TestCudaSimIssues.test_cuda_module_in_device_function-   sy     	566		 
	
 hhr*eS99SXXRXX6


#.r   zOnly works on CUDASIMc                 p     fd}t         j                  d        }t        j                  d      }t        j                  d      } |d   ||       t        j
                  j                  ||        |         j                  t              5   |d   ||       d d d         |        y # 1 sw Y   xY w)Nc                     g } t        j                         D ]]  }t        |t        j                  j
                        s(|j                  d       |j                         sJj                  d|z         _ j                  | g        y )Nr   zBlocked kernel thread: %s)
	threading	enumerate
isinstance	simulatorkernelBlockThreadjoinis_alivefailassertListEqual)blockthreadstr%   s     r   assert_no_blockthreadszLTestCudaSimIssues.test_deadlock_on_exception.<locals>.assert_no_blockthreadsD   ss    L((* ?!!Y%5%5%A%AB q	::<II9A=>?   r2r   c                     t        j                  d      }| |   ||<   t        j                          t        j                          y )Nr   )r   gridsyncthreads)xyr,   s      r   assign_with_synczFTestCudaSimIssues.test_deadlock_on_exception.<locals>.assign_with_syncR   s4    		!AQ4AaDr   r   )r   r   )r   r   )	rL   r!   r   rC   emptyr#   assert_array_equalassertRaises
IndexError)r%   rU   r[   rY   rZ   s   `    r   test_deadlock_on_exceptionz,TestCudaSimIssues.test_deadlock_on_exceptionB   s    	3 
	 
	 IIaLHHQKq!$


%%a+ z* 	)"T"1a(	) 	) 	)s   B,,B5N)__name__
__module____qualname__r*   r5   rF   r   r`   r0   r   r   r   r      s,    H*
:/* 01! 2!r   r   __main__)rI   numpyr   numbar   numba.cuda.testingr   r   numba.cuda.simulatorrL   unittestr   ra   mainr0   r   r   <module>rk      sC       @ ( ( W! W!t zHMMO r   