U
    hâËd.  ã                	   @   sŒ  d dl Zd dlmZmZmZmZ d dlmZm	Z	 d dl
mZ e g ¡Zejdejd�d Ze ejdejd� dd¡¡Zejd	ejd� d
d
d
¡d d Zejdejd�Zejg defdefgd�Zejddgdefdefgd�Zejddgejdejfdejfdejfdejfdejfgd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$G d)d*„ d*e	ƒZ%e&d+k�rˆe '¡  dS ),é    N)ÚcudaÚ	complex64Úint32Úfloat64)ÚunittestÚCUDATestCase)ÚENABLE_CUDASIMé
   ©Údtypeg       @éd   é}   é   y              ð?y               @é   ÚxÚy)ç      ð?é   )g      @é   )é   r   r   l   ï>[= é   )r   r   é   l   ­^ß} r	   ÚaÚbÚzT)r   Úalignc                 C   s&   t j t¡}t  d¡}t|ƒ| |< d S ©Nr   )r   ÚconstÚ
array_likeÚCONST_EMPTYÚgridÚlen©ÚAÚCÚi© r&   ú^/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/tests/cudapy/test_constmem.pyÚcuconstEmpty"   s    
r(   c                 C   s*   t j t¡}t  d¡}|| d | |< d S )Nr   r   )r   r   r   ÚCONST1Dr    r"   r&   r&   r'   Úcuconst(   s    
r*   c                 C   s2   t j t¡}t  d¡\}}|||f | ||f< d S )Nr   )r   r   r   ÚCONST2Dr    )r#   r$   r%   Újr&   r&   r'   Ú	cuconst2d0   s    r-   c                 C   s@   t j t¡}t jj}t jj}t jj}||||f | |||f< d S )N)r   r   r   ÚCONST3DZ	threadIdxr   r   r   )r#   r$   r%   r,   Úkr&   r&   r'   Ú	cuconst3d6   s
    r0   c                 C   s&   t j t¡}t  d¡}t|ƒ| |< d S r   )r   r   r   ÚCONST_RECORD_EMPTYr    r!   r"   r&   r&   r'   ÚcuconstRecEmpty>   s    
r2   c                 C   s:   t j t¡}t  d¡}|| d | |< || d ||< d S )Nr   r   r   )r   r   r   ÚCONST_RECORDr    )r#   ÚBr$   r%   r&   r&   r'   Ú
cuconstRecD   s    
r5   c                 C   sj   t j t¡}t  d¡}|| d | |< || d ||< || d ||< || d ||< || d ||< d S )Nr   r   r   r   r   r   )r   r   r   ÚCONST_RECORD_ALIGNr    )r#   r4   r$   ÚDÚEÚZr%   r&   r&   r'   ÚcuconstRecAlignK   s    
r:   c                 C   s:   t j t¡}t j t¡}t  d¡}|| ||  | |< d S r   )r   r   r   ÚCONST3BYTESr)   r    )r   r   r   r%   r&   r&   r'   ÚcuconstAlignU   s    
r<   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 )ÚTestCudaConstantMemoryc                 C   sf   t d d … f}t |¡tƒ}t t¡}|d |ƒ |  t |td k¡¡ t	sb|  
d| |¡d¡ d S )N)r   r   r   zld.const.f64z'as we're adding to it, load as a double)r   r   Újitr*   ÚnpÚ
zeros_liker)   Ú
assertTrueÚallr   ÚassertInÚinspect_asm)ÚselfÚsigÚjcuconstr#   r&   r&   r'   Útest_const_array]   s    
ýz'TestCudaConstantMemory.test_const_arrayc                 C   sD   t  d¡tƒ}tjddtjd�}|d |ƒ |  t |dk¡¡ d S ©Nzvoid(int64[:])r   éÿÿÿÿ©Z
fill_valuer   )r   r   r   )r   r>   r(   r?   ÚfullÚint64rA   rB   )rE   ZjcuconstEmptyr#   r&   r&   r'   Útest_const_emptyj   s    z'TestCudaConstantMemory.test_const_emptyc              	   C   sP   t  d¡tƒ}tjdtjtd�}|d |ƒ |  t |t	t
d d…  k¡¡ d S )Nzvoid(float64[:])r   rK   )r   r   )r   r>   r<   r?   rL   ÚnanÚfloatrA   rB   r;   r)   )rE   ZjcuconstAlignr#   r&   r&   r'   Útest_const_alignp   s    z'TestCudaConstantMemory.test_const_alignc                 C   sn   t d d …d d …f f}t |¡tƒ}tjtdd�}|d |ƒ |  t |tk¡¡ t	sj|  
d| |¡d¡ d S )Nr$   ©Úorder))r   r   )r   r   zld.const.u32zload the ints as ints)r   r   r>   r-   r?   r@   r+   rA   rB   r   rC   rD   )rE   rF   Z
jcuconst2dr#   r&   r&   r'   Útest_const_array_2dv   s    ýz*TestCudaConstantMemory.test_const_array_2dc                 C   s€   t d d …d d …d d …f f}t |¡tƒ}tjtdd�}|d |ƒ |  t |tk¡¡ t	s|| 
|¡}d}d}|  |||¡ d S )NÚFrR   )r   )r   r   r   zld.const.v2.f32z&Load the complex as a vector of 2x f32)r   r   r>   r0   r?   r@   r.   rA   rB   r   rD   rC   )rE   rF   Z
jcuconst3dr#   ÚasmZcomplex_loadÚdescriptionr&   r&   r'   Útest_const_array_3dƒ   s    
z*TestCudaConstantMemory.test_const_array_3dc                 C   sD   t  d¡tƒ}tjddtjd�}|d |ƒ |  t |dk¡¡ d S rI   )r   r>   r2   r?   rL   rM   rA   rB   )rE   ZjcuconstRecEmptyr#   r&   r&   r'   Útest_const_record_empty�   s    z.TestCudaConstantMemory.test_const_record_emptyc                 C   sd   t jdtd�}t jdtd�}t t¡ ||¡}|d ||ƒ t j 	|t
d ¡ t j 	|t
d ¡ d S )Nr   r
   ©r   r   r   r   )r?   ÚzerosrP   Úintr   r>   r5   Ú
specializeÚtestingÚassert_allcloser3   )rE   r#   r4   rG   r&   r&   r'   Útest_const_record–   s    z(TestCudaConstantMemory.test_const_recordc                 C   sÚ   t jdt jd�}t jdt jd�}t jdt jd�}t jdt jd�}t jdt jd�}t t¡ |||||¡}|d |||||ƒ t j |t	d ¡ t j |t	d ¡ t j |t	d ¡ t j |t	d ¡ t j |t	d ¡ d S )	Nr   r
   rZ   r   r   r   r   r   )
r?   r[   r   r   r>   r:   r]   r^   r_   r6   )rE   r#   r4   r$   r7   r8   rG   r&   r&   r'   Útest_const_record_alignŸ   s    z.TestCudaConstantMemory.test_const_record_alignN)Ú__name__Ú
__module__Ú__qualname__rH   rN   rQ   rT   rX   rY   r`   ra   r&   r&   r&   r'   r=   \   s   	r=   Ú__main__)(Únumpyr?   Znumbar   r   r   r   Znumba.cuda.testingr   r   Znumba.core.configr   Úarrayr   Zaranger)   ZasfortranarrayZreshaper+   r.   Zuint8r;   rP   r\   r1   r3   r   Zuint32r6   r(   r*   r-   r0   r2   r5   r:   r<   r=   rb   Úmainr&   r&   r&   r'   Ú<module>   sV   
ÿÿþþûøþ
S
