U
    hâËd   ã                   @   s    d dl Z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mZmZ dd	„ Zd
d„ Zdd„ ZedƒG dd„ deƒƒZedkrœe ¡  dS )é    N)ÚcudaÚint32Ú
complex128Úvoid)Útypes)ÚTypingError)ÚunittestÚCUDATestCaseÚskip_on_cudasimé   )Útest_struct_model_typeÚ
TestStructc                 C   sT   t jjdtd�}t|jd ƒD ]}| | ||< qt|jd ƒD ]}|| ||< q>d S )Néè  ©Údtyper   ©r   ÚlocalÚarrayr   ÚrangeÚshape©ÚAÚBÚCÚi© r   ú^/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/tests/cudapy/test_localmem.pyÚculocal
   s
    r   c                 C   sT   t jjdtd�}t|jd ƒD ]}| | ||< qt|jd ƒD ]}|| ||< q>d S )Néd   r   r   )r   r   r   r   r   r   r   r   r   r   Úculocalcomplex   s
    r   c                 C   sT   t jjdtd�}t|jd ƒD ]}| | ||< qt|jd ƒD ]}|| ||< q>d S )N)é   r   r   r   r   r   r   r   Úculocal1tuple   s
    r!   z'PTX inspection not available in cudasimc                   @   sŒ   e Zd Zdd„ Zdd„ Zdd„ Zdd„ Zed	ƒd
d„ ƒZed	ƒdd„ ƒZ	ed	ƒdd„ ƒZ
ed	ƒdd„ ƒZdd„ Zdd„ Zdd„ Zdd„ ZdS )ÚTestCudaLocalMemc                 C   sx   t d d … t d d … f}t |¡tƒ}|  d| |¡k¡ tjddd�}t |¡}|d ||ƒ |  t 	||k¡¡ d S )Nz.localr   r   r   ©r   r   )
r   r   Újitr   Ú
assertTrueZinspect_asmÚnpÚarangeÚ
zeros_likeÚall)ÚselfÚsigÚjculocalr   r   r   r   r   Útest_local_array$   s    
z!TestCudaLocalMem.test_local_arrayc                 C   sL   t  d¡tƒ}tjddd�}t |¡}|d ||ƒ |  t ||k¡¡ dS )zGEnsure that local arrays can be constructed with 1-tuple shape
        zvoid(int32[:], int32[:])r    r   r   r#   N)r   r$   r!   r&   r'   r(   r%   r)   )r*   r,   r   r   r   r   r   Útest_local_array_1_tuple-   s
    
z)TestCudaLocalMem.test_local_array_1_tuplec                 C   sX   d}t  |¡tƒ}tjddd�d d }t |¡}|d ||ƒ |  t ||k¡¡ d S )Nz"void(complex128[:], complex128[:])r   r   r   r   y               @r#   )r   r$   r   r&   r'   r(   r%   r)   )r*   r+   Zjculocalcomplexr   r   r   r   r   Útest_local_array_complex8   s    
z)TestCudaLocalMem.test_local_array_complexc                 C   s0   t t|j ¡ ƒƒj}|jd j}|  ||¡ d S )NÚl)ÚnextÚiterZ	overloadsÚvaluesZ_type_annotationÚtypemapr   ÚassertEqual)r*   Úfr   Ú
annotationZl_dtyper   r   r   Úcheck_dtype@   s    zTestCudaLocalMem.check_dtypezCan't check typing in simulatorc                 C   s0   t  ttd d d… ƒ¡dd„ ƒ}|  |t¡ d S )Nr   c                 S   s,   t jjdtd�}| d |d< |d | d< d S ©Né
   r   r   )r   r   r   r   ©Úxr0   r   r   r   r6   J   s    z,TestCudaLocalMem.test_numba_dtype.<locals>.f©r   r$   r   r   r8   ©r*   r6   r   r   r   Útest_numba_dtypeG   s    
