o
    i(                     @   s0  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mZmZ d dlZdd Zdd	 Zd
d Zdd ZeedededG dd deZdd Zdd ZeedededG dd deZeedG dd deZedkre  dS dS )    N)cuda)driver)skip_on_armskip_on_cudasimskip_under_cuda_memcheckContextResettingTestCaseForeignArray)
linux_onlywindows_onlyc                 C   s<   z|  }W n   d}t  }Y nd}|}|||f d S )NFT)	traceback
format_excput)the_workresult_queuearrsuccout r   ^/home/ubuntu/transcripts/venv/lib/python3.10/site-packages/numba/cuda/tests/cudapy/test_ipc.pycore_ipc_handle_test   s   
r   c                        fdd}t || d S )Nc                     sN   t t j} tj | j | d}| W  d    S 1 s w   Y  d S N)shapedtype)npr   intpr   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   
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   zQ| E}|  }z|  W d    n1 sw   Y  W tdW td ty@ } zt|dkr6tdW Y d }~nd }~ww W d    n1 sKw   Y  W n   d}t }Y n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_test9   s2   r0   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}	|	  | \}
}|
sN| | 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ipchhandle_bytesr"   r   r8   procr   r   r   r   r   test_ipc_handleX   s$   




zTestIpcMemory.test_ipc_handlec                 C   s.   d t dd t ddt d df}d}t||S )Nr9      )FT)slice	itertoolsproduct)rK   indicesforeignsr   r   r   variantsu   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rS| |	jj|jj n| t|	jt|j td}| }
||
f}|jt|d}|  |
 \}}|s| | nt j || |!d d S r2   )"r   r:   r   r   r;   as_cuda_arrayr   r   r&   r<   r=   pickledumpsloadsassertIsbaseassertEqualr"   r   r>   r!   r?   tuplerA   rB   rC   rD   r)   rE   rF   rG   rH   rI   rJ   )rK   	index_argforeignr   rL   expectrM   rN   buf
ipch_reconr   r8   rP   r   r   r   r   r   check_ipc_handle_serialization}   s4   



z,TestIpcMemory.check_ipc_handle_serializationc              	   C   P   |   D ]!\}}| j||d | || W d    n1 s w   Y  qd S N)indexrb   )rX   subTestrf   rK   ri   rb   r   r   r   test_ipc_handle_serialization      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}
|
  | \}}|sL| | nt j|| |
d d S r2   )r   r:   r   r   r;   rY   r   r   r<   rA   rB   rC   rD   r0   rE   rF   rG   rH   rI   rJ   )rK   ra   rb   r   rL   rc   rN   rM   r   r8   rP   r   r   r   r   r   check_ipc_array   s$   

zTestIpcMemory.check_ipc_arrayc              	   C   rg   rh   )rX   rj   rn   rk   r   r   r   test_ipc_array   rm   zTestIpcMemory.test_ipc_array)NF)	__name__
__module____qualname__rQ   rX   rf   rl   rn   ro   r   r   r   r   r1   R   s    
&
r1   c                    r   )Nc                     s   t j  3 t j } | }jttjj	 }tj
|tjd}t jj||jd   W d    |S 1 s;w   Y  |S )Nr4   )r"   )r   gpusdevicesrB   open_stagedr"   r   r   r   r   zerosr   device_to_hostr'   )this_ctx	deviceptrarrsize	hostarray
device_numr!   r   r   r      s   



		z(staged_ipc_handle_test.<locals>.the_workr#   )r!   r}   r   r   r   r|   r   staged_ipc_handle_test   s   r~   c                 C   s   zft j| W | E}| }z|  W d    n1 sw   Y  W tdW td tyF } zt|dkr<tdW Y d }~nd }~ww W d    n1 sQw   Y  W d    n1 s`w   Y  W n   d}t }Y nd}|}|||f d S r*   )	r   rs   r   r+   r,   r-   r   r   r   )r.   r}   r   r   r   r/   r   r   r   r   r   staged_ipc_array_test   s8   r   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 ],}	||	|f}
|jt|
d}|  | \}}|d |s| | qZt j !|| qZd S r2   )"r   r:   r   r   r;   rA   rB   rC   r&   r<   r=   rZ   r[   r\   r]   r^   r   r>   r_   r!   r?   r`   r"   rangelenrs   rD   r~   rE   rF   rJ   rG   rH   rI   )rK   r   rL   mpctxr   rM   rN   rd   re   r}   r8   rP   r   r   r   r   r   test_staged   s.   





zTestIpcStaged.test_stagedc                 C   s   t ttjD ]D}tjd}t|}| }t	d}|
 }|||f}|jt|d}|  | \}	}
|d |	sD| |
 qtj||
 qd S )Nr3   r5   r6   r9   )r   r   r   rs   r   randomr;   r<   rA   rB   rC   rD   r   rE   rF   rJ   rG   rH   rI   )rK   r}   r   rL   rN   rM   r   r8   rP   r   r   r   r   r   ro     s   



zTestIpcStaged.test_ipc_arrayN)rp   rq   rr   r   ro   r   r   r   r   r      s    "r   c                   @   s   e Zd Zdd ZdS )TestIpcNotSupportedc                 C   sf   t jdt jd}t|}| t}|  W d    n1 s!w   Y  t|j	}| 
d| d S )Nr3   r4   zOS does not support CUDA IPC)r   r:   r   r   r;   assertRaisesOSErrorr<   r,   	exceptionassertIn)rK   r   rL   raiseserrmsgr   r   r   test_unsupported0  s   


z$TestIpcNotSupported.test_unsupportedN)rp   rq   rr   r   r   r   r   r   r   -  s    r   __main__) multiprocessingrA   rT   r   rZ   numpyr   numbar   numba.cuda.cudadrvr   numba.cuda.testingr   r   r   r   r   numba.tests.supportr	   r
   unittestr   r$   r)   r0   r1   r~   r   r   r   rp   mainr   r   r   r   <module>   s@    q8
