ó
    ýZhÓ"  ã                   óð   • S SK Jr  S SKJr  S SKJr  S SKJrJ	r	J
r
   " S S\5      r " S S\
5      r " S	 S
\	5      r " S S\5      rSr " S S\R"                  5      rSr " S S\R(                  5      rg)é    )Úcuda)Úarray)Údeviceufunc)ÚUFuncMechanismÚGeneralizedUFuncÚGUFuncCallStepsc                   ó4   • \ rS rSrSrS rS rS	S jrS rSr	g)
ÚCUDAUFuncDispatcheré   z<
Invoke the CUDA ufunc specialization for the given inputs.
c                 ó2   • Xl         UR                  U l        g ©N)Ú	functionsÚ__name__)ÚselfÚtypes_to_retty_kernelsÚpyfuncs      Úe/home/gothic/public_html/Fooocus/fooocus_env/lib64/python3.13/site-packages/numba/cuda/vectorizers.pyÚ__init__ÚCUDAUFuncDispatcher.__init__   s   € Ø/ŒØŸ™ˆ�ó    c                 óB   • [         R                  U R                  X5      $ )af  
*args: numpy arrays or DeviceArrayBase (created by cuda.to_device).
       Cannot mix the two types in one call.

**kws:
    stream -- cuda stream; when defined, asynchronous mode is used.
    out    -- output array. Can be a numpy array or DeviceArrayBase
              depending on the input arguments.  Type must match
              the input arguments.
)ÚCUDAUFuncMechanismÚcallr   )r   ÚargsÚkwss      r   Ú__call__ÚCUDAUFuncDispatcher.__call__   s   € ô "×&Ñ& t§~¡~°tÓAÐAr   c                 ó˜  • [        [        U R                  R                  5       5      S   5      S:X  d   S5       eUR                  S:X  d   S5       eUR
                  S   n/ nUS:X  a  [        S5      eUS:X  a  US   $ U=(       d    [        R                  " 5       nUR                  5          [        R                  R                  R                  U5      (       a  UnO[        R                  " X5      nU R                  XTU5      n[        SUR                   S9nUR#                  XrS	9  S S S 5        US   $ ! , (       d  f       WS   $ = f)
Nr   é   zmust be a binary ufuncé   zmust use 1d arrayzReduction on an empty array.)r    )Údtype©Ústream)ÚlenÚlistr   ÚkeysÚndimÚshapeÚ	TypeErrorr   r#   Úauto_synchronizeÚcudadrvÚdevicearrayÚis_cuda_ndarrayÚ	to_deviceÚ_CUDAUFuncDispatcher__reduceÚnp_arrayr!   Úcopy_to_host)r   Úargr#   ÚnÚgpu_memsÚmemÚoutÚbufs           r   ÚreduceÚCUDAUFuncDispatcher.reduce   s'  € Ü”4˜Ÿ™×+Ñ+Ó-Ó.¨qÑ1Ó2°aÓ7ð 	Að :Aó 	AÐ7à�x‰x˜1‹}Ð1Ð1Ó1ˆ}à�I‰I�a‰LˆØˆà�‹6ÜÐ:Ó;Ð;Ø�!‹VØ�q‘6ˆMð ×(œ4Ÿ;š;›=ˆØ×$Ñ$Õ&ä�|‰|×'Ñ'×7Ñ7¸×<Ñ<Ø‘ä—n’n SÓ1�à—-‘- ¨vÓ6ˆCä˜4 s§y¡yÑ1ˆCØ×Ñ˜SÐÑ0÷ 'ð �1‰vˆ÷ 'Ô&ð �1‰vˆús   Â-A=D7Ä7
E	c                 ó®  • UR                   S   nUS-  S:w  ab  UR                  US-
  5      u  pVUR                  U5        UR                  U5        U R                  XRU5      nUR                  U5        U " XvXsS9$ UR                  US-  5      u  p‰UR                  U5        UR                  U	5        U " X‰XƒS9  US-  S:”  a  U R                  X‚U5      $ U$ )Nr   r   r    )r6   r#   )r(   ÚsplitÚappendr/   )
