&ThM:SSKrSSKrSSKrSSKrSSKrSSKrSSKJr SSKJ r SSK J r SSK J r SSKJr SSKJr \R$R'\R$R)\55r\R$R-\S5/r\R$R-\S 5rS /r\R4"5S 5r\R4"5S 5r\R4"5S 5rSr"SS\5r Sr!Sr""SS\5r#"SS\5r$g)N)Path)_build)get_cache_manager) _allocation) GPUTarget) GPUDriverincludelibcudac "[R"S5nU(aU/$[R"SS/5R 5nUR 5Vs/sHnSU;dM UR 5SPM nnUVs/sH"n[RRU5PM$ nn[R"S5nU(alU(deUR S5Vs/sHIn[RR[RRUS55(dMGUPMK nnSnU(aUS [U5-- nUS - nO US - nUS - n[S U55(dU5eU$s snfs snfs snf)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.zlibcuda_dirs..)s/Sddrww~~bggll4@AAdsA A ) rgetenv subprocess check_outputdecode splitlinessplitrdirnamerrstrany) env_libcuda_pathlibslinelocslocdirsenv_ld_library_pathdirmsgs r libcuda_dirsr-s\yy!67 !!  " "$4d#; < C C ED*.): U):nPT>T DJJL ):D U,0 1DSBGGOOC DD 1))$564288=s=PRPWPWP\P\]`bpPqAr=s &C  2SY>> KK MM >> SdS S SXUXX S K V 1ts F'F)F&AF 0F c$[/[5Q$N) libdevice_dirr-r library_dirsr3-s  +LN ++r2c^SSKJnJnJn SR U"5U"5/U"5Q5$)Nrmachinesystem architecture,)platformr6r7r8rr5s r platform_keyr;2s'66 88WY:<>: ;;r2c [R"U[5-RS55R 5n[ U5n[ R"S5RS5SnURUSU35nUc[R"5n[RRUS5n[US5nUR!U5 SSS5 [#XU[%5[&[(5n [U S5nUR+UR-5USU3SS 9nSSS5 SSS5 S SKn U R0R3X5n U R0R5U 5n U R6R9U 5 U $!,(df  N=f!,(df  N~=f!,(df  N=f) Nzutf-8 EXT_SUFFIX.rzmain.cwrbT)binaryr)hashlibsha256r;encode hexdigestr sysconfigget_config_varr get_filetempfileTemporaryDirectoryrrropenwriterr3 include_dir librariesputreadimportlib.utilutilspec_from_file_locationmodule_from_specloader exec_module) srcnamekeycacheext cache_pathtmpdirsrc_pathfso importlibspecmods rcompile_module_from_srcrd8s\ ..# .66w? @ J J LC c "E  " "< 0 6 6s ;B ?C4&#0J  ( ( *fww||FH5Hh$ %  YWBb$1"YYqvvxD63%YN  + >> 1 1$ CD .. ) )$ /CKKC J%$  + *s< -G F3G%F/7G F, (G/ F= 9G Gc.^\rSrSrU4SjrSrSrU=r$) CudaUtilsQcn>[US5(d[[U]U5UlUR$)Ninstance)hasattrsuperrf__new__ri)cls __class__s rrlCudaUtils.__new__Ss-sJ'' C8=CL||r2cZ[[[RR [ S55R 5S5nURUlURUlURUl URUl URUl URUl g)Nzdriver.c cuda_utils) rdrrrrr! read_text load_binaryget_device_propertiescuOccupancyMaxActiveClustersset_printf_fifo_sizefill_1d_tma_descriptorfill_2d_tma_descriptor)selfrcs r__init__CudaUtils.__init__Xs~%d277<<+L&M&W&W&Y[gh??%(%>%>",/,L,L)$'$<$<!&)&@&@#&)&@&@#r2)rurwrxrtrsrv)__name__ __module__ __qualname____firstlineno__rlrz__static_attributes__ __classcell__rns@rrfrfQs AAr2rfcUSS:Xag0SS_SS_SS _S S_S S _S S_SS_SS_SS_SS_SS_SS_SS_SS_SS_SS_U$)Nr* CUdeviceptri1int32_ti8int8_ti16int16_ti32i64int64_tu1uint32_tu8uint8_tu16uint16_tu32u64uint64_tfp16floatbf16fp32f32fp64double nvTmaDesc CUtensorMapr1)tys r ty_to_cpprgs !u|  i  h  y  y   y   j   i  z  z  z        w    ]! " #  r2c|^^^U4SjmU4SjmU4SjmSRUR5Vs/sH nT"U5PM sn5nSU-nSR[TUR555n[[ [ UR S555n[U5VVs0sHupVXV_M nnn[U5S:a)SSRS UR555-OSnSRS UR555n/n UR5H]upRUSS :XaU RS US 35 M%US:XaU RSU35 MAUS:wdMIU RSU35 M_ [[U55n Sn UR5VVs/sHupRUSS :XdMSUSUSUSUS3 PM! n nnUR5VVs/sHupRUS:XdM SUSUSUS3PM n nnUR5VVs/sHupRUS:wdM SU3PM n nnU RS5 S[U5S:aSU-OSSSRU 5SU RUR5VVs/sHupRT"U5S US!3PM snn5S"US#US$U RU 5SU RU 5S%[U 5S:aSSRU 5-OSS&3nU$s snfs snnfs snnfs snnfs snnfs snnf)'Nch>[U[5(aSR[TU55$U$)Nr9) isinstancetuplermap)sig_serialize_signatures rr+make_launcher.._serialize_signatures, c5 ! !88C 4c:; ; r2c>[U[5(a!SR[TU55nSUS3$USS:XagUS;ag[ U5$)Nr9[]rrz PyObject* constexprrrrrrr)rval_extracted_types rr&make_launcher.._extracted_typesW b% ((334Cse1:  a5C< + +}r2c >[U[5(a!SR[TU55nSUS3$USS:XagUS;agSS S S S S SSSSSS. [ U5$)N()rrOrr_dlbhiLBHIK) rrlongrrrrrrrrr)rr format_ofs rr make_launcher..format_ofs b% ''#i,-Cse1:  a5C< + +   B-  r2r iiiKKpOOOOOr9rz, c30# UH upSU3v M g7f)z&_argNr1rrrs rr make_launcher..s L:K5:Ksc3V# UHupUS:wdM [U5SU3v M! g7f)rz argN)rrs rrrs/hARVX\gVg3Yr]O4s3ARs ))rptr_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_declsrWrrrs @@@r make_launcherr~s ,''93C3C3EF3ER9R=3EFGK [ (F193C3C3EFGIVD)//#"678I"+I"67"6$!"6I7PST]P^abPbtyy L)//:K LLLhjI hARhhI" a5C<  % %8&< = ;   % %n 5 ;   % %QCj 1#3y> "FG__&&EA a5C< f #5aS1#=MaSPde& foetetevev\a\]   X qc!3A3oaSWev'0oo&7 M&7UQ2;LjQCj&7F M MM#$.v\KNOXKY\]K]w{~GwGcevffyy()x*p << @QR@QuqOB'(aS2@QRSTU &x(QR[P[\: <<  << |ehi{e|@e@}ADHDMDMN`Da}aFH|I'IWr Cf Ji G8$ NRSs5L0L =L&L&9 L, L,/ L2? L2$L8c \rSrSrSrSrSrg) CudaLauncheric^[TS5(a TRO [5nU4SjnUR5VVs0sHupVU"U5U_M nnnTRR5VVs0sHupVXV_M nnn[ X75m[ TS5nURUlURUlURUl URUl gs snnfs snnf)Nrc~>[U[5(a&TRRR U54$U$r/)rr"fn arg_namesindex)xrWs r'CudaLauncher.__init__..s2Z3=O=OSVV--33A69VUVVr2__triton_launcher) rjrdictrrrrdlaunchglobal_scratch_sizeglobal_scratch_alignlaunch_cooperative_grid) ryrWmetadatararg_idxidxvaluerrcs ` rrzCudaLauncher.__init__s%,S+%>%>CMMDF V;D??;LM;LZSWS\5(;L M25--2E2E2GH2GJCSZ2G HI1%c+>?jj #+#?#? $,$A$A!'/'G'G$NHs C$7C*c URS:a7X-U-nXpR-n[R"XRU5n OSn UR"XX4XPR U /UQ76 gNr)rr _allocatorrrr) rygridXgridYgridZstreamfunctionargs grid_size alloc_sizeglobal_scratchs r__call__CudaLauncher.__call__ sg  # #a ' -I"%=%==J(33J@Y@Y[abN!N E%;W;WYgojnor2)rrrrN)r|r}r~rrzrrr1r2rrrs Hpr2rc\^\rSrSrU4SjrSrSrSr\S5r Sr Sr S r S r U=r$) CudaDrivericV>[5Ul[Ul[TU]5 gr/)rfutilsr launcher_clsrkrz)ryrns rrzCudaDriver.__init__s[ ( r2c|UR5nURU5nUSS-US-nSn[SX#5$)Nr  r )get_current_deviceget_device_capabilityr)rydevice capability warp_sizes rget_current_targetCudaDriver.get_current_targetsI((*//7 ]R'*Q-7  77r2cJSSKnURSUR55$)Nrr )torchrrryrs rget_active_torch_device"CudaDriver.get_active_torch_device"s||FD$;$;$=>>r2c"SSKnUR$r)rr rs rget_device_interfaceCudaDriver.get_device_interface&szzr2cSSKnURR5=(a URRSL$![ a gf=f)NrF)rr is_availableversionhip ImportError)rs r is_activeCudaDriver.is_active*sC  ::**,L%--2C2Ct2K L  srDs  '2&., ''//"''**84 5ww||GY/0  We, H  0 ,, << 2AA, .{| p6p0,,r2