
    xKg                        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j2                  de      Z eej8                        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% Z1y)&    )cuda)CUDATestCaseNc                   .    e Zd ZdZddZd Zed        Zy)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                      || _         || _        y N)_func_retty)selffuncrettys      j/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/cache_usecases.py__init__zUseCase.__init__   s    
    c                    |D cg c]  }t        j                  |       }}| j                  r"t        j                  d| j                        }nt        j                  |d         } | j
                  |g|  |d   S c c}w )N dtyper   )npasarrayr
   ndarray
zeros_like_call)r   argsarg
array_argsarray_returns        r   __call__zUseCase.__call__   sn    156#bjjo
6;;::b<L==A7L

<-*-B 7s   Bc                     | j                   S r   r	   )r   s    r   r   zUseCase.func"   s    zzr   r   )__name__
__module____qualname____doc__r   r   propertyr   r   r   r   r   r      s%       r   r   c                       e Zd Zd Zy)CUDAUseCasec                 0     | j                   d   |g|  y )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<   y Nr   Zrxys      r   add_usecase_kernelr6   ,       bEAbEMAAbEr   Fc                 .    |d   |d   z   t         z   | d<   y r/   r0   r2   s      r   add_nocache_usecase_kernelr9   1   r7   r   r*   c                     | |z   t         z   S r   r0   )r4   r5   s     r   innerr;   >   s    q519r   c                 0    t        |d    |d         | d<   y r/   r;   r2   s      r   outer_kernelr>   C       1R5&!B% AbEr   c                 0    t        |d    |d         | d<   y r/   r=   r2   s      r   outer_uncached_kernelrA   H   r?   r   ab)align   r   g     @E@c                     ||   | d<   y r/   r   )r3   aryis      r   record_returnrI   _   s    FAbEr   )r   c                 V     t        j                  d       fd       }t        |      S )NTr,   c                     |d   z   | d<   y r/   r   )r3   r5   r4   s     r   closurezmake_closure.<locals>.closurek   s    AbE	"r   )r   jitr'   )r4   rL   s   ` r   make_closurerN   j   s+    	XXD  wr            	   c                     |d   dz   | d<   y )Nr   rE   r   r3   r4   s     r   ambiguous_functionrU   z       bEAIAbEr   c                     |d   dz   | d<   y )Nr      r   rT   s     r   rU   rU      rV   r   c                  	   t         j                  j                  dt        j                        } t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }	t         j                  j                  dt        j                        }
t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  j                  dt        j                        }t         j                  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  y )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_localsru      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                     |d   | d<   y r/   r   rT   s     r   simple_usecase_kernelrw      s    bEAbEr   c                 `    t         j                  j                         }|j                          y r   )r   cg	this_gridsync)r3   r4   grids      r   cg_usecase_kernelr}      s    77DIIKr   c                       e Zd ZdZd Zy)_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                    | j                  |j                  dd      d       | j                  |j                  dd      d       | j                  |j                  dd      d       |j	                  |j
                  d      }| j                  t        |      d       |j                  |j                  d      }| j                  t        |      d       |j                  d       y )NrE   rO   rX   r*   )rE   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 5q9 2 21a 8!<		!Q3--cnna@
j 19=//Ck 2I>!!!$r   N)r!   r"   r#   r$   r   r   r   r   r   r      s    
%r   r   c                  d    t         j                  t           } t               j	                  |        y 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'   rM   r6   r9   r   add_nocache_usecaser1   r;   r>   rA   r   r   r   int8r\   packed_record_typealigned_record_typeemptyr   rangesizerH   r[   r   rI   r   r   rN   closure1closure2closure3closure4rU   renamed_function1renamed_function2ru   rw   r   r}   
cg_usecaser   r   r   r   r   <module>r      s    +  
 @%' %
 
  
  ,-!"<= 
 
  
! ! 
! ! 	L!23 RXXRWW~RZZ/@AB bhhbggbjj0AB$O RXXa12
	z	AQJqM#TJqM# 
  bhhz)<= 
  #=8JK #M9LM 
  ????
 
    23  
    23  
1 1l 
  $$9: 
 
 
 *+
%, %($r   