
    xKg)                        d dl Zd dlmZmZ d dlmZmZ d dlmZm	Z	 ej                  re	j                  e	j                  fZn#e	j                  e	j                  e	j                  fZ G d de      Zedk(  r ej"                          yy)    N)unittestCUDATestCase)skip_on_cudasimskip_unless_cudasim)configcudac                       e Zd Zd Zd Zd Zd Zd Zd Zd Z	d Z
d	 Zd
 Zd Zd Zd Zd Zd Zd Zd Zd Zd Z ed      d        Z ed      d        Zd Z ed      d        Zy)TestCudaArrayc                    t        j                  d      }t        j                  |      }|j	                         }| j                  |j                  |j                         | j                  |j                  |j                         | j                  |j                  |j                         | j                  |j                  |j                         y )Nr   )nparanger   	to_devicecopy_to_hostassertEqualshapesize)selfxdxhxs       f/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_array.pytest_gpu_array_zero_lengthz(TestCudaArray.test_gpu_array_zero_length   s    IIaL^^A__"((+)"((+)    c                     d}t        j                  d      j                  }t        j                  t	        j
                  d            j                  }| j                  ||       | j                  ||       y )N )r   device_arrayr   device_array_liker   ndarrayr   )r   
null_shapeshape1shape2s       r   test_null_shapezTestCudaArray.test_null_shape   s[    
""2&,,''

27==,,r   c           
         t        j                  d      d        }t        j                  dt        j                        }t        j
                  d|t        j                        }t        j
                  d|dd	 t        j                  
      } |d   |       | j                  t        j                  |t        t        d                         y )Nzvoid(double[:])c                 ^    t        j                  d      }|| j                  d   k  r|| |<   y y N   r   r   gridr   )r   is     r   kernelz4TestCudaArray.test_gpu_array_strided.<locals>.kernel"   s,    		!A1771:~! r   
   dtypeP   )r   bufferr-   	      )r/   r-   )r+   r+   )r   jitr   r   doubler   byte
assertTrueallcloselistrange)r   r*   r   yzs        r   test_gpu_array_stridedz$TestCudaArray.test_gpu_array_strided    s    	#	$	 
%	
 IIb		*JJVARWW=JJq1R		:vqAtE!H~67r   c                    t        j                  d      d        }t        j                  dt        j                        }|d dd   }	 t         j
                  j                  |       t        d      # t        $ r Y y w xY w)Nzvoid(double[:], double[:])c                 h    t        j                  d      }|| j                  d   k  r|| |<   |||<   y y r%   r'   )r   r:   r)   s      r   
copykernelz<TestCudaArray.test_gpu_array_interleaved.<locals>.copykernel0   s5    		!A1771:~!! r   r+   r,      zDShould raise exception complaining the contiguous-ness of the array.)	r   r3   r   r   r4   devicearrayauto_deviceAssertionError
ValueError)r   r?   r   r:   s       r   test_gpu_array_interleavedz(TestCudaArray.test_gpu_array_interleaved.   s    	.	/	 
0	 IIb		*er!eH	B((+ ! "A B B  		s   
A4 4	B ?B c                     t         j                  j                  d      \  }}| j                  t	        j
                  |j                         t	        j                  d      k(               y )NrA   )r   rB   rC   r6   r   allr   array)r   d_s      r   test_auto_device_constz$TestCudaArray.test_auto_device_constI   sE    ++A.1q~~/288A;>?@r   c                     ||      }| j                  |j                  |j                         | j                  |j                  |j                         | j                  |j                  |j                         | j                  |j                  d   |j                  d          | j                  |j                  d   |j                  d          y)zk
        Tests of *_array_like where shape, strides, dtype, and flags should
        all be equal.
        C_CONTIGUOUSF_CONTIGUOUSN)r   r   stridesr-   flags)r   	like_funcrI   
array_likes       r   _test_array_like_samez#TestCudaArray._test_array_like_sameM   s    
 u%
