
    J/Phg                     8   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 Zd Z ed           G d d	e                      Z G d
 de          Z ed           G d de                      Zedk    r ej                     dS dS )    N)booleanconfigcudafloat32float64int32int64void)TypingError)skip_on_cudasimunittestCUDATestCasec                     | |z   S N xys     g/var/www/html/test/jupyter/venv/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_dispatcher.pyaddr   
   s    q5L    c                     ||z   | d<   d S Nr   r   )rr   r   s      r   
add_kernelr      s    q5AaDDDr   z/Specialization not implemented in the simulatorc                   8    e Zd Zd Zd Zd Zd Zd Zd Zd Z	dS )	TestDispatcherSpecializationc                     |                      t                    5 }|                    |           d d d            n# 1 swxY w Y   |                     dt	          |j                             d S )NzDispatcher already specialized)assertRaisesRuntimeError
specializeassertInstr	exception)self
dispatchertyes       r   _test_no_double_specializez7TestDispatcherSpecialization._test_no_double_specialize   s    |,, 	&!!"%%%	& 	& 	& 	& 	& 	& 	& 	& 	& 	& 	& 	& 	& 	& 	& 	6AK8H8HIIIIIs   =AAc                     t          j        d          d             }|                     |t          d d d                    d S )Nzvoid(float32[::1])c                     d S r   r   r   s    r   fzPTestDispatcherSpecialization.test_no_double_specialize_sig_same_types.<locals>.f       Dr      r   jitr)   r   r%   r-   s     r   (test_no_double_specialize_sig_same_typeszETestDispatcherSpecialization.test_no_double_specialize_sig_same_types   sQ     
&	'	'	 	 
(	'	 	''733Q3<88888r   c                     t           j        d             }|                    t          d d d                   }|                     |t          d d d                    d S )Nc                     d S r   r   r,   s    r   r-   zSTestDispatcherSpecialization.test_no_double_specialize_no_sig_same_types.<locals>.f'   r.   r   r/   )r   r1   r!   r   r)   r%   r-   f_specializeds      r   +test_no_double_specialize_no_sig_same_typeszHTestDispatcherSpecialization.test_no_double_specialize_no_sig_same_types$   sb     
	 	 
	 WSSqS\22''wsss|DDDDDr   c                     t          j        d          d             }|                     |t          d d d                    d S )Nzvoid(int32[::1])c                     d S r   r   r,   s    r   r-   zPTestDispatcherSpecialization.test_no_double_specialize_sig_diff_types.<locals>.f0   r.   r   r/   r0   r2   s     r   (test_no_double_specialize_sig_diff_typeszETestDispatcherSpecialization.test_no_double_specialize_sig_diff_types.   sO    	$	%	%	 	 
&	%	 	''733Q3<88888r   c                     t           j        d             }|                    t          d d d                   }|                     |t
          d d d                    d S )Nc                     d S r   r   r,   s    r   r-   zSTestDispatcherSpecialization.test_no_double_specialize_no_sig_diff_types.<locals>.f8   r.   r   r/   )r   r1   r!   r   r)   r   r6   s      r   +test_no_double_specialize_no_sig_diff_typeszHTestDispatcherSpecialization.test_no_double_specialize_no_sig_diff_types6   s`    		 	 
	 U33Q3Z00''wsss|DDDDDr   c                    t           j        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                     d S r   r   r,   s    r   r-   zBTestDispatcherSpecialization.test_specialize_cache_same.<locals>.fC   r.   r   r   r/      )
r   r1   assertEquallenspecializationsr!   r   assertIsr   assertIsNot)r%   r-   	f_float32f_float32_2f_int32s        r   test_specialize_cache_samez7TestDispatcherSpecialization.test_specialize_cache_same?   s&    
	 	 
	 	Q.//333LL1..	Q.//333ll733Q3<00Q.//333i---,,uSSqSz**Q.//333),,,,,r   c                    t           j        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                     d S r   r   r   s     r   r-   zPTestDispatcherSpecialization.test_specialize_cache_same_with_ordering.<locals>.fY   r.   r   r   r/   rA   )	r   r1   rB   rC   rD   r!   r   rF   rE   )r%   r-   f_f32a_f32af_f32c_f32cf_f32c_f32c_2s        r   (test_specialize_cache_same_with_orderingzETestDispatcherSpecialization.test_specialize_cache_same_with_orderingT   sL   
 
	 	 
	 	Q.//333 ll7111:wqqqz::Q.//333 ll733Q3<1>>Q.//333k222 WSSqS\733Q3<@@Q.//333k=11111r   N)
__name__
__module____qualname__r)   r3   r8   r;   r>   rJ   rP   r   r   r   r   r      s        J J J9 9 9E E E9 9 9E E E- - -*2 2 2 2 2r   r   c                      e Zd ZdZd Z ed          ej        d                         Z ed          d             Z	 ed          d             Z
 ed          d	             Zd
 Zd Zd Zd Zd Zd Z ed          d             Z ed          ej        d                         Zd Zd Zd Zd Z ed          d             Zd ZdS )TestDispatcherz9Most tests based on those in numba.tests.test_dispatcher.c                 ^   t          j        t                    }t          j        dt          j                  } |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          j        d          t                    }t          j        dt          j                  } |d         |dd           | 	                    |d         t          dd                     d S )Nr/   dtyper/   r/   {   i  r   皙(@F@        F@l    F: (i4[::1], i4, i4))
r   r1   r   npzeros
complex128rB   r   r   assertPreciseEqualr%   c_addr   s      r   test_coerce_input_typesz&TestDispatcher.test_coerce_input_typesq   s    $$ HQbm,,,dAsC   1s3}}---dAtT"""1s4///dAtU###1s4//000dA{C(((1s;44555 .,--j99HQbh'''dAsC   !c#smm44444r   zSimulator ignores signaturec                     t          j        d          t                    }t          j        dt          j                  } |d         |dd           |                     |d         t          dd	                     d S )
Nr^   r/   rW   rY   r[   r\   r      -   )r   r1   r   r_   r`   r   rb   r   rc   s      r   test_coerce_input_types_unsafez-TestDispatcher.test_coerce_input_types_unsafe   st     .,--j99HQbh'''dAtT"""!c"bkk22222r   c                     t          j        d          t                    }t          j        dt          j                  }|                     t                    5   |d         |dd           d d d            d S # 1 swxY w Y   d S )Nr^   r/   rW   rY   r[   r]   )r   r1   r   r_   r`   r   r   	TypeErrorrc   s      r   &test_coerce_input_types_unsafe_complexz5TestDispatcher.test_coerce_input_types_unsafe_complex   s     .,--j99HQbh'''y)) 	( 	(E$K4'''	( 	( 	( 	( 	( 	( 	( 	( 	( 	( 	( 	( 	( 	( 	( 	( 	( 	(s   A>>BBz"Simulator does not track overloadsc                 R   t          j        t                    }t          j        dt          j                  }d}d} |d         |||           |                     |d         ||z              |                     t          |j	                  d            |d         |||           |                     |d         ||z              |                     t          |j	                  d            |d         |||           |                     |d         ||z              |                     t          |j	                  d            |d         |dd           |                     |d         ||z              |                     t          |j	                  dd	           d
S )z8Test compiling new version in an ambiguous case
        r/   rW         ?rY   r   rA         zdidn't compile a new versionN)
r   r1   r   r_   r`   r   assertAlmostEqualrB   rC   	overloads)r%   rd   r   INTFLTs        r   test_ambiguous_new_versionz)TestDispatcher.test_ambiguous_new_version   s    $$HQbj)))dAsC   qtS3Y///U_--q111dAsC   qtS3Y///U_--q111dAsC   qtS3Y///U_--q111 	dAq!qtS3Y///U_--q 3< 	= 	= 	= 	= 	=r   z,Simulator doesn't support concurrent kernelsc                     g t           j        d              fdfdt          d          D             }|D ]}|                                 |D ]}|                                                                 dS )zz
        Test that (lazy) compiling from several threads at once doesn't
        produce errors (see issue #908).
        c                     |dz   | d<   d S )Nr/   r   r   )r   r   s     r   fooz%TestDispatcher.test_lock.<locals>.foo   s    q5AaDDDr   c                      	 t          j        dt           j                  }  d         | d                               | d         d           d S # t          $ r }                    |           Y d }~d S d }~ww xY w)Nr/   rW   rY   r   rA   )r_   r`   r	   rB   	Exceptionappend)r   r(   errorsrx   r%   s     r   wrapperz)TestDispatcher.test_lock.<locals>.wrapper   s    !HQbh///D	!Q  1q))))) ! ! !a         !s   AA 
A=A88A=c                 :    g | ]}t          j                   S ))target)	threadingThread).0ir}   s     r   
<listcomp>z,TestDispatcher.test_lock.<locals>.<listcomp>   s'    GGG9#7333GGGr      N)r   r1   rangestartjoinassertFalse)r%   threadstr|   rx   r}   s   `  @@@r   	test_lockzTestDispatcher.test_lock   s     		 	 
		! 	! 	! 	! 	! 	! 	! HGGGU2YYGGG 	 	AGGIIII 	 	AFFHHHH     r   c                     t          j        |          t                    }t          j        dt          j                  } |d         |dd           |                     |d         d           t          j        dt          j                  } |d         |dd           |                     |d         d	           t          j	        rd S | 
                    t                    5 }t          j        dt          j                  } |d         |d
d
           d d d            n# 1 swxY w Y   |                     dt          |j                             |                     t#          |j                  d|j                   d S )Nr/   rW   rY   rA   r   ro   rn         @      @              ?zNo matching definition)r   r1   r   r_   r`   r	   rb   r   r   ENABLE_CUDASIMr   rk   ra   r"   r#   r$   rB   rC   rr   )r%   sigsr-   r   cms        r   _test_explicit_signaturesz(TestDispatcher._test_explicit_signatures   s   DHTNN:&& HQbh'''$1a!a(((HQbj)))$3!c***  	F y)) 	R"-000AAdGAr2	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	.BL0A0ABBBQ[))1ak:::::s   )4D))D-0D-c                 8    ddg}|                      |           d S )N(int64[::1], int64, int64) (float64[::1], float64, float64))r   r%   r   s     r    test_explicit_signatures_stringsz/TestDispatcher.test_explicit_signatures_strings   s)    ,24&&t,,,,,r   c                     t           d d d         t           t           ft          d d d         t          t          fg}|                     |           d S Nr/   )r	   r   r   r   s     r   test_explicit_signatures_tuplesz.TestDispatcher.test_explicit_signatures_tuples   sH    sssUE*WSSqS\7G,LM&&t,,,,,r   c                     t          t          d d d         t          t                    t          t          d d d         t          t                    g}|                     |           d S r   )r
   r	   r   r   r   s     r   #test_explicit_signatures_signaturesz2TestDispatcher.test_explicit_signatures_signatures  sU    U33Q3Z..WSSqS\7G446&&t,,,,,r   c                    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   s     r   test_explicit_signatures_mixedz-TestDispatcher.test_explicit_signatures_mixed  s     sssUE*24&&t,,, sssUE*WSSqS\7G446&&t,,, U33Q3Z..24&&t,,,,,r   c                    ddg} t          j        |          t                    }t          j        dt          j                  } |d         |t          j        d          t          j        d                     |                     |d         d           t          j        dt          j                  } |d         |dd           |                     |d         d	           d S )
