§
    ƒÛgõ*  ã                   óü  — 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
„ 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'e	 	 d"d„¦   «         Z(	 	 d#d„Z)d„ Z*d„ Z+ G d„ d e,¦  «        Z-dS )$é    )Ú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_devicec                 ó<   — | €d S t          | t          ¦  «        sJ ‚| S ©N)Ú
isinstanceÚdict)Úxs    úI/granary/edev/quantum/lib/python3.11/site-packages/numba/cuda/compiler.pyÚ_nvvm_options_typer       s)   € Ø€yØˆtõ ˜!�TÑ"Ô"Ð"Ð"Ð"Øˆó    c                   óF   — e Zd Z eedd¬¦  «        Z eedd¬¦  «        ZdS )Ú	CUDAFlagsNzNVVM options)ÚtypeÚdefaultÚdoczCompute Capability)Ú__name__Ú
__module__Ú__qualname__r   r    Únvvm_optionsÚtupleÚcompute_capability© r!   r   r#   r#      sQ   € € € € € Ø�6ØØØðñ ô €Lð
  ˜ØØØ ðñ ô ÐÐÐr!   r#   c                   ó$   — e Zd Zed„ ¦   «         ZdS )ÚCUDACompileResultc                 ó    — t          | ¦  «        S r   )Úid©Úselfs    r   Úentry_pointzCUDACompileResult.entry_point7   s   € å�$‰xŒxˆr!   N)r'   r(   r)   Úpropertyr4   r-   r!   r   r/   r/   6   s-   € € € € € Øðð ñ „Xðð ð r!   r/   c                  ó8   — t          | ¦  «        } t          di | ¤ŽS )Nr-   )r
   r/   )Úentriess    r   Úcuda_compile_resultr8   <   s#   € Ý-¨gÑ6Ô6€GÝÐ'Ð'˜wÐ'Ð'Ð'r!   TF)Úmutates_CFGÚanalysis_onlyc                   ó   — e Zd ZdZd„ Zd„ ZdS )ÚCUDABackendÚcuda_backendc                 ó.   — t          j        | ¦  «         d S r   ©r   Ú__init__r2   s    r   r@   zCUDABackend.__init__F   ó   € ÝÔ˜dÑ#Ô#Ð#Ð#Ð#r!   c           
      óÖ   — |d         }t          j        |j        g|j        ¢R Ž }t	          |j        |j        |j        j        |j	        |j
        |j        ||j        ¬¦  «        |_        dS )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   rJ   Úreturn_typeÚargsr8   Ú	typingctxÚ	targetctxÚstatusÚfail_reasonrG   rH   rI   rK   rC   )r3   ÚstateÚloweredrJ   s       r   Úrun_passzCUDABackend.run_passI   sr   € ð ˜”+ˆÝÔ$ UÔ%6ÐD¸¼ÐDÐDÐDˆ	å&Ø œ?Ø œ?ØœÔ1Ø!Ô1Ø”MØÔ+ØØ”>ð	
ñ 	
ô 	
ˆŒð ˆtr!   N)r'   r(   r)   Ú_namer@   rT   r-   r!   r   r<   r<   A   s9   € € € € € ð €Eð$ð $ð $ðð ð ð ð r!   r<   c                   ó"   — e Zd ZdZdZd„ Zd„ ZdS )Ú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        | ¦  «         d S r   r?   r2   s    r   r@   zCreateLibrary.__init__g   rA   r!   c                 óÒ   — |j                              ¦   «         }|j        j        }|j        j        }|                     ||¬¦  «        |_        |j                             ¦   «          dS )N)r*   T)	rO   ÚcodegenÚfunc_idÚfunc_qualnameÚflagsr*   rX   rH   Úenable_object_caching)r3   rR   r[   Únamer*   s        r   rT   zCreateLibrary.run_passj   s\   € Ø”/×)Ò)Ñ+Ô+ˆØŒ}Ô*ˆØ”{Ô/ˆØ×.Ò.¨tÀ,Ð.ÑOÔOˆŒàŒ×+Ò+Ñ-Ô-Ð-àˆtr!   N)r'   r(   r)   Ú__doc__rU   r@   rT   r-   r!   r   rW   rW   ]   sC   € € € € € ðð ð €Eð$ð $ð $ðð ð ð ð r!   rW   c                   ó   — e Zd Zd„ Zd„ ZdS )ÚCUDACompilerc                 ó²  — t           }t          d¦  «        }|                     | j        ¦  «        }|j                             |j        ¦  «         |                     | j        ¦  «        }|j                             |j        ¦  «         |                      | j        ¦  «        }|j                             |j        ¦  «         |                     ¦   «          |gS )NÚcuda)	r   r   Údefine_untyped_pipelinerR   ÚpassesÚextendÚdefine_typed_pipelineÚdefine_cuda_lowering_pipelineÚfinalize)r3   ÚdpbÚpmÚuntyped_passesÚtyped_passesÚlowering_passess         r   Údefine_pipelineszCUDACompiler.define_pipelinesv   s²   € Ý ˆÝ˜Ñ Ô ˆà×4Ò4°T´ZÑ@Ô@ˆØ
Œ	×Ò˜Ô.Ñ/Ô/Ð/à×0Ò0°´Ñ<Ô<ˆØ
Œ	×Ò˜Ô,Ñ-Ô-Ð-à×<Ò<¸T¼ZÑHÔHˆØ
Œ	×Ò˜Ô/Ñ0Ô0Ð0à
�Š‰ŒˆØˆtˆr!   c                 óZ  — t          d¦  «        }|                     t          d¦  «         |                     t          d¦  «         |                     t          d¦  «         |                     t
          d¦  «         |                     t          d¦  «         |                     ¦   «          |S )NÚcuda_loweringz$ensure IR is legal prior to loweringzannotate typeszcreate libraryznative loweringzcuda backend)r   Úadd_passr   r   rW   r   r<   rk   )r3   rR   rm   s      r   rj   z*CUDACompiler.define_cuda_lowering_pipeline†   s‘   € Ý˜Ñ)Ô)ˆà
