
    xKgJ                        d dl Z d dlZd dlZd dlZd dlZd dlZd dlZd dlmZ d dl	m
Z
 d dlmZmZmZmZmZmZ d dlmZ d dlmZmZ  ed       G d d	ee             Z ed       G d
 dee             Zd Z ed       G d dee             Zd Z ed       G d dee             Z ed       G d de             Zy)    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                      e Zd Zej                  j                  e      Zej                  j                  ed      Z	dZ
d Zd Zd Zd Zd Zd Zd	 Zd
 Zd Zd Zee ed      d                      Zee ed      d                      Zd Ze ej:                  ej<                  dk(  d      d               Ze ej:                  ej<                  dk(  d      d               Z d Z!y)CUDACachingTestcache_usecases.pycuda_caching_test_fodderc                 X    t        j                  |        t        j                  |        y Nr   setUpr   selfs    h/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_caching.pyr   zCUDACachingTest.setUp       #))$/4     c                 X    t        j                  |        t        j                  |        y r   r   tearDownr   r   s    r   r   zCUDACachingTest.tearDown       d##,,T2r   c                    | j                  d       | j                         }| j                  d       |j                  }| j                   |dd      d       | j                  d       | j                   |dd      d       | j                  d       | j	                  |j
                  dd       |j                  } ||j                  d      }| j                  t        |      d       |j                  } ||j                  d      }| j                  t        |      d       | j                  d       | j	                  |j
                  dd       | j                          y )	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_cachingzCUDACachingTest.test_caching!   s   1  "1OO!Q+1#q	3/11%%%#c
I6$$"c
I611% 	$$&r   c                     | j                         }|j                  }| j                   |dd      d       | j                  d       y )Nr    r!   r"   r   )r(   add_nocache_usecaser*   r'   r   r3   r4   s      r   test_no_cachingzCUDACachingTest.test_no_caching:   s?      "##!Q+1r   c                     | j                  d       | j                         }|j                  } |d           | j                  d       y )Nr   )r%   r%   r    )r'   r(   many_localsr9   s      r   test_many_localsz CUDACachingTest.test_many_localsA   sB     	1  "OO$	1r   c                    | j                         }t        j                         5  t        j                  dt               |j
                  }| j                   |d      d       |j                  }| j                   |d      d       |j                  }| j                   |d      d       |j                  }| j                   |d      d       | j                  d       d d d        y # 1 sw Y   y xY w)Nerrorr!   r"      
         )r(   warningscatch_warningssimplefilterr   closure1r*   closure2closure3closure4r'   r9   s      r   test_closurezCUDACachingTest.test_closureM   s      "$$&!!'<8A##AaD!,A##AaD!,A##AaD"-A##AaD"-q! '&&s   B<C**C3c                 H   | j                         }|j                  dd       |j                  dd       |j                  dd       |j                  dd       |j	                  |j
                  d       |j                  |j                  d       |j                  d       | j                         }| j                  |j                  j                  dd       | j                         }| j                  ||       |j                  } |dd       | j                  |j                  dd        |dd       | j                  |j                  dd       | j                  | j                         |       | j                          | j                  | j                         |       y )Nr    r!   r#   g      @r   r%   )r(   r)   outer_uncachedouterr0   r1   r-   r.   simple_usecase_callerget_cache_mtimesr+   r,   assertIsNotassertEqualr2   )r   r3   mtimesmod2r4   s        r   test_cache_reusez CUDACachingTest.test_cache_reuse]   sJ     "1S!1a 		!Q  3!!#//15!!!$&&(,,a3!!#d#	!Q1%	#s1% 	..0&9$$&..0&9r   c                 P   | j                         }|j                  }| j                   |dd      d       t        | j                  d      5 }|j                  d       d d d        | j                         }|j                  }| j                   |dd      d       y # 1 sw Y   ?xY w)Nr    r!   r"   az
