U
    k/e"                  	   @   s  d dl Z d dlZd dlZdddg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d dddg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uint8Zint16Z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_cudaZcurrent_context
from_numbacontext_choices)moduleZctx1Znb_ctx1Znb_ctx2ctx2 r   I/tmp/pip-unpacked-wheel-seu8352k/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t|j| jjks4ttj|}|j|jksPtd}||}|j|j	jksptd S )N
   )
r   handlevalueAssertionErrorr	   r   r   r   
new_buffercontext)r   ctxnb_ctxr   sizebufr   r   r   test_context.   s    
r    hostc                 C   s   t |}|dkrb| dkstt| |j }t j||d}t jjdd| t j	d|dd< ||fS |dkrt
| d|d\}}|| |j }|j|d|jd	 ||fS td
dS )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#   r   paZallocate_bufferitemsize
frombufferr   randintr   make_random_bufferr   Zcopy_from_hostr   
ValueError)r   r(   r#   r   r   arrdbufr   r   r   r/   ;   s    
r/   r#   r         i  c              	   C   s  t |  \}}t|d||d\}}t|}||}|j|jksDttj|	 |d}	tj
||	 |dkrJt|d d d t|d |d  d fD ]6}
|||
 }tj|	 |d}	tj
||
 |	 qtjtdd ||d d d  W 5 Q R X |d }|| }|| |kst||||}|j|jks@ttj|	 |d}	tj
||	 tjtdd* ||||d d d d df  W 5 Q R X d}|d }|||  }|| | |kst|||||}|j|jksttj|	 |d}	tj
||	 tjtdd$ |||||d d d  W 5 Q R X G d	d
 d
}|||}|j|jksxttj|	 |d}	tj
||	 d S )Nr'   r(   r#   r   r"   r4      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)selfr9   r   r   r   __init__   s    z(test_from_object.<locals>.MyObj.__init__c                 S   s   | j jS r   )r9   __cuda_array_interface__)r:   r   r   r   r<      s    z8test_from_object.<locals>.MyObj.__cuda_array_interface__N)__name__
__module____qualname__r;   propertyr<   r   r   r   r   MyObj   s   rA   )r   r/   r
   Z	to_deviceZbuffer_from_objectr   r   r   r-   copy_to_hosttestingassert_equalslicepytestZraisesr0   Zreshape)r   r#   r   r   r   r1   cbufr9   Zcbuf2arr2ss1s2Zs3rA   r   r   r   test_from_objectN   s\    



.(rL   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   Zgpu_datac      X   r"   )r   r   r#   Zmemallocr,   r   rC   rD   rB   r   Z
CudaBufferr   r-   )	r   r#   r   r   r   memr9   rG   rH   r   r   r   test_numba_memalloc   s    
rR   c           	      C   sX   t |  \}}d}t|d||d\}}| }t|j|j|j|d}tj	|
 | d S )Nr   r'   r5   rM   )r   r/   r	   r   shapestridesr#   r   rC   rD   rB   )	r   r#   r   r   r   r1   rG   rQ   r9   r   r   r   test_pyarrow_memalloc   s    rU   c           
   	   C   s   t |  \}}d}tjd  t|d||d\}}|jj|jjksDt| }t	|j
|j|j|d}tj| | d|d< |j  tj| |d}	|	d dkstW 5 Q R X d S )Nr   r   r'   r5   rM   rN   r"   )r   r
   Zgpusr/   r   r   r   r   r	   r   rS   rT   r#   r   rC   rD   rB   synchronizer-   )
r   r#   r   r   r   r1   rG   rQ   r9   rH   r   r   r   test_numba_context   s     

rW   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 )Nr3   )r
   Zgridr   )Zan_arrayposr   r   r   increment_by_one   s    

z*test_pyarrow_jit.<locals>.increment_by_oner   r'   r5       r3   rM   r"   )r   r
   Zjitr/   r   r	   r   rS   rT   r#   r   rV   r   r-   rB   rC   rD   )r   r#   r   r   rY   r   r1   rG   ZthreadsperblockZblockspergridrQ   r9   Zarr1r   r   r   test_pyarrow_jit   s    

r[   )r!   r   N)rF   Zpyarrowr+   Znumpyr   ZdtypesZimportorskipr   r
   Znumba.cuda.cudadrv.devicearrayr   r   Zcontext_choice_idsr   r   markZparametrizerangelenr    r/   rL   rR   rU   rW   r[   r   r   r   r   <module>   sR   


	

K