RFC: Unified Profiling for Torch-TensorRT ExecuTorch Delegates

[apbose]

RFC: Unified Profiling for Torch-TensorRT ExecuTorch Delegates

Related design: RFC Discussion #4827: Global TensorRT engine profiling controls for opaque AOTInductor packages

Summary and recommendation

Torch-TensorRT currently has two deployment paths that can hide serialized TensorRT engines behind a package-level runner:

• A PyTorch AOTInductor .pt2 package executes TensorRT regions through torch.ops.tensorrt.execute_engine and core::runtime::TRTEngine.
• An ExecuTorch .pte package executes TensorRT regions through the independent native TensorRTBackend; fallback regions may execute through ExecuTorch's CudaBackend, which uses AOTInductor-generated CUDA code.

PR #4766 addresses the first path. It adds process-wide profiling controls to the core Torch-TensorRT runtime so an opaque .pt2 runner can enable profiling without exposing individual TRTEngine objects.

The recommended ExecuTorch design is not to reuse the LibTorch profiling implementation directly. Instead:

1. Use ExecuTorch EventTracer and ETDump as the common timeline for the complete .pte execution.
2. Instrument TensorRTBackend::execute() with named host-phase events.
3. Attach an optional nvinfer1::IProfiler to the execution context for TensorRT layer timings.
4. Configure detailed TensorRT profiling through ExecuTorch backend runtime options, with compile specs as an export-time fallback.
5. Treat internal CudaBackend and AOTInductor kernel profiling as an upstream ExecuTorch concern. Torch-TensorRT should validate that those events compose with its TensorRT events in a mixed-delegate ETDump.

This produces one top-level timeline without coupling the native ExecuTorch runner to LibTorch:

ExecuTorch method
    |
    |-- TensorRTBackend delegate
    |     |-- input binding and staging
    |     |-- output binding
    |     |-- enqueue host duration
    |     `-- optional TensorRT layer timings
    |
    |-- delegate boundary and copy overhead
    |
    `-- CudaBackend AOTInductor delegate
          `-- AOTInductor and CUDA events owned by ExecuTorch

Decisions requested

This RFC asks reviewers to decide the following before implementation begins:

1. Should the first implementation use ExecuTorch EventTracer and ETDump as the common timeline for all .pte execution?
2. Should profiling configuration be per-method load-time runtime options, with compile specs used only as embedded defaults?
3. Should Phase 1 be limited to stable host-phase events, leaving TensorRT layer callbacks for an independently opt-in Phase 2?
4. Should detailed TensorRT layer records be represented in ETDump, Perfetto/TREX sidecars, or both?
5. Is configuration after method load required, or can the initial design avoid process-global mutable state?
6. Which profiling levels must be supported by the native reference runner, Python runtime wheel, Android loaders, and DRIVE builds?

The recommendation in this document is: EventTracer/ETDump for the common timeline, per-method runtime options for configuration, host-phase events first, and TensorRT layer profiling as a separate opt-in phase.

Motivation

Per-engine profiling is straightforward when a caller owns a Torch-TensorRT module and can reach each engine object:

engine.enable_profiling()
engine.set_profile_format("perfetto")

Package-based deployments do not necessarily expose those objects. A PyTorch AOTInductor loader returns an opaque runner, while an ExecuTorch loader returns a method backed by native delegates. In both cases the serialized TensorRT plan exists, but profiling is runtime state attached after deserialization; it is not embedded as an active profiler in the plan.

Serialized TensorRT plan
        |
        | deserialize
        v
IExecutionContext
        |
        `-- setProfiler(...) is optional runtime configuration

The package runtime that creates the execution context must therefore attach and control the profiler.

Terminology

Term	Meaning in this document
TensorRT engine	A serialized and optimized TensorRT network for one compiled region.
TRTEngine	The Torch-TensorRT LibTorch runtime wrapper in core/runtime.
AOTInductor package	A PyTorch .pt2 artifact loaded through torch._inductor.aoti_load_package.
ExecuTorch package	A .pte program loaded by the native or Python ExecuTorch runtime.
TensorRTBackend	The standalone ExecuTorch delegate that deserializes and executes TensorRT plans without the legacy LibTorch runtime.
CudaBackend	ExecuTorch's CUDA delegate for AOTInductor-generated fallback regions.
ETDump	ExecuTorch's serialized event trace consumed by its inspection tools.

