Describe the feature
Hi neuron team:
Profiling cold process startup on trn2/trn3 for diffusion serving: ~6.7s per process is spent when the first device op lazily triggers nrt_init → tdrv_init (NeuronCore reset + ready-wait, a single blocking ncdev ioctl). It is per-process and one-time, and it is not weight load, allocation, fill, or DMA (those total <1s).
Impact: in a multi-process / per-stage serving topology and under autoscaling, every new worker pays this ~6.7s before it can serve — it dominates cold-start, scale-out, and crash-recovery latency.
We've mitigated host-side as far as we can — (1) overlapping it with weight load via a background device pre-touch, (2) amortizing it in a long-lived resident worker — so the residual is entirely in the runtime/driver.
Asks, in priority order:
- Can the initialized runtime be shared/persisted across processes (a resident runtime or pre-initialized core pool a new process attaches to, near-zero) instead of re-running bring-up? Biggest lever for multi-process serving + autoscaling.
- Can the core reset + ready-wait in nrt_init/tdrv_init be shortened or skipped when cores are already healthy?
- An official async/AOT "pre-init cores" API, so we don't rely on a throwaway torch.empty(device) on a background thread.
Yotta Labs & UC Merced
Use Case
We serve diffusion models (FLUX, Qwen-Image, and video generation models) on trn2. Profiling shows that ~6.7 seconds of every process's cold start is spent in the one-time NeuronCore runtime initialization (nrt_init → tdrv_init), rather than in compute or weight loading. This directly impacts us in three ways:
- Scale-out latency and cost: In a multi-process, per-stage serving architecture with autoscaling, every new worker incurs this ~6.7-second delay before it can serve its first request. This dominates cold-start latency and increases both the time and idle hardware cost required to add capacity or recover from failures.
- Serving architecture flexibility: The per-process initialization cost pushes us toward a single long-lived, monolithic worker purely to amortize startup overhead. This makes it difficult to adopt a multi-process, disjoint-NeuronCore architecture that would enable stage pipelining and independent scaling of different pipeline stages.
- Host-side optimizations are exhausted: We already overlap initialization with weight loading and amortize it using long-lived workers. The remaining ~6.7-second overhead is entirely inside the runtime/driver, so only a shared or persistent runtime initialization mechanism can eliminate it.
Removing or sharing this initialization would reduce cold-start latency by several seconds per worker and enable a more flexible, scalable, and cost-efficient serving architecture on Trainium.
Yotta Labs & UC Merced
Proposed Solution
No response
Other Information
No response
Acknowledgements
Describe the feature
Hi neuron team:
Profiling cold process startup on trn2/trn3 for diffusion serving: ~6.7s per process is spent when the first device op lazily triggers nrt_init → tdrv_init (NeuronCore reset + ready-wait, a single blocking ncdev ioctl). It is per-process and one-time, and it is not weight load, allocation, fill, or DMA (those total <1s).
Impact: in a multi-process / per-stage serving topology and under autoscaling, every new worker pays this ~6.7s before it can serve — it dominates cold-start, scale-out, and crash-recovery latency.
We've mitigated host-side as far as we can — (1) overlapping it with weight load via a background device pre-touch, (2) amortizing it in a long-lived resident worker — so the residual is entirely in the runtime/driver.
Asks, in priority order:
Yotta Labs & UC Merced
Use Case
We serve diffusion models (FLUX, Qwen-Image, and video generation models) on trn2. Profiling shows that ~6.7 seconds of every process's cold start is spent in the one-time NeuronCore runtime initialization (nrt_init → tdrv_init), rather than in compute or weight loading. This directly impacts us in three ways:
Removing or sharing this initialization would reduce cold-start latency by several seconds per worker and enable a more flexible, scalable, and cost-efficient serving architecture on Trainium.
Yotta Labs & UC Merced
Proposed Solution
No response
Other Information
No response
Acknowledgements