o
    8ήcp(                     @   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rDe  dS dS )    N)cudafloat32float64int32int64void)skip_on_cudasimunittestCUDATestCasec                 C   s   | | S N xyr   r   N/tmp/pip-target-vg8gfxp4/lib/python/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   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   z)TestDispatcher._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      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   r!   r   r   r"   r   r   r   r#   &   r$   zETestDispatcher.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   
z:TestDispatcher.test_no_double_specialize_no_sig_same_typesc                 C   r    )Nzvoid(int32[::1])c                 S   r!   r   r   r"   r   r   r   r#   /   r$   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   r!   r   r   r"   r   r   r   r#   7   r$   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   r!   r   r   r"   r   r   r   r#   B   r$   z4TestDispatcher.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   
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   r!   r   r   r   r   r   r   r#   X   r$   zBTestDispatcher.test_specialize_cache_same_with_ordering.<locals>.fr   r%   r/   )	r   r'   r0   r1   r2   r   r   r4   r3   )r   r#   f_f32a_f32af_f32c_f32c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%   dtyper%   r%   {   i  r   皙(@F@        F@l    F: (i4[::1], i4, i4))
r   r'   r   npzeros
complex128r0   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 )
NrD   r%   r=   r?   rA   rB   r      -   )r   r'   r   rE   rF   r   rH   r   rI   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 )NrD   r%   r=   r?   rA   rC   )r   r'   r   rE   rF   r   r   	TypeErrorrI   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%   r=   g      ?r?   r   r/         zdidn't compile a new versionN)
r   r'   r   rE   rF   r   assertAlmostEqualr0   r1   	overloads)r   rJ   r   INT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 ]}|  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%   r=   r?   r   r/   )rE   rF   r   r0   	Exceptionappend)r   r   )errorsrX   r   r   r   wrapper   s   z)TestDispatcher.test_lock.<locals>.wrapperc                    s   g | ]}t j d qS ))target)	threadingThread).0i)r\   r   r   
<listcomp>   s    z,TestDispatcher.test_lock.<locals>.<listcomp>   N)r   r'   rangestartjoinassertFalse)r   threadstr   )r[   rX   r   r\   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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   nra   r   r   r   pi_sin_array      
zKTestDispatcher.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'   rE   rF   r   r   r   r   get_regs_per_threadassertIsInstanceintassertGreaterr0   argsprintdetect)
r   rs   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   


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   rk   rl   rm   rq   r   r   r   rs     rt   zITestDispatcher.test_get_regs_per_thread_specialized.<locals>.pi_sin_arrayr   )	r   r'   r   r   r   rv   rw   rx   ry   )r   rs   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      dS ) Add two integers, kernel versionNr   abr   r   r   r         z<TestDispatcher.test_dispatcher_docstring.<locals>.add_kernelT)devicec                 S   r   ) Add two integers, device versionNr   r   r   r   r   
add_device  r   z<TestDispatcher.test_dispatcher_docstring.<locals>.add_devicer   r   )r   r'   r0   __doc__)r   r   r   r   r   r   test_dispatcher_docstring  s   


z(TestDispatcher.test_dispatcher_docstringN)__name__
__module____qualname__r   r)   r,   r-   r.   r8   r<   rK   r	   expectedFailurerN   rP   rW   rj   r   r   r   r   r   r   r   r      s"    	
	
/r   __main__)numpyrE   r^   numbar   r   r   r   r   r   numba.cuda.testingr   r	   r
   ro   r   r   r   r   mainr   r   r   r   <module>   s       