+
    :i8                     (   ^ RI HtHtHtHt ^ RIHt ^ RIHt ^ RI	H
t
HtHt ^ RIt^ RIHt ^RIHtHt ]P*                  ! R]P                  3R	]P,                  R3.4      t ! R
 R]4      t ! R R]4      t]R8X  d   ]
P6                  ! 4        R# R# )    )cudaint32float64void)TypingError)types)unittestCUDATestCaseskip_on_cudasimN)numpy_support)test_struct_model_type
TestStructijc                   P   a  ] tR t^t o R tR tR tR tR tR t	R t
R tR	tV tR
# )TestSharedMemoryIssuec                   a \         P                  ! R R7      R 4       o\         P                  V3R l4       pVR,          ! 4        R# )T)devicec                  P    \         P                  P                  ^\        R7      p R#    dtypeNr   sharedarrayr   )	inner_arrs    ]/var/www/html/photoedit/myenv/lib/python3.14/site-packages/numba/cuda/tests/cudapy/test_sm.pyinnerGTestSharedMemoryIssue.test_issue_953_sm_linkage_conflict.<locals>.inner   s    ))!5)9I    c                  `   < \         P                  P                  ^\        R7      p S! 4        R# r   r   )	outer_arrr   s    r   outerGTestSharedMemoryIssue.test_issue_953_sm_linkage_conflict.<locals>.outer   s!    ))!5)9IGr!   Nr   r   )r   jit)selfr$   r   s   & @r   "test_issue_953_sm_linkage_conflict8TestSharedMemoryIssue.test_issue_953_sm_linkage_conflict   sA    			: 
	: 
	 
	 	dr!   c                   a \         P                  V3R  l4       p\        P                  ! ^\        P                  R7      pVR,          ! V4       V P                  V^ ,          V4       R# )c                 p   < \         P                  P                  S\        R 7      pVP                  V ^ &   R# r   N)r   r   r   r   size)aarrshapes   & r   s9TestSharedMemoryIssue._check_shared_array_size.<locals>.s   s)    ++##E#7C88AaDr!   r   Nr&   )r   r'   npzerosr   assertEqual)r(   r1   expectedr2   results   &f&  r   _check_shared_array_size.TestSharedMemoryIssue._check_shared_array_size   sN    		 
	 !288,	$H-r!   c                *    V P                  ^^4       R#    Nr9   r(   s   &r   %test_issue_1051_shared_size_broken_1d;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_1d&   s    %%a+r!   c                *    V P                  R^4       R# )r=   N)r=      r>   r?   s   &r   %test_issue_1051_shared_size_broken_2d;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_2d)   s    %%fa0r!   c                *    V P                  R^4       R# )r=   N)r=   rC      r>   r?   s   &r   %test_issue_1051_shared_size_broken_3d;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_3d,   s    %%i4r!   c                   aa \         P                  VV3R  l4       p\        P                  ! ^\        P                  R7      pVR,          ! V4       V P                  V^ ,          V4       R# )c                 h   < \         P                  P                  SSR 7      pVP                  V ^ &   R# r-   )r   r   r   r.   )r/   r0   r1   tys   & r   r2   >TestSharedMemoryIssue._check_shared_array_size_fp16.<locals>.s1   s)    ++##E#4C88AaDr!   r   Nr&   )r   r'   r4   r5   float16r6   )r(   r1   r7   rL   r2   r8   s   &f&f  r   _check_shared_array_size_fp163TestSharedMemoryIssue._check_shared_array_size_fp160   sN    		 
	 !2::.	$H-r!   c                    V P                  ^^\        P                  4       V P                  ^^\        P                  4       R# r<   )rO   r   rN   r4   r?   s   &r   test_issue_fp16_support-TestSharedMemoryIssue.test_issue_fp16_support:   s.    **1a?**1a<r!   c                  aa ^o^0p^o^p\         P                  VV3R l4       p\        P                  ! V\        P                  R7      p\         P
                  ! V4      pW1V3,          ! V4       \         P                  ! 4        R# )z
