o hY@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_devicetorch._Cr ImportErrortorchcudacurrent_stream cuda_stream)idxrrrH/var/www/html/ai/venv/lib/python3.10/site-packages/triton/runtime/jit.pyget_cuda_streams   rcCsddl}|jSNr)rrcurrent_device)rrrrrs rcCsddl}|j|dSr)rr set_devicerrrrrset_current_device$sr"cCsddl}|j|Sr)rrget_device_capabilityr!rrrr#)s r#Tcs: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)selfr0src __class__rrr*<s  zDependenciesFinder.__init__cCs|j|jdSN)r0getid)r1noderrr visit_NameAszDependenciesFinder.visit_NamecCsj||j}t|tjr||j}t|tjs |dus-t|dddks-t|dddr/dSt||jS)N__name__tritonz.triton)visitvalue isinstanceast Attributegetattrendswithattr)r1r8lhsrrrvisit_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()r=funcinspect isbuiltin __module__ startswithr? JITFunctionr:hashr@parser2r% __globals__r/strrBr-r+r,r.)r1r8rHtreefinderrGrrr visit_CallLs     zDependenciesFinder.visit_Call)r&r') r:rK __qualname____doc__r*r9rFrT __classcell__rrr3rr%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)r\contentsf compiler_pathlib language_pathptxas ptxas_versionrrr version_keycs.   rsc@s eZdZUded<dddZdS)KernelInterfacer$runr&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$ functoolspartialrru)r1rvrrr __getitem__szKernelInterface.__getitem__N)r&r$)r:rKrU__annotations__ryrrrrrt}s rtcseZdZdZdZeddZeddZeddZed 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)4rMNcCst|dr|jSt|trdSt|tr*d|kr|dkrdSd|kr(|dkr(dSd St|tr1d S|dur7dStd t|d |) Ndtypei1iii32llu64i64fp32zUnsupported type z for )hasattrr|r?boolintfloat TypeErrortypeargrrr_key_ofs    zJITFunction._key_ofcCs"t|drt|jdr|jjSdS)Ndevicerr;)rrrrrrr _device_of  zJITFunction._device_ofcCs"t|drt|jtr|SdS)N is_pinnedF)rr?rrrrrr_pinned_memory_ofrzJITFunction._pinned_memory_ofcCsDt|dr|tjdkSt|tr|ddk|dkfS|dufS)Ndata_ptrrr{rrrM divisibilityr?rrrrr_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.0irrr1rr &z*JITFunction._get_config..cs:h|]\}}t|tst|tr|dkr|jvr|qS)r)r?rrrrr1rrrs:instance_descriptordivisible_by_16 equal_to_1) enumeratertuple)r1argsrrrrr _get_configszJITFunction._get_configcCs|durdSt|dd}idddddd d d d d ddddddddddddddddddd d!d"d#}t|D]}|||<qGt|trU|Sd$||S)%Nz*i8.rr}float8e4fp8e4float8e5fp8e5 float8e4b15fp8e4b15float16fp16bfloat16bf16float32rfloat64fp64int8i8int16i16int32r~int64ruint8u8uint16u16uint32u32uint64r*)rQsplitlistvaluesr?)key dtype_strtysvrrr_type_ofsN      zJITFunction._type_ofcs dfddt|D}|S)N,csg|] \}}|qSr)r)rrkrrr z/JITFunction._make_signature..)rbr)r1sig_key signaturerrr_make_signatureszJITFunction._make_signaturecCstt|j|}|Sr5)dictzip constexprs)r1 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: rrrgtyrrrrsz*JITFunction._call_hook..rz [num_warps=z , num_stages=]()c@seZdZddZdS)z.JITFunction._call_hook..LegacyCompilercSs||_||_dSr5)modulerg)r1rrgrrrr*sz7JITFunction._call_hook..LegacyCompiler.__init__N)r:rKrUr*rrrrLegacyCompilers r)rrr num_warps num_stages extern_libsconfigsr)rreprfncompileis_manual_warmupalready_compiled) rM cache_hookrr:rKrbr arg_namesrQr)r1rrrrrrrrrgr arg_reprsrrkwargsrrr _call_hooks  $zJITFunction._call_hookr&rQcCs|j|d}|dkr%d|dtjd|d|dtjd|d|d Sd |vr3d|dtjd S|d krDd|dtjd|d SdS)Nr;(z.data_ptr() % z == 0) if hasattr(z,, "data_ptr") else (z % z == 0, z == 1) if isinstance(z,, int) else (False,)Tensorz == 0)rz == 1)z(False,))rzr6rMrr1rarg_annotationrrr_get_arg_specialization_keys z'JITFunction._get_arg_specialization_keycCsD|j|d}d|vr|dS|dkrdS|dkrdSd|d S) Nr;rz.dtyperr}rrz_key_of(r)rzr6rrrr_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)r;rr device_typerrrrz5JITFunction._conclude_device_type..rrhipcss|]}|dkVqdS)cpuNrrrrr "sz4JITFunction._conclude_device_type..css|]}|VqdSr5r)rpinned_memory_flagrrrr#s)rversionrallanylen)r1rrris_cpuis_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 = NrI_empty)rrgdfltrrrr=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 ) rsrr1rrrrcache__spec__rrr")rrrbrrr arg_defaultsrr:rrsrrrrrrrrrr"exec)r1 regular_argsconstexpr_argsrsig_keysrrconstexpr_keysspecializationsrr spec_keys grid_argsargs_signaturer2scoperrr_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)rgrrrrrrrz(JITFunction.__init__..cSrr)defaultrrrrrrcss|]}|tjkVqdSr5r rrrrrsz'JITFunction.__init__..cs&h|]}t|trj|n|qSr)r?rQrindexrrrrrrz'JITFunction.__init__..def TRITON_DEBUG01TcSst|tr|jS|Sr5)r?rr:)rrrrsz&JITFunction.__init__..csi|] \}}||qSrrr) normalize_tyrr rz(JITFunction.__init__..cs$g|]\}}d|vrj|qS) constexpr)rr!rrrrrs$)%rrKrrrIr parametersrrrr has_defaultsrtextwrapdedent getsourcer2findrrrrNkernel_decoratorskernelr`environr6debugrGrzitemsrrrurVr:rP)r1rrrr3rGrr)r'r1rr*s2    zJITFunction.__init__cCs<|jdurt|j|jd}|||jt|_|jS)N)r0r2)rNr%rPr2r=rOr/rs)r1dependencies_finderrrr cache_keys zJITFunction.cache_keycOs |jttj|i|ddiS)NwarmupT)rumap MockTensor wrap_dtyper1rrrrrr7s zJITFunction.warmupcCsHt|j}t|tjsJt|jdksJt|jdtjs"J|S)Nrr)r@rOr2r?Modulerbody FunctionDef)r1rRrrrrOs zJITFunction.parsecOstd)Nz:Cannot call @triton.jit'd outside of the scope of a kernel) RuntimeErrorr;rrr__call__szJITFunction.__call__cs6|dkrd|_tt||||dkrd|_dSdS)Nr0r2)r1r)rM __setattr__rN)r1rgr>r3rrrAs  zJITFunction.__setattr__cCsd|jd|jjdS)Nz JITFunction(:r)rrr:rrrr__repr__szJITFunction.__repr__r&rQ)rrrrr&rQ)NNNN)r:rKrUrr staticmethodrrrrrrrrrrrrrr*propertyr6r7rOr@rArCrWrrr3rrMs:        i%   rMrr&JITFunction[T]cCdSr5r)rrrrjitrIrrr3rGrOptional[Iterable[int]]r3Optional[bool]rGCallable[[T], JITFunction[T]]cCrHr5rrKrrrrIs)rrr3rG interpret Optional[T]rO4Union[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 rr$r&rGcs6t|sJrddlm}||St|dS)Nr) GridSelectorrK)callableinterpreter.interpreterrRrM)rrRr3rrOrGrrr decorators  zjit..decoratorNrr$r&rGr)rrrr3rGrOrVrrUrrIs c@s0eZdZdZeddZddZeddZdS) r9zr Can be used in place of real tensors when calling: kernel.warmup(MockTensor(torch.float32), ...) cCs"|jjdkr|jdkrt|S|S)Nr|r)r4r:rKr9rrrrr:-s  zMockTensor.wrap_dtypecCs ||_dSr5)r|)r1r|rrrr*4 zMockTensor.__init__cCsdSrrrrrrr7rJzMockTensor.data_ptrN)r:rKrUrVrEr:r*rrrrrr9(s r9c@s.eZdZddZddZddZd d d Zd S) TensorWrappercCs*||_||_|j|_|j|_|jj|_dSr5)r|baseis_cudarshape)r1rZr|rrrr*=s zTensorWrapper.__init__cCs |jSr5)rZrrrrrrDrXzTensorWrapper.data_ptrcCs |j|Sr5)rZstride)r1rrrrr]Gs zTensorWrapper.strider&rQcCsd|jd|jdS)NzTensorWrapper[rr)r|rZrrrr__str__JszTensorWrapper.__str__NrD)r:rKrUr*rr]r^rrrrrY<s rYcCsPt|tr||jjkr|jSt|j|St|drt||Stdt|d)NrzCannot reinterpret a r)r?rYrZr|rrr)tensorr|rrr reinterpretNs     r`r5rW)rrLr3rMrGrMr&rN) rrPrrLr3rMrGrMrOrMr&rQ)/ __future__rrr@rwr+rIr`rir, collectionsrrtypingrrrr r r r r rcommon.backendrrradirnameabspathr^rcrkrrr"r#r$ NodeVisitorr% lru_cachersrtrMrIr9rYr`rrrrsX,  .  ^  3