i\ UddlZddlZddlZddlZddlZddlmZmZmZddl m Z m Z m Z ddl mZddlmZmZddlmZej&eZiZeeeefed<d ed eefd Zd ed e fd eefdZed dddd ed e ede eeefde ed ef dZej>Z dZ!ejDde#d edfdZ$d e#fdZ%ded e fdZ&edded e fdZ'y)N)Callable GeneratorIterable)AnyOptionalUnion) exposed_in) custom_op CustomOpDef) infer_schematriton_ops_to_kernelsnamereturnc.tj|gSN)rget)rs j/Volumes/fast/ai/experiments/voice-extract-mac/.venv/lib/python3.12/site-packages/torch/_library/triton.pyget_triton_kernels_for_oprs $ $T2 ..fn.c d d dtdtfdttdzdtdtt ffd |S) a{ Inspect the source of an arbitrary callable passed to torch._library.triton_op, and grab all of the triton kernels that are wrapped inside of it. This function traces local variable assignments to handle patterns like: kernel_fn = _my_kernel # global JITFunction wrapped = some_wrapper(kernel_fn) capture_triton(wrapped)[grid](...) It also recursively analyzes called functions to find triton kernels hidden behind helper function calls. That said, it is best effort. There are cases (e.g., recursion > MAX_RECURSION_DEPTH) that are not accounted for, so keep that in mind. Nr. visited_fnsdepthrc  ddlmddlmtj|} tt|}|vrgSkDrt jdgSj| tj|}ddlm}|}|j%|dt'j(|j+}Gd d t&j,}|} | j/|d t&j0d t2t4fd dt6d t6dzffd dt6d t8t4t:ffd} | |} ddt2t4dt8t4t:fdt8t4t2t&j0fdzdtt4dzd t2t6f  fd g} t} t3| j<}| j>D]}|jA||D]P}|g| | jB}|D]4}t|}|| vs| j|| jE|6R| jFD]}| jI|}|ddl%m&}||vr||jN}| tQ|s=tS|dsJ  |dz}|D]4}t|}|| vs| j|| jE|6| S#t$rt j dgcYSwxYw#ttf$rgcYSwxYw#tT$rt jd|dYwxYw)Nr Autotuner JITFunctionz.Triton not available, find_triton_kernels = []z7reached max recursion depth (%s) in find_triton_kernels)IndentedBufferT)stripceZdZddZdej ddfdZdejddfdZdejddfdZ y) Fget_inner_triton_kernels..find_triton_kernels..VisitorrNc<g|_i|_g|_g|_yr)triton_kernels assignmentscalled_functions return_exprs)selfs r__init__zOget_inner_triton_kernels..find_triton_kernels..Visitor.__init__Xs"13#>@ 35%46!rnodec|jD]\}t|tjs|jj |j gj|j^|j|yr) targets isinstanceastNamer' setdefaultidappendvalue generic_visit)r*r,targets r visit_AssignzSget_inner_triton_kernels..find_triton_kernels..Visitor.visit_AssignasW"llF!&#((3((33FIIrBII$**U+""4(rc|j%|jj|j|j|yr)r5r)r4r6r*r,s r visit_ReturnzSget_inner_triton_kernels..find_triton_kernels..Visitor.visit_Returngs2::)%%,,TZZ8""4(rc,d}t|jtjr |j}t|jtjrt|jjtj r|jjj dk(r|jjdk(rw|j|vri|jr&t|jdtj r|jj|jdj nt|jjtjrt|jjjtj rb|jjjj dk(r4|jjjdk(r|jj|jjd|jnt|jtj r|jj |vrf|jrt|jdtj rb|jj|jdj n/|jj|jj |j|y)N)capture_triton wrap_tritontorch_libraryropsz::) r/funcr0 Attributer5r1r3attrargsr&r4r(r6)r*r,triton_func_namesrDs r visit_CallzQget_inner_triton_kernels..find_triton_kernels..Visitor.visit_Callls$E!dii799D!$**cmm<&tzz'7'7B $ 0 0 3 3w > $ : = $ -> >#yyZ ! chh-O $ 3 3 : :499Q 1188#'::??"32dii[ A   3884yy||'8899DIIaL#(()K //66tyy|G--44TYY\\B""4(r)rN) __name__ __module__ __qualname__r+r0Assignr8Returnr;CallrGrrVisitorr$WsL 7 ) ) )  ) ) ) )sxx )D )rrOexprrcngGfddtj}|j|S)z3Extract all Name references from an AST expression.cZeZdZdejddffd Zdej ddfdZy)mget_inner_triton_kernels..find_triton_kernels..extract_names_from_expr..NameExtractorr,rNc<j|jyr)r4r3)r*r,namess r visit_Namezxget_inner_triton_kernels..find_triton_kernels..extract_names_from_expr..NameExtractor.visit_NamesLL)rc&|j|yr)r6r:s rrGzxget_inner_triton_kernels..find_triton_kernels..extract_names_from_expr..NameExtractor.visit_Calls&&t,r)rHrIrJr0r1rVrMrG)rUsr NameExtractorrSs1*388**-388--rrX)r0 NodeVisitorvisit)rPrXrUs @rextract_names_from_exprzVget_inner_triton_kernels..find_triton_kernels..extract_names_from_exprs/!E - - O ! !$ 'Lrobjct|fr|St|r&t|dr|j}t|r|Sy)zACheck if obj is a triton kernel or wrapper and return the kernel.rN)r/callablehasattrr)r\innerrr s rresolve_to_kernelzPget_inner_triton_kernels..find_triton_kernels..resolve_to_kernelsB# Y78 }d!3e[1 Lrfunc_objct|r tj|}t|r t |dsiStj |}i}|j |j|j |j|j |jt |dr|j |j|S#t$rYwxYw)zBBuild a combined namespace from a function's globals and closures.__code__ __globals__) r^inspectunwrap ValueErrorr_getclosurevarsupdatebuiltinsglobals nonlocalsre)rbfunc_closure_vars namespaces rbuild_namespacezNget_inner_triton_kernels..find_triton_kernels..build_namespaces!&~~h7HH%WXz-J ' 6 6x @ (*I   .77 8   .66 7   .88 9x/  !5!56 "sC CCrUror'visitedc*| t}g}|D]}||vr|j|||vr||}|}||j|>t|rC t j |}t|dr!| dz} | r|j| tjd||3||vr/||D]&} | } | |||} |j| (tjd||S#t $r|}YwxYw)z^ Resolve a list of names to triton kernels using the given namespace. rdr z'failed to resolve %s to a triton kernelz(%s not found in namespace or assignments) setaddr4r^rfrgrhr_extendloggerdebug)rUror'rqresultsrr\kernel unwrappednestedrhs_expr referencedtracedrr[find_triton_kernelsresolve_names_to_kernelsrars rrzWget_inner_triton_kernels..find_triton_kernels..resolve_names_to_kernelss=%$&G7? D!9$#D/C.s3F)v.  },(/s(;I#9j9%8 ); &F & 'v 6 (LL!JDQ ,1D$/$5%.find_triton_kernels,sc   : 6 ^^B   %K2 K I & & LLI# I &&r*F 9! fD )yy++-.5 )coo5 )nI  #(( tCy  6 ftm  f c3h &$B' =A'+ 2 92 CH~2 c4>12T92 X_ 2  &\ 2 2 j"$ U&*9+C+C&D**D  # #$;D$A B+%D5 9#8#8N&C)LL(OOC( & %#33I }}Y/H<&%i0==Hx'98Z0 !4X{ETUI!V,F "6 I 0 Y/ / -'4:_  NNK LI 2# I b  :IPT s; L 7L3!M %$M  L0/L03MM !M.-M.)Nr)rrrsintrr)rrrs @@rget_inner_triton_kernelsrsd$(,w S#X wX_ww f wr r ""rz torch.library)schema mutates_argsrcVdtdtfdtffd }||S||S)aCreate a custom operator whose implementation is backed by 1+ triton kernels. This is a more structured way of using triton kernels with PyTorch. Prefer using triton kernels with no ``torch.library`` custom operator wrappers (like :func:`torch.library.custom_op`, :func:`torch.library.triton_op`) because that is simpler; only use :func:`torch.library.custom_op`/:func:`torch.library.triton_op` if you want to create an operator that behaves like PyTorch built-in operators. For example, you may use a ``torch.library`` wrapper API to define the behavior of the triton kernel when passed a tensor subclass or under a TorchDispatchMode. Use :func:`torch.library.triton_op` instead of :func:`torch.library.custom_op` when the implementation consists of 1+ triton kernels. :func:`torch.library.custom_op` treats custom operators as opaque (:func:`torch.compile` and :func:`torch.export.export` will never trace into them), but ``triton_op`` makes the implementation visible to these subsystems, allowing them to optimize the triton kernel(s). Note that ``fn`` must only consist of calls to PyTorch-understood operators and triton kernels. Any triton kernels called inside ``fn`` must be wrapped in a call to :func:`torch.library.wrap_triton`. Args: name (str): A name for the custom op that looks like "{namespace}::{name}", e.g. "mylib::my_linear". The name is used as the op's stable identifier in PyTorch subsystems (e.g. torch.export, FX graphs). To avoid name collisions, please use your project name as the namespace; e.g. all custom ops in pytorch/fbgemm use "fbgemm" as the namespace. mutates_args (Iterable[str] or "unknown"): The names of args that the function mutates. This MUST be accurate, otherwise, the behavior is undefined. If "unknown", it pessimistically assumes that all inputs to the operator are being mutated. schema (str | None): A schema string for the operator. If None (recommended) we'll infer a schema for the operator from its type annotations. We recommend letting us infer a schema unless you have a specific reason not to. Example: "(Tensor x, int y) -> (Tensor, Tensor)". Example:: >>> # xdoctest: +REQUIRES(env:TORCH_DOCTEST_CUDA) >>> import torch >>> from torch.library import triton_op, wrap_triton >>> >>> import triton >>> from triton import language as tl >>> >>> @triton.jit >>> def add_kernel( >>> in_ptr0, >>> in_ptr1, >>> out_ptr, >>> n_elements, >>> BLOCK_SIZE: "tl.constexpr", >>> ): >>> pid = tl.program_id(axis=0) >>> block_start = pid * BLOCK_SIZE >>> offsets = block_start + tl.arange(0, BLOCK_SIZE) >>> mask = offsets < n_elements >>> x = tl.load(in_ptr0 + offsets, mask=mask) >>> y = tl.load(in_ptr1 + offsets, mask=mask) >>> output = x + y >>> tl.store(out_ptr + offsets, output, mask=mask) >>> >>> @triton_op("mylib::add", mutates_args={}) >>> def add(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor: >>> output = torch.empty_like(x) >>> n_elements = output.numel() >>> >>> def grid(meta): >>> return (triton.cdiv(n_elements, meta["BLOCK_SIZE"]),) >>> >>> # NB: we need to wrap the triton kernel in a call to wrap_triton >>> wrap_triton(add_kernel)[grid](x, y, output, n_elements, 16) >>> return output >>> >>> @torch.compile >>> def f(x, y): >>> return add(x, y) >>> >>> x = torch.randn(3, device="cuda") >>> y = torch.randn(3, device="cuda") >>> >>> z = f(x, y) >>> assert torch.allclose(z, x + y) r.rc fd}t|t}ddlm}|j fd}t }|t <|j|||S)NcVtd5|i|cdddS#1swYyxYw)NF)set_wrap_triton_enabled)rEkwargsrs r backend_fnz*triton_op..dec..backend_fns&)/4*6*0//s()r)rr)FunctionalTensorModecbddlm}|r|j||||Sddl}|Dcgc]R}t ||j j s0||j|j jjfvr|T}}|rtS|5 |i|cdddScc}w#1swYyxYw)Nr)(custom_triton_ops_decomposition_disabled) torch.export._tracer__torch_dispatch__torch._subclasses issubclass _subclasses FakeTensorTensorfunctional_tensorFunctionalTensorNotImplemented) modeoptypesrErrr?tunrecognized_typesrs rfunctional_decompz1triton_op..dec..functional_decomps& U79..r5$GG)# &"%a):):)E)EF ));;LL"# &&))t.v.T &TsAB B%%B.)r r _subclasses.functional_tensorr register_fakerrregister_torch_dispatch)rrresultrrr&rrs` rdecztriton_op..decsq +   %>   I R * /X2"5&4d#&&';=NO r)rrr )rrrrrs` ` r triton_opr(s9DGf%G+GR z 2wrTenabled)NNNc#vK t}|t_d|t_y#t_wxYww)aOIf triton kernels annotated with @wrap_triton should dispatch via HOP or go straight to the triton kernel execution. We have this switch because eager-mode performance of HOP dispatch is slow enough to matter (~1ms) and we know that wrap_triton isn't necessary in some situations (eager-mode with regular Tensors) N)is_wrap_triton_enabledwrap_triton_enabledr5)rprevs rrrs1)%'$+! $(!D!s9) 9 69c,ttdtS)Nr5)getattrrwrap_triton_enabled_defaultrNrrrrs &1L MMr triton_kernelct|S)z(This API has been renamed to wrap_triton)r>)rs rr=r=s } %%rcddlm}ddlm}ddlm}t |||fs tdts|S||ddS)aU Allows capture of a triton kernel into a graph via make_fx or non-strict ``torch.export``. These technologies perform Dispatcher-based tracing (via ``__torch_dispatch__``) and cannot see calls to raw triton kernels. The ``wrap_triton`` API wraps a triton kernel into a callable that can actually be traced into a graph. Please use this API together with :func:`torch.library.triton_op`. Examples: >>> # xdoctest: +SKIP >>> import torch >>> import triton >>> from triton import language as tl >>> from torch.fx.experimental.proxy_tensor import make_fx >>> from torch.library import wrap_triton >>> >>> @triton.jit >>> def add_kernel( >>> in_ptr0, >>> in_ptr1, >>> out_ptr, >>> n_elements, >>> BLOCK_SIZE: "tl.constexpr", >>> ): >>> pid = tl.program_id(axis=0) >>> block_start = pid * BLOCK_SIZE >>> offsets = block_start + tl.arange(0, BLOCK_SIZE) >>> mask = offsets < n_elements >>> x = tl.load(in_ptr0 + offsets, mask=mask) >>> y = tl.load(in_ptr1 + offsets, mask=mask) >>> output = x + y >>> tl.store(out_ptr + offsets, output, mask=mask) >>> >>> def add(x, y): >>> output = torch.empty_like(x) >>> n_elements = output.numel() >>> >>> def grid_fn(meta): >>> return (triton.cdiv(n_elements, meta["BLOCK_SIZE"]),) >>> >>> wrap_triton(add_kernel)[grid_fn](x, y, output, n_elements, 16) >>> return output >>> >>> x = torch.randn(3, device="cuda") >>> y = torch.randn(3, device="cuda") >>> gm = make_fx(add)(x, y) >>> print(gm.code) >>> # def forward(self, x_1, y_1): >>> # empty_like = torch.ops.aten.empty_like.default(x_1, pin_memory = False) >>> # triton_kernel_wrapper_mutation_proxy = triton_kernel_wrapper_mutation( >>> # kernel_idx = 0, constant_args_idx = 0, >>> # grid = [(1, 1, 1)], kwargs = { >>> # 'in_ptr0': x_1, 'in_ptr1': y_1, 'out_ptr': empty_like, >>> # 'n_elements': 3, 'BLOCK_SIZE': 16 >>> # }) >>> # return empty_like rrr)TraceableTritonKernelWrapperzPwrap_triton only works on functions annotated with triton.jit or triton.autotuneN) rrrr *torch._higher_order_ops.triton_kernel_wraprr/ RuntimeErrorr)rrr rs rr>r>sJ~3.W mk9%= > ^   " # ' tT BBrr)(r0 contextlibrflogging threadingcollections.abcrrrtypingrrrtorch.utils._exposed_inr custom_opsr r r getLoggerrHrvrrrrr__annotations__rrrlocalrrcontextmanagerboolrrr=r>rNrrrs 99''..&   8 $13tCf-.3/C/DL/M##s(!3M#V M#` O"m ! m mm Xc]*+ m SM mmm`&ioo'"  )T )i8H.I ) ) NN&(&#&  OICxICsICICr