Ë
    èÿæiŽ/  ã                   óÄ   — d dl mZ d dlmZmZ d dlmZmZ ddlm	Z	m
Z
mZmZ d dlmZ d dlZd dlZd dlZdZd	„ Zd
„ Zd„ Z G d„ dej,                  e«      Z G d„ de«      Zy)é    )Úir)ÚconfigÚ	serialize)ÚCodegenÚCodeLibraryé   )ÚdevicesÚdriverÚnvvmÚruntime)Úget_cudalibNznvptx64-nvidia-cudac                 óN  — d }d }	 t        j                  «       \  }}t        |d«      5 }|j                  | «       d d d «       	 t	        j
                  dg|¢|‘dt        j                  t        j                  ¬«      }|j                  j                  d«      |�t        j                  |«       |�t        j                  |«       S S # 1 sw Y   ŒŽxY w# t        $ r}d}t        |«      |‚d }~ww xY w# |�t        j                  |«       |�t        j                  |«       w w xY w)NÚwbÚnvdisasmT)ÚcheckÚstdoutÚstderrztnvdisasm has not been found. You may need to install the CUDA toolkit and ensure that it is available on your PATH.
zutf-8)ÚtempfileÚmkstempÚopenÚwriteÚ
subprocessÚrunÚPIPEÚFileNotFoundErrorÚRuntimeErrorr   ÚdecodeÚosÚcloseÚunlink)ÚcubinÚflagsÚfdÚfnameÚfÚcpÚeÚmsgs           úg/Volumes/fast/ai/experiments/voice-extract-mac/.venv/lib/python3.12/site-packages/numba/cuda/codegen.pyÚrun_nvdisasmr*      s  € ð 
€BØ€EðÜ×$Ñ$Ó&‰	ˆˆEÜ�%˜Ô !Ø�G‰G�EŒN÷ ð	+Ü—‘ Ð ;¨eÐ ;°UÐ ;À4Ü'1§¡Ü'1§¡ô8ˆBð �y‰y×Ñ Ó(àˆ>Ü�H‰H�RŒLØÐÜ�I‰I�eÕð ÷! Ðûô !ò 	+ð5ˆCô ˜sÓ#¨Ð*ûð		+ûð ˆ>Ü�H‰H�RŒLØÐÜ�I‰I�eÕð úsE   †#C3 ©C»C3 Á:C Á>C3 ÃCÃC3 Ã	C0ÃC+Ã+C0Ã0C3 Ã31D$c                 ó    — dg}t        | |«      S )Nz-gi©r*   ©r!   r"   s     r)   Údisassemble_cubinr.   +   s   € àˆG€EÜ˜˜uÓ%Ð%ó    c                 ó    — dg}t        | |«      S )Nz-cfgr,   r-   s     r)   Údisassemble_cubin_for_cfgr1   1   s   € àˆH€EÜ˜˜uÓ%Ð%r/   c                   óÌ   ‡ — e Zd ZdZ	 	 dˆ fd„	Zed„ «       Zd„ Zd„ Zdd„Z	dd„Z
dd„Zd	„ Zd
„ Zdd„Zdd„Zd„ Zd„ Zd„ Zd„ Zed„ «       Zed„ «       Zd„ Zd„ Zed„ «       Zˆ xZS )ÚCUDACodeLibraryzÙ
    The CUDACodeLibrary generates PTX, SASS, cubins for multiple different
    compute capabilities. It also loads cubins to multiple devices (via
    get_cufunc), which may be of different compute capabilities.
    c                 ó  •— t         ‰| �  ||«       d| _        t        «       | _        t        «       | _        d| _        d| _        i | _        i | _	        i | _
        i | _        i | _        || _        |€i }|| _        || _        y)aŸ  
        codegen:
            Codegen object.
        name:
            Name of the function in the source.
        entry_name:
            Name of the kernel function in the binary, if this is a global
            kernel and not a device function.
        max_registers:
            The maximum register usage to aim for when linking.
        nvvm_options:
                Dict of options to pass to NVVM.
        NF)ÚsuperÚ__init__Ú_moduleÚsetÚ_linking_librariesÚ_linking_filesÚneeds_cudadevrtÚ
