o
    8ήc=                     @   s   d dl Zd dlmZmZ d dlmZ d dlmZ d dlm	Z	m
Z
mZ d dlmZmZ d dlmZmZ d dlmZmZ ed	G d
d de
ZedkrPe	  dS dS )    N)	vectorizeguvectorize)cuda)driver)unittestContextResettingTestCaseForeignArray)skip_on_cudasimskip_if_external_memmgr)
linux_onlyoverride_config)callpatchz6CUDA Array Interface is not supported in the simulatorc                   @   s   e Zd Zdd Zdd Zdd Zeddd	 Zd
d Zdd Z	dd Z
dd Zdd Zdd Zdd Zdd Zdd Zdd Zdd Zed d! Zd"d# Zed$d% Zd&d' Zd(d) Zd*d+ Zd,d- Zd.d/ Zd0d1 Zd2d3 Zd4d5 Zd6d7 Zd8S )9TestCudaArrayInterfacec                 C   s&   t jr| t|jt|j d S d S N)r   USE_NV_BINDINGassertEqualintdevice_ctypes_pointer)selfab r   X/tmp/pip-target-vg8gfxp4/lib/python/numba/cuda/tests/cudapy/test_cuda_array_interface.pyassertPointersEqual   s
   z*TestCudaArrayInterface.assertPointersEqualc                 C   s   t d}| t| t|}| t| t|}| t| t|}| t| t j	
| | t j	
| | | || d S )N
   )nparangeassertFalser   is_cuda_array	to_device
assertTruer   as_cuda_arraytestingassert_array_equalcopy_to_hostr   )r   h_arrd_arrmy_arrwrappedr   r   r   test_as_cuda_array   s   


z)TestCudaArrayInterface.test_as_cuda_arrayc                 C   s   t jrt|jS |jjS r   )r   r   r   handlevalue)r   streamr   r   r   get_stream_value"   s   
z'TestCudaArrayInterface.get_stream_valuez+Ownership not relevant with external memmgrc                 C   s   t  }|jj}|  | t|d t t	d}t 
|}~| t|d tj| t	d ~| t|d |  d S )Nr   d      )r   current_contextmemory_managerdeallocationsclearr   lenr    r   r   r"   r#   assert_equalr%   )r   ctxdeallocsr'   cvtedr   r   r   test_ownership(   s   
z%TestCudaArrayInterface.test_ownershipc                 C   sz   t d}t|}t|}t|}tjdd }d}||j|| t j	
| ||  t j	
| ||  d S )Nr   c                 S   s.   t d}|t| krd S | |  |7  < d S )Nr0   )r   gridr5   )arrvalir   r   r   mutateG   s   
z6TestCudaArrayInterface.test_kernel_arg.<locals>.mutate   )r   r   r   r    r   r"   jitforallsizer#   r$   r%   )r   r&   r'   r(   r)   r?   r=   r   r   r   test_kernel_argA   s   



z&TestCudaArrayInterface.test_kernel_argc                 C   s   t dgdddd }tjd}tt|}d}|||}tj| ||  tt	|j
}||||d}tj| ||  d S )	Nz
f8(f8, f8)r   targetc                 S   s   | | S r   r   )r   r   r   r   r   vaddU   s   z3TestCudaArrayInterface.test_ufunc_arg.<locals>.vaddr      out)r   r   randomr   r   r    r#   r$   r%   device_arrayshaper   rG   r&   r<   r=   rJ   returnedr   r   r   test_ufunc_argT   s   

z%TestCudaArrayInterface.test_ufunc_argc                 C   s   t dgddddd }tjd}tt|}td}|||}tj|	 ||  tt
|j}||||d	}tj|	 ||  | ||j d S )
Nz(f8, f8, f8[:])z	(),()->()r   rE   c                 S   s   | | |d< d S )Nr   r   )inpr=   rJ   r   r   r   rG   f   s   z4TestCudaArrayInterface.test_gufunc_arg.<locals>.vaddr   r@   rI   )r   r   rK   r   r   r    float64r#   r$   r%   rL   rM   r   _arrrN   r   r   r   test_gufunc_arge   s   


