
    shJ                        S SK r S SKrS SKrS SKrS SKrS SKrS SK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  S SKJr  S SKJrJr  \" S5       " S S	\\5      5       r\" S5       " S
 S\\5      5       rS r\" S5       " S S\\5      5       rS r\" S5       " S S\\5      5       r\" S5       " S S\5      5       rg)    N)cuda)NumbaWarning)CUDATestCaseskip_on_cudasimskip_unless_cc_60skip_if_cudadevrt_missingskip_if_mvc_enabledtest_data_dir)SerialMixin)DispatcherCacheUsecasesTestskip_bad_accessz$Simulator does not implement cachingc                      \ rS rSr\R
                  R                  \5      r\R
                  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5      S 5       5       5       r\\\" S5      S 5       5       5       rS r\\R<                  " \R>                  S:H  S5      S 5       5       r \\R<                  " \R>                  S:H  S5      S 5       5       r!S r"Sr#g)CUDACachingTest   cache_usecases.pycuda_caching_test_fodderc                 \    [         R                  " U 5        [        R                  " U 5        g Nr   setUpr   selfs    x/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/tests/cudapy/test_caching.pyr   CUDACachingTest.setUp       #))$/4     c                 \    [         R                  " U 5        [        R                  " U 5        g r   r   tearDownr   r   s    r   r   CUDACachingTest.tearDown       d##,,T2r   c                    U R                  S5        U R                  5       nU R                  S5        UR                  nU R                  U" SS5      S5        U R                  S5        U R                  U" SS5      S5        U R                  S5        U R	                  UR
                  SS5        UR                  nU" UR                  S5      nU R                  [        U5      S5        UR                  nU" UR                  S5      nU R                  [        U5      S5        U R                  S5        U R	                  UR
                  SS5        U R                  5         g )	Nr                  @      @   r#   g     E@)check_pycacheimport_moduleadd_usecaseassertPreciseEqual
check_hitsfuncrecord_return_alignedaligned_arrtuplerecord_return_packed
packed_arrrun_in_separate_process)r   modfrecs       r   test_cachingCUDACachingTest.test_caching!   s   1  "1OO!Q+1#q	3/11%%%#c
I6$$"c
I611% 	$$&r   c                     U R                  5       nUR                  nU R                  U" SS5      S5        U R                  S5        g )Nr#   r$   r%   r   )r+   add_nocache_usecaser-   r*   r   r6   r7   s      r   test_no_cachingCUDACachingTest.test_no_caching:   s?      "##!Q+1r   c                     U R                  S5        U R                  5       nUR                  nUS   " 5         U R                  S5        g )Nr   )r(   r(   r#   )r*   r+   many_localsr=   s      r   test_many_locals CUDACachingTest.test_many_localsA   s@     	1  "OO	$	1r   c                    U R                  5       n[        R                  " 5          [        R                  " S[        5        UR
                  nU R                  U" S5      S5        UR                  nU R                  U" S5      S5        UR                  nU R                  U" S5      S5        UR                  nU R                  U" S5      S5        U R                  S5        S S S 5        g ! , (       d  f       g = f)Nerrorr$   r%      
         )r+   warningscatch_warningssimplefilterr   closure1r-   closure2closure3closure4r*   r=   s      r   test_closureCUDACachingTest.test_closureM   s      "$$&!!'<8A##AaD!,A##AaD!,A##AaD"-A##AaD"-q! '&&s   B=C,,
