U
    ,dp(                     @   s   d dl Zd dlZd dlmZmZmZmZmZ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edkre  dS )    N)cudafloat32float64int32int64void)skip_on_cudasimunittestCUDATestCasec                 C   s   | | S N xyr   r   K/tmp/pip-unpacked-wheel-eu7e0c37/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,Dispatcher objects not used in the simulatorc                   @   s   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
ejdd Zdd Zdd Zdd Zdd Zdd Zdd ZdS ) TestDispatcherc              	   C   s6   |  t}|| W 5 Q R X | dt|j d S )NzDispatcher already specialized)assertRaisesRuntimeError
specializeZassertInstr	exception)self
dispatchertyer   r   r   _test_no_double_specialize   s    z)TestDispatcher._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    zBTestDispatcher.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    
z7TestDispatcher.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    zETestDispatcher.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    
z:TestDispatcher.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    /   s    zBTestDispatcher.test_no_double_specialize_sig_diff_types.<locals>.fr!   r"   r$   r   r   r   (test_no_double_specialize_sig_diff_types-   s    
z7TestDispatcher.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    7   s    zETestDispatcher.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_types5   s    
z:TestDispatcher.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    B   s    z4TestDispatcher.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    
z)TestDispatcher.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    X   s    zBTestDispatcher.test_specialize_cache_same_with_ordering.<locals>.fr   r!   r*   )	r   r#   r+   r,   r-   r   r   r/   r.   )r   r    Zf_f32a_f32aZf_f32c_f32cZf_f32c_f32c_2r   r   r   (test_specialize_cache_same_with_orderingS   s    
  z7TestDispatcher.test_specialize_cache_same_with_orderingc                 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!   Zdtyper!   r!   {   i  r   皙(@F@        F@l    F: (i4[::1], i4, i4))
r   r#   r   npzerosZ
complex128r+   r   r   assertPreciseEqualr   c_addr   r   r   r   test_coerce_input_typesn   s    
z&TestDispatcher.test_coerce_input_typesc                 C   sH   t dt}tjdtjd}|d |dd | |d tdd	 d S )
Nr8   r!   r2   r3   r5   r6   r      -   )r   r#   r   r9   r:   r   r;   r   r<   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 )Nr8   r!   r2   r3   r5   r7   )r   r#   r   r9   r:   r   r   	TypeErrorr<   r   r   r   &test_coerce_input_types_unsafe_complex   s    z5TestDispatcher.test_coerce_input_types_unsafe_complexc                 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!   r2   g      ?r3   r   r*         zdidn't compile a new versionN)
r   r#   r   r9   r:   r   ZassertAlmostEqualr+   r,   Z	overloads)r   r=   r   INTZFLTr   r   r   test_ambiguous_new_version   s     
z)TestDispatcher.test_ambiguous_new_versionc                    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!   r2   r3   r   r*   )r9   r:   r   r+   	Exceptionappend)r   r   )errorsrH   r   r   r   wrapper   s    z)TestDispatcher.test_lock.<locals>.wrapperc                    s   g | ]}t j d qS ))target)	threadingThread).0i)rL   r   r   
<listcomp>   s     z,TestDispatcher.test_lock.<locals>.<listcomp>   N)r   r#   rangestartjoinZassertFalse)r   threadstr   )rK   rH   r   rL   r   	test_lock   s    


zTestDispatcher.test_lockc           
      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   Zgridmathsinr   nrQ   r   r   r   pi_sin_array   s    
zKTestDispatcher.test_get_regs_per_thread_unspecialized.<locals>.pi_sin_array
   r2   r!   r   z,f32 and f64 variant thread usages are equal.z-This may warrant some investigation. Devices:)r   r#   r9   r:   r   r   r   r   get_regs_per_threadassertIsInstanceintassertGreaterr+   argsprintdetect)
r   r`   NZarr_f32Zarr_f64Zsig_f32Zsig_f64Zregs_per_thread_f32Zregs_per_thread_f64Zregs_per_thread_allr   r   r   &test_get_regs_per_thread_unspecialized   s4    



z5TestDispatcher.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 rZ   r[   r^   r   r   r   r`     s    
zITestDispatcher.test_get_regs_per_thread_specialized.<locals>.pi_sin_arrayr   )	r   r#   r   r   r   rb   rc   rd   re   )r   r`   Zregs_per_threadr   r   r   $test_get_regs_per_thread_specialized  s
    
z3TestDispatcher.test_get_regs_per_thread_specializedc                 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_kernelT)Zdevicec                 S   s   dS ) Add two integers, device versionNr   rm   r   r   r   
add_device  s    z<TestDispatcher.test_dispatcher_docstring.<locals>.add_devicerl   rp   )r   r#   r+   __doc__)r   r   rq   r   r   r   test_dispatcher_docstring  s    


z(TestDispatcher.test_dispatcher_docstringN)__name__
__module____qualname__r   r%   r'   r(   r)   r0   r1   r>   r	   ZexpectedFailurerA   rC   rG   rY   rj   rk   rs   r   r   r   r   r      s    	
	
/r   __main__)Znumpyr9   rN   Znumbar   r   r   r   r   r   Znumba.cuda.testingr   r	   r
   r\   r   r   r   rt   mainr   r   r   r   <module>   s      