U
    9%eg                     @   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re  dS )    N)booleanconfigcudafloat32float64int32int64void)TypingError)skip_on_cudasimunittestCUDATestCasec                 C   s   | | S N xyr   r   f/var/www/html/Darija-Ai-API/env/lib/python3.8/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   s6   |  t}|| W 5 Q R X | 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   s,   t ddd }| |td d d  d S )Nzvoid(float32[::1])c                 S   s   d S r   r   r   r   r   r   f   s    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   s   d S r   r   r%   r   r   r   r&   '   s    zSTestDispatcherSpecialization.test_no_double_specialize_no_sig_same_types.<locals>.fr'   )r   r)   r   r   r$   r    r&   Zf_specializedr   r   r   +test_no_double_specialize_no_sig_same_types$   s    
zHTestDispatcherSpecialization.test_no_double_specialize_no_sig_same_typesc                 C   s,   t ddd }| |td d d  d S )Nzvoid(int32[::1])c                 S   s   d S r   r   r%   r   r   r   r&   0   s    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   s   d S r   r   r%   r   r   r   r&   8   s    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   s   d S r   r   r%   r   r   r   r&   C   s    zBTestDispatcherSpecialization.test_specialize_cache_same.<locals>.fr   r'      )
r   r)   assertEquallenspecializationsr   r   assertIsr   assertIsNot)r    r&   Z	f_float32Zf_float32_2Z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   s   d S r   r   r   r   r   r   r&   Y   s    zPTestDispatcherSpecialization.test_specialize_cache_same_with_ordering.<locals>.fr   r'   r0   )	r   r)   r1   r2   r3   r   r   r5   r4   )r    r&   Zf_f32a_f32aZf_f32c_f32cZ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+   r-   r.   r/   r6   r7   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
complex128r1   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 )
NrC   r'   r<   r>   r@   rA   r      -   )r   r)   r   rD   rE   r   rG   r   rH   r   r   r   test_coerce_input_types_unsafe   s    z-TestDispatcher.test_coerce_input_types_unsafec              	   C   sH   t dt}tjdtjd}| t |d |dd W 5 Q R X d S )NrC   r'   r<   r>   r@   rB   )r   r)   r   rD   rE   r   r   	TypeErrorrH   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'   r<         ?r>   r   r0         zdidn't compile a new versionN)
r   r)   r   rD   rE   r   ZassertAlmostEqualr1   r2   	overloads)r    rI   r   INTZ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 ]}|  qN  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               
      sd   z2t jdt jd} d | d | d d W n, tk
r^ } z | W 5 d }~X Y nX d S )Nr'   r<   r>   r   r0   )rD   rE   r   r1   	Exceptionappend)r   r#   )errorsrV   r    r   r   wrapper   s    z)TestDispatcher.test_lock.<locals>.wrapperc                    s   g | ]}t j d qS ))target)	threadingThread).0i)rZ   r   r   
<listcomp>   s     z,TestDispatcher.test_lock.<locals>.<listcomp>   N)r   r)   rangestartjoinZassertFalse)r    threadstr   )rY   rV   r    rZ   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	rxd S | 
t&}tjdtjd}|d |d
d
 W 5 Q R X | dt|j | t|jd|j d S )Nr'   r<   r>   r0   r   rQ   rP         @      @              ?zNo matching definition)r   r)   r   rD   rE   r   rG   r   r   ENABLE_CUDASIMr   rN   rF   r   r   r   r1   r2   rS   )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))rn   r    rl   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   rn   rq   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 rs   )r	   r   r   rn   rq   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'   rp   )r   rn   r	   r   rq   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)rp   r'   r<   r>         `>r         ?     ?)r   r)   r   rD   rE   r   r   rG   r    rl   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   sr   t dddgt}| t&}tjdtjd}|d |dd W 5 Q R X | t	|j
d	 | d
t	|j
 d S )Nz (float64[::1], float32, float64)z (float64[::1], float64, float32)z(float64[::1], int64, int64)r'   r<   r>   rx   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   rN   rD   rE   r   ZassertRegexpMatchesr   r   ZassertNotIn)r    r&   rm   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 )Nro   r'   r<   r>   rP   rh   r   rQ   rp         @)r   r)   r   rD   rE   r   rG   r1   r2   rS   r   r   )r    r&   r   rl   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Zdevicec                    s    ||| d< d S r   r   r   
