ȅiUZ %SSKrSSKrSSKrSSKrSSKrSSKJrJrJr SSK J r J r J r SSK Jr SSKJrJr SSKJr \R&"\5r0r\\\\4\S'S \S \\4S jrS \S \ 4S \\4Sjr\"S5SSS.S \S \ \S\ \\\4S\ \S \4 Sjjj5r\R>"5r Sr!\RDS\#S \S4Sj5r$S \#4Sjr%S\S \ 4Sjr&\"S5S\S \ 4Sj5r'g)N)Callable GeneratorIterable)AnyOptionalUnion) exposed_in) custom_op CustomOpDef) infer_schematriton_ops_to_kernelsnamereturnc.[RU/5$N)rget)rs O/home/james-whalen/.local/lib/python3.13/site-packages/torch/_library/triton.pyget_triton_kernels_for_oprs $ $T2 ..fn.c ^^SmS S[S[4S[[S-S[S[[ 4UU4SjjjmT"U5$) aO 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>^^^^^^^^SSKJm SSKJm Uc [5n[U5nX1;a/$UT:a[ RST5 /$URU5 [R"U5nSSKJn U"5nUR#USS9 [$R&"UR)55n"S S [$R*5nU"5mTR-U5 [R."U5m0mTR1TR25 [5US 5(aTR1UR65 S [$R8S [:[<4SjmS[>S [>S-4UU4SjjmSS[<S[[<S-S [:[>4UUUUUU4Sjjjm/n [5n TR@HHn T"U 5n U H7n [U 5nX;dMU RU5 U RCU 5 M9 MJ TRDHnSnUT;aTUnOUTRF;aTRFUnUcSSK$J%n UU;aUURLnUb[OU5(dMc[5US5(dMvT"UXS-5nUH8n[U5nUU ;dMU RU5 U RCU5 M: M U $![a [ R S5 /s$f=f![[4a /s$f=f![Pa [ RSUSS9 GM)f=f)Nr Autotuner JITFunctionz.Triton not available, find_triton_kernels = []z7reached max recursion depth (%s) in find_triton_kernels)IndentedBufferT)stripcj\rSrSrS SjrS\R SS4SjrS\RSS4Sjr Sr g) Fget_inner_triton_kernels..find_triton_kernels..VisitorTrNc./Ul0Ul/Ulgr)triton_kernels assignmentscalled_functions)selfs r__init__Oget_inner_triton_kernels..find_triton_kernels..Visitor.__init__Us13#>@ 35%rnodec URHcn[U[R5(dM$URR UR /5RUR5 Me URU5 gr) targets isinstanceastNamer( setdefaultidappendvalue generic_visit)r*r-targets r visit_AssignSget_inner_triton_kernels..find_triton_kernels..Visitor.visit_Assign\sZ"llF!&#((33((33FIIrBII$**U+""4(rcSn[UR[R5(Ga)URn[UR[R5(Ga[URR[R 5(aURRR S:XaURRS:XaURU;aqUR(a^[URS[R 5(a2URRURSR 5 GO[URR[R5(a[URRR[R 5(aURRRR S:XabURRRS:Xa>URRURRSUR35 O[UR[R 5(aURR U;apUR(a^[URS[R 5(a2URRURSR 5 O/URRURR 5 URU5 g)N)capture_triton wrap_tritontorch_libraryropsz::) r0funcr1 Attributer6r2r4attrargsr'r5r)r7)r*r-triton_func_namesrCs r visit_CallQget_inner_triton_kernels..find_triton_kernels..Visitor.visit_Callbs$E!dii7799D!$**cmm<<&tzz'7'7BB $ 0 0 3 3w > $ : = $ -> >#yyZ ! chh-O-O $ 3 3 : :499Q 1188#'::??"32dii[ A   38844yy||'8899DIIaL#(()K)K //66tyy|G--44TYY\\B""4(r)r(r)r')rN) __name__ __module__ __qualname____firstlineno__r+r1Assignr9CallrF__static_attributes__rrVisitorr$Ts5 6 ) ) ) )sxx )D )rrP __globals__exprrcp^/m"U4SjS[R5nU"5RU5 T$)z3Extract all Name references from an AST expression.ch>\rSrSrS\R SS4U4SjjrS\RSS4SjrSr g)mget_inner_triton_kernels..find_triton_kernels..extract_names_from_expr..NameExtractorr-rNc<>TRUR5 gr)r5r4)r*r-namess r visit_Namexget_inner_triton_kernels..find_triton_kernels..extract_names_from_expr..NameExtractor.visit_NamesLL)rc&URU5 gr)r7)r*r-s rrFxget_inner_triton_kernels..find_triton_kernels..extract_names_from_expr..NameExtractor.visit_Calls&&t,rrO) rHrIrJrKr1r2rYrMrFrN)rXsr NameExtractorrUs1*388**-388--rr])r1 NodeVisitorvisit)rRr]rXs @rextract_names_from_exprVget_inner_triton_kernels..find_triton_kernels..extract_names_from_exprs/!E - - O ! !$ 'Lrobjc>[UTT45(aU$[U5(a0[US5(aURn[UT5(aU$g)zACheck if obj is a triton kernel or wrapper and return the kernel.rN)r0callablehasattrr)rbinnerrr s rresolve_to_kernelPget_inner_triton_kernels..find_triton_kernels..resolve_to_kernelsM# Y788 }}d!3!3e[11 Lrrvisitedc>Uc [5nX;a/$URU5 UT;a)T "TU5nUc[RSU5 /$U/$UT R;a3T "T RU5nUc[RSU5 /$U/$UT R ;a3T "T R U5nUc[RSU5 /$U/$UT R ;a[RSU5 /$/nT R UH-nT "U5nUHnT "Xa5nURU5 M M/ U$)a Trace a name through local assignments back to global triton kernels. This handles patterns like: kernel_fn = _my_kernel # global wrapped = wrapper(kernel_fn) autotuned = autotune(wrapped) capture_triton(autotuned) # traces back to _my_kernel z4failed to resolve all_globals[%s] to a triton kernelz?failed to resolve closure_vars.nonlocals[%s] to a triton kernelz>failed to resolve closure_vars.builtins[%s] to a triton kernelz%s not in collector.assignments)setaddloggerwarning nonlocalsbuiltinsr(extend)rrikernelresultsrhs_expr referencedref_nametraced all_globals closure_vars collectorr`rgtrace_to_global_kernelss rr{Vget_inner_triton_kernels..find_triton_kernels..trace_to_global_kernelssj% KK {"*;t+<=>NNNPTIx|---*<+A+A$+GH>NNYIx|,,,*<+@+@+FG>NNXIx9000@$G %'G%11$74X> *H4XGFNN6*!+8 Nr)OPDEFS__code__r z$failed to analyze called function %s)exc_infor))triton.runtime.autotunerrtriton.runtime.jitr ImportErrorrmrnrkr4debugrlinspect getsourceOSError TypeErrortorch._inductor.utilsr!splicer1parse getrawvaluer^r_getclosurevarsupdateglobalsrerQrRliststrobjectr'r5r)rotorch._library.custom_opsr} _abstract_fnrd Exception)rrrfn_idsourcer!buffertreerPresolvedseen_idsrtraced_objectsrbobj_id func_namefunc_objr}nested_kernelsrr kernel_idrr rxryrzr`rgr{MAX_RECURSION_DEPTHfind_triton_kernelss @@@@@@@@rr5get_inner_triton_kernels..find_triton_kernels,s*   : 6  %K2  I & & LLI# I &&r*F 9! fD )yy++-.. )coo. )`I --b1 ') <//0 2} % %   r~~ . #(( tCy  6 ftm  37A A  #C4A  &\A A H"$ U,,D4T:N%C)LL(OOC( &-#33IHK'&y1l444'11)<<&%i0==Hx'9'98Z00 !4X{TUI!V,F "6 I 0 Y/ / -14DK  NNK LI ,# I T  :IPT s; K1&L! L0&L01!LLL-,L-0MM)Nr)rrrkintrr)rrrs @@rget_inner_triton_kernelsrsj$(,m S#X mX_mm f mm^ r ""rz torch.library)schema mutates_argsrc\^^S[S[4S[4UU4SjjnUcU$U"U5$)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 >^U4Sjn[TUT[TTS9S9nSSKJn UR T5 U4Sjn[ T5nU[ T'URX45 U$)Nc`>[S5 T"U0UD6sSSS5 $!,(df  g=f)NF)set_wrap_triton_enabled)rDkwargsrs r backend_fn*triton_op..dec..backend_fns&)/4*6*0//s -)r)rr)FunctionalTensorModec>SSKJn U"5(aURXX45$SSKnUVs/sH]n[ XvR R 5(aM)UURUR RR4;dM[UPM_ nnU(a[$U T "U0UD6sSSS5 $s snf!,(df  g=f)Nr)(custom_triton_ops_decomposition_disabled) torch.export._tracer__torch_dispatch__torch._subclasses issubclass _subclasses FakeTensorTensorfunctional_tensorFunctionalTensorNotImplemented) modeoptypesrDrrr>tunrecognized_typesrs rfunctional_decomp1triton_op..dec..functional_decomps& U799..r$GG)# &"%a):):)E)EF ));;LL"# &&))t.v.T &Ts(B5.B5 B5#B:: C)r r _subclasses.functional_tensorr register_fakerrregister_torch_dispatch)rrresultrrr'rrs` rdectriton_op..decsn +   %>   I R * /X2"5&4d#&&';O r)rrr )rrrrrs` ` r triton_oprs?DGf%G+GGR z 2wrTenabled)NNNc#v# [5nU[lSv U[lg!W[lf=f7f)a;If 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_enabledr6)rprevs rrrs0)%'$+! $(!D!s9) 9 69c,[[S[5$)Nr6)getattrrwrap_triton_enabled_defaultrOrrrrs &1L MMr triton_kernelc[U5$)z(This API has been renamed to wrap_triton)r=)rs rr<r<s } %%rcSSKJn SSKJn SSKJn [ XU45(d [S5e[5(dU$U"USS5$)ayAllows 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_wraprr0 RuntimeErrorr)rrr rs rr=r=sN~3.W m9%= > > ^   " # # ' tT BBrr)(r1 contextlibrlogging threadingcollections.abcrrrtypingrrrtorch.utils._exposed_inr custom_opsr r r getLoggerrHrmrdictrrr__annotations__rrrlocalrrcontextmanagerboolrrr<r=rOrrrs 99''..&   8 $13tCf-.3/C/DL/C##s(!3C#V C#L O"m ! m mm Xc]*+ m SM mmm` oo'"  )T )i8H.I ) ) NN&(&#&  OICxICsICICr