
    xKg(              	          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 Zd Zd Zd	 Ze ed
       ed       ed       G d de                                  Zd Zd Ze ed
       ed       ed       G d de                                  Ze ed       G d de                    Zedk(  r ej>                          yy)    N)cuda)driver)skip_on_armskip_on_cudasimskip_under_cuda_memcheckContextResettingTestCaseForeignArray)
linux_onlywindows_onlyc                 |    	  |        }d}|}|j                  ||f       y #  d}t        j                         }Y .xY w)NTF)	traceback
format_excput)the_workresult_queuearrsuccouts        d/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_ipc.pycore_ipc_handle_testr      sF    
j dC[!%""$s   ! ;c                 ,      fd}t        ||       y )Nc                      t        j                  t         j                        } t        j                  | j
                  z  |       5 }|j                         cd d d        S # 1 sw Y   y xY wN)shapedtype)npr   intpr   open_ipc_arrayitemsizecopy_to_host)r   darrhandlesizes     r   r   z&base_ipc_handle_test.<locals>.the_work!   sR    !  tu~~/E',.15$$&. . .s   
A$$A-r   )r"   r#   r   r   s   ``  r   base_ipc_handle_testr%       s    ' <0    c                 (      fd}t        ||       y )Nc                     t        j                  t         j                        } j                  t	        j
                         j                  | j                  z  |       }|j                         }j                          |S r   )
r   r   r   
open_arrayr   current_contextr#   r   r    close)r   r!   r   r"   s      r   r   z+serialize_ipc_handle_test.<locals>.the_work,   sf    !  !5!5!7'-{{enn'D', ! . !
r&   r$   )r"   r   r   s   `  r   serialize_ipc_handle_testr,   +   s     <0r&   c                 h   	 | 5 }|j                         }	 | 5  	 d d d        t        d      # 1 sw Y   xY w# t        $ r#}t        |      dk7  rt        d      Y d }~nd }~ww xY w	 d d d        n# 1 sw Y   nxY wd}}n#  d}t	        j
                         }Y nxY w|j                  ||f       y Nzdid not raise on reopenzIpcHandle is already openedzinvalid exception messageTF)r    AssertionError
ValueErrorstrr   r   r   )ipcarrr   r!   r   er   r   s          r   ipc_array_testr4   9   s    t##%C@  %%>?? V Fq6::()DEE ;F: VV$ %""$
 dC[!sX   B A29-9A26	9	A%A A2 A%%A2)	B 2A;7B BzHangs cuda-memcheckzIpc not available in CUDASIMz&CUDA IPC not supported on ARM in Numbac                   4    e Zd Zd Zd ZddZd ZddZd Zy)	TestIpcMemoryc                    t        j                  dt         j                        }t        j                  |      }t        j
                         }|j                  |j                        }t        j                  r|j                  j                  }nt        |j                        }|j                  }t        j                  d      }|j!                         }|||f}|j#                  t$        |      }	|	j'                          |j)                         \  }
}|
s| j+                  |       n t         j,                  j/                  ||       |	j1                  d       y 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   r>   procr   r   s               r   test_ipc_handlezTestIpcMemory.test_ipc_handleX   s    ii"''*$ ""$!!&//2   ;;//L -Lyy nnW%yy{dL1{{"6T{B

 $$&	cIIcNJJ##C-		!r&   c                 z    d t        dd       t        dd      t        d d      f}d}t        j                  ||      S )Nr?      )FT)slice	itertoolsproduct)rQ   indicesforeignss      r   variantszTestIpcMemory.variantsu   s<    q$q!eD!nE !  (33r&   Nc                 r   t        j                  dt         j                        }t        j                  |      }|||   }|rt        j
                  t        |            }|j                         }t        j                         }|j                  |j                        }t        j                  |      }t        j                  |      }	| j                  |	j                  d        | j!                  |	j"                  |j"                         t$        j&                  r;| j!                  |	j(                  j*                  |j(                  j*                         n8| j!                  t-        |	j(                        t-        |j(                               t/        j0                  d      }|j3                         }
||
f}|j5                  t6        |      }|j9                          |
j;                         \  }}|s| j=                  |       n t         j>                  jA                  ||       |jC                  d       y r8   )"r   r@   r   r   rA   as_cuda_arrayr	   r    r*   rB   rC   pickledumpsloadsassertIsbaseassertEqualr#   r   rD   r"   rE   tuplerG   rH   rI   rJ   r,   rK   rL   rM   rN   rO   rP   )rQ   	index_argforeignr   rR   expectrS   rT   buf
ipch_reconr   r>   rV   r   r   s                  r   check_ipc_handle_serializationz,TestIpcMemory.check_ipc_handle_serialization}   s   ii"''*$ I&F''V(<=F$$& ""$!!&//2 ll4 \\#&
joot,$))4  Z..779M9MNU:#4#45uT[[7IJ nnW%yy{l#{{";${G

 $$&	cIIcNJJ##FC0		!r&   c                     | j                         D ]3  \  }}| j                  ||      5  | j                  ||       d d d        5 y # 1 sw Y   @xY wN)indexrj   )r_   subTestrn   rQ   rq   rj   s      r   test_ipc_handle_serializationz+TestIpcMemory.test_ipc_handle_serialization   sE    #}}OE7E7;33E7C <;  /;;   AA	c                 V   t        j                  dt         j                        }t        j                  |      }|||   }|rt        j
                  t        |            }|j                         }|j                         }t        j                  d      }|j                         }||f}	|j                  t        |	      }
