
    xKgg                        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j2                          yy)    N)booleanconfigcudafloat32float64int32int64void)TypingError)skip_on_cudasimunittestCUDATestCasec                     | |z   S N xys     k/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_dispatcher.pyaddr   
   s    q5L    c                     ||z   | d<   y Nr   r   )rr   r   s      r   
add_kernelr      s    q5AaDr   z/Specialization not implemented in the simulatorc                   6    e Zd Zd Zd Zd Zd Zd Zd Zd Z	y)	TestDispatcherSpecializationc                     | j                  t              5 }|j                  |       d d d        | j                  dt	        j
                               y # 1 sw Y   /xY w)NzDispatcher already specialized)assertRaisesRuntimeError
specializeassertInstr	exception)self
dispatchertyes       r   _test_no_double_specializez7TestDispatcherSpecialization._test_no_double_specialize   sH    |,!!"% - 	6AKK8HI -,s   AAc                 r    t        j                  d      d        }| j                  |t        d d d          y )Nzvoid(float32[::1])c                      y r   r   r   s    r   fzPTestDispatcherSpecialization.test_no_double_specialize_sig_same_types.<locals>.f       r      r   jitr)   r   r%   r-   s     r   (test_no_double_specialize_sig_same_typeszETestDispatcherSpecialization.test_no_double_specialize_sig_same_types   s:     
&	'	 
(	 	''73Q3<8r   c                     t         j                  d        }|j                  t        d d d         }| j	                  |t        d d d          y )Nc                      y 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$   sJ     
	 
	 WSqS\2''wss|Dr   c                 r    t        j                  d      d        }| j                  |t        d d d          y )Nzvoid(int32[::1])c                      y 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.   s8    	$	%	 
&	 	''73Q3<8r   c                     t         j                  d        }|j                  t        d d d         }| j	                  |t
        d d d          y )Nc                      y 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   sH    		 
	 U3Q3Z0''wss|Dr   c                 B   t         j                  d        }| j                  t        |j                        d       |j                  t        d d d         }| j                  t        |j                        d       |j                  t        d d d         }| j                  t        |j                        d       | j                  ||       |j                  t        d d d         }| j                  t        |j                        d       | j                  ||       y )Nc                      y 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LL1.	Q../3ll73Q3<0Q../3i-,,uSqSz*Q../3),r   c                 x   t         j                  d        }| j                  t        |j                        d       |j                  t        d d  t        d d        }| j                  t        |j                        d       |j                  t        d d d   t        d d d         }| j                  t        |j                        d       | j                  ||       |j                  t        d d d   t        d d d         }| j                  t        |j                        d       | j                  ||       y )Nc                      y 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   s   
 
	 
	 	Q../3 ll71:wqz:Q../3 ll73Q3<1>Q../3k2 WSqS\73Q3<@Q../3k=1r   N)
__name__
__module____qualname__r)   r3   r8   r;   r>   rJ   rP   r   r   r   r   r      s(    J9E9E-*2r   r   c                   X   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y)TestDispatcherz9Most tests based on those in numba.tests.test_dispatcher.c                    t        j                  t              }t        j                  dt        j
                        } |d   |dd       | j                  |d   t        dd              |d   |dd       | j                  |d   t        dd              |d   |dd	       | j                  |d   t        dd	              |d   |d
