
    wgO                         d dl Z d dlZd dlmZ d dlmZ d dlmZm	Z	 d dlm
Z
  G d de	      Z G d d	e	      Z G d
 de	      Zedk(  r ej                          yy)    N)devicearray)cuda)unittestCUDATestCase)skip_on_cudasimc                      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 ed      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"d" Z# ed#      d$        Z$ ed#      d%        Z% ed#      d&        Z& ed#      d'        Z' ed#      d(        Z( ed#      d)        Z) ed#      d*        Z* ed#      d+        Z+ ed#      d,        Z,d- Z- ed       d.        Z.y/)0TestCudaNDArrayc                    t        j                  d      }t        j                  |       t	        j
                  d      }t        j                  |      }t        j                  |       t	        j                  d      }t        j                  |      }| j                  |j                  d       t        j                  |       y )Nd   )shapegX9v?r   )
r   device_arrayr   verify_cuda_ndarray_interfacenpempty	to_deviceasarrayassertEqualndim)selfdaryarys      o/home/mcse/projects/flask/flask-venv/lib/python3.12/site-packages/numba/cuda/tests/cudadrv/test_cuda_ndarray.pytest_device_array_interfacez+TestCudaNDArray.test_device_array_interface
   s      s+11$7hhsm~~c"11$7jj~~c"A&11$7    c                 D   t        j                  dt         j                        }d|j                  _        | j                  |j                  j                         t        j                  |      }|j                         }t         j                  j                  ||       y )Nr   dtypeF)r   arangefloat32flags	writeableassertFalser   r   copy_to_hosttestingassert_array_equal)r   r   r   retrs       r   test_device_array_from_readonlyz/TestCudaNDArray.test_device_array_from_readonly   sj    ii2::.#		,,-~~c"  "


%%dC0r   c                     t        j                  dd      }| j                  |j                  t	        j                  d             y )N)r   f4)r   r   )r   r   r   r   r   )r   r   s     r   test_devicearray_dtypez&TestCudaNDArray.test_devicearray_dtype!   s0      vT:RXXd^4r   c                 |    t        j                  dt         j                        }t        j                  |d       y )Nr   r   F)copy)r   r   r   r   r   )r   arrays     r   test_devicearray_no_copyz(TestCudaNDArray.test_devicearray_no_copy%   s$    		#RZZ0u5)r   c                    t        j                  d      j                  ddd      }t        j                  |      }| j                  |j                  |j                         | j                  |j                  dd  |j                  dd         y )N               )r   r   reshaper   r   r   r   r   r   r   s      r   test_devicearray_shapez&TestCudaNDArray.test_devicearray_shape)   si    ii	"**1a3~~c"DJJ/12

127r   c                    t        j                  dt         j                        }|j                         }t	        j
                  |      }d|d d  |j                  |       t         j                  j                  ||       y )Nr   r   r   	r   r   int32r,   r   r   r#   r$   r%   )r   r-   originalgpumems       r   test_devicearrayz TestCudaNDArray.test_devicearray/   sZ    		#RXX.::<&aE"


%%eX6r   c                 P   t        j                         }|j                         5  t        j                  dt        j
                  |      }| j                  |j                  |      j                  |       | j                  |j                  |       d d d        y # 1 sw Y   y xY w)N)r2   r2   )r   stream)r   r?   auto_synchronizer   r   float64r   bind)r   r?   arrs      r   test_stream_bindz TestCudaNDArray.test_stream_bind8   s    $$& 	1##jjC SXXf-44f=SZZ0	1 	1 	1s   A.BB%c                     t        j                  d      }t        j                  d      }| j	                  t        |      t        |             y )N)r2   r2   r   r   r   r   r   lenr6   s      r   test_len_1dzTestCudaNDArray.test_len_1dB   s7    hhtn  #S3t9-r   c                     t        j                  d      }t        j                  d      }| j	                  t        |      t        |             y )N)r2      rF   r6   s      r   test_len_2dzTestCudaNDArray.test_len_2dG   s8    hhv  (S3t9-r   c                     t        j                  d      }t        j                  d      }| j	                  t        |      t        |             y )N)r2   rJ      rF   r6   s      r   test_len_3dzTestCudaNDArray.test_len_3dL   s8    hhy!  +S3t9-r   c                    d}t        j                  |t         j                        }|j                         }t	        j
                  |      }|j                  |dz        \  }}d|d d  | j                  t        j                  |dk(               |j                  ||dz  d         |j                  |d |dz          | j                  t        j                  ||k(               y )Nr   r   r1   r   )