z&TestCudaArrayInterface.test_gufunc_argc              	   C   sR  t jd}t|}t|}t j| | t j|dd  | t j|dd  |dd  t j|ddd  |ddd  t|ddd }t j| |ddd  | |ddd j	|j	 | |ddd j
|j
 | |ddd jj|jj | |ddd j|j | |ddd j|j|jj  t j|dd< t j| t t dt j|dd f |dd |dd< t j| t |dd |dd f t||dd< t j| | t j|ddd< t j| ddd t dt j t j| ddd |ddd  dS )zmViews created via array interface support:
            - Strided slices
            - Strided slices
        r   N      r0   )r   rK   r   r    r"   r#   r$   r%   r   rM   stridesdtypeitemsize
alloc_sizenbytesrC   piconcatenatefull)r   r&   c_arrr<   arr_stridedr   r   r   test_array_viewsw   sJ   

"&z'TestCudaArrayInterface.test_array_viewsc                 C   s   t jd}t|}dd }|j}|j}|d d d j}|d d d j}| |||||| | |d |d  | |d |d  d S )Nr   c                 S   s   |d d | d d  S )Ndatar   r   )origslicedr   r   r   base_offset   s   zGTestCudaArrayInterface.test_negative_strided_issue.<locals>.base_offsetrM   rW   )r   rK   r   r    __array_interface____cuda_array_interface__r   )r   r&   r_   re   h_aic_aih_ai_slicedc_ai_slicedr   r   r   test_negative_strided_issue   s   
z2TestCudaArrayInterface.test_negative_strided_issuec                 C   sp   t jd}t|}|d d d }| t}|  W d    n1 s&w   Y  d}| |t|j	 d S )Nr   rf   z.D->H copy not implemented for negative strides)
r   rK   r   r    assertRaisesNotImplementedErrorr%   assertInstr	exception)r   r&   r_   rd   raisesexpected_msgr   r   r   "test_negative_strided_copy_to_host   s   

z9TestCudaArrayInterface.test_negative_strided_copy_to_hostc                 C   s   t jd}t jjdddd}t|}t|}|j }||d< | t}t	| W d    n1 s7w   Y  d}| 
|t|j d S )Nr   rV   bool)rC   rX   maskzMasked arrays are not supported)r   rK   randintr   r    rh   copyrn   ro   from_cuda_array_interfacerp   rq   rr   )r   r&   h_maskr_   c_maskmasked_cuda_array_interfacers   rt   r   r   r   test_masked_array   s   


z(TestCudaArrayInterface.test_masked_arrayc                 C   sF   t d}| |jd d d t jdd }t|}|d | d S )Nr   rb   c                 S   s4   t d}| jd }||k r| |  d7  < d S d S )Nr0   r   )r   r;   rM   )r<   xNr   r   r   add_one   s
   

z<TestCudaArrayInterface.test_zero_size_array.<locals>.add_one)r0   r   )r   rL   r   rh   rA   r   )r   r_   r   r'   r   r   r   test_zero_size_array   s   

z+TestCudaArrayInterface.test_zero_size_arrayc                 C   sH   t d}| |jd d  |d d dd d f }| |jd d  d S )N)rV         rW   r0   )r   rL   r   rh   assertNotEqual)r   r_   r   r   r   test_strides   s   
z#TestCudaArrayInterface.test_stridesc                 C   s   t ddd}t|}|j}| |d  t| }t j	
|| | |jd  |j|d< | |d  t| }t j	
|| | |jd  d S )Nr   rV   rU   rW   C_CONTIGUOUS)r   r   reshaper   r    rh   assertIsNonerz   r%   r#   r$   r!   flagsrW   assertIsNotNone)r   	hostarraydevarrayfacegotr   r   r   test_consuming_strides  s   

z-TestCudaArrayInterface.test_consuming_stridesc                 C   s8   t d}| |jd  t d}| |jd  d S Nr   r-   )r   rL   r   rh   mapped_array)r   r_   
mapped_arrr   r   r   test_produce_no_stream  s   

z-TestCudaArrayInterface.test_produce_no_streamc                 C   s   t d}| |jd  d S r   )r   managed_arrayr   rh   )r   managed_arrr   r   r   test_produce_managed_no_stream  s   
z5TestCudaArrayInterface.test_produce_managed_no_streamc                 C   sp   t  }t jd|d}|jd }| |}| || t  }t jd|d}|jd }| |}| || d S Nr   r-   r-   )r   r-   rL   rh   r.   r   r   )r   sr_   
cai_streamstream_valuer   r   r   r   test_produce_stream  s   



z*TestCudaArrayInterface.test_produce_streamc                 C   s:   t  }t jd|d}|jd }| |}| || d S r   )r   r-   r   rh   r.   r   )r   r   r   r   r   r   r   r   test_produce_managed_stream*  s
   

