§
    ƒÛgÝ%  ã                   óò   — d dl mZmZmZmZmZmZmZ d dlm	Z	 d dl
mZ d dlmZmZmZ d dlZd dlZd dlZd dlZeZeZ G d„ de	¦  «        Z ej        ¦   «         Z G d„ d	¦  «        Z G d
„ d¦  «        Zd„ ZdS )é    )ÚbyrefÚc_charÚc_char_pÚc_intÚc_size_tÚc_void_pÚPOINTER)ÚIntEnum)Úconfig)Ú
NvrtcErrorÚNvrtcCompilationErrorÚNvrtcSupportErrorNc                   ó>   — e Zd ZdZdZdZdZdZdZdZ	dZ
d	Zd
ZdZdZdS )ÚNvrtcResultr   é   é   é   é   é   é   é   é   é	   é
   é   N)Ú__name__Ú
__module__Ú__qualname__ÚNVRTC_SUCCESSÚNVRTC_ERROR_OUT_OF_MEMORYÚ$NVRTC_ERROR_PROGRAM_CREATION_FAILUREÚNVRTC_ERROR_INVALID_INPUTÚNVRTC_ERROR_INVALID_PROGRAMÚNVRTC_ERROR_INVALID_OPTIONÚNVRTC_ERROR_COMPILATIONÚ%NVRTC_ERROR_BUILTIN_OPERATION_FAILUREÚ1NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATIONÚ/NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATIONÚ%NVRTC_ERROR_NAME_EXPRESSION_NOT_VALIDÚNVRTC_ERROR_INTERNAL_ERROR© ó    úN/granary/edev/quantum/lib/python3.11/site-packages/numba/cuda/cudadrv/nvrtc.pyr   r      sW   € € € € € Ø€MØ !ÐØ+,Ð(Ø !ÐØ"#ÐØ!"ÐØÐØ,-Ð)Ø89Ð5Ø67Ð3Ø,.Ð)Ø!#ÐÐÐr,   r   c                   ó4   — e Zd ZdZd„ Zed„ ¦   «         Zd„ ZdS )ÚNvrtcProgramzá
    A class for managing the lifetime of nvrtcProgram instances. Instances of
    the class own an nvrtcProgram; when an instance is deleted, the underlying
    nvrtcProgram is destroyed using the appropriate NVRTC API.
    c                 ó"   — || _         || _        d S ©N)Ú_nvrtcÚ_handle)ÚselfÚnvrtcÚhandles      r-   Ú__init__zNvrtcProgram.__init__+   s   € ØˆŒØˆŒˆˆr,   c                 ó   — | j         S r1   )r3   ©r4   s    r-   r6   zNvrtcProgram.handle/   s
   € àŒ|Ðr,   c                 óL   — | j         r| j                             | ¦  «         d S d S r1   )r3   r2   Údestroy_programr9   s    r-   Ú__del__zNvrtcProgram.__del__3   s2   € ØŒ<ð 	.ØŒK×'Ò'¨Ñ-Ô-Ð-Ð-Ð-ð	.ð 	.r,   N)r   r   r   Ú__doc__r7   Úpropertyr6   r<   r+   r,   r-   r/   r/   %   sW   € € € € € ðð ð
ð ð ð ðð ñ „Xðð.ð .ð .ð .ð .r,   r/   c                   ó@  — e Zd ZdZe ee¦  «         ee¦  «        feeeee ee¦  «         ee¦  «        fe ee¦  «        feee ee¦  «        fee ee	¦  «        feeefee ee	¦  «        feeefee ee	¦  «        feeefdœ
Z
dZd„ Zd„ Zd„ Zd„ Zd„ Zd	„ Zd
„ ZdS )ÚNVRTCaB  
    Provides a Pythonic interface to the NVRTC APIs, abstracting away the C API
    calls.

    The sole instance of this class is a process-wide singleton, similar to the
    NVVM interface. Initialization is protected by a lock and uses the standard
    (for Numba) open_cudalib function to load the NVRTC library.
    )
