+
    :iA                     h   ^ RI t ^ RIHt ^ RIHt ^ RIHt ^ RIH	t	H
t
HtHtHtHtHt ^ RIHt ^ RIHt ^ RIHtHt ^ RIHt ^ R	IHt ^R
IHt ^ RIHtHtHt ^ RI H!t!  ! R R]P&                  4      t"] PF                  ! R] PH                  4      t% ! R R]4      t& ! R R]4      t' ! R R]4      t(R# )    N)cached_property)ir)cgutilsconfig	debuginfoitanium_manglertypestypingutils)
Dispatcher)BaseContext)BaseCallConvMinimalCallConv)	cmathdecl)	datamodel)nvvm)codegen	nvvmutilsufuncs)cuda_data_managerc                   8   a a ] tR t^t oR tV 3R ltRtVtV ;t# )CUDATypingContextc                   ^RI HpHpHpHp ^ RIHpHp V P                  VP                  4       V P                  VP                  4       V P                  VP                  4       V P                  \        P                  4       V P                  VP                  4       V P                  VP                  4       V P                  VP                  4       R# )   )cudadeclcudamathlibdevicedeclvector_types)enumdecl
cffi_utilsN) r   r   r   r   numba.core.typingr   r    install_registryregistryr   typing_registry)selfr   r   r   r   r   r    s   &      O/var/www/html/photoedit/myenv/lib/python3.14/site-packages/numba/cuda/target.pyload_additional_registries,CUDATypingContext.load_additional_registries   s    EE:h//0j112h//0i001m445h//0l::;    c                  < ^ RI Hp \        V\        4      '       d   \        W4      '       g    VP                  p\        \        V `;  V4      #   \
         d    TP                  '       g   \        R4      hTP                  P                  4       pRTR&   TP                  RR4      TR&   TP                  RR4      TR&   T! TP                  T4      pYAn        Tp Li ; i)r   )CUDADispatcherz<using cpu function on device but its compilation is disabledTdevicedebugFopt)numba.cuda.dispatcherr,   
isinstancer   _CUDATypingContext__dispatcherAttributeError_can_compile
ValueErrortargetoptionscopygetpy_funcsuperr   resolve_value_type)r&   valr,   r6   disp	__class__s   &&   r'   r;   $CUDATypingContext.resolve_value_type#   s    8sJ''3//&&  &@EE " '''$ &G H H # 1 1 6 6 8*.h')6):):7E)Jg&'4'8'8'Ee$%ckk=A $( s   A C,A/CC )	__name__
__module____qualname____firstlineno__r(   r;   __static_attributes____classdictcell____classcell__r>   __classdict__s   @@r'   r   r      s     
<F Fr*   r   z	[^a-z0-9]c                      a a ] tR t^Ct oRtRtRV 3R llt]R 4       t]R 4       t	R t
R tR tR t]R	 4       t]R
 4       t]R 4       tRRRR/R ltRR ltR tR tR tR tR tR tRtVtV ;t# )CUDATargetContextTc                x   < \         SV `  W4       \        P                  ! \        P
                  4      V n        R # N)r:   __init__r   chainr   default_managerdata_model_manager)r&   	typingctxtargetr>   s   &&&r'   rN   CUDATargetContext.__init__G   s,    +"3"9"9%%#
r*   c                "    \         P                  # rM   )r   	DIBuilderr&   s   &r'   rV   CUDATargetContext.DIBuilderM   s    """r*   c                    R # Fr@   rW   s   &r'   enable_boundscheck$CUDATargetContext.enable_boundscheckQ   s     r*   c                8    V P                   P                  V4      # rM   )_internal_codegen_create_empty_module)r&   names   &&r'   create_moduleCUDATargetContext.create_moduleW   s    %%::4@@r*   c                J    \         P                  ! R 4      V n        RV n        R# )znumba.cuda.jitN)r   JITCUDACodegenr^   _target_datarW   s   &r'   initCUDATargetContext.initZ   s    !(!7!78H!I r*   c                   ^ RI HpHpHp ^ RI HpHpHp ^ RI HpHp ^ RI H	p	 ^ RI
Hp
 ^ RIHp ^ RIHp ^RIHpHpHpHpHp ^ R	IHp V P/                  VP0                  4       V P/                  V
