U
    hâËdï8  ã                   @   s¸   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 d	ejfd
ejdfg¡ZG dd„ deƒZG dd„ deƒZedkr´e
 ¡  dS )é    )ÚcudaÚint32Úfloat64Úvoid)ÚTypingError)Útypes)ÚunittestÚCUDATestCaseÚskip_on_cudasimN)Únumpy_supporté   )Útest_struct_model_typeÚ
TestStructÚiÚj)é   é   c                   @   sL   e Zd Zdd„ Zdd„ Zdd„ Zdd„ Zd	d
„ Zdd„ Zdd„ Z	dd„ Z
dS )ÚTestSharedMemoryIssuec                    s4   t jdd�dd„ ƒ‰ t j‡ fdd„ƒ}|d ƒ  d S )NT)Zdevicec                  S   s   t jjdtd�} d S ©Nr   ©Údtype©r   ÚsharedÚarrayr   )Z	inner_arr© r   úX/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/tests/cudapy/test_sm.pyÚinner   s    zGTestSharedMemoryIssue.test_issue_953_sm_linkage_conflict.<locals>.innerc                     s   t jjdtd�} ˆ ƒ  d S r   r   )Z	outer_arr©r   r   r   Úouter   s    zGTestSharedMemoryIssue.test_issue_953_sm_linkage_conflict.<locals>.outer©r   r   )r   Újit)Úselfr   r   r   r   Ú"test_issue_953_sm_linkage_conflict   s
    

z8TestSharedMemoryIssue.test_issue_953_sm_linkage_conflictc                    sB   t j‡ fdd„ƒ}tjdtjd�}|d |ƒ |  |d |¡ d S )Nc                    s   t jjˆ td�}|j| d< d S ©Nr   r   )r   r   r   r   Úsize©ÚaÚarr©Úshaper   r   Ús   s    z9TestSharedMemoryIssue._check_shared_array_size.<locals>.sr   r   r   r   )r   r    ÚnpÚzerosr   ÚassertEqual)r!   r)   Úexpectedr*   Úresultr   r(   r   Ú_check_shared_array_size   s
    z.TestSharedMemoryIssue._check_shared_array_sizec                 C   s   |   dd¡ d S ©Nr   ©r0   ©r!   r   r   r   Ú%test_issue_1051_shared_size_broken_1d&   s    z;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_1dc                 C   s   |   dd¡ d S )N)r   r   é   r2   r3   r   r   r   Ú%test_issue_1051_shared_size_broken_2d)   s    z;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_2dc                 C   s   |   dd¡ d S )N)r   r   é   é   r2   r3   r   r   r   Ú%test_issue_1051_shared_size_broken_3d,   s    z;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_3dc                    sD   t j‡ ‡fdd„ƒ}tjdtjd�}|d |ƒ |  |d |¡ d S )Nc                    s   t jjˆ ˆd�}|j| d< d S r#   )r   r   r   r$   r%   ©r)   Útyr   r   r*   1   s    z>TestSharedMemoryIssue._check_shared_array_size_fp16.<locals>.sr   r   r   r   )r   r    r+   r,   Úfloat16r-   )r!   r)   r.   r;   r*   r/   r   r:   r   Ú_check_shared_array_size_fp160   s
    z3TestSharedMemoryIssue._check_shared_array_size_fp16c                 C   s$   |   ddtj¡ |   ddtj¡ d S r1   )r=   r   r<   r+   r3   r   r   r   Útest_issue_fp16_support:   s    z-TestSharedMemoryIssue.test_issue_fp16_supportc                    sZ   d‰d}d‰ d}t j‡ ‡fdd„ƒ}tj|tjd�}t  |¡}|||f |ƒ t  ¡  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   r7   r   c                    s^   t j ˆ ˆft¡}t j dt¡}t jj}d}tˆƒD ]}||||f 7 }q4|d | | d< d S )Né   r   )r   r   r   r   Ú	threadIdxÚxÚrange)Úd_block_costsZ
