
    J/PhJ                        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dS )    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                            e          Ze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!dS )CUDACachingTestcache_usecases.pycuda_caching_test_fodderc                 V    t          j        |            t          j        |            d S Nr   setUpr   selfs    d/var/www/html/test/jupyter/venv/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_caching.pyr   zCUDACachingTest.setUp   *    #)$///4         c                 V    t          j        |            t          j        |            d S r   r   tearDownr   r   s    r   r   zCUDACachingTest.tearDown   *    d####,T22222r   c                    |                      d           |                                 }|                      d           |j        }|                      |dd          d           |                      d           |                      |dd          d           |                      d           |                     |j        dd           |j        } ||j        d          }|                     t          |          d           |j	        } ||j
        d          }|                     t          |          d           |                      d           |                     |j        dd           |                                  d S )	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!   sv   1  ""1O!Q+++1#q		3///11%%%%a##c

I666$a""c

I66611%%% 	$$&&&&&r   c                     |                                  }|j        }|                      |dd          d           |                     d           d S )Nr    r!   r"   r   )r(   add_nocache_usecaser*   r'   r   r3   r4   s      r   test_no_cachingzCUDACachingTest.test_no_caching:   sU      ""#!Q+++1r   c                     |                      d           |                                 }|j        } |d                      |                      d           d S )Nr   )r%   r%   r    )r'   r(   many_localsr9   s      r   test_many_localsz CUDACachingTest.test_many_localsA   sY     	1  ""O$			1r   c                    |                                  }t          j                    5  t          j        dt                     |j        }|                      |d          d           |j        }|                      |d          d           |j        }|                      |d          d           |j	        }|                      |d          d           | 
                    d           d d d            d S # 1 swxY w Y   d S )Nerrorr!   r"      
         )r(   warningscatch_warningssimplefilterr   closure1r*   closure2closure3closure4r'   r9   s      r   test_closurezCUDACachingTest.test_closureM   sA     ""$&& 	" 	"!'<888A##AAaDD!,,,A##AAaDD!,,,A##AAaDD"---A##AAaDD"---q!!!	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	"s   CC==DDc                    |                                  }|                    dd           |                    dd           |                    dd           |                    dd           |                    |j        d           |                    |j        d           |                    d           | 	                                }| 
                    |j        j        dd           |                                  }|                     ||           |j        } |dd           | 
                    |j        dd            |dd           | 
                    |j        dd           |                     | 	                                |           |                                  |                     | 	                                |           d S )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]   s     ""1S!!!1a   		!Q  333!!#/1555!!!$$$&&((,a333!!##d###	!Q1%%%	#s1%%% 	..00&999$$&&&..00&99999r   c                 v   |                                  }|j        }|                      |dd          d           t          | j        d          5 }|                    d           d d d            n# 1 swxY w Y   |                                  }|j        }|                      |dd          d           d S )Nr    r!   r"   az
Z = 10
   )r(   r)   r*   openmodfilewriter9   s      r   test_cache_invalidatez%CUDACachingTest.test_cache_invalidatex   s      ""O!Q+++ $,$$ 	"GGL!!!	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	" 	"   ""O!Q,,,,,s   A33A7:A7c                    |                                  }|j        }|                      |dd          d           |                                  }|j        }d|_        |                      |dd          d           |j                                         |                      |dd          d           |                                  }|j        }|                      |dd          d           d S )Nr    r!   r"   rA   rX   )r(   r)   r*   Zr,   	recompiler9   s      r   test_recompilezCUDACachingTest.test_recompile   s      ""O!Q+++  ""O!Q+++	!Q,,,   ""O!Q,,,,,r   c                     |                                  }|j        }|                      |d          d           |j        }|                      |d          d           d S )Nr       r@   )r(   renamed_function1r*   renamed_function2r9   s      r   test_same_nameszCUDACachingTest.test_same_names   se      ""!!a(((!!a(((((r   zCG not supported with MVCc                     |                      d           |                                 }|                      d           |                    d           |                      d           |                                  d S )Nr   r    )r'   r(   
cg_usecaser2   )r   r3   s     r   test_cache_cgzCUDACachingTest.test_cache_cg   s}     	1  ""1q1 	$$&&&&&r   c           	         |                      d           dt          | j        | j                  z  }t	          j        t          j        d|gt          j        t          j                  }|	                    d          \  }}|j
        dk    rBt          d|j
        d	|                                d