j&6&67
(:(:;j&6&67^4#)).9	;^4#)).9	;r   c                     t        j                  dd      }t        D ]/  }| j                  |      5  | j	                  ||       d d d        1 y # 1 sw Y   <xY w)Nr+   CorderrR   r   r   ARRAY_LIKE_FUNCTIONSsubTestrT   r   d_arR   s      r   test_array_like_1dz TestCudaArray.test_array_like_1d[   K    #.-I	2**9c: 32 .22   AA	c                     t        j                  dd      }t        D ]/  }| j                  |      5  | j	                  ||       d d d        1 y # 1 sw Y   <xY wNr+      rV   rW   rY   rZ   r]   s      r   test_array_like_2dz TestCudaArray.test_array_like_2da   K    4-I	2**9c: 32 .22ra   c                     t        j                  dd      }t        D ]/  }| j                  |      5  | j	                  ||       d d d        1 y # 1 sw Y   <xY wrc   rZ   r]   s      r   test_array_like_2d_transposez*TestCudaArray.test_array_like_2d_transposeg   rg   ra   c                     t        j                  dd      }t        D ]/  }| j                  |      5  | j	                  ||       d d d        1 y # 1 sw Y   <xY w)Nr+   re      rV   rW   rY   rZ   r]   s      r   test_array_like_3dz TestCudaArray.test_array_like_3dm   K    C8-I	2**9c: 32 .22ra   c                     t        j                  dd      }t        D ]/  }| j                  |      5  | j	                  ||       d d d        1 y # 1 sw Y   <xY w)Nr+   FrW   rY   rZ   r]   s      r   test_array_like_1d_fz"TestCudaArray.test_array_like_1d_fs   r`   ra   c                     t        j                  dd      }t        D ]/  }| j                  |      5  | j	                  ||       d d d        1 y # 1 sw Y   <xY wNrd   rp   rW   rY   rZ   r]   s      r   test_array_like_2d_fz"TestCudaArray.test_array_like_2d_fy   rg   ra   c                     t        j                  dd      }t        D ]/  }| j                  |      5  | j	                  ||       d d d        1 y # 1 sw Y   <xY wrs   rZ   r]   s      r   test_array_like_2d_f_transposez,TestCudaArray.test_array_like_2d_f_transpose   rg   ra   c                     t        j                  dd      }t        D ]/  }| j                  |      5  | j	                  ||       d d d        1 y # 1 sw Y   <xY w)Nrk   rp   rW   rY   rZ   r]   s      r   test_array_like_3d_fz"TestCudaArray.test_array_like_3d_f   rn   ra   c                     ||      }| j                  |j                  |j                         | j                  |j                  |j                         t        j                  |      }| j                  |j
                  |j
                         | j                  |j                  d   |j                  d          | j                  |j                  d   |j                  d          y)z
        Tests of device_array_like where the original array is a view - the
        strides should not be equal because a contiguous array is expected.
        rN   rO   N)r   r   r-   r   