s_featuresZs_initialcostrA   Z
predictionr   ©Zexamples_per_blockZnum_weightsr   r   Ú
costs_funcH   s    ÿz9TestSharedMemoryIssue.test_issue_2393.<locals>.costs_funcr   N)r   r    r+   r,   r   Z	to_deviceZsynchronize)r!   Z
num_blocksZthreads_per_blockrF   Zblock_costsrD   r   rE   r   Útest_issue_2393>   s    
z%TestSharedMemoryIssue.test_issue_2393N)Ú__name__Ú
__module__Ú__qualname__r"   r0   r4   r6   r9   r=   r>   rG   r   r   r   r   r      s   

r   c                   @   s”   e Zd Zdd„ Zdd„ Zdd„ Zdd„ Zd	d
„ Zdd„ Zdd„ Z	dd„ Z
dd„ Zdd„ Zdd„ Zdd„ Zdd„ Zedƒdd„ ƒZedƒdd „ ƒZd!S )"ÚTestSharedMemoryc                    sn   t |ƒ}d‰t|ˆ ƒ}t |j¡‰ tj‡ ‡fdd„ƒ}t |¡}||ˆf ||ƒ | ¡ }t	j
 ||¡ d S )Né   c                    s‚   t jjˆˆ d�}t jj}t jj}t jj}|| | }|t| ƒk rL| | ||< t  ¡  |dkr~t	ˆƒD ]}|| ||| | < qdd S r#   ©
r   r   r   rA   rB   ZblockIdxZblockDimÚlenÚsyncthreadsrC   )rB   ÚyÚsmÚtxÚbxÚbdr   r   ©ÚdtÚnthreadsr   r   Úuse_sm_chunk_copyj   s    z8TestSharedMemory._test_shared.<locals>.use_sm_chunk_copy)rN   ÚintÚnpsÚ
from_dtyper   r   r    Údevice_array_likeÚcopy_to_hostr+   ÚtestingÚassert_array_equal)r!   r'   ÚnelemÚnblocksrX   Úd_resultÚhost_resultr   rU   r   Ú_test_shared_   s    
zTestSharedMemory._test_sharedc                 C   s^   t jdtd�}tt|ƒƒD ]4}||| _t jdt jd�}| dd¡| || _	q|  
|¡ d S )Né€   r   r5   r   r   )r+   ZrecarrayÚrecordwith2darrayrC   rN   r   ÚarangeÚfloat32Zreshaper   rd   )r!   r'   rB   r   r   r   r   Útest_shared_recarrayƒ   s    
z%TestSharedMemory.test_shared_recarrayc                 C   s"   t jjddt jd�}|  |¡ d S )Nr   )é   )r$   r   )r+   ÚrandomÚrandintZbool_rd   )r!   r'   r   r   r   Útest_shared_boolŒ   s    z!TestSharedMemory.test_shared_boolc                 C   s4   |j |jj }|ddd|f |ƒ tj ||¡ d S )Nr   r   )r$   r   Úitemsizer+   r^   r_   )r!   Úfuncr'   r.   Únsharedr   r   r   Ú_test_dynshared_slice�   s    z&TestSharedMemory._test_dynshared_slicec                 C   sD   t jdd„ ƒ}tjdtjd�}tjddgtjd�}|  |||¡ d S )Nc                 S   sT   t jjdtd�}|dd… }|dd… }d|d< d|d< |d | d< |d | d< d S ©Nr   r   r   r   r   ©rB   ÚdynsmemÚsm1Úsm2r   r   r   Úslice_write›   s    z@TestSharedMemory.test_dynshared_slice_write.<locals>.slice_writer   r   r   ©r   r    r+   r,   r   r   rq   )r!   rw   r'   r.   r   r   r   Útest_dynshared_slice_write™   s
    