r   r   r:   r,   r   r   split
assertTrueallr#   )r   Nr-   r;   r<   leftrights          r   test_devicearray_partitionz*TestCudaNDArray.test_devicearray_partitionQ   s    		!288,::<&ll16*eauz*+5a>*%a.)u012r   c                 8   d}t        j                  |t         j                        }|j                         }t	        j
                  |      }t	        j
                  |dz  |       |j                  |       t         j                  j                  ||dz         y )Nr   r   r1   )tor9   )r   rS   r-   r;   r<   s        r   test_devicearray_replacez(TestCudaNDArray.test_devicearray_replacea   sl    		!288,::<&uqyV,E"


%%eX\:r   zThis works in the simulatorc                 d   t        j                  t        j                  t        j                  d            j                  ddd            }| j                  t              5 }t        j                  |       d d d        | j                  dt        j                               y # 1 sw Y   /xY w)N   r2   r3   r4   z2transposing a non-2D DeviceNDArray isn't supported)r   r   r   r-   r   r5   assertRaisesNotImplementedError	transposer   str	exceptionr   r<   es      r   #test_devicearray_transpose_wrongdimz3TestCudaNDArray.test_devicearray_transpose_wrongdimj   s    2 7 ? ?1a HI23 	!qLL 	! 	@		! 	!s   "B&&B/c                 4   t        j                  t        j                  d            j                  ddd      }t        j                  t        j                  |      d      j                         }| j                  t        j                  ||k(               y )Nr0   r2   r3   r1   )r   r4   r1   axes)
r   r-   r   r5   r^   r   r   r#   rQ   rR   r   r;   r-   s      r   #test_devicearray_transpose_identityz3TestCudaNDArray.test_devicearray_transpose_identityu   sg    88BIIbM*221a;T^^H5"+--9\^ 	u012r   c                 l   t        j                  t        j                  t        j                  d            j                  dd            }| j                  t              5 }t        j                  |d       d d d        | j                  t        j                        ddg       y # 1 sw Y   2xY w)	Nr[   r2   r3   )r   r   re   zinvalid axes list (0, 0)zrepeated axis in transpose	containerr   r   r   r-   r   r5   r\   
ValueErrorr^   assertInr_   r`   ra   s      r   )test_devicearray_transpose_duplicatedaxisz9TestCudaNDArray.test_devicearray_transpose_duplicatedaxis|   s    2 7 ? ?1 EFz* 	.aLLf-	. 	*, 	 		. 	.   !B**B3c                 l   t        j                  t        j                  t        j                  d            j                  dd            }| j                  t              5 }t        j                  |d       d d d        | j                  t        j                        g d       y # 1 sw Y   2xY w)Nr[   r2   r3   )r   r1   re   )zinvalid axes list (0, 2)zinvalid axis for this arrayz0axis 2 is out of bounds for array of dimension 2rj   rl   ra   s      r   $test_devicearray_transpose_wrongaxisz4TestCudaNDArray.test_devicearray_transpose_wrongaxis   s    2 7 ? ?1 EFz* 	.aLLf-	. 	 	 		. 	.rp   c                    t        j                  t        j                  d      d      j                  dd      }t	        j
                  |      }dD ]i  }| j                  |      5  t         j                  j                  |j                  |      j                         |j                  |             d d d        k y # 1 sw Y   vxY w)Nr[   i2r   r2   r3   )i4u4i8f8)r   r-   r   r5   r   r   subTestr$   r%   viewr#   )r   r;   r-   r   s       r   test_devicearray_view_okz(TestCudaNDArray.test_devicearray_view_ok   s    88BIIbM6>>q!Dx(- 	EE* 

--JJu%224MM%( 	 s   'AB??C	c                 f   t        j                  t        j                  d      d      j                  dd      }t	        j
                  |      d d d d df   }|d d d d df   }t         j                  j                  |j                  d      j                         |j                  d             y )N    rt   r   r3      r1   u2)
r   r-   r   r5   r   r   r$   r%   rz   r#   rg   s      r   %test_devicearray_view_ok_not_c_contigz5TestCudaNDArray.test_devicearray_view_ok_not_c_contig   s    88BIIbM6>>q!Dx(CaC0AssF#


%%JJt))+MM$	
r   c                    t        j                  t        j                  d      d      j                  dd      }t	        j
                  |      d d d d df   }| j                  t              5 }|j                  d       d d d        t        j                        }| j                  d|       d	|v }d
|v }| j                  |xs |d       y # 1 sw Y   OxY w)Nr}   rt   r   r3   r~   r1   ru   z)To change to a dtype of a different size,zthe array must be C-contiguousz the last axis must be contiguousz&Expected message to mention contiguity)r   r-   r   r5   r   r   r\   rm   rz   r_   r`   rn   rQ   )r   r;   r-   rb   msgcontiguous_pre_np123contiguous_post_np123s          r   &test_devicearray_view_bad_not_c_contigz6TestCudaNDArray.test_devicearray_view_bad_not_c_contig   s    88BIIbM6>>q!Dx(CaC0z* 	aJJt	 !++A3G?3F Bc I,E0E@	B	 	s   /CCc                 b   t        j                  t        j                  d      d      j                  dd      }t	        j
                  |      }| j                  t              5 }|j                  d       d d d        | j                  dt        j                               y # 1 sw Y   /xY w)Nr[   rt   r   r3   r2   ru   zuWhen changing to a larger dtype, its size must be a divisor of the total size in bytes of the last axis of the array.)r   r-   r   r5   r   r   r\   rm   rz   r   r_   r`   )r   r;   r-   rb   s       r   "test_devicearray_view_bad_itemsizez2TestCudaNDArray.test_devicearray_view_bad_itemsize   s    88BIIbM6>>q!Dx(z* 	aJJt	. 			 	s   %B%%B.c                 B   t        j                  t        j                  d            j                  dd      }t        j                  t        j                  |            j                         }| j                  t        j                  ||j                  k(               y Nr[   r2   r3   )r   r-   r   r5   r^   r   r   r#   rQ   rR   Trg   s      r   test_devicearray_transpose_okz-TestCudaNDArray.test_devicearray_transpose_ok   sc    88BIIbM*221a8T^^H56CCEu

