o
    ig                     @   s   d dl Zd dlZd dlmZmZmZmZmZm	Z	m
Z
mZ d dlmZ d dlmZmZmZ d dlZdd Zdd Zed	G d
d deZG dd deZedG dd deZedkrbe  dS dS )    N)booleanconfigcudafloat32float64int32int64void)TypingError)skip_on_cudasimunittestCUDATestCasec                 C   s   | | S N xyr   r   d/home/ubuntu/veenaModal/venv/lib/python3.10/site-packages/numba/cuda/tests/cudapy/test_dispatcher.pyadd
   s   r   c                 C   s   || | d< d S Nr   r   rr   r   r   r   r   
add_kernel   s   r   z/Specialization not implemented in the simulatorc                   @   sD   e Zd Zdd Zdd Zdd Zdd Zd	d
 Zdd Zdd Z	dS )TestDispatcherSpecializationc                 C   sJ   |  t}|| W d    n1 sw   Y  | dt|j d S )NzDispatcher already specialized)assertRaisesRuntimeError
specializeassertInstr	exception)self
dispatchertyer   r   r   _test_no_double_specialize   s   z7TestDispatcherSpecialization._test_no_double_specializec                 C   ,   t ddd }| |td d d  d S )Nzvoid(float32[::1])c                 S      d S r   r   r   r   r   r   f      zPTestDispatcherSpecialization.test_no_double_specialize_sig_same_types.<locals>.f   r   jitr$   r   r    r(   r   r   r   (test_no_double_specialize_sig_same_types   s   
zETestDispatcherSpecialization.test_no_double_specialize_sig_same_typesc                 C   s<   t jdd }|td d d }| |td d d  d S )Nc                 S   r&   r   r   r'   r   r   r   r(   '   r)   zSTestDispatcherSpecialization.test_no_double_specialize_no_sig_same_types.<locals>.fr*   )r   r,   r   r   r$   r    r(   f_specializedr   r   r   +test_no_double_specialize_no_sig_same_types$   s   
zHTestDispatcherSpecialization.test_no_double_specialize_no_sig_same_typesc                 C   r%   )Nzvoid(int32[::1])c                 S   r&   r   r   r'   r   r   r   r(   0   r)   zPTestDispatcherSpecialization.test_no_double_specialize_sig_diff_types.<locals>.fr*   r+   r-   r   r   r   (test_no_double_specialize_sig_diff_types.   s   
zETestDispatcherSpecialization.test_no_double_specialize_sig_diff_typesc                 C   s<   t jdd }|td d d }| |td d d  d S )Nc                 S   r&   r   r   r'   r   r   r   r(   8   r)   zSTestDispatcherSpecialization.test_no_double_specialize_no_sig_diff_types.<locals>.fr*   )r   r,   r   r   r$   r   r/   r   r   r   +test_no_double_specialize_no_sig_diff_types6   s   
zHTestDispatcherSpecialization.test_no_double_specialize_no_sig_diff_typesc                 C   s   t jdd }| t|jd |td d d }| t|jd |td d d }| t|jd | || |td d d }| t|jd | 	|| d S )Nc                 S   r&   r   r   r'   r   r   r   r(   C   r)   zBTestDispatcherSpecialization.test_specialize_cache_same.<locals>.fr   r*      )
r   r,   assertEquallenspecializationsr   r   assertIsr   assertIsNot)r    r(   	f_float32f_float32_2f_int32r   r   r   test_specialize_cache_same?   s   
z7TestDispatcherSpecialization.test_specialize_cache_samec                 C   s   t jdd }| t|jd |td d  td d  }| t|jd |td d d td d d }| t|jd | || |td d d td d d }| t|jd | || d S )Nc                 S   r&   r   r   r   r   r   r   r(   Y   r)   zPTestDispatcherSpecialization.test_specialize_cache_same_with_ordering.<locals>.fr   r*   r4   )	r   r,   r5   r6   r7   r   r   r9   r8   )r    r(   f_f32a_f32af_f32c_f32cf_f32c_f32c_2r   r   r   (test_specialize_cache_same_with_orderingT   s   
  zETestDispatcherSpecialization.test_specialize_cache_same_with_orderingN)
__name__
__module____qualname__r$   r.   r1   r2   r3   r=   rA   r   r   r   r   r      s    	
	r   c                   @   s   e Zd ZdZdd Zedejdd Zeddd Z	ed	d
d Z
eddd Zdd Zdd Zdd Zdd Zdd Zdd Zeddd Zedejdd  Zd!d" Zd#d$ Zd%d& Zd'd( Zed)d*d+ Zd,d- Zd.S )/TestDispatcherz9Most tests based on those in numba.tests.test_dispatcher.c                 C   s   t t}tjdtjd}|d |dd | |d tdd |d |dd | |d tdd |d |dd	 | |d tdd	 |d |d
d | |d td
d t dt}tjdtjd}|d |dd | 	|d tdd d S )Nr*   dtyper*   r*   {   i  r   皙(@F@        F@l    F: (i4[::1], i4, i4))
r   r,   r   npzeros
complex128r5   r   r   assertPreciseEqualr    c_addr   r   r   r   test_coerce_input_typesq   s   
z&TestDispatcher.test_coerce_input_typeszSimulator ignores signaturec                 C   sH   t dt}tjdtjd}|d |dd | |d tdd	 d S )
NrM   r*   rF   rH   rJ   rK   r      -   )r   r,   r   rN   rO   r   rQ   r   rR   r   r   r   test_coerce_input_types_unsafe   s   z-TestDispatcher.test_coerce_input_types_unsafec                 C   s^   t dt}tjdtjd}| t |d |dd W d    d S 1 s(w   Y  d S )NrM   r*   rF   rH   rJ   rL   )r   r,   r   rN   rO   r   r   	TypeErrorrR   r   r   r   &test_coerce_input_types_unsafe_complex   s
   "z5TestDispatcher.test_coerce_input_types_unsafe_complexz"Simulator does not track overloadsc                 C   s   t t}tjdtjd}d}d}|d ||| | |d ||  | t|j	d |d ||| | |d ||  | t|j	d |d ||| | |d ||  | t|j	d |d |dd | |d ||  | t|j	dd	 d
