U h3@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}|rl|sldd|d D}d }|r|d t|7}|d 7}n|d 7}|d7}tdd|Dst ||S)NZTRITON_LIBCUDA_PATHz/sbin/ldconfigz-pcSs g|]}d|kr|dqS) libcuda.so.1)split).0linerO/var/www/html/venv/lib/python3.8/site-packages/triton/backends/nvidia/driver.py sz libcuda_dirs..cSsg|]}tj|qSr)ospathdirname)r locrrrrsZLD_LIBRARY_PATHcSs&g|]}tjtj|dr|qS)r rrexistsjoin)r dirrrrrs: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'szlibcuda_dirs..) rgetenv subprocess check_outputdecode splitlinesr stranyAssertionError)Zenv_libcuda_pathZlibslocsdirsZenv_ld_library_pathmsgrrr libcuda_dirss"   r'cCs tftSN) libdevice_dirr'rrrr library_dirs+sr*c Cst|d}t|}||d}|dkrtx}tj |d}t |d}| |W5QRXt |||ttt}t |d }|j||ddd}W5QRXW5QRXddl} | j||} | j| } | j| | S) Nzutf-8z.sozmain.cwrbT)binaryr)hashlibsha256encode hexdigestrZget_filetempfileTemporaryDirectoryrrropenwriterr* include_dir librariesputreadimportlib.utilutilspec_from_file_locationmodule_from_specloader exec_module) srcnamekeycache cache_pathZtmpdirZsrc_pathfso importlibspecmodrrrcompile_module_from_src0s    .  rJcs$eZdZfddZddZZS) CudaUtilscs"t|dstt|||_|jS)Ninstance)hasattrsuperrK__new__rL)cls __class__rrrOJs zCudaUtils.__new__cCsPtttjtdd}|j|_|j|_|j |_ |j |_ |j |_ |j |_ dS)Nzdriver.cZ cuda_utils) rJrrrrr read_textZ load_binaryZget_device_propertiesZcuOccupancyMaxActiveClustersZset_printf_fifo_sizeZfill_1d_tma_descriptorZfill_2d_tma_descriptor)selfrIrrr__init__OszCudaUtils.__init__)__name__ __module__ __qualname__rOrU __classcell__rrrQrrKHs rKcCs8|ddkrdSddddddd d dd d d d d d d|S)Nr*Z CUdeviceptrint32_tint8_tint16_tint64_tuint32_tuint8_tuint16_tuint64_tfloatdouble)i1i8Zi16Zi32Zi64u1u8u16u32Zu64Zfp16Zbf16Zfp32Zf32Zfp64rtyrrr ty_to_cpp^s( rmc s>ddd|D}dddddfd d |D}d |}t|d krvddd d|Dnd}fdd |D}dt|d krd|nddddd|Dddfdd |Dd|d|dddd |Ddt|d kr0dddd|Dndd}|S)N, css$|]\}}t|d|VqdS)z argNrmr irlrrrrwsz make_launcher..cSs|ddkrdSt|S)NrrZ PyObject*rorkrrr_extracted_typeys z&make_launcher.._extracted_typec Ss"dddddddddd d d d |S) NOrEdlbhrqBHIK) rrrcrdlongr\r]r[r^r`rar_rbrrkrrr format_of~s z make_launcher..format_ofcsg|]}|qSrr)r rl)rsr~rrrsz!make_launcher..Z iiiKKOOOOrcss|]\}}d|VqdS)z&_argNrrprrrrscsg|]}|kr|qSrrr rq) constantsrrrsa  #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; rrp)rsrrrsz 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|]6\}}|ddkr6d|d|d|d|d ndqS)rrZzDevicePtrInfo ptr_infoz = getPointer(_argrnz); if (!ptr_infoz.valid) return NULL;rrrprrrr1sz; Py_BEGIN_ALLOW_THREADS; _launch(gridX, gridY, gridZ, num_warps, num_ctas, clusterDimX, clusterDimY, clusterDimZ, shared_memory, (CUstream)_stream, (CUfunction)_functioncss4|],\}}|ddkr"d|dnd|VqdS)rrZZptr_infoz.dev_ptrZ_argNrrprrrr3saW); 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 signatureidsZ arg_declsZ args_formatformatZ args_listparamsr@r)rsrr~r make_launchertsD,./  ,Jrc@seZdZddZddZdS) CudaLaunchercsdtdrjjnti}tdr,jnt}fddfdd|D}fddjD}t|||t d }|j |_ dS) NZids_of_const_exprsfnrcst|trjj|S|Sr() isinstancer!r arg_namesindex)rq)r@rrez'CudaLauncher.__init__..csi|]\}}||qSrrr rBvaluecst_keyrr fsz)CudaLauncher.__init__..csi|]\}}||qSrrrrrrrgsZ__triton_launcher) rMrZ constexprstuplerdictrrrrJlaunch)rTr@metadatarrrrIr)rr@rrUbs   zCudaLauncher.__init__cOs|j||dSr()r)rTargskwargsrrr__call__lszCudaLauncher.__call__N)rVrWrXrUrrrrrr`s rcs0eZdZfddZddZeddZZS) CudaDrivercst|_t|_tdSr()rKutilsrZ launcher_clsrNrU)rTrQrrrUrszCudaDriver.__init__cCs6|}||}|dd|d}d}td||S)Nr  r )Zget_current_deviceZget_device_capabilityr)rTZdeviceZ capabilityZ warp_sizerrrget_current_targetws  zCudaDriver.get_current_targetcCsddl}|jo|jjdkS)Nr)torchr Z is_availableversionZhip)rrrr is_active~szCudaDriver.is_active)rVrWrXrUr staticmethodrrYrrrQrrps r)! functoolsrr.rr2pathlibrZtriton.runtime.buildrZtriton.runtime.cacherZtriton.backends.compilerrZtriton.backends.driverrrrrealpath__file__rr6r)r7 lru_cacher'r*rJobjectrKrmrrrrrrrs0       m