çg3ddlZddlZddlZddlZddlZddlmZddlmZddl m Z ddl m Z ddl mZejejeZejedgZejedZd gZejd Zejd Zd ZGd deZdZdZGddeZGddeZ dS)N)Path)_build)get_cache_manager) GPUTarget) GPUDriverincludelibcudactjd}|r|gStjddg}d|D}d|D}tjd}|r!|sd|dD}d }|r|d t|zz }|d z }n |d z }|d z }td|Ds J||S)NTRITON_LIBCUDA_PATHz/sbin/ldconfigz-pcJg|] }d|v|d!S) libcuda.so.1)split).0lines Y/var/www/html/ai-engine/env/lib/python3.11/site-packages/triton/backends/nvidia/driver.py z libcuda_dirs..s0 U U UnPT>T>TDJJLL >T>T>TcLg|]!}tj|"S)ospathdirname)rlocs rrz libcuda_dirs..s& 1 1 1SBGOOC 1 1 1rLD_LIBRARY_PATHcg|]A}tjtj|d?|BS)rrrexistsjoin)rdirs rrz libcuda_dirs..sAsssPRPWP\P\]`bpPqPqArArssssr:zlibcuda.so cannot found! z!Possible files are located at %s.z:Please create a symlink of libcuda.so to any of the files.zzlibcuda_dirs..'s@SSdrw~~bgll4@@AASSSSSSr) rgetenv subprocess check_outputdecode splitlinesrstrany)env_libcuda_pathlibslocsdirsenv_ld_library_pathmsgs r libcuda_dirsr2s4y!677" !!  "$4d#; < < C C E ED V U):): U U UD 1 1D 1 1 1D)$566t4tss288==sss &C ? 2SYY>> KK MM >> SSdSSS S SXXUXXXX Krc.tgtSN) libdevice_dirr2rrr library_dirsr6+s  +LNN ++rc tj|d}t |}||d}|t j5}tj |d}t|d5}| |dddn #1swxYwYt|||ttt }t|d5}|||dd}dddn #1swxYwYdddn #1swxYwYddl} | j||} | j| } | j| | S) Nzutf-8z.sozmain.cwrbT)binaryr)hashlibsha256encode hexdigestrget_filetempfileTemporaryDirectoryrrr openwriterr6 include_dir librariesputreadimportlib.utilutilspec_from_file_locationmodule_from_specloader exec_module) srcnamekeycache cache_pathtmpdirsrc_pathfso importlibspecmods rcompile_module_from_srcrZ0sO .G,, - - 7 7 9 9C c " "E4 --J  ( * * Lfw||FH55Hh$$                 h  YWWBb$ L1"YYqvvxxDdYKK  L L L L L L L L L L L L L L L  L L L L L L L L L L L L L L L > 1 1$ C CD . ) )$ / /CKC   JsZ71E (C > E  C E C =E .E = E  E E E E  E$'E$c$eZdZfdZdZxZS) CudaUtilsct|ds-tt|||_|jS)Ninstance)hasattrsuperr\__new__r^)cls __class__s rrazCudaUtils.__new__Js<sJ'' > C0088==CL|rc:tttjt dd}|j|_|j|_|j |_ |j |_ |j |_ |j |_ dS)Nzdriver.c cuda_utils) rZrrrr r read_text load_binaryget_device_propertiescuOccupancyMaxActiveClustersset_printf_fifo_sizefill_1d_tma_descriptorfill_2d_tma_descriptor)selfrYs r__init__zCudaUtils.__init__Os}%d27<<+L+L&M&M&W&W&Y&Y[ghh?%(%>",/,L)$'$<!&)&@#&)&@###r)__name__ __module__ __qualname__rarn __classcell__rcs@rr\r\HsN AAAAAAArr\cN|ddkrdSddddddd d dd d d d d d d|S)Nr* CUdeviceptrint32_tint8_tint16_tint64_tuint32_tuint8_tuint16_tuint64_tfloatdouble)i1i8i16i32i64u1u8u16u32u64fp16bf16fp32f32fp64rtys r ty_to_cppr^sZ !u||}    !  rc dd|D}d d d fd|D}d|z}t|dkr4ddd |Dznd}fd |D}d t|dkrd|zndd dd |Ddd fd|Dd|d|ddd|Ddt|dkr4ddd|Dzndd}|S)N, c3DK|]\}}t|d|VdS)z argNrrirs rr$z make_launcher..ws:SS2Yr]]3333SSSSSSrc<|ddkrdSt|S)Nrru PyObject*rrs r_extracted_typez&make_launcher.._extracted_typeys a5C<<;}}rc ,dddddddddd d d d |S) NOrUdlbhrBHIK) rrrlongrxryrwrzr|r}r{r~rrs r format_ofz make_launcher..format_of~s?      rc8g|]}|Srr)rrrrs rrz!make_launcher..s-WWWb99__R%8%899WWWr iiiKKOOOOrc3&K|] \}}d|V dS)z&_argNrrs rr$z make_launcher..s, L LB L L L L L Lrcg|]}|v| Srr)rr constantss rrz!make_launcher..s# @ @ @AQi-?-?a-?-?-?ra  #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[] = { c3 K|] }d|V dS)z&argNr)rrs rr$z make_launcher..s( < < < < < < < 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;  c8g|]\}}|d|dS)z _arg; r)rrrrs rrz!make_launcher..s8 O O OEAr## / /! / / / O O Orz 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 Ng|]"\}}|ddkrd|d|d|d|d nd#S)rruzDevicePtrInfo ptr_infoz = getPointer(_argrz); if (!ptr_infoz.valid) return NULL;rrrs rrz!make_launcher..1s~dddINIJLNoqrsotx{o{o{kqkkAkkkkTUkkkkBDdddrz; Py_BEGIN_ALLOW_THREADS; _launch(gridX, gridY, gridZ, num_warps, num_ctas, clusterDimX, clusterDimY, clusterDimZ, shared_memory, (CUstream)_stream, (CUfunction)_functionc3JK|]\}}|ddkrd|dnd|VdS)rruptr_infoz.dev_ptr_argNrrs rr$z make_launcher..3sdxdx]b]^`bABCDFIIIe{pqe{e{e{e{OYVWOYOYdxdxdxdxdxdxraW); 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; } )r itemsvalueslenkeys) r signatureids arg_decls args_formatformat args_listparamsrNrrs ` @@r make_launcherrts3 SSARARSSSSSI  ''WWWWWIDTDTDVDVWWWXXK ; &FPST]P^P^abPbPbtyy L L)//:K:K L L LLLLLhjIA @ @ @)) @ @ @FH \RUV_R`R`cdRdRd~BEN~N~Njl]H H ^yy < z'CudaLauncher.__init__..es/As9K9KRCF,221555QRrc.i|]\}}||SrrrrPvaluecst_keys r z)CudaLauncher.__init__..fs'MMMZS%WWS\\5MMMrc.i|]\}}||Srrrs rrz)CudaLauncher.__init__..gs'QQQZS%WWS\\5QQQr__triton_launcher) r_r constexprstuplerdictrrrrZlaunch)rmrNmetadatarrrrYrs ` @rrnzCudaLauncher.__init__bs#'#t:L:L%YSV%6%6RWRYRYZ%,S+%>%>JCMMDFF RRRRMMMM9??;L;LMMM QQQQ3=;N;N;P;PQQQ Iy#66%c+>??j rc |j|i|dSr4)r)rmargskwargss r__call__zCudaLauncher.__call__ls T$V$$$$$rN)rorprqrnrrrrrr`s2!!!%%%%%rrc:eZdZfdZdZedZxZS) CudaDriverct|_t|_t dSr4)r\utilsr launcher_clsr`rn)rmrcs rrnzCudaDriver.__init__rs2[[ ( rc|}||}|ddz|dz}d}td||S)Nr  r )get_current_deviceget_device_capabilityr)rmdevice capability warp_sizes rget_current_targetzCudaDriver.get_current_targetwsT((**//77 ]R'*Q-7  Y777rcXddl}|jo |jjduS)Nr)torchr is_availableversionhip)rs r is_activezCudaDriver.is_active~s. z&&((Hem.?4.GHr)rorprqrnr staticmethodrrrrss@rrrpsi 888II\IIIIIrr)! functoolsrr;r&r@pathlibrtriton.runtime.buildrtriton.runtime.cachertriton.backends.compilerrtriton.backends.driverrrrrealpath__file__r rDr5rE lru_cacher2r6rZobjectr\rrrrrrrrs ''''''222222......,,,,,, '//"'**844 5 5w||GY//0  We,, H 0,,,0AAAAAAAA,   ,iiiX % % % % %6 % % % IIIIIIIIIIr