|                                d          d S )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BBC  #.$!=(2(29 9 9 $$R$00Sq   . ###SZZ\\\\3::<<<<A   ! r   c                    |                                  }|j        }|                     t          j        |j        j        j        d           |                      |dd          d           | 	                    |j        dd           |                                  }|j        }|                      |dd          d           | 	                    |j        dd           | 
                    d           dS )	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    
   ""O 	qv|'>&* 	 	, 	, 	, 	!Q+++1%%% !!##!Q+++1%%% 	1r   ntz3cannot easily make a directory read-only on Windowsc                     t          j        | j                  j        }t          j        | j        d           |                     t           j        | j        |           |                                  d S )N@  )osstatrj   st_modechmodr   r   )r   	old_permss     r   test_non_creatable_pycachez*CUDACachingTest.test_non_creatable_pycache   s^    
 GDL))1	
u%%%$,	:::##%%%%%r   c                 >   t           j                            | j        d          }t          j        |           t          j        |          j        }t          j        |d           |                     t           j        ||           | 	                                 d S )N__pycache__r   )
r   pathjoinrj   mkdirr   r   r   r   r   )r   pycacher   s      r   test_non_writable_pycachez)CUDACachingTest.test_non_writable_pycache   s    
 ',,t|];;
GG$$,	
%   '9555##%%%%%r   c                     t          t          dz            }d}|                     t          |          5  t	          j        dd|g          d             }d d d            d S # 1 swxY w Y   d S )Nzjitlink.ptxz0Cannot pickle CUDACodeLibrary with linking fileszvoid()T)cachelinkc                      d S 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   s    ==011@##L#66 	 	Xhd$888  98	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	s   "A$$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       7??8$$DGLL':;;M(G! ! !3 3 3' ' '2  
 
 
" " " : : :6- - -- - -$) ) ) 455' ' 65  ' 455  65  8  4 X_RW_JL L& &L L _& X_RW_JL L& &L L _&    r   r   c                       e Zd Zej                            e          Zej                            ed          Z	dZ
d Zd Zd Zd ZdS )CUDAAndCPUCachingTestzcache_with_cpu_usecases.py cuda_and_cpu_caching_test_fodderc                 V    t          j        |            t          j        |            d S r   r   r   s    r   r   zCUDAAndCPUCachingTest.setUp  r   r   c                 V    t          j        |            t          j        |            d S r   r   r   s    r   r   zCUDAAndCPUCachingTest.tearDown  r   r   c                    |                      d           |                                 }|                      d           |j        }|j        }|                      |d          d           |                      d           |                      |d          d           |                      d           |                     |j        dd           |                     |j        dd           |                      |d          d           |                      d           |                      |d          d           |                      d           |                     |j        dd           |                     |j        dd           d S )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  sy    	1  ""1a!,,,1q		1---1
Aq)))Q***c

C0001sS1111
Aq)))Q*****r   c                    |                                  }|                    d           |                    d           |                    d           |                    d           |                                 }|                     |j        j        dd           |                     |j        j        dd           |                                  }|                     ||           |j        }|j        } |d           |                     |j        dd            |d           |                     |j        dd            |d           |                     |j        dd            |d           |                     |j        dd           |                     |                                 |           |                                  |                     |                                 |           d S )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  s     ""qs&&(( 	+Q222,a333!!##d###!a
Aq)))c