Nz (float64[::1], float32, float32)r   r/   rW   rY         `>r         ?     ?)r   r1   r   r_   r`   r   r   rb   r%   r   r-   r   s       r   (test_explicit_signatures_same_type_classz7TestDispatcher.test_explicit_signatures_same_type_class  s     324DHTNN:&&HQbj)))$2:a=="*V"4"4555!c***HQbj)))$1f!&899999r   z'No overload resolution in the simulatorc                     t          j        g d          t                    }|                     t                    5 }t          j        dt
          j                  } |d         |dd           d d d            n# 1 swxY 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/   rW   rY   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   r1   r   r   rk   r_   r`   r   assertRegexr#   r$   assertNotIn)r%   r-   r   r   s       r   -test_explicit_signatures_ambiguous_resolutionz<TestDispatcher.test_explicit_signatures_ambiguous_resolution+  s   7DH 6 6 6 7 77AC C y)) 	!R"*---AAdGAsC   	! 	! 	! 	! 	! 	! 	! 	! 	! 	! 	! 	! 	! 	! 	! 	"		
 		
 		
 	#bl"3"344444s   4A??BBz$Simulator does not use _prepare_argsc                 R    t          j        d          t                    }t          j        dt          j                  } |d         |dd           |                     |d         d           |                     t          |j	                  d|j	                   dd	g} t          j        |          t                    }t          j        dt          j
                  } |d         |t          j        d          d           |                     |d         d
           d S )Nr   r/   rW   rY   rn   r   r   ro   r         @)r   r1   r   r_   r`   r	   rb   rB   rC   rr   r   r   )r%   r-   r   r   s       r   test_explicit_signatures_unsafez.TestDispatcher.test_explicit_signatures_unsafeE  s    3DH122:>>HQbh''' 	$3!a(((Q[))1ak:::,24DHTNN:&&HQbj)))$28A;;$$$!c*****r   c                 ~     t          j        |d          t                    t           j        fd            }|S )NTdevicec                 &     ||          | d<   d S r   r   )r   r   r   
add_devices      r   r-   z,TestDispatcher.add_device_usecase.<locals>.f`  s    :a##AaDDDr   )r   r1   r   )r%   r   r-   r   s      @r   add_device_usecasez!TestDispatcher.add_device_usecase[  sO     1TXd400055
		$ 	$ 	$ 	$ 
	$ r   c                 6   ddg}|                      |          }t          j        dt          j                  } |d         |dd           |                     |d         d           t          j        dt          j                  } |d         |d	d
           |                     |d         d           t          j        rd S |                     t                    5 }t          j        dt          j
                  } |d         |dd           d d d            n# 1 swxY w Y   t          |j                  }|                     d|           |                     d|           |                     t          |j                  d|j                   d S )N(int64, int64)(float64, float64)r/   rW   rY   rA   r   ro   rn   r   r   r   zInvalid use of typez(with parameters (complex128, complex128))r   r_   r`   r	   rb   r   r   r   r   r   ra   r#   r$   r"   rB   rC   rr   )r%   r   r-   r   r   msgs         r   test_explicit_signatures_devicez.TestDispatcher.test_explicit_signatures_devicef  s    !"67##D)) HQbh'''$1a!a(((HQbj)))$3!c***  	F {++ 	r"-000AAdGAr2	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 ",+S111@#FFFQ[))1ak:::::s    4D  D$'D$c                    ddg}|                      |          }t          j        dt          j                  } |d         |t          j        d          t          j        d                     |                     |d         d           t          j        dt          j                  } |d         |dd           |                     |d         d	           d S )