r   r5   r4   r#   r3   ÚfatcutÚthincutr6   ÚleftÚrights
             r   Ú__reduceÚCUDAUFuncDispatcher.__reduce;   sÆ   € Ø�I‰I�a‰LˆØˆq‰5�A‹:Ø!Ÿi™i¨¨A©Ó.‰OˆFà�O‰O˜FÔ#Ø�O‰O˜GÔ$à—-‘- °&Ó9ˆCØ�O‰O˜CÔ Ù˜¨#Ñ=Ð=àŸ)™) A¨¡FÓ+‰KˆDà�O‰O˜DÔ!Ø�O‰O˜EÔ"á� $Ò6Ø�A‰v˜‹zØ—}‘} T°VÓ<Ð<à�r   )r   r   N©r   )
r   Ú
__module__Ú__qualname__Ú__firstlineno__Ú__doc__r   r   r8   r/   Ú__static_attributes__© r   r   r
   r
      s   † ñò(òBôõ:r   r
   c                   óR   ^ • \ rS rSrS/rU 4S jrS rS rS rS r	S r
S	 rS
rU =r$ )Ú_CUDAGUFuncCallStepséS   Ú_streamc                 óT   >• [         TU ]  XX45        UR                  SS5      U l        g )Nr#   r   )Úsuperr   ÚgetrM   )r   ÚninÚnoutr   ÚkwargsÚ	__class__s        €r   r   Ú_CUDAGUFuncCallSteps.__init__X   s$   ø€ Ü‰Ñ˜ DÔ1Ø—z‘z (¨AÓ.ˆ�r   c                 ó.   • [         R                  " U5      $ r   ©r   Úis_cuda_array©r   Úobjs     r   Úis_device_arrayÚ$_CUDAGUFuncCallSteps.is_device_array\   ó   € Ü×!Ò! #Ó&Ð&r   c                 óŽ   • [         R                  R                  R                  U5      (       a  U$ [         R                  " U5      $ r   ©r   r+   r,   r-   Úas_cuda_arrayrY   s     r   Úas_device_arrayÚ$_CUDAGUFuncCallSteps.as_device_array_   ó5   € ô �<‰<×#Ñ#×3Ñ3°C×8Ñ8ØˆJÜ×!Ò! #Ó&Ð&r   c                 ó>   • [         R                  " XR                  S9$ ©Nr"   )r   r.   rM   )r   Úhostarys     r   r.   Ú_CUDAGUFuncCallSteps.to_devicei   s   € Ü�~Š~˜g¯l©lÑ;Ð;r   c                 ó8   • UR                  X R                  S9nU$ re   )r1   rM   )r   Údevaryrf   r6   s       r   Úto_hostÚ_CUDAGUFuncCallSteps.to_hostl   s   € Ø×!Ñ! '·,±,Ð!Ð?ˆØˆ
r   c                 ó@   • [         R                  " XU R                  S9$ ©N)r(   r!   r#   )r   Údevice_arrayrM   )r   r(   r!   s      r   Úallocate_device_arrayÚ*_CUDAGUFuncCallSteps.allocate_device_arrayp   s   € Ü× Ò  uÀ$Ç,Á,ÑOÐOr   c                 ó<   • UR                  X R                  S9" U6   g re   )ÚforallrM   )r   ÚkernelÚnelemr   s       r   Úlaunch_kernelÚ"_CUDAGUFuncCallSteps.launch_kernels   s   € Ø�‰�e§L¡LˆÑ1°4Ò8r   )rM   )r   rD   rE   rF   Ú	__slots__r   r[   ra   r.   rj   ro   ru   rH   Ú__classcell__©rT   s   @r   rK   rK   S   s6   ø† àð€Iõ/ò'ò'ò<òòP÷9ð 9r   rK   c                   óD   ^ • \ rS rSrU 4S jr\S 5       rS rS rSr	U =r
$ )ÚCUDAGeneralizedUFuncéw   c                 óF   >• UR                   U l         [        TU ]	  X5        g r   )r   rO   r   )r   Ú	kernelmapÚenginer   rT   s       €r   r   ÚCUDAGeneralizedUFunc.__init__x   s   ø€ ØŸ™ˆŒÜ‰Ñ˜Õ+r   c                 ó   • [         $ r   )rK   ©r   s    r   Ú_call_stepsÚ CUDAGeneralizedUFunc._call_steps|   s   € ä#Ð#r   c                 ó~   • [         R                  R                  R                  USUR                  UR
                  S9$ ©NrC   ©r(   Ústridesr!   Úgpu_data)r   r+   r,   ÚDeviceNDArrayr!   r‰   )r   Úaryr(   s      r   Ú_broadcast_scalar_inputÚ,CUDAGeneralizedUFunc._broadcast_scalar_input€   s9   € Ü�|‰|×'Ñ'×5Ñ5¸EØ>BØ<?¿I¹IØ?B¿|¹|ð 6ð Mð 	Mr   c                 óä   • [        U5      [        UR                  5      -
  nSU-  UR                  -   n[        R                  R
                  R                  UUUR                  UR                  S9$ r†   )	r$   r(   rˆ   r   r+   r,   rŠ   r!   r‰   )r   r‹   ÚnewshapeÚnewaxÚ
newstridess        r   Ú_broadcast_add_axisÚ(CUDAGeneralizedUFunc._broadcast_add_axis†   sa   € Ü�H“¤ C§I¡I£Ñ.ˆà˜E‘\ C§K¡KÑ/ˆ
Ü�|‰|×'Ñ'×5Ñ5¸HØ>HØ<?¿I¹IØ?B¿|¹|ð 6ð Mð 	Mr   )r   )r   rD   rE   rF   r   Úpropertyrƒ   rŒ   r’   rH   rx   ry   s   @r   r{   r{   w   s.   ø† õ,ð ñ$ó ð$òM÷Mð Mr   r{   c                   óF   • \ rS rSrSrSrS rS rS rS r	S r
S	 rS
 rSrg)r   é�   z
Provide CUDA specialization
r   c                 ó(   • UR                  X#S9" U6   g re   )rr   )r   ÚfuncÚcountr#   r   s        r   ÚlaunchÚCUDAUFuncMechanism.launch–   s   € Ø�‰�EˆÑ)¨4Ò0r   c                 ó.   • [         R                  " U5      $ r   rW   rY   s     r   r[   Ú"CUDAUFuncMechanism.is_device_array™   r]   r   c                 óŽ   • [         R                  R                  R                  U5      (       a  U$ [         R                  " U5      $ r   r_   rY   s     r   ra   Ú"CUDAUFuncMechanism.as_device_arrayœ   rc   r   c                 ó*   • [         R                  " XS9$ re   )r   r.   )r   rf   r#   s      r   r.   ÚCUDAUFuncMechanism.to_device¦   s   € Ü�~Š~˜gÑ5Ð5r   c                 ó    • UR                  US9$ re   )r1   )r   ri   r#   s      r   rj   ÚCUDAUFuncMechanism.to_host©   s   € Ø×"Ñ"¨&Ð"Ð1Ð1r   c                 ó,   • [         R                  " XUS9$ rm   )r   rn   )r   r(   r!   r#   s       r   ro   Ú(CUDAUFuncMechanism.allocate_device_array¬   s   € Ü× Ò  uÀ&ÑIÐIr   c                 ó°  • [        [        U5      5       Vs/ s H+  nX1R                  :¼  d  UR                  U   X#   :w  d  M)  UPM-     nn[        U5      [        UR                  5      -
  nS/U-  [	        UR
                  5      -   nU H  nSXc'   M	     [        R                  R                  R                  UUUR                  UR                  S9$ s  snf )Nr   r‡   )Úranger$   r'   r(   r%   rˆ   r   r+   r,   rŠ   r!   r‰   )r   r‹   r(   ÚaxÚ
ax_differsÚ
missingdimrˆ   s          r   Úbroadcast_deviceÚ#CUDAUFuncMechanism.broadcast_device¯   sÄ   € Ü#(¬¨U«Ô#4ó 5Ò#4˜RØŸx™x›ØŸ™ 2™¨%©)Ñ3÷ Ñ#4ˆ
ð 5ô ˜“Z¤# c§i¡i£.Ñ0ˆ
Ø�#˜
Ñ"¤T¨#¯+©+Ó%6Ñ6ˆãˆBØˆG‹Kñ ô �|‰|×'Ñ'×5Ñ5¸EØ>EØ<?¿I¹IØ?B¿|¹|ð 6ð Mð 	Mùò5s   —(CÁCrI   N)r   rD   rE   rF   rG   ÚDEFAULT_STREAMrš   r[   ra   r.   rj   ro   r«   rH   rI   r   r   r   r   �   s3   † ñð €Nò1ò'ò'ò6ò2òJõMr   r   z�
def __vectorized_{name}({args}, __out__):
    __tid__ = __cuda__.grid(1)
    if __tid__ < __out__.shape[0]:
        __out__[__tid__] = __core__({argitems})
c                   ó<   • \ rS rSrS rS rS rS r\S 5       r	Sr
g)	ÚCUDAVectorizeéÈ   c                 óª   • [         R                  " USSS9" U R                  5      nX"R                  UR                     R
                  R                  4$ )NT)ÚdeviceÚinline)r   Újitr   Ú	overloadsr   Ú	signatureÚreturn_type)r   ÚsigÚcudevfns      r   Ú_compile_coreÚCUDAVectorize._compile_coreÉ   sA   € Ü—(’(˜3 t°DÒ9¸$¿+¹+ÓFˆØ×)Ñ)¨#¯(©(Ñ3×=Ñ=×IÑIÐIÐIr   c                 ó~   • U R                   R                  R                  5       nUR                  [        US.5        U$ )N©Ú__cuda__Ú__core__)r   Ú__globals__ÚcopyÚupdater   )r   ÚcorefnÚglbls      r   Ú_get_globalsÚCUDAVectorize._get_globalsÍ   s5   € Ø�{‰{×&Ñ&×+Ñ+Ó-ˆØ�‰¤Ø!'ñ)ô 	*àˆr   c                 ó.   • [         R                  " U5      $ r   ©r   r´   ©r   Úfnobjr¸   s      r   Ú_compile_kernelÚCUDAVectorize._compile_kernelÓ   s   € Ü�xŠx˜‹Ðr   c                 óB   • [        U R                  U R                  5      $ r   )r
   r~   r   r‚   s    r   Úbuild_ufuncÚCUDAVectorize.build_ufuncÖ   s   € Ü" 4§>¡>°4·;±;Ó?Ð?r   c                 ó   • [         $ r   )Úvectorizer_stager_sourcer‚   s    r   Ú_kernel_templateÚCUDAVectorize._kernel_templateÙ   s   € ä'Ð'r   rI   N)r   rD   rE   rF   rº   rÅ   rË   rÎ   r”   rÒ   rH   rI   r   r   r¯   r¯   È   s,   † òJòòò@ð ñ(ó ó(r   r¯   zy
def __gufunc_{name}({args}):
    __tid__ = __cuda__.grid(1)
    if __tid__ < {checkedarg}:
        __core__({argitems})
c                   ó6   • \ rS rSrS rS r\S 5       rS rSr	g)ÚCUDAGUFuncVectorizeéé   c                 ó–   • [         R                  " U R                  U R                  5      n[	        U R
                  UU R                  S9$ )N)r~   r   r   )r   ÚGUFuncEngineÚinputsigÚ	outputsigr{   r~   r   )r   r   s     r   rÎ   ÚCUDAGUFuncVectorize.build_ufuncê   s9   € Ü×)Ò)¨$¯-©-¸¿¹ÓHˆÜ#¨d¯n©nØ+1Ø+/¯;©;ñ8ð 	8r   c                 ó:   • [         R                  " U5      " U5      $ r   rÈ   rÉ   s      r   rË   Ú#CUDAGUFuncVectorize._compile_kernelð   s   € Ü�xŠx˜Œ}˜UÓ#Ð#r   c                 ó   • [         $ r   )Ú_gufunc_stager_sourcer‚   s    r   rÒ   Ú$CUDAGUFuncVectorize._kernel_templateó   s   € ä$Ð$r   c                 óÈ   • [         R                  " USS9" U R                  5      nU R                  R                  R                  5       nUR                  [         US.5        U$ )NT)r²   r½   )r   r´   r   Úpy_funcrÀ   rÁ   rÂ   )r   r¸   rÃ   Úglblss       r   rÅ   Ú CUDAGUFuncVectorize._get_globals÷   sN   € Ü—’˜# dÒ+¨D¯K©KÓ8ˆØ—‘×(Ñ(×-Ñ-Ó/ˆØ�‰¤$Ø"(ñ*ô 	+àˆr   rI   N)
r   rD   rE   rF   rÎ   rË   r”   rÒ   rÅ   rH   rI   r   r   rÕ   rÕ   é   s%   † ò8ò$ð ñ%ó ð%õr   rÕ   N)Únumbar   Únumpyr   r0   Ú
numba.cudar   Únumba.cuda.deviceufuncr   r   r   Úobjectr
   rK   r{   r   rÑ   ÚDeviceVectorizer¯   rß   ÚDeviceGUFuncVectorizerÕ   rI   r   r   Ú<module>rì      s†   ðÝ Ý #Ý "÷5ñ 5ôH˜&ô HôV!9˜?ô !9ôHMÐ+ô Mô2-M˜ô -Mð`Ð ô(�K×/Ñ/ô (ð2Ð ô˜+×;Ñ;õ r   