o
    i8                     @   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 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)devicec                  S   s   t jjdtd} d S Nr   dtyper   sharedarrayr   )	inner_arr r   \/home/ubuntu/veenaModal/venv/lib/python3.10/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   )	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      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   r4   r$   r   r   r   %test_issue_1051_shared_size_broken_1d&      z;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_1dc                 C      |  dd d S )N)r   r      r6   r7   r   r   r   %test_issue_1051_shared_size_broken_2d)   r9   z;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_2dc                 C   r:   )N)r   r         r6   r7   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   r.   z>TestSharedMemoryIssue._check_shared_array_size_fp16.<locals>.sr   r   r"   r   )r   r#   r/   r0   float16r1   )r$   r,   r2   rA   r-   r3   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 r5   )rC   r   rB   r/   r7   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   r=   r   c                    s^   t j ft}t jdt}t jj}d}tD ]
}||||f 7 }q|d | | d< d S )N   r   )r   r   r   r   	threadIdxxrange)d_block_costs
s_featuress_initialcostrG   
predictionr   examples_per_blocknum_weightsr   r   
costs_funcH   s   z9TestSharedMemoryIssue.test_issue_2393.<locals>.costs_funcr   N)r   r#   r/   r0   r   	to_devicesynchronize)r$   
num_blocksthreads_per_blockrQ   block_costsrJ   r   rN   r   test_issue_2393>   s   
z%TestSharedMemoryIssue.test_issue_2393N)__name__
__module____qualname__r%   r4   r8   r<   r?   rC   rD   rW   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 r&| | ||< t   |dkr?t	D ]}|| ||| | < q2d S d S r&   
r   r   r   rG   rH   blockIdxblockDimlensyncthreadsrI   )rH   ysmtxbxbdr   r   dtnthreadsr   r   use_sm_chunk_copyj   s   z8TestSharedMemory._test_shared.<locals>.use_sm_chunk_copy)r`   intnps
from_dtyper   r   r#   device_array_likecopy_to_hostr/   testingassert_array_equal)r$   r*   nelemnblocksrj   d_resulthost_resultr   rg   r   _test_shared_   s   
zTestSharedMemory._test_sharedc                 C   s^   t jdtd}tt|D ]}||| _t jdt jd}|dd| || _	q| 
| d S )N   r   r;   r   r   )r/   recarrayrecordwith2darrayrI   r`   r   arangefloat32reshaper   rv   )r$   r*   rH   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randintbool_rv   )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/   rp   rq   )r$   funcr*   r2   nsharedr   r   r   _test_dynshared_slice   s   z&TestSharedMemory._test_dynshared_slicec                 C   D   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   rH   dynsmemsm1sm2r   r   r   slice_write      z@TestSharedMemory.test_dynshared_slice_write.<locals>.slice_writer   r   r   r   r#   r/   r0   r   r   r   )r$   r   r*   r2   r   r   r   test_dynshared_slice_write   
   

z+TestSharedMemory.test_dynshared_slice_writec                 C   r   )Nc                 S   sT   t jjdtd}|dd }|dd }d|d< d|d< |d | d< |d | d< d S r   r   r   r   r   r   
slice_read   r   z>TestSharedMemory.test_dynshared_slice_read.<locals>.slice_readr   r   r   r   )r$   r   r*   r2   r   r   r   test_dynshared_slice_read   r   z*TestSharedMemory.test_dynshared_slice_readc                 C   D   t jdd }tjdtjd}tjg d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   r   r   r   r   slice_diff_sizes   s   zJTestSharedMemory.test_dynshared_slice_diff_sizes.<locals>.slice_diff_sizesr   r   )r   r   r   r   )r$   r   r*   r2   r   r   r   test_dynshared_slice_diff_sizes   s
   