Nz(float32, float32)r   r/   rW   rY   r   r   r   r   )r   r_   r`   r   r   rb   r   s       r   /test_explicit_signatures_device_same_type_classz>TestDispatcher.test_explicit_signatures_device_same_type_class  s     %&:;##D))HQbj)))$2:a=="*V"4"4555!c***HQbj)))$1f!&899999r   c                     g d}|                      |          }t          j        dt          j                  } |d         |dd           |                     |d         d           d S )	N)z(float32, float64)z(float64, float32)r   r/   rW   rY   rn   r   r   r   )r   r_   r`   r   rb   r   s       r   )test_explicit_signatures_device_ambiguousz8TestDispatcher.test_explicit_signatures_device_ambiguous  sq     NMM##D))HQbj)))$3!c*****r   z%CUDA Simulator does not force castingc                 $   dg}|                      |          }t          j        dt          j                  } |d         |dd           |                     |d         d           |                     t          |j                  d|j                   dd	g}|                      |          }t          j        dt          j                  } |d         |t          j	        d          d           |                     |d         d
           d S )Nr   r/   rW   rY   rn   r   r   ro   r   r   )
r   r_   r`   r	   rb   rB   rC   rr   r   r   r   s       r   &test_explicit_signatures_device_unsafez5TestDispatcher.test_explicit_signatures_device_unsafe  s    !!##D)) HQbh'''$3!a(((Q[))1ak::: "67##D)) HQbj)))$28A;;$$$!c*****r   c                     t           j        d             }t          j        d          d             }|                     d|j                   |                     d|j                   d S )Nc                     dS ) Add two integers, kernel versionNr   abs     r   r   z<TestDispatcher.test_dispatcher_docstring.<locals>.add_kernel        r   Tr   c                     dS ) Add two integers, device versionNr   r   s     r   r   z<TestDispatcher.test_dispatcher_docstring.<locals>.add_device  r   r   r   r   )r   r1   rB   __doc__)r%   r   r   s      r   test_dispatcher_docstringz(TestDispatcher.test_dispatcher_docstring  s     
	3 	3 
	3 
				3 	3 
		3 	;Z=OPPP;Z=OPPPPPr   N)rQ   rR   rS   r   re   r   r   expectedFailureri   rl   ru   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   rU   rU   n   s       CC5 5 58 _2333 3  433 _233( ( 43( _9::= = ;:=8 _CDD! ! ED!4; ; ;.- - -- - -
- - -- - -$: : :$ _>??5 5 @?52 _;<<+ +  =<+(	 	 	; ; ;:: : :"+ + + _<==+ + >=+,Q Q Q Q Qr   rU   z2CUDA simulator doesn't implement kernel propertiesc                   D    e Zd Zd Zd Zd Zd Zd Zd Zd Z	d Z
d	 Zd
S )TestDispatcherKernelPropertiesc                    t           j        d             }d}t          j        |t          j                  }t          j        |t          j                  } |d|f         ||            |d|f         ||           t          t          d d d         t                    }t          t
          d d d         t                    }|                    |          }|                    |          }| 	                    |t                     | 	                    |t                     |                     |d           |                     |d           |                                }	|                     |	|j                 |           |                     |	|j                 |           ||k    r3t          d           t          d           t          j                     d S d S )Nc                 ~    t          j        d          }||k     r"dt          j        | |                   z  | |<   d S d S Nr/   gQ	@r   gridmathsinr   nr   s      r   pi_sin_arrayz[TestDispatcherKernelProperties.test_get_regs_per_thread_unspecialized.<locals>.pi_sin_array  =    	!A1uudhqtnn,! ur   
   rW   r/   r   z,f32 and f64 variant thread usages are equal.z-This may warrant some investigation. Devices:)r   r1   r_   r`   r   r   r
   r	   get_regs_per_threadassertIsInstanceintassertGreaterrB   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_alls
             r   &test_get_regs_per_thread_unspecializedzETestDispatcherKernelProperties.test_get_regs_per_thread_unspecialized  s    
	- 	- 
	- (1BJ///(1BJ///QT7A&&&QT7A&&& wsss|U++wsss|U++*>>wGG*>>wGG1377713777.222.222
 +>>@@,W\:,	. 	. 	.,W\:,	. 	. 	. "555 @AAAABBBKMMMMM 65r   c                    t          j        t          t          d d d         t                              d             }|                                }|                     |t                     |                     |d           d S )Nr/   c                 ~    t          j        d          }||k     r"dt          j        | |                   z  | |<   d S d S r   r   r   s      r   r   zYTestDispatcherKernelProperties.test_get_regs_per_thread_specialized.<locals>.pi_sin_array  r   r   r   )	r   r1   r
   r   r	   r   r   r   r   )r%   r   regs_per_threads      r   $test_get_regs_per_thread_specializedzCTestDispatcherKernelProperties.test_get_regs_per_thread_specialized  s    	$wsss|U++	,	,	- 	- 