S )z8Test compiling new version in an ambiguous case
        r*   rF         ?rH   r   r4         zdidn't compile a new versionN)
r   r,   r   rN   rO   r   assertAlmostEqualr5   r6   	overloads)r    rS   r   INTFLTr   r   r   test_ambiguous_new_version   s    
z)TestDispatcher.test_ambiguous_new_versionz,Simulator doesn't support concurrent kernelsc                    sj   g  t jdd  fddfddtdD }|D ]}|  q|D ]}|  q'  dS )	zz
        Test that (lazy) compiling from several threads at once doesn't
        produce errors (see issue #908).
        c                 S   s   |d | d< d S )Nr*   r   r   )r   r   r   r   r   foo   s   z%TestDispatcher.test_lock.<locals>.fooc               
      sf   zt jdt jd} d | d | d d W d S  ty2 } z | W Y d }~d S d }~ww )Nr*   rF   rH   r   r4   )rN   rO   r   r5   	Exceptionappend)r   r#   )errorsrb   r    r   r   wrapper   s   z)TestDispatcher.test_lock.<locals>.wrapperc                    s   g | ]}t j d qS ))target)	threadingThread).0i)rf   r   r   
<listcomp>   s    z,TestDispatcher.test_lock.<locals>.<listcomp>   N)r   r,   rangestartjoinassertFalse)r    threadstr   )re   rb   r    rf   r   	test_lock   s   