ÚnvrtcVersionÚnvrtcCreateProgramÚnvrtcDestroyProgramÚnvrtcCompileProgramÚnvrtcGetPTXSizeÚnvrtcGetPTXÚnvrtcGetCUBINSizeÚnvrtcGetCUBINÚnvrtcGetProgramLogSizeÚnvrtcGetProgramLogNc                 ó  — t           5  | j        €Úddlm} t                               | ¦  «        x| _        }	  |d¦  «        }n)# t          $ r}d | _        t          d¦  «        |‚d }~ww xY w|j         	                    ¦   «         D ]c\  }}t          ||¦  «        }|d         |_        |dd …         |_        t          j        |¦  «        ||dœd„¦   «         }t          |||¦  «         Œdd d d ¦  «         n# 1 swxY w Y   | j        S )Nr   )Úopen_cudalibr5   zNVRTC cannot be loadedr   )ÚfuncÚnamec                 óò   —  | |Ž }|t           j        k    rt          ¦   «         ‚|t           j        k    rC	 t          |¦  «        j        }n# t
          $ r	 d|› d�}Y nw xY wd|› d|› �}t          |¦  «        ‚d S )Nz"Unknown nvrtc_result (error code: ú)zFailed to call z: )r   r%   r   r   rN   Ú
ValueErrorr   )rM   rN   ÚargsÚerrorÚ
error_nameÚmsgs         r-   Úchecked_callz#NVRTC.__new__.<locals>.checked_callx   sº   € à $  d ˜Ø ¥KÔ$GÒGÐGÝ"7Ñ"9Ô"9Ð9Ø"¥kÔ&?Ò?Ð?ðHÝ-8¸Ñ-?Ô-?Ô-D 
 
øÝ#-ð Hð Hð Hð/GØ>Cð/Gð /Gð /G 
 
 
ðHøøøð #I°DÐ"HÐ"H¸JÐ"HÐ"H˜CÝ",¨S¡/¤/Ð1ð @Ð?s   µA
 Á
AÁA)Ú_nvrtc_lockÚ_NVRTC__INSTANCEÚnumba.cuda.cudadrv.libsrL   ÚobjectÚ__new__ÚOSErrorr   Ú_PROTOTYPESÚitemsÚgetattrÚrestypeÚargtypesÚ	functoolsÚwrapsÚsetattr)	ÚclsrL   ÚinstÚlibÚerN   ÚprotorM   rV   s	            r-   r[   zNVRTC.__new__g   sƒ  € Ýð 	6ð 	6ØŒ~Ð%Ø@Ð@Ð@Ð@Ð@Ð@Ý(.¯ª°sÑ(;Ô(;Ð;�” ðMØ&˜, wÑ/Ô/�C�CøÝð Mð Mð MØ%)�C”NÝ+Ð,DÑEÔEÈ1ÐLøøøøðMøøøð
 $(Ô#3×#9Ò#9Ñ#;Ô#;ð 6ð 6‘K�D˜%Ý" 3¨Ñ-Ô-�DØ#(¨¤8�D”LØ$)¨!¨"¨"¤I�D”Må”_ TÑ*Ô*Ø15¸Dð 2ð 2ð 2ð 2ñ +Ô*ð2õ ˜D $¨Ñ5Ô5Ð5Ð5ð=	6ð 	6ð 	6ñ 	6ô 	6ð 	6ð 	6ð 	6ð 	6ð 	6ð 	6øøøð 	6ð 	6ð 	6ð 	6ð@ Œ~Ðs5   ˆ/C6¸AÁC6Á
A*ÁA%Á%A*Á*B C6Ã6C:Ã=C:c                 ó¶   — t          ¦   «         }t          ¦   «         }|                      t          |¦  «        t          |¦  «        ¦  «         |j        |j        fS )zB
        Get the NVRTC version as a tuple (major, minor).
        )r   rA   r   Úvalue)r4   ÚmajorÚminors      r-   Úget_versionzNVRTC.get_versionŠ   sG   € õ ‘”ˆÝ‘”ˆØ×Ò�% ™,œ,­¨e©¬Ñ5Ô5Ð5ØŒ{˜EœKÐ'Ð'r,   c                 ó0  — t          |t          ¦  «        r|                     ¦   «         }t          |t          ¦  «        r|                     ¦   «         }t          ¦   «         }|                      t          |¦  «        ||ddd¦  «         t          | |¦  «        S )z@
        Create an NVRTC program with managed lifetime.
        r   N)Ú
