§
    ƒÛgU  ã                   óø   — d dl Zd dlZd dlZd dlmZmZ d dlmZmZm	Z	 d dl
mZ d„ Zd„ Ze e	d¦  «         ed¦  «         G d	„ d
e¦  «        ¦   «         ¦   «         ¦   «         Zedk    r ej        ¦   «          dS dS )é    N)ÚunittestÚCUDATestCase)Úskip_on_cudasimÚskip_with_cuda_pythonÚskip_under_cuda_memcheck)Ú
linux_onlyc            	      óÌ  ‡‡‡‡‡‡‡‡‡‡‡‡— ddl mŠm} m} ddlm} dd l}dd l}dd lŠd|_	        | 
                    ¦   «         }t          j        |¦  «        }t          j        d¦  «        }|                     |¦  «         |                     t          j        ¦  «         d}d}	dŠ|j                             d¦  «         |j                             dd	||j        ¬
¦  «        Š|                     ‰¦  «        Šˆˆfd„t+          |	¦  «        D ¦   «         Šˆˆfd„t+          |	¦  «        D ¦   «         ŠdŠ|‰z  Š‰                     ¦   «         Š‰                      || d d d…         | d d d…         ¦  «        ¦  «        ˆˆfd„¦   «         Šˆˆˆˆˆˆfd„Šˆˆfd„t+          |	¦  «        D ¦   «         }
|
D ]}|                     ¦   «          Œ|
D ]}|                     ¦   «          Œ‰                     ¦   «          ‰‰z  }t+          |	¦  «        D ]5}|j                             ‰|                              ¦   «         |¦  «         Œ6|                     ¦   «          |                     ¦   «         S )Nr   )ÚcudaÚint32Úvoid)Úconfigé   znumba.cuda.cudadrv.driveri   é
   i   iè  )ÚlowÚhighÚsizeÚdtypec                 ó:   •— g | ]}‰                      ‰¦  «        ‘ŒS © ©Ú	to_device)Ú.0Ú_r
   Úxs     €€úX/granary/edev/quantum/lib/python3.11/site-packages/numba/cuda/tests/cudadrv/test_ptds.pyú
<listcomp>zchild_test.<locals>.<listcomp>)   ó%   ø€ Ð	6Ð	6Ð	6 ˆ$�.Š.˜Ñ
Ô
Ð	6Ð	6Ð	6ó    c                 ó:   •— g | ]}‰                      ‰¦  «        ‘ŒS r   r   )r   r   r
   Úrs     €€r   r   zchild_test.<locals>.<listcomp>*   r   r   é   c                 ó¬   •— ‰                      d¦  «        }|t          | ¦  «        k    rd S t          ‰¦  «        D ]}| |xx         ||         z  cc<   Œd S )Nr   )ÚgridÚlenÚrange)r    r   ÚiÚjÚN_ADDITIONSr
   s       €€r   Úfzchild_test.<locals>.f3   sd   ø€ à�IŠI�a‰LŒLˆà�s�1‰vŒvŠ:ˆ:ØˆFõ �{Ñ#Ô#ð 	ð 	ˆAØˆaˆDˆDŒD�A�a”D‰LˆDˆD‰DˆDð	ð 	r   c                 óJ   •—  ‰‰‰‰f         ‰|          ‰|          ¦  «         d S )Nr   )Únr)   Ún_blocksÚ	n_threadsÚrsÚstreamÚxss    €€€€€€r   Úkernel_threadz!child_test.<locals>.kernel_thread@   s/   ø€ Ø&ˆˆ(�I˜vÐ
%Ô& r¨!¤u¨b°¬eÑ4Ô4Ð4Ð4Ð4r   c                 ó@   •— g | ]}‰                      ‰|f¬ ¦  «        ‘ŒS )©ÚtargetÚargs)ÚThread)r   r&   r1   Ú	threadings     €€r   r   zchild_test.<locals>.<listcomp>D   s>   ø€ ð *ð *ð *Øð ×Ò }¸A¸4ÐÑ@Ô@ð *ð *ð *r   ) Únumbar
   r   r   Ú
numba.corer   ÚioÚnumpyr7   ÚCUDA_PER_THREAD_DEFAULT_STREAMÚStringIOÚloggingÚStreamHandlerÚ	getLoggerÚ
addHandlerÚsetLevelÚDEBUGÚrandomÚseedÚrandintÚ
zeros_liker%   Údefault_streamÚjitÚstartÚjoinÚsynchronizeÚtestingÚassert_equalÚcopy_to_hostÚflushÚgetvalue)r   r   r   r:   ÚnpÚlogbufÚhandlerÚcudadrv_loggerÚNÚ	N_THREADSÚthreadsÚthreadÚexpectedr&   r(   r
   r)   r1   r,   r-   r    r.   r/   r7   r   r0   s                 @@@@@@@@@@@@r   Ú
child_testr[   
   s  øøøøøøøøøøøø€ Ø'Ð'Ð'Ð'Ð'Ð'Ð'Ð'Ð'Ð'Ø!Ð!Ð!Ð!Ð!Ð!Ø€I€I€IØÐÐÐØÐÐÐð
 -.€FÔ)ð �[Š[‰]Œ]€FÝÔ# FÑ+Ô+€GÝÔ&Ð'BÑCÔC€NØ×Ò˜gÑ&Ô&Ð&Ø×Ò�GœMÑ*Ô*Ð*ð 	€AØ€IØ€Kð „I‡N‚N�1ÑÔÐØ