The AOTInductor .pt2 path and the ExecuTorch CudaBackend path both use AOTInductor technology, but they are not the same runtime integration.

Runtime architecture

PyTorch AOTInductor package

Torch-TensorRT GraphModule
        |
        v
torch_tensorrt.save(output_format="aot_inductor")
        |
        v
.pt2 package
        |
        v
torch._inductor.aoti_load_package
        |
        v
opaque runner
        |
        v
torch.ops.tensorrt.execute_engine
        |
        v
core::runtime::TRTEngine

This path uses the core Torch-TensorRT runtime and is the scope of PR #4766.

ExecuTorch package

Torch-TensorRT and ExecuTorch export
        |
        v
.pte program
        |
        v
ExecuTorch method
        |
        |-- TensorRT partition
        |      `-- TensorRTBackend::execute()
        |
        `-- supported fallback partition
               `-- CudaBackend and AOTInductor CUDA code

The TensorRTBackend implementation is in cpp/src/torch_tensorrt/executorch/TensorRTBackend.cpp. It owns an EngineHandle, deserializes the plan, creates its own IExecutionContext, binds ExecuTorch tensor addresses, and calls enqueueV3() directly. It does not instantiate core::runtime::TRTEngine and therefore does not inherit its profiling state.

PR #4735 enables this existing ExecuTorch deployment architecture on DRIVE OS. It does not unify the two runtime implementations.

Present implementation in PR #4766

User-facing API

PR #4766 registers process-wide C++ runtime operators:

torch.ops.tensorrt.set_profile_path_prefix("/tmp/profiles")
torch.ops.tensorrt.set_profile_format("perfetto")
torch.ops.tensorrt.set_profile_execution(True)

Matching getters expose the current state:

torch.ops.tensorrt.get_profile_path_prefix()
torch.ops.tensorrt.get_profile_format()
torch.ops.tensorrt.get_profile_execution()

The setting is global to the current process. It is not shared across processes or persisted into the package.

Configuration state

The runtime owns one snapshot:

struct GlobalProfilingConfig {
  bool enabled;
  std::string profile_format;
  std::string profile_path_prefix;
  uint64_t generation;
};

The strings and Boolean are protected by a global mutex. An atomic generation provides a low-cost unchanged path.

Setter called
    |
    |-- lock global configuration
    |-- update value
    |-- generation N -> N + 1
    `-- publish generation with release ordering

Lazy engine synchronization

Each TRTEngine stores the last generation it applied. Immediately before execution it compares that value with the global generation.

execute_engine
    |
    |-- acquire engine mutex
    |-- compare engine and global generations
    |      |
    |      |-- equal: one atomic load, continue
    |      |
    |      `-- different:
    |             |-- copy the protected global snapshot
    |             |-- update path and format
    |             |-- enable or disable the profiler
    |             `-- store the applied generation
    |
    `-- materialize and use the execution context

The engine registry is deliberately avoided. Settings work whether they are changed before or after an AOTInductor package is loaded because each engine synchronizes lazily on its next call.

Locking

PR #4766 moves the per-engine lock earlier in execute_engine() so global-state synchronization and execution-context access occur in one critical section.

engine mutex
    |-- synchronize profiling configuration
    |-- attach or remove profiler state
    |-- access execution context
    |-- bind tensors
    `-- enqueue

This prevents one execution from observing a context while another thread mutates profiler or context state.

Validation in the present PR

The C++ unit test verifies getters, setters, validation, and generation advancement. The Python integration test:

1. Compiles a Torch-TensorRT model.
2. Saves it as an AOTInductor package.
3. Loads the opaque package runner.
4. Changes the global profiling settings after loading.
5. Executes the runner.
6. Verifies that engine, input, output, and enqueue trace files are produced.

Scope and limitations

PR #4766 applies only to engines executed through core::runtime::TRTEngine. It does not configure:

• ExecuTorch's TensorRTBackend.
• ExecuTorch's CudaBackend or its AOTInductor-generated kernels.
• A Python-only Torch-TensorRT engine implementation unless equivalent operators are added there.
• Per-model filtering; every core TRTEngine in the process observes the same global snapshot.

It also does not turn asynchronous enqueueV3() host duration into GPU execution duration. Enqueue timing remains launch overhead unless an explicit synchronization or CUDA event measurement is introduced.

