o
    hZ"                     @   s  d dl Z d dlZzd dlZW n ey   e jjZY nw g dZe 	dZ
e 	dZd dlmZ dZddgZdd Zdd	 Ze jjd
eeeeddd ZdddZe jjd
eeeede jjdeede jdg ddd Ze jjd
eeeede jjdeeddd Ze jjd
eeeede jjdeeddd Ze jjd
eeeede jjdeeddd Ze jjd
eeeede jjdeeddd ZdS )     N)uint8int16float32zpyarrow.cudaz
numba.cuda)DeviceNDArrayc                 C   sF   t jd t }| }t }tj|}||f||fg| _	d S )Ni  )
nprandomseedcudaContextto_numbanb_cudacurrent_context
from_numbacontext_choices)modulectx1nb_ctx1nb_ctx2ctx2 r   Y/var/www/vscode/kcb/lib/python3.10/site-packages/pyarrow/tests/test_cuda_numba_interop.pysetup_module$   s   r   c                 C   s   | ` d S N)r   )r   r   r   r   teardown_module-   s   r   c)idsc                 C   st   t |  \}}|j|jjksJ |j| jjksJ tj|}|j|jks(J d}||}|j|jjks8J d S )N
   )	r   handlevaluer   r	   r
   r   
new_buffercontext)r   ctxnb_ctxr   sizebufr   r   r   test_context1   s   
r%   hostr   c                 C   s   t |}|dkr1| dksJ t| |j }t j||d}t jjdd| t jd|dd< ||fS |dkrSt	| d|d\}}|
| |j }|j|d|jd	 ||fS td
)z5Return a host or device buffer with random data.
    r&   r   dtype   )lowhighr#   r(   Ndevice)targetr(   )positionnbyteszinvalid target value)r   r(   paallocate_bufferitemsize
frombufferr   randintr   make_random_bufferr   copy_from_hostr#   
ValueError)r#   r-   r(   r!   r$   arrdbufr   r   r   r5   >   s   
r5   r(   r#   )r         i  c                 C   s  t |  \}}t|d||d\}}t|}||}|j|jks"J tj| |d}	tj	
||	 |dkrBt|d d d t|d |d  d fD ]}
|||
 }tj| |d}	tj	
||
 |	 qJ	 tjtdd	 ||d d d
  W d    n1 sw   Y  |d }|| }|| |ksJ ||||}|j|jksJ tj| |d}	tj	
||	 tjtdd	 ||||d d d d d
f  W d    n1 sw   Y  d}|d }|||  }|| | |ksJ |||||}|j|jksJ tj| |d}	tj	
||	 tjtdd	 |||||d d d
  W d    n	1 s=w   Y  G dd d}|||}|j|jksYJ tj| |d}	tj	
||	 d S )Nr,   r-   r(   r!   r'   r;      r   zarray data is non-contiguous)match   c                   @   s    e Zd Zdd Zedd ZdS )ztest_from_object.<locals>.MyObjc                 S   s
   || _ d S r   )darr)selfrA   r   r   r   __init__   s   
z(test_from_object.<locals>.MyObj.__init__c                 S   s   | j jS r   )rA   __cuda_array_interface__)rB   r   r   r   rD      s   z8test_from_object.<locals>.MyObj.__cuda_array_interface__N)__name__
__module____qualname__rC   propertyrD   r   r   r   r   MyObj   s    rI   )r   r5   r   	to_devicebuffer_from_objectr#   r   r3   copy_to_hosttestingassert_equalslicepytestraisesr7   reshape)r   r(   r#   r!   r"   r8   cbufrA   cbuf2arr2srdarrs1s2s3rI   r   r   r   test_from_objectQ   sd   


& r[   c           	      C   s   t |  \}}t|}d}|||j }t|f|jf||d}d|d d< d|dd < tj| d d d tj| dd  d t	j
|}tj| |d}tj||  d S )Nr   gpu_datac      X   r'   )r   r   r(   memallocr2   r   rM   rN   rL   r	   
CudaBufferr   r3   )	r   r(   r!   r"   r#   memrA   rS   rU   r   r   r   test_numba_memalloc   s   
rd   c           	      C   sX   t |  \}}d}t|d||d\}}| }t|j|j|j|d}tj	|
 | d S )Nr   r,   r<   r\   )r   r5   r   r   shapestridesr(   r   rM   rN   rL   )	r   r(   r!   r"   r#   r8   rS   rc   rA   r   r   r   test_pyarrow_memalloc   s   rg   c           
      C   s   t |  \}}d}tjd O t|d||d\}}|jj|jjks"J | }t|j	|j
|j|d}tj| | d|d< |j  tj| |d}	|	d dksTJ W d    d S 1 s_w   Y  d S )Nr   r   r,   r<   r\   r^   r'   )r   r   gpusr5   r    r   r   r   r   re   rf   r(   r   rM   rN   rL   synchronizer3   )
r   r(   r!   r"   r#   r8   rS   rc   rA   rU   r   r   r   test_numba_context   s   

"rj   c                 C   s   t |  \}}tjdd }d}t|d||d\}}d}|j|d  | }	| }
t|j|j|j	|
d}||	|f | |j
  tj| |j	d	}tj||d  d S )
Nc                 S   s,   t d}|| jk r| |  d7  < d S d S )Nr:   )r   gridr#   )an_arrayposr   r   r   increment_by_one   s   

z*test_pyarrow_jit.<locals>.increment_by_oner   r,   r<       r:   r\   r'   )r   r   jitr5   r#   r   r   re   rf   r(   r    ri   r   r3   rL   rM   rN   )r   r(   r!   r"   rn   r#   r8   rS   threadsperblockblockspergridrc   rA   arr1r   r   r   test_pyarrow_jit   s   

rt   )r&   r   N)rP   pyarrowr0   numpyr   ImportErrormark
pytestmarkdtypesimportorskipr	   r   numba.cuda.cudadrv.devicearrayr   r   context_choice_idsr   r   parametrizerangelenr%   r5   r[   rd   rg   rj   rt   r   r   r   r   <module>   s\   

	

K