§
    ƒÛgï8  ã                   ó  — d dl mZmZ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Zd dlmZ ddlmZmZ  ej        d	ej        fd
ej        dfg¦  «        Z G d„ de¦  «        Z G d„ de¦  «        Zedk    r e
j        ¦   «          dS dS )é    )ÚcudaÚint32Úfloat64Úvoid)ÚTypingError)Útypes)ÚunittestÚCUDATestCaseÚskip_on_cudasimN)Únumpy_supporté   )Útest_struct_model_typeÚ
TestStructÚiÚj)é   é   c                   ó>   — e Zd Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Z	d„ Z
d	S )
ÚTestSharedMemoryIssuec                 ó’   ‡— t          j        d¬¦  «        d„ ¦   «         Št           j        ˆfd„¦   «         } |d         ¦   «          d S )NT)Údevicec                  óR   — t           j                             dt          ¬¦  «        } d S ©Nr   ©Údtype©r   ÚsharedÚarrayr   )Ú	inner_arrs    úU/granary/edev/quantum/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_sm.pyÚinnerzGTestSharedMemoryIssue.test_issue_953_sm_linkage_conflict.<locals>.inner   s    € åœ×)Ò)¨!µ5Ð)Ñ9Ô9ˆIˆIˆIó    c                  óh   •— t           j                             dt          ¬¦  «        }  ‰¦   «          d S r   r   )Ú	outer_arrr!   s    €r    ÚouterzGTestSharedMemoryIssue.test_issue_953_sm_linkage_conflict.<locals>.outer   s,   ø€ åœ×)Ò)¨!µ5Ð)Ñ9Ô9ˆIØˆE‰GŒGˆGˆGˆGr"   ©r   r   )r   Újit)Úselfr%   r!   s     @r    Ú"test_issue_953_sm_linkage_conflictz8TestSharedMemoryIssue.test_issue_953_sm_linkage_conflict   sf   ø€ Ý	Œ˜Ð	Ñ	Ô	ð	:ð 	:ñ 
Ô	ð	:õ 
Œð	ð 	ð 	ð 	ñ 
Œð	ð 	ˆˆdŒ‰Œˆˆˆr"   c                 óÐ   ‡— t           j        ˆfd„¦   «         }t          j        dt          j        ¬¦  «        } |d         |¦  «         |                      |d         |¦  «         d S )Nc                 óh   •— t           j                             ‰t          ¬¦  «        }|j        | d<   d S ©Nr   r   )r   r   r   r   Úsize)ÚaÚarrÚshapes     €r    Úsz9TestSharedMemoryIssue._check_shared_array_size.<locals>.s   s,   ø€ å”+×#Ò# EµÐ#Ñ7Ô7ˆCØ”8ˆAˆa‰DˆDˆDr"   r   r   r&   r   )r   r'   ÚnpÚzerosr   ÚassertEqual)r(   r0   Úexpectedr1   Úresults    `   r    Ú_check_shared_array_sizez.TestSharedMemoryIssue._check_shared_array_size   sq   ø€ Ý	Œð	ð 	ð 	ð 	ñ 
Œð	õ ”˜!¥2¤8Ð,Ñ,Ô,ˆØˆˆ$Œ�‰ŒˆØ×Ò˜ œ HÑ-Ô-Ð-Ð-Ð-r"   c                 ó2   — |                       dd¦  «         d S ©Nr   ©r7   ©r(   s    r    Ú%test_issue_1051_shared_size_broken_1dz;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_1d&   s   € Ø×%Ò% a¨Ñ+Ô+Ð+Ð+Ð+r"   c                 ó2   — |                       dd¦  «         d S )N)r   r   é   r:   r;   s    r    Ú%test_issue_1051_shared_size_broken_2dz;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_2d)   s   € Ø×%Ò% f¨aÑ0Ô0Ð0Ð0Ð0r"   c                 ó2   — |                       dd¦  «         d S )N)r   r   é   é   r:   r;   s    r    Ú%test_issue_1051_shared_size_broken_3dz;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_3d,   s   € à×%Ò% i°Ñ4Ô4Ð4Ð4Ð4r"   c                 óÔ   ‡‡— t           j        ˆˆfd„¦   «         }t          j        dt          j        ¬¦  «        } |d         |¦  «         |                      |d         |¦  «         d S )Nc                 ó^   •— t           j                             ‰‰¬¦  «        }|j        | d<   d S r,   )r   r   r   r-   )r.   r/   r0   Útys     €€r    r1   z>TestSharedMemoryIssue._check_shared_array_size_fp16.<locals>.s1   s,   ø€ å”+×#Ò# E°Ð#Ñ4Ô4ˆCØ”8ˆAˆa‰DˆDˆDr"   r   r   r&   r   )r   r'   r2   r3   Úfloat16r4   )r(   r0   r5   rF   r1   r6   s    ` `  r    Ú_check_shared_array_size_fp16z3TestSharedMemoryIssue._check_shared_array_size_fp160   sw   øø€ Ý	Œð	ð 	ð 	ð 	ð 	ñ 
Œð	õ ”˜!¥2¤:Ð.Ñ.Ô.ˆØˆˆ$Œ�‰ŒˆØ×Ò˜ œ HÑ-Ô-Ð-Ð-Ð-r"   c                 óŠ   — |                       ddt          j        ¦  «         |                       ddt          j        ¦  «         d S r9   )rH   r   rG   r2   r;   s    r    Útest_issue_fp16_supportz-TestSharedMemoryIssue.test_issue_fp16_support:   s<   € Ø×*Ò*¨1¨aµ´Ñ?Ô?Ð?Ø×*Ò*¨1¨aµ´Ñ<Ô<Ð<Ð<Ð<r"   c                 óþ   ‡‡— dŠd}dŠd}t           j        ˆˆfd„¦   «         }t          j        |t          j        ¬¦  «        }t          j        |¦  «        } |||f         |¦  «         t          j        ¦   «          dS )z¡
        Test issue of warp misalign address due to nvvm not knowing the
        alignment(? but it should have taken the natural alignment of the type)
        r   é0   rA   r   c                 ó   •— t           j                             ‰‰ft          ¦  «        }t           j                             dt          ¦  «        }t           j        j        }d}t          ‰¦  «        D ]}||||f         z  }Œ|d         |z   | d<   d S )Né   r   )r   r   r   r   Ú	threadIdxÚxÚrange)Úd_block_costsÚ
