U yʹjN.@shddlZddlZddlZddlmZddlmZddlmZddZ dd Z d d Z d d Z ddZ dS)N)_build)get_cache_manager) version_keycCsddl}|jjdk S)Nr)torchversionZhip)rr\/var/www/html/bwave/python/venv/lib/python3.8/site-packages/triton/compiler/make_launcher.pyis_hip sr cCsFdd|D}|dd||}t|d}|S)NcSs&i|]\}}||ddkrdn|qS)r*Zptrr).0kvrrr sz%make_so_cache_key..-zutf-8)itemsjoinvalueshashlibmd5encode hexdigest)Z version_hash signature constantskeyrrr make_so_cache_keysrc Cstt||}t|}|d}||}|dkrt}t||}tj |d} t | d} | |W5QRXt || |} t | d.} |j | |ddW5QRW5QRSQRXW5QRXn|SdS)Nz.sozmain.cwrbT)binary)rrrZget_filetempfileTemporaryDirectorygenerate_launcherospathropenwriterputread) namerrZ so_cache_keyZso_cache_managerZso_name cache_pathZtmpdirsrcZsrc_pathfsorrr make_stubs       >r.c Cs<|ddkrtrdSdSdddddd d d d d d d d |S)Nrr ZhipDeviceptr_tZ CUdeviceptrint32_tZint8_tZint16_tint64_tuint32_tuint64_tfloatdouble) i1i8Zi16i32i64u32u64fp16bf16fp32f32fp64)r tyrrr ty_to_cpp0s"  rBcsddd|D}dddddd fd d |D}trd |d dfdd|Dddfdd |Dd|dddd|Ddddd |Ddddd|Dd}nd|ddfdd|Dddfdd |Dd |ddd!d|Dd"dd#d |Dd$dd%d|Dd&}|S)'N, css$|]\}}t|d|VqdS)z argN)rBr irArrr Dsz$generate_launcher..c Ss.|ddkrdSdddddddddd d |S) Nrr PyObject*r/r0r1r2r3r4) r5r7r8r9r:r;r<r=r>r?rr@rrr _extracted_typeFs  z*generate_launcher.._extracted_typec Ssddddddddd |S) NOr,dlIrEKL)rGr3r4longr1r/r2r0rr@rrr format_ofVs z$generate_launcher..format_ofZ iiiiiKKOOOrcsg|]}|qSrr)r rA)rHrPrr bsz%generate_launcher..a #define __HIP_PLATFORM_AMD__ #include #include #include static inline void gpuAssert(hipError_t code, const char *file, int line) { if (code != HIP_SUCCESS) { const char* prefix = "Triton Error [HIP]: "; const char* str = hipGetErrorString(code); char err[1024] = {0}; snprintf(err, 1024, "%s Code: %d, Messsage: %s", prefix, code, str ); PyErr_SetString(PyExc_RuntimeError, err); } } #define HIP_CHECK(ans) { gpuAssert((ans), __FILE__, __LINE__); } static void _launch(int gridX, int gridY, int gridZ, int num_warps, int shared_memory, hipStream_t stream, hipFunction_t function, z) { void *params[] = { c3s |]}|krd|VqdSz&argNrr rErrr rF{sa }; if (gridX*gridY*gridZ > 0) { HIP_CHECK(hipModuleLaunchKernel(function, gridX, gridY, gridZ, 64*num_warps, 1, 1, shared_memory, stream, params, 0)); } } typedef struct _DevicePtrInfo { hipDeviceptr_t 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 = (hipDeviceptr_t)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 = (hipDeviceptr_t)PyLong_AsUnsignedLongLong(ret); if (!ptr_info.dev_ptr) return ptr_info; uint64_t dev_ptr; hipError_t status = hipPointerGetAttribute(&dev_ptr, HIP_POINTER_ATTRIBUTE_DEVICE_POINTER, ptr_info.dev_ptr); if (status == hipErrorInvalidValue) { PyErr_Format(PyExc_ValueError, "Pointer argument (at %d) cannot be accessed from Triton (cpu tensor?)", idx); ptr_info.valid = false; } ptr_info.dev_ptr = (hipDeviceptr_t)dev_ptr; return ptr_info; } PyErr_SetString(PyExc_TypeError, "Pointer argument must be either uint64 or have data_ptr method"); return ptr_info; } static PyObject* launch(PyObject* self, PyObject* args) { int gridX, gridY, gridZ; uint64_t _stream; uint64_t _function; int num_warps; int shared_memory; PyObject *launch_enter_hook = NULL; PyObject *launch_exit_hook = NULL; PyObject *compiled_kernel = NULL;  cs$g|]\}}|d|dqSz _arg; rrDrHrr rQsz$ if (!PyArg_ParseTuple(args, "z", &gridX, &gridY, &gridZ, &num_warps, &shared_memory, &_stream, &_function, &launch_enter_hook, &launch_exit_hook, &compiled_kernel, css|]\}}d|VqdSz&_argNrrDrrr rFsz)) { return NULL; } if (launch_enter_hook != Py_None) { PyObject_CallObject(launch_enter_hook, args); } // raise exception asap rWc Ss>g|]6\}}|ddkr6d|d|d|d|d ndqSrr zDevicePtrInfo ptr_infoz = getPointer(_argrCz); if (!ptr_infoz.valid) return NULL;rrrDrrr rQszo; _launch(gridX, gridY, gridZ, num_warps, shared_memory, (hipStream_t)_stream, (hipFunction_t)_function, css4|],\}}|ddkr"d|dnd|VqdSrr Zptr_infoz.dev_ptrZ_argNrrDrrr rFsa5); if (launch_exit_hook != Py_None) { PyObject_CallObject(launch_exit_hook, args); } if (PyErr_Occurred()) { 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; } a_ #include "cuda.h" #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); PyErr_SetString(PyExc_RuntimeError, err); } } #define CUDA_CHECK(ans) { gpuAssert((ans), __FILE__, __LINE__); } static void _launch(int gridX, int gridY, int gridZ, int num_warps, int shared_memory, CUstream stream, CUfunction function, z) { void *params[] = { c3s |]}|krd|VqdSrRrrSrTrr rFsa }; if(gridX*gridY*gridZ > 0){ CUDA_CHECK(cuLaunchKernel(function, gridX, gridY, gridZ, 32*num_warps, 1, 1, shared_memory, stream, 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"); return ptr_info; } static PyObject* launch(PyObject* self, PyObject* args) { int gridX, gridY, gridZ; uint64_t _stream; uint64_t _function; int num_warps; int shared_memory; PyObject *launch_enter_hook = NULL; PyObject *launch_exit_hook = NULL; PyObject *compiled_kernel = NULL; cs$g|]\}}|d|dqSrVrrDrXrr rQEsz if(!PyArg_ParseTuple(args, "css|]\}}d|VqdSrYrrDrrr rFFsz)) { return NULL; } if (launch_enter_hook != Py_None) { PyObject_CallObject(launch_enter_hook, args); } // raise exception asap c Ss>g|]6\}}|ddkr6d|d|d|d|d ndqSrZrrDrrr rQPsze; _launch(gridX, gridY, gridZ, num_warps, shared_memory, (CUstream)_stream, (CUfunction)_function, css4|],\}}|ddkr"d|dnd|VqdSr[rrDrrr rFQsa); if (launch_exit_hook != Py_None) { PyObject_CallObject(launch_exit_hook, args); } if(PyErr_Occurred()) { 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; } )rrrr keys)rrZ arg_declsformatr+r)rHrrPr r"CsL "]^^gh  STT^_r")rr#r commonrZ runtime.cacherZ runtime.jitrr rr.rBr"rrrr s