§
    ƒÛg)  ã                   óÜ   — 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        ¦   «          dS dS )é    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dS )ÚTestCudaArrayc                 ó~  — t          j        d¦  «        }t          j        |¦  «        }|                     ¦   «         }|                      |j        |j        ¦  «         |                      |j        |j        ¦  «         |                      |j        |j        ¦  «         |                      |j        |j        ¦  «         d S )Nr   )ÚnpÚaranger   Ú	to_deviceÚcopy_to_hostÚassertEqualÚshapeÚsize)ÚselfÚxÚdxÚhxs       úX/granary/edev/quantum/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_array.pyÚtest_gpu_array_zero_lengthz(TestCudaArray.test_gpu_array_zero_length   s›   € ÝŒI�a‰LŒLˆÝŒ^˜AÑÔˆØ�_Š_ÑÔˆØ×Ò˜œ "¤(Ñ+Ô+Ð+Ø×Ò˜œ ¤Ñ)Ô)Ð)Ø×Ò˜œ "¤(Ñ+Ô+Ð+Ø×Ò˜œ ¤Ñ)Ô)Ð)Ð)Ð)ó    c                 óê   — d}t          j        d¦  «        j        }t          j        t	          j        d¦  «        ¦  «        j        }|                      ||¦  «         |                      ||¦  «         d S )N© )r   Údevice_arrayr   Údevice_array_liker   Úndarrayr   )r   Ú
null_shapeÚshape1Úshape2s       r   Útest_null_shapezTestCudaArray.test_null_shape   sg   € Øˆ
ÝÔ" 2Ñ&Ô&Ô,ˆÝÔ'­¬
°2©¬Ñ7Ô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         |¦  «         |                      t          j        |t          t          d¦  «        ¦  «        ¦  «        ¦  «         d S )Nzvoid(double[:])c                 ó^   — t          j        d¦  «        }|| j        d         k     r|| |<   d S d S ©Né   r   ©r   Úgridr   )r   Úis     r   Úkernelz4TestCudaArray.test_gpu_array_strided.<locals>.kernel"   s2   € å”	˜!‘”ˆAØ�1”7˜1”:Š~ˆ~Ø��!‘��ð ˆ~r   é
   ©ÚdtypeéP   )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»   € å	ŒÐ#Ñ	$Ô	$ð	ð 	ñ 
%Ô	$ð	õ
 ŒI�b¥¤	Ð*Ñ*Ô*ˆÝŒJ˜V¨AµR´WÐ=Ñ=Ô=ˆÝŒJ�q  1 R 4¤µ´	Ð:Ñ:Ô:ˆØˆˆvŒ�qÑÔÐØ�Š�œ A¥t­E°!©H¬H¡~¤~Ñ6Ô6Ñ7Ô7Ð7Ð7Ð7r   c                 ó  — t          j        d¦  «        d„ ¦   «         }t          j        dt          j        ¬¦  «        }|d dd…         }	 t           j                             |¦  «         t          d¦  «        ‚# t          $ r Y d S w xY w)Nzvoid(double[:], double[:])c                 óh   — t          j        d¦  «        }|| j        d         k     r|| |<   |||<   d S d S r%   r'   )r   r:   r)   s      r   Ú
copykernelz<TestCudaArray.test_gpu_array_interleaved.<locals>.copykernel0   s;   € å”	˜!‘”ˆAØ�1”7˜1”:Š~ˆ~Ø��!‘Ø��!‘��ð ˆ~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Ô	/ð	õ ŒI�b¥¤	Ð*Ñ*Ô*ˆØˆeˆr�!ˆeŒHˆð	BÝÔ×(Ò(¨Ñ+Ô+Ð+õ !ð "Añ Bô Bð Bøõ ð 	ð 	ð 	ØˆDˆDð	øøøs   Á
A8 Á8
BÂBc                 óè   — t           j                             d¦  «        \  }}|                      t	          j        |                     ¦   «         t	          j        d¦  «        k    ¦  «        ¦  «         d S )NrA   )r   rB   rC   r6   r   Úallr   Úarray)r   ÚdÚ_s      r   Útest_auto_device_constz$TestCudaArray.test_auto_device_constI   sU   € ÝÔ×+Ò+¨AÑ.Ô.‰ˆˆ1Ø�Š�œ˜qŸ~š~Ñ/Ô/µ2´8¸A±;´;Ò>Ñ?Ô?Ñ@Ô@Ð@Ð@Ð@r   c                 óŒ  —  ||¦  «        }|                       |j        |j        ¦  «         |                       |j        |j        ¦  «         |                       |j        |j        ¦  «         |                       |j        d         |j        d         ¦  «         |                       |j        d         |j        d         ¦  «         dS )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Æ   € ð
 �Y˜uÑ%Ô%ˆ
