
    J/Ph%                         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_FAILURE1NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION/NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION%NVRTC_ERROR_NAME_EXPRESSION_NOT_VALIDNVRTC_ERROR_INTERNAL_ERROR     X/var/www/html/test/jupyter/venv/lib/python3.11/site-packages/numba/cuda/cudadrv/nvrtc.pyr   r      sW        M !+,( !"#!",-)89567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"HJ"H"HC",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//CC M M M%)CN+,DEE1LM
 $(#3#9#9#;#; 6 6KD%"3--D#(8DL$)!""IDM_T**15D 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{EK''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     c3 	**,,CdC   	!;;==D
 	fsD!T4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     ;;;C3::<<;;;r,   c                 ,    g | ]}t          |          S r+   )r   rx   s     r-   r{   z)NVRTC.compile_program.<locals>.<listcomp>   s    DDDS8C==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DODDD"S\\1"NO4		$$W^S\\9MMM5$ 	 	 	44	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     ::##GNE(OO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^U8__===&))---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         &wwu~~wwu~~F  ,]Hh$ggh&7&79J9JL !-ggm.D.DE !-mU ' 1 1 3 )-9J9JK$mX> +M778;L;LM&x@ $0#*78#4#4#6  ,]HEA! !KH 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@    GGE""3--G LE57777D-6+--G7??27??8#<#<==Lgool33O***MWmVV<G ))'7;;M 


(
(C  MTMMMMoo  FdFFF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      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     				       $ $ $ $ $' $ $ $ in. . . . . . . .&Z" Z" Z" Z" Z" Z" Z" Z"z/ / / / /r,   