sg3ddlZddlZddlZddlZddlZddlmZddlmZddl m Z ddl m Z ddl mZejj!ejj#eZejj'edgZejj'edZd gZej.d Zej.d Zd ZGd deZdZdZGddeZGddeZ y)N)Path)_build)get_cache_manager) GPUTarget) GPUDriverincludelibcudac tjd}|r|gStjddgj }|j Dcgc]}d|vs|j d}}|Dcgc]!}tjj|#}}tjd}|r^|s\|j dDcgc]B}tjjtjj|dsA|D}}d}|r|d t|zz }|d z }n |d z }|d z }td |DsJ||Scc}wcc}wcc}w)NTRITON_LIBCUDA_PATHz/sbin/ldconfigz-p libcuda.so.1LD_LIBRARY_PATH: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,Sdrww~~bggll4@ASsAA ) rgetenv subprocess check_outputdecode splitlinessplitrdirnamerrstrany) env_libcuda_pathlibslinelocslocdirsenv_ld_library_pathdirmsgs r libcuda_dirsr+sOyy!67 !!  " "$4d#; < C C ED*.): UnPT>TDJJL  UD U,0 1SBGGOOC 1D 1))$564288=sPRPWPWP\P\]`bpPqArss &C  2SY>> KK MM >> SdS SXUXX S K V 1ts EE7&E AEEc$tgtSN) libdevice_dirr+r library_dirsr1+s  +LN ++r0c tj|jdj}t |}|j |d}|t j5}tjj|d}t|d5}|j|dddt|||ttt }t|d5}|j#|j%|dd}ddddddddl} | j(j+||} | j(j-| } | j.j1| | S#1swYxYw#1swYuxYw#1swYyxYw) Nzutf-8z.sozmain.cwrbT)binaryr)hashlibsha256encode hexdigestrget_filetempfileTemporaryDirectoryrrropenwriterr1 include_dir librariesputreadimportlib.utilutilspec_from_file_locationmodule_from_specloader exec_module) srcnamekeycache cache_pathtmpdirsrc_pathfso importlibspecmods rcompile_module_from_srcrU0sM ..G, - 7 7 9C c "E4& -J  ( ( * Lfww||FH5Hh$   h  YWBb$ L1"YYqvvxD6dYK  L  L >> 1 1$ CD .. ) )$ /CKKC J   L L  L Ls<(-FE)'4F&E5F)E2 .F5E> :FF c$eZdZfdZdZxZS) CudaUtilscdt|dstt|||_|jS)Ninstance)hasattrsuperrW__new__rY)cls __class__s rr\zCudaUtils.__new__Js*sJ' C8=CL||r0cZtttjj t dj d}|j|_|j|_|j|_ |j|_ |j|_ |j|_ y)Nzdriver.c cuda_utils) rUrrrrr read_text load_binaryget_device_propertiescuOccupancyMaxActiveClustersset_printf_fifo_sizefill_1d_tma_descriptorfill_2d_tma_descriptor)selfrTs r__init__zCudaUtils.__init__Os~%d277<<+L&M&W&W&Y[gh??%(%>%>",/,L,L)$'$<$<!&)&@&@#&)&@&@#r0)__name__ __module__ __qualname__r\ri __classcell__r^s@rrWrWHs Ar0rWc>|ddk(ryddddddd d dd d d d d d d|S)Nr* CUdeviceptrint32_tint8_tint16_tint64_tuint32_tuint8_tuint16_tuint64_tfloatdouble)i1i8i16i32i64u1u8u16u32u64fp16bf16fp32f32fp64r/tys r ty_to_cppr^sQ !u|   !  r0cdjd|jD}d}d}dj|jDcgc]}|||c}}d|z}t|dkDr)ddjd|jDznd} |j D cgc] } | |vs|  } } d t|dkDrd|zndd djd | Dd d j|jD cgc]\} }||d| dc}} d|d| ddj|jD cgc]\} }|ddk(rd| d| d| d| d nd c}} dt|dkDr)ddjd|jDzndd} | Scc}wcc} wcc}} wcc}} w)Nz, c3BK|]\}}t|d|yw)z argNrrirs rrz make_launcher..ws#S2Yr]O4s3Ssc*|ddk(ryt|S)Nrrp PyObject*rrs r_extracted_typez&make_launcher.._extracted_typeys a5C<}r0c &dddddddddd d d d |S) NOrPdlbhrBHIK) rrzr{longrsrtrrrurwrxrvryr/rs r format_ofz make_launcher..format_of~s:     r0 iiiKKOOOOrc3,K|] \}}d|yw)z&_argNr/rs rrz make_launcher..s LB5 Lsa  #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| yw)z&argNr/)rrs rrz make_launcher..s <4s 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;  z _argz; 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 rpzDevicePtrInfo ptr_infoz = getPointer(_argz); if (!ptr_infoz.valid) return NULL;z; Py_BEGIN_ALLOW_THREADS; _launch(gridX, gridY, gridZ, num_warps, num_ctas, clusterDimX, clusterDimY, clusterDimZ, shared_memory, (CUstream)_stream, (CUfunction)_functionc3HK|]\}}|ddk(rd|dnd|yw)rrpptr_infoz.dev_ptr_argNr/rs rrz make_launcher..3sdx]b]^`bABCDFIIgopqorrze{QUVWUXOYeYdxs "aW); 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) constants signatureids arg_declsrrr args_formatformat args_listrparamsrIs r make_launcherrts SARSSI  ''IDTDTDVWb9_R%89WXK ; &FPST]P^abPbtyy L)//:K LLLhjI#) @AQi-?a @F @.}\RUV_R`cdRd~BEN~Njl}mmyy n88Y__=N OEAr#$E!B / OPQR &x(RR[Q\](99R[RaRaRcdINIJLNoqrsotx{o{&qc);A3bCSTUSVVjkBDDdeffS|@I|JMN|NTX[_[d[ddxfofufufwdx[xTxTVSW)WH CR J_X AN P.ds G6 GGG #G ceZdZdZdZy) CudaLaunchercdtdrjjn ti}tdr jn t }fd}|j Dcic]\}}|||}}}jj Dcic]\}}|||}}}t|||td} | j|_ ycc}}wcc}}w)Nids_of_const_exprsfnrcrt|tr%jjj |S|Sr-) isinstancer r arg_namesindex)rrIs rz'CudaLauncher.__init__..es*As9KCFF,,2215QRr0__triton_launcher) rZr constexprstuplerdictrrrrUlaunch) rhrImetadatarrcst_keyrKvaluerrTs ` rrizCudaLauncher.__init__bs#'#t:LSVV%6%6RWRYZ%,S+%>CMMDF R;D??;LMZS%WS\5(M M;>==;N;N;PQZS%WS\5(Q QIy#6%c+>?jj NQs *CC$c(|j|i|yr-)r)rhargskwargss r__call__zCudaLauncher.__call__ls T$V$r0N)rjrkrlrirr/r0rrr`s !%r0rc4eZdZfdZdZedZxZS) CudaDrivercVt|_t|_t|yr-)rWutilsr launcher_clsr[ri)rhr^s rrizCudaDriver.__init__rs[ ( r0c~|j}|j|}|ddz|dz}d}td||S)Nr  r )get_current_deviceget_device_capabilityr)rhdevice capability warp_sizes rget_current_targetzCudaDriver.get_current_targetwsK((*//7 ]R'*Q-7  Y77r0crddl}|jjxr|jjduS)Nr)torchr is_availableversionhip)rs r is_activezCudaDriver.is_active~s,zz&&(Hemm.?.?4.GHr0)rjrkrlrir staticmethodrrmrns@rrrps# 8IIr0r)! functoolsrr6rr;pathlibrtriton.runtime.buildrtriton.runtime.cachertriton.backends.compilerrtriton.backends.driverrrrrealpath__file__rr?r.r@ lru_cacher+r1rUobjectrWrrrrr/r0rrs '2., ''//"''**84 5ww||GY/0  We, H 0,,0AA, ,iX %6 % IIr0