Z = 10
   )r(   r)   r*   openmodfilewriter9   s      r   test_cache_invalidatez%CUDACachingTest.test_cache_invalidatex   s      "OO!Q+ $,,$GGL! %   "OO!Q, %$s   BB%c                    | j                         }|j                  }| j                   |dd      d       | j                         }|j                  }d|_        | j                   |dd      d       |j                  j                          | j                   |dd      d       | j                         }|j                  }| j                   |dd      d       y )Nr    r!   r"   rA   rX   )r(   r)   r*   Zr,   	recompiler9   s      r   test_recompilezCUDACachingTest.test_recompile   s      "OO!Q+  "OO!Q+	!Q,   "OO!Q,r   c                     | j                         }|j                  }| j                   |d      d       |j                  }| j                   |d      d       y )Nr       r@   )r(   renamed_function1r*   renamed_function2r9   s      r   test_same_nameszCUDACachingTest.test_same_names   sN      "!!!a(!!!a(r   zCG not supported with MVCc                     | j                  d       | j                         }| j                  d       |j                  d       | j                  d       | j                          y )Nr   r    )r'   r(   
cg_usecaser2   )r   r3   s     r   test_cache_cgzCUDACachingTest.test_cache_cg   sV     	1  "1q1 	$$&r   c           	         | j                  d       dt        | j                  | j                        z  }t	        j
                  t        j                  d|gt        j                  t        j                        }|j                  d      \  }}|j                  dk7  r;t        d|j                  d	|j                         d
|j                         d      y )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'   dictrj   rk   
subprocessPopensys
executablePIPEcommunicate
returncodeAssertionErrordecode)r   codepopenouterrs        r   test_cache_cg_clean_runz'CUDACachingTest.test_cache_cg_clean_run   s     	1 t||T\\BC   #..$!=(2(29 $$R$0Sq   ##SZZ\3::<A  !r   c                    | j                         }|j                  }| j                  t        j                  |j
                  j                  j                  d       | j                   |dd      d       | j                  |j
                  dd       | j                         }|j                  }| j                   |dd      d       | j                  |j
                  dd       | j                  d       y)	zy
        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   r3   r4   rT   s       r   _test_pycache_fallbackz&CUDACachingTest._test_pycache_fallback   s    
   "OO 	qvv||'>'>&* 	 	, 	!Q+1% !!#!Q+1% 	1r   ntz3cannot easily make a directory read-only on Windowsc                    t        j                  | j                        j                  }t        j                  | j                  d       | j                  t         j                  | j                  |       | j                          y )N@  )osstatrj   st_modechmodr   r   )r   	old_permss     r   test_non_creatable_pycachez*CUDACachingTest.test_non_creatable_pycache   sR    
 GGDLL)11	
u%$,,	:##%r   c                 N   t         j                  j                  | j                  d      }t        j                  |       t        j
                  |      j                  }t        j                  |d       | j                  t         j                  ||       | j                          y )N__pycache__r   )
r   pathjoinrj   mkdirr   r   r   r   r   )r   pycacher   s      r   test_non_writable_pycachez)CUDACachingTest.test_non_writable_pycache   sj    
 '',,t||];
GGG$,,	
% '95##%r   c                     t        t        dz        }d}| j                  t        |      5  t	        j
                  dd|g      d        }d d d        y # 1 sw Y   y xY w)Nzjitlink.ptxz0Cannot pickle CUDACodeLibrary with linking fileszvoid()T)cachelinkc                       y r    r   r   r   r4   z>CUDACachingTest.test_cannot_cache_linking_libraries.<locals>.f  s    r   )strr
   assertRaisesRegexRuntimeErrorr   jit)r   r   msgr4   s       r   #test_cannot_cache_linking_librariesz3CUDACachingTest.test_cannot_cache_linking_libraries   sS    ==01@##L#6XXhd$8 9 766s    AAN)"__name__
