Ë
    èÿæi>  ã                   óî  — d dl mZ d dlmZ d dlmZmZmZmZm	Z	m
Z
 d dlmZmZmZmZmZmZ d dlmZ d dlmZmZmZ d dlmZ d dlmZmZmZ d d	lmZ d d
l m!Z! d dl"m#Z# d„ Z$ G d„ de«      Z% G d„ de«      Z&d„ Z' edd¬«       G d„ de«      «       Z( edd¬«       G d„ de«      «       Z) G d„ de«      Z*e	 	 	 d&d„«       Z+d„ Z,e	 	 	 d'd„«       Z-	 	 	 d(d„Z.	 	 d)d „Z/	 	 	 d*d!„Z0d"„ Z1d#„ Z2 G d$„ d%e3«      Z4y)+é    )Úir)ÚConcreteTemplate)ÚtypesÚtypingÚfuncdescÚconfigÚcompilerÚsigutils)Úsanitize_compile_result_entriesÚCompilerBaseÚDefaultPassBuilderÚFlagsÚOptionÚCompileResult)Úglobal_compiler_lock)ÚLoweringPassÚPassManagerÚregister_pass)ÚNumbaInvalidConfigWarning)ÚIRLegalizationÚNativeLoweringÚAnnotateTypes)Úwarn)Úget_current_device)ÚCUDACABICallConvc                 ó0   — | €y t        | t        «      sJ ‚| S ©N)Ú
isinstanceÚdict)Úxs    úh/Volumes/fast/ai/experiments/voice-extract-mac/.venv/lib/python3.12/site-packages/numba/cuda/compiler.pyÚ_nvvm_options_typer"      s!   € Ø€yØô ˜!œTÔ"Ð"Ð"Øˆó    c                   ó8   — e Zd Z eedd¬«      Z eedd¬«      Zy)Ú	CUDAFlagsNzNVVM options)ÚtypeÚdefaultÚdoczCompute Capability)Ú__name__Ú
__module__Ú__qualname__r   r"   Únvvm_optionsÚtupleÚcompute_capability© r#   r!   r%   r%      s+   „ ÙØØØô€Lñ
  ØØØ ôÑr#   r%   c                   ó   — e Zd Zed„ «       Zy)ÚCUDACompileResultc                 ó   — t        | «      S r   )Úid©Úselfs    r!   Úentry_pointzCUDACompileResult.entry_point9   s   € ä�$‹xˆr#   N)r)   r*   r+   Úpropertyr6   r/   r#   r!   r1   r1   8   s   „ Øñó ñr#   r1   c                  ó.   — t        | «      } t        di | ¤ŽS )Nr/   )r   r1   )Úentriess    r!   Úcuda_compile_resultr:   >   s   € Ü-¨gÓ6€GÜÑ'˜wÑ'Ð'r#   TF)Úmutates_CFGÚanalysis_onlyc                   ó   — e Zd ZdZd„ Zd„ Zy)ÚCUDABackendÚcuda_backendc                 ó.   — t        j                  | «       y r   ©r   Ú__init__r4   s    r!   rB   zCUDABackend.__init__H   ó   € Ü×Ñ˜dÕ#r#   c           
      ó2  — |d   }t        j                  |j                  g|j                  ¢­Ž }t	        |j
                  |j                  |j                  j                  |j                  |j                  |j                  ||j                  ¬«      |_        y)zH
        Back-end: Packages lowering output in a compile result
        Úcr)Útyping_contextÚtarget_contextÚtyping_errorÚtype_annotationÚlibraryÚcall_helperÚ	signatureÚfndescT)r   rL   Úreturn_typeÚargsr:   Ú	typingctxÚ	targetctxÚstatusÚfail_reasonrI   rJ   rK   rM   rE   )r5   ÚstateÚloweredrL   s       r!   Úrun_passzCUDABackend.run_passK   s{   € ð ˜‘+ˆÜ×$Ñ$ U×%6Ñ%6ÐD¸¿¹ÒDˆ	ä&Ø Ÿ?™?Ø Ÿ?™?ØŸ™×1Ñ1Ø!×1Ñ1Ø—M‘MØ×+Ñ+ØØ—>‘>ô	
ˆŒð r#   N)r)   r*   r+   Ú_namerB   rV   r/   r#   r!   r>   r>   C   s   „ ð €Eò$ór#   r>   c                   ó    — e Zd ZdZdZd„ Zd„ Zy)ÚCreateLibraryzå
    Create a CUDACodeLibrary for the NativeLowering pass to populate. The
    NativeLowering pass will create a code library if none exists, but we need
    to set it up with nvvm_options from the flags if they are present.
    Úcreate_libraryc                 ó.   — t        j                  | «       y r   rA   r4   s    r!   rB   zCreateLibrary.__init__i   rC   r#   c                 óô   — |j                   j                  «       }|j                  j                  }|j                  j
                  }|j                  ||¬«      |_        |j                  j                  «        y)N)r,   T)	rQ   ÚcodegenÚfunc_idÚfunc_qualnameÚflagsr,   rZ   rJ   Úenable_object_caching)r5   rT   r]   Únamer,   s        r!   rV   zCreateLibrary.run_passl   s\   € Ø—/‘/×)Ñ)Ó+ˆØ�}‰}×*Ñ*ˆØ—{‘{×/Ñ/ˆØ×.Ñ.¨tÀ,Ð.ÓOˆŒà�‰×+Ñ+Ô-àr#   N)r)   r*   r+   Ú__doc__rW   rB   rV   r/   r#   r!   rY   rY   _   s   „ ñð €Eò$ór#   rY   c                   ó   — e Zd Zd„ Zd„ Zy)ÚCUDACompilerc                 óÊ  — t         }t        d«      }|j                  | j                  «      }|j                  j                  |j                  «       |j                  | j                  «      }|j                  j                  |j                  «       | j                  | j                  «      }|j                  j                  |j                  «       |j                  «        |gS )NÚcuda)	r   r   Údefine_untyped_pipelinerT   ÚpassesÚextendÚdefine_typed_pipelineÚdefine_cuda_lowering_pipelineÚfinalize)r5   ÚdpbÚpmÚuntyped_passesÚtyped_passesÚlowering_passess         r!   Údefine_pipelineszCUDACompiler.define_pipelinesx   s¥   € Ü ˆÜ˜Ó ˆà×4Ñ4°T·Z±ZÓ@ˆØ
�	‰	×Ñ˜×.Ñ.Ô/à×0Ñ0°·±Ó<ˆØ
�	‰	×Ñ˜×,Ñ,Ô-à×<Ñ<¸T¿Z¹ZÓHˆØ
�	‰	×Ñ˜×/Ñ/Ô0à
�‰ŒØˆtˆr#   c                 ó  — t        d«      }|j                  t        d«       |j                  t        d«       |j                  t        d«       |j                  t
        d«       |j                  t        d«       |j                  «        |S )NÚcuda_loweringz$ensure IR is legal prior to loweringzannotate typeszcreate libraryznative loweringzcuda backend)r   Úadd_passr   r   rY   r   r>   rm   )r5   rT   ro   s      r!   rl   z*CUDACompiler.define_cuda_lowering_pipelineˆ   sh   € Ü˜Ó)ˆà