C:c                 F   U R                  5       nUR                  SS5        UR                  SS5        UR                  SS5        UR                  SS5        UR	                  UR
                  S5        UR                  UR                  S5        UR                  S5        U R                  5       nU R                  UR                  R                  SS5        U R                  5       nU R                  X5        UR                  nU" SS5        U R                  UR                  SS5        U" SS5        U R                  UR                  SS5        U R                  U R                  5       U5        U R                  5         U R                  U R                  5       U5        g )Nr#   r$   r&   g      @r   r(   )r+   r,   outer_uncachedouterr3   r4   r0   r1   simple_usecase_callerget_cache_mtimesr.   r/   assertIsNotassertEqualr5   )r   r6   mtimesmod2r7   s        r   test_cache_reuse CUDACachingTest.test_cache_reuse]   sH     "1S!1a 		!Q  3!!#//15!!!$&&(,,a3!!##	!Q1%	#s1% 	..0&9$$&..0&9r   c                 Z   U R                  5       nUR                  nU R                  U" SS5      S5        [        U R                  S5       nUR                  S5        S S S 5        U R                  5       nUR                  nU R                  U" SS5      S5        g ! , (       d  f       ND= f)Nr#   r$   r%   az
Z = 10
   )r+   r,   r-   openmodfilewriter=   s      r   test_cache_invalidate%CUDACachingTest.test_cache_invalidatex   s      "OO!Q+ $,,$GGL! %   "OO!Q, %$s   B
B*c                    U R                  5       nUR                  nU R                  U" SS5      S5        U R                  5       nUR                  nSUl        U R                  U" SS5      S5        UR                  R                  5         U R                  U" SS5      S5        U R                  5       nUR                  nU R                  U" SS5      S5        g )Nr#   r$   r%   rG   r`   )r+   r,   r-   Zr/   	recompiler=   s      r   test_recompileCUDACachingTest.test_recompile   s      "OO!Q+  "OO!Q+	!Q,   "OO!Q,r   c                     U R                  5       nUR                  nU R                  U" S5      S5        UR                  nU R                  U" S5      S5        g )Nr#      rF   )r+   renamed_function1r-   renamed_function2r=   s      r   test_same_namesCUDACachingTest.test_same_names   sN      "!!!a(!!!a(r   zCG not supported with MVCc                     U R                  S5        U R                  5       nU R                  S5        UR                  S5        U R                  S5        U R                  5         g )Nr   r#   )r*   r+   
cg_usecaser5   )r   r6   s     r   test_cache_cgCUDACachingTest.test_cache_cg   sV     	1  "1q1 	$$&r   c           	         U R                  S5        S[        U R                  U R                  S9-  n[        R
                  " [        R                  SU/[        R                  [        R                  S9nUR                  SS9u  p4UR                  S:w  a>  [        SUR                  < S	UR                  5       < S
UR                  5       < S35      eg )Nr   zif 1:
            import sys

            sys.path.insert(0, %(tempdir)r)
            mod = __import__(%(modname)r)
            mod.cg_usecase(0)
            )tempdirmodnamez-c)stdoutstderr<   )timeoutzprocess failed with code z: 
stdout follows
z
stderr follows

)r*   dictrv   rw   
subprocessPopensys
executablePIPEcommunicate
returncodeAssertionErrordecode)r   codepopenouterrs        r   test_cache_cg_clean_run'CUDACachingTest.test_cache_cg_clean_run   s     	1 t||T\\BC   #..$!=(2(29 $$R$0q   ##SZZ\3::<A  !r   c                    U R                  5       nUR                  nU R                  [        R                  UR
                  R                  R                  SS9  U R                  U" SS5      S5        U R                  UR
                  SS5        U R                  5       nUR                  nU R                  U" SS5      S5        U R                  UR
                  SS5        U R                  S5        g)	za
With a disabled __pycache__, test there is a working fallback
(e.g. on the user-wide cache dir)
T)ignore_errorsr#   r$   r%   r   r(   N)r+   r,   
addCleanupshutilrmtreer/   stats
cache_pathr-   r.   r*   )r   r6   r7   r[   s       r   _test_pycache_fallback&CUDACachingTest._test_pycache_fallback   s    
   "OO 	qvv||'>'>&* 	 	, 	!Q+1% !!#!Q+1% 	1r   ntz3cannot easily make a directory read-only on Windowsc                    [         R                  " U R                  5      R                  n[         R                  " U R                  S5        U R                  [         R                  U R                  U5        U R                  5         g )N@  )osstatrv   st_modechmodr   r   )r   	old_permss     r   test_non_creatable_pycache*CUDACachingTest.test_non_creatable_pycache   sR    
 GGDLL)11	
u%$,,	:##%r   c                 R   [         R                  R                  U R                  S5      n[         R                  " U5        [         R
                  " U5      R                  n[         R                  " US5        U R                  [         R                  X5        U R                  5         g )N__pycache__r   )
