U
    hâËdãg  ã                   @   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rÀe ¡  dS )é    N)ÚbooleanÚconfigÚcudaÚfloat32Úfloat64Úint32Úint64Úvoid)ÚTypingError)Úskip_on_cudasimÚunittestÚCUDATestCasec                 C   s   | | S ©N© ©ÚxÚyr   r   ú`/home/sam/Atlas/atlas_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 )Néd   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