z!TestCudaLocalMem.test_numba_dtypec                 C   s0   t  ttd d d… ƒ¡dd„ ƒ}|  |t¡ d S )Nr   c                 S   s.   t jjdtjd�}| d |d< |d | d< d S r9   )r   r   r   r&   r   r;   r   r   r   r6   U   s    z,TestCudaLocalMem.test_numpy_dtype.<locals>.fr=   r>   r   r   r   Útest_numpy_dtypeR   s    
z!TestCudaLocalMem.test_numpy_dtypec                 C   s0   t  ttd d d… ƒ¡dd„ ƒ}|  |t¡ d S )Nr   c                 S   s,   t jjddd�}| d |d< |d | d< d S )Nr:   r   r   r   ©r   r   r   r;   r   r   r   r6   `   s    z-TestCudaLocalMem.test_string_dtype.<locals>.fr=   r>   r   r   r   Útest_string_dtype]   s    
z"TestCudaLocalMem.test_string_dtypec              	   C   s@   d}|   t|¡�& t ttd d d… ƒ¡dd„ ƒ}W 5 Q R X d S )Nz*.*Invalid NumPy dtype specified: 'int33'.*r   c                 S   s,   t jjddd�}| d |d< |d | d< d S )Nr:   Zint33r   r   rA   r;   r   r   r   r6   m   s    z5TestCudaLocalMem.test_invalid_string_dtype.<locals>.f)ÚassertRaisesRegexr   r   r$   r   r   )r*   Úrer6   r   r   r   Útest_invalid_string_dtypeh   s    z*TestCudaLocalMem.test_invalid_string_dtypec                 C   s0   t  ttd d d… ƒ¡dd„ ƒ}|  |t¡ d S )Nr   c                 S   s,   t jjdtd�}| d |d< |d | d< d S r9   )r   r   r   r   r;   r   r   r   r6   t   s    z<TestCudaLocalMem.test_type_with_struct_data_model.<locals>.f)r   r$   r   r   r8   r>   r   r   r   Ú test_type_with_struct_data_models   s    
z1TestCudaLocalMem.test_type_with_struct_data_modelc                 C   sš   t  ttd d d… td d d… ƒ¡dd„ ƒ}tjddd�}tjddd�}|d ||ƒ t|ƒD ]\}}|  ||¡ q^t|ƒD ]\}}|  ||d ¡ q|d S )	Nr   c                 S   sr   t jjdtd�}tt|ƒƒD ]"}tt|ƒt|d ƒƒ}|||< qtt|ƒƒD ] }|| j| |< || j	||< qLd S )Nr:   r   é   )
r   r   r   r   r   Úlenr   r   r<   Úy)ZoutxZoutyÚarrr   Úobjr   r   r   r6   }   s    
z6TestCudaLocalMem.test_struct_model_type_arr.<locals>.f)r:   r   r   r#   rG   )r   r$   r   r   r&   r   Ú	enumerater5   )r*   r6   ZarrxZarryr   r<   rI   r   r   r   Útest_struct_model_type_arr|   s    "
z+TestCudaLocalMem.test_struct_model_type_arrc                    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 )Nr   r   )r   r   r   Úsize)ÚarJ   ©r   Útyr   r   Ús•   s    z8TestCudaLocalMem._check_local_array_size_fp16.<locals>.sr   r   r#   r   )r   r$   r&   ZzerosÚfloat16r5   )r*   r   ÚexpectedrQ   rR   Úresultr   rP   r   Ú_check_local_array_size_fp16”   s
    z-TestCudaLocalMem._check_local_array_size_fp16c                 C   s$   |   ddtj¡ |   ddtj¡ d S )NrG   )rV   r   rS   r&   )r*   r   r   r   Útest_issue_fp16_supportž   s    z(TestCudaLocalMem.test_issue_fp16_supportN)Ú__name__Ú
__module__Ú__qualname__r-   r.   r/   r8   r
   r?   r@   rB   rE   rF   rM   rV   rW   r   r   r   r   r"   "   s    	







	
r"   Ú__main__)Únumpyr&   Znumbar   r   r   r   Z
numba.corer   Znumba.core.errorsr   Znumba.cuda.testingr   r	   r
   Zextensions_usecasesr   r   r   r   r!   r"   rX   Úmainr   r   r   r   Ú<module>   s    