_llvm_strsÚ
_ptx_cacheÚ_ltoir_cacheÚ_cubin_cacheÚ_linkerinfo_cacheÚ_cufunc_cacheÚ_max_registersÚ_nvvm_optionsÚ_entry_name)ÚselfÚcodegenÚnameÚ
entry_nameÚmax_registersÚnvvm_optionsÚ	__class__s         €r)   r6   zCUDACodeLibrary.__init__>   s�   ø€ ô 	‰Ñ˜ $Ô'ð ˆŒô
 #&£%ˆÔô "›eˆÔà$ˆÔð ˆŒàˆŒàˆÔàˆÔà!#ˆÔàˆÔà+ˆÔØÐØˆLØ)ˆÔØ%ˆÕr/   c                 óŒ   — | j                   €(| j                  D �cg c]  }t        |«      ‘Œ c}| _         | j                   S c c}w ©N)r<   ÚmodulesÚstr©rE   Úmods     r)   Ú	llvm_strszCUDACodeLibrary.llvm_strso   s:   € à�?‰?Ð"Ø37·<²<Ó@±<¨Cœs 3�x°<Ñ@ˆDŒOØ�‰Ðùò As   ›Ac                 ó8   — dj                  | j                  «      S )Nz

)ÚjoinrR   )rE   s    r)   Úget_llvm_strzCUDACodeLibrary.get_llvm_stru   s   € Ø�{‰{˜4Ÿ>™>Ó*Ð*r/   c                 ó^   — |�|S t        j                  «       j                  }|j                  S rM   )r	   Úget_contextÚdeviceÚcompute_capability)rE   ÚccrX   s      r)   Ú
_ensure_cczCUDACodeLibrary._ensure_ccx   s,   € Øˆ>ØˆIä×$Ñ$Ó&×-Ñ-ˆØ×(Ñ(Ð(r/   c                 ó"  — | j                  |«      }| j                  j                  |d «      }|r|S t        j                  |Ž }| j
                  j                  «       }||d<   | j                  }t        j                  |fi |¤Ž}|j                  «       j                  d«      j                  «       }t        j                  r>t        d| j                  z  j                  dd«      «       t        |«       t        d«       || j                  |<   |S )NÚarchÚ zASSEMBLY %séP   Ú-zP================================================================================)r[   r=   Úgetr   Úget_arch_optionrC   ÚcopyrR   Ú
compile_irr   Ústripr   ÚDUMP_ASSEMBLYÚprintÚ_nameÚcenter)rE   rZ   Úptxesr]   ÚoptionsÚirsÚptxs          r)   Úget_asm_strzCUDACodeLibrary.get_asm_str   sà   € Ø�_‰_˜RÓ ˆà—‘×#Ñ# B¨Ó-ˆÙØˆLä×#Ñ# RÐ(ˆØ×$Ñ$×)Ñ)Ó+ˆØˆ�‰à�n‰nˆä�o‰o˜cÑ- WÑ-ˆð
 �j‰j‹l× Ñ  Ó(×.Ñ.Ó0ˆä×ÒÜ�= 4§:¡:Ñ-×5Ñ5°b¸#Ó>Ô?Ü�#ŒJÜ�(ŒOà!ˆ�‰˜Ñàˆ
r/   c                 ó6  — | j                  |«      }| j                  j                  |d «      }|�|S t        j                  |Ž }| j
                  j                  «       }||d<   d |d<   | j                  }t        j                  |fi |¤Ž}|| j                  |<   |S )Nr]   zgen-lto)	r[   r>   ra   r   rb   rC   rc   rR   rd   )rE   rZ   Últoirr]   rk   rl   s         r)   Ú	get_ltoirzCUDACodeLibrary.get_ltoirœ   s›   € Ø�_‰_˜RÓ ˆà×!Ñ!×%Ñ% b¨$Ó/ˆØÐØˆLä×#Ñ# RÐ(ˆØ×$Ñ$×)Ñ)Ó+ˆØˆ�‰Ø!ˆ�	Ñà�n‰nˆÜ—‘ Ñ/ wÑ/ˆØ %ˆ×Ñ˜"Ñàˆr/   c                 ó„  — | j                  |«      }| j                  j                  |d «      }|r|S t        j                  j                  | j                  |¬«      }|j                  r$| j                  |¬«      }|j                  |«       n1| j                  |¬«      }|j                  |j                  «       «       | j                  D ]  }|j                  |«       Œ | j                  r|j                  t!        dd¬«      «       |j#                  «       }|| j                  |<   |j$                  | j&                  |<   |S )N)rI   rZ   ©rZ   Ú	cudadevrtT)Ústatic)r[   r?   ra   r
   ÚLinkerÚnewrB   Últorq   Ú	add_ltoirrn   Úadd_ptxÚencoder:   Úadd_file_guess_extr;   r   ÚcompleteÚinfo_logr@   )rE   rZ   r!   Úlinkerrp   rm   Úpaths          r)   Ú	get_cubinzCUDACodeLibrary.get_cubin®   s
  € Ø�_‰_˜RÓ ˆà×!Ñ!×%Ñ% b¨$Ó/ˆÙØˆLä—‘×"Ñ"°×1DÑ1DÈÐ"ÓLˆà�:Š:Ø—N‘N b�NÓ)ˆEØ×Ñ˜UÕ#à×"Ñ" bÐ"Ó)ˆCØ�N‰N˜3Ÿ:™:›<Ô(à×'Ô'ˆDØ×%Ñ% dÕ+ð (à×ÒØ×%Ñ%¤k°+ÀdÔ&KÔLà—‘Ó!ˆØ %ˆ×Ñ˜"ÑØ%+§_¡_ˆ×Ñ˜rÑ"àˆr/   c                 óŽ  — | j                   €d}t        |«      ‚t        j                  «       }|j                  }| j
                  j                  |j                  d «      }|r|S | j                  |j                  ¬«      }|j                  |«      }|j                  | j                   «      }|| j
                  |j                  <   |S )NzLMissing entry_name - are you trying to get the cufunc for a device function?rs   )rD   r   r	   rW   rX   rA   ra   Úidr�   rY   Úcreate_module_imageÚget_function)rE   r(   ÚctxrX   Úcufuncr!   Úmodules          r)   Ú
