U
    ,d8                     @   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re
  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   C/tmp/pip-unpacked-wheel-eu7e0c37/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rt	|D ]0}
||
 ||| |
 < ||
 ||| |
 | < qd 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  qd S )N@   r   c                    s   t jj td}t d} | d }|t| k r|t|k rtt|t|d }|||< t   || j	| |< || j
||< d S )Nr   r   r   )r   r   r   r   Z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
   Znumpyr+   Znumba.npr   rZ   Zextensions_usecasesr   r   r   rh   rf   r   rK   rH   mainr   r   r   r   <module>   s   
O  _