
    jp"                        d dl Z d dlZ	 d dlZg dZ e j                  de      Z
 e j                  d      Zd dlmZ dZddgZd Zd Ze j                  j%                  d	 e ee            e
      d        ZddZe j                  j%                  d	 e ee            e
      e j                  j%                  dee
      e j                  j%                  dg d      d                      Ze j                  j%                  d	 e ee            e
      e j                  j%                  dee
      d               Ze j                  j%                  d	 e ee            e
      e j                  j%                  dee
      d               Ze j                  j%                  d	 e ee            e
      e j                  j%                  dee
      d               Ze j                  j%                  d	 e ee            e
      e j                  j%                  dee
      d               Zy# e$ r e j                  j                  ZY Ew xY w)    N)uint8int16float32zpyarrow.cuda)exc_typez
numba.cuda)DeviceNDArrayc                 
   t         j                  j                  d       t        j	                         }|j                         }t        j                         }t        j                  j                  |      }||f||fg| _	        y )Ni  )
nprandomseedcudaContextto_numbanb_cudacurrent_context
from_numbacontext_choices)modulectx1nb_ctx1nb_ctx2ctx2s        L/tmp/pip-target-jj7clzyk/lib/python/pyarrow/tests/test_cuda_numba_interop.pysetup_moduler   $   s_    IINN4<<>DmmoG%%'G<<""7+D#Wog?F    c                     | ` y N)r   )r   s    r   teardown_moduler   -   s    r   c)idsc                    t         |    \  }}|j                  |j                  j                  k(  sJ |j                  |j                         j                  j                  k(  sJ t        j
                  j                  |      }|j                  |j                  k(  sJ d}|j                  |      }|j                  |j                  j                  k(  sJ y )N
   )	r   handlevaluer   r   r   r   
new_buffercontext)r   ctxnb_ctxr   sizebufs         r   test_contextr*   1   s     "!$KC::,,,,,::..44444<<""6*D::$$$D
..
C::+++++r   c                    t        j                  |      }|dk(  rx| dk\  sJ t        j                  | |j                  z        }t        j
                  ||      }t         j                  j                  dd| t         j                        |dd ||fS |dk(  rQt        | d|      \  }}|j                  | |j                  z        }|j                  |d|j                  	       ||fS t        d
      )z5Return a host or device buffer with random data.
    hostr   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(   r3   r.   r&   r)   arrdbufs          r   r;   r;   >   s     HHUOEqyy  enn!45mmCu-""qs)+ # 3ACx	8	%d6GS~~dU^^34C!CHH=Dy
+
,,r   r.   r(   )r         i  c                    t         |    \  }}t        |d||      \  }}t        j                  |      }|j	                  |      }|j
                  |j
                  k(  sJ t        j                  |j                         |      }	t        j                  j                  ||	       |dk\  rt        |dz  d d       t        |dz  |dz   d       fD ]^  }