r   pathjoinrv   mkdirr   r   r   r   r   )r   pycacher   s      r   test_non_writable_pycache)CUDACachingTest.test_non_writable_pycache   sh    
 '',,t||];
GGG$,,	
% '5##%r   c                     [        [        S-  5      nSnU R                  [        U5         [        R
                  " SSU/S9S 5       nS S S 5        g ! , (       d  f       g = f)Nzjitlink.ptxz0Cannot pickle CUDACodeLibrary with linking fileszvoid()T)cachelinkc                      g r    r   r   r   r7   >CUDACachingTest.test_cannot_cache_linking_libraries.<locals>.f  s    r   )strr
   assertRaisesRegexRuntimeErrorr   jit)r   r   msgr7   s       r   #test_cannot_cache_linking_libraries3CUDACachingTest.test_cannot_cache_linking_libraries   sS    ==01@##L#6XXhd$8 9 766s   A
A r   N)$__name__
__module____qualname____firstlineno__r   r   dirname__file__herer   usecases_filerw   r   r   r9   r>   rB   rQ   r\   rd   ri   ro   r   r   r	   rs   r   r   r   unittestskipIfnamer   r   r   __static_attributes__r   r   r   r   r      s&   77??8$DGGLL':;M(G!3'2
" :6--$) 45' 6  ' 45 6  84 __RWW_JL&L & __RWW_JL&L &r   r   c                       \ rS rSr\R
                  R                  \5      r\R
                  R                  \S5      r
SrS rS rS rS rSrg	)
CUDAAndCPUCachingTesti	  zcache_with_cpu_usecases.py cuda_and_cpu_caching_test_fodderc                 \    [         R                  " U 5        [        R                  " U 5        g r   r   r   s    r   r   CUDAAndCPUCachingTest.setUp  r   r   c                 \    [         R                  " U 5        [        R                  " U 5        g r   r   r   s    r   r   CUDAAndCPUCachingTest.tearDown  r!   r   c                    U R                  S5        U R                  5       nU R                  S5        UR                  nUR                  nU R	                  U" S5      S5        U R                  S5        U R	                  U" S5      S5        U R                  S5        U R                  UR                  SS5        U R                  UR                  SS5        U R	                  U" S5      S5        U R                  S5        U R	                  U" S5      S5        U R                  S5        U R                  UR                  SS5        U R                  UR                  SS5        g )Nr   rI   r#   r$   r(         @rl   )r*   r+   
assign_cpuassign_cudar-   r.   r/   )r   r6   f_cpuf_cudas       r   test_cpu_and_cuda_targets/CUDAAndCPUCachingTest.test_cpu_and_cuda_targets  s    	1  "1a!,1q	1-1

Aq)Q*c
C01sS11

Aq)Q*r   c                    U R                  5       nUR                  S5        UR                  S5        UR                  S5        UR                  S5        U R                  5       nU R	                  UR                  R
                  SS5        U R	                  UR                  R
                  SS5        U R                  5       nU R                  X5        UR                  nUR                  nU" S5        U R	                  UR
                  SS5        U" S5        U R	                  UR
                  SS5        U" S5        U R	                  UR
                  SS5        U" S5        U R	                  UR
                  SS5        U R                  U R                  5       U5        U R                  5         U R                  U R                  5       U5        g )NrI   r   r   r#   r(   r&   )	r+   r   r   rW   r.   r/   rX   rY   r5   )r   r6   rZ   r[   r   r   s         r   test_cpu_and_cuda_reuse-CUDAAndCPUCachingTest.test_cpu_and_cuda_reuse0  s`     "qs&&( 	++Q2,,a3!!##!!a

