
    xKg%                         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y)    )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y)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_FAILURE1NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION/NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION%NVRTC_ERROR_NAME_EXPRESSION_NOT_VALIDNVRTC_ERROR_INTERNAL_ERROR     \/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/cudadrv/nvrtc.pyr   r      sI    M !+,( !"#!",-)895673,.)!#r,   r   c                   ,    e Zd ZdZd Zed        Zd Zy)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                      || _         || _        y 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                 T    | j                   r| j                  j                  |        y y r1   )r3   r2   destroy_programr9   s    r-   __del__zNvrtcProgram.__del__3   s     <<KK''- r,   N)r   r   r   __doc__r7   propertyr6   r<   r+   r,   r-   r/   r/   %   s%    
  .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y)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        j                  |       x| _        }	  |d      }|j                  j                         D ]Q  \  }}t        ||      }|d   |_        |dd  |_        t        j                  |      ||dd       }t        |||       S d d d        | j                  S # t        $ r}d | _        t        d      |d }~ww xY w# 1 sw Y   | j                  S xY w)Nr   )open_cudalibr5   zNVRTC cannot be loadedr   )funcnamec                      | | }|t         j                  k(  r
t               |t         j                  k7  r)	 t        |      j                  }d| d| }t        |      y # t
        $ r	 d| d}Y %w xY w)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"k&?&??H-8-?-D-D
 %4D6J<"HC",S/1 @ $. H/>>CWA/G
Hs   A   A21A2)_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    ~~%@(.s(;;M&w/C $(#3#3#9#9#;KD%"3-D#(8DL$)!"IDM__T*15D 2 +2 D$5' $< @ ~~5  M%)CN+,DE1LM @ ~~s.   /C%CA.C%	C"
CC""C%%C9c                     t               }t               }| j                  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   s=     %,e5{{EKK''r,   c                     t        |t              r|j                         }t        |t              r|j                         }t               }| j	                  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`     c3**,CdC ;;=D
 	fsD!T4HD&))r,   c                    |D cg c]  }|j                          }}|D cg c]  }t        |       }}t        t        |      z  } || }	 | j                  |j                  t        |      |       yc c}w c c}w # t
        $ r Y y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.
        FT)rr   r   lenrD   r6   r   )r4   programoptionsoptencoded_optionsoption_pointersc_options_type	c_optionss           r-   compile_programzNVRTC.compile_program   s     4;;7C3::<7;4CDOS8C=OD"S\1"O4		$$W^^S\9M <D % 		s   A5A:&A? ?	B
Bc                 L    | j                  t        |j                               y)z+
        Destroy an NVRTC program.
        N)rC   r   r6   )r4   rx   s     r-   r;   zNVRTC.destroy_program   s     	  w~~!67r,   c                     t               }| j                  |j                  t        |             t	        |j
                  z         }| j                  |j                  |       |j
                  j                         S )z9
        Get the compile log as a Python string.
        )r   rI   r6   r   r   rk   rJ   decode)r4   rx   log_sizelogs       r-   get_compile_logzNVRTC.get_compile_log   s[     :##GNNE(OD&)4yy!!r,   c                     t               }| j                  |j                  t        |             t	        |j
                  z         }| j                  |j                  |       |j
                  j                         S )z:
        Get the compiled PTX as a Python string.
        )r   rE   r6   r   r   rk   rF   r   )r4   rx   ptx_sizeptxs       r-   get_ptxzNVRTC.get_ptx   s[     :W^^U8_=&)-yy!!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   s     &wu~wu~F  ,]Hh$gh&79JL !-gm.DE !-mU ' 1 3 )-9JK$mX> +M78;LM&x@ $0#*8#4#6  ,]HEA!KH J!F(*"$8
"
"r,   r@   c                 &   t               }|j                  | |      }|\  }}d| | }dt        j                   }t        j
                  j                  t        j
                  j                  t                    }	t        j
                  j                  |	      }
d|
 }|||ddg}|j                  ||      }|j                  |      }|rd| d| }t        |      |rd| d| }t        j                  |       |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   rx   rl   rm   archincludecudadrv_pathnumba_cuda_pathnumba_includery   compile_errorr   rU   r   s                    r-   compiler      s    GE""3-G LE5(w7D6++,-G77??277??8#<=Lggool3O)*MWmVV<G ))'7;M 


(C <TF%uMo 5dV5Fc
--
 C8Or,   )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      sr    N N N  9 9  	    $' $ inn. .&Z" Z"z/r,   