
    J/Ph.                        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j        g           Z ej        dej                  dz  Z ej         ej        dej                                      dd                    Z ej        d	ej                                      d
d
d
          dz   dz  Z ej        dej                  Z ej        g defdefg          Z ej        ddgdefdefg          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                    Zd Zd Zd Zd Z d Z!d Z"d Z#d  Z$ G d! d"e	          Z%e&d#k    r ej'                     dS 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                     t           j                            t                    }t          j        d          }t          |          | |<   d S Nr   )r   const
array_likeCONST_EMPTYgridlenACis      e/var/www/html/test/jupyter/venv/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_constmem.pycuconstEmptyr)   "   s7    
k**A	!Aq66AaDDD    c                     t           j                            t                    }t          j        d          }||         dz   | |<   d S )Nr   r   )r   r   r    CONST1Dr"   r$   s      r(   cuconstr-   (   s;    
g&&A	!A Q4#:AaDDDr*   c                     t           j                            t                    }t          j        d          \  }}|||f         | ||f<   d S )Nr   )r   r   r    CONST2Dr"   )r%   r&   r'   js       r(   	cuconst2dr1   0   sA    
g&&A9Q<<DAq1gAadGGGr*   c                     t           j                            t                    }t           j        j        }t           j        j        }t           j        j        }||||f         | |||f<   d S )N)r   r   r    CONST3D	threadIdxr   r   r   )r%   r&   r'   r0   ks        r(   	cuconst3dr6   6   sR    
g&&AAAA1a7AaAgJJJr*   c                     t           j                            t                    }t          j        d          }t          |          | |<   d S r   )r   r   r    CONST_RECORD_EMPTYr"   r#   r$   s      r(   cuconstRecEmptyr9   >   s8    
011A	!Aq66AaDDDr*   c                     t           j                            t                    }t          j        d          }||         d         | |<   ||         d         ||<   d S )Nr   r   r   )r   r   r    CONST_RECORDr"   )r%   Br&   r'   s       r(   
cuconstRecr=   D   sJ    
l++A	!AQ49AaDQ49AaDDDr*   c                     t           j                            t                    }t          j        d          }||         d         | |<   ||         d         ||<   ||         d         ||<   ||         d         ||<   ||         d         ||<   d S )Nr   r   r   r   r   r   )r   r   r    CONST_RECORD_ALIGNr"   )r%   r<   r&   DEZr'   s          r(   cuconstRecAlignrC   K   s~    
011A	!AQ49AaDQ49AaDQ49AaDQ49AaDQ49AaDDDr*   c                     t           j                            t                    }t           j                            t                    }t          j        d          }||         ||         z   | |<   d S r   )r   r   r    CONST3BYTESr,   r"   )r   r   r   r'   s       r(   cuconstAlignrF   U   sR    
k**A
g&&A	!AQ4!A$;AaDDDr*   c                   >    e Zd Zd Zd Zd Zd Zd Zd Zd Z	d Z
d	S )
TestCudaConstantMemoryc                    t           d d          f} t          j        |          t                    }t	          j        t                    } |d         |           |                     t	          j        |t          dz   k                         t          s,| 
                    d|                    |          d           d S 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%   s       r(   test_const_arrayz'TestCudaConstantMemory.test_const_array]   s    qqqzm 48C==))M'""qqGaK/00111 	;MM$$S))9; ; ; ; ;	; 	;r*   c                     t          j        d          t                    }t          j        ddt          j                  } |d         |           |                     t          j        |dk                         d S Nzvoid(int64[:])r   