isinstanceÚstrÚencodeÚnvrtc_programrB   r   r/   )r4   ÚsrcrN   r6   s       r-   Úcreate_programzNVRTC.create_program“   s�   € õ �c�3ÑÔð 	Ø—*’*‘,”,ˆCÝ�d�CÑ Ô ð 	!Ø—;’;‘=”=ˆDå‘”ˆð
 	×Ò¥ f¡¤¨s°D¸!¸TÀ4ÑHÔHÐHÝ˜D &Ñ)Ô)Ð)r,   c                 óä   — d„ |D ¦   «         }d„ |D ¦   «         }t           t          |¦  «        z  } ||Ž }	 |                      |j        t          |¦  «        |¦  «         dS # t          $ r Y dS w xY w)zÌ
        Compile an NVRTC program. Compilation may fail due to a user error in
        the source; this function returns ``True`` if there is a compilation
        error and ``False`` on success.
        c                 ó6   — g | ]}|                      ¦   «         ‘ŒS r+   )rr   ©Ú.0Úopts     r-   ú
<listcomp>z)NVRTC.compile_program.<locals>.<listcomp>¬   s    € Ð;Ð;Ð;¨C˜3Ÿ:š:™<œ<Ð;Ð;Ð;r,   c                 ó,   — g | ]}t          |¦  «        ‘ŒS r+   )r   rx   s     r-   r{   z)NVRTC.compile_program.<locals>.<listcomp>­   s   € ÐDÐDÐD¨S�8 C™=œ=ÐDÐDÐDr,   FT)r   ÚlenrD   r6   r   )r4   ÚprogramÚoptionsÚencoded_optionsÚoption_pointersÚc_options_typeÚ	c_optionss          r-   Úcompile_programzNVRTC.compile_program¤   s’   € ð <Ð;°7Ð;Ñ;Ô;ˆØDÐD°OÐDÑDÔDˆÝ"¥S¨¡\¤\Ñ1ˆØ"�N OÐ4ˆ	ð	Ø×$Ò$ W¤^µS¸±\´\À9ÑMÔMÐMØ�5øÝ$ð 	ð 	ð 	Ø�4�4ð	øøøs   ¶)A! Á!
A/Á.A/c                 óT   — |                       t          |j        ¦  «        ¦  «         dS )z+
        Destroy an NVRTC program.
        N)rC   r   r6   )r4   r~   s     r-   r;   zNVRTC.destroy_program¶   s(   € ð 	× Ò ¥ w¤~Ñ!6Ô!6Ñ7Ô7Ð7Ð7Ð7r,   c                 ó  — t          ¦   «         }|                      |j        t          |¦  «        ¦  «         t	          |j        z  ¦   «         }|                      |j        |¦  «         |j                             ¦   «         S )z9
        Get the compile log as a Python string.
        )r   rI   r6   r   r   rk   rJ   Údecode)r4   r~   Úlog_sizeÚlogs       r-   Úget_compile_logzNVRTC.get_compile_log¼   si   € õ ‘:”:ˆØ×#Ò# G¤NµE¸(±O´OÑDÔDÐDå˜œÑ&Ñ)Ô)ˆØ×Ò ¤°Ñ4Ô4Ð4àŒy×ÒÑ!Ô!Ð!r,   c                 ó  — t          ¦   «         }|                      |j        t          |¦  «        ¦  «         t	          |j        z  ¦   «         }|                      |j        |¦  «         |j                             ¦   «         S )z:
        Get the compiled PTX as a Python string.
        )r   rE   r6   r   r   rk   rF   r‡   )r4   r~   Úptx_sizeÚptxs       r-   Úget_ptxzNVRTC.get_ptxÈ   si   € õ ‘:”:ˆØ×Ò˜Wœ^­U°8©_¬_Ñ=Ô=Ð=å˜œÑ&Ñ)Ô)ˆØ×Ò˜œ¨Ñ-Ô-Ð-àŒy×ÒÑ!Ô!Ð!r,   )r   r   r   r=   Únvrtc_resultr	   r   rs   r   r   r]   rX   r[   rn   ru   r„   r;   rŠ   rŽ   r+   r,   r-   r@   r@   8   sy  € € € € € ðð ð & w w¨u¡~¤~°w°w¸u±~´~ÐFð  ,¨]¸HÀhØ$ g g¨hÑ&7Ô&7¸¸ÀÑ9JÔ9JðLð !-¨g¨g°mÑ.DÔ.DÐEð !-¨m¸UØ ' ¨Ñ 1Ô 1ð 3ð )¨-¸¸ÀÑ9JÔ9JÐKà$ m°XÐ>ð +¨M¸7¸7À8Ñ;LÔ;LÐMà&¨°xÐ@ð $0°Ø#* 7¨8Ñ#4Ô#4ð#6ð  ,¨]¸HÐEðA!ð !€KðH €Jð!ð !ð !ðF(ð (ð (ð*ð *ð *ð"ð ð ð$8ð 8ð 8ð
