
    xKgU                         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                           yy)    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|_	        |j                         }t        j                  |      }t        j                  d      }|j                  |       |j                  t        j                          d}	d}
d|j"                  j%                  d       |j"                  j'                  dd	|	|j                  
      }|j)                  |      }t+        |
      D cg c]  }j-                  |       c}t+        |
      D cg c]  }j-                  |       c}d|	z  j/                         j1                   || d d d   | d d d               fd       fd}t+        |
      D cg c]  }|j3                  ||f       }}|D ]  }|j5                           |D ]  }|j7                           j9                          |z  }t+        |
      D ]/  }|j:                  j=                  |   j?                         |       1 |jA                          |jC                         S c c}w c c}w c c}w )Nr   )cudaint32void)config   znumba.cuda.cudadrv.driveri   
   i   i  )lowhighsizedtype   c                     j                  d      }|t        |       kD  ry t              D ]  }| |xx   ||   z  cc<    y )Nr   )gridlenrange)rxijN_ADDITIONSr
   s       f/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudadrv/test_ptds.pyfzchild_test.<locals>.f3   s@    IIaLs1v: {#AaDAaDLD $    c                 0     f   |    |           y )N )nr   n_blocks	n_threadsrsstreamxss    r   kernel_threadz!child_test.<locals>.kernel_thread@   s#    &(Iv
%&r!ube4r    targetargs)"numbar
   r   r   
numba.corer   ionumpy	threadingCUDA_PER_THREAD_DEFAULT_STREAMStringIOloggingStreamHandler	getLogger
addHandlersetLevelDEBUGrandomseedrandint
zeros_liker   	to_devicedefault_streamjitThreadstartjoinsynchronizetestingassert_equalcopy_to_hostflushgetvalue)r   r   r   r/   npr1   logbufhandlercudadrv_loggerN	N_THREADSr   r   _r)   r   threadsthreadexpectedr   r
   r   r$   r%   r&   r'   r(   s                      @@@@@@@@r   
child_testrT   
   s7   ''!
 -.F) [[]F##F+G&&'BCNg&GMM* 	AIK IINN1
		ad"((CA
aA &+9%5	6%5$..
%5	6B%*9%5	6%5$..
%5	6B II~H  "F 
XXd51:uSqSz*+ ,5 5
 i(*(1 }A4@(  *  
   	 ;H9


1 2 2 4h?  MMO??c 
7	64*s   >I&II#c                 ~    	 t               }d}| j                  ||f       y #  t        j                         }d}Y .xY w)NTF)rT   	traceback
format_excput)result_queueoutputsuccesss      r   child_test_wrapperr\   ]   sB     gv&'	%%'s   " <zHangs cuda-memcheckz&Streams not supported on the simulatorc                   (    e Zd Z ed      d        Zy)TestPTDSz1Function names unchanged for PTDS with NV Bindingc                     t        j                  d      }|j                         }|j                  t        |f      }|j                          |j                          |j                         \  }}|s| j                  |       d}|D ]0  }| j                  |d      5  | j                  ||       d d d        2 d}|D ]5  }| j                  |d      5  | d}	| j                  |	|       d d d        7 y # 1 sw Y   yxY w# 1 sw Y   NxY w)	Nspawnr*   )cuMemcpyHtoD_v2_ptdscuLaunchKernel_ptszcuMemcpyDtoH_v2_ptdsT)fnrS   )cuMemcpyHtoD_v2cuLaunchKernelcuMemcpyDtoH_v2F
)mpget_contextQueueProcessr\   rB   rC   getfailsubTestassertInassertNotIn)
selfctxrY   procr[   rZ   ptds_functionsrd   legacy_functions	fn_at_ends
             r   	test_ptdszTestPTDS.test_ptdso   s     nnW%yy{{{"4L?{K

		&**, IIf2 !Bd3b&) 43 !/ #Be4  "d"I	  F3 54 # 43 54s   C8D8D	D	N)__name__
__module____qualname__r   rx   r"   r    r   r^   r^   k   s     NO!4 P!4r    r^   __main__)multiprocessingri   r4   rV   numba.cuda.testingr   r   r   r   r   numba.tests.supportr   rT   r\   r^   ry   mainr"   r    r   <module>r      s|       5: : *Pf	( /09:#4| #4 ; 1 #4L zHMMO r    