Œ	×Ò˜a d°¸"¼(ÐÑCÔC€AØ
�Š�aÑÔ€Að 
7Ð	6Ð	6Ð	6Ð	6¥U¨9Ñ%5Ô%5Ð	6Ñ	6Ô	6€BØ	6Ð	6Ð	6Ð	6Ð	6¥U¨9Ñ%5Ô%5Ð	6Ñ	6Ô	6€Bð €IØ�I‰~€HØ× Ò Ñ"Ô"€Fð 
‡X‚Xˆdˆd�5˜˜˜1˜”:˜u S S q SœzÑ*Ô*Ñ+Ô+ðð ð ð ð ñ ,Ô+ðð5ð 5ð 5ð 5ð 5ð 5ð 5ð 5ð 5ð 5ð*ð *ð *ð *ð *Ý˜iÑ(Ô(ð*ñ *ô *€Gð ð ð ˆØ�Š‰Œˆˆð ð ð ˆØ�Š‰Œˆˆð 	×ÒÑÔÐð �;‰€HÝ�9ÑÔð @ð @ˆØ
Œ
×Ò  1¤× 2Ò 2Ñ 4Ô 4°hÑ?Ô?Ð?Ð?ð ‡M‚M�O„O€OØ�?Š?ÑÔÐr   c                 ó�   — 	 t          ¦   «         }d}n#  t          j        ¦   «         }d}Y nxY w|                      ||f¦  «         d S )NTF)r[   Ú	tracebackÚ
format_excÚput)Úresult_queueÚoutputÚsuccesss      r   Úchild_test_wrapperrc   ]   sV   € ðÝ‘”ˆØˆˆøðÝÔ%Ñ'Ô'ˆØˆˆˆøøøà×Ò�g˜vÐ&Ñ'Ô'Ð'Ð'Ð's   ‚ “,zHangs cuda-memcheckz&Streams not supported on the simulatorc                   ó6   — e Zd Z ed¦  «        d„ ¦   «         ZdS )ÚTestPTDSz1Function names unchanged for PTDS with NV Bindingc                 óv  — t          j        d¦  «        }|                     ¦   «         }|                     t          |f¬¦  «        }|                     ¦   «          |                     ¦   «          |                     ¦   «         \  }}|s|                      |¦  «         d}|D ]G}|  	                    |d¬¦  «        5  |  
                    ||¦  «         d d d ¦  «         n# 1 swxY w Y   ŒHd}|D ]L}|  	                    |d¬¦  «        5  |› d�}	|                      |	|¦  «         d d d ¦  «         n# 1 swxY w Y   ŒMd S )	NÚspawnr3   )ÚcuMemcpyHtoD_v2_ptdsÚcuLaunchKernel_ptszÚcuMemcpyDtoH_v2_ptdsT)ÚfnrZ   )ÚcuMemcpyHtoD_v2ÚcuLaunchKernelÚcuMemcpyDtoH_v2Fú
)ÚmpÚget_contextÚQueueÚProcessrc   rJ   rK   ÚgetÚfailÚsubTestÚassertInÚassertNotIn)
ÚselfÚctxr`   Úprocrb   ra   Úptds_functionsrk   Úlegacy_functionsÚ	fn_at_ends
             r   Ú	test_ptdszTestPTDS.test_ptdso   sÍ  € õ Œn˜WÑ%Ô%ˆØ—y’y‘{”{ˆØ�{Š{Õ"4¸L¸?ˆ{ÑKÔKˆØ�
Š
‰ŒˆØ�	Š	‰ŒˆØ&×*Ò*Ñ,Ô,‰ˆ�ð ð 	Ø�IŠI�fÑÔÐð2ˆð !ð 	*ð 	*ˆBØ—’ ¨d�Ñ3Ô3ð *ð *Ø—’˜b &Ñ)Ô)Ð)ð*ð *ð *ñ *ô *ð *ð *ð *ð *ð *ð *øøøð *ð *ð *ð *øð
/Ðð #ð 	4ð 	4ˆBØ—’ ¨e�Ñ4Ô4ð 4ð 4ð  "˜I˜I˜I�	Ø× Ò  ¨FÑ3Ô3Ð3ð4ð 4ð 4ñ 4ô 4ð 4ð 4ð 4ð 4ð 4ð 4øøøð 4ð 4ð 4ð 4øð	4ð 	4s$   Â9CÃC 	Ã#C 	ÄD-Ä-D1	Ä4D1	N)Ú__name__Ú
__module__Ú__qualname__r   r   r   r   r   re   re   k   s=   € € € € € ð ÐÐNÑOÔOð!4ð !4ñ PÔOð!4ð !4ð !4r   re   Ú__main__)Úmultiprocessingrp   r>   r]   Únumba.cuda.testingr   r   r   r   r   Únumba.tests.supportr   r[   rc   re   r€   Úmainr   r   r   ú<module>rˆ      s.  ðØ Ð Ð Ð Ø €€€Ø Ð Ð Ð Ø 5Ð 5Ð 5Ð 5Ð 5Ð 5Ð 5Ð 5ð:ð :ð :ð :ð :ð :ð :ð :ð :ð :à *Ð *Ð *Ð *Ð *Ð *ðPð Pð Pðf	(ð 	(ð 	(ð ØÐÐ/Ñ0Ô0Ø€Ð9Ñ:Ô:ð#4ð #4ð #4ð #4ð #4ˆ|ñ #4ô #4ñ ;Ô:ñ 1Ô0ñ „ð#4ðL ˆzÒÐØ€H„M�O„O€O€O€Oð Ðr   