
    xKg.                     z   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
                        j#                  dd            Z ej                  d	ej                        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j4                  dej(                  fdej(                  fdej(                  fdej6                  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jN                          yy)$    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                  j                  t              }t        j                  d      }t        |      | |<   y Nr   )r   const
array_likeCONST_EMPTYgridlenACis      i/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_constmem.pycuconstEmptyr)   "   s0    

k*A		!Aq6AaD    c                     t         j                  j                  t              }t        j                  d      }||   dz   | |<   y )Nr   r   )r   r   r    CONST1Dr"   r$   s      r(   cuconstr-   (   s6    

g&A		!A Q4#:AaDr*   c                     t         j                  j                  t              }t        j                  d      \  }}|||f   | ||f<   y )Nr   )r   r   r    CONST2Dr"   )r%   r&   r'   js       r(   	cuconst2dr1   0   s<    

g&A99Q<DAq1gAadGr*   c                    t         j                  j                  t              }t         j                  j
                  }t         j                  j                  }t         j                  j                  }||||f   | |||f<   y )N)r   r   r    CONST3D	threadIdxr   r   r   )r%   r&   r'   r0   ks        r(   	cuconst3dr6   6   s\    

g&AAAA1a7AaAgJr*   c                     t         j                  j                  t              }t        j                  d      }t        |      | |<   y r   )r   r   r    CONST_RECORD_EMPTYr"   r#   r$   s      r(   cuconstRecEmptyr9   >   s1    

01A		!Aq6AaDr*   c                     t         j                  j                  t              }t        j                  d      }||   d   | |<   ||   d   ||<   y )Nr   r   r   )r   r   r    CONST_RECORDr"   )r%   Br&   r'   s       r(   
cuconstRecr=   D   sE    

l+A		!AQ49AaDQ49AaDr*   c                     t         j                  j                  t              }t        j                  d      }||   d   | |<   ||   d   ||<   ||   d   ||<   ||   d   ||<   ||   d   ||<   y )Nr   r   r   r   r   r   )r   r   r    CONST_RECORD_ALIGNr"   )r%   r<   r&   DEZr'   s          r(   cuconstRecAlignrC   K   sy    

01A		!AQ49AaDQ49AaDQ49AaDQ49AaDQ49AaDr*   c                     t         j                  j                  t              }t         j                  j                  t              }t        j
                  d      }||   ||   z   | |<   y r   )r   r   r    CONST3BYTESr,   r"   )r   r   r   r'   s       r(   cuconstAlignrF   U   sL    

k*A

g&A		!AQ4!A$;AaDr*   c                   <    e Zd Zd Zd Zd Zd Zd Zd Zd Z	d Z
y	)
TestCudaConstantMemoryc                 L   t         d d  f} t        j                  |      t              }t	        j
                  t              } |d   |       | j                  t	        j                  |t        dz   k(               t        s#| j                  d|j                  |      d       y y )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    qzm 488C=)MM'"qqGaK/01MM$$S)9; r*   c                      t        j                  d      t              }t        j                  ddt        j
                        } |d   |       | j                  t        j                  |dk(               y 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   sU    2!12<@GGA"BHH5dAqAv'r*   c           	          t        j                  d      t              }t        j                  dt        j
                  t              } |d   |       | j                  t        j                  |t        t        d d z   k(               y )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   sb    4!34\BGGA"&&6dAq[72A;%>?@Ar*   c                 V   t         d d d d f   f} t        j                  |      t              }t	        j
                  t        d      } |d   |       | j                  t	        j                  |t        k(               t        s#| j                  d|j                  |      d       y y )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    QqSzm"TXXc]9-
MM'-"
>"1%qG|,-MM&&s+') r*   c                 h   t         d d d d d d f   f} t        j                  |      t              }t	        j
                  t        d      } |d   |       | j                  t	        j                  |t        k(               t        s)|j                  |      }d}d}| j                  |||       y y )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    1Q!"TXXc]9-
MM'- 
< #qG|,-((-C,LBKMM,[9	 r*   c                      t        j                  d      t              }t        j                  ddt        j
                        } |d   |       | j                  t        j                  |dk(               y rV   )r   rJ   r9   rK   rZ   r[   rM   rN   )rQ   jcuconstRecEmptyr%   s      r(   test_const_record_emptyz.TestCudaConstantMemory.test_const_record_empty   sW    5488$45oFGGA"BHH5q!qAv'r*   c                 v   t        j                  dt              }t        j                  dt              }t	        j
                  t              j                  ||      } |d   ||       t         j                  j                  |t        d          t         j                  j                  |t        d          y )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    HHQe$HHQc"88J'221a8q!


""1l3&78


""1l3&78r*   c                 r   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
              j                  |||||      } |d   |||||       t         j                  j                  |t        d          t         j                  j                  |t        d          t         j                  j                  |t        d          t         j                  j                  |t        d          t         j                  j                  |t        d          y )	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   s   HHQbjj)HHQbjj)HHQbjj)HHQbjj)HHQbjj)88O,771aA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):(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 5 ,bhhrl
"))Bbjj
)B
.
"

BIIc"**2r24BIIir||4<<Q1EJbii*RXX<#s
$&  rxxx<#s
$& RXX89
"(("((O"((O"((O"))"((O
 
 P?\ P?f zHMMO r*   