P0                  4       V P/                  VP0                  4       V P/                  VP0                  4       V P/                  V	P0                  4       V P/                  VP0                  4       V P/                  VP2                  4       R
# )r   )numberstupleobjslicing)rangeobj	iteratorsenumimpl)unicodecharseq)	cmathimpl)cffiimpl)arrayobj)
npdatetime)cudaimpl	printimpllibdeviceimplmathimplr   )ndarrayN)numba.cpythonri   rj   rk   rl   rm   rn   ro   rp   rq   
numba.miscrr   numba.nprs   rt   r!   ru   rv   rw   rx   r   numba.np.unsafery   r#   r$   impl_registry)r&   ri   rj   rk   rl   rm   rn   ro   rp   rq   rr   rs   rt   ru   rv   rw   rx   r   ry   s   &                  r'   r(   ,CUDATargetContext.load_additional_registries^   s     	=<??2+'%'	
 	
 	,h//0h//0i001m445i001h//0l889r*   c                    V P                   # rM   )r^   rW   s   &r'   r   CUDATargetContext.codegenv   s    %%%r*   c                    V P                   f9   \        P                  ! \        P                  ! 4       P
                  4      V n         V P                   # rM   )re   llcreate_target_datar   NVVMdata_layoutrW   s   &r'   target_dataCUDATargetContext.target_datay   s9    $ " 5 5diik6M6M ND   r*   c                    ^ RI Hp Rp\        V Uu. uF  p\        P                  ! V4      V3NK  	  up4      pV# u upi )z
Some CUDA intrinsics are at the module level, but cannot be treated as
constants, because they are loaded from a special register in the PTX.
These include threadIdx, blockDim, etc.
cuda)	threadIdxblockDimblockIdxgridDimlaneidwarpsize)numbar   tupler	   Module)r&   r   	nonconstsncnonconsts_with_mods   &    r'   nonconst_module_attrs'CUDATargetContext.nonconst_module_attrs   sM     	!	".7$9.7 &+\\$%7$<.7$9 :!!$9s   !<c                    \        V 4      # rM   )CUDACallConvrW   s   &r'   	call_convCUDATargetContext.call_conv   s    D!!r*   abi_tagsuidNc               4    \         P                  ! WVVR 7      # ))r   r   )r   mangle)r&   r`   argtypesr   r   s   &&&$$r'   manglerCUDATargetContext.mangler   s    %%dx*-/ 	/r*   c	           	         \         P                  ! VP                  RR7      p	V P                  4       P	                  VP
                   R2V	VVR7      p
V
P                  V4       V P                  WV	W4VV4      pW3# )a;  
Adapt a code library ``codelib`` with the numba compiled CUDA kernel
with name ``fname`` and arguments ``argtypes`` for NVVM.
A new library is created with a wrapper function that can be used as
the kernel entry point for the given kernel.

Returns the new code library and the wrapper function.

Parameters:

codelib:       The CodeLibrary containing the device function to wrap
               in a kernel call.