s_featuresÚs_initialcostrO   Ú
predictionr   Úexamples_per_blockÚnum_weightss         €€r    Ú
costs_funcz9TestSharedMemoryIssue.test_issue_2393.<locals>.costs_funcH   s‘   ø€ åœ×*Ò*Ð,>ÀÐ+LÝ+2ñ4ô 4ˆJå œK×-Ò-¨aµÑ9Ô9ˆMåœÔ(ˆIàˆJÝ˜;Ñ'Ô'ð 7ð 7�Ø˜j¨°A¨Ô6Ñ6�
�
à,¨QÔ/°*Ñ<ˆM˜!ÑÐÐr"   r   N)r   r'   r2   r3   r   Ú	to_deviceÚsynchronize)r(   Ú
num_blocksÚthreads_per_blockrX   Úblock_costsrR   rV   rW   s         @@r    Útest_issue_2393z%TestSharedMemoryIssue.test_issue_2393>   s�   øø€ ð
 ˆØˆ
ØÐØÐå	Œð	=ð 	=ð 	=ð 	=ð 	=ñ 
Œð	=õ ”h˜zµ´Ð<Ñ<Ô<ˆÝœ {Ñ3Ô3ˆà1ˆ
�:Ð0Ð0Ô1°-Ñ@Ô@Ð@åÔÑÔÐÐÐr"   N)Ú__name__Ú
__module__Ú__qualname__r)   r7   r<   r?   rC   rH   rJ   r^   © r"   r    r   r      sŒ   € € € € € ð
ð 
ð 
ð.ð .ð .ð,ð ,ð ,ð1ð 1ð 1ð5ð 5ð 5ð.ð .ð .ð=ð =ð =ðð ð ð ð r"   r   c                   ó¬   — 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 ed¦  «        d„ ¦   «         Z ed¦  «        d„ ¦   «         ZdS )ÚTestSharedMemoryc                 ój  ‡‡— t          |¦  «        }dŠt          |‰z  ¦  «        }t          j        |j        ¦  «        Št
          j        ˆˆfd„¦   «         }t          j        |¦  «        } ||‰f         ||¦  «         |                     ¦   «         }t          j
                             ||¦  «         d S )Né   c                 óx  •— t           j                             ‰	‰¬¦  «        }t           j        j        }t           j        j        }t           j        j        }||z  |z   }|t          | ¦  «        k     r| |         ||<   t          j        ¦   «          |dk    r#t          ‰	¦  «        D ]}||         |||z  |z   <   Œd S d S r,   ©
