
    shg                        S SK rS SKrS SKJrJrJrJrJrJ	r	J
r
Jr  S SKJr  S SKJrJrJr  S SKrS rS r\" S5       " S S	\5      5       r " S
 S\5      r\" S5       " S S\5      5       r\S:X  a  \R2                  " 5         gg)    N)booleanconfigcudafloat32float64int32int64void)TypingError)skip_on_cudasimunittestCUDATestCasec                 
    X-   $ N xys     {/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/tests/cudapy/test_dispatcher.pyaddr   
   s	    5L    c                     X-   U S'   g Nr   r   )rr   r   s      r   
add_kernelr      s    5AaDr   z/Specialization not implemented in the simulatorc                   >    \ rS rSrS rS rS rS rS rS r	S r
S	rg
)TestDispatcherSpecialization   c                     U R                  [        5       nUR                  U5        S S S 5        U R                  S[	        WR
                  5      5        g ! , (       d  f       N4= f)NzDispatcher already specialized)assertRaisesRuntimeError
specializeassertInstr	exception)self
dispatchertyes       r   _test_no_double_specialize7TestDispatcherSpecialization._test_no_double_specialize   sH    |,!!"% - 	6AKK8HI -,s   A
A$c                 t    [         R                  " S5      S 5       nU R                  U[        S S S2   5        g )Nzvoid(float32[::1])c                     g r   r   r   s    r   fPTestDispatcherSpecialization.test_no_double_specialize_sig_same_types.<locals>.f       r      r   jitr*   r   r&   r/   s     r   (test_no_double_specialize_sig_same_typesETestDispatcherSpecialization.test_no_double_specialize_sig_same_types   s:     
&	'	 
(	 	''73Q3<8r   c                     [         R                  S 5       nUR                  [        S S S2   5      nU R	                  U[        S S S2   5        g )Nc                     g r   r   r.   s    r   r/   STestDispatcherSpecialization.test_no_double_specialize_no_sig_same_types.<locals>.f'   r1   r   r2   )r   r4   r"   r   r*   r&   r/   f_specializeds      r   +test_no_double_specialize_no_sig_same_typesHTestDispatcherSpecialization.test_no_double_specialize_no_sig_same_types$   sJ     
	 
	 WSqS\2''wss|Dr   c                 t    [         R                  " S5      S 5       nU R                  U[        S S S2   5        g )Nzvoid(int32[::1])c                     g r   r   r.   s    r   r/   PTestDispatcherSpecialization.test_no_double_specialize_sig_diff_types.<locals>.f0   r1   r   r2   r3   r5   s     r   (test_no_double_specialize_sig_diff_typesETestDispatcherSpecialization.test_no_double_specialize_sig_diff_types.   s8    	$	%	 
&	 	''73Q3<8r   c                     [         R                  S 5       nUR                  [        S S S2   5      nU R	                  U[
        S S S2   5        g )Nc                     g r   r   r.   s    r   r/   STestDispatcherSpecialization.test_no_double_specialize_no_sig_diff_types.<locals>.f8   r1   r   r2   )r   r4   r"   r   r*   r   r;   s      r   +test_no_double_specialize_no_sig_diff_typesHTestDispatcherSpecialization.test_no_double_specialize_no_sig_diff_types6   sH    		 
	 U3Q3Z0''wss|Dr   c                 >   [         R                  S 5       nU R                  [        UR                  5      S5        UR                  [        S S S2   5      nU R                  [        UR                  5      S5        UR                  [        S S S2   5      nU R                  [        UR                  5      S5        U R                  X#5        UR                  [        S S S2   5      nU R                  [        UR                  5      S5        U R                  XB5        g )Nc                     g r   r   r.   s    r   r/   BTestDispatcherSpecialization.test_specialize_cache_same.<locals>.fC   r1   r   r   r2      )
r   r4   assertEquallenspecializationsr"   r   assertIsr   assertIsNot)r&   r/   	f_float32f_float32_2f_int32s        r   test_specialize_cache_same7TestDispatcherSpecialization.test_specialize_cache_same?   s     
	 
	 	Q../3LL1.	Q../3ll73Q3<0Q../3i-,,uSqSz*Q../3,r   c                 t   [         R                  S 5       nU R                  [        UR                  5      S5        UR                  [        S S  [        S S  5      nU R                  [        UR                  5      S5        UR                  [        S S S2   [        S S S2   5      nU R                  [        UR                  5      S5        U R                  X#5        UR                  [        S S S2   [        S S S2   5      nU R                  [        UR                  5      S5        U R                  X45        g )Nc                     g r   r   r   s     r   r/   PTestDispatcherSpecialization.test_specialize_cache_same_with_ordering.<locals>.fY   r1   r   r   r2   rL   )	r   r4   rM   rN   rO   r"   r   rQ   rP   )r&   r/   f_f32a_f32af_f32c_f32cf_f32c_f32c_2s        r   (test_specialize_cache_same_with_orderingETestDispatcherSpecialization.test_specialize_cache_same_with_orderingT   s    
 
	 
	 	Q../3 ll71:wqz:Q../3 ll73Q3<1>Q../32 WSqS\73Q3<@Q../3k1r   r   N)__name__
__module____qualname____firstlineno__r*   r6   r=   rB   rG   rU   r]   __static_attributes__r   r   r   r   r      s(    J9E9E-*2r   r   c                   `   \ rS rSrSrS r\" S5      \R                  S 5       5       r	\" S5      S 5       r
\" S5      S 5       r\" S	5      S
 5       rS rS rS rS rS rS r\" S5      S 5       r\" S5      \R                  S 5       5       rS rS rS rS r\" S5      S 5       rS rSrg)TestDispatchern   z9Most tests based on those in numba.tests.test_dispatcher.c                    [         R                  " [        5      n[        R                  " S[        R
                  S9nUS   " USS5        U R                  US   [        SS5      5        US   " USS5        U R                  US   [        SS5      5        US   " USS	5        U R                  US   [        SS	5      5        US   " US
S5        U R                  US   [        S
S5      5        [         R                  " S5      " [        5      n[        R                  " S[        R                  S9nUS   " USS5        U R                  US   [        SS5      5        g )Nr2   dtyper2   r2   {     r   皙(@F@        F@l    F: (i4[::1], i4, i4))
r   r4   r   npzeros
complex128rM   r   r   assertPreciseEqualr&   c_addr   s      r   test_coerce_input_types&TestDispatcher.test_coerce_input_typesq   s&    $ HHQbmm,dAsC 1s3}-dAtT"1s4/dAtU#1s4/0dA{C(1s;45 ,-j9HHQbhh'dAsC !c#sm4r   zSimulator ignores signaturec                     [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " USS5        U R                  US   [        SS	5      5        g )
Nrp   r2   rh   rj   rm   rn   r      -   )r   r4   r   rq   rr   r   rt   r   ru   s      r   test_coerce_input_types_unsafe-TestDispatcher.test_coerce_input_types_unsafe   sV     ,-j9HHQbhh'dAtT"!c"bk2r   c                    [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nU R                  [        5         US   " USS5        S S S 5        g ! , (       d  f       g = f)Nrp   r2   rh   rj   rm   ro   )r   r4   r   rq   rr   r   r    	TypeErrorru   s      r   &test_coerce_input_types_unsafe_complex5TestDispatcher.test_coerce_input_types_unsafe_complex   sV     ,-j9HHQbhh'y)$K4' *))s   A00
A>z"Simulator does not track overloadsc                    [         R                  " [        5      n[        R                  " S[        R
                  S9nSnSnUS   " X#U5        U R                  US   X4-   5        U R                  [        UR                  5      S5        US   " X$U5        U R                  US   XC-   5        U R                  [        UR                  5      S5        US   " X$U5        U R                  US   XD-   5        U R                  [        UR                  5      S5        US   " USS5        U R                  US   X3-   5        U R                  [        UR                  5      SS	5        g
)z8Test compiling new version in an ambiguous case
        r2   rh         ?rj   r   rL         zdidn't compile a new versionN)
r   r4   r   rq   rr   r   assertAlmostEqualrM   rN   	overloads)r&   rv   r   INTFLTs        r   test_ambiguous_new_version)TestDispatcher.test_ambiguous_new_version   s/    $HHQbjj)dAC qtSY/U__-q1dAC qtSY/U__-q1dAC qtSY/U__-q1 	dAq!qtSY/U__-q 3< 	=r   z,Simulator doesn't support concurrent kernelsc                 4  ^ ^^ / m[         R                  S 5       mUUU 4S jn[        S5       Vs/ s H  n[        R                  " US9PM     nnU H  nUR                  5         M     U H  nUR                  5         M     T R                  T5        gs  snf )zb
Test that (lazy) compiling from several threads at once doesn't
produce errors (see issue #908).
c                     US-   U S'   g )Nr2   r   r   )r   r   s     r   foo%TestDispatcher.test_lock.<locals>.foo   s    q5AaDr   c                     >  [         R                  " S[         R                  S9n TS   " U S5        TR                  U S   S5        g ! [         a  nTR                  U5         S nAg S nAff = f)Nr2   rh   rj   r   rL   )rq   rr   r	   rM   	Exceptionappend)r   r)   errorsr   r&   s     r   wrapper)TestDispatcher.test_lock.<locals>.wrapper   s\    !HHQbhh/D	!Q  1q) !a  !s   AA 
A-A((A-   )targetN)r   r4   range	threadingThreadstartjoinassertFalse)r&   r   ithreadstr   r   s   `    @@r   	test_lockTestDispatcher.test_lock   s     		 
		! >C2YGY9##73YGAGGI AFFH   Hs   Bc                    [         R                  " U5      " [        5      n[        R                  " S[        R
                  S9nUS   " USS5        U R                  US   S5        [        R                  " S[        R                  S9nUS   " USS5        U R                  US   S	5        [        R                  (       a  g U R                  [        5       n[        R                  " S[        R                  S9nUS   " US
S
5        S S S 5        U R                  S[        WR                  5      5        U R!                  [#        UR$                  5      SUR$                  5        g ! , (       d  f       Nd= f)Nr2   rh   rj   rL   r   r   r         @      @              ?zNo matching definition)r   r4   r   rq   rr   r	   rt   r   r   ENABLE_CUDASIMr    r   rs   r#   r$   r%   rM   rN   r   )r&   sigsr/   r   cms        r   _test_explicit_signatures(TestDispatcher._test_explicit_signatures   s   HHTN:& HHQbhh'	$1a!a(HHQbjj)	$3!c*   y)R"--0AdGAr2 * 	.BLL0ABQ[[)1akk:	 *)s   1E%%
E3c                 .    SS/nU R                  U5        g )N(int64[::1], int64, int64) (float64[::1], float64, float64))r   r&   r   s     r    test_explicit_signatures_strings/TestDispatcher.test_explicit_signatures_strings   s    ,24&&t,r   c                     [         S S S2   [         [         4[        S S S2   [        [        4/nU R                  U5        g Nr2   )r	   r   r   r   s     r   test_explicit_signatures_tuples.TestDispatcher.test_explicit_signatures_tuples   s8    ssUE*WSqS\7G,LM&&t,r   c                     [        [        S S S2   [        [        5      [        [        S S S2   [        [        5      /nU R                  U5        g r   )r
   r	   r   r   r   s     r   #test_explicit_signatures_signatures2TestDispatcher.test_explicit_signatures_signatures  s?    U3Q3Z.WSqS\7G46&&t,r   c                 J   [         S S S2   [         [         4S/nU R                  U5        [         S S S2   [         [         4[        [        S S S2   [        [        5      /nU R                  U5        [        [         S S S2   [         [         5      S/nU R                  U5        g )Nr2   r   )r	   r   r
   r   r   s     r   test_explicit_signatures_mixed-TestDispatcher.test_explicit_signatures_mixed  s     ssUE*24&&t, ssUE*WSqS\7G46&&t, U3Q3Z.24&&t,r   c                    SS/n[         R                  " U5      " [        5      n[        R                  " S[        R
                  S9nUS   " U[        R                  " S5      [        R                  " S5      5        U R                  US   S5        [        R                  " S[        R
                  S9nUS   " USS5        U R                  US   S	5        g )
Nz (float64[::1], float32, float32)r   r2   rh   rj         `>r         ?     ?)r   r4   r   rq   rr   r   r   rt   r&   r   r/   r   s       r   (test_explicit_signatures_same_type_class7TestDispatcher.test_explicit_signatures_same_type_class  s     324HHTN:&HHQbjj)	$2::a="**V"45!c*HHQbjj)	$1f!&89r   z'No overload resolution in the simulatorc                    [         R                  " / SQ5      " [        5      nU R                  [        5       n[
        R                  " S[
        R                  S9nUS   " USS5        S S S 5        U R                  [        WR                  5      S5        U R                  S[        UR                  5      5        g ! , (       d  f       NY= f)	N)z (float64[::1], float32, float64)z (float64[::1], float64, float32)z(float64[::1], int64, int64)r2   rh   rj   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   r4   r   r    r   rq   rr   r   assertRegexr$   r%   assertNotIn)r&   r/   r   r   s       r   -test_explicit_signatures_ambiguous_resolution<TestDispatcher.test_explicit_signatures_ambiguous_resolution+  s     HH 6 77AC y)R"**-AdGAsC  * 	"		
 	#bll"34# *)s   1B<<
C
z$Simulator does not use _prepare_argsc                 (   [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nUS   " USS5        U R                  US   S5        U R                  [        UR                  5      SUR                  5        SS	/n[         R                  " U5      " [        5      n[        R                  " S[        R                  S9nUS   " U[        R                  " S5      S5        U R                  US   S
5        g )Nr   r2   rh   rj   r   r   r   r   r         @)r   r4   r   rq   rr   r	   rt   rM   rN   r   r   r   )r&   r/   r   r   s       r   test_explicit_signatures_unsafe.TestDispatcher.test_explicit_signatures_unsafeE  s     HH12:>HHQbhh' 	
$3!a(Q[[)1akk:,24HHTN:&HHQbjj)	$288A;$!c*r   c                 x   ^ [         R                  " USS9" [        5      m[         R                  U4S j5       nU$ )NTdevicec                    > T" X5      U S'   g r   r   )r   r   r   
add_devices      r   r/   ,TestDispatcher.add_device_usecase.<locals>.f`  s    a#AaDr   )r   r4   r   )r&   r   r/   r   s      @r   add_device_usecase!TestDispatcher.add_device_usecase[  s7     XXd405
		$ 
	$ r   c                    SS/nU R                  U5      n[        R                  " S[        R                  S9nUS   " USS5        U R	                  US   S5        [        R                  " S[        R
                  S9nUS   " US	S
5        U R	                  US   S5        [        R                  (       a  g U R                  [        5       n[        R                  " S[        R                  S9nUS   " USS5        S S S 5        [        WR                  5      nU R                  SU5        U R                  SU5        U R                  [        UR                   5      SUR                   5        g ! , (       d  f       Nx= f)N(int64, int64)(float64, float64)r2   rh   rj   rL   r   r   r   r   r   r   zInvalid use of typez(with parameters (complex128, complex128))r   rq   rr   r	   rt   r   r   r   r    r   rs   r$   r%   r#   rM   rN   r   )r&   r   r/   r   r   msgs         r   test_explicit_signatures_device.TestDispatcher.test_explicit_signatures_devicef  s.    !"67##D) HHQbhh'	$1a!a(HHQbjj)	$3!c*   {+r"--0AdGAr2 , ",,+S1@#FQ[[)1akk: ,+s   1E..
E<c                    SS/nU R                  U5      n[        R                  " S[        R                  S9nUS   " U[        R                  " S5      [        R                  " S5      5        U R                  US   S5        [        R                  " S[        R                  S9nUS   " USS5        U R                  US   S	5        g )
Nz(float32, float32)r   r2   rh   rj   r   r   r   r   )r   rq   rr   r   r   rt   r   s       r   /test_explicit_signatures_device_same_type_class>TestDispatcher.test_explicit_signatures_device_same_type_class  s     %&:;##D)HHQbjj)	$2::a="**V"45!c*HHQbjj)	$1f!&89r   c                     / SQnU R                  U5      n[        R                  " S[        R                  S9nUS   " USS5        U R	                  US   S5        g )	N)z(float32, float64)z(float64, float32)r   r2   rh   rj   r   r   r   r   )r   rq   rr   r   rt   r   s       r   )test_explicit_signatures_device_ambiguous8TestDispatcher.test_explicit_signatures_device_ambiguous  sR     N##D)HHQbjj)	$3!c*r   z%CUDA Simulator does not force castingc                    S/nU R                  U5      n[        R                  " S[        R                  S9nUS   " USS5        U R	                  US   S5        U R                  [        UR                  5      SUR                  5        SS	/nU R                  U5      n[        R                  " S[        R                  S9nUS   " U[        R                  " S5      S5        U R	                  US   S
5        g )Nr   r2   rh   rj   r   r   r   r   r   r   )
r   rq   rr   r	   rt   rM   rN   r   r   r   r   s       r   &test_explicit_signatures_device_unsafe5TestDispatcher.test_explicit_signatures_device_unsafe  s     !!##D) HHQbhh'	$3!a(Q[[)1akk: "67##D) HHQbjj)	$288A;$!c*r   c                     [         R                  S 5       n[         R                  " SS9S 5       nU R                  SUR                  5        U R                  SUR                  5        g )Nc                     g) Add two integers, kernel versionNr   abs     r   r   <TestDispatcher.test_dispatcher_docstring.<locals>.add_kernel      r   Tr   c                     g) Add two integers, device versionNr   r   s     r   r   <TestDispatcher.test_dispatcher_docstring.<locals>.add_device  r   r   r   r   )r   r4   rM   __doc__)r&   r   r   s      r   test_dispatcher_docstring(TestDispatcher.test_dispatcher_docstring  sg     
	3 
	3 
		3 
	3 	;Z=O=OP;Z=O=OPr   r   N)r_   r`   ra   rb   r   rw   r   r   expectedFailurer|   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   rc   r   r   r   re   re   n   s   C58 233  43 23( 4( 9:= ;=8 CD! E!4;.--
--$:$ >?5 @52 ;<+  =+(	;::"+ <=+ >+,Qr   re   z2CUDA simulator doesn't implement kernel propertiesc                   J    \ rS rSrS rS rS rS rS rS r	S r
S	 rS
 rSrg)TestDispatcherKernelPropertiesrl   c                 :   [         R                  S 5       nSn[        R                  " U[        R                  S9n[        R                  " U[        R
                  S9nUSU4   " X25        USU4   " XB5        [        [        S S S2   [        5      n[        [
        S S S2   [        5      nUR                  U5      nUR                  U5      nU R                  U[        5        U R                  U[        5        U R                  US5        U R                  US5        UR                  5       n	U R                  XR                     U5        U R                  XR                     U5        Xx:X  a,  [        S5        [        S5        [         R                  " 5         g g )Nc                 v    [         R                  " S5      nX!:  a  S[        R                  " X   5      -  X'   g g Nr2   gQ	@r   gridmathsinr   nr   s      r   pi_sin_array[TestDispatcherKernelProperties.test_get_regs_per_thread_unspecialized.<locals>.pi_sin_array  /    		!Audhhqtn, r   
   rh   r2   r   z,f32 and f64 variant thread usages are equal.z-This may warrant some investigation. Devices:)r   r4   rq   rr   r   r   r
   r	   get_regs_per_threadassertIsInstanceintassertGreaterrM   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_unspecializedETestDispatcherKernelProperties.test_get_regs_per_thread_unspecialized  sl    
	- 
	- ((1BJJ/((1BJJ/QT7&QT7& wss|U+wss|U+*>>wG*>>wG137137.2.2
 +>>@,\\:,	.,\\:,	. 5 @AABKKM 6r   c                     [         R                  " [        [        S S S2   [        5      5      S 5       nUR                  5       nU R                  U[        5        U R                  US5        g )Nr2   c                 v    [         R                  " S5      nX!:  a  S[        R                  " X   5      -  X'   g g r   r   r   s      r   r   YTestDispatcherKernelProperties.test_get_regs_per_thread_specialized.<locals>.pi_sin_array  r   r   r   )	r   r4   r
   r   r	   r  r  r  r  )r&   r   regs_per_threads      r   $test_get_regs_per_thread_specializedCTestDispatcherKernelProperties.test_get_regs_per_thread_specialized  s\    	$wss|U+	,	- 
-	- '::<os3?A.r   c                 ,   [         R                  S 5       nUS   " SS5        US   " SS5        [        [        [        5      n[        [
        [        5      nUR                  U5      nUR                  U5      nU R                  U[        5        U R                  U[        5        U R                  US5        U R                  US5        UR                  5       nU R                  XbR                     U5        U R                  XcR                     U5        g )Nc                 *    U(       a  [        U 5        g g r   )r  )valto_prints     r   const_fmt_stringYTestDispatcherKernelProperties.test_get_const_mem_unspecialized.<locals>.const_fmt_string  s     c
 r   rj   r2   Fr      r   )r   r4   r
   r	   r   r   get_const_mem_sizer  r  assertGreaterEqualrM   r  )r&   r  sig_i64r  const_mem_size_i64const_mem_size_f64const_mem_size_alls          r    test_get_const_mem_unspecialized?TestDispatcherKernelProperties.test_get_const_mem_unspecialized  s    		 
	 	q%(sE* ug&w(-@@I-@@I0#60#6 	 2A6 2A6 .@@B+LL9;MN+LL9;MNr   c                 :  ^ [         R                  " S[         R                  S9m[        [        S S S2   5      n[        R
                  " U5      U4S j5       nUR                  U5      nU R                  U[        5        U R                  UTR                  5        g )N    rh   r2   c                 |   > [         R                  R                  T5      n[         R                  " S5      nX   X'   g r   )r   const
array_liker   )r   Cr   arrs      r   const_array_useVTestDispatcherKernelProperties.test_get_const_mem_specialized.<locals>.const_array_use,  s-    

%%c*A		!A4ADr   )rq   aranger	   r
   r   r4   r  r  r  r   nbytes)r&   sigr.  const_mem_sizer-  s       @r   test_get_const_mem_specialized=TestDispatcherKernelProperties.test_get_const_mem_specialized(  sy    ii"((+51:	#	 
	
 );;C@nc2

;r   c                   ^
 Sm
[         R                  U
4S j5       n[        R                  " T
[        R                  S9n[        R                  " T
[        R
                  S9nUS   " U5        US   " U5        [        [        S S S2   5      n[        [
        S S S2   5      nUR                  U5      nUR                  U5      nU R                  U[        5        U R                  U[        5        U R                  UT
S-  5        U R                  UT
S-  5        UR                  5       nUR                  5       n	U R                  XR                     U5        U R                  XR                     U5        g )Nr   c                    > [         R                  R                  TU R                  S9n[	        T5       H  nX!U'   M	     [	        T5       H	  nX   X'   M     g Nrh   )r   sharedarrayri   r   )arysmjr  s      r   simple_smem_TestDispatcherKernelProperties.test_get_shared_mem_per_block_unspecialized.<locals>.simple_smem;  sL    ""1CII"6B1X1 1X r   rh   rj   r2   r      )r   r4   rq   rr   r   r   r
   get_shared_mem_per_blockr  r  rM   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_unspecializedJTestDispatcherKernelProperties.test_get_shared_mem_per_block_unspecialized6  s<    
	 
	 ((1BJJ/((1BJJ/D'"D'"wss|$wss|$ 99'B
 99'B
j#.j#.QU+QU+
 %==?$==?5zB5zBr   c                     [         R                  " [        [        S S S2   5      5      S 5       nUR	                  5       nU R                  U[        5        U R                  US5        g )Nr2   c                     [         R                  R                  S[        S9n[         R                  " S5      nUS:X  a  [        S5       H  nX1U'   M	     [         R                  " 5         X   X'   g )Nd   rh   r2   r   )r   r9  r:  r   r   r   syncthreads)r;  r<  r   r=  s       r   r>  ]TestDispatcherKernelProperties.test_get_shared_mem_per_block_specialized.<locals>.simple_smem`  sY    ""3g"6B		!AAvsAqE $UCFr   i  )r   r4   r
   r   rA  r  r  rM   )r&   r>  shared_mem_per_blocks      r   )test_get_shared_mem_per_block_specializedHTestDispatcherKernelProperties.test_get_shared_mem_per_block_specialized_  s]    	$wss|$	%	 
&	  +CCE2C8-s3r   c                    Sn[         R                  S 5       n[        R                  " U[        R                  S9nUS   " U5        [        [        S S S2   5      nUR                  U5      nU R                  U[        5        U R                  US5        UR                  5       nU R                  XdR                     U5        g )Nr   c                 8    [         R                  " S5      nXU'   g r   )r   r   )r;  r   s     r   simple_maxthreadsfTestDispatcherKernelProperties.test_get_max_threads_per_block_unspecialized.<locals>.simple_maxthreadsq  s    		!AFr   rh   rj   r2   r   )r   r4   rq   rr   r   r
   get_max_threads_per_blockr  r  r  rM   r  )r&   r  rR  r	  r  max_threads_f32max_threads_f32_alls          r   ,test_get_max_threads_per_block_unspecializedKTestDispatcherKernelProperties.test_get_max_threads_per_block_unspecializedn  s    		 
	 ((1BJJ/$(wss|$+EEgNos3?A./IIK,\\:OLr   c                   ^	 Sm	[         R                  U	4S j5       n[        R                  " T	[        R                  S9n[        R                  " T	[        R
                  S9nUS   " U5        US   " U5        [        [        S S S2   5      n[        [
        S S S2   5      nUR                  U5      nUR                  U5      nU R                  U[        5        U R                  U[        5        U R                  UT	S-  5        U R                  UT	S-  5        UR                  5       nU R                  XR                     U5        U R                  XR                     U5        g )N  c                    > [         R                  R                  TU R                  S9n[	        T5       H  nX!U'   M	     [	        T5       H	  nX   X'   M     g r8  r   localr:  ri   r   r;  lmr=  r  s      r   simple_lmem_TestDispatcherKernelProperties.test_get_local_mem_per_thread_unspecialized.<locals>.simple_lmem  L    !!!399!5B1X1 1X r   rh   rj   r2   r   r@  )r   r4   rq   rr   r   r   r
   get_local_mem_per_threadr  r  r   rM   r  )
r&   r`  r	  r
  r  r  local_mem_f32local_mem_f64local_mem_allr  s
            @r   +test_get_local_mem_per_thread_unspecializedJTestDispatcherKernelProperties.test_get_local_mem_per_thread_unspecialized  s.    		 
	 ((1BJJ/((1BJJ/D'"D'"wss|$wss|$#<<WE#<<WEmS1mS1q1u5q1u5
 $<<>||4mD||4mDr   c                    ^ Sm[         R                  " [        [        S S S2   5      5      U4S j5       nUR	                  5       nU R                  U[        5        U R                  UTS-  5        g )NrZ  r2   c                    > [         R                  R                  TU R                  S9n[	        T5       H  nX!U'   M	     [	        T5       H	  nX   X'   M     g r8  r\  r^  s      r   r`  ]TestDispatcherKernelProperties.test_get_local_mem_per_thread_specialized.<locals>.simple_lmem  rb  r   r   )r   r4   r
   r   rc  r  r  r   )r&   r`  local_mem_per_threadr  s      @r   )test_get_local_mem_per_thread_specializedHTestDispatcherKernelProperties.test_get_local_mem_per_thread_specialized  si     	$wss|$	%	 
&	  +CCE2C8 4a!e<r   r   N)r_   r`   ra   rb   r  r  r%  r4  rF  rN  rW  rg  rm  rc   r   r   r   r   r     s7    -^
/!OF<'CR4M&%EN=r   r   __main__)numpyrq   r   numbar   r   r   r   r   r   r	   r
   numba.core.errorsr   numba.cuda.testingr   r   r   r   r   r   r   re   r   r_   mainr   r   r   <module>ru     s      M M M ) F F  BCX2< X2 DX2vWQ\ WQt
 EFo=\ o= Go=d zMMO r   