U
    -e.                  	   @   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re'  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&   f/var/www/html/Darija-Ai-Train/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