fill_valuer   )r   r   r   )r   rJ   r)   rK   fullint64rM   rN   )rQ   jcuconstEmptyr%   s      r(   test_const_emptyz'TestCudaConstantMemory.test_const_emptyj   sm    2!122<@@GA"BH555dAqAv'''''r*   c           	      8    t          j        d          t                    }t          j        dt          j        t                    } |d         |           |                     t          j        |t          t          d d         z   k                         d S )Nzvoid(float64[:])r   rX   )r   r   )r   rJ   rF   rK   rZ   nanfloatrM   rN   rE   r,   )rQ   jcuconstAlignr%   s      r(   test_const_alignz'TestCudaConstantMemory.test_const_alignp   s}    4!344\BBGA"&666dAq[72A2;%>?@@AAAAAr*   c                    t           d d d d f         f} t          j        |          t                    }t	          j        t          d          } |d         |           |                     t	          j        |t          k                         t          s,| 
                    d|                    |          d           d S d S )Nr&   order))r   r   )r   r   zld.const.u32zload the ints as ints)r   r   rJ   r1   rK   rL   r/   rM   rN   r	   rO   rP   )rQ   rR   
jcuconst2dr%   s       r(   test_const_array_2dz*TestCudaConstantMemory.test_const_array_2dv   s    QQQqqqSzm"TXc]]9--
M'---"
>"1%%%qG|,,--- 	)MM&&s++') ) ) ) )	) 	)r*   c                    t           d d d d d d f         f} t          j        |          t                    }t	          j        t          d          } |d         |           |                     t	          j        |t          k                         t          s2|
                    |          }d}d}|                     |||           d S d S )NFrd   )r   )r   r   r   zld.const.v2.f32z&Load the complex as a vector of 2x f32)r   r   rJ   r6   rK   rL   r3   rM   rN   r	   rP   rO   )rQ   rR   
jcuconst3dr%   asmcomplex_loaddescriptions          r(   test_const_array_3dz*TestCudaConstantMemory.test_const_array_3d   s    111QQQ!"TXc]]9--
M'--- 
< ###qG|,,--- 	:((--C,LBKMM,[99999		: 	:r*   c                     t          j        d          t                    }t          j        ddt          j                  } |d         |           |                     t          j        |dk                         d S rV   )r   rJ   r9   rK   rZ   r[   rM   rN   )rQ   jcuconstRecEmptyr%   s      r(   test_const_record_emptyz.TestCudaConstantMemory.test_const_record_empty   so    548$455oFFGA"BH555q!!!qAv'''''r*   c                    t          j        dt                    }t          j        dt                    }t	          j        t                                        ||          } |d         ||           t           j        	                    |t          d                    t           j        	                    |t          d                    d S )Nr   r   r   r   r   r   )rK   zerosr`   intr   rJ   r=   
specializetestingassert_allcloser;   )rQ   r%   r<   rS   s       r(   test_const_recordz(TestCudaConstantMemory.test_const_record   s    HQe$$$HQc"""8J''221a88q!

""1l3&7888

""1l3&788888r*   c                 ~   t          j        dt           j                  }t          j        dt           j                  }t          j        dt           j                  }t          j        dt           j                  }t          j        dt           j                  }t          j        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   rs   r   r   r   r   r   )
rK   rt   r   r   rJ   rC   rv   rw   rx   r?   )rQ   r%   r<   r&   r@   rA   rS   s          r(   test_const_record_alignz.TestCudaConstantMemory.test_const_record_align   sM   HQbj)))HQbj)))HQbj)))HQbj)))HQbj)))8O,,771aAFFq!Q1%%%

""1&8&=>>>

""1&8&=>>>

""1&8&=>>>

""1&8&=>>>

""1&8&=>>>>>r*   N)__name__
__module____qualname__rT   r]   rb   rg   rn   rq   ry   r{    r*   r(   rH   rH   \   s        ; ; ;( ( (B B B) ) ): : :( ( (9 9 9? ? ? ? ?r*   rH   __main__)(numpyrK   numbar   r   r   r   numba.cuda.testingr   r   numba.core.configr	   arrayr!   aranger,   asfortranarrayreshaper/   r3   uint8rE   r`   ru   r8   r;   r   uint32r?   r)   r-   r1   r6   r9   r=   rC   rF   rH   r|   mainr   r*   r(   <module>r      s       1 1 1 1 1 1 1 1 1 1 1 1 5 5 5 5 5 5 5 5 , , , , , ,bhrll
")Bbj
)
)
)B
.
"
BIc"""**2r224 4BIir|444<<Q1EEJbi***RX<#s
$& & &  rxx<#s
$& & & RX89
"("(O"(O"(O")"(O
   
 
 
                 P? P? P? P? P?\ P? P? P?f zHMOOOOO r*   