o ٰi9@s4dZddlmZddlmZdddddddZdS)zPNVRTC (NVIDIA Runtime Compilation) utilities for compiling CUDA source to CUBIN.) annotations)Sequencez kernel.cuN)namearch extra_optssourcestrrr str | NonerSequence[str] | Nonereturnbytesc Csz ddlm}m}Wnty}ztd|d}~ww|durzX|d\}||jjkr4td||\}}||jjkrBd}| |j j |\}} ||jjkrYtd|| |j j |\}} ||jjkrptd|d| | }Wnt y}z td |d |d}~ww|t|t|ddd\}} ||jjkrtd |d |d g} |r| dd|D|| t| | \}||jjkr|| \} }| |jjkr|dkrd|}|| |\} | |jjkrd|d}n d|}nd|}|| t||| \}}||jjkr,|| td|d|}|| |\}||jjkrJ|| td||| |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-devicecSs"g|] }t|tr |n|qS) isinstancerencode).0optrrE/home/ubuntu/.local/lib/python3.10/site-packages/tvm_ffi/cpp/nvrtc.py s"z!nvrtc_compile.. 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.bindingsr r ImportError RuntimeErrorcuInitCUresult CUDA_SUCCESScuCtxGetDevicecuDeviceGetAttributeCUdevice_attribute,CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR,CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR ExceptionnvrtcCreateProgramrr nvrtcResult NVRTC_SUCCESSextendnvrtcCompileProgramlennvrtcGetProgramLogSizenvrtcGetProgramLogdecodenvrtcDestroyProgramnvrtcGetCUBINSize nvrtcGetCUBIN)rrrrr reresultdevicemajorminorprogopts result_loglog_sizelog_buf error_msg cubin_size cubin_bufrrr nvrtc_compiles9       "         r=) rrrrrr rr r r )__doc__ __future__rtypingrr=rrrrs