234r   c                 0   t        j                  t        j                  d            j                  dd      }t	        j
                  |      j                  j                         }| j                  t        j                  ||j                  k(               y r   )
r   r-   r   r5   r   r   r   r#   rQ   rR   rg   s      r   test_devicearray_transpose_Tz,TestCudaNDArray.test_devicearray_transpose_T   s`    88BIIbM*221a8x(**779u

234r   c                    t        j                  d      j                  ddd      }t        j                  dd      }t	        j
                  |      }||d<   | j                  t              5 }|d   j                  |       d d d        | j                  t        j                  t        j                               y # 1 sw Y   =xY w)N   rJ   Forder)rJ   )
fill_valuer   r1   )r   r   r5   fullr   r   r\   rm   copy_to_devicer   r   errmsg_contiguous_bufferr_   r`   )r   asdrb   s        r   !test_devicearray_contiguous_slicez1TestCudaNDArray.test_devicearray_contiguous_slice   s     IIbM!!!Qc!2GGq-NN1! z* 	#aaD"	#00		# 	#s   .B??Cc                 Z   | j                  |j                  j                         | j                  |j                  j                         ||f||f||f||ffD ]  \  }}|j                  j                  rdndd|j                  j                  rdnd}t	        j
                  |      }|j                  |       | j                  t        j                  |j                         |k(        |       | j                  t        j                  |j                         |k(        |        y)z-
        Checks host->device memcpys
        Cr   z => )r   N)
rQ   r    c_contiguousf_contiguousr   r   r   r   rR   r#   )r   a_ca_fr;   r,   r   r   s          r   &_test_devicearray_contiguous_host_copyz6TestCudaNDArray._test_devicearray_contiguous_host_copy   s     			../		../ #J#J#J#J	
 	FNHd  ~~22;zz..C7C
 x(AT"OOBFF1>>#3s#:;OEOOBFF1>>#3s#:;OE	Fr   c                     t        j                  d      j                  ddd      }t        j                  |d      }| j	                  ||       y )N}   rJ   r   r   )r   r   r5   r-   r   r   r   r   s      r   (test_devicearray_contiguous_copy_host_3dz8TestCudaNDArray.test_devicearray_contiguous_copy_host_3d   s?    ii	"**1a3hhs#&33C=r   c                     t        j                  d      }t        j                  |d      }| j                  ||       y )NrJ   r   r   )r   r   r-   r   r   s      r   (test_devicearray_contiguous_copy_host_1dz8TestCudaNDArray.test_devicearray_contiguous_copy_host_1d   s/    iilhhs#&33C=r   c                 (   t        j                  d      j                  ddd      }t        j                  |d      }| j	                  |j
                  j                         | j	                  |j
                  j                         t        j                  |      }| j                  t              5 }|j                  t        j                  |             d d d        | j                  dj                  |j                  |j                        t!        j"                               |j                  t        j                  |             | j	                  t        j$                  |j'                         |k(               t        j                  |      }| j                  t              5 }|j                  t        j                  |             d d d        | j                  dj                  |j                  |j                        t!        |j"                               |j                  t        j                  |             | j	                  t        j$                  |j'                         |k(               y # 1 sw Y   xY w# 1 sw Y   xY w)Nr   rJ   r   r   zincompatible strides: {} vs. {})r   r   r5   r-   rQ   r    r   r   r   r   r\   rm   r   r   formatstridesr_   r`   rR   r#   )r   r   r   r   rb   s        r   'test_devicearray_contiguous_copy_devicez7TestCudaNDArray.test_devicearray_contiguous_copy_device  s   ii	"**1a3hhs#&		../		../NN3z* 	2aT^^C01	2-44S[[#++N	 	