"ð 
"ð 
"ð
"ð 
"ð 
"ð 
"ð 
"r,   r@   c                 óN  — t          ¦   «         }|                     | |¦  «        }|\  }}d|› |› �}dt          j        › �}t          j                             t          j                             t          ¦  «        ¦  «        }	t          j                             |	¦  «        }
d|
› �}|||ddg}| 	                    ||¦  «        }| 
                    |¦  «        }|rd|› d|› �}t          |¦  «        ‚|rd|› d|› �}t          j        |¦  «         |                     |¦  «        }||fS )a~  
    Compile a CUDA C/C++ source to PTX for a given compute capability.

    :param src: The source code to compile
    :type src: str
    :param name: The filename of the source (for information only)
    :type name: str
    :param cc: A tuple ``(major, minor)`` of the compute capability
    :type cc: tuple
    :return: The compiled PTX and compilation log
    :rtype: tuple
    z--gpu-architecture=compute_z-Iz-rdcÚtruez+NVRTC Compilation failure whilst compiling z:

z$NVRTC log messages whilst compiling )r@   ru   r   ÚCUDA_INCLUDE_PATHÚosÚpathÚdirnameÚabspathÚ__file__r„   rŠ   r   ÚwarningsÚwarnrŽ   )rt   rN   Úccr5   r~   rl   rm   ÚarchÚincludeÚcudadrv_pathÚnumba_cuda_pathÚnumba_includer   Úcompile_errorr‰   rU   r�   s                    r-   Úcompiler¡   Õ   s@  € õ ‰GŒG€EØ×"Ò" 3¨Ñ-Ô-€Gð �L€Eˆ5Ø7¨Ð7°Ð7Ð7€DØ-•6Ô+Ð-Ð-€Gå”7—?’?¥2¤7§?¢?µ8Ñ#<Ô#<Ñ=Ô=€LÝ”g—o’o lÑ3Ô3€OØ*˜Ð*Ð*€MØ�W˜m¨V°VÐ<€Gð ×)Ò)¨'°7Ñ;Ô;€Mð ×
Ò
 Ñ
(Ô
(€Cð ð ØM¸TÐMÐMÈÐMÐMˆÝ˜‰oŒoÐð ð ØF°dÐFÐFÀÐFÐFˆÝŒ�cÑÔÐà
�-Š-˜Ñ
 Ô
 €CØ�ˆ8€Or,   )Úctypesr   r   r   r   r   r   r	   Úenumr
   Ú
numba.corer   Únumba.cuda.cudadrv.errorr   r   r   rb   r“   Ú	threadingr˜   rs   r�   r   ÚLockrW   r/   r@   r¡   r+   r,   r-   ú<module>r¨      sx  ðØ NÐ NÐ NÐ NÐ NÐ NÐ NÐ NÐ NÐ NÐ NÐ NÐ NÐ NÐ NÐ NÐ NÐ NØ Ð Ð Ð Ð Ð Ø Ð Ð Ð Ð Ð ð9ð 9ð 9ð 9ð 9ð 9ð 9ð 9ð 9ð 9ð Ð Ð Ð Ø 	€	€	€	Ø Ð Ð Ð Ø €€€ð €ð €ð$ð $ð $ð $ð $�'ñ $ô $ð $ð ˆiŒnÑÔ€ð.ð .ð .ð .ð .ñ .ô .ð .ð&Z"ð Z"ð Z"ð Z"ð Z"ñ Z"ô Z"ð Z"ðz/ð /ð /ð /ð /r,   