Ϧi?SSKrSSKrSSKrSSKrSSKrSSKJr SSKJr SSK J r SSK J r SSK Jr \RR!\RR#\55r\RR'\S5/r\RR'\S5rS /r\R."5S 5r\R."5S 5rS r"S S\5rSrSr"SS\5r"SS\5r g)N)Path)_build)get_cache_manager) 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_dirsr2+s  +LN ++r1c 4[R"URS55R5n[ U5nUR US35nUc[ R"5n[RRUS5n[US5nURU5 SSS5 [XU[5[[ 5n[US5nUR#UR%5US3SS9nSSS5 SSS5 SSKn U R(R+X5n U R(R-U 5n U R.R1U 5 U $!,(df  N=f!,(df  N~=f!,(df  N=f) Nzutf-8z.sozmain.cwrbT)binaryr)hashlibsha256encode hexdigestrget_filetempfileTemporaryDirectoryrrropenwriterr2 include_dir librariesputreadimportlib.utilutilspec_from_file_locationmodule_from_specloader exec_module) srcnamekeycache cache_pathtmpdirsrc_pathfso importlibspecmods rcompile_module_from_srcrV0s+ ..G, - 7 7 9C c "E4& -J  ( ( *fww||FH5Hh$ %  YWBb$1"YYqvvxD6dYK  + >> 1 1$ CD .. ) )$ /CKKC J%$  + *s<+-F E'*3F #E8F ' E5 1F 8 F F  Fc.^\rSrSrU4SjrSrSrU=r$) CudaUtilsHcn>[US5(d[[U]U5UlUR$)Ninstance)hasattrsuperrX__new__r[)cls __class__s rr^CudaUtils.__new__Js-sJ'' C8=CL||r1cZ[[[RR [ S55R 5S5nURUlURUlURUl URUl URUl URUl g)Nzdriver.c cuda_utils) rVrrrrr read_text load_binaryget_device_propertiescuOccupancyMaxActiveClustersset_printf_fifo_sizefill_1d_tma_descriptorfill_2d_tma_descriptor)selfrUs r__init__CudaUtils.__init__Os~%d277<<+L&M&W&W&Y[gh??%(%>%>",/,L,L)$'$<$<!&)&@&@#&)&@&@#r1)rgrirjrfrerh)__name__ __module__ __qualname____firstlineno__r^rl__static_attributes__ __classcell__r`s@rrXrXHs AAr1rXcUSS: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 CUtensorMapr0tys r ty_to_cppr^s !u|  i  h  y  y   y   j   i  z  z  z        w    ]! " #  r1cSRSUR555nSnSnSRUR5Vs/sHoe"U"U55PM sn5nSU-n[U5S:a)SSRSUR555-OSn /n UR5HUupUSS :XaU R S U S 35 M%US :XaU R S U 35 MAU R SU 35 MW UR 5V s/sH oU;dM U PM n n S[U5S:aSU-OSSSRSU 55SSRUR5V Vs/sHupU"U5SU S3PM snn 5SUSU SSRUR5V Vs/sHupUSS :XaSU SU SU SU S3 OSPM! snn 5SSRUR5V Vs/sHupUS :Xa SU SU S U S!3OSPM snn 5S"[U 5S:aSSRU 5-OSS#3n U $s snfs sn fs snn fs snn fs snn f)$Nz, c3F# UHup[U5SU3v M g7f)z argNrrirs rr make_launcher..xs#SARYr]O4s3ARs!c:USS:XagUS:Xag[U5$)Nrrv PyObject*rrrs r_extracted_type&make_launcher.._extracted_typezs% a5C<  }r1c &SSSSSSSSSS S S S . U$) NOrQdlbhrBHIK) rrrlongr{r}ryrrrrrr0rs r format_of make_launcher..format_ofs:     r1 iiiKKOOOOrc30# UH upSU3v M g7f)z&_argNr0rs rrrs L:K5:Ksrvptr_infoz.dev_ptrrz*tma_ptr_arga  #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,# UH nSU3v M g7f)z&argNr0)rrs rrrs 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; } 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; 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 zDevicePtrInfo ptr_infoz = getPointer(_argz); if (!ptr_infoz.valid) return NULL;z; zCUtensorMap* tma_ptrz = getTmaDesc(_argz); if (!tma_ptrz) return NULL;z; Py_BEGIN_ALLOW_THREADS; _launch(gridX, gridY, gridZ, num_warps, num_ctas, clusterDimX, clusterDimY, clusterDimZ, shared_memory, (CUstream)_stream, (CUfunction)_functionaW); 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; } )ritemsvalueslenappendkeys) constants signatureids arg_declsrrr args_formatformat args_listinternal_args_listrparamsrJs r make_launcherrus SARSSI ''IDTDTDVWDVb9_R%89DVWXK ; &FPST]P^abPbtyy L)//:K LLLhjI" a5C<  % %8&< = ;   % %n 5  % %QCj 1##) @)Ai-?a)F @.}\RUV_R`cdRd~BEN~Njl}mmyy n88Y__=N O=NEA#$E!B /=N OPQR &x(RR[Q\](77PYP_P_Pa bPaGLGHmopqmrvymy$QC'9!BqcAQRSQTTh i@B BPa bcdd77GPGVGVGX YGX~C}~_aep_p"1#%7s/!N [vx xGX YZ[[S|@R|SVW|WTX[_[d[dew[xTx]_S`)`AI CT JuXAN P. b Ys$ I/$ I41I4I9&I?  J$c \rSrSrSrSrSrg) CudaLauncheric^S[TS5(aTRRO [50n[TS5(a TRO [ 5nU4SjnUR 5VVs0sHupgU"U5U_M nnnTRR 5VVs0sHupgU"U5U_M nnn[XHU5m[TS5n U RUl gs snnfs snnf)Nids_of_const_exprsfnrc|>[U[5(a%TRRR U5$U$r.) isinstancer!r arg_namesindex)rrJs r'CudaLauncher.__init__..s/As9K9KCFF,,2215RQRRr1__triton_launcher) r\r constexprstuplerdictrrrrVlaunch) rkrJmetadatarrcst_keyrLvaluerrUs ` rrlCudaLauncher.__init__s#'#t:L:LSVV%6%6RWRYZ%,S+%>%>CMMDF R;D??;LM;LZSWS\5(;L M;>==;N;N;PQ;PZSWS\5(;P QI#6%c+>?jj NQs 5C,*C2c(UR"U0UD6 gr.r)rkargskwargss r__call__CudaLauncher.__call__s T$V$r1rN)rnrorprqrlrrrr0r1rrrs !%r1rcP^\rSrSrU4SjrSrSr\S5rSr Sr Sr U=r $) CudaDrivericV>[5Ul[Ul[TU]5 gr.)rXutilsr launcher_clsr]rl)rkr`s rrlCudaDriver.__init__s[ ( r1c|UR5nURU5nUSS-US-nSn[SX#5$)Nr  r )get_current_deviceget_device_capabilityr)rkdevice capability warp_sizes rget_current_targetCudaDriver.get_current_targetsI((*//7 ]R'*Q-7  77r1c"SSKnUR$Nr)torchr )rkrs rget_device_interfaceCudaDriver.get_device_interfaceszzr1c|SSKnURR5=(a URRSL$r)rr is_availableversionhip)rs r is_activeCudaDriver.is_actives,zz&&(Hemm.?.?4.GHr1cSSKJn U$)Nr)do_bench)triton.testingr)rkrs rget_benchmarkerCudaDriver.get_benchmarkers +r1c\SSKnSnUR[US-5URSS9$)Nrir )dtyper)remptyint)rkr cache_sizes rget_empty_cache_for_benchmark(CudaDriver.get_empty_cache_for_benchmarks. ' {{3zQ/uyy{PPr1)rr) rnrorprqrlrr staticmethodrrrrrrsrts@rrrs9 8IIQQr1r)! functoolsrr7rr<pathlibrtriton.runtime.buildrtriton.runtime.cachertriton.backends.compilerrtriton.backends.driverrrr realpath__file__rr@r/rA lru_cacher,r2rVobjectrXrrrrr0r1rrs '2., ''//"''**84 5ww||GY/0  We, H  0 ,,0AA, .wt %6 % "Q"Qr1