
    i                        S SK Jr  S SKJr  S SKrS SKr " S S5      r " S S\5      r\R                  " SS	9S
 5       r
\R                  " SS	9S 5       r\" \
5      r\" \5      rSr\R                  " SS	9S 5       r\R                  " SS	9S 5       r\R                  " SS	9S 5       r\" \5      r\" \5      r\R(                  " S\R*                  4S\R,                  4/5      r\R(                  " S\R*                  4S\R,                  4/SS9r\R2                  " S\S9r\" \R8                  5       H  r\S-   \\   S'   \S-   \\   S'   M     \R<                  " \\S9r\R                  " SS	9S 5       r \" \ \S9r!\" \ \S9r"S r#\#" S5      r$\#" S5      r%\#" S5      r&\#" S5      r'\R                  " SS	9S 5       r(\" \(5      r)\R                  " SS	9S 5       r(\" \(5      r*\R                  " SS	9S  5       r+\R                  " SS	9S! 5       r,\" \,5      r-\R                  " SS	9S" 5       r.\" \.5      r/ " S# S$\5      r0S% r1g)&    )cuda)CUDATestCaseNc                   8    \ rS rSrSrSS jrS r\S 5       rSr	g)	UseCase   a  
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                     Xl         X l        g N_func_retty)selffuncrettys      a/var/www/html/trading/venv/lib/python3.13/site-packages/numba/cuda/tests/cudapy/cache_usecases.py__init__UseCase.__init__   s    
    c                    U Vs/ s H  n[         R                  " U5      PM     nnU R                  (       a   [         R                  " SU R                  S9nO[         R                  " US   5      nU R
                  " U/UQ76   US   $ s  snf )N dtyper   )npasarrayr   ndarray
zeros_like_call)r   argsarg
array_argsarray_returns        r   __call__UseCase.__call__   sl    156#bjjo
6;;::b<L==A7L

<-*-B 7s    B	c                     U R                   $ r	   r   )r   s    r   r   UseCase.func"   s    zzr   r
   r	   )
__name__
__module____qualname____firstlineno____doc__r   r!   propertyr   __static_attributes__r   r   r   r   r      s%       r   r   c                       \ rS rSrS rSrg)CUDAUseCase'   c                 0    U R                   S   " U/UQ76   g )N   r2   r$   )r   retr   s      r   r   CUDAUseCase._call(   s    

4$t$r   r   N)r&   r'   r(   r)   r   r,   r   r   r   r.   r.   '   s    %r   r.   Tcachec                 .    US   US   -   [         -   U S'   g Nr   Zrxys      r   add_usecase_kernelr?   ,       bEAbEMAAbEr   Fc                 .    US   US   -   [         -   U S'   g r8   r9   r;   s      r   add_nocache_usecase_kernelrB   1   r@   r   r2   c                     X-   [         -   $ r	   r9   )r=   r>   s     r   innerrD   >   s    519r   c                 0    [        US   * US   5      U S'   g r8   rD   r;   s      r   outer_kernelrG   C       1R5&!B% AbEr   c                 0    [        US   * US   5      U S'   g r8   rF   r;   s      r   outer_uncached_kernelrJ   H   rH   r   ab)align   r   g     @E@c                     X   U S'   g r8   r   )r<   aryis      r   record_returnrR   _   s    FAbEr   )r   c                 T   ^  [         R                  " SS9U 4S j5       n[        U5      $ )NTr5   c                    > TUS   -   U S'   g r8   r   )r<   r>   r=   s     r   closuremake_closure.<locals>.closurek   s    AbE	"r   )r   jitr.   )r=   rU   s   ` r   make_closurerX   j   s+    	XXD  wr         r   	   c                     US   S-   U S'   g )Nr   rN   r   r<   r=   s     r   ambiguous_functionr^   z       bEAIAbEr   c                     US   S-   U S'   g )Nr      r   r]   s     r   r^   r^      r_   r   c                  	   [         R                  R                  S[        R                  5      n [         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n	[         R                  R                  S[        R                  5      n
[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      n[         R                  R                  S[        R                  5      nSU S S & SUS S & SUS S & SUS S & SUS S & SUS S & SUS S & SUS S & SUS S & SU	S S & SU
S S & SUS S & SUS S & SUS S & SUS S & SUS S & SUS S & SUS S & SUS S & SUS S & SUS S & SUS S & SUS S & SUS S & g )Nr1   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_localsr~      sJ   			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-B			&"**	-BBqEBqEBqEBqEBqEBqEBqEBqEBqEBqEBqEBqEBqEBqEBqEBqEBqEBqEBqEBqEBqEBqEBqEBqEr   c                     US   U S'   g r8   r   r]   s     r   simple_usecase_kernelr      s    bEAbEr   c                 `    [         R                  R                  5       nUR                  5         g r	   )r   cg	this_gridsync)r<   r=   grids      r   cg_usecase_kernelr      s    77DIIKr   c                       \ rS rSrSrS rSrg)_TestModule   z
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                    U R                  UR                  SS5      S5        U R                  UR                  SS5      S5        U R                  UR                  SS5      S5        UR	                  UR
                  S5      nU R                  [        U5      S5        UR                  UR                  S5      nU R                  [        U5      S5        UR                  S5        g )NrN   rY   ra   r2   )rN   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_module_TestModule.check_module   s    1 5q9 2 21a 8!<		!Q3--cnna@
j 19=//Ck 2I>!!!$r   r   N)r&   r'   r(   r)   r*   r   r,   r   r   r   r   r      s    
%r   r   c                  d    [         R                  [           n [        5       R	                  U 5        g r	   )sysmodulesr&   r   r   )r   s    r   	self_testr      s     
++h
CMs#r   )2numbar   numba.cuda.testingr   numpyr   r   r   r.   rW   r?   rB   r   add_nocache_usecaser:   rD   rG   rJ   r   r   r   int8re   packed_record_typealigned_record_typeemptyr   rangesizerQ   rd   r   rR   r   r   rX   closure1closure2closure3closure4r^   renamed_function1renamed_function2r~   r   r   r   
cg_usecaser   r   r   r   r   <module>r      s    +  
 @%' %
     ,-!"<= 
   ! ! ! ! 	L!23 XXRWW~RZZ/@AB hhbggbjj0AB$O XXa12
	z	AQJqM#TJqM# 
  hhz)<=   #=8JK #M9LM 
  ????
     23      23  1 1l   $$9: 
  
 *+
%, %($r   