§
    1@Äj´  ã                   ó2  — d dl mZ d dlmZ d dlZd dlZ G d„ d¦  «        Z G d„ de¦  «        Z ej	        d¬	¦  «        d
„ ¦   «         Z
 ej	        d¬	¦  «        d„ ¦   «         Z ee
¦  «        Z ee¦  «        ZdZ ej	        d¬	¦  «        d„ ¦   «         Z ej	        d¬	¦  «        d„ ¦   «         Z ej	        d¬	¦  «        d„ ¦   «         Z ee¦  «        Z ee¦  «        Z ej        dej        fdej        fg¦  «        Z ej        dej        fdej        fgd¬¦  «        Z ej        de¬¦  «        Z eej        ¦  «        D ]Zedz   ee         d<   edz   ee         d<   Œ ej        ee¬¦  «        Z ej	        d¬	¦  «        d„ ¦   «         Z  ee e¬¦  «        Z! ee e¬¦  «        Z"d„ Z# e#d¦  «        Z$ e#d¦  «        Z% e#d¦  «        Z& e#d¦  «        Z' ej	        d¬	¦  «        d„ ¦   «         Z( ee(¦  «        Z) ej	        d¬	¦  «        d„ ¦   «         Z( ee(¦  «        Z* ej	        d¬	¦  «        d „ ¦   «         Z+ ej	        d¬	¦  «        d!„ ¦   «         Z, ee,¦  «        Z- ej	        d¬	¦  «        d"„ ¦   «         Z. ee.¦  «        Z/ G d#„ d$e¦  «        Z0d%„ Z1dS )&é    )Úcuda)ÚCUDATestCaseNc                   ó6   — e Zd ZdZdd„Zd„ Zed„ ¦   «         ZdS )ÚUseCasea2  
    Provide a way to call a kernel as if it were a function.

    This allows the CUDA cache tests to closely match the CPU cache tests, and
    also to support calling cache use cases as njitted functions. The class
    wraps a function that takes an array for the return value and arguments,
    and provides an interface that accepts arguments, launches the kernel
    appropriately, and returns the stored return value.

    The return type is inferred from the type of the first argument, unless it
    is explicitly overridden by the ``retty`` kwarg.
    Nc                 ó"   — || _         || _        d S ©N)Ú_funcÚ_retty)ÚselfÚfuncÚrettys      úWC:\Projects\Historical\venv\Lib\site-packages\numba/cuda/tests/cudapy/cache_usecases.pyÚ__init__zUseCase.__init__   s   € ØˆŒ
ØˆŒˆˆó    c                 óÀ   — d„ |D ¦   «         }| j         rt          j        d| j         ¬¦  «        }nt          j        |d         ¦  «        } | j        |g|¢R Ž  |d         S )Nc                 ó6   — g | ]}t          j        |¦  «        ‘ŒS © )ÚnpÚasarray)Ú.0Úargs     r   ú
<listcomp>z$UseCase.__call__.<locals>.<listcomp>   s    € Ð6Ð6Ð6¨#•b”j ‘o”oÐ6Ð6Ð6r   r   ©Údtyper   )r
   r   ÚndarrayÚ
zeros_likeÚ_call)r   ÚargsÚ
array_argsÚarray_returns       r   Ú__call__zUseCase.__call__   so   € Ø6Ð6°Ð6Ñ6Ô6ˆ
ØŒ;ð 	8Ýœ: b°´Ð<Ñ<Ô<ˆLˆLåœ=¨°A¬Ñ7Ô7ˆLàˆŒ
�<Ð- *Ð-Ð-Ð-Ð-Ø˜BÔÐr   c                 ó   — | j         S r   ©r	   )r   s    r   r   zUseCase.func"   s
   € àŒzÐr   r   )Ú__name__Ú
