
    J/Ph                     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      f/var/www/html/test/jupyter/venv/lib/python3.11/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#bjoo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<<<LL=A77L
<-*----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   ,        bEAbEMAAbEEEr   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    q519r   c                 F    t          |d          |d                   | d<   d S r2   r>   r5   s      r   outer_kernelrA   C   #    1R5&!B%  AbEEEr   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    FAbEEEr   )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    AbE	"r   )r   jitr*   )r7   rO   s   ` r   make_closurerQ   j   sD    	XD     wr            	   c                 "    |d         dz   | d<   d S )Nr   rH   r   r6   r7   s     r   ambiguous_functionrX   z       bEAIAbEEEr   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BqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEBqqqEEEr   c                     |d         | d<   d S r2   r   rW   s     r   simple_usecase_kernelrz      s    bEAbEEEr   c                 j    t           j                                        }|                                 d S r   )r   cg	this_gridsync)r6   r7   grids      r   cg_usecase_kernelr      s&    7DIIKKKKKr   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 5q999 2 21a 8 8!<<<		!Q333--cna@@
j 1 19===//CCk 2 2I>>>!!!$$$$$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MM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 RXRW~RZ/@ABB bhbgbj0AB$OOO RXa1222
	z		 " "AQJqM#TJqM#bhz)<=== 
   #{=8JKKK #M9LMMM 
      <??<??<??<??
 
    K 233  
    K 233  
1 1 1l 
   $$9:: 
 
  
 [*++
% % % % %, % % %($ $ $ $ $r   