Ë
    èÿæiÓ"  ã                   óð   — d dl mZ d dlmZ d dlmZ d dlmZm	Z	m
Z
  G d„ de«      Z G d„ de
«      Z G d	„ d
e	«      Z G d„ de«      ZdZ G d„ dej"                  «      ZdZ G d„ dej(                  «      Zy)é    )Úcuda)Úarray)Údeviceufunc)ÚUFuncMechanismÚGeneralizedUFuncÚGUFuncCallStepsc                   ó*   — e Zd ZdZd„ Zd„ Zdd„Zd„ Zy)ÚCUDAUFuncDispatcherzD
    Invoke the CUDA ufunc specialization for the given inputs.
    c                 ó4   — || _         |j                  | _        y ©N)Ú	functionsÚ__name__)ÚselfÚtypes_to_retty_kernelsÚpyfuncs      úk/Volumes/fast/ai/experiments/voice-extract-mac/.venv/lib/python3.12/site-packages/numba/cuda/vectorizers.pyÚ__init__zCUDAUFuncDispatcher.__init__   s   € Ø/ˆŒØŸ™ˆ�ó    c                 óD   — t         j                  | j                  ||«      S )a¦  
        *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__zCUDAUFuncDispatcher.__call__   s   € ô "×&Ñ& t§~¡~°t¸SÓAÐAr   c                 ó€  — t        t        | j                  j                  «       «      d   «      dk(  sJ d«       ‚|j                  dk(  sJ d«       ‚|j
                  d   }g }|dk(  rt        d«      ‚|dk(  r|d   S |xs t        j                  «       }|j                  «       5  t        j                  j                  j                  |«      r|}nt        j                  ||«      }| j                  |||«      }t        d|j                   ¬«      }|j#                  ||¬	«       d d d «       |d   S # 1 sw Y   d   S xY w)
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   Úreducez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�à—-‘-  X¨vÓ6ˆCä˜4 s§y¡yÔ1ˆCØ×Ñ˜S¨ÐÔ0÷ 'ð �1‰vˆ÷ 'ð �1‰vˆús   Â#B D0Ä0D=c                 óÆ  — |j                   d   }|dz  dk7  ri|j                  |dz
  «      \  }}|j                  |«       |j                  |«       | j                  |||«      }|j                  |«        | ||||¬«      S |j                  |dz  «      \  }}	|j                  |«       |j                  |	«        | ||	||¬«       |dz  dkD  r| j                  |||«      S |S )Nr   r   r   )r3   r    )r%   ÚsplitÚappendr,   )
r   r2   r1   r    r0   ÚfatcutÚthincutr3   ÚleftÚrights
             r   Ú__reducezCUDAUFuncDispatcher.__reduce;   sÖ   € Ø�I‰I�a‰LˆØˆq‰5�AŠ:Ø!Ÿi™i¨¨A©Ó.‰OˆF�Gà�O‰O˜FÔ#Ø�O‰O˜GÔ$à—-‘- ¨°&Ó9ˆCØ�O‰O˜CÔ Ù˜˜W¨#°fÔ=Ð=àŸ)™) A¨¡FÓ+‰KˆD�%à�O‰O˜DÔ!Ø�O‰O˜EÔ"á��u $¨vÕ6Ø�A‰v˜ŠzØ—}‘} T¨8°VÓ<Ð<à�r   N©r   )r   Ú