r   r   r   rO   rP   ÚblockIdxÚblockDimÚlenÚsyncthreadsrQ   )
rP   ÚyÚsmÚtxÚbxÚbdr   r   ÚdtÚnthreadss
           €€r    Úuse_sm_chunk_copyz8TestSharedMemory._test_shared.<locals>.use_sm_chunk_copyj   s»   ø€ å”×"Ò" 8°2Ð"Ñ6Ô6ˆBå”Ô!ˆBÝ””ˆBÝ””ˆBð �R‘˜"‘ˆAØ•3�q‘6”6ŠzˆzØ˜1œ��2‘åÔÑÔÐð �QŠwˆwÝ˜x™œð +ð +�AØ%'¨¤U�A�b˜2‘g ‘k‘N�Nð ˆwð+ð +r"   )rk   ÚintÚnpsÚ
from_dtyper   r   r'   Údevice_array_likeÚcopy_to_hostr2   ÚtestingÚassert_array_equal)	r(   r/   ÚnelemÚnblocksrt   Úd_resultÚhost_resultrr   rs   s	          @@r    Ú_test_sharedzTestSharedMemory._test_shared_   sÁ   øø€ õ �C‘”ˆØˆÝ�e˜hÑ&Ñ'Ô'ˆÝŒ^˜CœIÑ&Ô&ˆå	Œð	+ð 	+ð 	+ð 	+ð 	+ñ 
Œð	+õ& Ô)¨#Ñ.Ô.ˆØ,Ð˜' 8Ð+Ô,¨S°(Ñ;Ô;Ð;Ø×+Ò+Ñ-Ô-ˆÝ
Œ
×%Ò% c¨;Ñ7Ô7Ð7Ð7Ð7r"   c                 óF  — t          j        dt          ¬¦  «        }t          t	          |¦  «        ¦  «        D ]S}|||         _        t          j        dt           j        ¬¦  «        }|                     dd¦  «        |z  ||         _	        ŒT|  
                    |¦  «         d S )Né€   r   r>   r   r   )r2   ÚrecarrayÚrecordwith2darrayrQ   rk   r   ÚarangeÚfloat32Úreshaper   r€   )r(   r/   rP   r   s       r    Útest_shared_recarrayz%TestSharedMemory.test_shared_recarrayƒ   s‹   € ÝŒk˜#Õ%6Ð7Ñ7Ô7ˆÝ•s˜3‘x”x‘”ð 	+ð 	+ˆAØˆC�ŒFŒHÝ”	˜%¥r¤zÐ2Ñ2Ô2ˆAØ—y’y  A‘”¨Ñ*ˆC�ŒFŒHˆHà×Ò˜#ÑÔÐÐÐr"   c                 óˆ   — t           j                             ddt           j        ¬¦  «        }|                      |¦  «         d S )Nr   )é   )r-   r   )r2   ÚrandomÚrandintÚbool_r€   )r(   r/   s     r    Útest_shared_boolz!TestSharedMemory.test_shared_boolŒ   s:   € ÝŒi×Ò ¨µr´xÐÑ@Ô@ˆØ×Ò˜#ÑÔÐÐÐr"   c                 ó˜   — |j         |j        j        z  } |ddd|f         |¦  «         t          j                             ||¦  «         d S )Nr   r   )r-   r   Úitemsizer2   rz   r{   )r(   Úfuncr/   r5   Únshareds        r    Ú_test_dynshared_slicez&TestSharedMemory._test_dynshared_slice�   sQ   € ð
 ”(˜SœYÔ/Ñ/ˆØˆˆQ��1�gÐÔ˜sÑ#Ô#Ð#Ý
Œ
×%Ò% h°Ñ4Ô4Ð4Ð4Ð4r"   c                 óâ   — t           j        d„ ¦   «         }t          j        dt          j        ¬¦  «        }t          j        ddgt          j        ¬¦  «        }|                      |||¦  «         d S )Nc                 óº   — t           j                             dt          ¬¦  «        }|dd…         }|dd…         }d|d<   d|d<   |d         | d<   |d         | d<   d S ©Nr   r   r   r   r   ©rP   ÚdynsmemÚsm1Úsm2s       r    Úslice_writez@TestSharedMemory.test_dynshared_slice_write.<locals>.slice_write›   sf   € å”k×'Ò'¨µÐ'Ñ7Ô7ˆGØ˜!˜A˜#”,ˆCØ˜!˜A˜#”,ˆCàˆC�‰FØˆC�‰FØ˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DˆDˆDr"   r   r   r   ©r   r'   r2   r3   r   r   r“   )r(   r›   r/   r5   s       r    Útest_dynshared_slice_writez+TestSharedMemory.test_dynshared_slice_write™   sm   € å	Œð	ð 	ñ 
Œð	õ Œh�q¥¤Ð)Ñ)Ô)ˆÝ”8˜Q ˜F­"¬(Ð3Ñ3Ô3ˆØ×"Ò" ;°°XÑ>Ô>Ð>Ð>Ð>r"   c                 óâ   — t           j        d„ ¦   «         }t          j        dt          j        ¬¦  «        }t          j        ddgt          j        ¬¦  «        }|                      |||¦  «         d S )Nc                 óº   — t           j                             dt          ¬¦  «        }|dd…         }|dd…         }d|d<   d|d<   |d         | d<   |d         | d<   d S r–   r   r—   s       r    Ú
slice_readz>TestSharedMemory.test_dynshared_slice_read.<locals>.slice_read¬   sf   € å”k×'Ò'¨µÐ'Ñ7Ô7ˆGØ˜!˜A˜#”,ˆCØ˜!˜A˜#”,ˆCàˆG�A‰JØˆG�A‰JØ�q”6ˆAˆa‰DØ�q”6ˆAˆa‰DˆDˆDr"   r   r   r   rœ   )r(   r    r/   r5   s       r    Útest_dynshared_slice_readz*TestSharedMemory.test_dynshared_slice_readª   sm   € å	Œð	ð 	ñ 
Œð	õ Œh�q¥¤Ð)Ñ)Ô)ˆÝ”8˜Q ˜F­"¬(Ð3Ñ3Ô3ˆØ×"Ò" :¨s°HÑ=Ô=Ð=Ð=Ð=r"   c                 óâ   — t           j        d„ ¦   «         }t          j        dt          j        ¬¦  «        }t          j        g d¢t          j        ¬¦  «        }|                      |||¦  «         d S )Nc                 óÚ   — t           j                             dt          ¬¦  «        }|dd…         }|dd…         }d|d<   d|d<   d|d<   |d         | d<   |d         | d<   |d         | d<   d S )Nr   r   r   r   r   r   r—   s       r    Úslice_diff_sizeszJTestSharedMemory.test_dynshared_slice_diff_sizes.<locals>.slice_diff_sizes¾   s|   € å”k×'Ò'¨µÐ'Ñ7Ô7ˆGØ˜!˜A˜#”,ˆCØ˜!˜A˜#”,ˆCàˆG�A‰JØˆG�A‰JØˆG�A‰JØ�q”6ˆAˆa‰DØ�q”6ˆAˆa‰DØ�q”6ˆAˆa‰DˆDˆDr"   r   r   )r   r   r   rœ   )r(   r¤   r/   r5   s       r    Útest_dynshared_slice_diff_sizesz0TestSharedMemory.test_dynshared_slice_diff_sizes»   sp   € õ 
Œð
	ð 
	ñ 