Aq)c


Aq)q	Q*sQ* 	..0&9$$&..0&9r   r   N)r   r   r   r   r   r   r   r   r   r   r   rw   r   r   r   r   r   r   r   r   r   r   	  sB    77??8$DGGLL'CDM0G!3+2 :r   r   c                     [         R                  S   n U    [         R                  " 5       R                  R                  nS S S 5        [         R                  SS   HL  nU   [         R                  " 5       R                  R                  nUW:w  a  X4sS S S 5        s  $  S S S 5        MN     g ! , (       d  f       Nr= f! , (       d  f       Mq  = f)Nr   r(   )r   gpuscurrent_contextdevicecompute_capability)	first_gpufirst_ccgpuccs       r   get_different_cc_gpusr   S  s     		!I	'')00CC 
 yy}%%'..AABX~!' S S   
 Ss   *B, 2B=,
B:=
C	c                       \ rS rSr\R
                  R                  \5      r\R
                  R                  \S5      r
SrS rS rS rSrg)	TestMultiCCCachingid  r   !cuda_multi_cc_caching_test_fodderc                 \    [         R                  " U 5        [        R                  " U 5        g r   r   r   s    r   r   TestMultiCCCaching.setUpj  r   r   c                 \    [         R                  " U 5        [        R                  " U 5        g r   r   r   s    r   r   TestMultiCCCaching.tearDownn  r!   r   c                    [        5       nU(       d  U R                  S5        U R                  S5        U R                  5       nU R                  S5        US      UR                  nU R                  U" SS5      S5        U R                  S5        U R                  U" SS5      S5        U R                  S5        U R                  UR                  SS5        UR                  nU" UR                  S5      nU R                  [        U5      S	5        UR                  nU" UR                  S5      nU R                  [        U5      S	5        U R                  S5        U R                  UR                  SS5        S S S 5        US      UR                  nU R                  U" SS5      S5        U R                  S5        U R                  U" SS5      S5        U R                  S5        U R                  UR                  SS5        UR                  nU" UR                  S5      nU R                  [        U5      S	5        UR                  nU" UR                  S5      nU R                  [        U5      S	5        U R                  S5        U R                  UR                  SS5        S S S 5        U R                  5       nU R                  X%5        US      UR                  nU R                  U" SS5      S5        U R                  S
5        U R                  U" SS5      S5        U R                  S5        U R                  UR                  SS5        UR                  nU" UR                  S5      nU R                  [        U5      S	5        UR                  nU" UR                  S5      nU R                  [        U5      S	5        U R                  S5        U R                  UR                  SS5        S S S 5        U R                  5       nU R                  X&5        US      UR                  nU R                  U" SS5      S5        U R                  U" SS5      S5        UR                  nU" UR                  S5      nU R                  [        U5      S	5        UR                  nU" UR                  S5      nU R                  [        U5      S	5        S S S 5        US      UR                  nU R                  U" SS5      S5        U R                  U" SS5      S5        UR                  nU" UR                  S5      nU R                  [        U5      S	5        UR                  nU" UR                  S5      nU R                  [        U5      S	5        S S S 5        g ! , (       d  f       GN,= f! , (       d  f       GN= f! , (       d  f       GN= f! , (       d  f       GN= f! , (       d  f       g = f)Nz.Need two different CCs for multi-CC cache testr   r#   r$   r%   r&   r'   r(   r)      rF   rG   )r   skipTestr*   r+   r,   r-   r.   r/   r0   r1   r2   r3   r4   rX   )r   r   r6   r7   r8   r[   mod3s          r   
test_cacheTestMultiCCCaching.test_cacher  s   $&MMJK1  "1 !WA##AaGQ/q!##Ac1Is3q!OOAFFAq)))ACOOQ'C##E#J	:((ACNNA&C##E#J	:q!OOAFFAq)! ( !WA##AaGQ/q!##Ac1Is3q!OOAFFAq)))ACOOQ'C##E#J	:((ACNNA&C##E#J	:q!OOAFFAq)! ( !!##!W  A##AaGQ/q!##Ac1Is3q!OOAFFAq)**ACOOQ'C##E#J	:))ACNNA&C##E#J	:r"OOAFFAq)! 0 !!## !W  A##AaGQ/##Ac1Is3**ACOOQ'C##E#J	:))ACNNA&C##E#J	:  !W  A##AaGQ/##Ac1Is3**ACOOQ'C##E#J	:))ACNNA&C##E#J	: Wo W( W. W: W WsA   D VD V1D W#B3W#B3W'
V.1
W 
W
W$'
W5r   N)r   r   r   r   r   r   r   r   r   r   r   rw   r   r   r   r   r   r   r   r   r   d  s>    77??8$DGGLL':;M1G!3l;r   r   c                  ,    SSK Jn   SU l        SU l        g )Nr   config)
numba.corer   CUDA_LOW_OCCUPANCY_WARNINGSCUDA_WARN_ON_IMPLICIT_COPYr   s    r   child_initializerr     s     ")*F&()F%r   c                       \ rS rSrSr\R                  R                  \5      r	\R                  R                  \	S5      rSrS rS rS rSrg	)
TestMultiprocessCachei  Fr   cuda_mp_caching_test_fodderc                 \    [         R                  " U 5        [        R                  " U 5        g r   r   r   s    r   r   TestMultiprocessCache.setUp  r   r   c                 \    [         R                  " U 5        [        R                  " U 5        g r   r   r   s    r   r   TestMultiprocessCache.tearDown  r!   r   c                    U R                  5       nUR                  nSn [        R                  " S5      nUR                  U[        5      n [        UR                  U[        U5      5      5      nUR                  5         U R                  XcUS-
  -  S-  5        g ! [         a	    [        n Nxf = f! UR                  5         f = f)Nr$   spawnr(   r#   )r+   rV   multiprocessingget_contextAttributeErrorPoolr   sumimaprangecloserY   )r   r6   r7   nctxpoolress          r   test_multiprocessing*TestMultiprocessCache.test_multiprocessing  s      " %%	"!--g6C xx,-	dii58,-CJJL1q5kQ./  	"!C	" JJLs   B $B2 B/.B/2Cr   N)r   r   r   r   _numba_parallel_test_r   r   r   r   r   r   r   rw   r   r   r  r   r   r   r   r   r     sE    
 "77??8$DGGLL':;M+G!30r   r   z0Simulator does not implement the CUDACodeLibraryc                       \ rS rSrS rSrg)TestCUDACodeLibraryi  c                     SSK Jn  [        5       nSnU" X#5      nU R                  [        S5         UR                  5         S S S 5        g ! , (       d  f       g = f)Nr   )CUDACodeLibrarylibraryzCannot pickle unfinalized)numba.cuda.codegenr  objectr   r   _reduce_states)r   r  codegenr   cls        r   !test_cannot_serialize_unfinalized5TestCUDACodeLibrary.test_cannot_serialize_unfinalized  sI     	7 (W+##L2MN ONNs   A
Ar   N)r   r   r   r   r  r   r   r   r   r
  r
    s    
 r   r
  )r   r   r   r~   r   r   rJ   numbar   numba.core.errorsr   numba.cuda.testingr   r   r   r   r	   r
   numba.tests.supportr   numba.tests.test_cachingr   r   r   r   r   r   r   r   r
  r   r   r   <module>r     s     	   
    *D D ,7 78rk#> r 9rj 78F:K)D F: 9F:R" 78y;&A y; 9y;x* 78$0K)D $0 9$0N CD ,   E r   