-	,	- '::<<os333?A.....r   c                    t           j        d             } |d         dd            |d         dd           t          t          t                    }t          t
          t                    }|                    |          }|                    |          }|                     |t                     |                     |t                     | 	                    |d           | 	                    |d           |                                }| 
                    ||j                 |           | 
                    ||j                 |           d S )Nc                 ,    |rt          |            d S d S r   )r   )valto_prints     r   const_fmt_stringzYTestDispatcherKernelProperties.test_get_const_mem_unspecialized.<locals>.const_fmt_string  s%      c




 r   rY   r/   Fr      rp   )r   r1   r
   r	   r   r   get_const_mem_sizer   r   assertGreaterEqualrB   r   )r%   r   sig_i64r   const_mem_size_i64const_mem_size_f64const_mem_size_alls          r    test_get_const_mem_unspecializedz?TestDispatcherKernelProperties.test_get_const_mem_unspecialized  sM   		 	 
	 	q%(((sE*** ug&&w((-@@II-@@II0#6660#666 	 2A666 2A666 .@@BB+GL9;MNNN+GL9;MNNNNNr   c                 V   t          j        dt           j                  t          t          d d d                   }t	          j        |          fd            }|                    |          }|                     |t                     | 	                    |j
                   d S )N    rW   r/   c                     t           j                                      }t          j        d          }||         | |<   d S r   )r   const
array_liker   )r   Cr   arrs      r   const_array_usezVTestDispatcherKernelProperties.test_get_const_mem_specialized.<locals>.const_array_use,  s6    
%%c**A	!AQ4AaDDDr   )r_   aranger	   r
   r   r1   r   r   r   r   nbytes)r%   sigr  const_mem_sizer   s       @r   test_get_const_mem_specializedz=TestDispatcherKernelProperties.test_get_const_mem_specialized(  s    i"(+++51:	#	 	 	 	 
	
 );;C@@nc222
