o 灛i3@sddlZddlZddlZddlZddlZddlmZddlmZddl m Z ddl m Z ddl mZejejeZejedgZejedZd gZed d Zed d ZddZGdddeZddZddZGdddeZGdddeZ dS)N)Path)_build)get_cache_manager) GPUTarget) GPUDriverincludelibcudacCstd}|r |gStddg}dd|D}dd|D}td}|r6|s6dd|d D}d }|rG|d t|7}|d 7}n|d 7}|d7}tdd|Ds\J||S)NTRITON_LIBCUDA_PATHz/sbin/ldconfigz-pcSs g|] }d|vr|dqS) libcuda.so.1)split).0liner_/sda-disk/www/egybert/egybert_env/lib/python3.10/site-packages/triton/backends/nvidia/driver.py s z libcuda_dirs..cSsg|]}tj|qSr)ospathdirname)rlocrrrrLD_LIBRARY_PATHcSs&g|]}tjtj|dr|qS)r rrexistsjoin)rdirrrrrs&:zlibcuda.so cannot found! z!Possible files are located at %s.z:Please create a symlink of libcuda.so to any of the files.z'$zlibcuda_dirs..) rgetenv subprocess check_outputdecode splitlinesr strany)env_libcuda_pathlibslocsdirsenv_ld_library_pathmsgrrr libcuda_dirss"   r-cCs tgtSN) libdevice_dirr-rrrr library_dirs+s r0c Cs(t|d}t|}||d}|dur{tU}tj |d}t |d }| |Wdn1s;wYt |||ttt}t |d}|j||ddd}Wdn1sgwYWdn1svwYddl} | j||} | j| } | j| | S) Nzutf-8z.sozmain.cwrbT)binaryr)hashlibsha256encode hexdigestrget_filetempfileTemporaryDirectoryrrropenwriterr0 include_dir librariesputreadimportlib.utilutilspec_from_file_locationmodule_from_specloader exec_module) srcnamekeycache cache_pathtmpdirsrc_pathfso importlibspecmodrrrcompile_module_from_src0s(      rScs$eZdZfddZddZZS) CudaUtilscs"t|dstt|||_|jS)Ninstance)hasattrsuperrT__new__rU)cls __class__rrrXJs zCudaUtils.__new__cCsPtttjtdd}|j|_|j|_|j |_ |j |_ |j |_ |j |_ dS)Nzdriver.c cuda_utils) rSrrrrr read_text load_binaryget_device_propertiescuOccupancyMaxActiveClustersset_printf_fifo_sizefill_1d_tma_descriptorfill_2d_tma_descriptor)selfrRrrr__init__Os zCudaUtils.__init__)__name__ __module__ __qualname__rXre __classcell__rrrZrrTHs rTcCs8|ddkrdSddddddd d dd d d d d d d|S)Nr* CUdeviceptrint32_tint8_tint16_tint64_tuint32_tuint8_tuint16_tuint64_tfloatdouble)i1i8i16i32i64u1u8u16u32u64fp16bf16fp32f32fp64rtyrrr ty_to_cpp^s( rc s<ddd|D}dddddfd d |D}d |}t|d kr;ddd d|Dnd}fdd |D}dt|d krSd|nddddd|Dddfdd |Dd|d|dddd |Ddt|d krdddd|Dndd}|S)N, css&|]\}}t|d|VqdS)z argNrrirrrrrwrz make_launcher..cSs|ddkrdSt|S)Nrrj PyObject*rrrrr_extracted_typeys z&make_launcher.._extracted_typec Ss"dddddddddd d d d |S) NOrNdlbhrBHIK) rrtrulongrmrnrlrorqrrrprsrrrrr format_of~s z make_launcher..format_ofcsg|]}|qSrr)rr)rrrrrrz!make_launcher.. iiiKKOOOOrcss|] \}}d|VqdS)z&_argNrrrrrrscsg|]}|vr|qSrrrr) constantsrrrra  #include "cuda.h" #include #include #include static inline void gpuAssert(CUresult code, const char *file, int line) { if (code != CUDA_SUCCESS) { const char* prefix = "Triton Error [CUDA]: "; const char* str; cuGetErrorString(code, &str); char err[1024] = {0}; strcat(err, prefix); strcat(err, str); PyGILState_STATE gil_state; gil_state = PyGILState_Ensure(); PyErr_SetString(PyExc_RuntimeError, err); PyGILState_Release(gil_state); } } #define CUDA_CHECK(ans) { gpuAssert((ans), __FILE__, __LINE__); } typedef CUresult (*cuLaunchKernelEx_t)(const CUlaunchConfig* config, CUfunction f, void** kernelParams, void** extra); static cuLaunchKernelEx_t getLaunchKernelExHandle() { // Open the shared library void* handle = dlopen("libcuda.so.1", RTLD_LAZY); if (!handle) { PyErr_SetString(PyExc_RuntimeError, "Failed to open libcuda.so.1"); return NULL; } // Clear any existing error dlerror(); cuLaunchKernelEx_t cuLaunchKernelExHandle = (cuLaunchKernelEx_t)dlsym(handle, "cuLaunchKernelEx"); // Check for errors const char *dlsym_error = dlerror(); if (dlsym_error) { PyErr_SetString(PyExc_RuntimeError, "Failed to retrieve cuLaunchKernelEx from libcuda.so.1"); return NULL; } return cuLaunchKernelExHandle; } static void _launch(int gridX, int gridY, int gridZ, int num_warps, int num_ctas, int clusterDimX, int clusterDimY, int clusterDimZ, int shared_memory, CUstream stream, CUfunction functionz) { void *params[] = { css|]}d|VqdS)z&argNrrrrrrsa }; if (gridX*gridY*gridZ > 0) { if (num_ctas == 1) { CUDA_CHECK(cuLaunchKernel(function, gridX, gridY, gridZ, 32*num_warps, 1, 1, shared_memory, stream, params, 0)); } else { CUlaunchAttribute launchAttr[2]; launchAttr[0].id = CU_LAUNCH_ATTRIBUTE_CLUSTER_DIMENSION; launchAttr[0].value.clusterDim.x = clusterDimX; launchAttr[0].value.clusterDim.y = clusterDimY; launchAttr[0].value.clusterDim.z = clusterDimZ; launchAttr[1].id = CU_LAUNCH_ATTRIBUTE_CLUSTER_SCHEDULING_POLICY_PREFERENCE; launchAttr[1].value.clusterSchedulingPolicyPreference = CU_CLUSTER_SCHEDULING_POLICY_SPREAD; CUlaunchConfig config; config.gridDimX = gridX * clusterDimX; config.gridDimY = gridY * clusterDimY; config.gridDimZ = gridZ * clusterDimZ; config.blockDimX = 32 * num_warps; config.blockDimY = 1; config.blockDimZ = 1; config.sharedMemBytes = shared_memory; config.hStream = stream; config.attrs = launchAttr; config.numAttrs = 2; static cuLaunchKernelEx_t cuLaunchKernelExHandle = NULL; if (cuLaunchKernelExHandle == NULL) { cuLaunchKernelExHandle = getLaunchKernelExHandle(); } CUDA_CHECK(cuLaunchKernelExHandle(&config, function, params, 0)); } } } typedef struct _DevicePtrInfo { CUdeviceptr dev_ptr; bool valid; } DevicePtrInfo; static inline DevicePtrInfo getPointer(PyObject *obj, int idx) { DevicePtrInfo ptr_info; ptr_info.dev_ptr = 0; ptr_info.valid = true; if (PyLong_Check(obj)) { ptr_info.dev_ptr = PyLong_AsUnsignedLongLong(obj); return ptr_info; } if (obj == Py_None) { // valid nullptr return ptr_info; } PyObject *ptr = PyObject_GetAttrString(obj, "data_ptr"); if(ptr){ PyObject *empty_tuple = PyTuple_New(0); PyObject *ret = PyObject_Call(ptr, empty_tuple, NULL); Py_DECREF(empty_tuple); Py_DECREF(ptr); if (!PyLong_Check(ret)) { PyErr_SetString(PyExc_TypeError, "data_ptr method of Pointer object must return 64-bit int"); ptr_info.valid = false; return ptr_info; } ptr_info.dev_ptr = PyLong_AsUnsignedLongLong(ret); if(!ptr_info.dev_ptr) return ptr_info; uint64_t dev_ptr; int status = cuPointerGetAttribute(&dev_ptr, CU_POINTER_ATTRIBUTE_DEVICE_POINTER, ptr_info.dev_ptr); if (status == CUDA_ERROR_INVALID_VALUE) { PyErr_Format(PyExc_ValueError, "Pointer argument (at %d) cannot be accessed from Triton (cpu tensor?)", idx); ptr_info.valid = false; } ptr_info.dev_ptr = dev_ptr; Py_DECREF(ret); // Thanks ChatGPT! return ptr_info; } PyErr_SetString(PyExc_TypeError, "Pointer argument must be either uint64 or have data_ptr method"); ptr_info.valid = false; return ptr_info; } static PyObject* launch(PyObject* self, PyObject* args) { int gridX, gridY, gridZ; uint64_t _stream; uint64_t _function; PyObject *launch_enter_hook = NULL; PyObject *launch_exit_hook = NULL; PyObject *kernel_metadata = NULL; PyObject *launch_metadata = NULL;  cs$g|]\}}|d|dqS)z _arg; rr)rrrrs$z if(!PyArg_ParseTuple(args, "z", &gridX, &gridY, &gridZ, &_stream, &_function, &kernel_metadata, &launch_metadata, &launch_enter_hook, &launch_exit_hook ai)) { return NULL; } int num_warps, num_ctas, shared_memory, clusterDimX, clusterDimY, clusterDimZ; if (!PyArg_ParseTuple(kernel_metadata, "iiiiii", &num_warps, &num_ctas, &shared_memory, &clusterDimX, &clusterDimY, &clusterDimZ)) { PyErr_SetString(PyExc_TypeError, "kernel_metadata must be a tuple"); return NULL; } // extract launch metadata if (launch_enter_hook != Py_None){ PyObject* args = Py_BuildValue("(O)", launch_metadata); PyObject* ret = PyObject_CallObject(launch_enter_hook, args); Py_DECREF(args); if (!ret) return NULL; } // raise exception asap rc Ss>g|]\}}|ddkrd|d|d|d|d ndqS)rrjzDevicePtrInfo ptr_infoz = getPointer(_argrz); if (!ptr_infoz.valid) return NULL;rrrrrrr1s>z; Py_BEGIN_ALLOW_THREADS; _launch(gridX, gridY, gridZ, num_warps, num_ctas, clusterDimX, clusterDimY, clusterDimZ, shared_memory, (CUstream)_stream, (CUfunction)_functioncss6|]\}}|ddkrd|dnd|VqdS)rrjptr_infoz.dev_ptr_argNrrrrrr3s4aW); Py_END_ALLOW_THREADS; if (PyErr_Occurred()) { return NULL; } if(launch_exit_hook != Py_None){ PyObject* args = Py_BuildValue("(O)", launch_metadata); PyObject* ret = PyObject_CallObject(launch_exit_hook, args); Py_DECREF(args); if (!ret) return NULL; } // return None Py_INCREF(Py_None); return Py_None; } static PyMethodDef ModuleMethods[] = { {"launch", launch, METH_VARARGS, "Entry point for all kernels with this signature"}, {NULL, NULL, 0, NULL} // sentinel }; static struct PyModuleDef ModuleDef = { PyModuleDef_HEAD_INIT, "__triton_launcher", NULL, //documentation -1, //size ModuleMethods }; PyMODINIT_FUNC PyInit___triton_launcher(void) { PyObject *m = PyModule_Create(&ModuleDef); if(m == NULL) { return NULL; } PyModule_AddFunctions(m, ModuleMethods); return m; } )ritemsvalueslenkeys) r signatureids arg_decls args_formatformat args_listparamsrGr)rrrr make_launchertsD,./ * Jrc@seZdZddZddZdS) CudaLaunchercsdtdr jjnti}tdrjnt}fddfdd|D}fddjD}t|||t d }|j |_ dS) Nids_of_const_exprsfnrcst|tr jj|S|Sr.) isinstancer%r arg_namesindex)r)rGrresz'CudaLauncher.__init__..ci|] \}}||qSrrrrIvaluecst_keyrr fz)CudaLauncher.__init__..crrrrrrrrgr__triton_launcher) rVr constexprstuplerdictrrrrSlaunch)rdrGmetadatarrrrRr)rrGrrebs    zCudaLauncher.__init__cOs|j|i|dSr.)r)rdargskwargsrrr__call__lszCudaLauncher.__call__N)rfrgrhrerrrrrr`s rcs0eZdZfddZddZeddZZS) CudaDrivercst|_t|_tdSr.)rTutilsr launcher_clsrWre)rdrZrrrerszCudaDriver.__init__cCs6|}||}|dd|d}d}td||S)Nr  r )get_current_deviceget_device_capabilityr)rddevice capability warp_sizerrrget_current_targetws   zCudaDriver.get_current_targetcCsddl}|jo|jjduS)Nr)torchr is_availableversionhip)rrrr is_active~szCudaDriver.is_active)rfrgrhrer staticmethodrrirrrZrrps  r)! functoolsrr4r!r9pathlibrtriton.runtime.buildrtriton.runtime.cachertriton.backends.compilerrtriton.backends.driverrrrrealpath__file__rr=r/r> lru_cacher-r0rSobjectrTrrrrrrrrs2       m