add_devicer   r   r&   `  s    z,TestDispatcher.add_device_usecase.<locals>.f)r   r)   r   )r    rl   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 5 Q R X t|j}| d| | d| | t|jd|j d S )N(int64, int64)(float64, float64)r'   r<   r>   r0   r   rQ   rP   rh   ri   rj   zInvalid use of typez(with parameters (complex128, complex128))r   rD   rE   r   rG   r   r   rk   r   r
   rF   r   r   r   r1   r2   rS   )r    rl   r&   r   rm   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'   r<   r>   rw   r   rx   ry   )r   rD   rE   r   r   rG   rz   r   r   r   /test_explicit_signatures_device_same_type_class  s    
z>TestDispatcher.test_explicit_signatures_device_same_type_classc                 C   sH   dddg}|  |}tjdtjd}|d |dd | |d	 d
 d S )Nz(float32, float64)z(float64, float32)r   r'   r<   r>   rP   rh   r   ri   )r   rD   rE   r   rG   rz   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'   r<   r>   rP   rh   r   rQ   r   r}   )
r   rD   rE   r   rG   r1   r2   rS   r   r   rz   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   s   dS ) Add two integers, kernel versionNr   abr   r   r   r     s    z<TestDispatcher.test_dispatcher_docstring.<locals>.add_kernelTr   c                 S   s   dS ) Add two integers, device versionNr   r   r   r   r   r     s    z<TestDispatcher.test_dispatcher_docstring.<locals>.add_devicer   r   )r   r)   r1   __doc__)r    r   r   r   r   r   test_dispatcher_docstring  s    


z(TestDispatcher.test_dispatcher_docstringN)r8   r9   r:   r   rJ   r   r   ZexpectedFailurerM   rO   rU   rg   rn   rr   rt   ru   rv   r{   r|   r~   r   r   r   r   r   r   r   r   r   r   r;   n   s:   




r;   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rtd td t   d S )	Nc                 S   s,   t d}||k r(dt| |  | |< d S Nr'   gQ	@r   gridmathsinr   nr_   r   r   r   pi_sin_array  s    
z[TestDispatcherKernelProperties.test_get_regs_per_thread_unspecialized.<locals>.pi_sin_array
   r<   r'   r   z,f32 and f64 variant thread usages are equal.z-This may warrant some investigation. Devices:)r   r)   rD   rE   r   r   r	   r   get_regs_per_threadassertIsInstanceintassertGreaterr1   argsprintdetect)
r    r   Narr_f32arr_f64sig_f32sig_f64Zregs_per_thread_f32Zregs_per_thread_f64Zregs_per_thread_allr   r   r   &test_get_regs_per_thread_unspecialized  s4    



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   s,   t d}||k r(dt| |  | |< d S r   r   r   r   r   r   r     s    
zYTestDispatcherKernelProperties.test_get_regs_per_thread_specialized.<locals>.pi_sin_arrayr   )	r   r)   r	   r   r   r   r   r   r   )r    r   Z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 r   )r   )valZto_printr   r   r   const_fmt_string  s    zYTestDispatcherKernelProperties.test_get_const_mem_unspecialized.<locals>.const_fmt_stringr>   r'   Frx      rR   )r   r)   r	   r   r   r   get_const_mem_sizer   r   assertGreaterEqualr1   r   )r    r   Zsig_i64r   Zconst_mem_size_i64Zconst_mem_size_f64Z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    r<   r'   c                    s&   t j }t d}|| | |< d S rs   )r   constZ
array_liker   )r   Cr_   Zarrr   r   const_array_use,  s    
zVTestDispatcherKernelProperties.test_get_const_mem_specialized.<locals>.const_array_use)rD   Zaranger   r	   r   r)   r   r   r   r   nbytes)r    sigr   Z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                    sF   t jj | jd}t D ]}|||< qt D ]}|| | |< q0d S Nr<   )r   sharedarrayr=   rb   )arysmjr   r   r   simple_smem;  s
    
z_TestDispatcherKernelProperties.test_get_shared_mem_per_block_unspecialized.<locals>.simple_smemr<   r>   r'   rR      )r   r)   rD   rE   r   r   r	   get_shared_mem_per_blockr   r   r1   r   )
r    r   r   r   r   r   Z
sh_mem_f32Z
sh_mem_f64Zsh_mem_f32_allZ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r8tdD ]}|||< q*t   || | |< d S )Nd   r<   r'   r   )r   r   r   r   r   rb   Zsyncthreads)r   r   r_   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   r1   )r    r   Z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 rs   )r   r   )r   r_   r   r   r   simple_maxthreadsq  s    
zfTestDispatcherKernelProperties.test_get_max_threads_per_block_unspecialized.<locals>.simple_maxthreadsr<   r>   r'   r   )r   r)   rD   rE   r   r	   Zget_max_threads_per_blockr   r   r   r1   r   )r    r   r   r   r   Zmax_threads_f32Z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                    sF   t jj | jd}t D ]}|||< qt D ]}|| | |< q0d S r   r   localr   r=   rb   r   Zlmr   r   r   r   simple_lmem  s
    
z_TestDispatcherKernelProperties.test_get_local_mem_per_thread_unspecialized.<locals>.simple_lmemr<   r>   r'   rR   r   )r   r)   rD   rE   r   r   r	   get_local_mem_per_threadr   r   r   r1   r   )	r    r   r   r   r   r   Zlocal_mem_f32Zlocal_mem_f64Z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                    sF   t jj | jd}t D ]}|||< qt D ]}|| | |< q0d S r   r   r   r   r   r   r     s
    
z]TestDispatcherKernelProperties.test_get_local_mem_per_thread_specialized.<locals>.simple_lmemrR   )r   r)   r	   r   r   r   r   r   )r    r   Zlocal_mem_per_threadr   r   r   )test_get_local_mem_per_thread_specialized  s    zHTestDispatcherKernelProperties.test_get_local_mem_per_thread_specializedN)r8   r9   r:   r   r   r   r   r   r   r   r   r   r   r   r   r   r     s   /#)'r   __main__)numpyrD   r\   Znumbar   r   r   r   r   r   r   r	   Znumba.core.errorsr
   Znumba.cuda.testingr   r   r   r   r   r   r   r;   r   r8   mainr   r   r   r   <module>   s"   ([  \ s