zTestDispatcher.test_lockc                 C   s   t |t}tjdtjd}|d |dd | |d d tjdtjd}|d |dd | |d d	 tj	r<d S | 
t}tjdtjd}|d |d
d
 W d    n1 s\w   Y  | dt|j | t|jd|j d S )Nr*   rF   rH   r4   r   r[   rZ         @      @              ?zNo matching definition)r   r,   r   rN   rO   r   rQ   r   r   ENABLE_CUDASIMr   rX   rP   r   r   r   r5   r6   r^   )r    sigsr(   r   cmr   r   r   _test_explicit_signatures   s   z(TestDispatcher._test_explicit_signaturesc                 C   s   ddg}|  | d S )N(int64[::1], int64, int64) (float64[::1], float64, float64))r{   r    ry   r   r   r    test_explicit_signatures_strings   s   z/TestDispatcher.test_explicit_signatures_stringsc                 C   s6   t d d d t t ftd d d ttfg}| | d S Nr*   )r   r   r{   r~   r   r   r   test_explicit_signatures_tuples   s   (z.TestDispatcher.test_explicit_signatures_tuplesc                 C   s:   t td d d ttt td d d ttg}| | d S r   )r	   r   r   r{   r~   r   r   r   #test_explicit_signatures_signatures  s   z2TestDispatcher.test_explicit_signatures_signaturesc                 C   s~   t d d d t t fdg}| | t d d d t t fttd d d ttg}| | tt d d d t t dg}| | d S )Nr*   r}   )r   r{   r	   r   r~   r   r   r   test_explicit_signatures_mixed  s   

z-TestDispatcher.test_explicit_signatures_mixedc                 C   s   ddg}t |t}tjdtjd}|d |tdtd | |d d tjdtjd}|d |dd | |d d	 d S )
Nz (float64[::1], float32, float32)r}   r*   rF   rH         `>r         ?     ?)r   r,   r   rN   rO   r   r   rQ   r    ry   r(   r   r   r   r   (test_explicit_signatures_same_type_class  s   z7TestDispatcher.test_explicit_signatures_same_type_classz'No overload resolution in the simulatorc                 C   s   t g dt}| t}tjdtjd}|d |dd W d    n1 s)w   Y  | t	|j
d | dt	|j
 d S )	N)z (float64[::1], float32, float64)z (float64[::1], float64, float32)z(float64[::1], int64, int64)r*   rF   rH   r   g       @a  Ambiguous overloading for <function add_kernel [^>]*> \(Array\(float64, 1, 'C', False, aligned=True\), float64, float64\):\n\(Array\(float64, 1, 'C', False, aligned=True\), float32, float64\) -> none\n\(Array\(float64, 1, 'C', False, aligned=True\), float64, float32\) -> noner   )r   r,   r   r   rX   rN   rO   r   assertRegexr   r   assertNotIn)r    r(   rz   r   r   r   r   -test_explicit_signatures_ambiguous_resolution+  s   z<TestDispatcher.test_explicit_signatures_ambiguous_resolutionz$Simulator does not use _prepare_argsc                 C   s   t dt}tjdtjd}|d |dd | |d d | t|j	d|j	 dd	g}t |t}tjdtj
d}|d |tdd | |d d
 d S )Nr|   r*   rF   rH   rZ   ru   r   r[   r}         @)r   r,   r   rN   rO   r   rQ   r5   r6   r^   r   r   )r    r(   r   ry   r   r   r   test_explicit_signatures_unsafeE  s   z.TestDispatcher.test_explicit_signatures_unsafec                    s(   t j|ddt t j fdd}|S )NTdevicec                    s    ||| d< d S r   r   r   
add_devicer   r   r(   `  s   z,TestDispatcher.add_device_usecase.<locals>.f)r   r,   r   )r    ry   r(   r   r   r   add_device_usecase[  s   z!TestDispatcher.add_device_usecasec                 C   s  ddg}|  |}tjdtjd}|d |dd | |d d tjdtjd}|d |d	d
 | |d d tjr>d S | t	}tjdtj
d}|d |dd W d    n1 s^w   Y  t|j}| d| | d| | t|jd|j d S )N(int64, int64)(float64, float64)r*   rF   rH   r4   r   r[   rZ   ru   rv   rw   zInvalid use of typez(with parameters (complex128, complex128))r   rN   rO   r   rQ   r   r   rx   r   r
   rP   r   r   r   r5   r6   r^   )r    ry   r(   r   rz   msgr   r   r   test_explicit_signatures_devicef  s$   

z.TestDispatcher.test_explicit_signatures_devicec                 C   s   ddg}|  |}tjdtjd}|d |tdtd | |d d tjdtjd}|d |dd | |d d	 d S )
Nz(float32, float32)r   r*   rF   rH   r   r   r   r   )r   rN   rO   r   r   rQ   r   r   r   r   /test_explicit_signatures_device_same_type_class  s   
z>TestDispatcher.test_explicit_signatures_device_same_type_classc                 C   sF   g d}|  |}tjdtjd}|d |dd | |d d d S )	N)z(float32, float64)z(float64, float32)r   r*   rF   rH   rZ   ru   r   rv   )r   rN   rO   r   rQ   r   r   r   r   )test_explicit_signatures_device_ambiguous  s
   
