`i94dZddlmZddlmZddddddZdS)zPNVRTC (NVIDIA Runtime Compilation) utilities for compiling CUDA source to CUBIN.) annotations)Sequencez kernel.cuN)namearch extra_optssourcestrrr str | NonerSequence[str] | Nonereturnbytesc  ddlm}m}n"#t$r}t d|d}~wwxYw| |d\}||jjkrt d||\}}||jjkrd}| |j j |\}} ||jjkrt d|| |j j |\}} ||jjkrt d|d| | }n&#t$r}t d |d |d}~wwxYw|t|t|ddd\}} ||jjkrt d |d |zd g} |r| d|D|| t+| | \}||jjkr|| \} }| |jjkrR|dkrLd|z}|| |\} | |jjkrd|d}n d|}nd|}|| t ||| \}}||jjkr'|| t d|d|z}|| |\}||jjkr'|| t d||| |S)aCompile CUDA source code to CUBIN using NVRTC. This function uses the NVIDIA Runtime Compilation (NVRTC) library to compile CUDA C++ source code into a CUBIN binary that can be loaded and executed using the CUDA Driver API. Parameters ---------- source : str The CUDA C++ source code to compile. name : str, optional The name to use for the source file (for error messages). Default: "kernel.cu" arch : str, optional The target GPU architecture (e.g., "sm_75", "sm_80", "sm_89"). If not specified, attempts to auto-detect from the current GPU. extra_opts : Sequence[str], optional Additional compilation options to pass to NVRTC (e.g., ["-I/path/to/include", "-DDEFINE=1"]). Returns ------- bytes The compiled CUBIN binary data. Raises ------ RuntimeError If NVRTC compilation fails or CUDA bindings are not available. Example ------- .. code-block:: python from tvm_ffi.cpp import nvrtc cuda_source = ''' extern "C" __global__ void add_one(float* x, float* y, int n) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx < n) { y[idx] = x[idx] + 1.0f; } } ''' cubin_bytes = nvrtc.nvrtc_compile(cuda_source) # Use cubin_bytes with tvm_ffi.cpp.load_inline and embed_cubin parameter r)drivernvrtczBCUDA bindings not available. Install with: pip install cuda-pythonNz"Failed to initialize CUDA driver: z(Failed to get compute capability major: z(Failed to get compute capability minor: sm_z(Failed to auto-detect GPU architecture: z-. Please specify 'arch' parameter explicitly.z Failed to create NVRTC program: s--gpu-architecture=s-default-devicecdg|]-}t|tr|n|.S) isinstancer encode).0opts e/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/tvm_ffi/cpp/nvrtc.py z!nvrtc_compile..s3YYYsZS%9%9BSZZ\\\sYYY zNVRTC compilation failed: zutf-8z-NVRTC compilation failed (couldn't get log): zNVRTC compilation failed: z%Failed to get CUBIN size from NVRTC: z Failed to get CUBIN from NVRTC: ) cuda.bindingsrr ImportError RuntimeErrorcuInitCUresult CUDA_SUCCESScuCtxGetDevicecuDeviceGetAttributeCUdevice_attribute,CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR,CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR ExceptionnvrtcCreateProgramr r nvrtcResult NVRTC_SUCCESSextendnvrtcCompileProgramlennvrtcGetProgramLogSizenvrtcGetProgramLogdecodenvrtcDestroyProgramnvrtcGetCUBINSize nvrtcGetCUBIN)rrrrrreresultdevicemajorminorprogopts result_loglog_sizelog_buf error_msg cubin_size cubin_bufs r nvrtc_compilerAs2r/////////  P     |  a((IV555"#P#P#PQQQ$2244NFF555#77)VX^MFE555"#Vf#V#VWWW"77)VX^MFE555"#Vf#V#VWWW''''DD   >1>>>  ++CJJv,>,> 4@P@PRSUY[_``LFD "000FfFFGGG . D [ YYjYYYZZZ))$D 4@@IV "000$;;DAA H *8 8 8X\\XoG!44T7CCMZU.<<<S'..:Q:QSS TFTT =V==I !!$'''9%%%0066FJ "000 !!$'''K6KKLLLz!I##D)44IV "000 !!$'''FfFFGGG d### s( *%*C3D%% E/EE) rr rr rr rr r r )__doc__ __future__rtypingrrArrrrEsm"WV"""""" '+ SSSSSSSSr