z0TestSharedMemory.test_dynshared_slice_diff_sizesc                 C   r   )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=   r   r   r   r   r   r   slice_overlap   s   zDTestSharedMemory.test_dynshared_slice_overlap.<locals>.slice_overlap   r   )r   r   r   r   r=   r   )r$   r   r*   r2   r   r   r   test_dynshared_slice_overlap   s
   
z-TestSharedMemory.test_dynshared_slice_overlapc                 C   r   )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   r=   r;   c   r   r   r   r   r   r   r   
slice_gaps   s*   z>TestSharedMemory.test_dynshared_slice_gaps.<locals>.slice_gapsrF   r   )r   r   r   r   r   r=   r   r   )r$   r   r*   r2   r   r   r   test_dynshared_slice_gaps   s
   
z*TestSharedMemory.test_dynshared_slice_gapsc                 C   r   )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   r=   r   r   r   r   r   slice_write_backwards  s   zTTestSharedMemory.test_dynshared_slice_write_backwards.<locals>.slice_write_backwardsr=   r   )r   r   r=   r   r   )r$   r   r*   r2   r   r   r   $test_dynshared_slice_write_backwards  s
   
z5TestSharedMemory.test_dynshared_slice_write_backwardsc                 C   r   )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=   r   r   rH   r   r   r   r   r   slice_nonunit_stride!  "   zRTestSharedMemory.test_dynshared_slice_nonunit_stride.<locals>.slice_nonunit_strider;   r   )r   r   r   r   r   r   r   )r$   r   r*   r2   r   r   r   #test_dynshared_slice_nonunit_stride  
   
z4TestSharedMemory.test_dynshared_slice_nonunit_stridec                 C   r   )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   r=   r   r   r   r   r   r   slice_nonunit_reverse_stride@  r   zbTestSharedMemory.test_dynshared_slice_nonunit_reverse_stride.<locals>.slice_nonunit_reverse_strider;   r   )r   r   r   r   r   r   r   )r$   r   r*   r2   r   r   r   +test_dynshared_slice_nonunit_reverse_stride=  r   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 )Nr~   r\   r   c                    s   t jjd d}|d| }|||d  }t jj}t jj}t jj}|| | }	|	t| k rA||k r9| |	 ||< n| |	 ||| < t   |dkrft	|D ]}
||
 ||| |
 < ||
 ||| |
 | < qMd S d S )Nr   r   r   r]   )rH   rb   	chunksizer   r   r   rd   re   rf   r   r   rh   r   r   sm_slice_copyk  s$   z7TestSharedMemory.test_issue_5073.<locals>.sm_slice_copyr   )r/   rz   r`   rk   rl   rm   r   r   r   r#   rn   ro   rp   rq   )
r$   r*   rr   ri   rs   r   r   r   rt   ru   r   r   r   test_issue_5073\  s   

z TestSharedMemory.test_issue_5073zCan't check typing in simulatorc                 C   s   d}dd }|  t| tt | W d    n1 sw   Y  d}dd }|  t| tt | W d    d S 1 sDw   Y  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   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$   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  q7t|D ]\}}| | | d d  qJd S )N@   r   c                    s   t jj td}t d} | d }|t| k rB|t|k rDtt|t|d }|||< t   || j	| |< || j
||< d S d S d S )Nr   r   r   )r   r   r   r   gridr`   r   r   ra   rH   rb   )outxoutyr*   r   riobjri   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/   r0   	enumerater1   )r$   r   arrxarryr   rH   rb   r   r   r   test_struct_model_type_static  s   "z.TestSharedMemory.test_struct_model_type_staticN)rX   rY   rZ   rv   r}   r   r   r   r   r   r   r   r   r   r   r   r
   r   r   r   r   r   r   r[   ^   s$    $		#.
r[   __main__)numbar   r   r   r   numba.core.errorsr   
numba.corer   numba.cuda.testingr   r	   r
   numpyr/   numba.npr   rl   extensions_usecasesr   r   r   r{   ry   r   r[   rX   mainr   r   r   r   <module>   s"    
O  _