�Š•NØ:ñ	<ô 	<ð 	<à
�Š•MÐ#3Ñ4Ô4Ð4ð 	�Š•MÐ#3Ñ4Ô4Ð4Ø
�Š•NÐ$5Ñ6Ô6Ð6Ø
�Š•K Ñ0Ô0Ð0à
�Š‰ŒˆØˆ	r!   N)r'   r(   r)   rq   rj   r-   r!   r   rc   rc   u   s2   € € € € € ðð ð ð ð ð ð ð r!   rc   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 ¦  «         n# 1 swxY w Y   |j        }|                     ¦   «          |S )Nz#Compute Capability must be suppliedé   ©Úcuda_targetTÚpythonÚnumpyr   )Útarget_overridere   )rN   rO   ÚfuncrM   rL   r^   ÚlocalsÚpipeline_class)Ú
ValueErrorÚ
descriptorrx   rD   rE   r#   Ú
no_compileÚno_cpython_wrapperÚno_cfunc_wrapperÚ	debuginfoÚdbg_directives_onlyÚerror_modelÚforceinlineÚfastmathr*   r,   Únumba.core.target_extensionr{   r   Úcompile_extrarc   rH   rk   )ÚpyfuncrL   rM   ÚdebugÚlineinfoÚinlinerˆ   r*   Úccrx   rN   rO   r^   r{   ÚcresrH   s                   r   Úcompile_cudar‘   –   s°  € ð 
€zÝÐ>Ñ?Ô?Ð?à'Ð'Ð'Ð'Ð'Ð'ØÔ*€IØÔ*€Iå‰KŒK€Eà€EÔØ#€EÔØ!€EÔð ð �ð ØˆŒàð )Ø$(ˆÔ!àð $Ø$ˆÔÐà#ˆÔàð !Ø ˆÔØð ØˆŒØð *Ø)ˆÔØ!€EÔð <Ð;Ð;Ð;Ð;Ð;Ø	ˆ˜Ñ	 Ô	 ð Cð CÝÔ%°	Ø09Ø+1Ø+/Ø2=Ø,1Ø-/Ý5AðCñ Cô CˆðCð Cð Cñ Cô Cð Cð Cð Cð Cð Cð Cøøøð Cð Cð Cð Cð Œl€GØ×ÒÑÔÐà€Ks   Â!"CÃCÃCc           
      óþ  — |r |rd}t          t          |¦  «        ¦  «         ||rdnddœ}	t          j        |¦  «        \  }
}|pt          j        }t          | ||
||||	|¬¦  «        }|j        j        }|r!|s|t          j
        k    rt          d¦  «        ‚|r|j        }nD|j        }| j        }|j        }|j        }|                     |j        |j        |||	||¦  «        \  }}|                     |¬¦  «        }||fS )a  Compile a Python function to PTX 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.
    :param debug: Whether to include debug info in the generated PTX.
    :type debug: bool
    :param lineinfo: Whether to include a line mapping from the generated PTX
                     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 but only the line
                     mapping in the final PTX.
    :type lineinfo: bool
    :param device: Whether to compile a device function. Defaults to ``False``,
                   to compile global kernel functions.
    :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
    :return: (ptx, resty): The PTX code and inferred return type
    :rtype: tuple
    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ˆ   Úopt)rŒ   r�   rˆ   r*   r�   z'CUDA kernel must have void return type.)r�   )r   r   r	   Únormalize_signaturer   ÚCUDA_DEFAULT_PTX_CCr‘   rJ   rL   r   ÚvoidÚ	TypeErrorrH   rE   Ú__code__Úco_filenameÚco_firstlinenoÚprepare_cuda_kernelrK   Úget_asm_str)r‹   ÚsigrŒ   r�   Údevicerˆ   r�   r”   Úmsgr*   rM   rL   r�   ÚrestyÚlibÚtgtÚcodeÚfilenameÚlinenumÚkernelÚptxs                        r   Úcompile_ptxr©   Ñ   sM  € ð< ð -�ð -ð2ˆõ 	Õ& sÑ+Ô+Ñ,Ô,Ð,ð Øˆˆqˆq˜Qðð €Lõ
 !Ô4°SÑ9Ô9Ñ€Dˆ+à	Ð	)�vÔ)€BÝ˜ ¨T¸Ø!)°HØ%1°bð:ñ :ô :€Dð ŒNÔ&€Eàð C�Vð C ­¬Ò 3Ð 3ÝÐAÑBÔBÐBàð 