get_cufunczCUDACodeLibrary.get_cufuncÉ   s´   € Ø×ÑÐ#ð+ˆCä˜sÓ#Ð#ä×!Ñ!Ó#ˆØ—‘ˆà×#Ñ#×'Ñ'¨¯	©	°4Ó8ˆÙØˆMà—‘ &×";Ñ";�Ó<ˆØ×(Ñ(¨Ó/ˆð ×$Ñ$ T×%5Ñ%5Ó6ˆð )/ˆ×Ñ˜6Ÿ9™9Ñ%àˆr/   c                 óX   — 	 | j                   |   S # t        $ r t        d|› �«      ‚w xY w)NzNo linkerinfo for CC )r@   ÚKeyError©rE   rZ   s     r)   Úget_linkerinfozCUDACodeLibrary.get_linkerinfoá   s:   € ð	9Ø×)Ñ)¨"Ñ-Ð-øÜò 	9ÜÐ2°2°$Ð7Ó8Ð8ð	9ús   ‚ ‘)c                 ó8   — t        | j                  |¬«      «      S ©Nrs   )r.   r�   rŒ   s     r)   Úget_sasszCUDACodeLibrary.get_sassç   s   € Ü  §¡°2 Ó!6Ó7Ð7r/   c                 ó8   — t        | j                  |¬«      «      S r�   )r1   r�   rŒ   s     r)   Úget_sass_cfgzCUDACodeLibrary.get_sass_cfgê   s   € Ü(¨¯©¸2¨Ó)>Ó?Ð?r/   c                 ó`   — | j                  «        | j                  �t        d«      ‚|| _        y )Nz(CUDACodeLibrary only supports one module)Ú_raise_if_finalizedr7   r   rP   s     r)   Úadd_ir_modulezCUDACodeLibrary.add_ir_moduleí   s+   € Ø× Ñ Ô"Ø�<‰<Ð#ÜÐIÓJÐJØˆ�r/   c                 óz   — |j                  «        | j                  «        | j                  j                  |«       y rM   )Ú_ensure_finalizedr”   r9   Úadd)rE   Úlibrarys     r)   Úadd_linking_libraryz#CUDACodeLibrary.add_linking_libraryó   s0   € Ø×!Ñ!Ô#ð
 	× Ñ Ô"à×Ñ×#Ñ# GÕ,r/   c                 ó:   — | j                   j                  |«       y rM   )r:   r˜   )rE   Úfilepaths     r)   Úadd_linking_filez CUDACodeLibrary.add_linking_fileý   s   € Ø×Ñ×Ñ Õ)r/   c                 ó|   — | j                   j                  D ]  }|j                  |k(  sŒ|c S  t        d|› d�«      ‚)Nz	Function z
 not found)r7   Ú	functionsrG   r‹   )rE   rG   Úfns      r)   r…   zCUDACodeLibrary.get_function   s<   € Ø—,‘,×(Ô(ˆBØ�w‰w˜$‹Ø’	ð )ô ˜ 4 &¨
Ð3Ó4Ð4r/   c                 ó„   — | j                   g| j                  D ��cg c]  }|j                  D ]  }|‘Œ Œ c}}z   S c c}}w rM   )r7   r9   rN   )rE   ÚlibrQ   s      r)   rN   zCUDACodeLibrary.modules  sK   € à—‘ˆ~°×0GÒ0Gô !9Ñ0G¨Ø,/¯K¬K Sò "%Ø,7ð "%Ð0Gò !9ñ 9ð 	9ùó !9s   œ<c                 ó„   — g }| j                   D ].  }|j                  |j                  «       |j                  |«       Œ0 |S rM   )r9   ÚextendÚlinking_librariesÚappend)rE   Úlibsr¢   s      r)   r¥   z!CUDACodeLibrary.linking_libraries  s>   € ð
 ˆØ×*Ô*ˆCØ�K‰K˜×-Ñ-Ô.Ø�K‰K˜Õð +ð ˆr/   c                 óÀ   — | j                  «        | j                  D ]8  }|j                  D ]'  }|j                  D ]  }|j                  rŒd|_        Œ Œ) Œ: d| _        y )NÚlinkonce_odrT)r”   r9   rN   rŸ   Úis_declarationÚlinkageÚ