fndesc:        The FunctionDescriptor of the source function.
debug:         Whether to compile with debug.
lineinfo:      Whether to emit line info.
nvvm_options:  Dict of NVVM options used when compiling the new library.
filename:      The source filename that the function is contained in.
linenum:       The source line that the function is on.
max_registers: The max_registers argument for the code library.
cudapyns_kernel_)
entry_namenvvm_optionsmax_registers)r   prepend_namespacellvm_func_namer   create_libraryr`   add_linking_librarygenerate_kernel_wrapper)r&   codelibfndescr.   lineinfor   filenamelinenumr   kernel_namelibrarywrappers   &&&&&&&&&   r'   prepare_cuda_kernel%CUDATargetContext.prepare_cuda_kernel   s    . &77!!h
 ,,.//7<<.0I;F=I>K 0 M 	##G,..w/4/68 r*   c           	     	  a!a" VP                   pV P                  V4      p	\        V	P                  4      p
\        P
                  ! \        P                  ! 4       V
4      pV P                  R4      o"\        P
                  ! \        P                  ! ^ 4      V P                  P                  \        P                  4      .V
,           4      p\        P                  ! S"WP                  4      p\        P                   ! VP"                  RR7      p\        P                  ! S"W4      o!\        P$                  ! S!P'                  R4      4      pV'       g	   V'       dW   T;'       d    V'       * pV P)                  S"VV VR7      pVP*                  ! S!W2P,                  W4       VP.                  ! W4       V!V"3R lpV! R4      p. p. pR F?  pVP1                  V! R	V,          4      4       VP1                  V! R
V,          4      4       KA  	  V	P3                  VS!P,                  4      pV P                  P5                  W\        P6                  VV4      w  ppV'       Ed   \8        P:                  ! VVP<                  4      ;_uu_ 4        VP?                  4        RRR4       VPA                  VPC                  VPD                  4      4      ;_uu_ 4        \        PF                  ! VPH                  PJ                  R4      pVPM                  VVVPN                  RR4      pVPQ                  V^4      p\R        PT                  ! V4      pVPA                  V4      ;_uu_ 4        \W        RV4       F)  w  ppVPY                  V4      p VP[                  V V4       K+  	  \W        RV4       F)  w  ppVP]                  V4      p VP[                  V V4       K+  	  RRR4       RRR4       VP?                  4        \^        P`                  ! S!4       VPc                  S"4       V'       g	   V'       d   XPd                  ! 4        VPe                  4        \f        Ph                  '       d   \j        Pl                  ! VS"4       VPo                  S!P"                  4      #   + '       g   i     EL; i  + '       g   i     L; i  + '       g   i     L; i)z
Generate the kernel wrapper in the given ``library``.
The function being wrapped is described by ``fndesc``.
The wrapper function is returned.
zcuda.kernel.wrapperr   r   r!   )modulefilepathcgctxdirectives_onlyc                    < SP                   V ,           p\        P                  ! S\        P                  ! ^ 4      V4      p\        P
                  ! VP                  P                  R4      Vn        V# )    N)	r`   r   add_global_variabler   IntTypeConstanttypepointeeinitializer)postfixr`   gvwrapfnwrapper_modules   &  r'   define_error_gvBCUDATargetContext.generate_kernel_wrapper.<locals>.define_error_gv   sO    ;;(D,,^RZZ^-13B[[$?BNIr*   __errcode__xyzz	__tid%s__z__ctaid%s__N	monotonic)8r   get_arg_packerlistargument_typesr   FunctionTypeVoidTypera   r   r   get_return_typer	   pyobjectFunctionr   r   r   r`   	IRBuilderappend_basic_blockrV   mark_subprogramargsmark_locationappendfrom_argumentscall_functionvoidr   	if_likelyis_okret_voidif_thennot_is_python_excr   r   r   cmpxchgcodeextract_valuer   SRegBuilderziptidstorectaidr   set_cuda_kerneladd_ir_modulefinalizer   	DUMP_LLVMr   	dump_llvmget_function)#r&   r   r   r   r.   r   r   r   r   arginfoargtyswrapfntyfntyfuncprefixedbuilderr   r   r   gv_excgv_tidgv_ctaidicallargsstatus_oldxchgchangedsregdimptrr<   r   r   s#   &&&&&&&&                         @@r'   r   )CUDATargetContext.generate_kernel_wrapper   s    ??%%h/g,,-??2;;=&9++,ABrzz"~ $ > >u~~ NO!' () {{>41F1FG"44TYY8L^X@,,v88<=H&44u9On08-17F ' HI %%[[( ##G5	 !/AMM/+/:;OOOMA,=>?  ))'6;;?NN005::x;	 5""7FLL99  " : f.B.B!CDDkk&++"5"5t<
 vsFKK'2KA!//a8 !,,W5__W--%(%7	S"hhsmc3/ &8 &)%9	S"jjoc3/ &: . E, 	V$n-H OOFN3##FKK00I :99  .- EDs2   RBR5A3R"	R5R	"R2-R55S	c           
        VP                   p\        VP                  RR7      4       Uu. uF#  pV P                  \        P
                  V4      NK%  	  pp\        P                  ! \        P                  ! ^4      \        V4      4      p\        P                  ! Wv4      p\        P                  p	\        P                  ! WHP                  RV	R7      p