Proposed ExecuTorch design

Goals

• Produce a coherent profile for mixed TensorRT and AOTInductor .pte programs.
• Preserve the native runner's lack of LibTorch and libc10 dependencies.
• Make profiling opt-in and avoid synchronization solely for measurement on the normal fast path.
• Associate delegate events with the exported program where ExecuTorch debug handles permit it.
• Support native C++ runners and the ExecuTorch Python runtime wheel.
• Keep profiling configuration independent for separately loaded methods when possible.

Non-goals

• Reimplement AOTInductor kernel profiling inside Torch-TensorRT.
• Make enqueueV3() host duration appear to be GPU duration.
• Change TensorRT engine numerics, memory planning, or caller-stream behavior.
• Require LibTorch in the native ExecuTorch runtime.
• Make profiling state global across operating-system processes.

Common timeline through EventTracer and ETDump

ExecuTorch supplies an EventTracer* through BackendExecutionContext::event_tracer(). Delegate hooks compile conditionally behind ET_EVENT_TRACER_ENABLED and can remain in production source without imposing trace work when no tracer is present.

TensorRTBackend::execute() currently discards its execution context. It should instead obtain the tracer and emit stable named events:

auto* tracer = context.event_tracer();

auto event = event_tracer_start_profiling_delegate(
    tracer,
    "torchtrt.input_binding",
    kUnsetDebugHandle);

// Existing input validation, shape binding, and host-to-device staging.

event_tracer_end_profiling_delegate(tracer, event);

Recommended first-level events are:

Event	Included work
torchtrt.delegate	Entire host-side TensorRTBackend::execute() call.
torchtrt.input_binding	Shape validation, input binding, and required host-to-device staging.
torchtrt.shape_inference	inferShapes() and related validation.
torchtrt.output_binding	Output sizing, allocation selection, aliases, and binding.
torchtrt.enqueue_host	Host entry to and return from enqueueV3().
torchtrt.output_staging	Required device-to-host or reflect-copy work.
torchtrt.synchronize	Synchronization performed by the existing execution contract.

The event names explicitly distinguish host duration from GPU duration.

ETDump
  method.forward
    TensorRT delegate
      torchtrt.input_binding
      torchtrt.shape_inference
      torchtrt.output_binding
      torchtrt.enqueue_host
      torchtrt.synchronize
    CudaBackend delegate
      upstream AOTInductor and CUDA events

TensorRT layer profiling

Host-phase events do not show which TensorRT layers dominate engine execution. For that detail, EngineHandle should optionally own a runtime-neutral nvinfer1::IProfiler implementation.

TensorRTBackend::init
    |-- deserialize engine
    |-- create execution context
    |-- create profiler when requested
    `-- execution_context->setProfiler(profiler.get())

TensorRTBackend::execute
    |-- enqueue
    |-- obtain completed profiler records when safe
    `-- serialize records to ETDump metadata or a sidecar trace

The existing TRTEngineProfiler is mostly TensorRT-specific, but the core runtime's host-phase tracing uses torch::autograd::profiler::RecordProfile, which is not suitable for a no-LibTorch ExecuTorch runner. Reuse should therefore be limited to a small extracted library containing the TensorRT IProfiler implementation and trace serialization, with no ATen, TorchScript, or autograd dependency.

Layer profiling must remain independently opt-in because TensorRT layer callbacks can add overhead and can change scheduling characteristics.

Configuration

The preferred ExecuTorch control plane is backend runtime options supplied while loading the method. TensorRTBackend::init() already reads a runtime option for weight streaming through BackendInitContext::get_runtime_spec(), so profiling can follow the same precedence:

load-time backend option
        |
        | present
        v
use runtime value

otherwise
        |
        v
use compile spec embedded in .pte

otherwise
        |
        v
profiling disabled

Proposed keys are:

Key	Values	Purpose
torchtrt_profile_level	off, delegate, phases, layers	Selects profiling detail and overhead.
torchtrt_profile_format	etdump, perfetto, trex	Selects the primary or sidecar representation.
torchtrt_profile_path_prefix	non-empty path	Selects a directory for sidecar traces.

ETDump ownership remains with the runner. A backend path is needed only for optional sidecar formats that ETDump cannot represent conveniently.