d       | j                  |d   t        d
d              t        j                  d      t              }t        j                  dt        j                        } |d   |dd       | j                  |d   t        dd             y )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   s2    $ 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                      t        j                  d      t              }t        j                  dt        j
                        } |d   |dd       | j                  |d   t        dd	             y )
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   sZ     .,-j9HHQbhh'dAtT"!c"bk2r   c                      t        j                  d      t              }t        j                  dt        j
                        }| j                  t              5   |d   |dd       d d d        y # 1 sw Y   y xY w)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   sZ     .,-j9HHQbhh'y)E$K4' *))s   A11A:z"Simulator does not track overloadsc                    t        j                  t              }t        j                  dt        j
                        }d}d} |d   |||       | j                  |d   ||z          | j                  t        |j                        d        |d   |||       | j                  |d   ||z          | j                  t        |j                        d        |d   |||       | j                  |d   ||z          | j                  t        |j                        d        |d   |dd       | j                  |d   ||z          | j                  t        |j                        dd	       y
)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   sE    $HHQbjj)dAsC qtS3Y/U__-q1dAsC qtS3Y/U__-q1dAsC qtS3Y/U__-q1 	dAq!qtS3Y/U__-q 3< 	=r   z,Simulator doesn't support concurrent kernelsc                 (    g t         j                  d         fd}t        d      D cg c]  }t        j                  |       }}|D ]  }|j                           |D ]  }|j                            j                         yc c}w )zz
        Test that (lazy) compiling from several threads at once doesn't
        produce errors (see issue #908).
        c                     |dz   | d<   y )Nr/   r   r   )r   r   s     r   fooz%TestDispatcher.test_lock.<locals>.foo   s    q5AaDr   c                      	 t        j                  dt         j                        }  d   | d       j                  | d   d       y # t        $ r}j                  |       Y d }~y 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^    !HHQbhh/D	!Q  1q) !a  !s   AA
 
	A.A))A.   )targetN)r   r1   range	threadingThreadstartjoinassertFalse)r%   r}   ithreadstr|   rx   s   `    @@r   	test_lockzTestDispatcher.test_lock   s     		 
		! >C2YGY9##73YGAGGI AFFH   Hs   Bc                     t        j                  |      t              }t        j                  dt        j
                        } |d   |dd       | j                  |d   d       t        j                  dt        j                        } |d   |dd       | j                  |d   d	       t        j                  ry | j                  t              5 }t        j                  dt        j                        } |d   |d
d
       d d d        | j                  dt        j                               | j!                  t#        |j$                        d|j$                         y # 1 sw Y   _xY w)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   DHHTN:& HHQbhh'$1a!a(HHQbjj)$3!c*   y)R"--0AAdGAr2 * 	.BLL0ABQ[[)1akk:	 *)s   3E%%E.c                 .    ddg}| j                  |       y )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}| j                  |       y Nr/   )r	   r   r   r   s     r   test_explicit_signatures_tuplesz.TestDispatcher.test_explicit_signatures_tuples   s8    ssUE*WSqS\7G,LM&&t,r   c                     t        t        d d d   t        t              t        t        d d d   t        t              g}| j                  |       y r   )r
   r	   r   r   r   s     r   #test_explicit_signatures_signaturesz2TestDispatcher.test_explicit_signatures_signatures  s?    U3Q3Z.WSqS\7G46&&t,r   c                 J   t         d d d   t         t         fdg}| j                  |       t         d d d   t         t         ft        t        d d d   t        t              g}| j                  |       t        t         d d d   t         t               dg}| j                  |       y )Nr/   r   )r	   r   r
   r   r   s     r   test_explicit_signatures_mixedz-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                    ddg} t        j                  |      t              }t        j                  dt        j
                        } |d   |t        j                  d      t        j                  d             | j                  |d   d       t        j                  dt        j
                        } |d   |dd       | j                  |d   d	       y )
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DHHTN:&HHQbjj)$2::a="**V"45!c*HHQbjj)$1f!&89r   z'No overload resolution in the simulatorc                     t        j                  g d      t              }| j                  t              5 }t        j                  dt
        j                        } |d   |dd       d d d        | j                  t        j                        d       | j                  dt        |j                               y # 1 sw Y   TxY w)	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    7DHH 6 77AC y)R"**-AAdGAsC  * 	"		
 	#bll"34# *)s   3B==Cz$Simulator does not use _prepare_argsc                 *    t        j                  d      t              }t        j                  dt        j
                        } |d   |dd       | j                  |d   d       | j                  t        |j                        d|j                         dd	g} t        j                  |      t              }t        j                  dt        j                        } |d   |t        j                  d      d       | j                  |d   d
       y )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     3DHH12:>HHQbhh' 	$3!a(Q[[)1akk:,24DHHTN:&HHQbjj)$288A;$!c*r   c                 z     t        j                  |d      t              t         j                  fd       }|S )NTdevicec                      ||      | d<   y r   r   )r   r   r   
add_devices      r   r-   z,TestDispatcher.add_device_usecase.<locals>.f`  s    a#AaDr   )r   r1   r   )r%   r   r-   r   s      @r   add_device_usecasez!TestDispatcher.add_device_usecase[  s9     1TXXd405
		$ 
	$ r   c                    ddg}| j                  |      }t        j                  dt        j                        } |d   |dd       | j	                  |d   d       t        j                  dt        j
                        } |d   |d	d
       | j	                  |d   d       t        j                  ry | j                  t              5 }t        j                  dt        j                        } |d   |dd       d d d        t        j                        }| j                  d|       | j                  d|       | j                  t        |j                         d|j                          y # 1 sw Y   sxY w)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  s4    !"67##D) HHQbhh'$1a!a(HHQbjj)$3!c*   {+r"--0AAdGAr2 , ",,+S1@#FQ[[)1akk: ,+s   
