+ jZNRt^RIHt^RIt^RIt^RIHtHt^RIH t H t H t ^RI t ] 'd^RIHt.R/OtR t] P$P']]4t] !R 4t]!R R 7!R R44tRRltRRltRRltRRltRRltRRltRR/RRlltRRlt!RR4t RR/R R!llt!R"R#lt"!R$R4t#RR/R%R&llt$RR/R'R(llt%RR/R)R*llt&R+R,RR/R-R.llt'R#)0a Structured compile/cache instrumentation for ``torch._native`` ops. Native DSL ops compile device kernels lazily on first call. Those compiles are the dominant first-call latency, and a silent cache miss is a common "why is this slow again?" question. This module surfaces both, with no runtime cost unless ``TORCH_LOGS=+native_dsl_compile`` is set. That artifact is off by default to avoid log spew: these wrappers sit on the compile *cache*, so they run on every op call, cache hits included. Two sinks, both gated on that artifact and fed by a single :class:`CompileEvent`: * The ``native_dsl_compile`` artifact logger: a one-line human-readable summary per compile -- outcome, wall time, and running hit/miss totals. * ``trace_structured`` artifacts (tlparse): a JSON record per compile, for production jobs where only the structured trace is retrievable. The DSLs compile through different machinery, so each has an entry point. All reduce to the same shared core (:func:`_make_wrapper`): snapshot the cache, time the call, snapshot again, and flag ``compiled`` when the miss counter advanced. They differ only in how a snapshot is sampled: * :func:`instrument_cutedsl_compile` -- for CuTeDSL, stacked *above* the vendored ``quack`` ``@jit_cache`` decorator. It reads the cache wrapper's ``cache_info()`` and times the wrapped ``cute.compile`` call:: @instrument_cutedsl_compile("aten::topk") @jit_cache def _compile_topk_radix(N, K, deterministic): ... A ``cache_info().misses`` delta means the cache ran a real ``cute.compile``, so the measured wall time *is* the compile time; otherwise the key was served from the in-memory or on-disk ``.o`` cache. * :func:`instrument_triton_kernel` -- for Triton ``@triton.jit`` kernels, which compile *and* launch in one ``kernel[grid](...)`` call and keep their own per-kernel cache (``JITFunction.device_caches``). Stacked above ``@triton.jit``, it wraps the kernel in a transparent proxy that watches *that one kernel's* variant count; a launch that grows it means a fresh compile. Scoping to the single kernel gives clean per-kernel attribution -- two kernels in one module never collide. Because compile and launch are fused, ``wall_ms`` on a miss is compile + host-launch latency (compile dominates); on a hit it is just host-launch latency. * :func:`instrument_helion_kernel` -- for regular and AOT Helion kernels. It tracks compiled configurations across the kernel's bound specialization cache; a growing count means the call compiled a new configuration. The DSLs only expose miss-side signal directly (CuTeDSL's vendored cache reports aggregate counters; Triton's cache only grows), so finer reasons (disk-hit vs lock-timeout, Triton's on-disk cache) are not distinguished here -- the boolean ``compiled`` flag plus wall time covers the common case. Neither entry point touches the underlying DSL/vendored code, and neither hijacks Triton's process-global ``knobs.runtime`` hooks (those would also capture Inductor's unrelated Triton compiles). ) annotationsN)asdict dataclass)Any TYPE_CHECKINGTypeVar)Callable CompileEventInstrumentedHelionKernelInstrumentedTritonKernelnative_dsl_compileRT)frozencx]tRt^_t$RtR]R&R]R&R]R&R]R&R]R &R]R &R ]R &R ]R &RRltRtR#)r aOne compile-function invocation, as recorded to logs and tlparse. ``compiled`` is the ground truth (did the cache run a real compile); ``outcome`` is its human-readable form. ``hits`` / ``misses`` are the cache's running totals *after* this call, useful for spotting churn (e.g. misses climbing across calls of the same shape => keys not stable, or the persistent cache is disabled). stropdsloutcomeboolcompiledfloatwall_mskeyinthitsmissescV^8dQhRR/#)returnzdict[str, Any])formats"^/home/luca/screen-checking/.venv/lib/python3.14/site-packages/torch/_native/instrumentation.py __annotate__CompileEvent.__annotate__ssc \V4#N)rselfs&r!as_dictCompileEvent.as_dictss d|r$rN)__name__ __module__ __qualname____firstlineno____doc____annotations__r)__static_attributes__rr$r!r r _s; G H LN N H I Kr$cV^8dQhRR/#)rrrr)r s"r!r"r"wsCCDCr$cL^RIHpVPP\4#)aRTrue if the ``native_dsl_compile`` artifact is enabled. Lets the wrapper skip all instrumentation work (cache sampling, timing, key formatting, event construction) on the hot path when nothing is listening -- important because the CuTeDSL wrapper sits on the compile *cache* and so runs on every op call, cache hits included. Both sinks share this one gate; tlparse emission cannot be left to ``trace_structured``'s internal ``trace_log.handlers`` check, which is true in any job collecting a structured trace. ``log_state`` is read off the module because ``TORCH_LOGS`` parsing rebinds it.  _internal)torch._loggingr5 log_stateis_artifact_enabled_ARTIFACT_NAMEr4s r! _listeningr:ws)    2 2> BBr$c V^8dQhRRRR/#)reventr rNoner)r s"r!r"r"s$r$c a\PRSPSPSPSP SP SPSP4^RI H pV!RRRV3RlR7R#) z5Fan the event out to the artifact logger and tlparse.z6%s [%s] %s in %.1fms (key=%s, cache hits=%d misses=%d))trace_structuredartifactcR\RR/#)nameencodingjson)r9rr$r!_emit..s V^ZHr$Fc<\S4#r&) _json_payload)r<sr!rErFs =/r$) metadata_fnexpect_trace_id payload_fnN) loginforrrrrrrtorch._logging._internalr?)r<r?sf r!_emitrOs^HH@        :H/ r$c V^8dQhRRRR/#)rr<r rrr)r s"r!r"r"s777#7r$cL^RIpVPVP4RR7#)NT) sort_keys)rDdumpsr))r<rDs& r!rHrHs ::emmo: 66r$c(V^8dQhRRRRRRRR/#) rargstuplekwargsdictkey_fnzCallable | Nonerrr)r s"r!r"r"s(((e(T(?(s(r$c:Ve V!V/VB#VUu.uFp\V4NK ppT\VP44UUu.uF wrVV RV: 2NK upp, pRRP V4,R,# \dLi;iuupiuuppi)N=(z, )) Exceptionreprsorteditemsjoin)rVrXrZapartskvs&&& r! _format_keyrhs  4*6* *# #dT!WdE # vflln'= >'=tq1QEl'= >>E 5! !C ''     # >sBBB BBc 0V^8dQhRRRRRRRRRR R R/#) rfnCallable[..., R]rrrrZCallable[..., str] | Nonesample+Callable[[], tuple[int | None, int | None]]rr)r s"r!r"r"sD::: : : & : 8 :  :r$c\aaaaa\P!S4RVVVVV3Rll4pV#)aIShared instrumentation core for both DSL entry points. ``sample()`` returns a ``(hits, misses)`` snapshot of the relevant cache; a ``misses`` increase across the call means a real compile fired. Timing, error handling, classification, and emission are identical across DSLs -- only ``sample`` (and the reported ``dsl``) differ. c$V^8dQhRRRRRR/#)rrVrrXrr r)r s"r!r"#_make_wrapper..__annotate__s!((s(c(a(r$cZ<\4'g S !V/VB#S!4wr#\P!4pS !V/VBpRp\P!4T, R,pS!4wrT'*;'dTRJ;'dT RJ;'dY8p T'dRp M T 'dRMRp \ \ SS T T T\ YS4T;'g^T ;'g^R74T# \dRphi;i \P!4T, R,pS!4wrX'*;'dTRJ;'dT RJ;'dY8p T'dRp M T 'dRMRp \ \ SS T T T\ YS4T;'g^T ;'g^R74i;i)FTg@@Nerrorr cache_hit)rrrrrrrr)r:time perf_counterr_rOr rh)rVrX_ misses_beforestartresultoutcome_is_errorr hits_after misses_afterrrrrjrZrrms*, r!wrapper_make_wrapper..wrappers ||t&v& &!8!!# ((F$  ((*U2c9G'-x $J$$11!-11 ,11!0   !(0*k #%##D&9#q',,1   9 #   ((*U2c9G'-x $J$$11!-11 ,11!0   !(0*k #%##D&9#q',,1  s6 C..C==DAF* F* F* F*+1F* F*) functoolswraps)rjrrrZrmr~sfffff r! _make_wrapperrs-__R(((T Nr$c V^8dQhRRRR/#)rrjrrrnr)r s"r!r"r"sC$Or$caRV3RllpV#)aSampler reading ``(hits, misses)`` from a ``jit_cache`` wrapper. Defensive: the wrapped callable may not be ``jit_cache``-decorated (e.g. a plain function in tests). Then we report ``(None, None)`` and the core still times the call, just without compile/hit classification. cV^8dQhRR/#rrztuple[int | None, int | None]r)r s"r!r")_cache_info_sampler..__annotate__s1r$c<\SRR4pVfR#V!4pVPVP3# \dRu#i;i) cache_infoN)NN)getattrrrr_)rMcirjs r!rm#_cache_info_sampler..samplesPr<. <  B77BII% %   s5 AAr)rjrmsf r!_cache_info_samplerrs Mr$rZc$V^8dQhRRRRRR/#rrrrZrlrz.Callable[[Callable[..., R]], Callable[..., R]]r)r s"r!r"r" s(  &4 r$caaRVV3RllpV#)a^Instrument a CuTeDSL (``@jit_cache``-decorated) compile function. Args: op: Operator symbol being compiled for, e.g. ``"aten::topk"``. key_fn: Optional callable with the wrapped function's signature returning a short string describing the compile key for logs. Defaults to a repr of the args/kwargs. Returns a decorator. The decorated function behaves identically to the original (same return value, same caching); it only adds a log line and a tlparse artifact per call. Errors raised by the wrapped compile are timed, reported with ``outcome="error"``, and re-raised unchanged. c V^8dQhRRRR/#rrjrkrr)r s"r!r"0instrument_cutedsl_compile..__annotate__ s&+;r$c <\VSRS\V44pRF+p\W4'gK\W\ W44K- V#)cutedsl)cache cache_clearr)rrhasattrsetattrr)rjr~attrrZrs& r! decorator-instrument_cutedsl_compile..decorator sHB 6;Nr;RS;Dr  wr'89;r$rrrZrsfd r!instrument_cutedsl_compiler s& r$c V^8dQhRRRR/#rkernelrrz int | Noner)r s"r!r"r"-sszr$c\VRR4pVfR#\RVP444# \dR#i;i)a Compiled-variant count across one JITFunction's per-device caches. Triton stores one entry per specialized variant in ``JITFunction.device_caches[device]``, a ``(kernel_cache, ...)`` tuple whose first element is the dict of compiled kernels. The count grows by one each time a new (signature, constexpr, options) variant compiles, so a delta across a launch tells us whether *this* kernel compiled. Returns None if the object doesn't expose ``device_caches`` (a future Triton that renames this, or a non-kernel in tests). device_cachesNc3F"TFp\V^,4xK R#5i)rRN)len).0 per_devices& r! %_triton_cache_size..<sH*3z!}%%s!rsumvaluesr_)rcachess& r!_triton_cache_sizer-sHV_d 3F ~H HHH 5 AAc`]tRtRtRtRRlt]RRl4tRRltR R lt R R lt R t R#)r iAa&Transparent proxy over a single ``@triton.jit`` kernel. Watches only this kernel's ``device_caches``, so a launch that grows the variant count is unambiguously *this* kernel compiling -- two kernels in one module never collide. All attribute access except the ``kernel[grid]`` launch path is delegated to the wrapped kernel, so the proxy is a drop-in for the original ``JITFunction`` for attribute/method use (``.warmup``, ``.cache_key``). Composition with ``torch.library.wrap_triton``: that API does a strict ``isinstance(JITFunction)`` check and would reject this proxy, but it targets a different path (export/tracing, where the compile does not run through our eager launch). When a native op needs ``wrap_triton``, pass the raw kernel via :attr:`jit_kernel`: ``wrap_triton(my_kernel.jit_kernel)``. c(V^8dQhRRRRRRRR/# rrrrrrZrlrr=r)r s"r!r"%InstrumentedTritonKernel.__annotate__Rs2 *   r$c *WnW nW0nR#r&)_kernel_op_key_fnr(rrrZs&&&&r!__init__!InstrumentedTritonKernel.__init__Rs   r$cV^8dQhRR/#rrrr)r s"r!r"r]sCr$c VP#)z;The wrapped raw ``JITFunction`` (e.g. for ``wrap_triton``).rr's&r! jit_kernel#InstrumentedTritonKernel.jit_kernel\s||r$cV^8dQhRR/#rr)r s"r!r"ras6666r$c 0R\VP43#r&)rrr's&r!_sample InstrumentedTritonKernel._sampleas' 555r$c V^8dQhRRRR/#)rgridrrzCallable[..., Any]r)r s"r!r"rfsWWW(:Wr$c VPV,p\W PRVPVP4#triton)rrrrr)r(rlaunchers&& r! __getitem__$InstrumentedTritonKernel.__getitem__fs/<<%Xxx4<<VVr$c V^8dQhRRRR/#rrBrrrr)r s"r!r"rls++++r$c .\VPV4#r&rrr(rBs&&r! __getattr__$InstrumentedTritonKernel.__getattr__lst||T**r$)rrrN) r+r,r-r.r/rpropertyrrrrr1rr$r!r r As5 6 W ++r$c$V^8dQhRRRRRR/#)rrrrZrlrz)Callable[[Any], InstrumentedTritonKernel]r)r s"r!r"r"rs(  &/ r$caaRVV3RllpV#)a}Instrument a single ``@triton.jit`` kernel, stacked above the jit:: @instrument_triton_kernel("aten::bmm") @triton.jit def _bmm_kernel(...): ... Unlike CuTeDSL, a Triton kernel compiles lazily *inside* its ``kernel[grid](...)`` launch and caches variants on the kernel object. So rather than wrap a separate compile fn, wrap the kernel itself: the proxy watches that one kernel's variant count and records a compile when the launch grows it. Because compile and launch are fused, ``wall_ms`` on a miss is compile + host-launch latency (compile dominates); on a hit it is just host-launch latency. Args: op: Operator symbol being compiled for, e.g. ``"aten::bmm"``. key_fn: Optional callable with the kernel's launch-arg signature returning a short string for logs (the constexprs it sees, e.g. ``BLOCK``, are the actual compile key). Defaults to a repr. c V^8dQhRRRR/#)rrrrr r)r s"r!r".instrument_triton_kernel..__annotate__<<#<":.decorator'F;;r$rrsfd r!instrument_triton_kernelrrs4<< r$c V^8dQhRRRR/#rr)r s"r!r"r"sszr$c\VRR4pVfR#\RVP444# \dR#i;i)aCompiled-config count across a Helion Kernel's bound specializations. Helion stores compiled configs in ``Kernel._bound_kernels[*]._compile_cache``; the total grows by one each time a new configuration compiles, so a delta across a call tells us whether *this* kernel compiled. Returns None if the object doesn't expose these private attributes (a future Helion that renames them, or a non-kernel in tests). _bound_kernelsNc3N"TFp\\VRR44xK R#5i)_compile_cacheNr)rr)rbounds& r!r%_helion_cache_size..s) / /4 5 5/s#%r)r bound_kernelss& r!_helion_cache_sizersXF$4d;M &--/    rc`]tRtRtRtRRlt]RRl4tRRltR R lt R R lt R t R#)r iz9Transparent proxy over a Helion ``Kernel`` compile cache.c(V^8dQhRRRRRRRR/#rr)r s"r!r"%InstrumentedHelionKernel.__annotate__s8RRR R* R  Rr$c JWn\WRW0P4VnR#)helionN)rrr_wrappedrs&&&&r!r!InstrumentedHelionKernel.__init__s  %f(FLLQ r$cV^8dQhRR/#rr)r s"r!r"rssr$c VP#r&rr's&r! helion_kernel&InstrumentedHelionKernel.helion_kernels ||r$cV^8dQhRR/#rr)r s"r!r"rs6666r$c 0R\VP43#r&)rrr's&r!r InstrumentedHelionKernel._samples' 555r$c$V^8dQhRRRRRR/#)rrVrrXrr)r s"r!r"rs!..c.S.S.r$c &VP!V/VB#r&)r)r(rVrXs&*,r!__call__!InstrumentedHelionKernel.__call__s}}d-f--r$c V^8dQhRRRR/#rr)r s"r!r"rs++++r$c .\VPV4#r&rrs&&r!r$InstrumentedHelionKernel.__getattr__st||T**r$)rrN) r+r,r-r.r/rrrrrrr1rr$r!r r s3CR6.++r$c$V^8dQhRRRRRR/#)rrrrZrlrz)Callable[[Any], InstrumentedHelionKernel]r)r s"r!r"r"s(     & / r$caaRVV3RllpV#)z@Instrument a Helion ``Kernel`` without importing Helion eagerly.c V^8dQhRRRR/#)rrrrr r)r s"r!r".instrument_helion_kernel..__annotate__rr$c<\VSS4#r&)r rs&r!r+instrument_helion_kernel..decoratorrr$rrsfd r!instrument_helion_kernelrs<< r$c$V^8dQhRRRRRR/#rr)r s"r!r"r"s(  &4 r$c.aaa^RIHoRVVV3RllpV#)aCache + instrument a CuTeDSL compile function in one decorator:: @instrumented_cutedsl_cache("aten::topk") def _compile_topk_radix(N, K, deterministic): return cute.compile(...) Equivalent to ``instrument_cutedsl_compile(op) `` stacked above the vendored ``@jit_cache``. ``jit_cache`` is imported lazily so plain ``import torch._native`` doesn't pull in the CuTeDSL runtime. ) jit_cachec V^8dQhRRRR/#rr)r s"r!r"0instrumented_cutedsl_cache..__annotate__s LL&L+;Lr$c6<\SSR7!S!V44#rZ)r)rjrrZrs&r!r-instrumented_cutedsl_cache..decorators)"V.__annotate__s KK&K+CKr$cH<\SSR7!SPV44#r )rjit)rjrZrrs&r!r,instrumented_triton_cache..decorators'6:6::b>JJr$r)rrZrrsfd @r!instrumented_triton_cachers KK r$aotFc ,V^8dQhRRRRRRRRR R /#) rrrrrrZrlsettingsobjectrz6Callable[[Callable[..., R]], InstrumentedHelionKernel]r)r s"r!r"r"s<   &    < r$c aaaa^RIp^RIpV'dVPPM VPoRVVVV3RllpV#)z5Create and instrument a regular or AOT Helion kernel.Nc V^8dQhRRRR/#)rrjrkrr r)r s"r!r"0instrumented_helion_kernel..__annotate__s  & +C r$cB<\SSR7!S!R/SB!V44#)r r)r)rjkernel_decoratorrZrrs&r!r-instrumented_helion_kernel..decorators''6:  (x ( ,  r$)rhelion.experimental experimental aot_kernelr)rrrZrrrrsf$dl @r!instrumented_helion_kernelr$s79r-s8t# )..( & nn&&x@ CL $.C(>7 (:z*)-@(.+.+b)-@*++4 )- ,)-0)-0)- r$