U yʹjY @sddlmZmZddlZddlZddlZddlZddlZddlZddl Z ddl m Z m Z ddl mZmZmZmZmZmZmZmZmZddlmZmZejejejeZdZd.dd Z d d Z!d d Z"ddZ#edZ$Gdddej%Z&e'ddZ(Gdddee$Z)Gddde)e$Z*edddddZ+eddddddddd d!d"dZ+d/dddddd#d$ddddd%d&d'dZ+Gd(d)d)Z,Gd*d+d+Z-d,d-Z.dS)0) annotationsdivisionN) defaultdict namedtuple) CallableGenericIterableListOptionalTypeVarUnioncastoverload) get_backend path_to_ptxasz2.1.0cCsT|dkrt}zddlm}||WStk rNddl}|j|jYSXdS)Nr)_cuda_getCurrentRawStream)get_current_deviceZtorch._Cr ImportErrortorchcudaZcurrent_streamZ cuda_stream)idxrrrQ/var/www/html/bwave/python/venv/lib/python3.8/site-packages/triton/runtime/jit.pyget_cuda_streams  rcCsddl}|jSNr)rrZcurrent_device)rrrrrsrcCsddl}|j|dSr)rrZ set_devicerrrrrset_current_device$srcCsddl}|j|Sr)rrget_device_capabilityrrrrr)srTcs>eZdZdZddfdd 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. Nonereturncs*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_NamecCs`||j}t|tjr&||j}q |dksPt|dddksPt|dddrTdSt||jS)N__name__Ztritonz.triton)visitvalue isinstanceast Attributegetattrendswithattr)r-r4lhsrrrvisit_AttributeDs   *z"DependenciesFinder.visit_AttributecCs||j}|dkrdSt|r&dS|jrF|jdsBd|jkrFdSt|tsbtd|j d|j dkrt |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$)r8funcinspect isbuiltin __module__ startswithr: JITFunctionAssertionErrorr6hashr;parser.r __globals__r+strr=r)r'r(r*)r-r4rCtreefinderrBrrr visit_CallLs      zDependenciesFinder.visit_Call) r6rF __qualname____doc__r&r5rArP __classcell__rrr/rr 5s r c Cshddl}g}ttd}|t|g7}W5QRXtj t d}| |gD]>}t|j |jjd}|t|g7}W5QRXqTttj t dd}|t|g7}W5QRXtj t d}| |gD]>}t|j |jjd}|t|g7}W5QRXqtd}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)rXcontentsfZ compiler_pathlibZ language_pathZptxasZ ptxas_versionrrr version_keycs" "$"$ rkc@s$eZdZUded<ddddZdS)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 functoolspartialrrm)r-rnrrr __getitem__szKernelInterface.__getitem__N)r6rFrQ__annotations__rqrrrrrl}s rlcseZdZdZdZeddZeddZeddZed d Z d d Z ed dZ ddZ ddZ ddZddddZddddZddddddZd d!Zd0d"d#Zed$d%Zd&d'Zd(d)Zd*d+Zfd,d-Zd.d/ZZS)1rHNcCst|dr|jSt|trdSt|trVd|kr<|dkr.is_divisible_by_16cs&h|]\}}|r|jkr|qSr)do_not_specialize.0irrr-rr s z*JITFunction._get_config..cs:h|]2\}}t|tst|tr|dkr|jkr|qS)r)r:r{r|rrr-rrrs   Zinstance_descriptordivisible_by_16 equal_to_1) enumeratertuple)r-argsrrrrr _get_configszJITFunction._get_configcCsz|dkr dSt|dd}dddddd d d d d ddddddd}t|D] }|||<qPt|trl|Sd||S)Nz*i8.ruZfp8e4Zfp8e5Zfp8e4b15Zfp16Zbf16ryZfp64i8Zi16rvrxu8u16u32rw)r{Zfloat8e4Zfloat8e5Z float8e4b15Zfloat16Zbfloat16Zfloat32Zfloat64Zint8Zint16Zint32Zint64Zuint8Zuint16Zuint32Zuint64*)rMsplitlistvaluesr:)keyZ dtype_strZtysvrrr_type_ofs. zJITFunction._type_ofcs dfddt|D}|S)N,csg|]\}}|qSr)r)rrkrrr sz/JITFunction._make_signature..)r^r)r-Zsig_key signaturerrr_make_signatureszJITFunction._make_signaturecCstt|j|}|Sr1)dictzip constexprs)r-Z constexpr_key constantsrrr_make_constantsszJITFunction._make_constantsc  CstjdkrdS|jj} |jj} dddt|j|dD} | d|d|d| d } t|}Gd d d } t |||||||d }tj|| | | | d |i|dddS)NF, cSsg|]\}}|d|qS)z: rrrctyrrrrsz*JITFunction._call_hook..rz [num_warps=z , num_stages=]()c@seZdZddZdS)z.JITFunction._call_hook..LegacyCompilercSs||_||_dSr1)modulerc)r-rrcrrrr&sz7JITFunction._call_hook..LegacyCompiler.__init__N)r6rFrQr&rrrrLegacyCompilersr)rrr num_warps num_stages extern_libsconfigsr)rreprfncompileZis_manual_warmupZalready_compiled) rH cache_hookrr6rFr^r arg_namesrMr)r-rrrrrrrrrcrZ arg_reprsrrkwargsrrr _call_hooks  zJITFunction._call_hookrMr"cCs|j|d}|dkrJd|dtjd|d|dtjd|d|d Sd |krfd|dtjd S|d krd|dtjd|d SdSdS)Nr7(z.data_ptr() % z == 0) if hasattr(z,, "data_ptr") else (z % z == 0, z == 1) if isinstance(z,, int) else (False,)Tensorz == 0)r|z == 1)z(False,))rrr2rHrr-rZarg_annotationrrr_get_arg_specialization_keys4z'JITFunction._get_arg_specialization_keycCsH|j|d}d|kr |dS|dkr,dS|dkr8dSd|d SdS) Nr7rz.dtyper{rur}ryz_key_of(r)rrr2rrrr_get_arg_sig_keys zJITFunction._get_arg_sig_keyz List[str]z List[bool]) device_typespinned_memory_flagsr#cCsvdd|D}d|kr.ddl}|jjr*dSdStdd|D}tdd|D}|r^|r^dSt|dkrr|dSdS) NcSsg|]}|dkr|qS)r7rrZ device_typerrrrsz5JITFunction._conclude_device_type..rrhipcss|]}|dkVqdS)cpuNrrrrr "sz4JITFunction._conclude_device_type..css|] }|VqdSr1r)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]$\} } | jkrq|| g7}qd|} d d djD} dd d tjjD} djjd| d|dt |dkr4|d nddt | dkrR| d ndd| d|d|d|d|ddddjDd}t t j j jjjttttd }t|||jjS)Ncs g|]\}}|jkr|qSrrrrrrr+s z.JITFunction._make_launcher..cs g|]\}}|jkr|qSrrrrrrr,s rcsg|]}|qSr)rrrrrrr/s[cSsg|]}d|dqS)z _device_of(rrrrrrr0s]cSsg|]}d|dqS)z_pinned_memory_of(rrrrrrr1srcSsg|]}d|d|qS)"z": rrrrrr<scss.|]&\}}|tjkr|n |d|VqdS)z = NrD_empty)rrcZdfltrrrr=sz-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|] }|qSrrrrrrrqsa, 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 ) rkrr-rrrrcache__spec__rrr)rrr^rrr arg_defaultsrr6rrkrrrrrrrrrrexec)r-Z regular_argsZconstexpr_argsrZsig_keysrrZconstexpr_keysZspecializationsrrZ spec_keysZ grid_argsZargs_signaturer.scoperrr_make_launcher*sj     ,,2E zJITFunction._make_launchercsZ|_|j_|_t|}dd|jD_dd|jD_ t ddj D_ |dkrlgn|_ fddj D_ t t|_jjdd_tt_d_g_d_tjd d d krd n|_|_d dfdd|jD_fddjD_ _!|j"_"|j#_#|j$_$|j_dS)NcSsg|] }|jqSr)rcrrrrrrsz(JITFunction.__init__..cSsg|] }|jqSr)defaultrrrrrscss|]}|tjkVqdSr1rrrrrrsz'JITFunction.__init__..cs&h|]}t|trj|n|qSr)r:rMrindexrrrrrsz'JITFunction.__init__..defZ TRITON_DEBUG01TcSst|tr|jS|Sr1)r:rr6)rrrrz&JITFunction.__init__..csi|]\}}||qSrrr) normalize_tyrr sz(JITFunction.__init__..cs$g|]\}}d|krj|qS)Z constexpr)rrrrrrrs)%rrFrrrDr parametersrrrrZ has_defaultsrtextwrapdedent getsourcer.findrrrrJkernel_decoratorskernelr\environr2debugrBrritemsrrrmrRr6rL)r-rrrrrBrr)rr-rr&s2   zJITFunction.__init__cCs<|jdkr6t|j|jd}|||jt|_|jS)N)r,r.)rJr rLr.r8rKr+rk)r-Zdependencies_finderrrr cache_keys  zJITFunction.cache_keycOs|jttj||ddiS)NwarmupT)rmmap MockTensor wrap_dtyper-rrrrrrszJITFunction.warmupcCsHt|j}t|tjstt|jdks.tt|jdtjsDt|S)Nrr) r;rKr.r:ModulerIrbody FunctionDef)r-rNrrrrKs  zJITFunction.parsecOs tddS)Nz:Cannot call @triton.jit'd outside of the scope of a kernel) RuntimeErrorrrrr__call__szJITFunction.__call__cs2|dkrd|_tt||||dkr.d|_dS)Nrr.)rr%rH __setattr__rJ)r-rcr9r/rrrs zJITFunction.__setattr__cCsd|jd|jjdS)Nz JITFunction(:r)rrr6rrrr__repr__szJITFunction.__repr__)NNNN)r6rFrQrr staticmethodrrrrrrrrrrrrrr&propertyrrrKrrrrSrrr/rrHs8       i %  rHJITFunction[T]rr#cCsdSr1r)rrrrjitsrrrrrBzOptional[Iterable[int]]zOptional[bool]zCallable[[T], JITFunction[T]])rrrBr#cCsdSr1rrrrrrs)rrrrB interpretz Optional[T]z4Union[JITFunction[T], Callable[[T], JITFunction[T]]])rrrrBr r#cs4dddfdd }|dk r,||S|SdS)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 rrrcs:t|s tr$ddlm}||St|dSdS)Nr) GridSelectorr)callablerIZinterpreter.interpreterr rH)rr rrr rBrrr decorators  zjit..decoratorNr)rrrrrBr r rr rrs c@s0eZdZdZeddZddZeddZdS) rzr Can be used in place of real tensors when calling: kernel.warmup(MockTensor(torch.float32), ...) cCs"|jjdkr|jdkrt|S|S)Nrtr)r0r6rFrrrrrr-s  zMockTensor.wrap_dtypecCs ||_dSr1)rt)r-rtrrrr&4szMockTensor.__init__cCsdSrrrrrrr7szMockTensor.data_ptrN)r6rFrQrRrrr&rrrrrr(s  rc@s2eZdZddZddZddZddd d Zd S) TensorWrappercCs*||_||_|j|_|j|_|jj|_dSr1)rtbaseZis_cudarshape)r-rrtrrrr&=s zTensorWrapper.__init__cCs |jSr1)rrrrrrrDszTensorWrapper.data_ptrcCs |j|Sr1)rstride)r-rrrrrGszTensorWrapper.striderMr"cCsd|jd|jdS)NzTensorWrapper[rr)rtrrrrr__str__JszTensorWrapper.__str__N)r6rFrQr&rrrrrrrr<srcCsVt|tr*||jjkr|jSt|j|Sn(t|dr>t||Stdt|ddS)NrzCannot reinterpret a r)r:rrrtrzr~r)Ztensorrtrrr reinterpretNs    r)N)N)/ __future__rrr;ror'rDr\rer collectionsrrtypingrrrr r r r r rZcommon.backendrrr]dirnameabspathrZr_rgrrrrr NodeVisitorr lru_cacherkrlrHrrrrrrrrsV, .  ^ 3