__module__Ú__qualname__Ú__doc__r   r   r5   r,   © r   r   r
   r
      s   „ ñò(òBóó:r   r
   c                   óH   ‡ — e Zd ZdgZˆ fd„Zd„ Zd„ Zd„ Zd„ Zd„ Z	d„ Z
ˆ xZS )	Ú_CUDAGUFuncCallStepsÚ_streamc                 óX   •— t         ‰| �  ||||«       |j                  dd«      | _        y )Nr    r   )Úsuperr   ÚgetrE   )r   ÚninÚnoutr   ÚkwargsÚ	__class__s        €r   r   z_CUDAGUFuncCallSteps.__init__X   s(   ø€ Ü‰Ñ˜˜d D¨&Ô1Ø—z‘z (¨AÓ.ˆ�r   c                 ó,   — t        j                  |«      S r   ©r   Úis_cuda_array©r   Úobjs     r   Úis_device_arrayz$_CUDAGUFuncCallSteps.is_device_array\   ó   € Ü×!Ñ! #Ó&Ð&r   c                 ó‚   — t         j                  j                  j                  |«      r|S t        j                  |«      S r   ©r   r(   r)   r*   Úas_cuda_arrayrP   s     r   Úas_device_arrayz$_CUDAGUFuncCallSteps.as_device_array_   ó2   € ô �<‰<×#Ñ#×3Ñ3°CÔ8ØˆJÜ×!Ñ! #Ó&Ð&r   c                 óD   — t        j                  || j                  ¬«      S ©Nr   )r   r+   rE   )r   Úhostarys     r   r+   z_CUDAGUFuncCallSteps.to_devicei   s   € Ü�~‰~˜g¨d¯l©lÔ;Ð;r   c                 ó@   — |j                  || j                  ¬«      }|S rZ   )r.   rE   )r   Údevaryr[   r3   s       r   Úto_hostz_CUDAGUFuncCallSteps.to_hostl   s    € Ø×!Ñ! '°$·,±,Ð!Ó?ˆØˆ
r   c                 óF   — t        j                  ||| j                  ¬«      S ©N)r%   r   r    )r   Údevice_arrayrE   )r   r%   r   s      r   Úallocate_device_arrayz*_CUDAGUFuncCallSteps.allocate_device_arrayp   s   € Ü× Ñ  u°EÀ$Ç,Á,ÔOÐOr   c                 óD   —  |j                  || j                  ¬«      |Ž  y rZ   )ÚforallrE   )r   ÚkernelÚnelemr   s       r   Úlaunch_kernelz"_CUDAGUFuncCallSteps.launch_kernels   s   € Ø1ˆ�‰�e D§L¡LˆÓ1°4Ò8r   )r   r?   r@   Ú	__slots__r   rR   rW   r+   r^   rb   rg   Ú__classcell__©rL   s   @r   rD   rD   S   s1   ø„ àð€Iô/ò'ò'ò<òòPö9r   rD   c                   ó:   ‡ — e Zd Zˆ fd„Zed„ «       Zd„ Zd„ Zˆ xZS )ÚCUDAGeneralizedUFuncc                 óH   •— |j                   | _         t        ‰| �	  ||«       y r   )r   rG   r   )r   Ú	kernelmapÚenginer   rL   s       €r   r   zCUDAGeneralizedUFunc.__init__x   s   ø€ ØŸ™ˆŒÜ‰Ñ˜ FÕ+r   c                 ó   — t         S r   )rD   ©r   s    r   Ú_call_stepsz CUDAGeneralizedUFunc._call_steps|   s   € ä#Ð#r   c                 ó„   — t         j                  j                  j                  |d|j                  |j
                  ¬«      S ©Nr>   ©r%   Ústridesr   Úgpu_data)r   r(   r)   ÚDeviceNDArrayr   rw   )r   Úaryr%   s      r   Ú_broadcast_scalar_inputz,CUDAGeneralizedUFunc._broadcast_scalar_input€   s9   € Ü�|‰|×'Ñ'×5Ñ5¸EØ>BØ<?¿I¹IØ?B¿|¹|ð 6ó Mð 	Mr   c                 óê   — t        |«      t        |j                  «      z
  }d|z  |j                  z   }t        j                  j
                  j                  |||j                  |j                  ¬«      S rt   )	r!   r%   rv   r   r(   r)   rx   r   rw   )r   ry   ÚnewshapeÚnewaxÚ
newstridess        r   Ú_broadcast_add_axisz(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?   r@   r   Úpropertyrr   rz   r   ri   rj   s   @r   rl   rl   w   s(   ø„ ô,ð ñ$ó ð$òMöMr   rl   c                   ó>   — e Zd ZdZdZd„ Zd„ Zd„ Zd„ Zd„ Z	d„ Z
d	„ Zy
)r   z%
    Provide CUDA specialization
    r   c                 ó0   —  |j                  ||¬«      |Ž  y rZ   )rd   )r   ÚfuncÚcountr    r   s        r   ÚlaunchzCUDAUFuncMechanism.launch–   s   € Ø)ˆ�‰�E &ˆÓ)¨4Ò0r   c                 ó,   — t        j                  |«      S r   rN   rP   s     r   rR   z"CUDAUFuncMechanism.is_device_array™   rS   r   c                 ó‚   — t         j                  j                  j                  |«      r|S t        j                  |«      S r   rU   rP   s     r   rW   z"CUDAUFuncMechanism.as_device_arrayœ   rX   r   c                 ó0   — t        j                  ||¬«      S rZ   )r   r+   )r   r[   r    s      r   r+   zCUDAUFuncMechanism.to_device¦   s   € Ü�~‰~˜g¨fÔ5Ð5r   c                 ó&   — |j                  |¬«      S rZ   )r.   )r   r]   r    s      r   r^   zCUDAUFuncMechanism.to_host©   s   € Ø×"Ñ"¨&Ð"Ó1Ð1r   c                 ó2   — t        j                  |||¬«      S r`   )r   ra   )r   r%   r   r    s       r   rb   z(CUDAUFuncMechanism.allocate_device_array¬   s   € Ü× Ñ  u°EÀ&ÔIÐIr   c                 óª  — t        t        |«      «      D �cg c](  }||j                  k\  s|j                  |   ||   k7  r|‘Œ* }}t        |«      t        |j                  «      z
  }dg|z  t	        |j
                  «      z   }|D ]  }d||<   Œ	 t        j                  j                  j                  |||j                  |j                  ¬«      S c c}w )Nr   ru   )Úranger!   r$   r%   r"   rv   r   r(   r)   rx   r   rw   )r   ry   r%   ÚaxÚ