RV
n        RV
n        Wn        V P'                  VP(                  4      pV P+                  V4      p^V^,
          P-                  4       ,          V
n        \        P0                  ! \        P                  ! ^4      4      pVP3                  WR4      pV P5                  V4      ! W4      pVP6                   Uu. uF#  pV P                  \        P8                  V4      NK%  	  ppVP:                   Uu. uF#  pV P                  \        P8                  V4      NK%  	  ppV P=                  WP?                  WP@                  P                  4      VVVPB                  VPD                  RR	7       VPG                  4       # u upi u upi u upi )
Q
Unlike the parent version.  This returns a a pointer in the constant
addrspace.
A)order_cudapy_cmem	addrspaceinternalTgenericN)datashapestridesitemsizeparentmeminfo)$r   itertobytesget_constantr	   byter   	ArrayTyper   lenr   r   ADDRSPACE_CONSTANTr   r   r   linkageglobal_constantr   get_data_typedtypeget_abi_sizeof
bit_lengthalignPointerTypeaddrspacecast
make_arrayr  intpr  populate_arraybitcastr  r  r  	_getvalue)r&   r   arytyarrlmodr  	constvals
constarytyconstaryr  r   lldtyper(  ptrtygenptraryskshapekstridess   &&&&               r'   make_constant_array%CUDATargetContext.make_constant_array  s    ~~ #++C+01
1 ejj!,1 	 
 \\"**Q-Y@
;;z5++	((}}n3<>
!! $$U[[1##G,..00 rzz!}-&&r)< ooe$T3<?IIFIq$##EJJ2IF>AkkJkD%%ejj!4kJCoofhhmm&L"($,%(\\#**$(	 	 	* }}C
2 GJs   )I)I)I"c                   \         P                  ! VP                  R4      R,           4      pRP                  R\        P
                  ! V4      .4      pVP                  P                  V4      pVfF   \         P                  ! WP                  V\        P                  R7      pRVn        RVn        W5n        VP                  P                  P                   pVP#                  VP%                  \        P                  4      4      # )r  zutf-8    $__conststring__r  r  T)r   make_bytearrayencodejoinr   mangle_identifierglobalsr8   r   r   r   r!  r"  r#  r   r   elementr.  
