§
    ƒÛg¦E  ã                   óÌ  — 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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    úZ/granary/edev/quantum/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_caching8   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_locals?   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_closureK   sA  € Ø× Ò Ñ"Ô"ˆåÔ$Ñ&Ô&ð 	"ð 	"ÝÔ! '­<Ñ8Ô8Ð8à”ˆAØ×#Ò# A A a¡D¤D¨!Ñ,Ô,Ð,Ø”ˆAØ×#Ò# A A a¡D¤D¨!Ñ,Ô,Ð,Ø”ˆAØ×#Ò# A A a¡D¤D¨"Ñ-Ô-Ð-Ø”ˆAØ×#Ò# A A a¡D¤D¨"Ñ-Ô-Ð-Ø×Ò˜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Ñ!Ô!Ð!Ø×Ò˜1˜aÑ Ô Ð Ø�	Š	�!�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_invalidatev   sò   € Ø× Ò Ñ"Ô"ˆØŒOˆØ×Ò   ! Q¡¤¨Ñ+Ô+Ð+õ �$”, Ñ$Ô$ð 	"¨Ø�GŠG�LÑ!Ô!Ð!ð	"ð 	"ð 	"ñ 	"ô 	"ð 	"ð 	"ð 	"ð 	"ð 	"ð 	"øøøð 	"ð 	"ð 	"ð 	"ð × Ò Ñ"Ô"ˆØŒ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                 óî  — |                       ¦   «         }|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ˆð 	�Š�œ q¤v¤|Ô'>Ø&*ð 	ñ 	,ô 	,ð 	,ð 	×Ò   ! 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ÚstatÚtempdirÚst_modeÚchmodrk   rp   )r   Ú	old_permss     r   Útest_non_creatable_pycachez*CUDACachingTest.test_non_creatable_pycacheÇ   s^   € õ
 ”G˜DœLÑ)Ô)Ô1ˆ	Ý
Œ�”˜uÑ%Ô%Ð%Ø�Š�œ $¤,°	Ñ:Ô:Ð:à×#Ò#Ñ%Ô%Ð%Ð%Ð%r   c                 ó>  — t           j                             | j        d¦  «        }t          j        |¦  «         t          j        |¦  «        j        }t          j        |d¦  «         |                      t           j        ||¦  «         |  	                    ¦   «          d S )NÚ__pycache__rs   )
rt   ÚpathÚjoinrv   Úmkdirru   rw   rx   rk   rp   )r   Úpycachery   s      r   Útest_non_writable_pycachez)CUDACachingTest.test_non_writable_pycacheÒ   s   € õ
 ”'—,’,˜tœ|¨]Ñ;Ô;ˆÝ