z+TestSharedMemory.test_dynshared_slice_writec                 C   sD   t jdd„ ƒ}tjdtjd�}tjddgtjd�}|  |||¡ d S )Nc                 S   sT   t jjdtd�}|dd… }|dd… }d|d< d|d< |d | d< |d | d< d S rr   r   rs   r   r   r   Ú
slice_read¬   s    z>TestSharedMemory.test_dynshared_slice_read.<locals>.slice_readr   r   r   rx   )r!   rz   r'   r.   r   r   r   Útest_dynshared_slice_readª   s
    

z*TestSharedMemory.test_dynshared_slice_readc                 C   sF   t jdd„ ƒ}tjdtjd�}tjdddgtjd�}|  |||¡ d S )Nc                 S   sh   t jjdt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   rs   r   r   r   Úslice_diff_sizes¾   s    zJTestSharedMemory.test_dynshared_slice_diff_sizes.<locals>.slice_diff_sizesr   r   r   r   rx   )r!   r|   r'   r.   r   r   r   Útest_dynshared_slice_diff_sizes»   s
    
z0TestSharedMemory.test_dynshared_slice_diff_sizesc                 C   sJ   t jdd„ ƒ}tjdtjd�}tjdddddgtjd�}|  |||¡ d S )	Nc                 S   sˆ   t j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 S )Nr   r   r   r   r7   r   r   rs   r   r   r   Úslice_overlapÑ   s    zDTestSharedMemory.test_dynshared_slice_overlap.<locals>.slice_overlapé   r   r   r   r   r7   rx   )r!   r~   r'   r.   r   r   r   Útest_dynshared_slice_overlapÏ   s
    
z-TestSharedMemory.test_dynshared_slice_overlapc                 C   sN   t jdd„ ƒ}tjdtjd�}tjdddddd	dgtjd�}|  |||¡ d S )
Nc                 S   sØ   t j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< d S )
Nr   r   r   r   r7   r5   éc   r   r   r   rs   r   r   r   Ú
slice_gapsè   s*    z>TestSharedMemory.test_dynshared_slice_gaps.<locals>.slice_gapsr@   r   r�   r   r   r   r7   rx   )r!   r‚   r'   r.   r   r   r   Útest_dynshared_slice_gapså   s
    
z*TestSharedMemory.test_dynshared_slice_gapsc                 C   sH   t jdd„ ƒ}tjdtjd�}tjddddgtjd�}|  |||¡ d S )Nc                 S   s€   t j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 S )Nr   r   r   éÿÿÿÿr   r   r7   r   rs   r   r   r   Úslice_write_backwards  s    zTTestSharedMemory.test_dynshared_slice_write_backwards.<locals>.slice_write_backwardsr7   r   r   r   r   rx   )r!   r…   r'   r.   r   r   r   Ú$test_dynshared_slice_write_backwards  s
    
z5TestSharedMemory.test_dynshared_slice_write_backwardsc                 C   sL   t jdd„ ƒ}tjdtjd�}tjddddddgtjd�}|  |||¡ d S )	Nc                 S   s²   t j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 S )	Nr   r   r   r�   r   r   r7   r   r   ©rB   rt   ru   r   r   r   Úslice_nonunit_stride!  s"    zRTestSharedMemory.test_dynshared_slice_nonunit_stride.<locals>.slice_nonunit_strider5   r   r   r�   r   r   rx   )r!   rˆ   r'   r.   r   r   r   Ú#test_dynshared_slice_nonunit_stride  s
    
z4TestSharedMemory.test_dynshared_slice_nonunit_stridec                 C   sL   t jdd„ ƒ}tjdtjd�}tjddddddgtjd�}|  |||¡ d S )	Nc                 S   s²   t j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 S )Nr   r   r„   éþÿÿÿr�   r   r   r   r7   r   r   r‡   r   r   r   Úslice_nonunit_reverse_stride@  s"    zbTestSharedMemory.test_dynshared_slice_nonunit_reverse_stride.<locals>.slice_nonunit_reverse_strider5   r   r�   r   r   r   rx   )r!   r‹   r'   r.   r   r   r   Ú+test_dynshared_slice_nonunit_reverse_stride=  s
    