7ØŒlˆˆàÔ!ˆØŒˆØÔ#ˆØÔ%ˆà×-Ò-¨d¬l¸D¼KÈØ.6¸ÀhØ.5ñ7ô 7‰ˆˆVð �/Š/˜Rˆ/Ñ
 Ô
 €CØ�ˆ:Ðr!   c           
      óV   — t          ¦   «         j        }t          | ||||||d¬¦  «        S )zÑCompile a Python function to PTX for a given set of argument types for
    the current device's compute capabilility. This calls :func:`compile_ptx`
    with an appropriate ``cc`` value for the current device.T)rŒ   r�   rŸ   rˆ   r�   r”   )r   r,   r©   )r‹   rž   rŒ   r�   rŸ   rˆ   r”   r�   s           r   Úcompile_ptx_for_current_devicer«     s>   € õ
 
Ñ	Ô	Ô	0€BÝ�v˜s¨%¸(Ø$¨x¸BÀDðJñ Jô Jð Jr!   c                 ó.   — t          | ||¦  «        j        S r   )Ú declare_device_function_templateÚkey©r`   ÚrestypeÚargtypess      r   Údeclare_device_functionr²     s   € Ý+¨D°'¸8ÑDÔDÔHÐHr!   c                 ó(  ‡‡	— ddl m} |j        }|j        }t	          j        |g|¢R Ž Š	t          | ‰	¦  «        Š G ˆˆ	fd„dt          ¦  «        }t          j	        | ||¬¦  «        }| 
                    ‰|¦  «         | 
                    ‰|¦  «         |S )Nrv   rw   c                   ó   •— e Zd Z” Z”gZdS )úBdeclare_device_function_template.<locals>.device_function_templateN)r'   r(   r)   r®   Úcases)Úextfnrž   s   €€r   Údevice_function_templaterµ   *  s   ø€ € € € € ØˆØ�ˆˆˆr!   r¸   r¯   )r€   rx   rD   rE   r   rJ   ÚExternFunctionr   r   ÚExternalFunctionDescriptorÚinsert_user_function)
r`   r°   r±   rx   rN   rO   r¸   rK   r·   rž   s
           @@r   r­   r­   #  sÚ   øø€ Ø'Ð'Ð'Ð'Ð'Ð'ØÔ*€IØÔ*€IÝ
Ô
˜7Ð
. XÐ
.Ð
.Ð
.€CÝ˜4 Ñ%Ô%€Eðð ð ð ð ð ð ð Õ#3ñ ô ð õ Ô0Ø˜7¨Xð7ñ 7ô 7€Fà×"Ò" 5Ð*BÑCÔCÐCØ×"Ò" 5¨&Ñ1Ô1Ð1à#Ð#r!   c                   ó   — e Zd Zd„ ZdS )r¹   c                 ó"   — || _         || _        d S r   )r`   rž   )r3   r`   rž   s      r   r@   zExternFunction.__init__7  s   € ØˆŒ	ØˆŒˆˆr!   N)r'   r(   r)   r@   r-   r!   r   r¹   r¹   6  s#   € € € € € ðð ð ð ð r!   r¹   )FFFFNN)FFFFNT)FFFFT).Ú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   r    r#   r/   r8   r<   rW   rc   r‘   r©   r«   r²   r­   Úobjectr¹   r-   r!   r   ú<module>rÈ      sF  ðØ 8Ð 8Ð 8Ð 8Ð 8Ð 8Ø JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ Jð0ð 0ð 0ð 0ð 0ð 0ð 0ð 0ð 0ð 0ð 0ð 0ð 0ð 0ð 0ð 0ð :Ð 9Ð 9Ð 9Ð 9Ð 9ðGð Gð Gð Gð Gð Gð Gð Gð Gð Gà 7Ð 7Ð 7Ð 7Ð 7Ð 7ð4ð 4ð 4ð 4ð 4ð 4ð 4ð 4ð 4ð 4à Ð Ð Ð Ð Ð Ø -Ð -Ð -Ð -Ð -Ð -ðð ð ð
ð 
ð 
ð 
ð 
�ñ 
ô 
ð 
ð:ð ð ð ð ˜ñ ô ð ð(ð (ð (ð
 €˜4¨uÐ5Ñ5Ô5ðð ð ð ð �,ñ ô ñ 6Ô5ðð6 €˜5°Ð6Ñ6Ô6ðð ð ð ð �Lñ ô ñ 7Ô6ðð.ð ð ð ð �<ñ ô ð ðB ØBGØ<@Øð7ð 7ð 7ñ Ôð7ðt ØAFØ-1ð@ð @ð @ñ Ôð@ðF GLØEIðJð Jð Jð JðIð Ið Ið$ð $ð $ð&ð ð ð ð �Vñ ô ð ð ð r!   