Œ�ÑÔÐÝ”G˜GÑ$Ô$Ô,ˆ	Ý
Œ�˜%Ñ Ô Ð Ø�Š�œ '¨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ð 	ð 	ÝŒX�h d°$°Ð8Ñ8Ô8ðð ñ 9Ô8ðð	ð 	ð 	ñ 	ô 	ð 	ð 	ð 	ð 	ð 	ð 	ð 	øøøð 	ð 	ð 	ð 	ð 	ð 	s   µ"A$Á$A(Á+A(N)!Ú__name__Ú
__module__Ú__qualname__rt   r}   ÚdirnameÚ__file__Úherer~   Úusecases_fileÚmodnamer   r   r6   r:   r=   rK   rU   r\   r`   re   r   r   r	   rh   rp   r   ÚunittestÚskipIfÚnamerz   r�   rŒ   r†   r   r   r   r      s¾  € € € € € àŒ7�?Š?˜8Ñ$Ô$€DØ”G—L’L Ð':Ñ;Ô;€MØ(€Gð!ð !ð !ð3ð 3ð 3ð'ð 'ð 'ð2ð ð ð
ð 
ð 
ð"ð "ð "ð :ð :ð :ð6-ð -ð -ð-ð -ð -ð$)ð )ð )ð ØØÐÐ4Ñ5Ô5ð'ð 'ñ 6Ô5ñ Ôñ Ôð'ðð ð ð4 Ø€X„_�R”W ’_ØJñLô Lð&ð &ñLô Lñ „_ð&ð Ø€X„_�R”W ’_Ø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ÑÔÐà�Š˜œ
 A qÑ)Ô)Ð)Ø�Š˜œ Q¨Ñ*Ô*Ð*à×Ò   c¡
¤
¨CÑ0Ô0Ð0Ø×Ò˜1ÑÔÐØ×Ò   s¡¤¨SÑ1Ô1Ð1Ø×Ò˜1ÑÔÐà�Š˜œ
 A qÑ)Ô)Ð)Ø�Š˜œ 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_reuse  sÌ  € à× Ò Ñ"Ô"ˆØ�Š�qÑÔÐØ�Š�sÑÔÐØ�Š˜ÑÔÐØ�Š˜ÑÔÐà×&Ò&Ñ(Ô(ˆð 	�Š˜œÔ+¨Q°Ñ2Ô2Ð2Ø�Š˜œÔ,¨a°Ñ3Ô3Ð3à×!Ò!Ñ#Ô#ˆØ×Ò˜˜dÑ#Ô#Ð#Ø”ˆØÔ!ˆàˆˆa‰ŒˆØ�Š˜œ
 A qÑ)Ô)Ð)Øˆˆc‰
Œ
ˆ
Ø�Š˜œ
 A qÑ)Ô)Ð)Øˆˆq‰	Œ	ˆ	Ø�Š˜œ Q¨Ñ*Ô*Ð*Øˆˆs‰ŒˆØ�Š˜œ Q¨Ñ*Ô*Ð*ð 	×Ò˜×.Ò.Ñ0Ô0°&Ñ9Ô9Ð9à×$Ò$Ñ&Ô&Ð&Ø×Ò˜×.Ò.Ñ0Ô0°&Ñ9Ô9Ð9Ð9Ð9r   N)r�   rŽ   r�   rt   r}   r�   r‘   r’   r~   r“   r”   r   r   r£   r¥   r†   r   r   r™   r™   è   s~   € € € € € àŒ7�?Š?˜8Ñ$Ô$€DØ”G—L’L Ð'CÑDÔD€MØ0€Gð!ð !ð !ð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¯   2  sg  € õ ”	˜!”€IØ	ð Dð DÝÔ'Ñ)Ô)Ô0ÔCˆðDð Dð Dñ Dô Dð Dð Dð Dð Dð Dð Døøøð Dð Dð Dð Dõ Œy˜˜˜Œ}ð (ð (ˆØð 	(ð 	(ÝÔ%Ñ'Ô'Ô.ÔAˆBØ�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.setUpI  r   r   c                 óV   — t          j        | ¦  «         t          j        | ¦  «         d S r   r   r   s    r   r   zTestMultiCCCaching.tearDownM  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_cacheQ  sv  € Ý$Ñ&Ô&ˆØð 	LØ�MŠMÐJÑKÔKÐKà×Ò˜1ÑÔÐØ× Ò Ñ"Ô"ˆØ×Ò˜1ÑÔÐð �!ŒWð 	*ð 	*Ø”ˆAØ×#Ò# A A a¨¡G¤G¨QÑ/Ô/Ð/Ø×Ò˜qÑ!Ô!Ð!Ø×#Ò# A A c¨1¡I¤I¨sÑ3Ô3Ð3Ø×Ò˜qÑ!Ô!Ð!Ø�OŠO˜AœF A qÑ)Ô)Ð)àÔ)ˆAØ�!�C”O QÑ'Ô'ˆCØ×#Ò#¥E¨#¡J¤J°	Ñ:Ô:Ð:àÔ(ˆAØ�!�C”N AÑ&Ô&ˆCØ×#Ò#¥E¨#¡J¤J°	Ñ:Ô:Ð:Ø×Ò˜qÑ!Ô!Ð!Ø�OŠO˜AœF A qÑ)Ô)Ð)ð!	*ð 	*ð 	*ñ 	*ô 	*ð 	*ð 	*ð 	*ð 	*ð 	*ð 	*øøøð 	*ð 	*ð 	*ð 	*ð( �!ŒWð 	*ð 	*Ø”ˆAØ×#Ò# A A a¨¡G¤G¨QÑ/Ô/Ð/Ø×Ò˜qÑ!Ô!Ð!Ø×#Ò# A A c¨1¡I¤I¨sÑ3Ô3Ð3Ø×Ò˜qÑ!Ô!Ð!Ø�OŠO˜AœF A qÑ)Ô)Ð)àÔ)ˆAØ�!�C”O QÑ'Ô'ˆCØ×#Ò#¥E¨#¡J¤J°	Ñ:Ô:Ð:àÔ(ˆAØ�!�C”N AÑ&Ô&ˆCØ×#Ò#¥E¨#¡J¤J°	Ñ:Ô:Ð:Ø×Ò˜qÑ!Ô!Ð!Ø�OŠO˜AœF A qÑ)Ô)Ð)ð!	*ð 	*ð 	*ñ 	*ô 	*ð 	*ð 	*ð 	*ð 	*ð 	*ð 	*øøøð 	*ð 	*ð 	*ð 	*ð( ×!Ò!Ñ#Ô#ˆØ×Ò˜˜dÑ#Ô#Ð#à�!ŒWð 	*ð 	*ØÔ ˆAØ×#Ò# A A a¨¡G¤G¨QÑ/Ô/Ð/Ø×Ò˜qÑ!Ô!Ð!Ø×#Ò# A A c¨1¡I¤I¨sÑ3Ô3Ð3Ø×Ò˜qÑ!Ô!Ð!Ø�OŠO˜AœF A qÑ)Ô)Ð)àÔ*ˆAØ�!�C”O QÑ'Ô'ˆCØ×#Ò#¥E¨#¡J¤J°	Ñ:Ô:Ð:àÔ)ˆAØ�!�C”N AÑ&Ô&ˆCØ×#Ò#¥E¨#¡J¤J°	Ñ:Ô:Ð:Ø×Ò˜rÑ"Ô"Ð"Ø�OŠO˜AœF A qÑ)Ô)Ð)ð!	*ð 	*ð 	*ñ 	*ô 	*ð 	*ð 	*ð 	*ð 	*ð 	*ð 	*øøøð 	*ð 	*ð 	*ð 	*ð0 ×!Ò!Ñ#Ô#ˆØ×Ò˜˜dÑ#Ô#Ð#ð �!ŒWð 	;ð 	;ØÔ ˆAØ×#Ò# A A a¨¡G¤G¨QÑ/Ô/Ð/Ø×#Ò# A A c¨1¡I¤I¨sÑ3Ô3Ð3àÔ*ˆAØ�!�C”O QÑ'Ô'ˆCØ×#Ò#¥E¨#¡J¤J°	Ñ:Ô:Ð:àÔ)ˆAØ�!�C”N AÑ&Ô&ˆCØ×#Ò#¥E¨#¡J¤J°	Ñ:Ô:Ð:ð	;ð 	;ð 	;ñ 	;ô 	;ð 	;ð 	;ð 	;ð 	;ð 	;ð 	;øøøð 	;ð 	;ð 	;ð 	;ð �!ŒWð 	;ð 	;ØÔ ˆAØ×#Ò# A A a¨¡G¤G¨QÑ/Ô/Ð/Ø×#Ò# A A c¨1¡I¤I¨sÑ3Ô3Ð3àÔ*ˆAØ�!�C”O QÑ'Ô'ˆCØ×#Ò#¥E¨#¡J¤J°	Ñ:Ô:Ð:àÔ)ˆAØ�!�C”N AÑ&Ô&ˆCØ×#Ò#¥E¨#¡J¤J°	Ñ:Ô:Ð:ð	;ð 	;ð 	;ñ 	;ô 	;ð 	;ð 	;ð 	;ð 	;ð 	;ð 	;ð 	;øøøð 	;ð 	;ð 	;ð 	;ð 	;ð 	;s_   Á,D5F-Æ-F1Æ4F1Ç D5LÌLÌLÌ>D5Q?Ñ?RÒRÒ<B>VÖV
ÖV
ÖB>Y$Ù$Y(Ù+Y(N)r�   rŽ   r�   rt   r}   r�   r‘   r’   r~   r“   r”   r   r   r¹   r†   r   r   r±   r±   C  st   € € € € € àŒ7�?Š?˜8Ñ$Ô$€DØ”G—L’L Ð':Ñ;Ô;€MØ1€Gð!ð !ð !ð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Ô6ˆCˆCøÝð 	"ð 	"ð 	"Ý!ˆCˆCˆCð	"øøøð �xŠx˜Õ,Ñ-Ô-ˆð	Ý�d—i’i ¥5¨¡8¤8Ñ,Ô,Ñ-Ô-ˆCà�JŠJ‰LŒLˆLˆLøˆD�JŠJ‰LŒLˆLˆLøøøØ×Ò˜˜a 1 q¡5™k¨QÑ.Ñ/Ô/Ð/Ð/Ð/s   Ÿ4 ´AÁAÁ'0B, Â,CN)r�   rŽ   r�   Ú_numba_parallel_test_rt   r}   r�   r‘   r’   r~   r“   r”   r   r   rÔ   r†   r   r   rÂ   rÂ   È  sw   € € € € € ð
 "ÐàŒ7�?Š?˜8Ñ$Ô$€DØ”G—L’L Ð':Ñ;Ô;€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õ ‘(”(ˆØˆØˆ_˜W dÑ+Ô+ˆØ×#Ò#¥LÐ2MÑNÔNð 	 ð 	 Ø×ÒÑÔÐð	 ð 	 ð 	 ñ 	 ô 	 ð 	 ð 	 ð 	 ð 	 ð 	 ð 	 ð 	 øøøð 	 ð 	 ð 	 ð 	 ð 	 ð 	 s   ¾A Á A$Á'A$N)r�   rŽ   r�   rà   r†   r   r   r×   r×   ð  s#   € € € € € ð
 ð  ð  ð  ð  r   r×   )rÈ   rt   rl   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ðSð Sð Sð Sð S�kÐ#>ñ Sô Sñ 9Ô8ðSðl €Ð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ð $0˜KÐ)Dñ $0ô $0ñ 9Ô8ð$0ðN €ÐCÑDÔDð ð  ð  ð  ð  ˜,ñ  ô  ñ EÔDð ð  ð  r   