a
    ò.;bÖ&  ã                   @   sø   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 d dlZdd„ Zdd„ Zd	d
„ Zdd„ ZedƒedƒedƒG dd„ deƒƒƒƒZdd„ Zdd„ ZedƒedƒedƒG dd„ deƒƒƒƒZedkrôe ¡  dS )é    N)Úcuda)Údriver)Úskip_on_armÚskip_on_cudasimÚskip_under_cuda_memcheckÚContextResettingTestCaseÚForeignArrayc                 C   s>   z
| ƒ }W n   d}t  ¡ }Y n
0 d}|}| ||f¡ d S )NFT)Ú	tracebackÚ
format_excÚput)Úthe_workÚresult_queueÚarrÚsuccÚout© r   ú?lib/python3.9/site-packages/numba/cuda/tests/cudapy/test_ipc.pyÚcore_ipc_handle_test   s    
r   c                    s   ‡ ‡fdd„}t ||ƒ d S )Nc                     sN   t  t j¡} tjˆ ˆ| j | d}| ¡ W  d   ƒ S 1 s@0    Y  d S ©N)ÚshapeÚdtype)Únpr   Úintpr   Zopen_ipc_arrayÚitemsizeÚcopy_to_host)r   Údarr©ÚhandleÚsizer   r   r       s    ÿz&base_ipc_handle_test.<locals>.the_work©r   )r   r   r   r   r   r   r   Úbase_ipc_handle_test   s    r    c                    s   ‡ fdd„}t ||ƒ d S )Nc                     s<   t  t j¡} ˆ jt ¡ ˆ j| j | d}| ¡ }ˆ  	¡  |S r   )
r   r   r   Z
open_arrayr   Úcurrent_contextr   r   r   Úclose)r   r   r   ©r   r   r   r   +   s    

þz+serialize_ipc_handle_test.<locals>.the_workr   )r   r   r   r   r#   r   Úserialize_ipc_handle_test*   s    
r$   c                 C   sÎ   zš| ‚}|  ¡ }z(|  W d   ƒ n1 s,0    Y  W n8 typ } z t|ƒdkr\tdƒ‚W Y d }~nd }~0 0 tdƒ‚W d   ƒ n1 sŽ0    Y  W n   d}t ¡ }Y n
0 d}|}| ||f¡ d S ©NzIpcHandle is already openedzinvalid exception messagezdid not raise on reopenFT)r   Ú
ValueErrorÚstrÚAssertionErrorr	   r
   r   )Úipcarrr   r   r   Úer   r   r   r   r   Úipc_array_test8   s     "*r+   zHangs cuda-memcheckzIpc not available in CUDASIMz&CUDA IPC not supported on ARM in Numbac                   @   s@   e Zd Zdd„ Zdd„ Zddd„Zd	d
„ Zddd„Zdd„ ZdS )ÚTestIpcMemoryc                 C   s¸   t jdt jd}t |¡}t ¡ }| |j¡}tj	r>|j
j}n
t|j
ƒ}|j}t d¡}| ¡ }|||f}|jt|d}	|	 ¡  | ¡ \}
}|
sœ|  |¡ nt j ||¡ |	 d¡ d S ©Né
   ©r   Úspawn©ÚtargetÚargsé   )r   Úaranger   r   Ú	to_devicer!   Úget_ipc_handleÚgpu_datar   ÚUSE_NV_BINDINGr   ÚreservedÚbytesr   ÚmpÚget_contextÚQueueÚProcessr    ÚstartÚgetÚfailÚtestingÚassert_equalÚjoin)Úselfr   ÚdevarrÚctxÚipchZhandle_bytesr   r   r3   Úprocr   r   r   r   r   Útest_ipc_handleV   s$    




zTestIpcMemory.test_ipc_handlec                 C   s.   d t dd ƒt ddƒt d dƒf}d}t ||¡S )Nr4   é   )FT)ÚsliceÚ	itertoolsÚproduct)rF   ÚindicesZforeignsr   r   r   Úvariantss   s    zTestIpcMemory.variantsNFc                 C   s(  t jdt jd}t |¡}|d ur*|| }|r<t t|ƒ¡}| ¡ }t ¡ }| 	|j
¡}t |¡}t |¡}	|  |	jd ¡ |  |	j|j¡ tjr¦|  |	jj|jj¡ n|  t|	jƒt|jƒ¡ t d¡}| ¡ }
||
f}|jt|d}| ¡  |
 ¡ \}}|s|  |¡ nt j  ||¡ | !d¡ d S r-   )"r   r5   r   r   r6   Úas_cuda_arrayr   r   r!   r7   r8   ÚpickleÚdumpsÚloadsÚassertIsÚbaseÚassertEqualr   r   r9   r   r:   Útupler<   r=   r>   r?   r$   r@   rA   rB   rC   rD   rE   )rF   Ú	index_argÚforeignr   rG   ÚexpectrH   rI   ÚbufÚ