z8TestDispatcher.test_explicit_signatures_device_ambiguousz%CUDA Simulator does not force castingc                 C   s   dg}|  |}tjdtjd}|d |dd | |d d | t|jd|j dd	g}|  |}tjdtjd}|d |t	dd | |d d
 d S )Nr   r*   rF   rH   rZ   ru   r   r[   r   r   )
r   rN   rO   r   rQ   r5   r6   r^   r   r   r   r   r   r   &test_explicit_signatures_device_unsafe  s   

z5TestDispatcher.test_explicit_signatures_device_unsafec                 C   sB   t jdd }t jdddd }| d|j | d|j d S )	Nc                 S      dS ) Add two integers, kernel versionNr   abr   r   r   r         z<TestDispatcher.test_dispatcher_docstring.<locals>.add_kernelTr   c                 S   r   ) Add two integers, device versionNr   r   r   r   r   r     r   z<TestDispatcher.test_dispatcher_docstring.<locals>.add_devicer   r   )r   r,   r5   __doc__)r    r   r   r   r   r   test_dispatcher_docstring  s   


z(TestDispatcher.test_dispatcher_docstringN)rB   rC   rD   r   rT   r   r   expectedFailurerW   rY   ra   rt   r{   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   rE   n   s<    




rE   z2CUDA simulator doesn't implement kernel propertiesc                   @   sT   e Zd Zdd Zdd Zdd Zdd Zd	d
 Zdd Zdd Z	dd Z
dd ZdS )TestDispatcherKernelPropertiesc           
      C   s  t jdd }d}tj|tjd}tj|tjd}|d|f || |d|f || ttd d d t}ttd d d t}||}||}| 	|t
 | 	|t
 | |d | |d | }	| |	|j | | |	|j | ||krtd td t   d S d S )	Nc                 S   0   t d}||k rdt| |  | |< d S d S Nr*   gQ	@r   gridmathsinr   nrk   r   r   r   pi_sin_array     
z[TestDispatcherKernelProperties.test_get_regs_per_thread_unspecialized.<locals>.pi_sin_array
   rF   r*   r   z,f32 and f64 variant thread usages are equal.z-This may warrant some investigation. Devices:)r   r,   rN   rO   r   r   r	   r   get_regs_per_threadassertIsInstanceintassertGreaterr5   argsprintdetect)
r    r   Narr_f32arr_f64sig_f32sig_f64regs_per_thread_f32regs_per_thread_f64regs_per_thread_allr   r   r   &test_get_regs_per_thread_unspecialized  s6   


zETestDispatcherKernelProperties.test_get_regs_per_thread_unspecializedc                 C   sF   t ttd d d tdd }| }| |t | |d d S )Nr*   c                 S   r   r   r   r   r   r   r   r     r   zYTestDispatcherKernelProperties.test_get_regs_per_thread_specialized.<locals>.pi_sin_arrayr   )	r   r,   r	   r   r   r   r   r   r   )r    r   regs_per_threadr   r   r   $test_get_regs_per_thread_specialized  s
   
zCTestDispatcherKernelProperties.test_get_regs_per_thread_specializedc                 C   s   t jdd }|d dd |d dd ttt}ttt}||}||}| |t | |t | 	|d | 	|d | }| 
||j | | 
||j | d S )	Nc                 S   s   |rt |  d S d S r   )r   )valto_printr   r   r   const_fmt_string  s   zYTestDispatcherKernelProperties.test_get_const_mem_unspecialized.<locals>.const_fmt_stringrH   r*   Fr      r\   )r   r,   r	   r   r   r   get_const_mem_sizer   r   assertGreaterEqualr5   r   )r    r   sig_i64r   const_mem_size_i64const_mem_size_f64const_mem_size_allr   r   r    test_get_const_mem_unspecialized  s   