as_pointer)r&   modstringtextr`   r   chartys   &&&    r'   insert_const_string%CUDATargetContext.insert_const_string@  s    
 %%fmmG&<w&FGxx*(::6BD E [[__T":,,S))T7;7N7NPB#BJ!%B!N ((zz&++D,C,CDEEr*   c                    VP                   pV P                  W24      p\        P                  ! \        P                  ! ^4      4      pVP                  WER4      # )z
Insert a constant string in the constant addresspace and return a
generic i8 pointer to the data.

This function attempts to deduplicate.
r  )r   rN  r   r)  r   r*  )r&   r   rK  r2  r   	charptrtys   &&&   r'   insert_string_const_addrspace/CUDATargetContext.insert_string_const_addrspaceV  sF     ~~%%d3NN2::a=1	$$RI>>r*   c                    R# )zRun O1 function passes
        Nr@   r&   r   s   &&r'   optimize_function#CUDATargetContext.optimize_functionb  s     	r*   c                .    \         P                  ! V4      # rM   )r   get_ufunc_info)r&   	ufunc_keys   &&r'   rY   CUDATargetContext.get_ufunc_infoo  s    $$Y//r*   )r^   re   rQ   r   r@   rM   )rA   rB   rC   rD   implement_powi_as_math_callstrict_alignmentrN   propertyrV   r[   ra   rf   r(   r   r   r   r   r   r   r   r   r=  rN  rR  rV  rY  rE   rF   rG   rH   s   @@r'   rK   rK   C   s     "&
 # #  
A!:0& ! !
 " " " "/" /$ /" HZ1x)VF,
?0 0r*   rK   c                       ] tR tRtRtR# )r   is  r@   N)rA   rB   rC   rD   rE   r@   r*   r'   r   r   s  s    r*   r   c                   b   a  ] tR tRt o RtR tR tRR ltR tR t	RR	 lt
R
 tR tR tRtV tR# )CUDACABICallConviw  z
Calling convention aimed at matching the CUDA C/C++ ABI. The implemented
function signature is:

    <Python return type> (<Python arguments>)

Exceptions are unsupported in this convention.
c                    R # rM   r@   )r&   r   s   &&r'   _make_call_helper"CUDACABICallConv._make_call_helper  s     r*   c                $    VP                  V4      # rM   )ret)r&   r   retvals   &&&r'   return_valueCUDACABICallConv.return_value  s    {{6""r*   Nc                    R p\        V4      h)z7Python exceptions are unsupported in the CUDA C/C++ ABINotImplementedError)r&   r   excexc_argsloc	func_namemsgs   &&&&&& r'   return_user_exc CUDACABICallConv.return_user_exc  s    G!#&&r*   c                    R p\        V4      h)z2Return status is unsupported in the CUDA C/C++ ABIrk  )r&   r   r  rq  s   &&& r'   return_status_propagate(CUDACABICallConv.return_status_propagate  s    B!#&&r*   c                    V P                  V4      p\        VP                  4      p\        P                  ! V P                  V4      V4      pV# )z=
Get the LLVM IR Function type for *restype* and *argtypes*.
)_get_arg_packerr   r   r   r   r   )r&   restyper   r   r   s   &&&  r'   get_function_type"CUDACABICallConv.get_function_type  sD     &&x0../t33G<hGr*   c                    V'       d   Q hV P                  V4      pTP                  V P                  V4      V Uu. uF  pRV,           NK  	  up4       R# u upi )z1
Set names and attributes of function arguments.
zarg.N)rx  assign_namesget_arguments)r&   fnr   fe_argtypesnoaliasr   as   &&&&&  r'   decorate_function"CUDACABICallConv.decorate_function  sR     {&&{3T//3267$Qfqjj$7	97s   A
c                    VP                   # )z0
Get the Python-level arguments of LLVM *func*.
)r   rU  s   &&r'   r~  CUDACABICallConv.get_arguments  s     yyr*   c                    V P                  V4      pVP                  W4      pVP                  W'4      pRp	V P                  P	                  WV4      p
W3# )z#
Call the Numba-compiled *callee*.
N)rx  as_argumentscallcontextget_returned_value)r&   r   calleerestyr   r   r   realargsr   r  outs   &&&&&&     r'   r   CUDACABICallConv.call_function  sW     &&v.''6||F- ll--gdC{r*   c                X    V P                   P                  V,          P                  4       # rM   )r  rQ   r   )r&   tys   &&r'   r    CUDACABICallConv.get_return_type  s     ||..r2BBDDr*   r@   )NNNrZ   )rA   rB   rC   rD   __doc__rc  rh  rr  ru  rz  r  r~  r   r   rE   rF   )rI   s   @r'   ra  ra  w  s>     
#'
'9E Er*   ra  ))re	functoolsr   llvmlite.bindingbindingr   llvmliter   
numba.corer   r   r   r   r	   r
   r   numba.core.dispatcherr   numba.core.baser   numba.core.callconvr   r   r"   r   r   cudadrvr   
numba.cudar   r   r   numba.cuda.modelsr   r   compileIVALID_CHARSrK   r   ra  r@   r*   r'   <module>r     s    	 %  ' ' ' , ' = '    1 1 /$F** $FT jjrtt,m0 m0`		? 	AE| AEr*   