,-q~~/3678NN3z* 	2aT^^C01	2-44S[[#++N	 	
,-q~~/3678%	2 	2	2 	2s   2%I;+%J;JJc                    d}d}t        j                  |      }t        j                  |      j                  |d      }t        j                  |      j                  |d      }t	        t        |            D ]i  }t        d       f|z  t         j                  fz   }|d | |fz   ||d  z   }t        j                  ||   |      }	t        j                  ||   |      }
t        j                  |	      }t        j                  |
      }t         j                  j                  |j                         |	       t         j                  j                  |j                         |
       |j                  |
       |j                  |	       t         j                  j                  |j                         |
       t         j                  j                  |j                         |	       l y )Nr3   )r1   r2   r   r   r   )r   prodr   r5   rangerG   slicenewaxisbroadcast_tor   r   r$   r%   r#   r   )r   	broadsize	coreshapecoresizecore_ccore_fdimnewindex
broadshapebroad_cbroad_fdbroad_cdbroad_fs                r   $test_devicearray_broadcast_host_copyz4TestCudaNDArray.test_devicearray_broadcast_host_copy  s~   		779%8$,,Yc,B8$,,Yc,BY( 	LCd~+rzzm;H"4CI<7)CD/IJoofX&6
CGoofX&6
CG~~g.H~~g.HJJ))(*?*?*A7KJJ))(*?*?*A7K##G,##G,JJ))(*?*?*A7KJJ))(*?*?*A7K	Lr   c                    t        j                  d      }t        j                  |      }t        j                  d      d d d   }|j	                  |       t         j
                  j                  |j                         |       y )N
      r1   )r   r   r   r   r   r$   r%   r#   )r   r   r   rC   s       r   (test_devicearray_contiguous_host_stridedz8TestCudaNDArray.test_devicearray_contiguous_host_strided3  s\    iimNN3iimCaC 	