__module____qualname__r   r   dirname__file__herer   usecases_filerk   r   r   r6   r:   r=   rK   rU   r\   r`   re   r   r   r	   rh   r   r   r   unittestskipIfnamer   r   r   r   r   r   r   r      s*   77??8$DGGLL':;M(G!3'2
" :6--$) 45' 6  ' 45 6  84 X__RWW_JL&L & X__RWW_JL&L &r   r   c                       e Zd Zej                  j                  e      Zej                  j                  ed      Z	dZ
d Zd Zd Zd Zy)CUDAAndCPUCachingTestzcache_with_cpu_usecases.py cuda_and_cpu_caching_test_fodderc                 X    t        j                  |        t        j                  |        y r   r   r   s    r   r   zCUDAAndCPUCachingTest.setUp  r   r   c                 X    t        j                  |        t        j                  |        y r   r   r   s    r   r   zCUDAAndCPUCachingTest.tearDown  r   r   c                    | j                  d       | j                         }| j                  d       |j                  }|j                  }| j	                   |d      d       | j                  d       | j	                   |d      d       | j                  d       | j                  |j                  dd       | j                  |j                  dd       | j	                   |d      d       | j                  d       | j	                   |d      d       | j                  d       | j                  |j                  dd       | j                  |j                  dd       y )Nr   rC   r    r!   r%         @rb   )r'   r(   
assign_cpuassign_cudar*   r+   r,   )r   r3   f_cpuf_cudas       r   test_cpu_and_cuda_targetsz/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                    | j                         }|j                  d       |j                  d       |j                  d       |j                  d       | j                         }| j	                  |j                  j
                  dd       | j	                  |j                  j
                  dd       | j                         }| j                  ||       |j                  }|j                  } |d       | j	                  |j
                  dd        |d       | j	                  |j
                  dd        |d       | j	                  |j
                  dd        |d       | j	                  |j
                  dd       | j                  | j                         |       | j                          | j                  | j                         |       y )NrC   r   r   r    r%   r#   )	r(   r   r   rP   r+   r,   rQ   rR   r2   )r   r3   rS   rT   r   r   s         r   test_cpu_and_cuda_reusez-CUDAAndCPUCachingTest.test_cpu_and_cuda_reuse0  sb     "qs&&( 	++Q2,,a3!!#d#!!a

Aq)c


Aq)q	Q*sQ* 	..0&9$$&..0&9r   N)r   r   r   r   r   r   r   r   r   r   rk   r   r   r   r   r   r   r   r   r   	  sB    77??8$DGGLL'CDM0G!3+2 :r   r   c                     t         j                  d   } | 5  t        j                         j                  j                  }d d d        t         j                  dd  D ]J  }|5  t        j                         j                  j                  }|k7  r| |fcd d d        c S 	 d d d        L y # 1 sw Y   jxY w# 1 sw Y   cxY w)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~!3' S S   
 Ss   )B(1B4(B14B=	c                       e Zd Zej                  j                  e      Zej                  j                  ed      Z	dZ
d Zd Zd Zy)TestMultiCCCachingr   !cuda_multi_cc_caching_test_fodderc                 X    t        j                  |        t        j                  |        y r   r   r   s    r   r   zTestMultiCCCaching.setUpj  r   r   c                 X    t        j                  |        t        j                  |        y r   r   r   s    r   r   zTestMultiCCCaching.tearDownn  r   r   c                    t               }|s| j                  d       | j                  d       | j                         }| j                  d       |d   5  |j                  }| j                   |dd      d       | j                  d       | j                   |dd      d       | j                  d       | j                  |j                  dd       |j                  } ||j                  d      }| j                  t        |      d	       |j                  } ||j                  d      }| j                  t        |      d	       | j                  d       | j                  |j                  dd       d d d        |d   5  |j                  }| j                   |dd      d       | j                  d       | j                   |dd      d       | j                  d       | j                  |j                  dd       |j                  } ||j                  d      }| j                  t        |      d	       |j                  } ||j                  d      }| j                  t        |      d	       | j                  d       | j                  |j                  dd       d d d        | j                         }| j                  ||       |d   5  |j                  }| j                   |dd      d       | j                  d
       | j                   |dd      d       | j                  d       | j                  |j                  dd       |j                  } ||j                  d      }| j                  t        |      d	       |j                  } ||j                  d      }| j                  t        |      d	       | j                  d       | j                  |j                  dd       d d d        | j                         }| j                  ||       |d   5  |j                  }| j                   |dd      d       | j                   |dd      d       |j                  } ||j                  d      }| j                  t        |      d	       |j                  } ||j                  d      }| j                  t        |      d	       d d d        |d   5  |j                  }| j                   |dd      d       | j                   |dd      d       |j                  } ||j                  d      }| j                  t        |      d	       |j                  } ||j                  d      }| j                  t        |      d	       d d d        y # 1 sw Y   )xY w# 1 sw Y   	xY w# 1 sw Y   xY w# 1 sw Y   xY w# 1 sw Y   y xY w)Nz.Need two different CCs for multi-CC cache testr   r    r!   r"   r#   r$   r%   r&      r@   rA   )r   skipTestr'   r(   r)   r*   r+   r,   r-   r.   r/   r0   r1   rQ   )r   r   r3   r4   r5   rT   mod3s          r   
test_cachez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)! ( !!#d#!W  A##AaGQ/q!##Ac1Is3q!OOAFFAq)**ACOOQ'C##E#J	:))ACNNA&C##E#J	:r"OOAFFAq)! 0 !!#d# !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 V)D V6 B3W B3WV&)V36W WWN)r   r   r   r   r   r   r   r   r   r   rk   r   r   r   r   r   r   r   r   d  s>    77??8$DGGLL':;M1G!3l;r   r   c                  ,    ddl m}  d| _        d| _        y )Nr   config)
numba.corer   CUDA_LOW_OCCUPANCY_WARNINGSCUDA_WARN_ON_IMPLICIT_COPYr   s    r   child_initializerr     s     ")*F&()F%r   c                       e Zd ZdZej
                  j                  e      Zej
                  j                  ed      Z
dZd Zd Zd Zy)TestMultiprocessCacheFr   cuda_mp_caching_test_fodderc                 X    t        j                  |        t        j                  |        y r   r   r   s    r   r   zTestMultiprocessCache.setUp  r   r   c                 X    t        j                  |        t        j                  |        y r   r   r   s    r   r   zTestMultiprocessCache.tearDown  r   r   c                    | j                         }|j                  }d}	 t        j                  d      }|j                  |t              }	 t        |j                  |t        |                  }|j                          | j                  |||dz
  z  dz         y # t        $ r	 t        }Y xw xY w# |j                          w xY w)Nr!   spawnr%   r    )r(   rO   multiprocessingget_contextAttributeErrorPoolr   sumimaprangecloserR   )r   r3   r4   nctxpoolress          r   test_multiprocessingz*TestMultiprocessCache.test_multiprocessing  s      " %%	"!--g6C xx,-	dii58,-CJJLa1q5kQ./  	"!C	" JJLs   B $B1 B.-B.1CN)r   r   r   _numba_parallel_test_r   r   r   r   r   r   r   rk   r   r   r   r   r   r   r   r     sE    
 "77??8$DGGLL':;M+G!30r   r   z0Simulator does not implement the CUDACodeLibraryc                       e Zd Zd Zy)TestCUDACodeLibraryc                     ddl m} t               }d} |||      }| j                  t        d      5  |j                          d d d        y # 1 sw Y   y xY w)Nr   )CUDACodeLibrarylibraryzCannot pickle unfinalized)numba.cuda.codegenr   objectr   r   _reduce_states)r   r   codegenr   cls        r   !test_cannot_serialize_unfinalizedz5TestCUDACodeLibrary.test_cannot_serialize_unfinalized  sK     	7 (Wd+##L2MN ONNs   AAN)r   r   r   r   r   r   r   r   r     s    
 r   r   )r   r   r   rr   rt   r   rD   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   