Ø×Ò˜œ jÔ&6Ñ7Ô7Ð7Ø×Ò˜œ¨
Ô(:Ñ;Ô;Ð;Ø×Ò˜œ jÔ&6Ñ7Ô7Ð7Ø×Ò˜œ ^Ô4Ø#Ô)¨.Ô9ñ	;ô 	;ð 	;à×Ò˜œ ^Ô4Ø#Ô)¨.Ô9ñ	;ô 	;ð 	;ð 	;ð 	;r   c                 óÎ   — t          j        dd¬¦  «        }t          D ]F}|                      |¬¦  «        5  |                      ||¦  «         d d d ¦  «         n# 1 swxY w Y   ŒGd S )Nr+   ÚC©Úorder©rR   ©r   r   ÚARRAY_LIKE_FUNCTIONSÚsubTestrT   ©r   Úd_arR   s      r   Útest_array_like_1dz TestCudaArray.test_array_like_1d[   ó°   € ÝÔ ¨#Ð.Ñ.Ô.ˆÝ-ð 	;ð 	;ˆIØ—’¨	�Ñ2Ô2ð ;ð ;Ø×*Ò*¨9°cÑ:Ô:Ð:ð;ð ;ð ;ñ ;ô ;ð ;ð ;ð ;ð ;ð ;ð ;øøøð ;ð ;ð ;ð ;øð	;ð 	;ó   ¶AÁA	Á A	c                 óÎ   — t          j        dd¬¦  «        }t          D ]F}|                      |¬¦  «        5  |                      ||¦  «         d d d ¦  «         n# 1 swxY w Y   ŒGd S ©N©r+   é   rV   rW   rY   rZ   r]   s      r   Útest_array_like_2dz TestCudaArray.test_array_like_2da   ó°   € ÝÔ °Ð4Ñ4Ô4ˆÝ-ð 	;ð 	;ˆIØ—’¨	�Ñ2Ô2ð ;ð ;Ø×*Ò*¨9°cÑ:Ô:Ð:ð;ð ;ð ;ñ ;ô ;ð ;ð ;ð ;ð ;ð ;ð ;øøøð ;ð ;ð ;ð ;øð	;ð 	;ra   c                 óÎ   — t          j        dd¬¦  «        }t          D ]F}|                      |¬¦  «        5  |                      ||¦  «         d d d ¦  «         n# 1 swxY w Y   ŒGd S 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 ]F}|                      |¬¦  «        5  |                      ||¦  «         d d d ¦  «         n# 1 swxY w Y   ŒGd S )N©r+   re   é   rV   rW   rY   rZ   r]   s      r   Útest_array_like_3dz TestCudaArray.test_array_like_3dm   ó°   € ÝÔ °CÐ8Ñ8Ô8ˆÝ-ð 	;ð 	;ˆIØ—’¨	�Ñ2Ô2ð ;ð ;Ø×*Ò*¨9°cÑ:Ô:Ð:ð;ð ;ð ;ñ ;ô ;ð ;ð ;ð ;ð ;ð ;ð ;øøøð ;ð ;ð ;ð ;øð	;ð 	;ra   c                 óÎ   — t          j        dd¬¦  «        }t          D ]F}|                      |¬¦  «        5  |                      ||¦  «         d d d ¦  «         n# 1 swxY w Y   ŒGd S )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 ]F}|                      |¬¦  «        5  |                      ||¦  «         d d d ¦  «         n# 1 swxY w Y   ŒGd S ©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 ]F}|                      |¬¦  «        5  |                      ||¦  «         d d d ¦  «         n# 1 swxY w Y   ŒGd S 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 ]F}|                      |¬¦  «        5  |                      ||¦  «         d d d ¦  «         n# 1 swxY w Y   ŒGd S )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        ¦  «         t          j        |¦  «        }|                       |j        |j        ¦  «         |                       |j        d         |j        d         ¦  «         |                       |j        d         |j        d         ¦  «         dS )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Ô5Ð5Ø×Ò˜œ w¤}Ñ5Ô5Ð5õ ”- Ñ%Ô%ˆØ×Ò˜œ¨'¬/Ñ:Ô:Ð:Ø×Ò˜œ ~Ô6Ø œ ~Ô6ñ	8ô 	8ð 	8à×Ò˜œ ~Ô6Ø œ ~Ô6ñ	8ô 	8ð 	8ð 	8ð 	8r   c                 ó  — d}t          j        |¦  «        d d d…         }t          j        |¦  «        d d d…         }t          D ]G}|                      |¬¦  «        5  |                      |||¦  «         d d d ¦  «         n# 1 swxY w Y   ŒHd S )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œ   sê   € ØˆÝŒx˜‰Œ˜s˜s ˜sÔ#ˆÝÔ" 5Ñ)Ô)¨#¨#¨A¨#Ô.ˆÝ-ð 	Dð 	DˆIØ—’¨	�Ñ2Ô2ð Dð DØ×*Ò*¨9°d¸FÑCÔCÐCðDð Dð Dñ Dô Dð Dð Dð Dð Dð Dð Døøøð Dð Dð Dð Døð	Dð 	Ds   ÁB Â B	ÂB	c                 ó$  — d}t          j        |d¬¦  «        d d d…         }t          j        |d¬¦  «        d d d…         }t          D ]G}|                      |¬¦  «        5  |                      |||¦  «         d d d ¦  «         n# 1 swxY w Y   ŒHd S )Nr+   rp   rW   rA   rY   r�   rƒ   s        r   Útest_array_like_1d_view_fz'TestCudaArray.test_array_like_1d_view_f¤   sö   € ØˆÝŒx˜ SÐ)Ñ)Ô)¨#¨#¨A¨#Ô.ˆÝÔ" 5°Ð4Ñ4Ô4°S°S°q°SÔ9ˆÝ-ð 	Dð 	DˆIØ—’¨	�Ñ2Ô2ð Dð DØ×*Ò*¨9°d¸FÑCÔCÐCðDð Dð Dñ Dô Dð Dð Dð Dð Dð Dð Døøøð Dð Dð Dð Døð	Dð 	Ds   Á BÂB	ÂB	c                 ó0  — d}t          j        |¦  «        d d d…d d d…f         }t          j        |¦  «        d d d…d d d…f         }t          D ]G}|                      |¬¦  «        5  |                      |||¦  «         d d d ¦  «         n# 1 swxY w Y   ŒHd S )Nrd   rA   rY   r�   rƒ   s        r   Útest_array_like_2d_viewz%TestCudaArray.test_array_like_2d_view¬   sþ   € ØˆÝŒx˜‰Œ˜s˜s ˜s C C a C˜xÔ(ˆÝÔ" 5Ñ)Ô)¨#¨#¨A¨#¨s¨s°¨s¨(Ô3ˆÝ-ð 	Dð 	DˆIØ—’¨	�Ñ2Ô2ð Dð DØ×*Ò*¨9°d¸FÑCÔCÐCðDð Dð Dñ Dô Dð Dð Dð Dð Dð Dð Døøøð Dð Dð Dð Døð	Dð 	Ds   Á&B
Â
B	ÂB	c                 ó8  — d}t          j        |d¬¦  «        d d d…d d d…f         }t          j        |d¬¦  «        d d d…d d d…f         }t          D ]G}|                      |¬¦  «        5  |                      |||¦  «         d d d ¦  «         n# 1 swxY w Y   ŒHd S ©Nrd   rp   rW   rA   rY   r�   rƒ   s        r   Útest_array_like_2d_view_fz'TestCudaArray.test_array_like_2d_view_f´   s
  € ØˆÝŒx˜ SÐ)Ñ)Ô)¨#¨#¨A¨#¨s¨s°¨s¨(Ô3ˆÝÔ" 5°Ð4Ñ4Ô4°S°S°q°S¸#¸#¸A¸#°XÔ>ˆÝ-ð 	Dð 	DˆIØ—’¨	�Ñ2Ô2ð Dð DØ×*Ò*¨9°d¸FÑCÔCÐCðDð Dð Dñ Dô Dð Dð Dð Dð Dð Dð Døøøð Dð Dð Dð Døð	Dð 	Ds   Á*BÂB	ÂB	z5Numba and NumPy stride semantics differ for transposec                 ó  — d}t          j        |¦  «        d d d…d d d…f         j        }t          D ]Ö}|                      |¬¦  «        5   ||¦  «        }|                      |j        |j        ¦  «         |                      |j        |j        ¦  «         |                      d|j        ¦  «         |  	                    |j
        d         ¦  «         |                      |j
        d         ¦  «         d d d ¦  «         n# 1 swxY w Y   Œ×d S )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¼   sC  € àˆÝÔ" 5Ñ)Ô)¨#¨#¨A¨#¨s¨s°¨s¨(Ô3Ô5ˆÝ-ð 	=ð 	=ˆIØ—’¨	�Ñ2Ô2ð 