|j	                  ||
         }t        j                  |j                         |      }	t        j                  j                  ||
   |	       ` 	 t        j                  t        d      5  |j	                  |d d d          d d d        |dz  }||z  }||z  |k(  sJ |j	                  |j                  ||            }|j
                  |j
                  k(  sJ t        j                  |j                         |      }	t        j                  j                  ||	       t        j                  t        d      5  |j	                  |j                  ||      d d d d df          d d d        d}|dz  }|||z  z  }||z  |z  |k(  sJ |j	                  |j                  |||            }|j
                  |j
                  k(  sJ t        j                  |j                         |      }	t        j                  j                  ||	       t        j                  t        d      5  |j	                  |j                  |||      d d d          d d d         G d	 d
      }|j	                   ||            }|j
                  |j
                  k(  sJ t        j                  |j                         |      }	t        j                  j                  ||	       y # 1 sw Y   ZxY w# 1 sw Y   xY w# 1 sw Y   xY w)Nr2   r3   r.   r&   r-   rA      zarray data is non-contiguous)match   c                   "    e Zd Zd Zed        Zy)test_from_object.<locals>.MyObjc                     || _         y r   )darr)selfrJ   s     r   __init__z(test_from_object.<locals>.MyObj.__init__   s	    DIr   c                 .    | j                   j                  S r   )rJ   __cuda_array_interface__)rK   s    r   rN   z8test_from_object.<locals>.MyObj.__cuda_array_interface__   s    99555r   N)__name__
__module____qualname__rL   propertyrN    r   r   MyObjrH      s    	 
	6 
	6r   rT   )r   r;   r   	to_devicebuffer_from_objectr(   r	   r9   copy_to_hosttestingassert_equalslicepytestraisesr=   reshape)r   r.   r(   r&   r'   r>   cbufrJ   cbuf2arr2srdarrs1s2s3rT   s                   r   test_from_objectrf   Q   sm   
 "!$KC"43OIC S!D""4(E::"""==++-U;DJJC& qyat,a47T24 	2A**473E==!3!3!5UCDJJ##CFD1		2  ]]:"@C 	.""4!9-	.
 1W2XBw$&&t||B';<zzTYY&&&}}U//1?


T*]]:"@C 	A""4<<B#73Q3#?@	A
 1WBrE]Bw|t###&&t||BB'?@zzTYY&&&}}U//1?


T*]]:"@C 	B""4<<B#;CaC#@A	B6 6 ""5;/E::"""==++-U;DJJC&Y	. 	.	A 	A	B 	Bs$   N<,O		)O<O	OOc                 x   t         |    \  }}t        j                  |      }d}|j                  ||j                  z        }t        |f|j                  f||      }d|d d d|dd  t        j                  j                  |j                         d d d       t        j                  j                  |j                         dd  d       t        j                  j                  |      }t        j                  |j                         |      }t        j                  j                  ||j                                y )Nr!   gpu_datac      X   r-   )r   r	   r.   memallocr8   r   rX   rY   rW   r   
CudaBufferr   r9   )	r   r.   r&   r'   r(   memrJ   r^   r`   s	            r   test_numba_memallocrp      s     "!$KCHHUOE D
//$/
0C$5>>"3USIDD!HDHJJD--/3R8JJD--/3R8 ??%%c*D==**,E:DJJD$"3"3"56r   c                    t         |    \  }}d}t        |d||      \  }}|j                         }t        |j                  |j
                  |j                  |      }t        j                  j                  |j                         |       y )Nr!   r2   rC   rh   )r   r;   r   r   shapestridesr.   r	   rX   rY   rW   )	r   r.   r&   r'   r(   r>   r^   ro   rJ   s	            r   test_pyarrow_memallocrt      sq     "!$KCD"43OIC --/CCKKSIDJJD--/5r   c                 f   t         |    \  }}d}t        j                  d   5  t        |d||      \  }}|j                  j
                  |j
                  j                  k(  sJ |j                         }t        |j                  |j                  |j                  |      }t        j                  j                  |j                         |       d|d<   |j                  j!                          t        j"                  |j                         |      }	|	d   dk(  sJ 	 d d d        y # 1 sw Y   y xY w)Nr!   r   r2   rC   rh   rj   r-   )r   r   gpusr;   r%   r"   r#   r   r   rr   rs   r.   r	   rX   rY   rW   synchronizer9   )
r   r.   r&   r'   r(   r>   r^   ro   rJ   r`   s
             r   test_numba_contextrx      s     "!$KCD	a 
&tH-2=	T||""fmm&9&9999mmoSYYSYYM


 1 1 3S9Q  "}}T..0>Aw"}}
 
 
s   C;D''D0c                 
   t         |    \  }}t        j                  d        }d}t        |d||      \  }}d}|j                  |dz
  z   |z  }	|j                         }
t        |j                  |j                  |j                  |
      } ||	|f   |       |j                  j                          t        j                  |j                         |j                        }t        j                  j!                  ||dz          y )	Nc                 h    t         j                  d      }|| j                  k  r| |xx   dz  cc<   y y )Nr@   )r   gridr(   )an_arrayposs     r   increment_by_onez*test_pyarrow_jit.<locals>.increment_by_one   s.    ll1oSMQM r   r!   r2   rC       r@   rh   r-   )r   r   jitr;   r(   r   r   rr   rs   r.   r%   rw   r	   r9   rW   rX   rY   )r   r.   r&   r'   r~   r(   r>   r^   threadsperblockblockspergridro   rJ   arr1s                r   test_pyarrow_jitr      s     "!$KC[[  D"43OICOXX1!45/IM
--/CCKKSID4]O34T:LL==**,CII>DJJD#'*r   )r,   r   N)r[   pyarrowr6   numpyr	   ImportErrormark
pytestmarkdtypesimportorskipr   r   numba.cuda.cudadrv.devicearrayr   r   context_choice_idsr   r   parametrizerangelenr*   r;   rf   rp   rt   rx   r   rS   r   r   <module>r      s  $  # 
'v>K@
&

l
+ 8 $l3 @ eC(:$;</  1,1,-& eC(:$;</  1&f51H' 2 61H'V eC(:$;</  1&f57 617( eC(:$;</  1&f56 616 eC(:$;</  1&f5 61  eC(:$;</  1&f5+ 61+K  #""J#s   H5 5II