o i%@sddlmZmZmZmZmZmZmZddlm Z ddl m Z ddl m Z mZmZddlZddlZddlZddlZeZeZGddde ZeZGdd d ZGd d d Zd d ZdS))byrefc_charc_char_pc_intc_size_tc_void_pPOINTER)IntEnum)config) NvrtcErrorNvrtcCompilationErrorNvrtcSupportErrorNc@s<eZdZdZdZdZdZdZdZdZ dZ d Z d Z d Z d Zd S) NvrtcResultr N)__name__ __module__ __qualname__ NVRTC_SUCCESSNVRTC_ERROR_OUT_OF_MEMORY$NVRTC_ERROR_PROGRAM_CREATION_FAILURENVRTC_ERROR_INVALID_INPUTNVRTC_ERROR_INVALID_PROGRAMNVRTC_ERROR_INVALID_OPTIONNVRTC_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_VALIDNVRTC_ERROR_INTERNAL_ERRORr)r)V/home/ubuntu/transcripts/venv/lib/python3.10/site-packages/numba/cuda/cudadrv/nvrtc.pyrsrc@s,eZdZdZddZeddZddZdS) 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. cCs||_||_dSN)_nvrtc_handle)selfnvrtchandler)r)r*__init__+s zNvrtcProgram.__init__cCs|jSr,)r.r/r)r)r*r1/szNvrtcProgram.handlecCs|jr |j|dSdSr,)r.r-destroy_programr3r)r)r*__del__3szNvrtcProgram.__del__N)rrr__doc__r2propertyr1r5r)r)r)r*r+%s   r+c @seZdZdZeeeeefeeeeeeeeefeeefeeeeefeeee feeefeeee feeefeeee feeefd Z dZ ddZ ddZ dd Zd d Zd d ZddZddZdS)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. ) nvrtcVersionnvrtcCreateProgramnvrtcDestroyProgramnvrtcCompileProgramnvrtcGetPTXSize nvrtcGetPTXnvrtcGetCUBINSize nvrtcGetCUBINnvrtcGetProgramLogSizenvrtcGetProgramLogNc Csto|jdur_ddlm}t||_}z|d}Wnty0}zd|_td|d}~ww|j D]1\}}t ||}|d|_ |dd|_ t |||ddd}t|||q6Wd|jSWd|jS1stwY|jS) Nr) open_cudalibr0zNVRTC cannot be loadedr)funcnamecWsl||}|tjkr t|tjkr4zt|j}Wnty'd|d}Ynwd|d|}t|dS)Nz"Unknown nvrtc_result (error code: )zFailed to call z: )rr#r rrE ValueErrorr )rDrEargserror error_namemsgr)r)r* checked_callxs    z#NVRTC.__new__..checked_call) _nvrtc_lock_NVRTC__INSTANCEnumba.cuda.cudadrv.libsrCobject__new__OSErrorr _PROTOTYPESitemsgetattrrestypeargtypes functoolswrapssetattr) clsrCinstliberEprotorDrLr)r)r*rQgs4            z NVRTC.__new__cCs,t}t}|t|t||j|jfS)zB Get the NVRTC version as a tuple (major, minor). )rr9rvalue)r/majorminorr)r)r* get_versions zNVRTC.get_versioncCsLt|tr |}t|tr|}t}|t|||dddt||S)z@ Create an NVRTC program with managed lifetime. rN) isinstancestrencode nvrtc_programr:rr+)r/srcrEr1r)r)r*create_programs   zNVRTC.create_programcCs`dd|D}dd|D}tt|}||}z ||jt||WdSty/YdSw)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. cSsg|]}|qSr))rf.0optr)r)r* z)NVRTC.compile_program..cSsg|]}t|qSr))rrjr)r)r*rmrnFT)rlenr<r1r )r/programoptionsencoded_optionsoption_pointersc_options_type c_optionsr)r)r*compile_programs  zNVRTC.compile_programcCs|t|jdS)z+ Destroy an NVRTC program. N)r;rr1)r/rpr)r)r*r4szNVRTC.destroy_programcC<t}||jt|t|j}||j||jS)z9 Get the compile log as a Python string. )rrAr1rrr`rBdecode)r/rplog_sizelogr)r)r*get_compile_log   zNVRTC.get_compile_logcCrw)z: Get the compiled PTX as a Python string. )rr=r1rrr`r>rx)r/rpptx_sizeptxr)r)r*get_ptxr|z NVRTC.get_ptx)rrrr6 nvrtc_resultrrrgrrrSrNrQrcrirvr4r{rr)r)r)r*r88s6    $#  r8cCst}|||}|\}}d||}dtj}tjtjt} tj| } d| } ||| ddg} | || } | |}| rOd|d|}t ||r^d|d|}t |||}||fS)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-rdctruez+NVRTC Compilation failure whilst compiling z: z$NVRTC log messages whilst compiling )r8rir CUDA_INCLUDE_PATHospathdirnameabspath__file__rvr{r warningswarnr)rhrEccr0rprarbarchinclude cudadrv_pathnumba_cuda_path numba_includerq compile_errorrzrKr~r)r)r*compiles&        r)ctypesrrrrrrrenumr numba.corer numba.cuda.cudadrv.errorr r r rXr threadingrrgrrLockrMr+r8rr)r)r)r*s $