+
    :ig                        ^ RI t^ RIt^ RIHtHtHtHtHtH	t	H
t
Ht ^ RIHt ^ RIHtHtHt ^ RItR tR t]! R4       ! R R	]4      4       t ! R
 R]4      t]! R4       ! R R]4      4       t]R8X  d   ]P2                  ! 4        R# R# )    N)booleanconfigcudafloat32float64int32int64void)TypingError)skip_on_cudasimunittestCUDATestCasec                     W,           # N xys   &&e/var/www/html/photoedit/myenv/lib/python3.14/site-packages/numba/cuda/tests/cudapy/test_dispatcher.pyaddr   
   s	    5L    c                     W,           V ^ &   R# r   Nr   )rr   r   s   &&&r   
add_kernelr      s    5AaDr   z/Specialization not implemented in the simulatorc                   J   a  ] tR t^t o R tR tR tR tR tR t	R t
RtV tR	# )
TestDispatcherSpecializationc                    V P                  \        4      ;_uu_ 4       pVP                  V4       R R R 4       V P                  R\	        XP
                  4      4       R #   + '       g   i     L7; i)NzDispatcher already specialized)assertRaisesRuntimeError
specializeassertInstr	exception)self
dispatchertyes   &&& r   _test_no_double_specialize7TestDispatcherSpecialization._test_no_double_specialize   sK    |,,!!"% - 	6AKK8HI -,s   AA/	c                x    \         P                  ! R 4      R 4       pV P                  V\        R,          4       R# )zvoid(float32[::1])c                     R # r   r   r   s   &r   fPTestDispatcherSpecialization.test_no_double_specialize_sig_same_types.<locals>.f       r   NN   Nr   jitr)   r   r%   r.   s   & r   (test_no_double_specialize_sig_same_typesETestDispatcherSpecialization.test_no_double_specialize_sig_same_types   s6     
&	'	 
(	 	''73<8r   c                    \         P                  R  4       pVP                  \        R,          4      pV P	                  V\        R,          4       R# )c                     R # r   r   r-   s   &r   r.   STestDispatcherSpecialization.test_no_double_specialize_no_sig_same_types.<locals>.f'   r0   r   r1   N)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$   sB     
	 
	 WS\2''ws|Dr   c                x    \         P                  ! R 4      R 4       pV P                  V\        R,          4       R# )zvoid(int32[::1])c                     R # r   r   r-   s   &r   r.   PTestDispatcherSpecialization.test_no_double_specialize_sig_diff_types.<locals>.f0   r0   r   r1   Nr3   r5   s   & r   (test_no_double_specialize_sig_diff_typesETestDispatcherSpecialization.test_no_double_specialize_sig_diff_types.   s4    	$	%	 
&	 	''73<8r   c                    \         P                  R  4       pVP                  \        R,          4      pV P	                  V\
        R,          4       R# )c                     R # r   r   r-   s   &r   r.   STestDispatcherSpecialization.test_no_double_specialize_no_sig_diff_types.<locals>.f8   r0   r   r1   N)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   s@    		 
	 U3Z0''ws|Dr   c                F   \         P                  R  4       pV P                  \        VP                  4      ^ 4       VP                  \        R,          4      pV P                  \        VP                  4      ^4       VP                  \        R,          4      pV P                  \        VP                  4      ^4       V P                  W#4       VP                  \        R,          4      pV P                  \        VP                  4      ^4       V P                  WB4       R# )c                     R # r   r   r-   s   &r   r.   BTestDispatcherSpecialization.test_specialize_cache_same.<locals>.fC   r0   r   r1   N)
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.	Q../3ll73<0Q../3i-,,uSz*Q../3,r   c                   \         P                  R  4       pV P                  \        VP                  4      ^ 4       VP                  \        R,          \        R,          4      pV P                  \        VP                  4      ^4       VP                  \        R,          \        R,          4      pV P                  \        VP                  4      ^4       V P                  W#4       VP                  \        R,          \        R,          4      pV P                  \        VP                  4      ^4       V P                  W44       R# )c                     R # r   r   r   s   &&r   r.   PTestDispatcherSpecialization.test_specialize_cache_same_with_ordering.<locals>.fY   r0   r   :NNNr1   N)	r   r4   rL   rM   rN   r!   r   rP   rO   )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 ll73<>Q../32 WS\73<@Q../3k1r   r   N)__name__
__module____qualname____firstlineno__r)   r6   r=   rB   rG   rT   r\   __static_attributes____classdictcell____classdict__s   @r   r   r      s0     J9E9E-*2 2r   r   c                   l  a  ] tR t^nt o RtR t]! R4      ]P                  R 4       4       t	]! R4      R 4       t
]! R4      R 4       t]! R4      R	 4       tR
 tR tR tR tR tR t]! R4      R 4       t]! R4      ]P                  R 4       4       tR tR tR tR t]! R4      R 4       tR tRtV tR# )TestDispatcherz9Most tests based on those in numba.tests.test_dispatcher.c                   \         P                  ! \        4      p\        P                  ! ^\        P
                  R7      pVR	,          ! V^{R4       V P                  V^ ,          \        ^{R4      4       VR	,          ! VRR4       V P                  V^ ,          \        RR4      4       VR	,          ! VRR4       V P                  V^ ,          \        RR4      4       VR	,          ! VRR4       V P                  V^ ,          \        RR4      4       \         P                  ! R4      ! \        4      p\        P                  ! ^\        P                  R7      pVR	,          ! V^{R4       V P                  V^ ,          \        ^{R4      4       R# )
r2   dtype  皙(@F@        F@l    F: (i4[::1], i4, i4)Nr2   r2   )
r   r4   r   npzeros
complex128rL   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                    \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! VRR4       V P                  V^ ,          \        ^^-4      4       R# )ro   ri   rl   rm   Nrp   )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                (   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pV P                  \        4      ;_uu_ 4        VR,          ! VRR4       RRR4       R#   + '       g   i     R# ; i)ro   ri   rl   rn   Nrp   )r   r4   r   rq   rr   r   r   	TypeErrorru   s   &  r   &test_coerce_input_types_unsafe_complex5TestDispatcher.test_coerce_input_types_unsafe_complex   s\     ,-j9HHQbhh'y))$K4' *)))s   $B  B	z"Simulator does not track overloadsc                2   \         P                  ! \        4      p\        P                  ! ^\        P
                  R7      p^pRpVR,          ! W#V4       V P                  V^ ,          W4,           4       V P                  \        VP                  4      ^4       VR,          ! W$V4       V P                  V^ ,          WC,           4       V P                  \        VP                  4      ^4       VR,          ! W$V4       V P                  V^ ,          WD,           4       V P                  \        VP                  4      ^4       VR,          ! V^^4       V P                  V^ ,          W3,           4       V P                  \        VP                  4      ^R4       R# )z8Test compiling new version in an ambiguous case
        ri         ?zdidn't compile a new versionNrp   )
r   r4   r   rq   rr   r   assertAlmostEqualrL   rM   	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                :  a aa . o\         P                  R 4       oVVV 3R lp\        ^4       Uu. uF  p\        P                  ! VR7      NK  	  ppV F  pVP                  4        K  	  V F  pVP                  4        K  	  S P                  S4       R# u upi )zb
Test that (lazy) compiling from several threads at once doesn't
produce errors (see issue #908).
c                     V^,           V ^ &   R# r2   Nr   )r   r   s   &&r   foo%TestDispatcher.test_lock.<locals>.foo   s    q5AaDr   c                     <  \         P                  ! ^\         P                  R7      p SR,          ! V ^4       SP                  V ^ ,          ^4       R#   \         d   pSP                  T4        Rp?R# Rp?ii ; i)r2   ri   Nrp   )rq   rr   r	   rL   	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;A66A;)targetN)r   r4   range	threadingThreadstartjoinassertFalse)r%   r   ithreadstr   r   s   f    @@r   	test_lockTestDispatcher.test_lock   s     		 
		! >C2YGY9##73YGAGGI AFFH   Hs    Bc                @   \         P                  ! V4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! V^^4       V P                  V^ ,          ^4       \        P                  ! ^\        P                  R7      pVR,          ! VRR4       V P                  V^ ,          R4       \        P                  '       d   R# V P                  \        4      ;_uu_ 4       p\        P                  ! ^\        P                  R7      pVR,          ! VRR4       RRR4       V P                  R\        XP                  4      4       V P!                  \#        VP$                  4      ^VP$                  4       R#   + '       g   i     Lg; i)	r2   ri   r         @      @N              ?zNo matching definitionrp   )r   r4   r   rq   rr   r	   rt   r   r   ENABLE_CUDASIMr   r}   rs   r"   r#   r$   rL   rM   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   68FF	c                0    R R.pV P                  V4       R# )(int64[::1], int64, int64) (float64[::1], float64, float64)N)r   r%   r   s   & r    test_explicit_signatures_strings/TestDispatcher.test_explicit_signatures_strings   s    ,24&&t,r   c                    \         R ,          \         \         3\        R ,          \        \        3.pV P                  V4       R# r1   N)r	   r   r   r   s   & r   test_explicit_signatures_tuples.TestDispatcher.test_explicit_signatures_tuples   s0    sUE*WS\7G,LM&&t,r   c                    \        \        R ,          \        \        4      \        \        R ,          \        \        4      .pV P                  V4       R# r   )r
   r	   r   r   r   s   & r   #test_explicit_signatures_signatures2TestDispatcher.test_explicit_signatures_signatures  s7    U3Z.WS\7G46&&t,r   c                T   \         R ,          \         \         3R.pV P                  V4       \         R ,          \         \         3\        \        R ,          \        \        4      .pV P                  V4       \        \         R ,          \         \         4      R.pV P                  V4       R# )r1   r   N)r	   r   r
   r   r   s   & r   test_explicit_signatures_mixed-TestDispatcher.test_explicit_signatures_mixed  s     sUE*24&&t, sUE*WS\7G46&&t, U3Z.24&&t,r   c                   R R.p\         P                  ! V4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! V\        P                  ! ^4      \        P                  ! R4      4       V P                  V^ ,          R4       \        P                  ! ^\        P
                  R7      pVR,          ! V^R4       V P                  V^ ,          R4       R# )z (float64[::1], float32, float32)r   ri         ?     ?Nrp         `>)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                   \         P                  ! . RO4      ! \        4      pV P                  \        4      ;_uu_ 4       p\
        P                  ! ^\
        P                  R7      pVR,          ! VRR4       RRR4       V P                  \        XP                  4      R4       V P                  R\        VP                  4      4       R#   + '       g   i     L\; i)	 (float64[::1], float32, float64)ri   r   g       @Na  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   z (float64[::1], float64, float32)z(float64[::1], int64, int64)rp   )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    8CC	z$Simulator does not use _prepare_argsc                V   \         P                  ! R 4      ! \        4      p\        P                  ! ^\        P
                  R7      pVR,          ! VRR4       V P                  V^ ,          ^4       V P                  \        VP                  4      ^VP                  4       R R.p\         P                  ! V4      ! \        4      p\        P                  ! ^\        P                  R7      pVR,          ! V\        P                  ! ^4      R4       V P                  V^ ,          R4       R# )r   ri   r   r   r         @Nrp   )r   r4   r   rq   rr   r	   rt   rL   rM   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                ~   a \         P                  ! VR R7      ! \        4      o\         P                  V3R l4       pV# )Tdevicec                    < S! W4      V ^ &   R# 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                R   R R.pV P                  V4      p\        P                  ! ^\        P                  R7      pVR
,          ! V^^4       V P	                  V^ ,          ^4       \        P                  ! ^\        P
                  R7      pVR
,          ! VRR4       V P	                  V^ ,          R4       \        P                  '       d   R# V P                  \        4      ;_uu_ 4       p\        P                  ! ^\        P                  R7      pVR
,          ! VRR4       RRR4       \        XP                  4      pV P                  RV4       V P                  R	V4       V P                  \        VP                   4      ^VP                   4       R#   + '       g   i     L{; i)(int64, int64)(float64, float64)ri   r   r   r   Nr   zInvalid use of typez(with parameters (complex128, complex128)rp   )r   rq   rr   r	   rt   r   r   r   r   r   rs   r#   r$   r"   rL   rM   r   )r%   r   r.   r   r   msgs   &     r   test_explicit_signatures_device.TestDispatcher.test_explicit_signatures_devicef  s4    !"67##D) HHQbhh'	$1a!a(HHQbjj)	$3!c*    {++r"--0AdGAr2 , ",,+S1@#FQ[[)1akk: ,+s   +8FF&	c                   R R.pV P                  V4      p\        P                  ! ^\        P                  R7      pVR,          ! V\        P                  ! ^4      \        P                  ! R4      4       V P                  V^ ,          R4       \        P                  ! ^\        P                  R7      pVR,          ! V^R4       V P                  V^ ,          R4       R# )z(float32, float32)r   ri   r   r   Nrp   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                    . ROpV P                  V4      p\        P                  ! ^\        P                  R7      pVR,          ! VRR4       V P	                  V^ ,          R4       R# )(float32, float64)ri   r   r   r   N)r   z(float64, float32)r   rp   )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                    R .pV P                  V4      p\        P                  ! ^\        P                  R7      pVR,          ! VRR4       V P	                  V^ ,          ^4       V P                  \        VP                  4      ^VP                  4       R R.pV P                  V4      p\        P                  ! ^\        P                  R7      pVR,          ! V\        P                  ! ^4      R4       V P	                  V^ ,          R4       R# )r   ri   r   r   r   r   Nrp   )
r   rq   rr   r	   rt   rL   rM   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                    \         P                  R  4       p\         P                  ! RR7      R 4       pV P                  RVP                  4       V P                  RVP                  4       R# )c                    R# ) Add two integers, kernel versionNr   abs   &&r   r   <TestDispatcher.test_dispatcher_docstring.<locals>.add_kernel      r   Tr   c                    R# ) Add two integers, device versionNr   r   s   &&r   r   <TestDispatcher.test_dispatcher_docstring.<locals>.add_device  r   r   r   r   N)r   r4   rL   __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_   r`   ra   r   rw   r   r   expectedFailurerz   r~   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   rb   rc   rd   s   @r   rg   rg   n   s    C58 233  43 23( 4( 9:= ;=8 CD! E!4;.--
--$:$ >?5 @52 ;<+  =+(	;::"+ <=+ >+,Q Qr   rg   z2CUDA simulator doesn't implement kernel propertiesc                   V   a  ] tR tRt o R tR tR tR tR tR t	R t
R	 tR
 tRtV tR# )TestDispatcherKernelPropertiesrk   c                p   \         P                  R  4       p^
p\        P                  ! V\        P                  R7      p\        P                  ! V\        P
                  R7      pV^V3,          ! W24       V^V3,          ! WB4       \        \        R,          \        4      p\        \
        R,          \        4      pVP                  V4      pVP                  V4      pV P                  V\        4       V P                  V\        4       V P                  V^ 4       V P                  V^ 4       VP                  4       p	V P                  WP                  ,          V4       V P                  WP                  ,          V4       Wx8X  d.   \        R4       \        R4       \         P                  ! 4        R# R# )c                     \         P                  ! ^4      pW!8  d(   R\        P                  ! W,          4      ,          W&   R# R# r2   gQ	@N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   ri   r1   z,f32 and f64 variant thread usages are equal.z-This may warrant some investigation. Devices:N)r   r4   rq   rr   r   r   r
   r	   get_regs_per_threadassertIsInstanceintassertGreaterrL   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  sd    
	- 
	- ((1BJJ/((1BJJ/QT7&QT7& ws|U+ws|U+*>>wG*>>wG137137.2.2
 +>>@,\\:,	.,\\:,	. 5 @AABKKM 6r   c                    \         P                  ! \        \        R ,          \        4      4      R 4       pVP                  4       pV P                  V\        4       V P                  V^ 4       R# )r1   c                     \         P                  ! ^4      pW!8  d(   R\        P                  ! W,          4      ,          W&   R# R# r   r   r   s   && r   r   YTestDispatcherKernelProperties.test_get_regs_per_thread_specialized.<locals>.pi_sin_array  r   r   N)	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  sX    	$ws|U+	,	- 
-	- '::<os3?A.r   c                N   \         P                  R  4       pVR,          ! ^R4       VR,          ! RR4       \        \        \        4      p\        \
        \        4      pVP                  V4      pVP                  V4      pV P                  V\        4       V P                  V\        4       V P                  V^4       V P                  V^4       VP                  4       pV P                  WbP                  ,          V4       V P                  WcP                  ,          V4       R# )c                 0    V'       d   \        V 4       R # R # r   )r  )valto_prints   &&r   const_fmt_stringYTestDispatcherKernelProperties.test_get_const_mem_unspecialized.<locals>.const_fmt_string  s     c
 r   Fr   Nrp   )r   r4   r
   r	   r   r   get_const_mem_sizer   r   assertGreaterEqualrL   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                D  a \         P                  ! ^ \         P                  R7      o\        \        R,          4      p\        P
                  ! V4      V3R l4       pVP                  V4      pV P                  V\        4       V P                  VSP                  4       R# )    ri   r1   c                    < \         P                  P                  S4      p\         P                  ! ^4      pW,          W&   R# 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   N)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(  su    ii"((+5:	#	 
	
 );;C@nc2

;r   c                2  a
 ^
o
\         P                  V
3R l4       p\        P                  ! S
\        P                  R7      p\        P                  ! S
\        P
                  R7      pVR,          ! V4       VR,          ! V4       \        \        R,          4      p\        \
        R,          4      pVP                  V4      pVP                  V4      pV P                  V\        4       V P                  V\        4       V P                  VS
^,          4       V P                  VS
^,          4       VP                  4       pVP                  4       p	V P                  WP                  ,          V4       V P                  WP                  ,          V4       R# )
   c                    < \         P                  P                  SV P                  R 7      p\	        S4       F  pW!V&   K	  	  \	        S4       F  pW,          W&   K  	  R# ri   N)r   sharedarrayrj   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   ri   r1   Nrp   )r   r4   rq   rr   r   r   r
   get_shared_mem_per_blockr   r   rL   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  s4    
	 
	 ((1BJJ/((1BJJ/D'"D'"ws|$ws|$ 99'B
 99'B
j#.j#.QU+QU+
 %==?$==?5zB5zBr   c                    \         P                  ! \        \        R ,          4      4      R 4       pVP	                  4       pV P                  V\        4       V P                  VR4       R# )r1   c                     \         P                  P                  ^d\        R7      p\         P                  ! ^4      pV^ 8X  d   \        ^d4       F  pW1V&   K	  	  \         P                  ! 4        W,          W&   R# )d   ri   N)r   r6  r7  r   r   r   syncthreads)r8  r9  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  N)r   r4   r
   r   r=  r   r   rL   )r%   r;  shared_mem_per_blocks   &  r   )test_get_shared_mem_per_block_specializedHTestDispatcherKernelProperties.test_get_shared_mem_per_block_specialized_  sY    	$ws|$	%	 
&	  +CCE2C8-s3r   c                   ^
p\         P                  R 4       p\        P                  ! V\        P                  R7      pVR,          ! V4       \        \        R,          4      pVP                  V4      pV P                  V\        4       V P                  V^ 4       VP                  4       pV P                  WdP                  ,          V4       R# )r3  c                 :    \         P                  ! ^4      pWV&   R# r   )r   r   )r8  r   s   & r   simple_maxthreadsfTestDispatcherKernelProperties.test_get_max_threads_per_block_unspecialized.<locals>.simple_maxthreadsq  s    		!AFr   ri   r1   Nrp   )r   r4   rq   rr   r   r
   get_max_threads_per_blockr   r   r  rL   r  )r%   r  rN  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|$+EEgNos3?A./IIK,\\:OLr   c                  a	 R o	\         P                  V	3R l4       p\        P                  ! S	\        P                  R7      p\        P                  ! S	\        P
                  R7      pVR,          ! V4       VR,          ! V4       \        \        R,          4      p\        \
        R,          4      pVP                  V4      pVP                  V4      pV P                  V\        4       V P                  V\        4       V P                  VS	^,          4       V P                  VS	^,          4       VP                  4       pV P                  WP                  ,          V4       V P                  WP                  ,          V4       R# )  c                    < \         P                  P                  SV P                  R 7      p\	        S4       F  pW!V&   K	  	  \	        S4       F  pW,          W&   K  	  R# r5  r   localr7  rj   r   r8  lmr:  r  s   &  r   simple_lmem_TestDispatcherKernelProperties.test_get_local_mem_per_thread_unspecialized.<locals>.simple_lmem  L    !!!399!5B1X1 1X r   ri   r1   Nrp   )r   r4   rq   rr   r   r   r
   get_local_mem_per_threadr   r   r  rL   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|$ws|$#<<WE#<<WEmS1mS1q1u5q1u5
 $<<>||4mD||4mDr   c                   a R o\         P                  ! \        \        R,          4      4      V3R l4       pVP	                  4       pV P                  V\        4       V P                  VS^,          4       R# )rV  r1   c                    < \         P                  P                  SV P                  R 7      p\	        S4       F  pW!V&   K	  	  \	        S4       F  pW,          W&   K  	  R# r5  rX  rZ  s   &  r   r\  ]TestDispatcherKernelProperties.test_get_local_mem_per_thread_specialized.<locals>.simple_lmem  r^  r   N)r   r4   r
   r   r_  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  se     	$ws|$	%	 
&	  +CCE2C8 4a!e<r   r   N)r^   r_   r`   ra   r  r  r!  r0  rB  rJ  rS  rc  ri  rb   rc   rd   s   @r   r   r     s?     -^
/!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   rg   r   r^   mainr   r   r   <module>rq     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   