%%ann&6<r   c                    t        j                  t        j                  d            }t        j                  d      }| j	                  t
              5 }|j                  t        j                  |      d d d          d d d        | j                  t        j                  t        j                               y # 1 sw Y   =xY w)Nr   r1   )r   r   r   r   r\   rm   r   r   r   r   r_   r`   )r   r   rC   rb   s       r   *test_devicearray_contiguous_device_stridedz:TestCudaNDArray.test_devicearray_contiguous_device_strided:  s    NN299R=)iimz* 	7aT^^C0156	700		7 	7s   +B::Cz,DeviceNDArray class not present in simulatorc                     t        j                  ddt        j                        }| j	                  |j
                  d          | j	                  |j
                  d          y )N)r4   r   )i   r~   C_CONTIGUOUSF_CONTIGUOUS)r   DeviceNDArrayr   rA   rQ   r    )r   rC   s     r    test_devicearray_relaxed_stridesz0TestCudaNDArray.test_devicearray_relaxed_stridesD  sH     ''2::F 			.12		.12r   c                 P   d}d}t        j                  ||      D ]  \  }}t        j                  ||      }t	        j
                  |      }| j                  |j                  d   |j                  d          | j                  |j                  d   |j                  d           y )N))r4   r3   )r3   r4   )r   r   r   r   r   )	itertoolsproductr   ndarrayr   r   r   r    )r   shapesordersr   r   rC   d_arrs          r   !test_c_f_contiguity_matches_numpyz1TestCudaNDArray.test_c_f_contiguity_matches_numpyS  s     "%--ff= 	:LE5**U%0CNN3'ESYY~6"[[8:SYY~6"[[8:	:r   z Typing not done in the simulatorc                     t        j                  dd      }t        j                  |      }| j	                  |j
                  j                  d       y Nr   r   r   r   zerosr   r   r   _numba_type_layoutr   r   r   s      r   &test_devicearray_typing_order_simple_cz6TestCudaNDArray.test_devicearray_typing_order_simple_ca  <     HHRs#NN1..4r   c                     t        j                  dd      }t        j                  |      }| j	                  |j
                  j                  d       y )Nr   r   r   r   r   r   s      r   &test_devicearray_typing_order_simple_fz6TestCudaNDArray.test_devicearray_typing_order_simple_fh  r   r   c                     t        j                  dd      }t        j                  |      }| j	                  |j
                  j                  d       y )Nr1   r   r   r   r   r   s      r   "test_devicearray_typing_order_2d_cz2TestCudaNDArray.test_devicearray_typing_order_2d_co  <     HHWC(NN1..4r   c                     t        j                  dd      }t        j                  |      }| j	                  |j
                  j                  d       y )Nr   r   r   r   r   s      r   "test_devicearray_typing_order_2d_fz2TestCudaNDArray.test_devicearray_typing_order_2d_fv  r   r   c                     t        j                  dd      }t        j                  |      d d df   }| j	                  |j
                  j                  d       y )NrJ   rJ   r   r   r1   Ar   r   s      r   /test_devicearray_typing_order_noncontig_slice_cz?TestCudaNDArray.test_devicearray_typing_order_noncontig_slice_c}  sE     HHV3'NN1ac"..4r   c                     t        j                  dd      }t        j                  |      dd d f   }| j	                  |j
                  j                  d       y )Nr   r   r   r1   r   r   r   s      r   /test_devicearray_typing_order_noncontig_slice_fz?TestCudaNDArray.test_devicearray_typing_order_noncontig_slice_f  E     HHV3'NN1ac"..4r   c                     t        j                  dd      }t        j                  |      dd d f   }| j	                  |j
                  j                  d       y )Nr   r   r   r1   r   r   s      r   ,test_devicearray_typing_order_contig_slice_cz<TestCudaNDArray.test_devicearray_typing_order_contig_slice_c  r   r   c                     t        j                  dd      }t        j                  |      d d df   }| j	                  |j
                  j                  d       y )Nr   r   r   r1   r   r   r   s      r   ,test_devicearray_typing_order_contig_slice_fz<TestCudaNDArray.test_devicearray_typing_order_contig_slice_f  sE     HHV3'NN1ac"..4r   c                     t        j                  t        j                  dg      d      }t        j                  |      }| j                  |j                  j                  d       y )Nr4   )r   r   )r   r   r-   r   r   r   r   r   r   s      r   )test_devicearray_typing_order_broadcastedz9TestCudaNDArray.test_devicearray_typing_order_broadcasted  sF     OOBHHaSM51NN1..4r   c                     t        j                  dt         j                        }t        j                  |      }t        j
                  |      }| j                  |j                  |j                         y )Nr   r   )r   r   int16r   r   r   r   r   )r   r   r   gots       r   test_bug6697zTestCudaNDArray.test_bug6697  sJ    ii"((+~~c"jjDJJ/r   c                    t        j                  ddt        j                        }t        j                  ddt        j                        }| j                  |j                  d       |j                  |       |j                  |       t        j                  |      }| j                  |j                  d       | j                  |j                  d       |j                  |       |j                  |       y )N)r   )r~   )r   r   r   )r   r   r   int8r   r   r   r#   r   r   r   r   )r   	dev_array
host_arraydev_array_from_hosts       r   test_issue_8477zTestCudaNDArray.test_issue_8477  s      --D$46GG=	 ZZdDH
 	++T2 	z*  ,
 #nnZ8,22D9,44d;  !45**95r   N)/__name__