Aq)))q			Q***sQ*** 	..00&999$$&&&..00&99999r   N)r   r   r   r   r   r   r   r   r   r   rk   r   r   r   r   r   r   r   r   r   	  s~        7??8$$DGLL'CDDM0G! ! !3 3 3+ + +2 :  :  :  :  :r   r   c                  f   t           j        d         } | 5  t          j                    j        j        }d d d            n# 1 swxY w Y   t           j        dd          D ]R}|5  t          j                    j        j        }||k    r| |fcd d d            c S 	 d d d            n# 1 swxY w Y   Sd S )Nr   r%   )r   gpuscurrent_contextdevicecompute_capability)	first_gpufirst_ccgpuccs       r   get_different_cc_gpusr   S  sg    	!I	 D D'))0CD D D D D D D D D D D D D D D y} ( ( 	( 	(%''.ABX~~!3'	( 	( 	( 	( 	( 	( 	( 	( 	( 	(	( 	( 	( 	( 	( 	( 	( 	( 	( 	( 	( 	( 	( 	( 	(
 4s!   ?AA"'B%%B)	,B)	c                       e Zd Zej                            e          Zej                            ed          Z	dZ
d Zd Zd ZdS )TestMultiCCCachingr   !cuda_multi_cc_caching_test_fodderc                 V    t          j        |            t          j        |            d S r   r   r   s    r   r   zTestMultiCCCaching.setUpj  r   r   c                 V    t          j        |            t          j        |            d S r   r   r   s    r   r   zTestMultiCCCaching.tearDownn  r   r   c                    t                      }|s|                     d           |                     d           |                                 }|                     d           |d         5  |j        }|                      |dd          d           |                     d           |                      |dd          d           |                     d           |                     |j        dd           |j        } ||j	        d          }|                     t          |          d	           |j        } ||j        d          }|                     t          |          d	           |                     d           |                     |j        dd           d d d            n# 1 swxY w Y   |d         5  |j        }|                      |dd          d           |                     d           |                      |dd          d           |                     d           |                     |j        dd           |j        } ||j	        d          }|                     t          |          d	           |j        } ||j        d          }|                     t          |          d	           |                     d           |                     |j        dd           d d d            n# 1 swxY w Y   |                                 }|                     ||           |d         5  |j        }|                      |dd          d           |                     d
           |                      |dd          d           |                     d           |                     |j        dd           |j        } ||j	        d          }|                     t          |          d	           |j        } ||j        d          }|                     t          |          d	           |                     d           |                     |j        dd           d d d            n# 1 swxY w Y   |                                 }|                     ||           |d         5  |j        }|                      |dd          d           |                      |dd          d           |j        } ||j	        d          }|                     t          |          d	           |j        } ||j        d          }|                     t          |          d	           d d d            n# 1 swxY w Y   |d         5  |j        }|                      |dd          d           |                      |dd          d           |j        } ||j	        d          }|                     t          |          d	           |j        } ||j        d          }|                     t          |          d	           d d d            d S # 1 swxY w Y   d S )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  sv   $&& 	LMMJKKK1  ""1 !W 	* 	*A##AAaGGQ///q!!!##AAc1IIs333q!!!OOAFAq))))A!COQ''C##E#JJ	:::(A!CNA&&C##E#JJ	:::q!!!OOAFAq)))!	* 	* 	* 	* 	* 	* 	* 	* 	* 	* 	* 	* 	* 	* 	*( !W 	* 	*A##AAaGGQ///q!!!##AAc1IIs333q!!!OOAFAq))))A!COQ''C##E#JJ	:::(A!CNA&&C##E#JJ	:::q!!!OOAFAq)))!	* 	* 	* 	* 	* 	* 	* 	* 	* 	* 	* 	* 	* 	* 	*( !!##d###!W 	* 	* A##AAaGGQ///q!!!##AAc1IIs333q!!!OOAFAq)))*A!COQ''C##E#JJ	:::)A!CNA&&C##E#JJ	:::r"""OOAFAq)))!	* 	* 	* 	* 	* 	* 	* 	* 	* 	* 	* 	* 	* 	* 	*0 !!##d### !W 	; 	; A##AAaGGQ///##AAc1IIs333*A!COQ''C##E#JJ	:::)A!CNA&&C##E#JJ	:::	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; !W 	; 	; A##AAaGGQ///##AAc1IIs333*A!COQ''C##E#JJ	:::)A!CNA&&C##E#JJ	:::	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	;s_   ,D5F--F14F1 D5LLL>D5Q??RR<B>VV
V
B>Y$$Y(+Y(N)r   r   r   r   r   r   r   r   r   r   rk   r   r   r   r   r   r   r   r   d  st        7??8$$DGLL':;;M1G! ! !3 3 3l; l; l; l; l;r   r   c                  .    ddl m}  d| _        d| _        d S )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                            e          Zej        	                    ed          Z
dZd Zd Zd ZdS )TestMultiprocessCacheFr   cuda_mp_caching_test_fodderc                 V    t          j        |            t          j        |            d S r   r   r   s    r   r   zTestMultiprocessCache.setUp  r   r   c                 V    t          j        |            t          j        |            d S r   r   r   s    r   r   zTestMultiprocessCache.tearDown  r   r   c                    |                                  }|j        }d}	 t          j        d          }n# t          $ r
 t          }Y nw xY w|                    |t                    }	 t          |                    |t          |                              }|
                                 n# |
                                 w xY w|                     |||dz
  z  dz             d S )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66CC 	" 	" 	"!CCC	" xx,--	dii588,,--CJJLLLLDJJLLLLa1q5kQ./////s   4 AA'0B, ,CN)r   r   r   _numba_parallel_test_r   r   r   r   r   r   r   rk   r   r   r   r   r   r   r   r     sw        
 "7??8$$DGLL':;;M+G! ! !3 3 30 0 0 0 0r   r   z0Simulator does not implement the CUDACodeLibraryc                       e Zd Zd ZdS )TestCUDACodeLibraryc                     ddl m} t                      }d} |||          }|                     t          d          5  |                                 d d d            d S # 1 swxY w Y   d S )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  s     	766666 ((_Wd++##L2MNN 	  	 	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	 s   A  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 D D D D D D D D D D D D D D , + + + + +7 7 7 7 7 7 7 7 788r r r r rk#> r r 98rj 788F: F: F: F: F:K)D F: F: 98F:R  " 788y; y; y; y; y;&A y; y; 98y;x* * * 788$0 $0 $0 $0 $0K)D $0 $0 98$0N CDD         ,     ED     r   