z<TestSharedMemory.test_dynshared_slice_nonunit_reverse_stridec           
         s”   t  d¡}t|ƒ}d}t|| ƒ}t |j¡‰ ||jj }t|d ƒ}tj	‡ fdd„ƒ}t 
|¡}|||d|f |||ƒ | ¡ }	t j ||	¡ d S )Nrj   rL   r   c                    sÐ   t jjdˆ d�}|d|… }|||d … }t jj}t jj}t jj}|| | }	|	t| ƒk r‚||k rr| |	 ||< n| |	 ||| < t  ¡  |dkrÌt	|ƒD ]0}
||
 ||| |
 < ||
 ||| |
 | < qšd S )Nr   r   r   rM   )rB   rP   Ú	chunksizert   ru   rv   rR   rS   rT   r   r   ©rV   r   r   Úsm_slice_copyk  s     z7TestSharedMemory.test_issue_5073.<locals>.sm_slice_copyr   )r+   rg   rN   rY   rZ   r[   r   rn   r   r    r\   r]   r^   r_   )
r!   r'   r`   rW   ra   rp   r�   r�   rb   rc   r   rŽ   r   Útest_issue_5073\  s    

z TestSharedMemory.test_issue_5073zCan't check typing in simulatorc              	   C   sl   d}dd„ }|   t|¡� t tƒ ¡|ƒ W 5 Q R X d}dd„ }|   t|¡� t tƒ ¡|ƒ W 5 Q R X d S )Nz+.*Cannot infer the type of variable 'arr'.*c                  S   s   t jjdt d¡d�} d S )Né
   ÚOr   )r   r   r   r+   r   ©r'   r   r   r   Úunsupported_typeŽ  s    zBTestSharedMemory.test_invalid_array_type.<locals>.unsupported_typez*.*Invalid NumPy dtype specified: 'int33'.*c                  S   s   t jjddd�} d S )Nr‘   Zint33r   )r   r   r   r“   r   r   r   Úinvalid_string_type•  s    zETestSharedMemory.test_invalid_array_type.<locals>.invalid_string_type)ÚassertRaisesRegexr   r   r    r   )r!   Zrgxr”   r•   r   r   r   Útest_invalid_array_typeŠ  s    z(TestSharedMemory.test_invalid_array_typez+Struct model array unsupported in simulatorc                    sº   d‰ t  ttd d d… td d d… ƒ¡‡ fdd„ƒ}tjˆ fdd�}tjˆ fdd�}|dˆ f ||ƒ t|ƒD ]\}}|  |ˆ | d ¡ qnt|ƒD ] \}}|  |ˆ | d d ¡ q”d S )Né@   r   c                    s„   t jjˆ td�}t  d¡}ˆ | d }|t| ƒk r€|t|ƒk r€tt|ƒt|d ƒƒ}|||< t  ¡  || j	| |< || j
||< d S )Nr   r   r   )r   r   r   r   ÚgridrN   r   r   rO   rB   rP   )ZoutxZoutyr'   r   ÚriÚobj©rW   r   r   Úwrite_then_reverse_read_staticž  s    
zVTestSharedMemory.test_struct_model_type_static.<locals>.write_then_reverse_read_staticr   r   r   )r   r    r   r   r+   r,   Ú	enumerater-   )r!   r�   ZarrxZarryr   rB   rP   r   rœ   r   Útest_struct_model_type_staticš  s    "z.TestSharedMemory.test_struct_model_type_staticN)rH   rI   rJ   rd   ri   rm   rq   ry   r{   r}   r€   rƒ   r†   r‰   rŒ   r�   r
   r—   rŸ   r   r   r   r   rK   ^   s"   $		#.
rK   Ú__main__)Znumbar   r   r   r   Znumba.core.errorsr   Z
numba.corer   Znumba.cuda.testingr   r	   r
   Únumpyr+   Znumba.npr   rZ   Zextensions_usecasesr   r   r   rh   rf   r   rK   rH   Úmainr   r   r   r   Ú<module>   s   
ÿO  _