6^(hM:ddlZddlZddlZddlZddlZddlZddlmZddlm Z ddl m Z ddl m Z ddlmZddlmZej$j'ej$j)eZej$j-edgZej$j-ed Zd gZej4d Zej4d Zej4d ZdZGddeZ dZ!dZ"GddeZ#GddeZ$y)N)Path)_build)get_cache_manager) _allocation) 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_dirsr2-s  +LN ++r1c^ddlm}m}m}dj ||g|S)Nrmachinesystem architecture,)platformr5r6r7rr4s r platform_keyr:2s'66 88WY:<>: ;;r1c tj|tzjdj }t |}t jdjdd}|j|d|}|tj5}tjj|d}t|d5}|j!|dddt#|||t%t&t(} t| d5}|j+|j-|d|d }ddddddd dl} | j0j3||} | j0j5| } | j6j9| | S#1swYxYw#1swYuxYw#1swYyxYw) Nzutf-8 EXT_SUFFIX.rzmain.cwrbT)binaryr)hashlibsha256r:encode hexdigestr sysconfigget_config_varrget_filetempfileTemporaryDirectoryrrropenwriterr2 include_dir librariesputreadimportlib.utilutilspec_from_file_locationmodule_from_specloader exec_module) srcnamekeycacheext cache_pathtmpdirsrc_pathfso importlibspecmods rcompile_module_from_srcrc8s~ ..# .66w? @ J J LC c "E  " "< 0 6 6s ;B ?C4&#0J  ( ( * Ofww||FH5Hh$   h  YWBb$ O1"YYqvvxD63%YN  O  O >> 1 1$ CD .. ) )$ /CKKC J   O O  O Os<-F7 F4F7(F+7F7F( $F7+F4 0F77Gc$eZdZfdZdZxZS) CudaUtilscdt|dstt|||_|jS)Ninstance)hasattrsuperre__new__rg)cls __class__s rrjzCudaUtils.__new__Ss*sJ' C8=CL||r1cZtttjj t dj d}|j|_|j|_|j|_ |j|_ |j|_ |j|_ y)Nzdriver.c cuda_utils) rcrrrrr read_text load_binaryget_device_propertiescuOccupancyMaxActiveClustersset_printf_fifo_sizefill_1d_tma_descriptorfill_2d_tma_descriptor)selfrbs r__init__zCudaUtils.__init__Xs~%d277<<+L&M&W&W&Y[gh??%(%>%>",/,L,L)$'$<$<!&)&@&@#&)&@&@#r1)__name__ __module__ __qualname__rjrw __classcell__rls@rrereQs Ar1rec~|ddk(ryidddddd d dd d d ddddddddddddddddddddd|S)Nr* CUdeviceptri1int32_ti8int8_ti16int16_ti32i64int64_tu1uint32_tu8uint8_tu16uint16_tu32u64uint64_tfp16floatbf16fp32f32fp64double nvTmaDesc CUtensorMapr0)tys r ty_to_cpprgs !u|  i  h  y  y   y   j   i  z  z  z        w    ]! " #  r1c:fdfdfddj|jDcgc] }| c}}d|z}djt|j}tt t |j d}t|Dcic]\}}|| }}}t|dkDr)ddjd |jDznd}djd |jD}g} |jD]W\}}|dd k(r| jd |d $|dk(r| jd|>|dk7sD| jd|Ytt|} d} |jDcgc]\}}|dd k(rd|d|d|d|d } }}|jDcgc]\}}|dk(r d|d|d|d} }}|jDcgc]\}}|dk7s d|} }}| jddt|dkDrd|zndddj| d| j|jDcgc]\}}|d |d!c}}d"|d#|d$| j| d| j| d%t| dkDrddj| zndd&}|Scc}wcc}}wcc}}wcc}}wcc}}wcc}}w)'Nc^t|trdjt|S|S)Nr8) isinstancetuplermap)sig_serialize_signatures rrz+make_launcher.._serialize_signatures) c5 !88C 4c:; ; r1ct|tr!djt|}d|dS|ddk(ry|dvryt |S)Nr8[]rr~z PyObject* constexprrrrrrr)rval_extracted_types rrz&make_launcher.._extracted_typesT b% ((334Cse1:  a5C< + +}r1c t|tr!djt|}d|dS|ddk(ry|dvrydd d d d d dddddd t |S)N()rr~Orr^dlbhiLBHIK) rrlongrrrrrrrrr)rr format_ofs rrz make_launcher..format_ofs b% ''#i,-Cse1:  a5C< + +   B-  r1r iiiKKpOOOOOr8rz, c3,K|] \}}d|yw)z&_argNr0rrrs rrz make_launcher..s LB5 Lsc3NK|]\}}|dk7s t|d|yw)rz argN)rrs rrz make_launcher..s,h2VX\gVgYr]O4s3hs %%r~ptr_infoz.dev_ptrrz*tma_ptrr_argz zDevicePtrInfo ptr_infoz = getPointer(_argz); if (!ptr_infoz.valid) return NULL;zCUtensorMap* tma_ptrz = getTmaDesc(_argz); if (!tma_ptrz) return NULL;z&argz&global_scratchaB #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 launch_cooperative_grid, int clusterDimX, int clusterDimY, int clusterDimZ, int shared_memory, CUstream stream, CUfunction function, CUdeviceptr global_scratchz) { void *params[] = { aO }; if (gridX*gridY*gridZ > 0) { if ((num_ctas == 1) && (0 == launch_cooperative_grid)) { CUDA_CHECK(cuLaunchKernel(function, gridX, gridY, gridZ, 32*num_warps, 1, 1, shared_memory, stream, params, 0)); } else if ((num_ctas == 1) && (0 != launch_cooperative_grid)) { CUlaunchAttribute launchAttr[1]; CUlaunchAttribute coopAttr = { .id = CU_LAUNCH_ATTRIBUTE_COOPERATIVE, .value = 1}; launchAttr[0] = coopAttr; CUlaunchConfig config; config.gridDimX = gridX; config.gridDimY = gridY; config.gridDimZ = gridZ; config.blockDimX = 32 * num_warps; config.blockDimY = 1; config.blockDimZ = 1; config.sharedMemBytes = shared_memory; config.hStream = stream; config.attrs = launchAttr; config.numAttrs = 1; static cuLaunchKernelEx_t cuLaunchKernelExHandle = NULL; if (cuLaunchKernelExHandle == NULL) { cuLaunchKernelExHandle = getLaunchKernelExHandle(); } CUDA_CHECK(cuLaunchKernelExHandle(&config, function, params, 0)); } else { CUlaunchAttribute launchAttr[3]; 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; unsigned numAttrs = 2; if (0 != launch_cooperative_grid) { CUlaunchAttribute coopAttr = { .id = CU_LAUNCH_ATTRIBUTE_COOPERATIVE, .value = 1}; launchAttr[2] = coopAttr; numAttrs = 3; } 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 = numAttrs; 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; } else if (status != CUDA_SUCCESS) { CUDA_CHECK(status); // Catch any other cuda API errors 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 inline CUtensorMap* getTmaDesc(PyObject *obj) { if (sizeof(CUtensorMap*) != 8) { PyErr_SetString(PyExc_SystemError, "getTmaDesc() requires 64-bit compilation"); return NULL; } PyObject *method_handle = PyObject_GetAttrString(obj, "tma_desc_cpu_ptr"); if (!method_handle) { PyErr_SetString(PyExc_TypeError, "tma_desc_cpu_ptr() method does not exist"); return NULL; } PyObject *empty_tuple = PyTuple_New(0); if (!empty_tuple) { Py_DECREF(method_handle); PyErr_SetString(PyExc_SystemError, "Internal Python error!"); return NULL; } PyObject *method_ret = PyObject_Call(method_handle, empty_tuple, NULL); Py_DECREF(empty_tuple); Py_DECREF(method_handle); if (!method_ret) { PyErr_SetString(PyExc_SystemError, "Internal Python error!"); return NULL; } if (!PyLong_Check(method_ret)) { PyErr_SetString(PyExc_TypeError, "tma_desc_cpu_ptr() must return 64-bit int"); Py_DECREF(method_ret); return NULL; } uint64_t ptr_as_uint = PyLong_AsUnsignedLongLong(method_ret); Py_DECREF(method_ret); if (!ptr_as_uint) { PyErr_SetString(PyExc_ValueError, "received NULL ptr from tma_desc_cpu_ptr()"); return NULL; } if (ptr_as_uint % 64 != 0) { PyErr_SetString(PyExc_ValueError, "tma_desc_cpu_ptr() must be 64-byte aligned"); return NULL; } return (CUtensorMap*)(ptr_as_uint); } static void ensureCudaContext() { CUcontext pctx; CUDA_CHECK(cuCtxGetCurrent(&pctx)); if (!pctx) { // Ensure device context. CUdevice device; CUDA_CHECK(cuDeviceGet(&device, 0)); CUDA_CHECK(cuDevicePrimaryCtxRetain(&pctx, device)); CUDA_CHECK(cuCtxSetCurrent(pctx)); } } static PyObject* launch(PyObject* self, PyObject* args) { // ensure cuda context is valid before calling any CUDA APIs, e.g. before getPointer calls cuPointerGetAttributes ensureCudaContext(); int gridX, gridY, gridZ; uint64_t _stream; uint64_t _function; int launch_cooperative_grid; PyObject *launch_enter_hook = NULL; PyObject *launch_exit_hook = NULL; PyObject *kernel_metadata = NULL; PyObject *launch_metadata = NULL; PyObject *global_scratch_obj = NULL; z _arg;z if(!PyArg_ParseTuple(args, "a*", &gridX, &gridY, &gridZ, &_stream, &_function, &launch_cooperative_grid, &global_scratch_obj, &kernel_metadata, &launch_metadata, &launch_enter_hook, &launch_exit_hookat)) { 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; } CUdeviceptr global_scratch = 0; if (global_scratch_obj != Py_None) { DevicePtrInfo global_scratch_info = getPointer(global_scratch_obj, -1); if (!global_scratch_info.valid) { return NULL; } global_scratch = global_scratch_info.dev_ptr; } // raise exception asap z Py_BEGIN_ALLOW_THREADS; _launch(gridX, gridY, gridZ, num_warps, num_ctas, launch_cooperative_grid, clusterDimX, clusterDimY, clusterDimZ, shared_memory, (CUstream)_stream, (CUfunction)_function, global_scratcha0); 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; } Py_RETURN_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; } ) rvaluesrlistfilterboolr enumeratelenitemsappendrange) constants signaturer args_formatformatrs args_list arg_declsinternal_args_listparamsnewline ptr_decls tma_declsrVrrrs @@@r make_launcherr~s ,''93C3C3EFR9R=FGK [ (F193C3C3EFGIVD)//#"678I"+I"67$!QA7I7PST]P^abPbtyy L)//:K LLLhjI hARhhI"22 a5C<  % %8&< = ;   % %n 5 ;   % %QCj 123y> "FG__& Ar a5C< !#5aS1#=MaSPdeI foetetev\a\]_a   qc!3A3oaSWI'0oo&7 MUQ2;LQCj MF M MM#$.v\KNOXKY\]K]w{~GwGcevffyy()x*p << @QRuq"OB'(aS2RSTU &x(QR[P[\: <<  << |ehi{e|@e@}ADHDMDMN`Da}aFH|I'IWr Cf Ji G8$ NRSs)K:* K?.!L%L  L$LLceZdZdZdZy) CudaLauncherctdr jn t}fd}|jDcic]\}}|||}}}jjDcic]\}}|| }}}t ||t d}|j|_|j|_|j|_ |j|_ ycc}}wcc}}w)Nrctt|tr&jjj |fS|Sr.)rr!fn arg_namesindex)xrVs rz'CudaLauncher.__init__..s-Z3=OSVV--33A69UVr1__triton_launcher) rhrdictrrrrclaunchglobal_scratch_sizeglobal_scratch_alignlaunch_cooperative_grid) rvrVmetadatararg_idxidxvaluerrbs ` rrwzCudaLauncher.__init__s%,S+%>CMMDF V;D??;LMZS%WS\5(M M25--2E2E2GHJCS%ZH HIy1%c+>?jj #+#?#? $,$A$A!'/'G'G$NHs C/ C"c |jdkDr9||z|z}||jz}tj||j|} nd} |j||||||j | g|yNr)rr _allocatorrrr) rvgridXgridYgridZstreamfunctionargs grid_size alloc_sizeglobal_scratchs r__call__zCudaLauncher.__call__ su  # #a ' -I"T%=%==J(33J@Y@Y[abN!N E5%4;W;WYgojnor1N)rxryrzrwrr0r1rrrs Hpr1rcReZdZfdZdZdZdZedZdZ dZ dZ xZ S) CudaDrivercVt|_t|_t|yr.)reutilsr launcher_clsrirw)rvrls rrwzCudaDriver.__init__s[ ( r1c~|j}|j|}|ddz|dz}d}td||S)Nr  r )get_current_deviceget_device_capabilityr)rvdevice capability warp_sizes rget_current_targetzCudaDriver.get_current_targetsK((*//7 ]R'*Q-7  Y77r1cJddl}|jd|jS)Nrr )torchrrrvr s rget_active_torch_devicez"CudaDriver.get_active_torch_device"s||FD$;$;$=>>r1c"ddl}|jSr)r r r s rget_device_interfacezCudaDriver.get_device_interface&szzr1c ddl}|jjxr|jjduS#t $rYywxYw)NrF)r r is_availableversionhip ImportError)r s r is_activezCudaDriver.is_active*sC  ::**,L%--2C2Ct2K L  s7: AAcddlm}|S)Nr)do_bench)triton.testingr)rvrs rget_benchmarkerzCudaDriver.get_benchmarker2s +r1cbddl}d}|jt|dz|jdS)Nrir )dtyper)r emptyint)rvr  cache_sizes rget_empty_cache_for_benchmarkz(CudaDriver.get_empty_cache_for_benchmark6s. ' {{3zQ/uyy{PPr1c$|jyr.)zero_)rvrYs r clear_cachezCudaDriver.clear_cache?s  r1) rxryrzrwrr r staticmethodrrrr"r{r|s@rrrs; 8?Qr1r)% functoolsrrErArrHpathlibrtriton.runtime.buildrtriton.runtime.cachertriton.runtimertriton.backends.compilerrtriton.backends.driverrrr realpath__file__rrLr/rM lru_cacher,r2r:rcobjectrerrrrr0r1rr/s '2&., ''//"''**84 5ww||GY/0  We, H 0,,<< 2AA, .{| p6p0,,r1