§
    ƒÛg.  ã                   ó  — 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      ú[/granary/edev/quantum/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_constmem.pyÚcuconstEmptyr)   "   s7   € ÝŒ
×Ò�kÑ*Ô*€AÝŒ	�!‰Œ€AÝˆq‰6Œ6€A€a�D€D€Dó    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ð ˆQŒ4�#‰:€A€a�D€D€Dr*   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ÝŒ9�Q‰<Œ<�D€A€qØ��1�Œg€A€aˆ€d�G€G€Gr*   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Ø�1�a˜�7”€A€aˆˆA€g�J€J€Jr*   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Ô1€AÝŒ	�!‰Œ€AÝˆq‰6Œ6€A€a�D€D€Dr*   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ØˆQŒ4�Œ9€A€a�DØˆQŒ4�Œ9€A€a�D€D€Dr*   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Ô1€AÝŒ	�!‰Œ€AØˆQŒ4�Œ9€A€a�DØˆQŒ4�Œ9€A€a�DØˆQŒ4�Œ9€A€a�DØˆQŒ4�Œ9€A€a�DØˆQŒ4�Œ9€A€a�D€D€Dr*   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ØˆQŒ4�!�A”$‰;€A€a�D€D€Dr*   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¹   € Ý�q�q�qŒzˆmˆØ •4”8˜C‘=”=¥Ñ)Ô)ˆÝŒM�'Ñ"Ô"ˆØˆ�Œ�qÑÔÐØ�Š�œ˜q¥G¨a¡KÒ/Ñ0Ô0Ñ1Ô1Ð1åð 	;Ø�MŠMØØ×$Ò$ 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µ<Ñ@Ô@ˆÝŒG�A "­B¬HÐ5Ñ5Ô5ˆØˆ�dÔ˜AÑÔÐØ�Š�œ˜q Ašv™œÑ'Ô'Ð'Ð'Ð'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ˆÝŒG�A¥"¤&µÐ6Ñ6Ô6ˆØˆ�dÔ˜AÑÔÐØ�Š�œ˜q¥[µ7¸2¸A¸2´;Ñ%>Ò?Ñ@Ô@Ñ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Â   € Ý�Q�Q�Q�q�q�q�SŒzˆmˆØ"•T”X˜c‘]”]¥9Ñ-Ô-ˆ
ÝŒM�'¨Ð-Ñ-Ô-ˆØ"ˆ
�>Ô" 1Ñ%Ô%Ð%Ø�Š�œ˜q¥Gš|Ñ,Ô,Ñ-Ô-Ð-åð 	)Ø�MŠMØØ×&Ò& 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Ì   € Ý˜˜˜˜1˜1˜1˜Q˜Q˜Q˜ÔÐ!ˆØ"•T”X˜c‘]”]¥9Ñ-Ô-ˆ
ÝŒM�'¨Ð-Ñ-Ô-ˆØ ˆ
�<Ô  Ñ#Ô#Ð#Ø�Š�œ˜q¥Gš|Ñ,Ô,Ñ-Ô-Ð-åð 	:Ø×(Ò(¨Ñ-Ô-ˆCØ,ˆLØBˆKØ�MŠM˜,¨¨[Ñ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   € Ø5�4œ8Ð$4Ñ5Ô5µoÑFÔFÐÝŒG�A "­B¬HÐ5Ñ5Ô5ˆØÐ˜Ô˜qÑ!Ô!Ð!Ø�Š�œ˜q Ašv™œÑ'Ô'Ð'Ð'Ð'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Ÿ   € ÝŒH�Q�eÐ$Ñ$Ô$ˆÝŒH�Q�cÐ"Ñ"Ô"ˆÝ”8�JÑ'Ô'×2Ò2°1°aÑ8Ô8ˆàˆ�Œ�q˜!ÑÔÐÝ
Œ
×"Ò" 1¥l°3Ô&7Ñ8Ô8Ð8Ý
Œ
×"Ò" 1¥l°3Ô&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  € ÝŒH�Q�bœjÐ)Ñ)Ô)ˆÝŒH�Q�bœjÐ)Ñ)Ô)ˆÝŒH�Q�bœjÐ)Ñ)Ô)ˆÝŒH�Q�bœjÐ)Ñ)Ô)ˆÝŒH�Q�bœjÐ)Ñ)Ô)ˆÝ”8�OÑ,Ô,×7Ò7¸¸1¸aÀÀ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Ø ,Ð ,Ð ,Ð ,Ð ,Ð ,àˆbŒh�r‰lŒl€Ø
ˆ"Œ)�B˜bœjÐ
)Ñ
)Ô
)¨BÑ
.€Ø
ˆ"Ô
Ø€B„Iˆc˜œÐ"Ñ"Ô"×*Ò*¨2¨rÑ2Ô2ñ4ô 4€àˆBŒI�i r¤|Ð4Ñ4Ô4×<Ò<¸QÀÀ1ÑEÔEÈÑJØñ€àˆbŒi˜ ¤Ð*Ñ*Ô*€à�R”XØØ�ˆ<˜#˜s˜Ð
$ð&ñ &ô &Ð ð ˆrŒxØˆxÐØ�ˆ<˜#˜s˜Ð
$ð&ñ &ô &€ð �R”XØÐ8Ð9Ø
ˆ"Œ(à�"”(ˆOØ�"”(ˆOØ�"”(ˆOØ�"”)ÐØ�"”(ˆOð
ð ðñ ô ð
ñ 
ô 
Ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðP?ð P?ð P?ð P?ð P?˜\ñ P?ô P?ð P?ðf ˆzÒÐØ€H„M�O„O€O€O€Oð Ðr*   