Test issue of warp misalign address due to nvvm not knowing the
alignment(? but it should have taken the natural alignment of the type)
c                 >  < \         P                  P                  SS3\        4      p\         P                  P                  ^\        4      p\         P                  P
                  p^ p\        S4       F  pWAW53,          ,          pK  	  V^ ,          V,           V ^ &   R# )   N)r   r   r   r   	threadIdxxrange)d_block_costs
s_featuress_initialcostrW   
predictionr   examples_per_blocknum_weightss   &     r   
costs_func9TestSharedMemoryIssue.test_issue_2393.<locals>.costs_funcH   s    **,>+L+24J KK--a9M((IJ;'66
 (  -Q/*<M!r!   r   N)r   r'   r4   r5   r   	to_devicesynchronize)r(   
num_blocksthreads_per_blockr`   block_costsrZ   r^   r_   s   &     @@r   test_issue_2393%TestSharedMemoryIssue.test_issue_2393>   sr    
 
		= 
	= hhz<{3001-@r!    N)__name__
__module____qualname____firstlineno__r)   r9   r@   rD   rH   rO   rR   rg   __static_attributes____classdictcell____classdict__s   @r   r   r      s2     
.,15.= r!   r   c                      a  ] tR t^^t o R tR tR tR tR tR t	R t
R tR	 tR
 tR tR tR t]! R4      R 4       t]! R4      R 4       tRtV tR# )TestSharedMemoryc                f  aa \        V4      p^o\        VS,          4      p\        P                  ! VP                  4      o\
        P                  VV3R l4       p\
        P                  ! V4      pWCS3,          ! W4       VP                  4       p\        P                  P                  W4       R# )   c                   < \         P                  P                  S	SR 7      p\         P                  P                  p\         P
                  P                  p\         P                  P                  pWE,          V,           pV\        V 4      8  d   W,          W#&   \         P                  ! 4        V^ 8X  d-   \        S	4       F  pW',          WV,          V,           &   K  	  R# R# r-   
r   r   r   rW   rX   blockIdxblockDimlensyncthreadsrY   )
rX   ysmtxbxbdr   r   dtnthreadss
   &&      r   use_sm_chunk_copy8TestSharedMemory._test_shared.<locals>.use_sm_chunk_copyj   s    ""82"6B!!BBB "A3q6z QwxA%'UA2gkN ) r!   N)rz   intnps
from_dtyper   r   r'   device_array_likecopy_to_hostr4   testingassert_array_equal)	r(   r0   nelemnblocksr   d_resulthost_resultr   r   s	   &&     @@r   _test_sharedTestSharedMemory._test_shared_   s     Ceh&'^^CII&		+ 
	+& ))#.8+,S;++-


%%c7r!   c                D   \         P                  ! ^\        R7      p\        \	        V4      4       FZ  pW!V,          n        \         P                  ! ^\         P                  R7      pVP                  ^^4      V,          W,          n	        K\  	  V P                  V4       R# )   r   N)r4   recarrayrecordwith2darrayrY   rz   r   arangefloat32reshaper   r   )r(   r0   rX   r   s   &   r   test_shared_recarray%TestSharedMemory.test_shared_recarray   sj    kk#%67s3xAFH		%rzz2AyyA*CFH !
 	#r!   c                    \         P                  P                  ^R\         P                  R7      pV P	                  V4       R# )r=   )r.   r   N)   )r4   randomrandintbool_r   )r(   r0   s   & r   test_shared_bool!TestSharedMemory.test_shared_bool   s/    iirxx@#r!   c                    VP                   VP                  P                  ,          pV^^^ V3,          ! V4       \        P                  P                  W24       R# )r   N)r.   r   itemsizer4   r   r   )r(   funcr0   r7   nshareds   &&&& r   _test_dynshared_slice&TestSharedMemory._test_dynshared_slice   sD    
 ((SYY///Q1gs#


%%h4r!   c                    \         P                  R  4       p\        P                  ! ^\        P                  R7      p\        P
                  ! ^^.\        P                  R7      pV P                  WV4       R# )c                     \         P                  P                  ^ \        R7      pVR,          pVR,          p^V^ &   ^V^ &   V^ ,          V ^ &   V^,          V ^&   R# r   r   r   r   N:r   r=   NNr   rX   dynsmemsm1sm2s   &   r   slice_write@TestSharedMemory.test_dynshared_slice_write.<locals>.slice_write   sY    kk'''7G#,C#,CCFCF1:AaD1:AaDr!   r   Nr   r'   r4   r5   r   r   r   )r(   r   r0   r7   s   &   r   test_dynshared_slice_write+TestSharedMemory.test_dynshared_slice_write   sV    		 
	 hhq)88QF"((3"";X>r!   c                    \         P                  R  4       p\        P                  ! ^\        P                  R7      p\        P
                  ! ^^.\        P                  R7      pV P                  WV4       R# )c                     \         P                  P                  ^ \        R7      pVR,          pVR,          p^V^ &   ^V^&   V^ ,          V ^ &   V^ ,          V ^&   R# r   r   r   s   &   r   
slice_read>TestSharedMemory.test_dynshared_slice_read.<locals>.slice_read   sY    kk'''7G#,C#,CGAJGAJq6AaDq6AaDr!   r   Nr   )r(   r   r0   r7   s   &   r   test_dynshared_slice_read*TestSharedMemory.test_dynshared_slice_read   sV    		 
	 hhq)88QF"((3"":H=r!   c                    \         P                  R  4       p\        P                  ! ^\        P                  R7      p\        P
                  ! . RO\        P                  R7      pV P                  WV4       R# )c                     \         P                  P                  ^ \        R7      pVR,          pVR,          p^V^ &   ^V^&   ^V^&   V^ ,          V ^ &   V^ ,          V ^&   V^,          V ^&   R# )r   r   r   r   rC   NNr   r   s   &   r   slice_diff_sizesJTestSharedMemory.test_dynshared_slice_diff_sizes.<locals>.slice_diff_sizes   so    kk'''7G#,C#,CGAJGAJGAJq6AaDq6AaDq6AaDr!   r   N)r   r=   rC   r   )r(   r   r0   r7   s   &   r   test_dynshared_slice_diff_sizes0TestSharedMemory.test_dynshared_slice_diff_sizes   sU     

	 

	 hhq)88IRXX6""#3(Cr!   c                    \         P                  R  4       p\        P                  ! ^\        P                  R7      p\        P
                  ! . RO\        P                  R7      pV P                  WV4       R# )c                    \         P                  P                  ^ \        R7      pVR,          pVR,          p^V^ &   ^V^&   ^V^&   ^V^&   V^ ,          V ^ &   V^,          V ^&   V^ ,          V ^&   V^,          V ^&   V^,          V ^&   R# )r   r   :r   r=   N:r   rG   NNr   r   s   &   r   slice_overlapDTestSharedMemory.test_dynshared_slice_overlap.<locals>.slice_overlap   s    kk'''7G#,C#,CGAJGAJGAJGAJq6AaDq6AaDq6AaDq6AaDq6AaDr!   r   N)r   r=   r=   rC   rG   r   )r(   r   r0   r7   s   &   r   test_dynshared_slice_overlap-TestSharedMemory.test_dynshared_slice_overlap   sR    		 
	 hhq)88O288<""=x@r!   c                    \         P                  R  4       p\        P                  ! ^\        P                  R7      p\        P
                  ! . RO\        P                  R7      pV P                  WV4       R# )c                    \         P                  P                  ^ \        R7      pVR,          pVR,          p^cV^ &   ^cV^&   ^cV^&   ^cV^&   ^cV^&   ^cV^&   ^cV^&   ^V^ &   ^V^&   ^V^ &   ^V^&   V^ ,          V ^ &   V^,          V ^&   V^,          V ^&   V^,          V ^&   V^,          V ^&   V^,          V ^&   V^,          V ^&   R# )r   r   r   :rG      NNr   r   s   &   r   
slice_gaps>TestSharedMemory.test_dynshared_slice_gaps.<locals>.slice_gaps   s    kk'''7G#,C#,C GAJGAJGAJGAJGAJGAJGAJCFCFCFCF1:AaD1:AaD1:AaD1:AaD1:AaD1:AaD1:AaDr!   r   N)c   r   r=   r   rC   rG   r   r   )r(   r   r0   r7   s   &   r   test_dynshared_slice_gaps*TestSharedMemory.test_dynshared_slice_gaps   sU     
	 
	6 hhq)884BHHE"":H=r!   c                    \         P                  R  4       p\        P                  ! ^\        P                  R7      p\        P
                  ! . RO\        P                  R7      pV P                  WV4       R# )c                    \         P                  P                  ^ \        R7      pV^RR1,          pV^^R1,          p^V^ &   ^V^&   ^V^ &   ^V^&   V^ ,          V ^ &   V^,          V ^&   V^,          V ^&   V^,          V ^&   R# )r   r   Nr   r   s   &   r   slice_write_backwardsTTestSharedMemory.test_dynshared_slice_write_backwards.<locals>.slice_write_backwards  s    kk'''7G!%R%.C!Ab&/CCFCFCFCF1:AaD1:AaD1:AaD1:AaDr!   r   N)r=   r   rG   rC   r   )r(   r   r0   r7   s   &   r   $test_dynshared_slice_write_backwards5TestSharedMemory.test_dynshared_slice_write_backwards  sU     
	 
	 hhq)88L9""#8xHr!   c                    \         P                  R  4       p\        P                  ! ^\        P                  R7      p\        P
                  ! . RO\        P                  R7      pV P                  WV4       R# )c                 L   \         P                  P                  ^ \        R7      pVR,          p^cV^ &   ^cV^&   ^cV^&   ^cV^&   ^cV^&   ^cV^&   ^V^ &   ^V^&   ^V^&   V^ ,          V ^ &   V^,          V ^&   V^,          V ^&   V^,          V ^&   V^,          V ^&   V^,          V ^&   R# )r   r   :NNr=   Nr   rX   r   r   s   &  r   slice_nonunit_strideRTestSharedMemory.test_dynshared_slice_nonunit_stride.<locals>.slice_nonunit_stride!  s    kk'''7G#,C GAJGAJGAJGAJGAJGAJCFCFCF1:AaD1:AaD1:AaD1:AaD1:AaD1:AaDr!   r   N)r   r   r=   r   rC   r   r   )r(   r   r0   r7   s   &   r   #test_dynshared_slice_nonunit_stride4TestSharedMemory.test_dynshared_slice_nonunit_stride  sV     
	 
	. hhq)881B""#7hGr!   c                    \         P                  R  4       p\        P                  ! ^\        P                  R7      p\        P
                  ! . RO\        P                  R7      pV P                  WV4       R# )c                 R   \         P                  P                  ^ \        R7      pVRRR1,          p^cV^ &   ^cV^&   ^cV^&   ^cV^&   ^cV^&   ^cV^&   ^V^ &   ^V^&   ^V^&   V^ ,          V ^ &   V^,          V ^&   V^,          V ^&   V^,          V ^&   V^,          V ^&   V^,          V ^&   R# )r   r   Nr   r   r   s   &  r   slice_nonunit_reverse_stridebTestSharedMemory.test_dynshared_slice_nonunit_reverse_stride.<locals>.slice_nonunit_reverse_stride@  s    kk'''7G"&b&/C GAJGAJGAJGAJGAJGAJCFCFCF1:AaD1:AaD1:AaD1:AaD1:AaD1:AaDr!   r   N)r   rC   r   r=   r   r   r   )r(   r   r0   r7   s   &   r   +test_dynshared_slice_nonunit_reverse_stride<TestSharedMemory.test_dynshared_slice_nonunit_reverse_stride=  sV     
	 
	. hhq)881B""#?hOr!   c                  a
 \         P                  ! R 4      p\        V4      p^p\        W#,          4      p\        P
                  ! VP                  4      o
W1P                  P                  ,          p\        V^,          4      p\        P                  V
3R l4       p\        P                  ! V4      pWtV^ V3,          ! WV4       VP                  4       p	\         P                  P                  W4       R# )r   c                 V  < \         P                  P                  ^ SR7      pV^ V pW2V^,           p\         P                  P                  p\         P
                  P                  p\         P                  P                  pWx,          V,           p	V	\        V 4      8  d#   Wb8  d   W	,          WF&   MW	,          WVV,
          &   \         P                  ! 4        V^ 8X  dL   \        V4       F:  p
WJ,          WV,          V
,           &   WZ,          WV,          V
,           V,           &   K<  	  R# R# )r   r   Nrw   )rX   r|   	chunksizer   r   r   r~   r   r   r   r   r   s   &&&        r   sm_slice_copy7TestSharedMemory.test_issue_5073.<locals>.sm_slice_copyk  s    kk'''4G!I&CIM2C!!BBB "A3q6z>dCG*+$CY' Qwy)A%(VA2gkN14A2gkI-. * r!   N)r4   r   rz   r   r   r   r   r   r   r'   r   r   r   r   )r(   r0   r   r   r   r   r   r   r   r   r   s   &         @r   test_issue_5073 TestSharedMemory.test_issue_5073\  s     iioCe&'^^CII&YY///1%			8 
	82 ))#.xG34SIN++-


%%c7r!   zCan't check typing in simulatorc                   R pR pV P                  \        V4      ;_uu_ 4        \        P                  ! \	        4       4      ! V4       RRR4       RpR pV P                  \        V4      ;_uu_ 4        \        P                  ! \	        4       4      ! V4       RRR4       R#   + '       g   i     Lb; i  + '       g   i     R# ; i)z+.*Cannot infer the type of variable 'arr'.*c                  p    \         P                  P                  ^
\        P                  ! R4      R7      p R# )
   Or   N)r   r   r   r4   r   r0   s    r   unsupported_typeBTestSharedMemory.test_invalid_array_type.<locals>.unsupported_type  s#    ++##Bbhhsm#<Cr!   Nz*.*Invalid NumPy dtype specified: 'int33'.*c                  H    \         P                  P                  ^
RR7      p R# )r   int33r   N)r   r   r   r   s    r   invalid_string_typeETestSharedMemory.test_invalid_array_type.<locals>.invalid_string_type  s    ++##Bg#6Cr!   )assertRaisesRegexr   r   r'   r   )r(   rgxr   r   s   &   r   test_invalid_array_type(TestSharedMemory.test_invalid_array_type  s    ;	=##K55HHTV-. 6 ;	7##K55HHTV01 65 65 655s   %B#4%B6#B3	6C	z+Struct model array unsupported in simulatorc                  a ^@o\         P                  ! \        \        R,          \        R,          4      4      V3R l4       p\        P
                  ! S3RR7      p\        P
                  ! S3RR7      pV^S3,          ! W#4       \        V4       F%  w  rEV P                  VSV,
          ^,
          4       K'  	  \        V4       F,  w  rFV P                  VSV,
          ^,
          ^,          4       K.  	  R# )@   :NNr   c                   < \         P                  P                  S\        R 7      p\         P                  ! ^4      pSV,
          ^,
          pV\        V 4      8  d{   V\        V4      8  di   \        \        V4      \        V^,          4      4      pWRV&   \         P                  ! 4        W$,          P                  W&   W$,          P                  W&   R# R# R# r-   )r   r   r   r   gridrz   r   r   r{   rX   r|   )outxoutyr0   r   riobjr   s   &&    r   write_then_reverse_read_staticVTestSharedMemory.test_struct_model_type_static.<locals>.write_then_reverse_read_static  s     ++##H4J#KC		!AA!B3t9}SY q5Q<8A  "'))')) "/}r!   r   r   N)r   r'   r   r   r4   r5   	enumerater6   )r(   r  arrxarryr   rX   r|   r   s   &      @r   test_struct_model_type_static.TestSharedMemory.test_struct_model_type_static  s    	$uSz5:.	/	$ 
0	$" xx73xx73&q({3D?dODAQ1q 01 $dODAQA!1Q 67 $r!   ri   N)rj   rk   rl   rm   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r  rn   ro   rp   s   @r   rs   rs   ^   s     "8H5?">"D(A,!>FI,H>P>,8\ 672 82 BC8 D8r!   rs   __main__)rC   r=   )numbar   r   r   r   numba.core.errorsr   
numba.corer   numba.cuda.testingr	   r
   r   numpyr4   numba.npr   r   extensions_usecasesr   r   r   r   r   r   rs   rj   mainri   r!   r   <module>r     s    , , )  F F  ) CHHsBHHo"BJJ79 : LL L^Z8| Z8z
 zMMO r!   