z2TestCudaArrayInterface.test_produce_managed_streamc                 C   s*   t td}t|}| |jd d S )Nr   r   )r   r   rL   r"   r   r-   )r   f_arrr_   r   r   r   test_consume_no_stream2  s   
z-TestCudaArrayInterface.test_consume_no_streamc                 C   sX   t  }tt jd|d}t |}| |jj | |}| |j}| || d S )Nr   r   )	r   r-   r   rL   r"   r!   externalr.   r   )r   r   r   r_   r   imported_stream_valuer   r   r   test_consume_stream:  s   

z*TestCudaArrayInterface.test_consume_streamc                 C   sZ   t td}tjtjjjdd d}t| W d    n1 s"w   Y  |	  d S )Nr   synchronizereturn_value)
r   r   rL   r   objectcudadrvr   Streamr"   assert_not_called)r   r   	mock_syncr   r   r   test_consume_no_syncF  s   z+TestCudaArrayInterface.test_consume_no_syncc                 C   sf   t  }tt jd|d}tjt jjjdd d}t 	| W d    n1 s(w   Y  |
  d S )Nr   r   r   r   )r   r-   r   rL   r   r   r   r   r   r"   assert_called_once_withr   r   r   r   r   r   r   test_consume_syncQ  s   z(TestCudaArrayInterface.test_consume_syncc              	   C   s   t  }tt jd|d}tdd- tjt jjj	dd d}t 
| W d    n1 s.w   Y  |  W d    d S 1 sBw   Y  d S )Nr   r   CUDA_ARRAY_INTERFACE_SYNCFr   r   )r   r-   r   rL   r   r   r   r   r   r   r"   r   r   r   r   r   test_consume_sync_disabled]  s   
"z1TestCudaArrayInterface.test_consume_sync_disabledc                 C   sj   t td}tjdd }tjtjjjdd d}|d | W d    n1 s*w   Y  |	  d S )Nr   c                 S      d S r   r   r   r   r   r   fq     z5TestCudaArrayInterface.test_launch_no_sync.<locals>.fr   r   r0   r0   )
r   r   rL   rA   r   r   r   r   r   r   )r   r   r   r   r   r   r   test_launch_no_syncm  s   
z*TestCudaArrayInterface.test_launch_no_syncc                 C   sv   t  }tt jd|d}t jdd }tjt jjj	dd d}|d | W d    n1 s0w   Y  |
  d S )Nr   r   c                 S   r   r   r   r   r   r   r   r     r   z2TestCudaArrayInterface.test_launch_sync.<locals>.fr   r   r   )r   r-   r   rL   rA   r   r   r   r   r   r   )r   r   r   r   r   r   r   r   test_launch_sync|  s   
z'TestCudaArrayInterface.test_launch_syncc                 C   s   t  }t  }tt jd|d}tt jd|d}t jdd }tjt jjj	dd d}|d || W d    n1 s>w   Y  |
t t g d S )Nr   r   c                 S   r   r   r   r   yr   r   r   r     r   z>TestCudaArrayInterface.test_launch_sync_two_streams.<locals>.fr   r   r   )r   r-   r   rL   rA   r   r   r   r   r   assert_has_callsr   r   s1s2f_arr1f_arr2r   r   r   r   r   test_launch_sync_two_streams  s   
z3TestCudaArrayInterface.test_launch_sync_two_streamsc              	   C   s   t  }t  }tt jd|d}tt jd|d}tdd6 t jdd }tjt jj	j
dd d}|d	 || W d    n1 sDw   Y  |  W d    d S 1 sXw   Y  d S )
Nr   r   r   Fc                 S   r   r   r   r   r   r   r   r     r   z;TestCudaArrayInterface.test_launch_sync_disabled.<locals>.fr   r   r   )r   r-   r   rL   r   rA   r   r   r   r   r   r   r   r   r   r   test_launch_sync_disabled  s   

"z0TestCudaArrayInterface.test_launch_sync_disabledN) __name__
__module____qualname__r   r*   r.   r
   r:   rD   rP   rT   ra   rm   ru   r~   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r      s>    
B



r   __main__)numpyr   numbar   r   r   numba.cuda.cudadrvr   numba.cuda.testingr   r   r   r	   r
   numba.tests.supportr   r   unittest.mockr   r   r   r   mainr   r   r   r   <module>   s        (