o iY@sddlmZmZddlZddlZddlZddlZddlZddlZddl Z ddl m Z m Z ddl mZmZmZmZmZmZmZmZmZddlmZmZejejejeZdZd1dd Z d d Z!d d Z"ddZ#edZ$Gdddej%Z&e'ddZ(Gdddee$Z)Gddde)e$Z*ed2ddZ+edddddd3d%dZ+ d1dddddd&d4d*dZ+Gd+d,d,Z,Gd-d.d.Z-d/d0Z.dS)5) annotationsdivisionN) defaultdict namedtuple) CallableGenericIterableListOptionalTypeVarUnioncastoverload) get_backend path_to_ptxasz2.1.0cCsN|durt}z ddlm}||WSty&ddl}|j|jYSw)Nr)_cuda_getCurrentRawStream)get_current_deviceZtorch._Cr ImportErrortorchcudaZcurrent_streamZ cuda_stream)idxrrra/var/www/html/eduruby.in/lip-sync/lip-sync-env/lib/python3.10/site-packages/triton/runtime/jit.pyget_cuda_streams   rcCsddl}|jSNr)rrZcurrent_device)rrrrrs rcCsddl}|j|dSr)rrZ set_devicerrrrrset_current_device$srcCsddl}|j|Sr)rrget_device_capabilityrrrrr)s rTcs:eZdZdZd fdd ZddZdd Zd d ZZS) DependenciesFinderz This AST visitor is used to find dependencies of a JITFunction. This can be used to invalidate a JITFunction's hash when its source code -- or that of its dependencies -- changes. returnNonecs*tt|d|_||_dS)Nutf-8)super__init__hashlibmd5encode hexdigestretglobals)selfr+src __class__rrr%<s  zDependenciesFinder.__init__cCs|j|jdSN)r+getid)r,noderrr visit_NameAszDependenciesFinder.visit_NamecCsj||j}t|tjr||j}t|tjs |dus-t|dddks-t|dddr/dSt||jS)N__name__Ztritonz.triton)visitvalue isinstanceast Attributegetattrendswithattr)r,r3lhsrrrvisit_AttributeDs    * z"DependenciesFinder.visit_AttributecCs||j}|dur dSt|rdS|jr#|jds!d|jvr#dSt|ts1Jd|jd|j durLt |j }t |j|j }|||j|_ tt|dd}|j|j |d|_t|j|_dS)Nztriton.z.triton.z Function "zv" is being called from a Triton function but is not a Triton function itself. Decorate it with @triton.jit to fix thisnoinlineFr#)r7funcinspect isbuiltin __module__ startswithr9 JITFunctionr5hashr:parser-r __globals__r*strr<r(r&r'r))r,r3rBtreefinderrArrr visit_CallLs     zDependenciesFinder.visit_Call)r!r") r5rE __qualname____doc__r%r4r@rN __classcell__rrr.rr 5s r c Csddl}g}ttd}|t|g7}Wdn1s"wYtj t d}| |gD])}t|j |jjd}|t|g7}Wdn1sXwYq4ttj t dd}|t|g7}Wdn1swYtj t d}| |gD])}t|j |jjd}|t|g7}Wdn1swYqtd}tt|dg}d td|dd |S)Nrrbcompilerz_C/libtriton.solanguagez --version-)pkgutilopen__file__r&r'readr)ospathjoin TRITON_PATH iter_modules module_finder find_specnameoriginr subprocess check_outputTRITON_VERSION)rVcontentsfZ compiler_pathlibZ language_pathZptxasZ ptxas_versionrrr version_keycs.   ric@s eZdZUded<dddZdS)KernelInterfacerrunr!cCstttjtt|j|dS)z A JIT function is launched with: fn[grid](*args, **kwargs). Hence JITFunction.__getitem__ returns a callable proxy that memorizes the grid. )grid)r r functoolspartialrrk)r,rlrrr __getitem__szKernelInterface.__getitem__N)r!r)r5rErO__annotations__rorrrrrj}s rjcseZdZdZdZeddZeddZeddZed d Z d d Z ed dZ ddZ ddZ ddZd1ddZd1ddZd2dd Zd!d"Zd3d#d$Zed%d&Zd'd(Zd)d*Zd+d,Zfd-d.Zd/d0ZZS)4rGNcCst|dr|jSt|trdSt|tr*d|kr|dkrdSd|kr(|dkr(dSd St|tr1d S|dur7dStd t|d |) Ndtypei1iii32llu64i64fp32zUnsupported type z for )hasattrrrr9boolintfloat TypeErrortypeargrrr_key_ofs    zJITFunction._key_ofcCs"t|drt|jdr|jjSdS)Ndevicer}r6)rxrr}r~rrr _device_of  zJITFunction._device_ofcCs"t|drt|jtr|SdS)N is_pinnedF)rxr9rrr~rrr_pinned_memory_ofrzJITFunction._pinned_memory_ofcCsDt|dr|tjdkSt|tr|ddk|dkfS|dufS)Ndata_ptrrrqrxrrG divisibilityr9rzr~rrr_spec_ofs   zJITFunction._spec_ofcsRddfddt|D}fddt|D}tdddgt|t|S) NcSsDt|dr|tjdkSt|tr|tjdkS|dur dSdS)NrrTFr)xrrris_divisible_by_16s  z3JITFunction._get_config..is_divisible_by_16cs&h|]\}}|r|jvr|qSr)do_not_specialize.0irrr,rr &z*JITFunction._get_config..cs:h|]\}}t|tst|tr|dkr|jvr|qS)r)r9ryrzrrr,rrrs:Zinstance_descriptordivisible_by_16 equal_to_1) enumeratertuple)r,argsrrrrr _get_configszJITFunction._get_configcCs|durdSt|dd}idddddd d d d d ddddddddddddddddddd d!d"d#}t|D]}|||<qGt|trU|Sd$||S)%Nz*i8.ryrsZfloat8e4Zfp8e4Zfloat8e5Zfp8e5Z float8e4b15Zfp8e4b15Zfloat16Zfp16Zbfloat16Zbf16Zfloat32rwZfloat64Zfp64Zint8i8Zint16Zi16Zint32rtZint64rvZuint8u8Zuint16u16Zuint32u32Zuint64ru*)rKsplitlistvaluesr9)keyZ dtype_strZtysvrrr_type_ofsN      zJITFunction._type_ofcs dfddt|D}|S)N,csg|] \}}|qSr)r)rrkrrr z/JITFunction._make_signature..)r\r)r,Zsig_key signaturerrr_make_signatureszJITFunction._make_signaturecCstt|j|}|Sr0)dictzip constexprs)r,Z constexpr_key constantsrrr_make_constantsszJITFunction._make_constantsc  CstjdurdS|jj} |jj} dddt|j|dD} | d|d|d| d } t|}Gd d d } t |||||||d }tj|| | | | d |i|dddS)NF, cSsg|] \}}|d|qS)z: rrratyrrrrsz*JITFunction._call_hook..rz [num_warps=z , num_stages=]()c@seZdZddZdS)z.JITFunction._call_hook..LegacyCompilercSs||_||_dSr0)modulera)r,rrarrrr%sz7JITFunction._call_hook..LegacyCompiler.__init__N)r5rErOr%rrrrLegacyCompilers r)rrr num_warps num_stages extern_libsconfigsr)rreprfncompileZis_manual_warmupZalready_compiled) rG cache_hookrr5rEr\r arg_namesrKr)r,rrrrrrrrrarZ arg_reprsrrkwargsrrr _call_hooks  $zJITFunction._call_hookr!rKcCs|j|d}|dkr%d|dtjd|d|dtjd|d|d Sd |vr3d|dtjd S|d krDd|dtjd|d SdS)Nr6(z.data_ptr() % z == 0) if hasattr(z,, "data_ptr") else (z % z == 0, z == 1) if isinstance(z,, int) else (False,)Tensorz == 0)rzz == 1)z(False,))rpr1rGrr,rZarg_annotationrrr_get_arg_specialization_keys z'JITFunction._get_arg_specialization_keycCsD|j|d}d|vr|dS|dkrdS|dkrdSd|d S) Nr6rz.dtyperyrsr{rwz_key_of(r)rpr1rrrr_get_arg_sig_keys  zJITFunction._get_arg_sig_key device_types List[str]pinned_memory_flags List[bool]cCsvdd|D}d|vrddl}|jjrdSdStdd|D}tdd|D}|r/|r/dSt|dkr9|dSdS) NcSsg|]}|dkr|qS)r6rrZ device_typerrrrz5JITFunction._conclude_device_type..rrhipcss|]}|dkVqdS)cpuNrrrrr "sz4JITFunction._conclude_device_type..css|]}|VqdSr0r)rZpinned_memory_flagrrrr#s)rversionrallanylen)r,rrrZis_cpuZis_pinned_memoryrrr_conclude_device_typesz!JITFunction._conclude_device_typecsfddtjD}fddtjD}d|}dfdd|D}dddd|Dd}ddd d|Dd}d|}g}t|D]\} } | jvrZqP|| g7}qPd|} d d djD} dd d tjjD} djjd| d|dt |dkr|d nddt | dkr| d ndd| d|d|d|d|ddddjDd}t t j j jjjttttd }t|||jjS)Ncs g|] \}}|jvr|qSrrrrrrr+ z.JITFunction._make_launcher..cs g|] \}}|jvr|qSrrrrrrr,rrcsg|]}|qSr)rrrrrrr/s[cSg|]}d|dqS)z _device_of(rrrrrrr0r]cSr)z_pinned_memory_of(rrrrrrr1rrcSsg|] }d|d|qS)"z": rrrrrr<scss0|]\}}|tjkr |n|d|VqdS)z = NrC_empty)rraZdfltrrrr=s.z-JITFunction._make_launcher..z def rz, grid=None, num_warps=4, num_stages=3, extern_libs=None, stream=None, warmup=False, device=None, device_type=None): from ..compiler import compile, CompiledKernel sig_key = z, constexpr_key = rrz spec_key = aV key = (version_key, sig_key, constexpr_key, spec_key, num_warps, num_stages, self.debug) if not extern_libs is None: key = (key, tuple(extern_libs.items())) assert num_warps > 0 and (num_warps & (num_warps - 1)) == 0, "num_warps must be a power of 2" assert grid is not None if callable(grid): grid = grid({z}) grid_size = len(grid) grid_0 = grid[0] grid_1 = grid[1] if grid_size > 1 else 1 grid_2 = grid[2] if grid_size > 2 else 1 if device_type is None: device_types = [_device_type for _device_type in zW if _device_type != ''] device_type = self._conclude_device_type(device_types, a) device_backend = None if device_type not in ['cuda', 'hip']: device_backend = get_backend(device_type) if device_backend is None: raise ValueError('Cannot find backend for ' + device_type) if device is None: if device_type in ['cuda', 'hip']: device = get_current_device() set_current_device(device) else: device = device_backend.get_current_device() device_backend.set_current_device(device) if stream is None and not warmup: if device_type in ['cuda', 'hip']: stream = get_cuda_stream(device) else: stream = device_backend.get_stream() bin = cache[device].get(key, None) if bin is not None: if not warmup: bin.c_wrapper(grid_0, grid_1, grid_2, bin.num_warps, bin.shared, stream, bin.cu_function, CompiledKernel.launch_enter_hook, CompiledKernel.launch_exit_hook, bin, zt) return bin # kernel not cached -- compile else: # build dict of constant values args = [z] all_args = cSsg|]}|qSrrrrrrrqa, configs = self._get_config(*all_args), constants = self._make_constants(constexpr_key) constants.update({i: None for i, arg in enumerate(all_args) if arg is None}) constants.update({i: 1 for i in configs[0].equal_to_1}) # build kernel signature -- doesn't include specialized arguments signature = { i: self._type_of(_key_of(arg)) for i, arg in enumerate(all_args) if i not in self.constexprs } # build stub signature -- includes arguments that are specialized for i, arg in constants.items(): if callable(arg): raise TypeError(f"Callable constexpr at index {i} is not supported") if not self._call_hook(key, signature, device, constants, num_warps, num_stages, extern_libs, configs): bin = compile(self, signature=signature, device=device, constants=constants, num_warps=num_warps, num_stages=num_stages, extern_libs=extern_libs, configs=configs, debug=self.debug, device_type=device_type) if not warmup: bin.c_wrapper(grid_0, grid_1, grid_2, bin.num_warps, bin.shared, stream, bin.cu_function, CompiledKernel.launch_enter_hook, CompiledKernel.launch_exit_hook, bin, *args) self.cache[device][key] = bin return bin return None ) rirr,rrrrcache__spec__rrr)rrr\rrr arg_defaultsrr5rrirrrrrrrrrrexec)r,Z regular_argsZconstexpr_argsrZsig_keysrrZconstexpr_keysZspecializationsrrZ spec_keysZ grid_argsZargs_signaturer-scoperrr_make_launcher*sj     ,12E zJITFunction._make_launchercsZ|_|j_|_t|}dd|jD_dd|jD_ t ddj D_ |dur6gn|_ fddj D_ t t|_jjdd_tt_d_g_d_tjd d d krrd n|_|_d dfdd|jD_fddjD_ _!|j"_"|j#_#|j$_$|j_dS)NcSg|]}|jqSr)rarrrrrrrz(JITFunction.__init__..cSrr)defaultrrrrrrcss|]}|tjkVqdSr0rrrrrrsz'JITFunction.__init__..cs&h|]}t|trj|n|qSr)r9rKrindexrrrrrrz'JITFunction.__init__..defZ TRITON_DEBUG01TcSst|tr|jS|Sr0)r9r}r5)rrrrsz&JITFunction.__init__..csi|] \}}||qSrrr) normalize_tyrr rz(JITFunction.__init__..cs$g|]\}}d|vrj|qS)Z constexpr)rrrrrrrs$)%rrErrrCr parametersrrrrZ has_defaultsrtextwrapdedent getsourcer-findrrrrHkernel_decoratorskernelrZenvironr1debugrArpitemsrrrkrPr5rJ)r,rrrrrArr)rr,rr%s2    zJITFunction.__init__cCs<|jdurt|j|jd}|||jt|_|jS)N)r+r-)rHr rJr-r7rIr*ri)r,Zdependencies_finderrrr cache_keys zJITFunction.cache_keycOs |jttj|i|ddiS)NwarmupT)rkmap MockTensor wrap_dtyper,rrrrrrs zJITFunction.warmupcCsHt|j}t|tjsJt|jdksJt|jdtjs"J|S)Nrr)r:rIr-r9Modulerbody FunctionDef)r,rLrrrrIs zJITFunction.parsecOstd)Nz:Cannot call @triton.jit'd outside of the scope of a kernel) RuntimeErrorrrrr__call__szJITFunction.__call__cs6|dkrd|_tt||||dkrd|_dSdS)Nrr-)rr$rG __setattr__rH)r,rar8r.rrrs  zJITFunction.__setattr__cCsd|jd|jjdS)Nz JITFunction(:r)rrr5rrrr__repr__szJITFunction.__repr__r!rK)rrrrr!rK)NNNN)r5rErOrr staticmethodrrrrrrrrrrrrrr%propertyrrrIrrr rQrrr.rrGs:        i%   rGrr!JITFunction[T]cCdSr0r)rrrrjitrrrrrArOptional[Iterable[int]]rOptional[bool]rACallable[[T], JITFunction[T]]cCrr0rrrrrrs)rrrrA interpret Optional[T]r4Union[JITFunction[T], Callable[[T], JITFunction[T]]]cs*dfdd }|dur||S|S) a< Decorator for JIT-compiling a function using the Triton compiler. :note: When a jit'd function is called, arguments are implicitly converted to pointers if they have a :code:`.data_ptr()` method and a `.dtype` attribute. :note: This function will be compiled and run on the GPU. It will only have access to: * python primitives, * builtins within the triton package, * arguments to this function, * other jit'd functions :param fn: the function to be jit-compiled :type fn: Callable rrr!r cs6t|sJrddlm}||St|dS)Nr) GridSelectorr)callableZinterpreter.interpreterrrG)rrrrrrArrr decorators  zjit..decoratorNrrr!r r)rrrrrArrrrrrs c@s0eZdZdZeddZddZeddZdS) rzr Can be used in place of real tensors when calling: kernel.warmup(MockTensor(torch.float32), ...) cCs"|jjdkr|jdkrt|S|S)Nrrr)r/r5rErr~rrrr-s  zMockTensor.wrap_dtypecCs ||_dSr0)rr)r,rrrrrr%4 zMockTensor.__init__cCsdSrrrrrrr7rzMockTensor.data_ptrN)r5rErOrPr rr%rrrrrr(s rc@s.eZdZddZddZddZd d d Zd S) TensorWrappercCs*||_||_|j|_|j|_|jj|_dSr0)rrbaseZis_cudarshape)r,rrrrrrr%=s zTensorWrapper.__init__cCs |jSr0)rrrrrrrDrzTensorWrapper.data_ptrcCs |j|Sr0)rstride)r,rrrrr!Gs zTensorWrapper.strider!rKcCsd|jd|jdS)NzTensorWrapper[rr)rrrrrrr__str__JszTensorWrapper.__str__Nr )r5rErOr%rr!r"rrrrr<s rcCsPt|tr||jjkr|jSt|j|St|drt||Stdt|d)NrzCannot reinterpret a r)r9rrrrrxr|r})Ztensorrrrrr reinterpretNs     r#r0r)rrrrrArr!r) rrrrrrrArrrr!r)/ __future__rrr:rmr&rCrZrcr collectionsrrtypingrrrr r r r r rZcommon.backendrrr[dirnameabspathrXr]rerrrrr NodeVisitorr lru_cacherirjrGrrrr#rrrrsX,  .  ^  3