|
j                          |j                         \  }}|s| j!                  |       n t         j"                  j%                  ||       |
j'                  d       y r8   )r   r@   r   r   rA   ra   r	   r    rB   rG   rH   rI   rJ   r4   rK   rL   rM   rN   rO   rP   )rQ   ri   rj   r   rR   rk   rT   rS   r   r>   rV   r   r   s                r   check_ipc_arrayzTestIpcMemory.check_ipc_array   s    ii"''*$ I&F''V(<=F$$&$$& nnW%yy{l#{{.t{<

 $$&	cIIcNJJ##FC0		!r&   c                     | j                         D ]3  \  }}| j                  ||      5  | j                  ||       d d d        5 y # 1 sw Y   @xY wrp   )r_   rr   rw   rs   s      r   test_ipc_arrayzTestIpcMemory.test_ipc_array   sE    #}}OE7E7;$$UG4 <;  /;;ru   )NF)	__name__
__module____qualname__rW   r_   rn   rt   rw   ry    r&   r   r6   r6   R   s"    :4$LD
25r&   r6   c                 ,      fd}t        ||       y )Nc                     t         j                     5  t         j                  j                         } j	                  |       }j
                  t        j                  t        j                        j                  z  }t        j                  |t        j                        }t         j                  j                  ||j
                         j                          d d d        |S # 1 sw Y   S xY w)Nr:   )r#   )r   gpusdevicesrH   open_stagedr#   r   r   r   r   zerosr   device_to_hostr+   )this_ctx	deviceptrarrsize	hostarray
device_numr"   s       r   r   z(staged_ipc_handle_test.<locals>.the_work   s    YYz"||//1H**84IkkRXXbgg%6%?%??G8IKK&&96;; '  LLN #  # s   CC**C4r$   )r"   r   r   r   s   ``  r   staged_ipc_handle_testr      s    
 <0r&   c                    	 t         j                  |   5  | 5 }|j                         }	 | 5  	 d d d        t        d      # 1 sw Y   xY w# t        $ r#}t        |      dk7  rt        d      Y d }~nd }~ww xY w	 d d d        n# 1 sw Y   nxY wd d d        n# 1 sw Y   nxY wd}}n#  d}t        j                         }Y nxY w|j                  ||f       y r.   )	r   r   r    r/   r0   r1   r   r   r   )r2   r   r   r!   r   r3   r   r   s           r   staged_ipc_array_testr      s    YYz"4'')D   ))BCC  ! J1v!>>,-HII ?J>  #""& %""$
 dC[!sw   B, BBAAABA
A	A9A4/B4A99B=	BB	B	B, B$ B, ,Cc                       e Zd Zd Zd Zy)TestIpcStagedc                 T   t        j                  dt         j                        }t        j                  |      }t        j                  d      }|j                         }t        j                         }|j                  |j                        }t        j                  |      }t        j                  |      }| j                  |j                  d        t         j"                  r;| j%                  |j&                  j(                  |j&                  j(                         n8| j%                  t+        |j&                        t+        |j&                               | j%                  |j,                  |j,                         t/        t1        t        j2                              D ]  }	||	|f}
|j5                  t6        |
      }|j9                          |j;                         \  }}|j=                  d       |s| j?                  |       gt         j@                  jC                  ||        y r8   )"r   r@   r   r   rA   rG   rH   rI   r*   rB   rC   rb   rc   rd   re   rf   r   rD   rg   r"   rE   rh   r#   rangelenr   rJ   r   rK   rL   rP   rM   rN   rO   )rQ   r   rR   mpctxr   rS   rT   rl   rm   r   r>   rV   r   r   s                 r   test_stagedzTestIpcStaged.test_staged   s|   ii"''*$ w'{{} ""$!!&//2ll4 \\#&
joot,  Z..779M9MNU:#4#45uT[[7IJ$))4  DII/J*l3D==(>T=JDJJL$((*ID#IIaL		#

''S1 0r&   c                 ,   t        t        t        j                              D ]  }t        j
                  j                  d      }t        j                  |      }|j                         }t        j                  d      }|j                         }|||f}|j                  t        |      }|j                          |j                         \  }	}
|j                  d       |	s| j!                  |
       t        j"                  j%                  ||
        y )Nr9   r;   r<   r?   )r   r   r   r   r   randomrA   rB   rG   rH   rI   rJ   r   rK   rL   rP   rM   rN   rO   )rQ   r   r   rR   rT   rS   r   r>   rV   r   r   s              r   ry   zTestIpcStaged.test_ipc_array  s    DII/J))""2&C^^C(F((*D ..)C99;L*l3D;;&;$;GDJJL$((*ID#IIaL		#

''S1# 0r&   N)rz   r{   r|   r   ry   r}   r&   r   r   r      s    
 2D2r&   r   c                       e Zd Zd Zy)TestIpcNotSupportedc                 :   t        j                  dt         j                        }t        j                  |      }| j                  t              5 }|j                          d d d        t        j                        }| j                  d|       y # 1 sw Y   1xY w)Nr9   r:   zOS does not support CUDA IPC)r   r@   r   r   rA   assertRaisesOSErrorrB   r1   	exceptionassertIn)rQ   r   rR   raiseserrmsgs        r   test_unsupportedz$TestIpcNotSupported.test_unsupported0  sm    ii"''*$w'6!!# (V%%&4f= ('s   BBN)rz   r{   r|   r   r}   r&   r   r   r   -  s    >r&   r   __main__) multiprocessingrG   r[   r   rb   numpyr   numbar   numba.cuda.cudadrvr   numba.cuda.testingr   r   r   r   r	   numba.tests.supportr
   r   unittestr   r%   r,   r4   r6   r   r   r   r   rz   mainr}   r&   r   <module>r      s!         %H H 9 "11"2 /0/056n5, n5 7 1 1 n5b1 "4 /0/05652, 52 7 1 1 52p /0>2 > 1 > zHMMO r&   