Œð
	õ Œh�q¥¤Ð)Ñ)Ô)ˆÝ”8˜I˜I˜I­R¬XÐ6Ñ6Ô6ˆØ×"Ò"Ð#3°S¸(ÑCÔCÐCÐCÐCr"   c                 óâ   — t           j        d„ ¦   «         }t          j        dt          j        ¬¦  «        }t          j        g d¢t          j        ¬¦  «        }|                      |||¦  «         d S )Nc                 ó  — t           j                             dt          ¬¦  «        }|dd…         }|dd…         }d|d<   d|d<   d|d<   d|d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   d S )Nr   r   r   r   rA   r   r   r—   s       r    Úslice_overlapzDTestSharedMemory.test_dynshared_slice_overlap.<locals>.slice_overlapÑ   sŸ   € å”k×'Ò'¨µÐ'Ñ7Ô7ˆGØ˜!˜A˜#”,ˆCØ˜!˜A˜#”,ˆCàˆG�A‰JØˆG�A‰JØˆG�A‰JØˆG�A‰JØ�q”6ˆAˆa‰DØ�q”6ˆAˆa‰DØ�q”6ˆAˆa‰DØ�q”6ˆAˆa‰DØ�q”6ˆAˆa‰DˆDˆDr"   é   r   )r   r   r   r   rA   rœ   )r(   r¨   r/   r5   s       r    Útest_dynshared_slice_overlapz-TestSharedMemory.test_dynshared_slice_overlapÏ   sm   € å	Œð	ð 	ñ 
