Trace¶

本页暂无中文版。以下为英文原文。

tileops.trace.trace — the in-kernel timeline tracer: annotate a kernel body, wire the markers into a @tilelang.jit builder, run, and export a timeline.

For a walkthrough, see the In-Kernel Timeline Trace guide.

The trace namespace (use the singleton trace, do not instantiate).

Stateless except for the process-local run switch (enabled / output); all annotation state lives in the per-build registry.

enabled `property` ¶

enabled

Whether tracing is on for this process (default False).

Example

trace.enabled         # False

output `property` ¶

output

Output directory for dumped artifacts (default "debug").

Example

trace.output          # 'debug'

enable ¶

enable(
    output="debug",
)

Turn tracing on for this process and set the output directory.

Call once at startup, before any traced kernel is built — a builder reads enabled at build time and is typically cached.

Parameters:

output (str, default: 'debug' ) –

Directory for dumped artifacts (created on first dump). Defaults to "debug" (gitignored).

Example

trace.enable()              # dumps under debug/
trace.enable("out/traces")  # or a directory of your choosing

disable ¶

disable()

Turn tracing off for this process.

Example

trace.disable()

group ¶

group(
    name,
    lead,
)

Declare the logical work-group enclosed markers belong to.

Governs only who records (the elected writer is tx == lead); the enclosed compute still runs on all threads.

Parameters:

name (str) –

Work-group name, interned to a stable group id.
lead (int) –

Branch-baseline thread id of the electing writer.

Returns:

GroupScope –

A with context manager.

Example

with trace.group("producer", lead=0):
    with trace.range("tma"):
        ...  # only thread 0 records

range ¶

range(
    name,
    lane=DEFAULT_LANE,
    payload=None,
)

Time a span: RANGE_BEGIN on enter, RANGE_END on exit.

Parameters:

name (str) –

Range name, interned to a stable event id.
lane (str, default: DEFAULT_LANE ) –

Render sub-lane, interned dynamically (default "main").
payload –

Optional 32-bit i32 payload — a Python int or a runtime PrimExpr (e.g. a loop index). Records in both BEGIN and END events. Useful for explicitly tagging iterations in a loop. None defaults to 0, representing "no label". The payload is a user-provided tag, NOT an implicit value derived from thread/block indices.

Returns:

RangeScope –

A with context manager.

Example

with trace.range("mma", lane="wgmma"):
    T.wgmma_gemm(...)
# Explicitly tag each iteration with its index:
for i in range(N):
    with trace.range("iteration", payload=i):
        work()

range_start ¶

range_start(
    name,
    lane=DEFAULT_LANE,
    payload=None,
)

Open a range explicitly (use when with does not fit the control flow).

Parameters:

name (str) –

Range name, interned to a stable event id.
lane (str, default: DEFAULT_LANE ) –

Render sub-lane, interned dynamically (default "main").
payload –

Optional 32-bit i32 payload — a Python int or a runtime PrimExpr (e.g. a loop index). None defaults to 0, representing "no label". The payload is a user-provided tag, NOT an implicit value derived from thread/block indices.

Returns:

AnnoToken –

A token to pass to range_end.

Example

tok = trace.range_start("phase")
...  # work
trace.range_end(tok)
# Explicitly tag with payload:
tok = trace.range_start("iteration", payload=i)
...
trace.range_end(tok)

range_end ¶

range_end(
    tok,
)

Close the range opened for tok (None no-ops).

Parameters:

tok (AnnoToken | None) –

The token returned by range_start.

Example

trace.range_end(tok)

record ¶

record(
    name,
    payload=None,
    lane=DEFAULT_LANE,
)

Emit a single instant event (a zero-width mark).

Parameters:

name (str) –

Event name, interned to a stable event id.
payload –

Optional 32-bit i32 payload — a Python int or a runtime PrimExpr (e.g. a loop index). None records 0.
lane (str, default: DEFAULT_LANE ) –

Render sub-lane, interned dynamically (default "main").

Example

trace.record("iter", payload=ki)  # tag the mark with the loop index

dag ¶

dag(
    src_name,
    dst_name,
)

Declare a dependency arrow from one named range to another.

Call once in the kernel body. Emits no device marker; at render time each CTA's src_name slices are paired with its dst_name slices in timestamp order, drawing one arrow per pair.

Parameters:

src_name (str) –

Source (producer-side) range name.
dst_name (str) –

Destination (consumer-side) range name.

Example

trace.dag("arrive", "wait")  # producer "arrive" -> consumer "wait"

out_idx ¶

out_idx(
    n_outputs,
    traced=None,
)

Return the out_idx for a kernel with n_outputs real outputs.

Grows by one (for the trailing slots output that finalize appends) when traced, so the same builder works either way.

Parameters:

n_outputs (int) –

Number of real (non-slots) outputs.
traced (bool | None, default: None ) –

Whether this build is traced. None (default) reads the process switch enabled. A cached builder must pass an explicit value and include it in its cache key — otherwise the out_idx baked into the cached kernel ignores later switch flips.