�‰”NØ:ô	<à
�‰”MÐ#3Ô4ð 	�‰”MÐ#3Ô4Ø
�‰”NÐ$5Ô6Ø
�‰”K Ô0à
�‰ŒØˆ	r#   N)r)   r*   r+   rs   rl   r/   r#   r!   re   re   w   s   „ òó r#   re   Nc	                 óè  — |€t        d«      ‚ddlm}	 |	j                  }
|	j                  }t        «       }d|_        d|_        d|_        |s|rd|_	        |rd|_
        |rd|_        nd|_        |rd|_        |rd|_        |r||_        ||_        ddlm}  |d	«      5  t%        j&                  |
|| |||i t(        ¬
«      }d d d «       j*                  }|j-                  «        |S # 1 sw Y   Œ'xY w)Nz#Compute Capability must be suppliedé   ©Úcuda_targetTÚpythonÚnumpyr   )Útarget_overriderg   )rP   rQ   ÚfuncrO   rN   r`   ÚlocalsÚpipeline_class)Ú
ValueErrorÚ
descriptorrz   rF   rG   r%   Ú
no_compileÚno_cpython_wrapperÚno_cfunc_wrapperÚ	debuginfoÚdbg_directives_onlyÚerror_modelÚforceinlineÚfastmathr,   r.   Únumba.core.target_extensionr}   r	   Úcompile_extrare   rJ   rm   )ÚpyfuncrN   rO   ÚdebugÚlineinfoÚinlinerŠ   r,   Úccrz   rP   rQ   r`   r}   ÚcresrJ   s                   r!   Úcompile_cudar“   ˜   s   € ð 
€zÜÐ>Ó?Ð?å'Ø×*Ñ*€IØ×*Ñ*€Iä‹K€Eà€EÔØ#€EÔØ!€EÔñ ‘ØˆŒáØ$(ˆÔ!áØ$ˆÕà#ˆÔáØ ˆÔÙØˆŒÙØ)ˆÔØ!€EÔõ <Ù	˜Õ	 Ü×%Ñ%°	Ø09Ø+1Ø+/Ø2=Ø,1Ø-/Ü5AôCˆ÷ 
!ð �l‰l€GØ×ÑÔà€K÷ 
!Ð	 ús   Â "C(Ã(C1c                 ó  — |j                   j                  |j                  › d�||¬«      }|j                  |«       |j                  }|j
                  }t        | «      }|j                  ||«      }	| j                  j                  |j
                  |«      }
| j                  d«      }t        j                  ||
|j                  «      }t        j                  ||	|«      }t        j                  |j                  d«      «      }| j                  |«      }|j!                  ||j"                  «      }| j                  j%                  |||||«      \  }}|j'                  |«       |j)                  |«       |j+                  «        |S )zš
    Wrap a Numba ABI function in a C ABI wrapper at the NVVM IR level.

    The C ABI wrapper will have the same name as the source Python function.
    Ú
_function_)Ú
entry_namer,   zcuda.cabi.wrapperÚ )r]   rZ   rb   Úadd_linking_libraryÚargtypesÚrestyper   Úget_function_typeÚ	call_convÚcreate_moduler   ÚFunctionÚllvm_func_nameÚ	IRBuilderÚappend_basic_blockÚget_arg_packerÚfrom_argumentsrO   Úcall_functionÚretÚadd_ir_modulerm   )ÚcontextÚlibrM   Úwrapper_function_namer,   rJ   r™   rš   Úc_call_convÚwrapfntyÚfntyÚwrapper_moduler~   ÚwrapfnÚbuilderÚarginfoÚcallargsÚ_Úreturn_values                      r!   Úcabi_wrap_functionr´   Ó   s^  € ð �k‰k×(Ñ(¨C¯H©H¨:°ZÐ)@Ø4IØ6Bð )ó D€Gð ×Ñ Ô$ð �‰€HØ�n‰n€GÜ" 7Ó+€KØ×,Ñ,¨W°hÓ?€HØ×Ñ×.Ñ.¨v¯~©~¸xÓH€Dð ×*Ñ*Ð+>Ó?€NÜ�;‰;�~ t¨V×-BÑ-BÓC€Dô
 �[‰[˜¨Ð3HÓI€FÜ�l‰l˜6×4Ñ4°RÓ8Ó9€Gà×$Ñ$ XÓ.€GØ×%Ñ% g¨v¯{©{Ó;€Hð ×'Ñ'×5Ñ5Ø��w ¨(ó4�O€A€|à‡K�K�Ôà×Ñ˜.Ô)Ø×ÑÔØ€Nr#   c           
      ó€  — |dvrt        d|› �«      ‚|dk(  r|st        d«      ‚|
dvrt        d|
› �«      ‚|r|rd}t        t        |«      «       |
dk(  }|	xs
 t        «       }	||rd	nd
dœ}|rd|d<   t	        j
                  |«      \  }}|xs t        j                  }t        | |||||||¬«      }|j                  j                  }|r |s|t        j                  k7  rt        d«      ‚|j                  }|rG|j                  }|dk(  rˆ|	j!                  d| j"                  «      }t%        |||j&                  ||«      }nR| j(                  }|j*                  }|j,                  }|j/                  |j                  |j&                  |||||«      \  }}|r|j1                  |¬«      }||fS |j3                  |¬«      }||fS )aÃ  Compile a Python function to PTX or LTO-IR for a given set of argument
    types.

    :param pyfunc: The Python function to compile.
    :param sig: The signature representing the function's input and output
                types. If this is a tuple of argument types without a return
                type, the inferred return type is returned by this function. If
                a signature including a return type is passed, the compiled code
                will include a cast from the inferred return type to the
                specified return type, and this function will return the
                specified return type.
    :param debug: Whether to include debug info in the compiled code.
    :type debug: bool
    :param lineinfo: Whether to include a line mapping from the compiled code
                     to the source code. Usually this is used with optimized
                     code (since debug mode would automatically include this),
                     so we want debug info in the LLVM IR but only the line
                     mapping in the final output.
    :type lineinfo: bool
    :param device: Whether to compile a device function.
    :type device: bool
    :param fastmath: Whether to enable fast math flags (ftz=1, prec_sqrt=0,
                     prec_div=, and fma=1)
    :type fastmath: bool
    :param cc: Compute capability to compile for, as a tuple
               ``(MAJOR, MINOR)``. Defaults to ``(5, 0)``.
    :type cc: tuple
    :param opt: Enable optimizations. Defaults to ``True``.
    :type opt: bool
    :param abi: The ABI for a compiled function - either ``"numba"`` or
                ``"c"``. Note that the Numba ABI is not considered stable.
                The C ABI is only supported for device functions at present.
    :type abi: str
    :param abi_info: A dict of ABI-specific options. The ``"c"`` ABI supports
                     one option, ``"abi_name"``, for providing the wrapper
                     function's name. The ``"numba"`` ABI has no options.
    :type abi_info: dict
    :param output: Type of output to generate, either ``"ptx"`` or ``"ltoir"``.
    :type output: str
    :return: (code, resty): The compiled code and inferred return type
    :rtype: tuple
    )ÚnumbaÚczUnsupported ABI: r·   z&The C ABI is not supported for kernels)ÚptxÚltoirzUnsupported output type: z{debug=True with opt=True (the default) is not supported by CUDA. This may result in a crash - set debug=False or opt=False.r¹   é   r   )rŠ   ÚoptNzgen-lto)rŽ   r�   rŠ   r,   r‘   z'CUDA kernel must have void return type.Úabi_name)r‘   )ÚNotImplementedErrorr   r   r   r
   Únormalize_signaturer   ÚCUDA_DEFAULT_PTX_CCr“   rL   rN   r   ÚvoidÚ	TypeErrorrG   rJ   Úgetr)   r´   rM   Ú__code__Úco_filenameÚco_firstlinenoÚprepare_cuda_kernelÚ	get_ltoirÚget_asm_str)r�   ÚsigrŽ   r�   ÚdevicerŠ   r‘   r»   ÚabiÚabi_infoÚoutputÚmsgÚltor,   rO   rN   r’   ÚrestyÚtgtr¨   Úwrapper_nameÚcodeÚfilenameÚlinenumÚkernels                            r!   Úcompiler×   ÿ   sè  € ð\ �.Ñ Ü!Ð$5°c°UÐ";Ó<Ð<à
ˆc‚z™&Ü!Ð"JÓKÐKàÐ%Ñ%Ü!Ð$=¸f¸XÐ"FÓGÐGá‘ð2ˆô 	Ô& sÓ+Ô,à�WÑ€CØÒ!œ4›6€Hð Ù‰q˜Qñ€Lñ
 Ø"&ˆ�YÑä ×4Ñ4°SÓ9Ñ€Dˆ+à	Ò	)Œv×)Ñ)€BÜ˜ ¨T¸Ø!)°HØ%1°bô:€Dð �N‰N×&Ñ&€Eá‘V ¬¯©Ò 3ÜÐAÓBÐBà
×
Ñ
€CáØ�l‰lˆØ�#Š:Ø#Ÿ<™<¨
°F·O±OÓDˆLÜ$ S¨#¨t¯{©{¸LØ%1ó3‰Cð �‰ˆØ×#Ñ#ˆØ×%Ñ%ˆà×-Ñ-¨d¯l©l¸D¿K¹KÈØ.6¸ÀhØ.5ó7‰ˆˆVñ Ø�}‰} ˆ}Ó#ˆð �ˆ;Ðð �‰ "ˆÓ%ˆØ�ˆ;Ðr#   c
                 óV   — t        «       j                  }
t        | ||||||
||||	¬«      S )zËCompile a Python function to PTX or LTO-IR for a given signature for the
    current device's compute capabilility. This calls :func:`compile` with an
    appropriate ``cc`` value for the current device.©	rŽ   r�   rÊ   rŠ   r‘   r»   rË   rÌ   rÍ   )r   r.   r×   )r�   rÉ   rŽ   r�   rÊ   rŠ   r»   rË   rÌ   rÍ   r‘   s              r!   Úcompile_for_current_devicerÚ   j  s7   € ô 
Ó	×	0Ñ	0€BÜ�6˜3 e°hÀvØ$¨°¸#Ø$¨Vô5ð 5r#   c
                 ó.   — t        | |||||||||	d¬«      S )zøCompile a Python function to PTX for a given signature. See
    :func:`compile`. The defaults for this function are to compile a kernel
    with the Numba ABI, rather than :func:`compile`'s default of compiling a
    device function with the C ABI.r¸   rÙ   )r×   )
r�   rÉ   rŽ   r�   rÊ   rŠ   r‘   r»   rË   rÌ   s
             r!   Úcompile_ptxrÜ   v  s)   € ô �6˜3 e°hÀvØ$¨°¸#Ø$¨Uô4ð 4r#   c	                 óT   — t        «       j                  }	t        | ||||||	|||¬«
      S )z‚Compile a Python function to PTX for a given signature for the current
    device's compute capabilility. See :func:`compile_ptx`.)rŽ   r�   rÊ   rŠ   r‘   r»   rË   rÌ   )r   r.   rÜ   )
r�   rÉ   rŽ   r�   rÊ   rŠ   r»   rË   rÌ   r‘   s
             r!   Úcompile_ptx_for_current_devicerÞ   �  s5   € ô
 
Ó	×	0Ñ	0€BÜ�v˜s¨%¸(Ø$¨x¸BÀCØ¨ô3ð 3r#   c                 ó0   — t        | ||«      j                  S r   )Ú declare_device_function_templateÚkey©rb   rš   r™   s      r!   Údeclare_device_functionrã   Œ  s   € Ü+¨D°'¸8ÓD×HÑHÐHr#   c                 ó&  ‡‡	— ddl m} |j                  }|j                  }t	        j
                  |g|¢­Ž Š	t        | ‰	«      Š G ˆˆ	fd„dt        «      }t        j                  | ||¬«      }|j                  ‰|«       |j                  ‰|«       |S )Nrx   ry   c                   ó   •— e Zd ZW ° ZW °gZy)úBdeclare_device_function_template.<locals>.device_function_templateN)r)   r*   r+   rá   Úcases)ÚextfnrÉ   s   €€r!   Údevice_function_templateræ   —  s   ø„ ÙˆÙ�‰r#   ré   râ   )r‚   rz   rF   rG   r   rL   ÚExternFunctionr   r   ÚExternalFunctionDescriptorÚinsert_user_function)
rb   rš   r™   rz   rP   rQ   ré   rM   rè   rÉ   s
           @@r!   rà   rà   �  sŠ   ù€ Ý'Ø×*Ñ*€IØ×*Ñ*€IÜ
×
Ñ
˜7Ð
. XÒ
.€CÜ˜4 Ó%€E÷Ô#3ô ô ×0Ñ0Ø˜7¨Xô7€Fà×"Ñ" 5Ð*BÔCØ×"Ñ" 5¨&Ô1à#Ð#r#   c                   ó   — e Zd Zd„ Zy)rê   c                 ó    — || _         || _        y r   )rb   rÉ   )r5   rb   rÉ   s      r!   rB   zExternFunction.__init__¤  s   € ØˆŒ	Øˆ�r#   N)r)   r*   r+   rB   r/   r#   r!   rê   rê   £  s   „ ór#   rê   )FFFFNN)	FFTFNTr·   Nr¸   )FFTFTr·   Nr¸   )FFFFNTr¶   N)FFFFTr¶   N)5Úllvmliter   Únumba.core.typing.templatesr   Ú
numba.corer   r   r   r   r	   r
   Únumba.core.compilerr   r   r   r   r   r   Únumba.core.compiler_lockr   Únumba.core.compiler_machineryr   r   r   Únumba.core.errorsr   Únumba.core.typed_passesr   r   r   Úwarningsr   Únumba.cuda.apir   Únumba.cuda.targetr   r"   r%   r1   r:   r>   rY   re   r“   r´   r×   rÚ   rÜ   rÞ   rã   rà   Úobjectrê   r/   r#   r!   Ú<module>rû      sC  ðÝ Ý 8ß J× J÷0÷ 0õ :÷Gñ Gå 7÷4ñ 4å Ý -Ý .òô
�ô 
ô:˜ô ò(ñ
 ˜4¨uÔ5ô�,ó ó 6ðñ6 ˜5°Ô6ô�Ló ó 7ðô.�<ô ðB ØBGØ<@Øò7ó ð7òt)ðX Ø=AØAEØògó ðgðT CHØ@DØ>Có	5ð BGØIMó4ð GLØEIØ9=ó3òIò$ô&�Võ r#   