__module__Ú__qualname__Ú__doc__r   r!   Úpropertyr   r   r   r   r   r      s\   € € € € € ðð ðð ð ð ð ð  ð  ð ðð ñ „Xðð ð r   r   c                   ó   — e Zd Zd„ ZdS )ÚCUDAUseCasec                 ó.   —  | j         d         |g|¢R Ž  d S )N©é   r-   r#   )r   Úretr   s      r   r   zCUDAUseCase._call(   s'   € ØˆŒ
�4Ô˜Ð$˜tÐ$Ð$Ð$Ð$Ð$Ð$r   N)r$   r%   r&   r   r   r   r   r*   r*   '   s#   € € € € € ð%ð %ð %ð %ð %r   r*   T©Úcachec                 ó>   — |d         |d         z   t           z   | d<   d S ©Nr   ©ÚZ©ÚrÚxÚys      r   Úadd_usecase_kernelr9   ,   ó    € àˆbŒE�A�b”E‰M�AÑ€A€b�E€E€Er   Fc                 ó>   — |d         |d         z   t           z   | d<   d S r2   r3   r5   s      r   Úadd_nocache_usecase_kernelr<   1   r:   r   r-   c                 ó   — | |z   t           z   S r   r3   )r7   r8   s     r   Úinnerr>   >   s   € àˆq‰5•1‰9Ðr   c                 óF   — t          |d          |d         ¦  «        | d<   d S r2   ©r>   r5   s      r   Úouter_kernelrA   C   ó#   € å�1�R”5�&˜!˜Bœ%Ñ Ô €A€b�E€E€Er   c                 óF   — t          |d          |d         ¦  «        | d<   d S r2   r@   r5   s      r   Úouter_uncached_kernelrD   H   rB   r   ÚaÚb)Úaligné   r   g     @E@c                 ó   — ||         | d<   d S r2   r   )r6   ÚaryÚis      r   Úrecord_returnrL   _   s   € à�ŒF€A€b�E€E€Er   )r   c                 ób   ‡ — t          j        d¬¦  «        ˆ fd„¦   «         }t          |¦  «        S )NTr/   c                 ó$   •— ‰|d         z   | d<   d S r2   r   )r6   r8   r7   s     €r   Úclosurezmake_closure.<locals>.closurek   s   ø€ à�A�b”E‘	ˆˆ"‰ˆˆr   )r   Újitr*   )r7   rO   s   ` r   Úmake_closurerQ   j   sD   ø€ Ý	„X�DÐÑÔðð ð ð ñ Ôðõ �wÑÔÐr   é   é   é   é	   c                 ó"   — |d         dz   | d<   d S )Nr   rH   r   ©r6   r7   s     r   Úambiguous_functionrX   z   ó   € àˆbŒE�A‰I€A€b�E€E€Er   c                 ó"   — |d         dz   | d<   d S )Nr   é   r   rW   s     r   rX   rX   ‚   rY   r   c                  ó6	  — 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                             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                             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                             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                             dt          j        ¦  «        }t           j                             dt          j        ¦  «        }t           j                             dt          j        ¦  «        }t           j                             dt          j        ¦  «        }d| d d …<   d|d d …<   d|d d …<   d|d d …<   d|d d …<   d|d d …<   d|d d …<   d|d d …<   d|d d …<   d|	d d …<   d|
d d …<   d|d d …<   d|d d …<   d|d d …<   d|d d …<   d|d d …<   d|d d …<   d|d d …<   d|d d …<   d|d d …<   d|d d …<   d|d d …<   d|d d …<   d|d d …<   d S )Nr,   r   )r   ÚlocalÚarrayr   Úfloat64)ÚaaÚabÚacÚadÚaeÚafÚagÚahÚaiÚajÚakÚalÚamÚanÚaoÚapÚarÚatÚauÚavÚawÚaxÚayÚazs                           r   Úmany_localsrx   Š   s–  € å	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€BÝ	Œ×	Ò	˜&¥"¤*Ñ	-Ô	-€Bà€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�EØ€B€q€q€q�E€E€Er   c                 ó   — |d         | d<   d S r2   r   rW   s     r   Úsimple_usecase_kernelrz   Á   s   € àˆbŒE€A€b�E€E€Er   c                 ój   — t           j                             ¦   «         }|                     ¦   «          d S r   )r   ÚcgÚ	this_gridÚsync)r6   r7   Úgrids      r   Úcg_usecase_kernelr€   Ë   s&   € åŒ7×ÒÑÔ€DØ‡I‚I�K„K€K€K€Kr   c                   ó   — e Zd ZdZd„ ZdS )Ú_TestModulez¥
    Tests for functionality of this module's functions.
    Note this does not define any "test_*" method, instead check_module()
    should be called by hand.
    c                 ó$  — |                       |                     dd¦  «        d¦  «         |                       |                     dd¦  «        d¦  «         |                       |                     dd¦  «        d¦  «         |                     |j        d¦  «        }|                       t          |¦  «        d¦  «         |                     |j        d¦  «        }|                       t          |¦  «        d¦  «         | 	                    d¦  «         d S )NrH   rR   r[   r-   )rH   g     ÀE@)
