
    sh8                     "   S SK JrJrJrJr  S SKJr  S SKJr  S SK	J
r
JrJr  S SKrS SKJr  SSKJrJr  \R*                  " S	\R                  4S
\R,                  S4/5      r " S S\5      r " S S\5      r\S:X  a  \
R6                  " 5         gg)    )cudaint32float64void)TypingError)types)unittestCUDATestCaseskip_on_cudasimN)numpy_support   )test_struct_model_type
TestStructij)      c                   D    \ rS rSrS rS rS rS rS rS r	S r
S	 rS
rg)TestSharedMemoryIssue   c                    ^ [         R                  " SS9S 5       m[         R                  U4S j5       nUS   " 5         g )NT)devicec                  H    [         R                  R                  S[        S9n g Nr   dtyper   sharedarrayr   )	inner_arrs    s/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/tests/cudapy/test_sm.pyinnerGTestSharedMemoryIssue.test_issue_953_sm_linkage_conflict.<locals>.inner   s    ))!5)9I    c                  X   > [         R                  R                  S[        S9n T" 5         g r   r   )	outer_arrr"   s    r!   outerGTestSharedMemoryIssue.test_issue_953_sm_linkage_conflict.<locals>.outer   s!    ))!5)9IGr$   r   r   )r   jit)selfr'   r"   s     @r!   "test_issue_953_sm_linkage_conflict8TestSharedMemoryIssue.test_issue_953_sm_linkage_conflict   sA    			: 
	: 
	 
	 	dr$   c                    ^ [         R                  U4S j5       n[        R                  " S[        R                  S9nUS   " U5        U R                  US   U5        g )Nc                 h   > [         R                  R                  T[        S9nUR                  U S'   g Nr   r   )r   r   r   r   size)aarrshapes     r!   s9TestSharedMemoryIssue._check_shared_array_size.<locals>.s   s)    ++##E#7C88AaDr$   r   r   r)   r   )r   r*   npzerosr   assertEqual)r+   r4   expectedr5   results    `   r!   _check_shared_array_size.TestSharedMemoryIssue._check_shared_array_size   sN    		 
	 !288,	$H-r$   c                 (    U R                  SS5        g Nr   r<   r+   s    r!   %test_issue_1051_shared_size_broken_1d;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_1d&   s    %%a+r$   c                 (    U R                  SS5        g )N)r   r      r@   rA   s    r!   %test_issue_1051_shared_size_broken_2d;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_2d)   s    %%fa0r$   c                 (    U R                  SS5        g )N)r   r         r@   rA   s    r!   %test_issue_1051_shared_size_broken_3d;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_3d,   s    %%i4r$   c                    ^^ [         R                  UU4S j5       n[        R                  " S[        R                  S9nUS   " U5        U R                  US   U5        g )Nc                 `   > [         R                  R                  TTS9nUR                  U S'   g r0   )r   r   r   r1   )r2   r3   r4   tys     r!   r5   >TestSharedMemoryIssue._check_shared_array_size_fp16.<locals>.s1   s)    ++##E#4C88AaDr$   r   r   r)   r   )r   r*   r7   r8   float16r9   )r+   r4   r:   rO   r5   r;   s    ` `  r!   _check_shared_array_size_fp163TestSharedMemoryIssue._check_shared_array_size_fp160   sN    		 
	 !2::.	$H-r$   c                     U R                  SS[        R                  5        U R                  SS[        R                  5        g r?   )rR   r   rQ   r7   rA   s    r!   test_issue_fp16_support-TestSharedMemoryIssue.test_issue_fp16_support:   s.    **1a?**1a<r$   c                    ^^ SmSnSmSn[         R                  UU4S j5       n[        R                  " U[        R                  S9n[         R
                  " U5      nX1U4   " U5        [         R                  " 5         g)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)
r   0   rI   r   c                   > [         R                  R                  TT4[        5      n[         R                  R                  S[        5      n[         R                  R
                  nSn[        T5       H  nXAX54   -  nM     US   U-   U S'   g )N   r   )r   r   r   r   	threadIdxxrange)d_block_costs
s_featuress_initialcostr[   
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*   r7   r8   r   	to_devicesynchronize)r+   
num_blocksthreads_per_blockrd   block_costsr^   rb   rc   s         @@r!   test_issue_2393%TestSharedMemoryIssue.test_issue_2393>   sr    
 
		= 
	= hhz<{3001-@r$    N)__name__
__module____qualname____firstlineno__r,   r<   rB   rF   rK   rR   rU   rk   __static_attributes__rm   r$   r!   r   r      s*    
.,15.=r$   r   c                       \ rS rSrS rS rS rS rS rS r	S r
S	 rS
 rS rS rS rS r\" S5      S 5       r\" S5      S 5       rSrg)TestSharedMemory^   c                 T  ^^ [        U5      nSm[        UT-  5      n[        R                  " UR                  5      m[
        R                  UU4S j5       n[
        R                  " U5      nXCT4   " X5        UR                  5       n[        R                  R                  X5        g )N   c                   > [         R                  R                  T	TS9n[         R                  R                  n[         R
                  R                  n[         R                  R                  nXE-  U-   nU[        U 5      :  a  X   X#'   [         R                  " 5         US:X  a  [        T	5       H  nX'   XU-  U-   '   M     g g r0   
r   r   r   r[   r\   blockIdxblockDimlensyncthreadsr]   )
r\   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$   )r|   intnps
from_dtyper   r   r*   device_array_likecopy_to_hostr7   testingassert_array_equal)	r+   r3   nelemnblocksr   d_resulthost_resultr   r   s	          @@r!   _test_sharedTestSharedMemory._test_shared_   s     Ceh&'^^CII&		+ 
	+& ))#.8+,S;++-


%%c7r$   c                    [         R                  " S[        S9n[        [	        U5      5       HK  nX!U   l        [         R                  " S[         R                  S9nUR                  SS5      U-  X   l	        MM     U R                  U5        g )N   r   rE   r   r   )r7   recarrayrecordwith2darrayr]   r|   r   arangefloat32reshaper   r   )r+   r3   r\   r   s       r!   test_shared_recarray%TestSharedMemory.test_shared_recarray   sj    kk#%67s3xAFH		%rzz2AyyA*CFH !
 	#r$   c                     [         R                  R                  SS[         R                  S9nU R	                  U5        g )Nr   )   )r1   r   )r7   randomrandintbool_r   )r+   r3   s     r!   test_shared_bool!TestSharedMemory.test_shared_bool   s/    iirxx@#r$   c                     UR                   UR                  R                  -  nUSSSU4   " U5        [        R                  R                  X25        g )Nr   r   )r1   r   itemsizer7   r   r   )r+   funcr3   r:   nshareds        r!   _test_dynshared_slice&TestSharedMemory._test_dynshared_slice   sD    
 ((SYY///Q1gs#


%%h4r$   c                     [         R                  S 5       n[        R                  " S[        R                  S9n[        R
                  " SS/[        R                  S9nU R                  XU5        g )Nc                     [         R                  R                  S[        S9nUSS nUSS nSUS'   SUS'   US   U S'   US   U S'   g Nr   r   r   r   r   r\   dynsmemsm1sm2s       r!   slice_write@TestSharedMemory.test_dynshared_slice_write.<locals>.slice_write   s]    kk'''7G!A,C!A,CCFCF1:AaD1:AaDr$   r   r   r   r   r*   r7   r8   r   r   r   )r+   r   r3   r:   s       r!   test_dynshared_slice_write+TestSharedMemory.test_dynshared_slice_write   sV    		 
	 hhq)88QF"((3"";X>r$   c                     [         R                  S 5       n[        R                  " S[        R                  S9n[        R
                  " SS/[        R                  S9nU R                  XU5        g )Nc                     [         R                  R                  S[        S9nUSS nUSS nSUS'   SUS'   US   U S'   US   U S'   g r   r   r   s       r!   
slice_read>TestSharedMemory.test_dynshared_slice_read.<locals>.slice_read   s]    kk'''7G!A,C!A,CGAJGAJq6AaDq6AaDr$   r   r   r   r   )r+   r   r3   r:   s       r!   test_dynshared_slice_read*TestSharedMemory.test_dynshared_slice_read   sV    		 
	 hhq)88QF"((3"":H=r$   c                     [         R                  S 5       n[        R                  " S[        R                  S9n[        R
                  " / SQ[        R                  S9nU R                  XU5        g )Nc                     [         R                  R                  S[        S9nUSS nUSS nSUS'   SUS'   SUS'   US   U S'   US   U S'   US   U S'   g )Nr   r   r   r   r   r   r   s       r!   slice_diff_sizesJTestSharedMemory.test_dynshared_slice_diff_sizes.<locals>.slice_diff_sizes   ss    kk'''7G!A,C!A,CGAJGAJGAJq6AaDq6AaDq6AaDr$   r   r   )r   r   r   r   )r+   r   r3   r:   s       r!   test_dynshared_slice_diff_sizes0TestSharedMemory.test_dynshared_slice_diff_sizes   sU     

	 

	 hhq)88IRXX6""#3(Cr$   c                     [         R                  S 5       n[        R                  " S[        R                  S9n[        R
                  " / SQ[        R                  S9nU R                  XU5        g )Nc                     [         R                  R                  S[        S9nUSS nUSS nSUS'   SUS'   SUS'   SUS'   US   U S'   US   U S'   US   U S'   US   U S'   US   U S'   g )Nr   r   r   r   rI   r   r   r   s       r!   slice_overlapDTestSharedMemory.test_dynshared_slice_overlap.<locals>.slice_overlap   s    kk'''7G!A,C!A,CGAJGAJGAJGAJq6AaDq6AaDq6AaDq6AaDq6AaDr$      r   )r   r   r   r   rI   r   )r+   r   r3   r:   s       r!   test_dynshared_slice_overlap-TestSharedMemory.test_dynshared_slice_overlap   sR    		 
	 hhq)88O288<""=x@r$   c                     [         R                  S 5       n[        R                  " S[        R                  S9n[        R
                  " / SQ[        R                  S9nU R                  XU5        g )Nc                 :   [         R                  R                  S[        S9nUSS nUSS nSUS'   SUS'   SUS'   SUS'   SUS'   SUS	'   SUS'   SUS'   SUS'   SUS'   SUS'   US   U S'   US   U S'   US   U S'   US   U S'   US   U S'   US	   U S	'   US   U S'   g )
Nr   r   r   r   rI   rE   c   r   r   r   r   s       r!   
slice_gaps>TestSharedMemory.test_dynshared_slice_gaps.<locals>.slice_gaps   s    kk'''7G!A,C!A,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$   rZ   r   )r   r   r   r   r   rI   r   r   )r+   r   r3   r:   s       r!   test_dynshared_slice_gaps*TestSharedMemory.test_dynshared_slice_gaps   sU     
	 
	6 hhq)884BHHE"":H=r$   c                     [         R                  S 5       n[        R                  " S[        R                  S9n[        R
                  " / SQ[        R                  S9nU R                  XU5        g )Nc                     [         R                  R                  S[        S9nUSS S2   nUSSS2   nSUS'   SUS'   SUS'   SUS'   US   U S'   US   U S'   US   U S'   US   U S'   g )Nr   r   r   r   r   rI   r   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$   rI   r   )r   r   rI   r   r   )r+   r   r3   r:   s       r!   $test_dynshared_slice_write_backwards5TestSharedMemory.test_dynshared_slice_write_backwards  sU     
	 
	 hhq)88L9""#8xHr$   c                     [         R                  S 5       n[        R                  " S[        R                  S9n[        R
                  " / SQ[        R                  S9nU R                  XU5        g )Nc                    [         R                  R                  S[        S9nUS S S2   nSUS'   SUS'   SUS'   SUS'   SUS'   SUS'   SUS'   SUS'   SUS'   US   U S'   US   U S'   US   U S'   US   U S'   US   U S'   US   U S'   g )	Nr   r   r   r   r   r   rI   r   r   r\   r   r   s      r!   slice_nonunit_strideRTestSharedMemory.test_dynshared_slice_nonunit_stride.<locals>.slice_nonunit_stride!  s    kk'''7G#A#,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$   rE   r   )r   r   r   r   r   r   r   )r+   r   r3   r:   s       r!   #test_dynshared_slice_nonunit_stride4TestSharedMemory.test_dynshared_slice_nonunit_stride  sV     
	 
	. hhq)881B""#7hGr$   c                     [         R                  S 5       n[        R                  " S[        R                  S9n[        R
                  " / SQ[        R                  S9nU R                  XU5        g )Nc                    [         R                  R                  S[        S9nUSS S2   nSUS'   SUS'   SUS'   SUS'   SUS	'   SUS
'   SUS'   SUS'   SUS'   US   U S'   US   U S'   US   U S'   US   U S'   US	   U S	'   US
   U S
'   g )Nr   r   r   r   r   r   r   rI   r   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$   rE   r   )r   r   r   r   r   r   r   )r+   r   r3   r:   s       r!   +test_dynshared_slice_nonunit_reverse_stride<TestSharedMemory.test_dynshared_slice_nonunit_reverse_stride=  sV     
	 
	. hhq)881B""#?hOr$   c                   ^
 [         R                  " S5      n[        U5      nSn[        X#-  5      n[        R
                  " UR                  5      m
X1R                  R                  -  n[        US-  5      n[        R                  U
4S j5       n[        R                  " U5      nXtUSU4   " XU5        UR                  5       n	[         R                  R                  X5        g )Nr   rw   r   c                   > [         R                  R                  STS9nUSU nX2US-   n[         R                  R                  n[         R
                  R                  n[         R                  R                  nXx-  U-   n	U	[        U 5      :  a  Xb:  a  X	   XF'   O	X	   XVU-
  '   [         R                  " 5         US:X  a.  [        U5       H  n
XJ   XU-  U
-   '   XZ   XU-  U
-   U-   '   M      g g )Nr   r   r   ry   )r\   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$   r   )r7   r   r|   r   r   r   r   r   r   r*   r   r   r   r   )r+   r3   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                 h   SnS nU R                  [        U5         [        R                  " [	        5       5      " U5        S S S 5        SnS nU R                  [        U5         [        R                  " [	        5       5      " U5        S S S 5        g ! , (       d  f       NW= f! , (       d  f       g = f)Nz+.*Cannot infer the type of variable 'arr'.*c                  h    [         R                  R                  S[        R                  " S5      S9n g )N
   Or   )r   r   r   r7   r   r3   s    r!   unsupported_typeBTestSharedMemory.test_invalid_array_type.<locals>.unsupported_type  s#    ++##Bbhhsm#<Cr$   z*.*Invalid NumPy dtype specified: 'int33'.*c                  @    [         R                  R                  SSS9n g )Nr   int33r   )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HHTV-. 6 ;	7##K5HHTV01 65 65 65s   %B$%B#
B #
B1z+Struct model array unsupported in simulatorc           	        ^ Sm[         R                  " [        [        S S S2   [        S S S2   5      5      U4S j5       n[        R
                  " T4SS9n[        R
                  " T4SS9nUST4   " X#5        [        U5       H  u  pEU R                  UTU-
  S-
  5        M     [        U5       H   u  pFU R                  UTU-
  S-
  S-  5        M"     g )N@   r   c                 z  > [         R                  R                  T[        S9n[         R                  " S5      nTU-
  S-
  nU[        U 5      :  ak  U[        U5      :  a[  [        [        U5      [        US-  5      5      nXRU'   [         R                  " 5         X$   R                  X'   X$   R                  X'   g g g )Nr   r   r   )r   r   r   r   gridr|   r   r   r}   r\   r~   )outxoutyr3   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   r   )r   r*   r   r   r7   r8   	enumerater9   )r+   r  arrxarryr   r\   r~   r   s          @r!   test_struct_model_type_static.TestSharedMemory.test_struct_model_type_static  s    	$uSqSz51:.	/	$ 
0	$" xx73xx73&q({3D?dODAQ1q 01 $dODAQA!1Q 67 $r$   rm   N)rn   ro   rp   rq   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r  rr   rm   r$   r!   rt   rt   ^   s    "8H5?">"D(A,!>FI,H>P>,8\ 672 82 BC8 D8r$   rt   __main__)numbar   r   r   r   numba.core.errorsr   
numba.corer   numba.cuda.testingr	   r
   r   numpyr7   numba.npr   r   extensions_usecasesr   r   r   r   r   r   rt   rn   mainrm   r$   r!   <module>r     s    , , )  F F  ) CHHsBHHo"BJJ79 : LL L^Z8| Z8z
 zMMO r$   