__module____qualname__r   r'   r*   r.   r7   r=   rD   rH   rK   rN   rV   rY   r   rc   rh   ro   rr   r{   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r    r   r   r	   r	   	   s   815*871.
.
.
3 ; 23 43
B	5
5
(F.>
>
98L*= CD3 E3: 785 95 785 95 785 95 785 95 785 95 785 95 785 95 785 95 785 950 CD 6 E 6r   r	   c                       e Zd Zd Zy)TestRecarrayc                    t        j                  ddt         j                  fdt         j                  fg      }t        j                  |j
                  t         j                        |_        t        j                  |j
                  t         j                        dz  |_        |j                  }|j                  }d }t        j                  |      }t        j                  |      } t        j                  |      d|j
                  f   |||       t         j                  j                  ||       t         j                  j                  ||       y )N)   value1value2r   r   c                     t        j                  d      }|| j                  k  r%| j                  |   ||<   | j                  |   ||<   y y )Nr4   )r   gridsizer  r  )xout1out2is       r   testz(TestRecarray.test_recarray.<locals>.test  sA    		!A166z((1+Q((1+Q r   r4   )r   recarrayint64rA   r   r  r  r  
zeros_liker   jitr$   r%   )r   r   expect1expect2r  got1got2s          r   test_recarrayzTestRecarray.test_recarray  s    KKrxx rzz"&
  99QVV288499QVV2::6<((((	& }}W%}}W%!q!&&y!!T40


%%gt4


%%gt4r   N)r   r   r   r  r   r   r   r   r     s    5r   r   c                   Z    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y)TestCoreContiguousc                     | j                  t        j                  |      t        j                  |      j                  d          y )Nr   )r   r   is_contiguous
array_corer    )r   rz   s     r   _test_against_array_corez+TestCoreContiguous._test_against_array_core  s7    %%d+""4(..~>	
r   c                 T    t        j                  dd      }| j                  |       y r   r   r   r  r   d_as     r   test_device_array_like_1dz,TestCoreContiguous.test_device_array_like_1d  "    #.%%c*r   c                 T    t        j                  dd      }| j                  |       y Nr   r[   r   r   r  r  s     r   test_device_array_like_2dz,TestCoreContiguous.test_device_array_like_2d  "    4%%c*r   c                 h    t        j                  dd      }| j                  |j                         y r"  r   r   r  r   r  s     r   #test_device_array_like_2d_transposez6TestCoreContiguous.test_device_array_like_2d_transpose  &    4%%cee,r   c                 T    t        j                  dd      }| j                  |       y )Nr   r[      r   r   r  r  s     r   test_device_array_like_3dz,TestCoreContiguous.test_device_array_like_3d  "    C8%%c*r   c                 T    t        j                  dd      }| j                  |       y )Nr   r   r   r  r  s     r   test_device_array_like_1d_fz.TestCoreContiguous.test_device_array_like_1d_f  r   r   c                 T    t        j                  dd      }| j                  |       y Nr#  r   r   r  r  s     r   test_device_array_like_2d_fz.TestCoreContiguous.test_device_array_like_2d_f  r%  r   c                 h    t        j                  dd      }| j                  |j                         y r2  r'  r  s     r   %test_device_array_like_2d_f_transposez8TestCoreContiguous.test_device_array_like_2d_f_transpose  r)  r   c                 T    t        j                  dd      }| j                  |       y )Nr+  r   r   r  r  s     r   test_device_array_like_3d_fz.TestCoreContiguous.test_device_array_like_3d_f	  r.  r   c                 `    d}t        j                  |      d d d   }| j                  |       y )Nr   r1   r   r   r  r   r   rz   s      r   test_1d_viewzTestCoreContiguous.test_1d_view  s+    xxss#%%d+r   c                 d    d}t        j                  |d      d d d   }| j                  |       y )Nr   r   r   r1   r9  r:  s      r   test_1d_view_fz!TestCoreContiguous.test_1d_view_f  s.    xxS)#A#.%%d+r   c                 j    d}t        j                  |      d d dd d df   }| j                  |       y )Nr#  r1   r9  r:  s      r   test_2d_viewzTestCoreContiguous.test_2d_view  s3    xxssCaCx(%%d+r   c                 n    d}t        j                  |d      d d dd d df   }| j                  |       y )Nr#  r   r   r1   r9  r:  s      r   test_2d_view_fz!TestCoreContiguous.test_2d_view_f  s6    xxS)#A#ss(3%%d+r   N)r   r   r   r  r  r$  r(  r-  r0  r3  r5  r7  r;  r=  r?  rA  r   r   r   r  r    sC    
++-+++-+,
,
,
,r   r  __main__)r   numpyr   numba.cuda.cudadrvr   numbar   numba.cuda.testingr   r   r   r	   r   r  r   mainr   r   r   <module>rH     sY      *  5 .6l 6D5< 569, 9,x zHMMO r   