An exact analogue of #4766's change-after-load semantics can be added later as a native C++ process-global configuration with a generation counter. That should not be the first choice for ExecuTorch because runtime options provide better method-level isolation and are already part of its backend contract. If Python must change configuration after loading, the runtime wheel would need a small binding for the native setter.

Debug handle mapping

TensorRTBackend.preprocess() currently returns the serialized backend blob. It should also produce a delegate debug-handle map using ExecuTorch's DelegateMappingBuilder and return it in PreprocessResult.

runtime event name or delegate identifier
        |
        v
delegate debug-handle map
        |
        v
execute_engine node and original graph region

The minimum mapping associates the delegate execution event with the execute_engine node represented by the partition. If layer-to-original-node provenance is available from TensorRT metadata, a later extension may add finer-grained mappings. Layer names alone must not be claimed to map to source operators unless that relationship is explicitly recorded during conversion.

AOTInductor and CudaBackend profiling

The CudaBackend does not execute TensorRT and never enters TensorRTBackend. Its internal profiling belongs in ExecuTorch:

CudaBackend
    |-- generated AOTInductor wrapper
    |-- CUDA kernel launches
    |-- memory and stream operations
    `-- synchronization

Torch-TensorRT's responsibility is integration validation:

• A mixed .pte produces both TensorRT and CudaBackend events in one ETDump.
• Caller-stream behavior remains correct when tracing is enabled.
• Boundary copies and synchronization remain visible at the ExecuTorch level.
• Enabling TensorRT detail does not suppress or corrupt AOTInductor events.

Internal AOTInductor kernel events, NVTX ranges, or CUPTI integration should be implemented or reused upstream rather than duplicated in Torch-TensorRT.

Async execution and measurement semantics

The ExecuTorch TensorRT backend can return without synchronizing when it uses device-resident buffers and a caller-owned CUDA stream. Profiling must preserve that contract.

enqueueV3 returns
        |
        | GPU work may still be in flight
        v
TensorRTBackend::execute returns

Consequences:

• torchtrt.enqueue_host measures launch overhead, not GPU execution.
• torchtrt.delegate measures host-visible backend duration and may end before GPU completion.
• Profiling must not add cudaDeviceSynchronize() merely to make a duration easier to interpret.
• GPU duration requires CUDA events recorded on the selected stream and consumed only after existing synchronization, a later call's in-flight wait, or teardown.
• Trace metadata should label host and GPU measurements explicitly.

The synchronized staging path can report a complete host duration because execution already waits for copies or output visibility. Results from that path must not be generalized to the asynchronous path.

Thread safety and lifetime

EngineHandle already owns a mutex and serializes access to its execution context. Profiler attachment, replacement, and result extraction must occur while holding that mutex.

EngineHandle mutex
    |-- wait for prior in-flight enqueue when required
    |-- apply profiling configuration
    |-- configure execution context
    |-- bind inputs and outputs
    |-- enqueue
    `-- record completion state

The profiler must outlive every execution that references it. Destruction must wait for the existing in-flight event before destroying the profiler or execution context. If process-global dynamic configuration is added later, it should use the same snapshot-plus-generation model as #4766 and apply snapshots under the engine mutex.

Trace filenames must include a stable engine or delegate identifier. Multi-engine programs must not have several delegates append incompatible data to the same sidecar file.

Proposed implementation phases

Phase 1: common ExecuTorch timeline

• Enable ETDump in the reference runner.
• Stop discarding BackendExecutionContext in TensorRTBackend::execute().
• Add whole-delegate and host-phase EventTracer ranges.
• Add a delegate debug-handle map in Python preprocessing.
• Validate a TensorRT-only .pte and a mixed TensorRT plus CudaBackend .pte.

This phase provides actionable host-level attribution without attaching a TensorRT layer profiler.

Phase 2: TensorRT layer detail

• Extract or implement a no-LibTorch nvinfer1::IProfiler.
• Store it on EngineHandle.
• Add runtime and compile-spec configuration.
• Export layer records as ETDump metadata or Perfetto/TREX sidecars.
• Verify disabled mode has no callbacks or file output.

Phase 3: dynamic configuration if required

• Add a native process-global setter and generation counter only if changing profiling after method load is a confirmed requirement.
• Expose that setter through the Python runtime wheel if needed.
• Preserve runtime-option precedence or document that a dynamic global update intentionally overrides it.