3E//E8c                    ddg}| j                  |      }t        j                  dt        j                        } |d   |t        j                  d      t        j                  d             | j                  |d   d       t        j                  dt        j                        } |d   |dd       | j                  |d   d	       y )
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)HHQbjj)$2::a="**V"45!c*HHQbjj)$1f!&89r   c                     g d}| j                  |      }t        j                  dt        j                        } |d   |dd       | j	                  |d   d       y )	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  sT     N##D)HHQbjj)$3!c*r   z%CUDA Simulator does not force castingc                    dg}| j                  |      }t        j                  dt        j                        } |d   |dd       | j	                  |d   d       | j                  t        |j                        d|j                         dd	g}| j                  |      }t        j                  dt        j                        } |d   |t        j                  d      d       | j	                  |d   d
       y )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) HHQbhh'$3!a(Q[[)1akk: "67##D) HHQbjj)$288A;$!c*r   c                     t         j                  d        }t        j                  d      d        }| j                  d|j                         | j                  d|j                         y )Nc                      y) Add two integers, kernel versionNr   abs     r   r   z<TestDispatcher.test_dispatcher_docstring.<locals>.add_kernel      r   Tr   c                      y) 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  sg     
	3 
	3 
		3 
	3 	;Z=O=OP;Z=O=O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58 233  43 23( 4( 9:= ;=8 CD! E!4;.--
--$:$ >?5 @52 ;<+  =+(	;::"+ <=+ >+,Qr   rU   z2CUDA simulator doesn't implement kernel propertiesc                   B    e Zd Zd Zd Zd Zd Zd Zd Zd Z	d Z
d	 Zy
)TestDispatcherKernelPropertiesc                 H   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              }|j                  |      }|j                  |      }| j                  |t               | j                  |t               | j                  |d       | j                  |d       |j                         }	| j                  |	|j                     |       | j                  |	|j                     |       ||k(  r+t        d       t        d       t        j                          y y )Nc                 v    t        j                  d      }||k  rdt        j                  | |         z  | |<   y y 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  5    		!A1udhhqtn,! r   
   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{    
	- 
	- ((1BJJ/((1BJJ/QT7A&QT7A& wss|U+wss|U+*>>wG*>>wG137137.2.2
 +>>@,W\\:,	.,W\\:,	. "55 @AABKKM 6r   c                     t        j                  t        t        d d d   t                    d        }|j                         }| j                  |t               | j                  |d       y )Nr/   c                 v    t        j                  d      }||k  rdt        j                  | |         z  | |<   y y 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\    	$wss|U+	,	- 
-	- '::<os3?A.r   c                 0   t         j                  d        } |d   dd        |d   dd       t        t        t              }t        t
        t              }|j                  |      }|j                  |      }| j                  |t               | j                  |t               | j                  |d       | j                  |d       |j                         }| j                  ||j                     |       | j                  ||j                     |       y )Nc                      |rt        |        y y 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  s   		 
	 	q%(sE* ug&w(-@@I-@@I0#60#6 	 2A6 2A6 .@@B+GLL9;MN+GLL9;MNr   c                 :   t        j                  dt         j                        t        t        d d d         }t	        j
                  |      fd       }|j                  |      }| j                  |t               | j                  |j                         y )N    rW   r/   c                 ~    t         j                  j                        }t        j                  d      }||   | |<   y r   )r   const