Œð	õ Œh�q¥¤Ð)Ñ)Ô)ˆÝ”8˜O˜O˜Oµ2´8Ð<Ñ<Ô<ˆØ×"Ò" =°#°xÑ@Ô@Ð@Ð@Ð@r"   c                 óâ   — t           j        d„ ¦   «         }t          j        dt          j        ¬¦  «        }t          j        g d¢t          j        ¬¦  «        }|                      |||¦  «         d S )Nc                 ó‚  — t           j                             dt          ¬¦  «        }|dd…         }|dd…         }d|d<   d|d<   d|d<   d|d<   d|d<   d|d	<   d|d<   d|d<   d|d<   d|d<   d|d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   |d	         | d	<   |d         | d<   d S )
Nr   r   r   r   rA   r>   éc   r   r©   r   r—   s       r    Ú
slice_gapsz>TestSharedMemory.test_dynshared_slice_gaps.<locals>.slice_gapsè   sú   € å”k×'Ò'¨µÐ'Ñ7Ô7ˆGØ˜!˜A˜#”,ˆCØ˜!˜A˜#”,ˆCð ˆG�A‰JØˆG�A‰JØˆG�A‰JØˆG�A‰JØˆG�A‰JØˆG�A‰JØˆG�A‰JàˆC�‰FØˆC�‰FØˆC�‰FØˆC�‰Fà˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DˆDˆDr"   rN   r   )r­   r   r   r­   r   rA   r­   rœ   )r(   r®   r/   r5   s       r    Útest_dynshared_slice_gapsz*TestSharedMemory.test_dynshared_slice_gapså   sr   € õ 
Œð	ð 	ñ 
Œð	õ6 Œh�q¥¤Ð)Ñ)Ô)ˆÝ”8Ð4Ð4Ð4½B¼HÐEÑEÔEˆØ×"Ò" :¨s°HÑ=Ô=Ð=Ð=Ð=r"   c                 óâ   — t           j        d„ ¦   «         }t          j        dt          j        ¬¦  «        }t          j        g d¢t          j        ¬¦  «        }|                      |||¦  «         d S )Nc                 óþ   — t           j                             dt          ¬¦  «        }|dd d…         }|ddd…         }d|d<   d|d<   d|d<   d|d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   d S )Nr   r   r   éÿÿÿÿr   r   rA   r   r—   s       r    Úslice_write_backwardszTTestSharedMemory.test_dynshared_slice_write_backwards.<locals>.slice_write_backwards  s–   € å”k×'Ò'¨µÐ'Ñ7Ô7ˆGØ˜!˜%˜R˜%”.ˆCØ˜!˜A˜b˜&”/ˆCàˆC�‰FØˆC�‰FØˆC�‰FØˆC�‰FØ˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DˆDˆDr"   rA   r   )r   r   rA   r   rœ   )r(   r³   r/   r5   s       r    Ú$test_dynshared_slice_write_backwardsz5TestSharedMemory.test_dynshared_slice_write_backwards  sp   € õ 
Œð	ð 	ñ 
Œð	õ Œh�q¥¤Ð)Ñ)Ô)ˆÝ”8˜L˜L˜Lµ´Ð9Ñ9Ô9ˆØ×"Ò"Ð#8¸#¸xÑHÔHÐHÐHÐHr"   c                 óâ   — t           j        d„ ¦   «         }t          j        dt          j        ¬¦  «        }t          j        g d¢t          j        ¬¦  «        }|                      |||¦  «         d S )Nc                 óF  — t           j                             dt          ¬¦  «        }|d d d…         }d|d<   d|d<   d|d<   d|d<   d|d<   d|d<   d|d<   d|d<   d|d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   d S )	Nr   r   r   r­   r   r   rA   r©   r   ©rP   r˜   r™   s      r    Úslice_nonunit_stridezRTestSharedMemory.test_dynshared_slice_nonunit_stride.<locals>.slice_nonunit_stride!  sÐ   € å”k×'Ò'¨µÐ'Ñ7Ô7ˆGØ˜#˜#˜A˜#”,ˆCð ˆG�A‰JØˆG�A‰JØˆG�A‰JØˆG�A‰JØˆG�A‰JØˆG�A‰JàˆC�‰FØˆC�‰FØˆC�‰Fà˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DˆDˆDr"   r>   r   )r   r­   r   r­   r   r­   rœ   )r(   r¸   r/   r5   s       r    Ú#test_dynshared_slice_nonunit_stridez4TestSharedMemory.test_dynshared_slice_nonunit_stride  ss   € õ 
Œð	ð 	ñ 
Œð	õ. Œh�q¥¤Ð)Ñ)Ô)ˆÝ”8Ð1Ð1Ð1½¼ÐBÑBÔBˆØ×"Ò"Ð#7¸¸hÑGÔGÐGÐGÐGr"   c                 óâ   — t           j        d„ ¦   «         }t          j        dt          j        ¬¦  «        }t          j        g d¢t          j        ¬¦  «        }|                      |||¦  «         d S )Nc                 óF  — t           j                             dt          ¬¦  «        }|dd d…         }d|d<   d|d<   d|d<   d|d<   d|d	<   d|d
<   d|d<   d|d<   d|d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   |d	         | d	<   |d
         | d
