
    sh'                     V   S SK rS SKrS SKJr  S SKJrJr  S SKJrJr  S SK	J
r
JrJr  S SKJr  S SKJr  S SKJr  S S	KJrJrJrJrJrJrJr  \R6                  " S
\R,                  S9rS rS rS rS r S r!Sr"S r#\" S5       " S S\5      5       r$\%S:X  a  \RL                  " 5         gg)    N)unittest)skip_on_cudasimskip_if_cuda_includes_missing)CUDATestCasetest_data_dir)CudaAPIErrorLinkerLinkerError)
NvrtcError)require_context)ignore_internal_warnings)cudavoidfloat64int64int32typeoffloat32
   dtypec                     [         R                  R                  [        5      n[         R                  " S5      nX   S-   X'   g )N         ?)r   const
array_likeCONST1Dgrid)ACis      x/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/tests/cudadrv/test_linker.pysimple_const_memr#      s0    

g&A		!A4#:AD    c                    SnSnSn	Sn
SnSnSnSnSnSnSnSnSnSnSnSnSnSnSnSn[        U5       H^  nXr-  nX-  nX-  n	X-  n
X-  nX-  nX-  nX-  nX-  nUU-  nUU-  nUU-  nUU-  nUU-  nUU-  nUU-  nUU-  nUU-  nUU-  nUU-  nM`     Xx-   U	-   U
-   U-   U [        R                  " S5      '   U [        R                  " S5      ==   X-   U-   U-   U-   -  ss'   U [        R                  " S5      ==   UU-   U-   U-   U-   -  ss'   U [        R                  " S5      ==   UU-   U-   U-   U-   -  ss'   g )Nr   r   r   )ranger   r   )xabcdefa1a2a3a4a5b1b2b3b4b5c1c2c3c4c5d1d2d3d4d5r!   s                               r"   func_with_lots_of_registersrB      s   	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B1X









a
a
a
a
a
a
q
q
q
q
q) * glR'",AdiilOdiilOrw|b(2--OdiilOrBw|b(2--OdiilOrBw|b(2--Or$   c                     [         R                  R                  SU5      n[         R                  " S5      nUS:X  a  [	        S5       H  nXBU'   M	     [         R
                  " 5         X#   X'   g )Nd   r   r   )r   sharedarrayr   r&   syncthreads)arydtysmr!   js        r"   simple_smemrL   H   sV    			3	$B		!AAvsAqE UCFr$   c                     [         R                  " S5      u  p[         R                  R                  S[        5      nUS-   US-   -  X1U4'   [         R
                  " 5         X1U4   XU4'   g )N   )r      r   )r   r   rE   rF   r   rG   )rH   r!   rK   rJ   s       r"   coop_smem2drP   R   s\    99Q<DA			8W	-BA!a% B!tHa4C1Ir$   c                 8    [         R                  " S5      nXU'   g Nr   )r   r   )rH   r!   s     r"   simple_maxthreadsrS   Z   s    		!AFr$   i  c                     [         R                  R                  [        U5      n[	        UR
                  S   5       H	  nX   X4'   M     [	        UR
                  S   5       H	  nX4   X'   M     g Nr   )r   localrF   	LMEM_SIZEr&   shape)r   BrI   r    r!   s        r"   simple_lmemrZ   b   sX    

C(A1771:t 1771:t r$   z$Linking unsupported in the simulatorc                       \ rS rSrSS0r\S 5       rS rS rS r	S r
S	 rS
 rS rS r\S 5       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g)
TestLinkerj   NUMBA_CUDA_USE_NVIDIA_BINDING0c                 .    [         R                  " SS9nAg)z9Simply go through the constructor and destructor
        )      )ccN)r	   new)selflinkers     r"   test_linker_basicTestLinker.test_linker_basicn   s     v&r$   c                 ~   [         R                  " SS5      q[        [        S-  5      nU(       a  S/nO/ n[         R
                  " USU/06S 5       n[        R                  " S/[        R                  S9n[        R                  " S	/[        R                  S9nUS
   " XV5        U R                  US   S:H  5        g )Nbarint32(int32)zjitlink.ptxzvoid(int32[:], int32[:])linkc                 ^    [         R                  " S5      nX==   [        X   5      -  ss'   g rR   )r   r   rj   )r'   yr!   s      r"   foo%TestLinker._test_linking.<locals>.foo   s!    		!ADCIDr$   {   r   iA  )r   r   r   i  )
r   declare_devicerj   strr   jitnprF   r   
assertTrue)re   eagerrl   argsro   r   rY   s          r"   _test_linkingTestLinker._test_linkingu   s    !!%8==01./DD	4	%tf	%	 
&	 HHcU"((+HHcU"((+D	!!-.r$   c                 "    U R                  SS9  g )NFrw   ry   re   s    r"   test_linking_lazy_compile$TestLinker.test_linking_lazy_compile   s    'r$   c                 "    U R                  SS9  g )NTr|   r}   r~   s    r"   test_linking_eager_compile%TestLinker.test_linking_eager_compile   s    &r$   c                 d  ^ [         R                  " SS5      m[        [        S-  5      n[         R                  " U/S9U4S j5       n[
        R                  " S[
        R                  S9n[
        R                  " U5      nUS   " XC5        US	-  n[
        R                  R                  XE5        g )
Nrj   rk   z
jitlink.curl   c                 j   > [         R                  " S5      nU[        U 5      :  a  T" X   5      X'   g g rR   )r   r   len)rr'   r!   rj   s      r"   kernel*TestLinker.test_linking_cu.<locals>.kernel   s-    		!A3q6z14y r$   r   r   )r       rN   )r   rr   rs   r   rt   ru   aranger   
zeros_liketestingassert_array_equal)re   rl   r   r'   r   expectedrj   s         @r"   test_linking_cuTestLinker.test_linking_cu   s    !!%8=</0	v		! 
	! IIb)MM!ua q5


%%a2r$   c                   ^ [         R                  " SS5      m[        [        S-  5      n[        R
                  " SS9 n[        5         [         R                  " SU/S9U4S j5       nS S S 5        U R                  [        W5      S	S
5        U R                  S[        US   R                  5      5        U R                  S[        US   R                  5      5        g ! , (       d  f       N{= f)Nrj   rk   zwarn.cuT)recordvoid(int32)r   c                    > T" U 5        g N r'   rj   s    r"   r   6TestLinker.test_linking_cu_log_warning.<locals>.kernel   
    Ar$   r   zExpected warnings from NVRTCzNVRTC log messagesr   zdeclared but never referenced)r   rr   rs   r   warningscatch_warningsr   rt   assertEqualr   assertInmessage)re   rl   wr   rj   s       @r"   test_linking_cu_log_warning&TestLinker.test_linking_cu_log_warning   s    !!%8=9,-$$D1Q$&XXm4&1 2	 2 	Q$BC*C!,=>5s1Q4<<7HI 21s   *C
C,c                   ^ [         R                  " SS5      m[        [        S-  5      nU R	                  [
        5       n[         R                  " SU/S9U4S j5       nS S S 5        WR                  R                  S   nU R                  SU5        U R                  S	U5        U R                  S
U5        g ! , (       d  f       N^= f)Nrj   rk   zerror.cur   r   c                    > T" U 5        g r   r   r   s    r"   r   0TestLinker.test_linking_cu_error.<locals>.kernel   r   r$   r   zNVRTC Compilation failurez identifier "SYNTAX" is undefinedz in the compilation of "error.cu")
r   rr   rs   r   assertRaisesr   rt   	exceptionrx   r   )re   rl   r,   r   msgrj   s        @r"   test_linking_cu_error TestLinker.test_linking_cu_error   s    !!%8=:-.z*aXXm4&1 2 +
 kkq!1378#>8#> +*s     B88
Cc                     SnU R                  [        U5         [        R                  " SS/S9S 5       nS S S 5        g ! , (       d  f       g = f)Nz/Don't know how to link file with extension .cuhvoid()z
header.cuhr   c                      g r   r   r   r$   r"   r   >TestLinker.test_linking_unknown_filetype_error.<locals>.kernel       r$   assertRaisesRegexRuntimeErrorr   rt   re   expected_errr   s      r"   #test_linking_unknown_filetype_error.TestLinker.test_linking_unknown_filetype_error   sC    H##L,?XXhl^4 5 @??	   ?
Ac                     SnU R                  [        U5         [        R                  " SS/S9S 5       nS S S 5        g ! , (       d  f       g = f)Nz-Don't know how to link file with no extensionr   datar   c                      g r   r   r   r$   r"   r   DTestLinker.test_linking_file_with_no_extension_error.<locals>.kernel   r   r$   r   r   s      r"   )test_linking_file_with_no_extension_error4TestLinker.test_linking_file_with_no_extension_error   sC    F##L,?XXhfX. / @??r   c                 `    [        [        S-  5      n[        R                  " SU/S9S 5       ng )Nzcuda_include.cur   r   c                      g r   r   r   r$   r"   r   7TestLinker.test_linking_cu_cuda_include.<locals>.kernel   s    r$   )rs   r   r   rt   )re   rl   r   s      r"   test_linking_cu_cuda_include'TestLinker.test_linking_cu_cuda_include   s3    =#445 
($	(	 
)	r$   c                     U R                  [        5       n[        R                  " SS/S9S 5       nS S S 5        U R	                  SWR
                  R                  5        g ! , (       d  f       N5= f)Nvoid(int32[::1])znonexistent.ar   c                     SU S'   g rU   r   )r'   s    r"   r-   2TestLinker.test_try_to_link_nonexistent.<locals>.f   s    !r$   znonexistent.a not found)r   r
   r   rt   r   r   rx   )re   r,   r-   s      r"   test_try_to_link_nonexistent'TestLinker.test_try_to_link_nonexistent   s]    {+qXX(/@A B , 	/1A1AB	 ,+s   A""
A0c                     [         R                  " [        5      nUR                  " [        R
                  " S5      /[        S5      Q76 nU R                  UR                  5       S5        g)zEnsure that the jitted kernel used in the test_set_registers_* tests
uses more than 57 registers - this ensures that test_set_registers_*
are really checking that they reduced the number of registers used from
something greater than the maximum.r      9   N)	r   rt   rB   
specializeru   emptyr&   assertGreaterget_regs_per_threadre   compileds     r"   test_set_registers_no_max$TestLinker.test_set_registers_no_max   sM    
 8878&&rxx|?eAh?87792>r$   c                     [         R                  " SS9" [        5      nUR                  " [        R
                  " S5      /[        S5      Q76 nU R                  UR                  5       S5        g )Nr   max_registersr   r   	r   rt   rB   r   ru   r   r&   assertLessEqualr   r   s     r"   test_set_registers_57 TestLinker.test_set_registers_57   P    88"-.IJ&&rxx|?eAh?X99;R@r$   c                     [         R                  " SS9" [        5      nUR                  " [        R
                  " S5      /[        S5      Q76 nU R                  UR                  5       S5        g )N&   r   r   r   r   r   s     r"   test_set_registers_38 TestLinker.test_set_registers_38   r   r$   c           	          [        [        S S S2   [        [        [        [        [        [        5      n[        R                  " USS9" [
        5      nU R                  UR                  5       S5        g )Nr   r   r   )r   r   r   r   rt   rB   r   r   )re   sigr   s      r"   test_set_registers_eager#TestLinker.test_set_registers_eager   sO    73Q3<ueUEJ88Cr23NOX99;R@r$   c                     [        [        S S S2   5      n[        R                  " U5      " [        5      nUR                  5       nU R                  U[        R                  5        g rR   )	r   r   r   rt   r#   get_const_mem_sizeassertGreaterEqualr   nbytes)re   r   r   const_mem_sizes       r"   test_get_const_mem_size"TestLinker.test_get_const_mem_size  sI    73Q3< 88C=!12!446?r$   c                     [         R                  " [        5      nUR                  " [        R
                  " S5      /[        S5      Q76 nUR                  5       nU R                  US5        g )Nr   r   r   )	r   rt   rB   r   ru   r   r&   get_shared_mem_per_blockr   )re   r   shared_mem_sizes      r"   test_get_no_shared_memory$TestLinker.test_get_no_shared_memory  sP    8878&&rxx|?eAh?";;=!,r$   c                     [        [        S S S2   [        [        R                  5      5      n[        R
                  " U5      " [        5      nUR                  5       nU R                  US5        g )Nr   i  )	r   r   r   ru   r   rt   rL   r   r   )re   r   r   r   s       r"   test_get_shared_mem_per_block(TestLinker.test_get_shared_mem_per_block  sO    51:vbhh/088C=-";;=#.r$   c                     [         R                  " [        5      nUR                  [        R
                  " S[        R                  S9[        R                  5      nUR                  5       nU R                  US5        g )NrD   r   i   )
r   rt   rL   r   ru   zerosr   r   r   r   )re   r   compiled_specializedr   s       r"   #test_get_shared_mem_per_specialized.TestLinker.test_get_shared_mem_per_specialized  sW    88K('22HHS)2:: 7.GGI#.r$   c                     [         R                  " S5      " [        5      nUR                  5       nU R	                  US5        g )Nzvoid(float32[:,::1])r   )r   rt   rP   get_max_threads_per_blockr   )re   r   max_threadss      r"   test_get_max_threads_per_block)TestLinker.test_get_max_threads_per_block  s4    8823K@88:;*r$   c                 6   [         R                  " S5      " [        5      nUR                  5       nUS-   n[        R
                  " U[        R                  S9n USU4   " U5        g ! [         a&  nU R                  SUR                  5         S nAg S nAff = f)Nr   r   r   cuLaunchKernel)
r   rt   rS   r   ru   r   r   r   r   r   )re   r   r   nelemrH   r,   s         r"   test_max_threads_exceeded$TestLinker.test_max_threads_exceeded   s~    88./0AB88:ahhuBHH-	3QXs# 	3MM*AEE22	3s   A( (
B2BBc                 `   [        [        S S S2   [        S S S2   [        [        R                  5      5      n[        R
                  " U5      " [        5      nUR                  5       n[        R                  " [        R                  5      R                  [        -  nU R                  X45        g rR   )r   r   r   ru   r   rt   rZ   get_local_mem_per_threadr   itemsizerW   r   )re   r   r   local_mem_size	calc_sizes        r"   test_get_local_mem_per_thread(TestLinker.test_get_local_mem_per_thread*  st    51:uSqSz6"((+;<88C=-!::<HHRXX&//);	:r$   c                    [         R                  " [        5      nUR                  [        R
                  " [        [        R                  S9[        R
                  " [        [        R                  S9[        R                  5      nUR                  5       n[        R                  " [        R                  5      R                  [        -  nU R                  X45        g )Nr   )r   rt   rZ   r   ru   r   rW   r   r   r  r   r  r   )re   r   r   r  r  s        r"   "test_get_local_mem_per_specialized-TestLinker.test_get_local_mem_per_specialized1  s    88K('22HHYbhh/HHYbhh/JJ  .FFHHHRZZ(11I=	:r$   r   N)__name__
__module____qualname____firstlineno___NUMBA_NVIDIA_BINDING_0_ENVr   rg   ry   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r  r  r  __static_attributes__r   r$   r"   r\   r\   j   s    #BC"H /.('3*J$?$ # #C?A
A
A
@-//+
3;;r$   r\   __main__)'numpyru   r   numba.cuda.testingr   r   r   r   r   numba.cuda.cudadrv.driverr   r	   r
   numba.cuda.cudadrv.errorr   
numba.cudar   numba.tests.supportr   numbar   r   r   r   r   r   r   r   r   r#   rB   rL   rP   rS   rW   rZ   r\   r  mainr   r$   r"   <module>r     s      ' O :4 4 / & 8 D D D ))Bbjj
)-.`
 	 78N; N; 9N;b zMMO r$   