z?TestDispatcherKernelProperties.test_get_const_mem_unspecializedc                    s`   t jdt jd ttd d d }t| fdd}||}| |t | 	| j
 d S )N    rF   r*   c                    s&   t j }t d}|| | |< d S r   )r   const
array_liker   )r   Crk   arrr   r   const_array_use,  s   
zVTestDispatcherKernelProperties.test_get_const_mem_specialized.<locals>.const_array_use)rN   aranger   r	   r   r,   r   r   r   r   nbytes)r    sigr   const_mem_sizer   r   r   test_get_const_mem_specialized(  s   
z=TestDispatcherKernelProperties.test_get_const_mem_specializedc           
         s   d t j fdd}tj tjd}tj tjd}|d | |d | ttd d d }ttd d d }||}||}| |t	 | |t	 | 
| d  | 
| d  | }| }	| 
||j | | 
|	|j | d S )	Nr   c                    F   t jj | jd}t D ]}|||< qt D ]}|| | |< qd S NrF   )r   sharedarrayrG   rn   )arysmjr   r   r   simple_smem;     
z_TestDispatcherKernelProperties.test_get_shared_mem_per_block_unspecialized.<locals>.simple_smemrF   rH   r*   r\      )r   r,   rN   rO   r   r   r	   get_shared_mem_per_blockr   r   r5   r   )
r    r   r   r   r   r   
sh_mem_f32
sh_mem_f64sh_mem_f32_allsh_mem_f64_allr   r   r   +test_get_shared_mem_per_block_unspecialized6  s&   	

zJTestDispatcherKernelProperties.test_get_shared_mem_per_block_unspecializedc                 C   sD   t ttd d d dd }| }| |t | |d d S )Nr*   c                 S   sP   t jjdtd}t d}|dkrtdD ]}|||< qt   || | |< d S )Nd   rF   r*   r   )r   r   r   r   r   rn   syncthreads)r   r   rk   r   r   r   r   r   `  s   

z]TestDispatcherKernelProperties.test_get_shared_mem_per_block_specialized.<locals>.simple_smemi  )r   r,   r	   r   r   r   r   r5   )r    r   shared_mem_per_blockr   r   r   )test_get_shared_mem_per_block_specialized_  s
   
	zHTestDispatcherKernelProperties.test_get_shared_mem_per_block_specializedc                 C   s   d}t jdd }tj|tjd}|d | ttd d d }||}| |t | 	|d | }| 
||j | d S )Nr   c                 S   s   t d}|| |< d S r   )r   r   )r   rk   r   r   r   simple_maxthreadsq  s   
zfTestDispatcherKernelProperties.test_get_max_threads_per_block_unspecialized.<locals>.simple_maxthreadsrF   rH   r*   r   )r   r,   rN   rO   r   r	   get_max_threads_per_blockr   r   r   r5   r   )r    r   r   r   r   max_threads_f32max_threads_f32_allr   r   r   ,test_get_max_threads_per_block_unspecializedn  s   

zKTestDispatcherKernelProperties.test_get_max_threads_per_block_unspecializedc           	         s   d t j fdd}tj tjd}tj tjd}|d | |d | ttd d d }ttd d d }||}||}| |t	 | |t	 | 
| d  | 
| d  | }| ||j | | ||j | d S )	N  c                    r   r   r   localr   rG   rn   r   lmr   r   r   r   simple_lmem  r   z_TestDispatcherKernelProperties.test_get_local_mem_per_thread_unspecialized.<locals>.simple_lmemrF   rH   r*   r\   r   )r   r,   rN   rO   r   r   r	   get_local_mem_per_threadr   r   r   r5   r   )	r    r   r   r   r   r   local_mem_f32local_mem_f64local_mem_allr   r   r   +test_get_local_mem_per_thread_unspecialized  s$   	

zJTestDispatcherKernelProperties.test_get_local_mem_per_thread_unspecializedc                    sP   d t ttd d d  fdd}| }| |t | | d  d S )Nr   r*   c                    r   r   r   r   r   r   r   r     r   z]TestDispatcherKernelProperties.test_get_local_mem_per_thread_specialized.<locals>.simple_lmemr\   )r   r,   r	   r   r   r   r   r   )r    r   local_mem_per_threadr   r   r   )test_get_local_mem_per_thread_specialized  s   zHTestDispatcherKernelProperties.test_get_local_mem_per_thread_specializedN)rB   rC   rD   r   r   r   r   r   r   r   r   r   r   r   r   r   r     s    /#)'r   __main__)numpyrN   rh   numbar   r   r   r   r   r   r   r	   numba.core.errorsr
   numba.cuda.testingr   r   r   r   r   r   r   rE   r   rB   mainr   r   r   r   <module>   s&    ([  \ s