a
    
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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   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%   n/var/www/html/stable-diffusion-webui/venv/lib/python3.9/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
    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	s||
|}d}d}| ||| d S )NFrQ   )r   )r   r   r   zld.const.v2.f32z&Load the complex as a vector of 2x f32)r   r   r=   r/   r>   r?   r-   r@   rA   r   rC   rB   )rD   rE   Z
jcuconst3dr"   Z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 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
   rX   r   r   r   r   r   )
r>   rY   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   rV   rW   r^   r_   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:   rO   rZ   r0   r2   r   Zuint32r5   r'   r)   r,   r/   r1   r4   r9   r;   r<   r`   mainr%   r%   r%   r&   <module>   sV   

S