_finalized)rE   r™   rQ   r    s       r)   ÚfinalizezCUDACodeLibrary.finalize  sW   € ð 	× Ñ Ô"ð ×.Ô.ˆGØ—”�ØŸ-œ-�BØ×,Ó,Ø%3˜�
ñ (ñ 'ð /ð ˆ�r/   c                 ó@  — | j                   rd}t        |«      ‚| j                  st        d«      ‚t        d| j                  | j
                  | j                  | j                  | j                  | j                  | j                  | j                  | j                  ¬«
      S )zÀ
        Reduce the instance for serialization. We retain the PTX and cubins,
        but loaded functions are discarded. They are recreated when needed
        after deserialization.
        z0Cannot pickle CUDACodeLibrary with linking filesz)Cannot pickle unfinalized CUDACodeLibraryN)
rF   rG   rH   rR   Ú	ptx_cacheÚcubin_cacheÚlinkerinfo_cacherI   rJ   r;   )r:   r   r¬   ÚdictrG   rD   rR   r=   r?   r@   rB   rC   r;   )rE   r(   s     r)   Ú_reduce_stateszCUDACodeLibrary._reduce_states0  s‹   € ð ×ÒØDˆCÜ˜sÓ#Ð#Ø�ŠÜÐJÓKÐKÜØØ—‘Ø×'Ñ'Ø—n‘nØ—o‘oØ×)Ñ)Ø!×3Ñ3Ø×-Ñ-Ø×+Ñ+Ø ×0Ñ0ô
ð 	
r/   c                 óŒ   —  | |||¬«      }||_         ||_        ||_        ||_        ||_        |	|_        |
|_        d|_        |S )z&
        Rebuild an instance.
        )rH   T)r<   r=   r?   r@   rB   rC   r;   r¬   )ÚclsrF   rG   rH   rR   r¯   r°   r±   rI   rJ   r;   Úinstances               r)   Ú_rebuildzCUDACodeLibrary._rebuildH  sW   € ñ �w °Ô<ˆà'ˆÔØ'ˆÔØ +ˆÔØ%5ˆÔ"à"/ˆÔØ!-ˆÔØ#2ˆÔ à"ˆÔàˆr/   )NNNrM   )Ú__name__Ú
