h5cSSKrSSKrSSKrSSKrSSKrSSKrSSKJr SSKJr SSK J r SSK J r SSK Jr SSKJr \R"R%\R"R'\55r\R"R+\S5/r\R"R+\S 5rS /r\R2"5S 5r\R2"5S 5r"S S\5rSrSSSSSS.rSSSSSS.r Sr!Sr""SS5r#\$"S\%"S555r&S\&S'S\&S 'S \&S'S!r'S"r("S#S$\5r)"S%S&\5r*g)'N)Path)knobs)compile_module_from_src) _allocation) GPUTarget) GPUDriverincludelibcudac *[RR=n(aU/$[R"SS/5R 5nUR 5Vs/sHnSU;dM UR5SPM nnUVs/sH"n[RRU5PM$ nn[R"S5nU(alU(deURS5Vs/sHIn[RR[RRUS55(dMGUPMK nnSnU(aUS[U5-- nUS - nO US - nUS - n[S U55(dU5eU$s snfs snfs snf) Nz/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 )rnvidia libcuda_path subprocess check_outputdecode splitlinessplitrrdirnamegetenvrrstrany) env_libcuda_pathlibslinelocslocdirsenv_ld_library_pathdirmsgs r libcuda_dirsr.s_ <<4444 !!  " "$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 *AF4Fc$[/[5Q$N) libdevice_dirr.r library_dirsr4,s  +LN ++r3c.^\rSrSrU4SjrSrSrU=r$) CudaUtils6cn>[US5(d[[U]U5UlUR$)Ninstance)hasattrsuperr6__new__r9)cls __class__s rr<CudaUtils.__new__8s-sJ'' C8=CL||r3cZ[[[RR [ S55R 5S[5[[S9nURUl URUl URUl URUl URUlg)Nzdriver.c cuda_utilssrcnamer4 include_dirs libraries)rrrrrr! read_textr4rErF load_binaryget_device_propertiescuOccupancyMaxActiveClustersset_printf_fifo_sizefill_tma_descriptor)selfmods r__init__CudaUtils.__init__=s%RWW\\':67AAC%%  ??%(%>%>",/,L,L)$'$<$<!#&#:#: r3)rJrLrIrHrK)__name__ __module__ __qualname____firstlineno__r<rO__static_attributes__ __classcell__r>s@rr6r66s ; ;r3r6cUSS:XagURS5(ag0SS_SS _S S _S S_S S_SS_SS_SS_SS_SS_SS_SS_SS_SS_SS_SS_U$)Nr* CUdeviceptr tensordesc CUtensorMapi1int32_ti8int8_ti16int16_ti32i64int64_tu1uint32_tu8uint8_tu16uint16_tu32u64uint64_tfp16doublebf16fp32f32fp64 nvTmaDesc) startswith)tys r ty_to_cpprxQs !u| }}\""  i  h  y  y   y   j   i  z  z  z        x    ]! " #  r3rkrgrn)rorqrrrsrt pack_fp16 pack_bf16 pack_fp32 pack_fp64 iiiKKppOOOOOcZ^^^^U4SjnU4SjmU4SjmU4SjmU"UR55n[U5VVs0sHupVXV_M nnnSRUR5Vs/sH nT"U5PM sn5n[U-n /n UR5H n T"X5 M [U 5VVs0sHupVXV_M nnn[ U5S:a)SSRSUR 555-OSn /n UR 5HUupWUS :XaM U[ ;aU R[ US U35 M6U R[U5S U35 MW SRU 5n/nUR 5H~upWUSS :XaURS US 35 M%U[ ;aURSUS35 MFUS:XaURSU35 MbUS :wdMjURSU35 M [[ U55nSnUR 5VVs/sHupWUSS :XdMSUSUSUSUS3 PM! nnnUR 5VVs/sHupWUS:XdM SUSUSUS3PM nnnUR 5VVs/sH-upWU[ ;dM[ USUS[USUS3PM/ nnnUR 5VVs/sHupWUS :wdM SU3PM nnnURS 5 S![ U5S:aSU-OSS"SRU5S#URUR 5VVs/sHupWT"U5SUS$3PM snn5S%U S&U S'URU5SURU5SURU5S([ U5S:aSSRU5-OSS)3nU$s snnfs snfs snnfs snnfs snnfs snnfs snnfs snnf)*Nc>/nSnUGH9n[U[5(GaURS5(aT (aT UOSnUS- n[R"SU5nUR S5nUR S5nUR S5S-nUc;URSU-5 [SU-5Hn URS5 M OURS 5 [U5Hn URS 5 M [U5Hn URS5 M GM(URU5 GM< T (aU[T 5:XdeU$) Nrr[ztensordesc<([^[>]*)\[([^]]*)\],rYrdrurc) isinstancer#rvrematchgroupcountappendrangelen) signatureoutputtensordesc_idxsigmetardtypeshapendim_tensordesc_metas r_expand_signature(make_launcher.._expand_signature~s+C#s##|(D(D:I~6t!#!CSI A A{{3'!+<MM#+. #1t8_ e,-MM+.tAMM%(%tAMM%(% c"7:#nO8L&LLL r3cv>[U[5(aUH nT"X!5 M gURU5 gr0)rtupler)rrx_flatten_signatures rr)make_launcher.._flatten_signatures0 c5 ! !"1- MM# r3c>[U[5(a!SR[TU55nSUS3$USS:XagUS;ag[ U5$)Nr[]rrYz PyObject* constexprru)rrrmaprx)rwval_extracted_types rr&make_launcher.._extracted_typesW b% ((334Cse1:  a5C< + +}r3c >[U[5(a!SR[TU55nSUS3$USS:XagUS;agUR S5(agS S S S S SSSSSS. [ U5$)N()rrYOrr[dlbhiLBHIK) rplongr`rbr^rerirkrgrn)rrrrrvrx)rwr format_ofs rr make_launcher..format_ofs b% ''#i,-Cse1:  a5C< + + == & &   B-  r3rrz, c30# UH upSU3v M g7f)z&_argNr2)rrrws rr make_launcher..s L:K5:Ksrz argrYptr_infoz.dev_ptr_arg_storageruz*tma_ptrz zDevicePtrInfo ptr_infoz = getPointer(_argz); if (!ptr_infoz.valid) return NULL;zCUtensorMap* tma_ptrz = getTmaDesc(_argz); if (!tma_ptrz) return NULL;z _argz _storage = z(_argz);z&argz&global_scratchaR #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 launch_pdl, int clusterDimX, int clusterDimY, int clusterDimZ, int shared_memory, CUstream stream, CUfunction function, CUdeviceptr global_scratchz) { void *params[] = { a, }; if (gridX*gridY*gridZ > 0) { // 4 attributes that we can currently pass maxmimum CUlaunchAttribute launchAttr[4]; static cuLaunchKernelEx_t cuLaunchKernelExHandle = NULL; if (cuLaunchKernelExHandle == NULL) { cuLaunchKernelExHandle = getLaunchKernelExHandle(); } CUlaunchConfig config; config.gridDimX = gridX; config.gridDimY = gridY; config.gridDimZ = gridZ; if (num_ctas != 1) { config.gridDimX *= clusterDimX; config.gridDimY *= clusterDimY; config.gridDimZ *= clusterDimZ; } config.blockDimX = 32 * num_warps; config.blockDimY = 1; config.blockDimZ = 1; config.sharedMemBytes = shared_memory; config.hStream = stream; config.attrs = launchAttr; int num_attrs = 0; if (launch_pdl != 0) { CUlaunchAttribute pdlAttr = { .id = CU_LAUNCH_ATTRIBUTE_PROGRAMMATIC_STREAM_SERIALIZATION, .value = 1}; launchAttr[num_attrs] = pdlAttr; ++num_attrs; } if (launch_cooperative_grid != 0) { CUlaunchAttribute coopAttr = { .id = CU_LAUNCH_ATTRIBUTE_COOPERATIVE, .value = 1}; launchAttr[num_attrs] = coopAttr; ++num_attrs; } if (num_ctas != 1) { CUlaunchAttribute clusterAttr = {}; clusterAttr.id = CU_LAUNCH_ATTRIBUTE_CLUSTER_DIMENSION; clusterAttr.value.clusterDim.x = clusterDimX; clusterAttr.value.clusterDim.y = clusterDimY; clusterAttr.value.clusterDim.z = clusterDimZ; launchAttr[num_attrs] = clusterAttr; ++num_attrs; CUlaunchAttribute clusterSchedulingAttr = {}; clusterSchedulingAttr.id = CU_LAUNCH_ATTRIBUTE_CLUSTER_SCHEDULING_POLICY_PREFERENCE; clusterSchedulingAttr.value.clusterSchedulingPolicyPreference = CU_CLUSTER_SCHEDULING_POLICY_SPREAD; launchAttr[num_attrs] = clusterSchedulingAttr; ++num_attrs; } config.numAttrs = num_attrs; 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 uint16_t pack_fp16(double f) { uint16_t result; // from https://github.com/python/pythoncapi-compat #if 0x030600B1 <= PY_VERSION_HEX && PY_VERSION_HEX <= 0x030B00A1 && !defined(PYPY_VERSION) _PyFloat_Pack2(f, (unsigned char*)&result, 1); #else PyFloat_Pack2(f, (unsigned char*)&result, 1); #endif return result; } static uint16_t pack_bf16(double f) { float f32 = (float)f; uint32_t u32 = *(uint32_t*)&f32; return (uint16_t)(u32 >> 16); } static uint32_t pack_fp32(double f) { float f32 = (float)f; return *(uint32_t*)&f32; } static uint64_t pack_fp64(double f) { return *(uint64_t*)&f; } 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; int launch_pdl; PyObject *launch_enter_hook = NULL; PyObject *launch_exit_hook = NULL; PyObject *kernel_metadata = NULL; PyObject *launch_metadata = NULL; PyObject *global_scratch_obj = NULL; ;z if(!PyArg_ParseTuple(args, "a7", &gridX, &gridY, &gridZ, &_stream, &_function, &launch_cooperative_grid, &launch_pdl, &global_scratch_obj, &kernel_metadata, &launch_metadata, &launch_enter_hook, &launch_exit_hooka)) { 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; Py_DECREF(ret); } 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, launch_pdl, clusterDimX, clusterDimY, clusterDimZ, shared_memory, (CUstream)_stream, (CUfunction)_function, global_scratchaC); 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_DECREF(ret); } 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; } ) values enumerater_BASE_ARGS_FORMATritemsFLOAT_STORAGE_TYPErrxrFLOAT_PACK_FUNCTION) constantsrrrexpand_signaturersrw args_formatformatflat_signaturer args_list arg_decl_list arg_declsinternal_args_listparamsnewline ptr_decls tma_declsfloat_storage_declsrCrrrs ` @@@r make_launcherr|s $L.))9)9);<"+,<"=>"=$!"=I>''93C3C3EF3ER9R=3EFGK  ,FN!3/""+N";<";$!";I "FG__&&EA a5C< f #5aS1#=MaSPde& foetetevev\a\]   X qc!3A3oaSWev __&&EA # # Z b ! "%s+6I"6M5NeTUSVVXY& '0oo&7 M&7UQ2;LjQCj&7F M MM#$.F\[^_h[ilm[mGKNWGWsuFvvyy()Q*b << @QR@QuqOB'(aS2@QRSTU &x(QR[P[\< <<  <<  <<#$%&HqtuGqHKLqLIMPTPYPYZlPmImRTHU'UM M C\ JE ?F =8  NDSsHO>/PP PP PP6P P P! P!5P'c$\rSrSrSrSrSrSrg)TmaDescKernelParamiSc`SSKnURURURSS9Ulg)Nrcpurdevice)torchempty TMA_DESC_SIZEuint8descrMrs rrOTmaDescKernelParam.__init__Vs%KK 2 2%++eKT r3c6URR5$r0)rdata_ptr)rMs rtma_desc_cpu_ptr#TmaDescKernelParam.tma_desc_cpu_ptr[syy!!##r3)rN)rQrRrSrTrrOrrUr2r3rrrSsMU $r3rc#(# UHoU4v M g7fr0r2)rrs rrr`s: 1A s  c DUc=UR/URQURQURQURQ$USnUSnUSnUSnUSnURR5nURnURn U SS:Xde[ 5n U /UQU Qn U(a[ U5nUS==S-ss'[ RRRRRU R5UUU[UUUU 5 U $) Nswizzle elem_size elem_type block_size fp4_paddedrrr)baserstridesrrlisttritonruntimedriveractiveutilsrLrTMA_DTYPE_DEVICE_TO_HOST) argmetadatarrrrrrrrrresults rmake_tensordesc_argrfs'M399Ms{{MSYYMMMy!G%I%I,'J,'Jxx  "H IIEkkG 2;!    D %U %W %FU  b Q  NN  &&::  +   Mr3c8^^^^SSKJm SSKJm UUUU4SjnU$)Nr)TensorDescriptorcf>US[[5nU[[5SnSn/n[U5HXupV[UT T45(a/T (aT UOSnUS- nUR [ Xg55 MGUR U5 MZ T (aU[T 5:XdeT "/UQUQ76$)Nrr)rrrrextendrr) args meta_argsraw_kernel_argsr final_argsrrrGluonTensorDescriptorrlauncherrs rinner%wrap_handle_tensordesc..inners0#/01 s#4567 0FA# 02GHII:I~6t!#!!"5c"@A!!#& 1#nO8L&LLL00Z00r3)triton.tools.tensor_descriptorr'triton.experimental.gluon.nvidia.hopper)rrrrrs`` @@rwrap_handle_tensordescr s?a 1 1 Lr3c \rSrSrSrSrSrg) CudaLauncheric$^[TS5(a TRO [5nU4SjnUR5VVs0sHupVU"U5U_M nnnTRR5VVs0sHupVXV_M nnn[ USS5n[ X7U5m[TS[5[[S9n [SUR555n [R"[R UR"S5UlU (a['U R(U5O U R(UlUR*UlUR,UlUR.UlUR0Ulgs snnfs snnf)Nrc~>[U[5(a&TRRR U54$U$r0)rr#fn arg_namesindex)rrCs r'CudaLauncher.__init__..s2Z3=O=OSVV--33A69VUVVr3r__triton_launcherrBc3r# UH-n[U[5=(a URS5v M/ g7f)r[N)rr#rv)rrs rr(CudaLauncher.__init__..s+!vcu\_*S#"6"W3>>,;W"Wcus57r)r:rdictrrgetattrrrr4rErFr$r functoolsreduceoperatormul cluster_dimsnum_ctasr launchglobal_scratch_sizeglobal_scratch_alignlaunch_cooperative_grid launch_pdl) rMrCrrarg_idxidxvaluerrrNhas_tensor_desc_args ` rrOCudaLauncher.__init__sA%,S+%>%>CMMDF V;D??;LM;LZSWS\5(;L M25--2E2E2GH2GJCSZ2G H!(,=tDI/B%$%%  "!vclcscscu!vv!((x7L7LaP M`,SZZIfifpfp #+#?#? $,$A$A!'/'G'G$"--%NHs F7F c URS:aDX-U-nXpR-UR-n[R"XRU5n OSn UR "XX4XPR URU /UQ76 gNr)r!rr _allocatorr"r r#r$) rMgridXgridYgridZstreamfunctionr grid_size alloc_sizeglobal_scratchs r__call__CudaLauncher.__call__s}  # #a ' -I"]]2T5M5MMJ(33J@Y@Y[abN!N E%;W;WY]YhYh" +%) +r3)r"r!r r#r$rN)rQrRrSrTrOr5rUr2r3rr r s ..+r3r c\^\rSrSrU4SjrSrSrSr\S5r Sr Sr S r S r U=r$) CudaDrivericV>[5Ul[Ul[TU]5 gr0)r6rr  launcher_clsr;rO)rMr>s rrOCudaDriver.__init__s[ ( r3c|UR5nURU5nUSS-US-nSn[SX#5$)Nrrr r )get_current_deviceget_device_capabilityr)rMr capability warp_sizes rget_current_targetCudaDriver.get_current_targetsI((*//7 ]R'*Q-7  77r3cJSSKnURSUR55$)Nrr )rrr>rs rget_active_torch_device"CudaDriver.get_active_torch_devices||FD$;$;$=>>r3c"SSKnUR$r+)rr rs rget_device_interfaceCudaDriver.get_device_interfaceszzr3cSSKnURR5=(a URRSL$![ a gf=f)NrF)rr is_availableversionhip ImportError)rs r is_activeCudaDriver.is_activesC  ::**,L%--2C2Ct2K L  srksn  8&., ''//"''**84 5 Wi01  We, H  . ,,;;6 4            #Tn $ $ :b ::  %P,!+6!+H,,r3