Phase 4: AOTInductor depth

• Reuse or add upstream CudaBackend and AOTInductor internal events.
• Validate event correlation and caller-stream behavior in the mixed model.
• Keep Torch-TensorRT changes limited to compatibility and integration tests.

Test plan

Unit tests

• Runtime options override compile specs.
• Compile specs provide defaults when runtime options are absent.
• Invalid levels, formats, and empty paths fail with actionable errors.
• No EventTracer produces no backend event work beyond compiled-out or null checks.
• phases emits the expected stable event names.
• layers attaches a profiler exactly once per context.
• Engine destruction is safe after an asynchronous enqueue.
• Multiple engines produce unique trace identifiers.

Native integration tests

• A TensorRT-only .pte emits delegate and phase events and preserves numerical output.
• A mixed TensorRT plus CudaBackend .pte emits both delegate types in one ETDump.
• Device-resident caller-stream execution remains asynchronous when profiling phases are enabled.
• Host-staged execution reports its existing synchronization without adding another one.
• Aliased KV-cache inputs continue to update caller-owned storage correctly.

Packaging and platform tests

• The native reference runner links EventTracer and ETDump without LibTorch.
• The ExecuTorch Python runtime wheel can enable the selected profile level.
• DRIVE builds select the same profiling sources and retain the dedicated CUDA and TensorRT dependencies introduced by PR Add DRIVE OS 7.2.4 build support for ExecuTorch deployment #4735.
• Profiling-disabled builds remain available when ET_EVENT_TRACER_ENABLED is off.

Alternatives considered

Reuse torch.ops.tensorrt in ExecuTorch

Rejected for the native runner. torch.ops is a LibTorch registration mechanism, while the ExecuTorch deployment intentionally avoids LibTorch. The Python runtime wheel could expose a binding, but the underlying implementation still needs a native ExecuTorch control path.

Link the complete core Torch-TensorRT runtime into ExecuTorch

Rejected. It would defeat the standalone backend architecture, increase dependency size, and risk pulling in ATen, TorchScript, and autograd profiling dependencies. A small runtime-neutral TensorRT profiler may be shared instead.

Always attach the TensorRT layer profiler

Rejected. Profiling has overhead and may affect execution scheduling. Layer detail must be opt-in.

Force synchronization to obtain simple GPU durations

Rejected. It would materially change the asynchronous fast path being measured. CUDA-event results should be collected only at a point where the runtime already knows the work is complete.

Implement AOTInductor profiling inside Torch-TensorRT

Rejected. The CudaBackend owns that execution path. Torch-TensorRT should compose with its events and test the mixed workflow rather than fork its profiling implementation.

Open questions

1. Is change-after-load configuration a requirement for ExecuTorch, or are load-time runtime options sufficient?
2. Should TensorRT layer records be represented as ETDump delegate metadata, a Perfetto/TREX sidecar, or both?
3. Does the pinned CudaBackend already emit sufficient kernel-level events for the target performance investigation?
4. What stable engine identifier should be used for filenames and cross-run comparison?
5. Is per-engine filtering required, or is per-method profiling sufficient?
6. Which profiling levels must be supported in the Python runtime wheel, Android loaders, and the native reference runner?
7. Should detailed profiling be available in production builds or only builds with ET_EVENT_TRACER_ENABLED?

Success criteria

The design is complete when a mixed .pte program can be run through the native reference runner and produce a trace that answers all of the following without changing model results or normal stream semantics:

• Which TensorRT and AOTInductor delegate regions ran?
• How much host time did each delegate call consume?
• Within each TensorRT delegate, how much time was spent on binding, staging, enqueue, and synchronization?
• When layer profiling is enabled, which TensorRT layers dominate reported engine time?
• What time is spent between delegates on copies or runtime orchestration?
• Can those events be associated with the corresponding exported graph regions?

PR #4766 satisfies the opaque .pt2 control problem for the core runtime. The proposed ExecuTorch phases provide the corresponding .pte observability while respecting the separate runtime architecture.

References

• PR #4766: global engine profiling controls
• Issue #4679: profiling engines hidden inside an AOTInductor package
• PR #4735: DRIVE OS ExecuTorch deployment support
• ExecuTorch delegate debugging
• ExecuTorch BackendExecutionContext
• ExecuTorch delegate EventTracer hooks