<   d S )Nr   r   r²   éþÿÿÿr­   r   r   r   rA   r©   r   r·   s      r    Úslice_nonunit_reverse_stridezbTestSharedMemory.test_dynshared_slice_nonunit_reverse_stride.<locals>.slice_nonunit_reverse_stride@  sÐ   € å”k×'Ò'¨µÐ'Ñ7Ô7ˆGØ˜"˜&˜b˜&”/ˆCð ˆG�A‰JØˆG�A‰JØˆG�A‰JØˆG�A‰JØˆG�A‰JØˆG�A‰JàˆC�‰FØˆC�‰FØˆC�‰Fà˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DØ˜1”:ˆAˆa‰DˆDˆDr"   r>   r   )r­   r   r­   r   r­   r   rœ   )r(   r½   r/   r5   s       r    Ú+test_dynshared_slice_nonunit_reverse_stridez<TestSharedMemory.test_dynshared_slice_nonunit_reverse_stride=  ss   € õ 
Œð	ð 	ñ 
Œð	õ. Œh�q¥¤Ð)Ñ)Ô)ˆÝ”8Ð1Ð1Ð1½¼ÐBÑBÔBˆØ×"Ò"Ð#?ÀÀhÑOÔOÐOÐOÐOr"   c                 óÖ  ‡
— t          j        d¦  «        }t          |¦  «        }d}t          ||z  ¦  «        }t	          j        |j        ¦  «        Š
||j        j        z  }t          |dz  ¦  «        }t          j	        ˆ
fd„¦   «         }t          j
        |¦  «        } |||d|f         |||¦  «         |                     ¦   «         }	t           j                             ||	¦  «         d S )NrŠ   rf   r   c                 óø  •— t           j                             d‰¬¦  «        }|d|…         }|||dz  …         }t           j        j        }t           j        j        }t           j        j        }||z  |z   }	|	t          | ¦  «        k     r ||k     r| |	         ||<   n| |	         |||z
  <   t          j        ¦   «          |dk    r7t          |¦  «        D ])}
||
         |||z  |
z   <   ||
         |||z  |