;;;;;r   c                 j  
 d
t           j        
fd            }t          j        
t          j                  }t          j        
t          j                  } |d         |            |d         |           t          t          d d d                   }t          t
          d d d                   }|                    |          }|                    |          }|                     |t                     |                     |t                     | 
                    |
dz             | 
                    |
dz             |                                }|                                }	| 
                    ||j                 |           | 
                    |	|j                 |           d S )Nr   c                     t           j                            | j                  }t	                    D ]}|||<   t	                    D ]}||         | |<   d S NrW   )r   sharedarrayrX   r   )arysmjr   s      r   simple_smemz_TestDispatcherKernelProperties.test_get_shared_mem_per_block_unspecialized.<locals>.simple_smem;  sm    ""1CI"66B1XX  11XX  AA r   rW   rY   r/   rp      )r   r1   r_   r`   r   r   r
   get_shared_mem_per_blockr   r   rB   r   )r%   r  r   r   r   r   
sh_mem_f32
sh_mem_f64sh_mem_f32_allsh_mem_f64_allr   s             @r   +test_get_shared_mem_per_block_unspecializedzJTestDispatcherKernelProperties.test_get_shared_mem_per_block_unspecialized6  s    
	 	 	 	 
	 (1BJ///(1BJ///D'"""D'"""wsss|$$wsss|$$ 99'BB
 99'BB
j#...j#...QU+++QU+++
 %==??$==??5zBBB5zBBBBBr   c                     t          j        t          t          d d d                             d             }|                                }|                     |t                     |                     |d           d S )Nr/   c                     t           j                            dt                    }t          j        d          }|dk    rt          d          D ]}|||<   t          j                     ||         | |<   d S )Nd   rW   r/   r   )r   r
  r  r   r   r   syncthreads)r  r  r   r  s       r   r  z]TestDispatcherKernelProperties.test_get_shared_mem_per_block_specialized.<locals>.simple_smem`  sr    ""3g"66B	!AAvvs  ABqEEUCFFFr   i  )r   r1   r
   r   r  r   r   rB   )r%   r  shared_mem_per_blocks      r   )test_get_shared_mem_per_block_specializedzHTestDispatcherKernelProperties.test_get_shared_mem_per_block_specialized_  s    	$wsss|$$	%	%	 	 
&	%	  +CCEE2C888-s33333r   c                    d}t           j        d             }t          j        |t          j                  } |d         |           t          t          d d d                   }|                    |          }|                     |t                     | 	                    |d           |                                }| 
                    ||j                 |           d S )Nr   c                 8    t          j        d          }|| |<   d S r   )r   r   )r  r   s     r   simple_maxthreadszfTestDispatcherKernelProperties.test_get_max_threads_per_block_unspecialized.<locals>.simple_maxthreadsq  s    	!ACFFFr   rW   rY   r/   r   )r   r1   r_   r`   r   r
   get_max_threads_per_blockr   r   r   rB   r   )r%   r   r  r   r   max_threads_f32max_threads_f32_alls          r   ,test_get_max_threads_per_block_unspecializedzKTestDispatcherKernelProperties.test_get_max_threads_per_block_unspecializedn  s    		 	 
	 (1BJ///$(((wsss|$$+EEgNNos333?A.../IIKK,W\:OLLLLLr   c                 B  	 d	t           j        	fd            }t          j        	t          j                  }t          j        	t          j                  } |d         |            |d         |           t          t          d d d                   }t          t
          d d d                   }|                    |          }|                    |          }|                     |t                     |                     |t                     | 
                    |	dz             | 
                    |	dz             |                                }|                     ||j                 |           |                     ||j                 |           d S )N  c                     t           j                            | j                  }t	                    D ]}|||<   t	                    D ]}||         | |<   d S r	  r   localr  rX   r   r  lmr  r   s      r   simple_lmemz_TestDispatcherKernelProperties.test_get_local_mem_per_thread_unspecialized.<locals>.simple_lmem  m    !!!39!55B1XX  11XX  AA r   rW   rY   r/   rp   r  )r   r1   r_   r`   r   r   r
   get_local_mem_per_threadr   r   r   rB   r   )
r%   r+  r   r   r   r   local_mem_f32local_mem_f64local_mem_allr   s
            @r   +test_get_local_mem_per_thread_unspecializedzJTestDispatcherKernelProperties.test_get_local_mem_per_thread_unspecialized  s    		 	 	 	 
	 (1BJ///(1BJ///D'"""D'"""wsss|$$wsss|$$#<<WEE#<<WEEmS111mS111q1u555q1u555
 $<<>>w|4mDDDw|4mDDDDDr   c                    dt          j        t          t          d d d                             fd            }|                                }|                     |t                     |                     |dz             d S )Nr%  r/   c                     t           j                            | j                  }t	                    D ]}|||<   t	                    D ]}||         | |<   d S r	  r'  r)  s      r   r+  z]TestDispatcherKernelProperties.test_get_local_mem_per_thread_specialized.<locals>.simple_lmem  r,  r   rp   )r   r1   r
   r   r-  r   r   r   )r%   r+  local_mem_per_threadr   s      @r   )test_get_local_mem_per_thread_specializedzHTestDispatcherKernelProperties.test_get_local_mem_per_thread_specialized  s     	$wsss|$$	%	%	 	 	 	 
&	%	  +CCEE2C888 4a!e<<<<<r   N)rQ   rR   rS   r   r   r   r  r  r  r#  r1  r5  r   r   r   r   r     s        - - -^
/ 
/ 
/!O !O !OF< < <'C 'C 'CR4 4 4M M M&%E %E %EN= = = = =r   r   __main__)numpyr_   r   numbar   r   r   r   r   r   r	   r
   numba.core.errorsr   numba.cuda.testingr   r   r   r   r   r   r   rU   r   rQ   mainr   r   r   <module>r<     s           M M M M M M M M M M M M M M M M M M M M ) ) ) ) ) ) F F F F F F F F F F      BCCX2 X2 X2 X2 X2< X2 X2 DCX2vWQ WQ WQ WQ WQ\ WQ WQ WQt
 EFFo= o= o= o= o=\ o= o= GFo=d zHMOOOOO r   