Returns:

list[int] –

Negative output indices for @tilelang.jit.

Example

@tilelang.jit(out_idx=trace.out_idx(1))          # follows the switch
def factory(): ...
@tilelang.jit(out_idx=trace.out_idx(1, traced))  # cached builder
def factory(): ...

finalize ¶

finalize(
    primfunc,
    traced=None,
    max_events=MAX_EVENTS_DEFAULT,
)

Lower the markers when traced, else strip them — return either.

The one line a builder returns instead of branching. Pairs with out_idx; pass the same traced to both.

Parameters:

primfunc –

The built kernel; its body carries trace.* markers.
traced (bool | None, default: None ) –

Whether this build is traced. None (default) reads the process switch enabled. A cached builder must pass an explicit value (see out_idx).
max_events (int, default: MAX_EVENTS_DEFAULT ) –

Per-slot event capacity when lowering.

Returns:

–

A PrimFunc ready for @tilelang.jit (with a trailing slots
–

output when traced).

Example

def factory():
    @T.prim_func
    def main(...): ...
    return trace.finalize(main, traced, max_events=1024)

lower ¶

lower(
    primfunc,
    max_events=MAX_EVENTS_DEFAULT,
)

Primitive: materialize the markers and append the slots output.

Prefer finalize unless you need to lower unconditionally.

Parameters:

primfunc –

The built kernel; its body carries trace.* markers.
max_events (int, default: MAX_EVENTS_DEFAULT ) –

Per-slot event capacity (compile-time cursor bound).

Returns:

–

The lowered PrimFunc with a trailing slots int64 output.

Example

return trace.lower(main, max_events=1024)

strip ¶

strip(
    primfunc,
)

Primitive: no-op every marker so the kernel compiles without tracing.

Prefer finalize unless you need to strip unconditionally. The generated CUDA is identical to an un-instrumented build.

Parameters:

primfunc –

The built kernel; its body may carry trace.* markers.

Returns:

–

A PrimFunc with markers no-opped, signature unchanged.

Example

return trace.strip(main)

run ¶

run(
    compiled,
    inputs,
    *,
    stem
)

Run a kernel and, when tracing is on, dump its timeline.

The one call a forward needs — it branches on the switch internally, so the caller does not. With tracing off it just returns the kernel's outputs unchanged. With tracing on the kernel is the traced build (returns (*real_outputs, slots)): this splits the trailing slots tensor off, decodes it, writes the timeline via dump (a fresh file each call), and returns the real outputs only.

Parameters:

compiled –

A compiled kernel. Build it with traced=trace.enabled (via out_idx / finalize) so its outputs match the switch.
inputs (tuple) –

Tuple of positional tensors to pass to compiled.
stem (str) –

Descriptive file stem (op name + shape); see dump.

Returns:

–

The kernel's real outputs: a single tensor, or a tuple in order.

Example

compiled = build()(block=128)
c = trace.run(compiled, (a, b), stem="gemm_128x256x512")

decode ¶

decode(
    compiled,
    slots,
)

Decode a returned slots tensor into render-ready events.

Resolves the host decode maps from the compiled kernel.

Parameters:

compiled –

The compiled traced kernel (carries the host maps).
slots –

The trailing slots int64 tensor the kernel returned.

Returns:

list –

A flat list of Slice / Instant events.

Example

*_, slots = compiled(a, b)
events = trace.decode(compiled, slots)

dump ¶

dump(
    events,
    compiled,
    *,
    stem
)

Write events as an HTML timeline under output, never overwriting.

Picks a collision-free base name: {output}/{stem} when free, else {output}/{stem}_1, _2, ... so repeated dumps accumulate.

Parameters:

events (list) –

Event list from decode.
compiled –

The compiled traced kernel (carries the host maps).
stem (str) –

File stem (no directory, no extension), e.g. op name + shape.

Returns:

str –

The base path written (no extension); {base}.html exists.

Example

base = trace.dump(events, compiled, stem="gemm_128x256x512")
base                  # 'debug/gemm_128x256x512'

export_html ¶

export_html(
    events,
    path,
    *,
    compiled,
    title="",
    sm_clock_ghz=1.5
)

Write events as a self-contained Plotly HTML timeline.

Parameters:

events (list) –

Event list from decode.
path (str) –

Destination .html path.
compiled –

The compiled traced kernel (carries the host maps).
title (str, default: '' ) –

Base figure title; the CTA index is appended per tab.
sm_clock_ghz (float, default: 1.5 ) –

Locked SM clock in GHz for the ~ns hover column.

Example

trace.export_html(events, "timeline.html", compiled=compiled)

Trace¶

enabled property ¶

output property ¶

enable ¶

disable ¶

group ¶

range ¶

range_start ¶

range_end ¶

record ¶

dag ¶

out_idx ¶

finalize ¶

lower ¶

strip ¶

run ¶

decode ¶

dump ¶

export_html ¶

enabled `property` ¶

output `property` ¶