array_liker   )r   Cr   arrs      r   const_array_usezVTestDispatcherKernelProperties.test_get_const_mem_specialized.<locals>.const_array_use,  s1    

%%c*A		!AQ4AaDr   )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(  sy    ii"((+51:	#	 
	
 );;C@nc2

;r   c                   
 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         }|j                  |      }|j                  |      }| j                  |t               | j                  |t               | j                  |
dz         | j                  |
dz         |j                         }|j                         }	| j                  ||j                     |       | j                  |	|j                     |       y )Nr   c                     t         j                  j                  | j                        }t	              D ]  }|||<   	 t	              D ]
  }||   | |<    y NrW   )r   sharedarrayrX   r   )arysmjr   s      r   simple_smemz_TestDispatcherKernelProperties.test_get_shared_mem_per_block_unspecialized.<locals>.simple_smem;  sR    ""1CII"6B1X1 1X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  sD    
	 
	 ((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                     t        j                  t        t        d d d               d        }|j	                         }| j                  |t               | j                  |d       y )Nr/   c                     t         j                  j                  dt              }t        j                  d      }|dk(  rt        d      D ]  }|||<   	 t        j                          ||   | |<   y )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`  s_    ""3g"6B		!AAvsABqE $UCFr   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]    	$wss|$	%	 
&	  +CCE2C8-s3r   c                    d}t         j                  d        }t        j                  |t        j                        } |d   |       t        t        d d d         }|j                  |      }| j                  |t               | j                  |d       |j                         }| j                  ||j                     |       y )Nr   c                 8    t        j                  d      }|| |<   y r   )r   r   )r	  r   s     r   simple_maxthreadszfTestDispatcherKernelProperties.test_get_max_threads_per_block_unspecialized.<locals>.simple_maxthreadsq  s    		!ACFr   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    		 
	 ((1BJJ/$(wss|$+EEgNos3?A./IIK,W\\:OLr   c                   	 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         }|j                  |      }|j                  |      }| j                  |t               | j                  |t               | j                  |	dz         | j                  |	dz         |j                         }| j                  ||j                     |       | j                  ||j                     |       y )N  c                     t         j                  j                  | j                        }t	              D ]  }|||<   	 t	              D ]
  }||   | |<    y r  r   localr  rX   r   r	  lmr  r   s      r   simple_lmemz_TestDispatcherKernelProperties.test_get_local_mem_per_thread_unspecialized.<locals>.simple_lmem  R    !!!399!5B1X1 1X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  s6    		 
	 ((1BJJ/((1BJJ/D'"D'"wss|$wss|$#<<WE#<<WEmS1mS1q1u5q1u5
 $<<>w||4mDw||4mDr   c                     dt        j                  t        t        d d d               fd       }|j	                         }| j                  |t               | j                  |dz         y )Nr"  r/   c                     t         j                  j                  | j                        }t	              D ]  }|||<   	 t	              D ]
  }||   | |<    y 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  si     	$wss|$	%	 
&	  +CCE2C8 4a!e<r   N)rQ   rR   rS   r   r   r   r  r  r  r   r.  r2  r   r   r   r   r     s7    -^
/!OF<'CR4M&%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>r9     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HMMO r   