ipch_reconr   r3   rJ   r   r   r   r   r   Úcheck_ipc_handle_serialization{   s4    



z,TestIpcMemory.check_ipc_handle_serializationc              	   C   sP   |   ¡ D ]B\}}| j||d |  ||¡ W d   ƒ q1 s@0    Y  qd S ©N)Úindexr[   )rQ   ÚsubTestr_   ©rF   ra   r[   r   r   r   Útest_ipc_handle_serialization¡   s    z+TestIpcMemory.test_ipc_handle_serializationc                 C   s´   t jdt jd}t |¡}|d ur*|| }|r<t t|ƒ¡}| ¡ }| ¡ }t	 
d¡}| ¡ }||f}	|jt|	d}
|
 ¡  | ¡ \}}|s˜|  |¡ nt j ||¡ |
 d¡ d S r-   )r   r5   r   r   r6   rR   r   r   r7   r<   r=   r>   r?   r+   r@   rA   rB   rC   rD   rE   )rF   rZ   r[   r   rG   r\   rI   rH   r   r3   rJ   r   r   r   r   r   Úcheck_ipc_array¦   s$    

zTestIpcMemory.check_ipc_arrayc              	   C   sP   |   ¡ D ]B\}}| j||d |  ||¡ W d   ƒ q1 s@0    Y  qd S r`   )rQ   rb   re   rc   r   r   r   Útest_ipc_array¿   s    zTestIpcMemory.test_ipc_array)NF)NF)	Ú__name__Ú
__module__Ú__qualname__rK   rQ   r_   rd   re   rf   r   r   r   r   r,   Q   s   
&
r,   c                    s   ‡ ‡fdd„}t ||ƒ d S )Nc                     s‚   t jˆ  d t j ¡ } ˆ | ¡}ˆjt tj¡j	 }tj
|tjd}t jj||ˆjd ˆ ¡  W d   ƒ n1 st0    Y  |S )Nr/   )r   )r   ÚgpusZdevicesr=   Zopen_stagedr   r   r   r   r   Zzerosr   Zdevice_to_hostr"   )Zthis_ctxZ	deviceptrZarrsizeZ	hostarray©Ú
device_numr   r   r   r   Æ   s    

ÿ&z(staged_ipc_handle_test.<locals>.the_workr   )r   rl   r   r   r   rk   r   Ústaged_ipc_handle_testÅ   s    rm   c                 C   sø   zÄt j| ¦ | ‚}| ¡ }z(|  W d   ƒ n1 s80    Y  W n8 ty| } z t|ƒdkrhtdƒ‚W Y d }~nd }~0 0 tdƒ‚W d   ƒ n1 sš0    Y  W d   ƒ n1 s¸0    Y  W n   d}t ¡ }Y n
0 d}|}| ||f¡ d S r%   )	r   rj   r   r&   r'   r(   r	   r
   r   )r)   rl   r   r   r   r*   r   r   r   r   r   Ústaged_ipc_array_testÕ   s"    "Hrn   c                   @   s   e Zd Zdd„ Zdd„ ZdS )ÚTestIpcStagedc                 C   s  t jdt jd}t |¡}t d¡}| ¡ }t ¡ }| 	|j
¡}t |¡}t |¡}|  |jd ¡ tjr~|  |jj|jj¡ n|  t|jƒt|jƒ¡ |  |j|j¡ tttjƒƒD ]X}	||	|f}
|jt|
d}| ¡  | ¡ \}}| d¡ |sþ|  |¡ q´t j  !||¡ q´d S r-   )"r   r5   r   r   r6   r<   r=   r>   r!   r7   r8   rS   rT   rU   rV   rW   r   r9   rX   r   r:   rY   r   ÚrangeÚlenrj   r?   rm   r@   rA   rE   rB   rC   rD   )rF   r   rG   Zmpctxr   rH   rI   r]   r^   rl   r3   rJ   r   r   r   r   r   Útest_stagedó   s,    





zTestIpcStaged.test_stagedc                 C   sœ   t ttjƒƒD ]ˆ}tj d¡}t |¡}| ¡ }t 	d¡}| 
¡ }|||f}|jt|d}| ¡  | ¡ \}	}
| d¡ |	sˆ|  |
¡ qtj ||
¡ qd S )Nr.   r0   r1   r4   )rp   rq   r   rj   r   Zrandomr6   r7   r<   r=   r>   r?   rn   r@   rA   rE   rB   rC   rD   )rF   rl   r   rG   rI   rH   r   r3   rJ   r   r   r   r   r   rf     s    



zTestIpcStaged.test_ipc_arrayN)rg   rh   ri   rr   rf   r   r   r   r   ro   ï   s   "ro   Ú__main__)Zmultiprocessingr<   rN   r	   rS   Znumpyr   Znumbar   Znumba.cuda.cudadrvr   Znumba.cuda.testingr   r   r   r   r   Zunittestr   r    r$   r+   r,   rm   rn   ro   rg   Úmainr   r   r   r   Ú<module>   s0   q8