zeros_likerP   rQ   )r   rR   viewd_viewnb_likenp_likes         r   _test_array_like_viewz#TestCudaArray._test_array_like_view   s    
 F#w}}5w}}5 --%'//:~6 ~6	8~6 ~6	8r   c                     d}t        j                  |      d d d   }t        j                  |      d d d   }t        D ]0  }| j                  |      5  | j                  |||       d d d        2 y # 1 sw Y   =xY w)Nr+   rA   rY   r   zerosr   r   r[   r\   r   r   r   r{   r|   rR   s        r   test_array_like_1d_viewz%TestCudaArray.test_array_like_1d_view   so    xxss#""5)#A#.-I	2**9dFC 32 .22s   A33A<	c                    d}t        j                  |d      d d d   }t        j                  |d      d d d   }t        D ]0  }| j                  |      5  | j                  |||       d d d        2 y # 1 sw Y   =xY w)Nr+   rp   rW   rA   rY   r   r   s        r   test_array_like_1d_view_fz'TestCudaArray.test_array_like_1d_view_f   st    xxS)#A#.""54SqS9-I	2**9dFC 32 .22s   A77B 	c                    d}t        j                  |      d d dd d df   }t        j                  |      d d dd d df   }t        D ]0  }| j                  |      5  | j                  |||       d d d        2 y # 1 sw Y   =xY w)Nrd   rA   rY   r   r   s        r   test_array_like_2d_viewz%TestCudaArray.test_array_like_2d_view   s    xxssCaCx(""5)#A#ss(3-I	2**9dFC 32 .22s   A==B	c                    d}t        j                  |d      d d dd d df   }t        j                  |d      d d dd d df   }t        D ]0  }| j                  |      5  | j                  |||       d d d        2 y # 1 sw Y   =xY wNrd   rp   rW   rA   rY   r   r   s        r   test_array_like_2d_view_fz'TestCudaArray.test_array_like_2d_view_f   s    xxS)#A#ss(3""54SqS#A#X>-I	2**9dFC 32 .22s   "BB
	z5Numba and NumPy stride semantics differ for transposec                    d}t        j                  |      d d dd d df   j                  }t        D ]  }| j	                  |      5   ||      }| j                  |j                  |j                         | j                  |j                  |j                         | j                  d|j                         | j                  |j                  d          | j                  |j                  d          d d d         y # 1 sw Y   xY w)Nrd   rA   rY   )(      rN   rO   )r   r   Tr[   r\   r   r   r-   rP   r6   rQ   assertFalse)r   r   r|   rR   likes        r   (test_array_like_2d_view_transpose_devicez6TestCudaArray.test_array_like_2d_view_transpose_device   s    ""5)#A#ss(355-I	2
 !(  tzz:  tzz:  $,,7

> :;  N!;< 32 .22s   B-D  D		c                    d}t        j                  |      d d dd d df   j                  }t        j                  |      d d dd d df   j                  }t
        D ]  }| j                  |      5  t        j                  |      } ||      }| j                  |j                  |j                         | j                  |j                  |j                         | j                  |j                  |j                         | j                  |j                  d   |j                  d          | j                  |j                  d   |j                  d          d d d         y # 1 sw Y   xY w)Nrd   rA   rY   rN   rO   )r   r   r   r   r   r[   r\   rz   r   r   r-   rP   rQ   )r   r   r{   r|   rR   r~   r}   s          r   +test_array_like_2d_view_transpose_simulatorz9TestCudaArray.test_array_like_2d_view_transpose_simulator   s,    xxssCaCx(**""5)#A#ss(355--I	2 ---#F+  w}}=  w}}=  '//B  ~!>!(~!>@  ~!>!(~!>@ 32 .22s   3C(E''E1	c                 B   d}t        j                  |d      d d dd d df   j                  }t        j                  |d      d d dd d df   j                  }t
        D ]0  }| j                  |      5  | j                  |||       d d d        2 y # 1 sw Y   =xY wr   )r   r   r   r   r   r[   r\   r   r   s        r   #test_array_like_2d_view_f_transposez1TestCudaArray.test_array_like_2d_view_f_transpose   s    xxS)#A#ss(355""54SqS#A#X>@@-I	2**9dFC 32 .22s   6BB	z-Kernel overloads not created in the simulatorc                 0   t         j                  d        }d}t        j                  |f      }t        j                  |      }t        j
                  |f      } |d   ||        |d   ||       | j                  dt        |j                               y )Nc                 D    t        j                  d      }| |   dz  ||<   y )Nr&   rA   )r   r(   )Aoutr)   s      r   funcz+TestCudaArray.test_issue_4628.<locals>.func   s     		!AqTAXCFr      )r&   r   r&   )	r   r3   r   onesr   r   r   len	overloads)r   r   nar^   results         r   test_issue_4628zTestCudaArray.test_issue_4628   s     
	 
	 GGQDMnnQ1$VQVS&!C/0r   N)__name__
__module____qualname__r   r"   r<   rF   rL   rT   r_   rf   ri   rm   rq   rt   rv   rx   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r
   r
      s    *-8B6A;;;;;;;;;8"DDDD LM= N=   % &@&@&D DE1 F1r   r
   __main__)numpyr   numba.cuda.testingr   r   r   r   numbar   r   ENABLE_CUDASIMr   pinned_array_liker[   mapped_array_liker
   r   mainr   r   r   <module>r      s~     5 C  
 22D4J4JK 22D4J4J 224q1L q1h zHMMO r   