z   |z   <   Œ(d S d S )Nr   r   r   rh   )rP   rm   Ú	chunksizer˜   r™   rš   ro   rp   rq   r   r   rr   s              €r    Úsm_slice_copyz7TestSharedMemory.test_issue_5073.<locals>.sm_slice_copyk  s  ø€ å”k×'Ò'¨°Ð'Ñ4Ô4ˆGØ˜!˜I˜+Ô&ˆCØ˜) I°¡MÐ1Ô2ˆCå”Ô!ˆBÝ””ˆBÝ””ˆBð �R‘˜"‘ˆAØ•3�q‘6”6ŠzˆzØ˜	’>�>Ø œd�C˜‘G�Gà*+¨A¬$�C˜˜Y™Ñ'åÔÑÔÐð �QŠwˆwÝ˜yÑ)Ô)ð 8ð 8�AØ%(¨¤V�A�b˜2‘g ‘k‘NØ14°Q´�A�b˜2‘g ‘k IÑ-Ñ.Ð.ð ˆwð8ð 8r"   r   )r2   r…   rk   ru   rv   rw   r   r�   r   r'   rx   ry   rz   r{   )r(   r/   r|   rs   r}   r’   rÁ   rÂ   r~   r   rr   s             @r    Útest_issue_5073z TestSharedMemory.test_issue_5073\  sí   ø€ õ Œi˜‰oŒoˆÝ�C‘”ˆØˆÝ�e˜hÑ&Ñ'Ô'ˆÝŒ^˜CœIÑ&Ô&ˆØ˜SœYÔ/Ñ/ˆÝ˜ 1™Ñ%Ô%ˆ	å	Œð	8ð 	8ð 	8ð 	8ñ 
Œð	8õ2 Ô)¨#Ñ.Ô.ˆØ4ˆ�g˜x¨¨GÐ3Ô4°S¸(ÀIÑNÔNÐNØ×+Ò+Ñ-Ô-ˆÝ
Œ
×%Ò% c¨;Ñ7Ô7Ð7Ð7Ð7r"   zCan't check typing in simulatorc                 óŒ  — d}d„ }|                       t          |¦  «        5   t          j        t	          ¦   «         ¦  «        |¦  «         d d d ¦  «         n# 1 swxY w Y   d}d„ }|                       t          |¦  «        5   t          j        t	          ¦   «         ¦  «        |¦  «         d d d ¦  «         d S # 1 swxY w Y   d S )Nz+.*Cannot infer the type of variable 'arr'.*c                  ól   — t           j                             dt          j        d¦  «        ¬¦  «        } d S )Né
   ÚOr   )r   r   r   r2   r   ©r/   s    r    Úunsupported_typezBTestSharedMemory.test_invalid_array_type.<locals>.unsupported_typeŽ  s(   € Ý”+×#Ò# B­b¬h°s©m¬mÐ#Ñ<Ô<ˆCˆCˆCr"   z*.*Invalid NumPy dtype specified: 'int33'.*c                  óH   — t           j                             dd¬¦  «        } d S )NrÆ   Úint33r   )r   r   r   rÈ   s    r    Úinvalid_string_typezETestSharedMemory.test_invalid_array_type.<locals>.invalid_string_type•  s    € Ý”+×#Ò# B¨gÐ#Ñ6Ô6ˆCˆCˆCr"   )ÚassertRaisesRegexr   r   r'   r   )r(   ÚrgxrÉ   rÌ   s       r    Útest_invalid_array_typez(TestSharedMemory.test_invalid_array_typeŠ  sO  € à;ˆð	=ð 	=ð 	=à×#Ò#¥K°Ñ5Ô5ð 	/ð 	/Ø�DŒH•T‘V”VÑÔÐ-Ñ.Ô.Ð.ð	/ð 	/ð 	/ñ 	/ô 	/ð 	/ð 	/ð 	/ð 	/ð 	/ð 	/øøøð 	/ð 	/ð 	/ð 	/ð ;ˆð	7ð 	7ð 	7à×#Ò#¥K°Ñ5Ô5ð 	2ð 	2Ø�DŒH•T‘V”VÑÔÐ0Ñ1Ô1Ð1ð	2ð 	2ð 	2ñ 	2ô 	2ð 	2ð 	2ð 	2ð 	2ð 	2ð 	2ð 	2øøøð 	2ð 	2ð 	2ð 	2ð 	2ð 	2s#   ¡*AÁAÁAÂ*B9Â9B=Ã B=z+Struct model array unsupported in simulatorc           	      óì  ‡— dŠt          j        t          t          d d d…         t          d d d…         ¦  «        ¦  «        ˆfd„¦   «         }t	          j        ‰fd¬¦  «        }t	          j        ‰fd¬¦  «        } |d‰f         ||¦  «         t          |¦  «        D ]!\  }}|                      |‰|z
  dz
  ¦  «         Œ"t          |¦  «        D ]$\  }}|                      |‰|z
  dz
  dz  ¦  «         Œ%d S )Né@   r   c                 óª  •— t           j                             ‰t          ¬¦  «        }t          j        d¦  «        }‰|z
  dz
  }|t          | ¦  «        k     rz|t          |¦  «        k     rit          t          |¦  «        t          |dz  ¦  «        ¦  «        }|||<   t          j        ¦   «          ||         j	        | |<   ||         j
        ||<   d S d S d S )Nr   r   r   )r   r   r   r   Úgridrk   r   r   rl   rP   rm   )ÚoutxÚoutyr/   r   ÚriÚobjrs   s         €r    Úwrite_then_reverse_read_staticzVTestSharedMemory.test_struct_model_type_static.<locals>.write_then_reverse_read_staticž  s¼   ø€ õ ”+×#Ò# HÕ4JÐ#ÑKÔKˆCå”	˜!‘”ˆAØ˜A‘ Ñ!ˆBà•3�t‘9”9Š}ˆ} ¥S¨¡Y¤Y¢ å ¥ q¡¤­5°°Q±©<¬<Ñ8Ô8�Ø��A‘åÔ Ñ"Ô"Ð"à˜bœ'œ)��Q‘Ø˜bœ'œ)��Q‘��ð ˆ}  r"   r   r   r   )r   r'   r   r   r2   r3   Ú	enumerater4   )r(   rØ   ÚarrxÚarryr   rP   rm   rs   s          @r    Útest_struct_model_type_staticz.TestSharedMemory.test_struct_model_type_staticš  s#  ø€ àˆå	Œ•$•u˜S˜S˜q˜S”z¥5¨¨¨1¨¤:Ñ.Ô.Ñ	/Ô	/ð	$ð 	$ð 	$ð 	$ñ 