ÚassertPreciseEqualÚadd_usecaseÚouter_uncachedÚouterÚrecord_return_packedÚ
packed_arrÚtupleÚrecord_return_alignedÚaligned_arrÚsimple_usecase_caller)r   ÚmodÚ
packed_recÚaligned_recs       r   Úcheck_modulez_TestModule.check_moduleÛ   sò   € Ø×Ò §¢°°1Ñ 5Ô 5°qÑ9Ô9Ð9Ø×Ò × 2Ò 2°1°aÑ 8Ô 8¸!Ñ<Ô<Ð<Ø×Ò §	¢	¨!¨Q¡¤°Ñ3Ô3Ð3à×-Ò-¨c¬n¸aÑ@Ô@ˆ
Ø×Ò¥ jÑ 1Ô 1°9Ñ=Ô=Ð=Ø×/Ò/°´ÀÑCÔCˆØ×Ò¥ kÑ 2Ô 2°IÑ>Ô>Ð>à×!Ò! !Ñ$Ô$Ð$Ð$Ð$r   N)r$   r%   r&   r'   r‘   r   r   r   r‚   r‚   Ô   s-   € € € € € ðð ð
%ð 
%ð 
%ð 
%ð 
%r   r‚   c                  óv   — t           j        t                   } t          ¦   «                              | ¦  «         d S r   )ÚsysÚmodulesr$   r‚   r‘   )rŽ   s    r   Ú	self_testr•   è   s,   € Ý
Œ+•hÔ
€CÝ�M„M×Ò˜sÑ#Ô#Ð#Ð#Ð#r   )2Únumbar   Únumba.cuda.testingr   Únumpyr   r“   r   r*   rP   r9   r<   r…   Úadd_nocache_usecaser4   r>   rA   rD   r‡   r†   r   Úint8r_   Úpacked_record_typeÚaligned_record_typeÚemptyr‰   ÚrangeÚsizerK   r^   rŒ   rL   rˆ   r‹   rQ   Úclosure1Úclosure2Úclosure3Úclosure4rX   Úrenamed_function1Úrenamed_function2rx   rz   r�   r€   Ú
cg_usecaser‚   r•   r   r   r   ú<module>r§      st  ðØ Ð Ð Ð Ð Ð Ø +Ð +Ð +Ð +Ð +Ð +Ø Ð Ð Ð Ø 
€
€
€
ðð ð ð ð ñ ô ð ð@%ð %ð %ð %ð %�'ñ %ô %ð %ð
 
€„�ÐÑÔðð ñ Ôðð 
€„�ÐÑÔðð ñ Ôðð ˆkÐ,Ñ-Ô-€Ø!�kÐ"<Ñ=Ô=Ð à€ð
 
€„�ÐÑÔðð ñ Ôðð 
€„�ÐÑÔð!ð !ñ Ôð!ð 
€„�ÐÑÔð!ð !ñ Ôð!ð 	ˆ�LÑ!Ô!€Ø�Ð2Ñ3Ô3€ð �R”X  R¤W˜~°°R´ZÐ/@ÐAÑBÔBÐ Ø�b”h  b¤g °°b´jÐ0AÐBÈ$ÐOÑOÔOÐ àˆRŒX�aÐ1Ð2Ñ2Ô2€
Ø	ˆˆzŒÑ	Ô	ð "ð "€AØ˜Q™€Jˆq„M�#ÑØ˜T™€Jˆq„M�#ÑÐàˆbŒh�zÐ)<Ð=Ñ=Ô=€ð 
€„�ÐÑÔðð ñ Ôðð #�{ =Ð8JÐKÑKÔKÐ Ø#˜ MÐ9LÐMÑMÔMÐ ð
 ð  ð  ð ˆ<˜‰?Œ?€Øˆ<˜‰?Œ?€Øˆ<˜‰?Œ?€Øˆ<˜‰?Œ?€ð
 
€„�ÐÑÔðð ñ Ôðð  �KÐ 2Ñ3Ô3Ð ð 
€„�ÐÑÔðð ñ Ôðð  �KÐ 2Ñ3Ô3Ð ð 
€„�ÐÑÔð1ð 1ñ Ôð1ðl 
€„�ÐÑÔðð ñ Ôðð $˜Ð$9Ñ:Ô:Ð ð
 
€„�ÐÑÔðð ñ Ôðð
 ˆ[Ð*Ñ+Ô+€
ð%ð %ð %ð %ð %�,ñ %ô %ð %ð($ð $ð $ð $ð $r   