=ð 
=ð
 !�y Ñ(Ô(�Ø× Ò  ¤¨t¬zÑ:Ô:Ð:Ø× Ò  ¤¨t¬zÑ:Ô:Ð:Ø× Ò  ¨$¬,Ñ7Ô7Ð7Ø—’ ¤
¨>Ô :Ñ;Ô;Ð;Ø× Ò  ¤¨NÔ!;Ñ<Ô<Ð<ð
=ð 
=ð 
=ñ 
=ô 
=ð 
=ð 
=ð 
=ð 
=ð 
=ð 
=øøøð 
=ð 
=ð 
=ð 
=øð	=ð 	=s   Á	B'C<Ã<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 �]}|                      |¬¦  «        5  t          j        |¦  «        } ||¦  «        }|                      |j	        |j	        ¦  «         |                      |j
        |j
        ¦  «         |                      |j        |j        ¦  «         |                      |j        d         |j        d         ¦  «         |                      |j        d         |j        d         ¦  «         d d d ¦  «         n# 1 swxY w Y   �Œ
d S )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¸  € ð ˆÝŒx˜‰Œ˜s˜s ˜s C C a C˜xÔ(Ô*ˆÝÔ" 5Ñ)Ô)¨#¨#¨A¨#¨s¨s°¨s¨(Ô3Ô5ˆÝ-ð 	@ñ 	@ˆIØ—’¨	�Ñ2Ô2ð @ð @õ œ-¨Ñ-Ô-�Ø#˜) FÑ+Ô+�Ø× Ò  ¤¨w¬}Ñ=Ô=Ð=Ø× Ò  ¤¨w¬}Ñ=Ô=Ð=Ø× Ò  ¤°'´/ÑBÔBÐBØ× Ò  ¤¨~Ô!>Ø!(¤¨~Ô!>ñ@ô @ð @à× Ò  ¤¨~Ô!>Ø!(¤¨~Ô!>ñ@ô @ð @ð@ð @ð @ñ @ô @ð @ð @ð @ð @ð @ð @øøøð @ð @ð @ð @ùð	@ð 	@s   Á1CEÅE	ÅE	c                 óL  — 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 ]G}|                      |¬¦  «        5  |                      |||¦  «         d d d ¦  «         n# 1 swxY w Y   ŒHd S 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  € ØˆÝŒx˜ SÐ)Ñ)Ô)¨#¨#¨A¨#¨s¨s°¨s¨(Ô3Ô5ˆÝÔ" 5°Ð4Ñ4Ô4°S°S°q°S¸#¸#¸A¸#°XÔ>Ô@ˆÝ-ð 	Dð 	DˆIØ—’¨	�Ñ2Ô2ð Dð DØ×*Ò*¨9°d¸FÑCÔCÐCðDð Dð Dñ Dô Dð Dð Dð Dð Dð Dð Døøøð Dð Dð Dð Døð	Dð 	Ds   Á4BÂB	ÂB	z-Kernel overloads not created in the simulatorc                 óH  — t           j        d„ ¦   «         }d}t          j        |f¦  «        }t          j        |¦  «        }t          j        |f¦  «        } |d         ||¦  «          |d         ||¦  «         |                      dt          |j        ¦  «        ¦  «         d S )Nc                 óJ   — t          j        d¦  «        }| |         dz  ||<   d S )Nr&   rA   )r   r(   )ÚAÚoutr)   s      r   Úfuncz+TestCudaArray.test_issue_4628.<locals>.funcó   s$   € å”	˜!‘”ˆAØ�q”T˜A‘XˆC�‰FˆFˆ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¥   € õ 
Œð	ð 	ñ 
Œð	ð ˆÝŒG�Q�D‰MŒMˆÝŒn˜QÑÔˆÝ”˜1˜$‘”ˆàˆˆVŒ�Q˜ÑÔÐØˆˆVŒ�S˜&Ñ!Ô!Ð!à×Ò˜�C ¤Ñ/Ô/Ñ0Ô0Ð0Ð0Ð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ð 8ð 8ðBð Bð Bð6Að Að Að;ð ;ð ;ð;ð ;ð ;ð;ð ;ð ;ð;ð ;ð ;ð;ð ;ð ;ð;ð ;ð ;ð;ð ;ð ;ð;ð ;ð ;ð;ð ;ð ;ð8ð 8ð 8ð"Dð Dð DðDð Dð DðDð Dð DðDð Dð Dð €_ÐLÑMÔMð=ð =ñ NÔMð=ð  Ðð %ñ &ô &ð@ð @ñ&ô &ð@ð&Dð Dð Dð €_ÐDÑEÔEð1ð 1ñ FÔEð1ð 1ð 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Ð 5Ð 5Ð 5Ð 5Ð 5Ð 5Ð 5Ø CÐ CÐ CÐ CÐ CÐ CÐ CÐ CØ Ð Ð Ð Ð Ð Ð Ð ð 
Ôð 4Ø Ô2°DÔ4JÐKÐÐà Ô2°DÔ4JØ Ô2ð4Ððq1ð q1ð q1ð q1ð q1�Lñ q1ô q1ð q1ðh ˆzÒÐØ€H„M�O„O€O€O€Oð Ðr   