U
    ,‰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   Zalignc                 C   s&   t j t¡}t  d¡}t|ƒ| |< d S ©Nr   )r   ÚconstÚ
array_likeÚCONST_EMPTYÚgridÚlen©ÚAÚCÚi© r%   úI/tmp/pip-unpacked-wheel-eu7e0c37/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
    r/   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    
r1   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    
r4   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"   r3   r#   ÚDÚEÚZr$   r%   r%   r&   ÚcuconstRecAlignK   s    
r9   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Úint64r@   rA   )rD   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   rJ   )r   r   )r   r=   r;   r>   rK   ÚnanÚfloatr@   rA   r:   r(   )rD   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*   r@   rA   r   rB   rC   )rD   rE   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	stj
 ¡ dkrtd}d}nd}d}|  || |¡|¡ d S )	NÚFrQ   )r   )r   r   r   )é   r   zld.const.v2.f32z&Load the complex as a vector of 2x f32zld.const.f32z$load each half of the complex as f32)r   r   r=   r/   r>   r?   r-   r@   rA   r   ZruntimeÚget_versionrB   rC   )rD   rE   Z
jcuconst3dr"   Z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 rH   )r   r=   r1   r>   rK   rL   r@   rA   )rD   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>   ÚzerosrO   Úintr   r=   r4   Ú
specializeÚtestingÚassert_allcloser2   )rD   r"   r3   rF   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=   r9   r]   r^   r_   r5   )rD   r"   r3   r#   r6   r7   rF   r%   r%   r&   Útest_const_record_align¨   s    z.TestCudaConstantMemory.test_const_record_alignN)Ú__name__Ú
__module__Ú__qualname__rG   rM   rP   rS   rX   rY   r`   ra   r%   r%   r%   r&   r<   \   s   	r<   Ú__main__)(Znumpyr>   Znumbar   r   r   r   Znumba.cuda.testingr   r   Znumba.core.configr   Úarrayr   Zaranger(   ZasfortranarrayZreshaper*   r-   Zuint8r:   rO   r\   r0   r2   r   Zuint32r5   r'   r)   r,   r/   r1   r4   r9   r;   r<   rb   Úmainr%   r%   r%   r&   Ú<module>   sV   
ÿÿþþûøþ
\