0Ô	/ð	$õ" Œx˜˜¨7Ð3Ñ3Ô3ˆÝŒx˜˜¨7Ð3Ñ3Ô3ˆà3Ð& q¨( {Ô3°D¸$Ñ?Ô?Ð?å˜d‘O”Oð 	2ð 	2‰DˆAˆqØ×Ò˜Q ¨1¡¨qÑ 0Ñ1Ô1Ð1Ð1Ý˜d‘O”Oð 	8ð 	8‰DˆAˆqØ×Ò˜Q ¨A¡°Ñ!1°QÑ 6Ñ7Ô7Ð7Ð7ð	8ð 	8r"   N)r_   r`   ra   r€   rˆ   rŽ   r“   r�   r¡   r¥   rª   r¯   r´   r¹   r¾   rÃ   r   rÏ   rÜ   rb   r"   r    rd   rd   ^   s8  € € € € € ð"8ð "8ð "8ðHð ð ðð ð ð5ð 5ð 5ð?ð ?ð ?ð">ð >ð >ð"Dð Dð Dð(Að Að Að,!>ð !>ð !>ðFIð Ið Ið,Hð Hð Hð>Pð Pð Pð>,8ð ,8ð ,8ð\ €_Ð6Ñ7Ô7ð2ð 2ñ 8Ô7ð2ð €_ÐBÑCÔCð8ð 8ñ DÔCð8ð 8ð 8r"   rd   Ú__main__)Únumbar   r   r   r   Únumba.core.errorsr   Ú
numba.corer   Únumba.cuda.testingr	   r
   r   Únumpyr2   Únumba.npr   rv   Úextensions_usecasesr   r   r   r†   r„   r   rd   r_   Úmainrb   r"   r    ú<module>ræ      sZ  ðØ ,Ð ,Ð ,Ð ,Ð ,Ð ,Ð ,Ð ,Ð ,Ð ,Ð ,Ð ,Ø )Ð )Ð )Ð )Ð )Ð )Ø Ð Ð Ð Ð Ð Ø FÐ FÐ FÐ FÐ FÐ FÐ FÐ FÐ FÐ Fà Ð Ð Ð Ø )Ð )Ð )Ð )Ð )Ð )à CÐ CÐ CÐ CÐ CÐ CÐ CÐ Cà�B”H˜s B¤H˜oØ" B¤J°Ð7ð9ñ :ô :Ð ðLð Lð Lð Lð L˜Lñ Lô Lð Lð^Z8ð Z8ð Z8ð Z8ð Z8�|ñ Z8ô Z8ð Z8ðz
 ˆzÒÐØ€H„M�O„O€O€O€Oð Ðr"   