__module__Ú__qualname__Ú__doc__r6   ÚpropertyrR   rU   r[   rn   rq   r�   r‰   r�   r�   r’   r•   rš   r�   r…   rN   r¥   r­   r³   Úclassmethodr·   Ú__classcell__)rK   s   @r)   r3   r3   7   s´   ø„ ñð FJØ"õ/&ðb ñó ðò
+ò)óó:ó$ò6ò09ó8ó@òò-ò*ò5ð ñ9ó ð9ð ñó ðòò4
ð0 ñó ôr/   r3   c                   ó,   — e Zd ZdZeZd„ Zd„ Zd„ Zd„ Z	y)ÚJITCUDACodegenz™
    This codegen implementation for CUDA only generates optimized LLVM IR.
    Generation of PTX code is done separately (see numba.cuda.compiler).
    c                  ó   — y rM   © )rE   Úmodule_names     r)   r6   zJITCUDACodegen.__init__g  ó   € Ør/   c                 ó¶   — t        j                  |«      }t        |_        t	        j
                  «       j                  |_        t	        j                  |«       |S rM   )r   ÚModuleÚCUDA_TRIPLEÚtripler   ÚNVVMÚdata_layoutÚadd_ir_version)rE   rG   Ú	ir_modules      r)   Ú_create_empty_modulez#JITCUDACodegen._create_empty_modulej  s?   € Ü—I‘I˜d“Oˆ	Ü&ˆ	ÔÜ $§	¡	£× 7Ñ 7ˆ	ÔÜ×Ñ˜IÔ&ØÐr/   c                  ó   — y rM   rÂ   )rE   rˆ   s     r)   Ú_add_modulezJITCUDACodegen._add_moduleq  rÄ   r/   c                 ó–   — t        j                  «       }|j                  j                  }t        j                  j                  «       |fS )zP
        Return a tuple unambiguously describing the codegen behaviour.
        )r	   rW   rX   rY   r   Úget_version)rE   r†   rZ   s      r)   Úmagic_tuplezJITCUDACodegen.magic_tuplet  s9   € ô ×!Ñ!Ó#ˆØ�Z‰Z×*Ñ*ˆÜ—‘×+Ñ+Ó-¨rÐ2Ð2r/   N)
r¸   r¹   rº   r»   r3   Ú_library_classr6   rÍ   rÏ   rÒ   rÂ   r/   r)   rÀ   rÀ   _  s"   „ ñð
 %€Nòòòó3r/   rÀ   )Úllvmliter   Ú
numba.corer   r   Únumba.core.codegenr   r   Úcudadrvr	   r
   r   r   Únumba.cuda.cudadrv.libsr   r   r   r   rÇ   r*   r.   r1   ÚReduceMixinr3   rÀ   rÂ   r/   r)   Ú<module>rÚ      sY   ðÝ ç (ß 3ß 3Ó 3Ý /ã 	Û Û ð $€òò6&ò&ôe�i×+Ñ+¨[ô eôP	3�Wõ 3r/   