ax_differsÚ
missingdimrv   s          r   Úbroadcast_devicez#CUDAUFuncMechanism.broadcast_device¯   sÊ   € Ü#(¬¨U«Ô#4ó 5Ñ#4˜RØ˜sŸx™xšØŸ™ 2™¨%°©)Ò3ò Ð#4ˆ
ð 5ô ˜“Z¤# c§i¡i£.Ñ0ˆ
Ø�#˜
Ñ"¤T¨#¯+©+Ó%6Ñ6ˆãˆBØˆG�BŠKð ô �|‰|×'Ñ'×5Ñ5¸EØ>EØ<?¿I¹IØ?B¿|¹|ð 6ó Mð 	Mùò5s   —-CN)r   r?   r@   rA   ÚDEFAULT_STREAMr…   rR   rW   r+   r^   rb   r�   rB   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                   ó4   — e Zd Zd„ Zd„ Zd„ Zd„ Zed„ «       Zy)ÚCUDAVectorizec                 ó°   —  t        j                  |dd¬«      | j                  «      }||j                  |j                     j
                  j                  fS )NT)ÚdeviceÚinline)r   Újitr   Ú	overloadsr   Ú	signatureÚreturn_type)r   ÚsigÚcudevfns      r   Ú_compile_corezCUDAVectorize._compile_coreÉ   sE   € Ø9”$—(‘(˜3 t°DÔ9¸$¿+¹+ÓFˆØ˜×)Ñ)¨#¯(©(Ñ3×=Ñ=×IÑIÐIÐIr   c                 ó~   — | j                   j                  j                  «       }|j                  t        |dœ«       |S )N©Ú__cuda__Ú__core__)r   Ú__globals__ÚcopyÚupdater   )r   ÚcorefnÚglbls      r   Ú_get_globalszCUDAVectorize._get_globalsÍ   s5   € Ø�{‰{×&Ñ&×+Ñ+Ó-ˆØ�‰¤Ø!'ñ)ô 	*àˆr   c                 ó,   — t        j                  |«      S r   ©r   r—   ©r   Úfnobjr›   s      r   Ú_compile_kernelzCUDAVectorize._compile_kernelÓ   s   € Ü�x‰x˜‹Ðr   c                 óB   — t        | j                  | j                  «      S r   )r
   rn   r   rq   s    r   Úbuild_ufunczCUDAVectorize.build_ufuncÖ   s   € Ü" 4§>¡>°4·;±;Ó?Ð?r   c                 ó   — t         S r   )Úvectorizer_stager_sourcerq   s    r   Ú_kernel_templatezCUDAVectorize._kernel_templateÙ   s   € ä'Ð'r   N)	r   r?   r@   r�   r§   r¬   r®   r€   r±   rB   r   r   r“   r“   È   s,   „ òJòòò@ð ñ(ó ñ(r   r“   zy
def __gufunc_{name}({args}):
    __tid__ = __cuda__.grid(1)
    if __tid__ < {checkedarg}:
        __core__({argitems})
c                   ó.   — e Zd Zd„ Zd„ Zed„ «       Zd„ Zy)ÚCUDAGUFuncVectorizec                 óš   — t        j                  | j                  | j                  «      }t	        | j
                  || j                  ¬«      S )N)rn   ro   r   )r   ÚGUFuncEngineÚinputsigÚ	outputsigrl   rn   r   )r   ro   s     r   r®   zCUDAGUFuncVectorize.build_ufuncê   s9   € Ü×)Ñ)¨$¯-©-¸¿¹ÓHˆÜ#¨d¯n©nØ+1Ø+/¯;©;ô8ð 	8r   c                 ó8   —  t        j                  |«      |«      S r   r©   rª   s      r   r¬   z#CUDAGUFuncVectorize._compile_kernelð   s   € ØŒt�x‰x˜‹}˜UÓ#Ð#r   c                 ó   — t         S r   )Ú_gufunc_stager_sourcerq   s    r   r±   z$CUDAGUFuncVectorize._kernel_templateó   s   € ä$Ð$r   c                 óÌ   —  t        j                  |d¬«      | j                  «      }| j                  j                  j                  «       }|j                  t         |dœ«       |S )NT)r•   rŸ   )r   r—   r   Úpy_funcr¢   r£   r¤   )r   r›   r¥   Úglblss       r   r§   z CUDAGUFuncVectorize._get_globals÷   sP   € Ø+”—‘˜# dÔ+¨D¯K©KÓ8ˆØ—‘×(Ñ(×-Ñ-Ó/ˆØ�‰¤$Ø"(ñ*ô 	+àˆr   N)r   r?   r@   r®   r¬   r€   r±   r§   rB   r   r   r³   r³   é   s%   „ ò8ò$ð ñ%ó ð%ór   r³   N)Únumbar   Únumpyr   r-   Ú
numba.cudar   Únumba.cuda.deviceufuncr   r   r   Úobjectr
   rD   rl   r   r°   ÚDeviceVectorizer“   rº   ÚDeviceGUFuncVectorizer³   rB   r   r   Ú<module>rÅ      s†   ðÝ Ý #Ý "÷5ñ 5ôH˜&ô HôV!9˜?ô !9ôHMÐ+ô Mô2-M˜ô -Mð`Ð ô(�K×/Ñ/ô (ð2Ð ô˜+×;Ñ;õ r   