剖析火炬内龙(第一部分):AWS列车的系统级剖面。
许多ML从业者对如何让模型运行得很快有些直觉，但较少人知道如何分析一个模型为何缓慢。 在这个系列的第一部分中，你将学习如何系统地分析推论的性能，并找出时间花在什么地方。
这个帖子从零开始，并假定没有之前的剖析经验。 行走道使用AWSTrainium，这是专为深层学习培训和推论工作量而专门制造的AI芯片，以及火炬内龙，它的开源本土PyTorch后端。 火炬内龙支持Trainium上的标准PyTorch工作流程，包括急速执行，火炬。 编译，以及标准PyTorch API，最小代码更改。 此外，火炬内龙与Neuron Kernel接口(NKI)集成，使Trainium硬件的自定义操作器和性能优化成为可能。 到最后，你应该知道：
如何配置火炬。 profile 与 Neuron 后端并捕捉干净的追踪。
- 如何读取按键平均()表，以识别哪些操作消耗了Neuron设备时间，并比较其Self CPU时间。
- 如何通过Neuron Explorer跟踪一个操作，从框架发送到运行时间调度，到NeuronCore上的硬件执行。
- 同一行动的三个执行阶段如何改变这些信号。
PyTorch 后端的神经元
在描述工作量之前，它有助于理解Neuron如何与PyTorch融合。
Neuron后端使用PyTorch的PrivateUse1扩展点。 这种机制让一个加速器与框架整合，而不增加一个专门的发送键，如cuda. 火炬 neuronx包与上游PyTorch合作，将Neuron注册为PyTorch代码的本土后端。 关于这一扩展机制的详细情况，请参见PrivateUse1的促进新后端集成。
设置配置器
您使用两个对象来配置 Neuron 工作量：
- NeuronConfig控制哪条Neuron数据流和NeuronCores来捕捉和写入输出。 你把它传递给火炬。 profile. profile 作为实验 config.
- NeuronProfiler取一个NeuronConfig，并提供一个导出追寻回调。 召回书写了框架跟踪，track.json，与Neuron系统跟踪并列。
nc = NeuronConfig(
模式= [ProfileMode.DEVICE, ProfileMode.RUNTIME] (英语).
配置文件输出 dir=UTPUT DIR
抓取启用 for nc='0'
关联 device time= true
)
出口国 = Neuron Profiler(nc)
模式参数选择要捕获的数据流，例如设备指令和运行时间痕迹。 抓取启用 for nc 参数选择NeuronCores来跟踪。 设置关联 device time= True 将设备时间从微量传播到密钥平均()表，您可以在此比较CPU和Neuron设备时间。 关于完整的参数引用和可用的ProfileMode值，参见NeuronConfig Profile Modes文档。
工具：Neuron 探索器
Neuron Explorer是一个基于浏览器的浏览器，用于系统跟踪，设备时间表，以及每引擎崩溃。 你可以通过SSH隧道或Visual Studio代码扩展访问。 在此文章中，您使用 Explorer 来确认哪些密钥平均值( ) 以数字方式报告，并检查发送之间发生的情况。
剖析后，输出目录包含系统跟踪，当您获取设备数据时，匹配.neff和.ntff文件。 Neuron可执行文件格式(NEFF)包含编译的Neuron可执行文件，而匹配的NTFF包含其硬件指令跟踪。 文件共享一个数字散列以识别它们为对。 要在Neuron Explorer中打开配置，上传系统或组合配置的完整输出目录。 对于只使用设备配置文件，请单独上传匹配的.neff和.ntff文件。 参见步骤5：上传您对UI和命令行指令的描述文件。
配置文件 RMSNorm
对于这第一种分析，你描述RMSNorm，一个有明确的算术和还原操作序列的小型操作。
您描述同一操作的三项执行：
- Naive RMSNorm 带有单个 PyTorch 操作：乘法，表示法，加法，对等方根，和乘法。
- 火炬.rms norm(),PyTorch内置降格为单一引信的Neuron内核。
- Naive RMSNorm配有火炬。 编译，使用与第一次执行相同的操作。
脚本使用相同的输入，剖面配置，以及排程，因此结果可以直接比较。
复制此作品的所有脚本都可以在火炬-neuronx实例目录中找到。
简介1:Naive PyTorch操作
开始写 RMSNorm 与单个 PyTorch 操作。 每行映射到一个或多个ATen操作，并可以创建中间载荷。
设备 = 火炬.device('neuron')
x = 火炬.randn (4, 512, dtype=torch.bfloat16，设备=device).
重量 = 火炬.ones(512,dtype=torch.bfloat16，设备=device)
def rms norm fn(t): (法语).
x sq = t * t # 平方边
平均值 sq = x sq.mean(dim=-1, keepdim=True) # 正方形平均值
rms = 火炬.rsqrt(平均值+ 1e-6) # 1 / sqrt(平均值+sq + eps)
返回 t * rms * 重量 # 正常 + 比例
使用 Neuron 后端所需的两个参数调用配置器：
- 包括剖析器活动。 Private Useus1用于从Neuron设备中收集活动。
- 通过NeuronConfig作为实验配置，以启用Neuron特有的Kineto特性。
该表使用wait=2, warmup=1，以及活动=1横跨四个 prof.step() 调用。 这记录了一个温暖的迭代，但不包括初始设置的间接费用。
调用火炬。 neuron. 同步( ) 配置器循环。 没有这种同步，配置器可以在设备完成前导出并留下Neuron时间列空。
带火炬。 profile. profile(
活动=[Profiler Activity.CPU, Professor Activity.Private Uservation1] (中文(简体) ).
实验配置=nc
记录( S)
运行( trace ready=exporter.export trace)
调度=torch.profile.schedule(等待=2，温暖=1，活动=1)
作为prof:
范围(4):
带有火炬.no grad():
=rms norm fn(x) (中文(简体) ).
火炬。 neuron. 同步( )
步骤( )
剖析完成后，请检查密钥平均()表。 按自定义的私人使用1 时间排序，将操作放在设备成本最高的地方。 在此显示的表格被过滤到报告Neuron设备时间的 RMSNorm 操作。 完整的密钥平均()输出还包括CPU侧事件，如分配和主机到设备传输。
打印(prof.key perors (.). table(sort by="self privateuse1 time total，行 limit=4)).
键平均值表显示为3 aten::mul调用占Neuron设备时间的最大份额。
ATen操作运行的表显示，运行频率，以及使用CPU和Neuron设备的时间。 aten:::mul的总设备时间最高，因为函数执行三个乘法。
本国合计：2.355毫秒
每个行直接映射到函数。 aten::mul 运行于 x * x, x * rms，以及最终乘法按重量计算。 aten:：意为减小最后一个维度。 aten: add 新增 epsilon值并触发主机到设备复制，因为scalar开始于CPU. aten::rsqrt计算对等的方根。
Neuron Explorer遵循 aten::mul 操作从框架发送到运行时间调度到Neuron硬件车道的执行。
在探索器中读取路径
在Explorer中，时间从左到右流，事件生命周期从上到下流。 每个ATen操作都遵循相同的模式：在框架车道内创建，运行时车道内排程为nrta execut schedule，在硬件车道内设备执行为nc exec running.
系统跟踪视图将事件组织成车道。 在默认分组模式中，标签遵循模式实例/PID/子系统/识别器：
在急切模式下，每个操作独立运行。 使用依赖链查看器跟踪 aten:: mul. 从框架层的tensor通过运行时间调度到设备执行。
依赖链查看器突出显示一个操作在整个框架，运行时间和硬件车道的生命周期。
简介2：火炬.rms 规范
将单个操作替换为 PyTorch 的内置rms norm 函数。 Neuron将操作降低到一个引信内核：
def rms norm fn(t): (法语).
返回火炬.rms norm(t, [512]，重量)
引信的RMSNorm在密钥平均值表中作为一个设备操作出现。
本国总产量：403.3 μs
引信的RMSNorm遵循一个框架到硬件路径，而不是分别发送每次算术操作。
比较下移路径
在天真的执行中，每次行动都独立跨越调度边界，并遵循自己的框架-运行时间-硬件路径。 在引信执行中，一个Python呼叫覆盖了完整的计算。 追踪包含一个发送和一个硬件事件，组件操作之间没有间隙。
配置 3：有火炬的Naive PyTorch. 编译
最后的例子，用火炬编译原始的天真RMSNorm函数。 编译。
通过@torch.compile(后端="neuron")，火炬迪纳摩将Python字节代码作为FX图抓取，AOTAutograd将其降为ATen操作。 Neuron后端接收完整的图表，而不是单独发送每个操作。 它将操作整合，优化内存访问，并将图集成一个单一的NEFF文物。 文物是作为硬件发送的
@torch.compile(后端='neuron'，全图=True) 互联网档案馆的存檔，存档日期2011-09-02.
def 天真编译( t):
x_sq = t * t
平均值 sq = x sq.mean(dim=-1, keepdim=True)
rms = 火炬.rsqrt(平均值+ 1e-6)
返回 t * rms * 重量
编译的功能作为火炬编译区出现在关键平均值表中。
本国总产量：375.5 微克
编译的图表通过运行时间调度和硬件执行遵循一个路径。
表格显示 1设备执行，如Profile 2，但代码走的路径不同。 编译器没有调用目的建造的引信操作，而是从原来的Python函数中发现了聚变。 在本次运行中，Neuron时间与火炬.rms norm相当，为375 μs，而403 μs.
侧式摘要
双火炬.rms norm和火炬。 将引信 RMSNorm 编成一个内核调度。 在这些运行中，任一方法都占用了天真的执行装置时间的六分之一。 Fusion可以消除硬件执行之间的空隙，避免在单独发送的操作之间实现中间值如x sq和平均 sq.
结论
key 平均 () 显示操作会消耗设备时间，而其全输出也会是揭示 CPU 侧设置和传输。 探索者提供相应的视觉跟踪。 您可以跟踪一个操作，从框架发送到运行时间调度到硬件执行，并看到核聚变消除的漏洞。
在更大的模型上可以使用相同的工作流程：在表格中找到热点，在跟踪中验证，然后确定是引信操作还是火炬。 汇编可以降低成本。
在第二部分中，我们将进入设备级别配置。 您打开 NEFF，检查 TensorEngine, VectorEngine，和直接内存访问( DMA) 分解，并将内核放置在屋顶线上，以确定它是计算键还是内存键。
特别感谢Aritra Roy Gosthipaty的灵感开始这个系列，感谢我的同行Pinak Panigrahi、Samir Araujo、Akhil Azhikodan、Esha Lakhotia及其他评论！
本文由 AI图鉴自动抓取并翻译，版权归原作者所有；如需查看原文请点下方链接。
Profiling TorchNeuron (Part 1): System-level profiling on AWS Trainium.
Many ML practitioners have some intuition about how to make models run fast, but fewer know how to analyze why a model is slow. In Part 1 of this series, you will learn how to systematically analyze inference performance and identify where the time is spent.
This post starts from scratch and assumes no prior profiling experience. The walkthrough uses AWS Trainium, a purpose-built AI chip for deep learning training and inference workloads, and TorchNeuron, its open source native PyTorch backend. TorchNeuron supports standard PyTorch workflows on Trainium, including eager execution, torch.compile, and standard PyTorch APIs, with minimal code changes. Additionally, TorchNeuron integrates with Neuron Kernel Interface (NKI) to enable custom operators and performance optimization for Trainium hardware. By the end, you should know:
- How to configure torch.profiler with the Neuron backend and capture a clean trace.
- How to read the key_averages() table to identify which operations consume Neuron device time and compare their Self CPU time.
- How to follow one operation through Neuron Explorer, from framework dispatch through runtime scheduling to hardware execution on a NeuronCore.
- How these signals change across three implementations of the same operation.
Neuron as a PyTorch Backend
Before you profile a workload, it helps to understand how Neuron integrates with PyTorch.
The Neuron backend uses PyTorch's PrivateUse1 extension point. This mechanism lets an accelerator integrate with the framework without adding a dedicated dispatch key such as cuda. The torch_neuronx package works with upstream PyTorch to register Neuron as a native backend for PyTorch code. For details about this extension mechanism, see Facilitating New Backend Integration by PrivateUse1.
Set Up the Profiler
You use two objects to profile a Neuron workload:
- NeuronConfig controls which Neuron data streams and NeuronCores to capture and where to write the output. You pass it to torch.profiler.profile as experimental_config.
- NeuronProfiler takes a NeuronConfig and provides an export_trace callback. The callback writes the framework trace, trace.json, alongside the Neuron system traces.
nc = NeuronConfig(
modes=[ProfileMode.DEVICE, ProfileMode.RUNTIME],
profile_output_dir=OUTPUT_DIR,
capture_enabled_for_nc='0',
correlate_device_time=True,
)
exporter = NeuronProfiler(nc)
The modes parameter selects the data streams to capture, such as device instructions and runtime traces. The capture_enabled_for_nc parameter selects the NeuronCores to trace. Setting correlate_device_time=True propagates device time from the trace into the key_averages() table, where you can compare CPU and Neuron device time. For the complete parameter reference and available ProfileMode values, see the NeuronConfig Profile Modes documentation.
Tooling: Neuron Explorer
Neuron Explorer is a browser-based viewer for the system trace, device timeline, and per-engine breakdown. You can access it through an SSH tunnel or the Visual Studio Code extension. In this post, you use Explorer to confirm what key_averages() reports numerically and inspect what happens between dispatches.
After profiling, the output directory contains the system trace and, when you capture device data, matched .neff and .ntff files. The Neuron Executable File Format (NEFF) contains the compiled Neuron executable, while the matching NTFF contains its hardware instruction trace. The files share a numeric hash that identifies them as a pair. To open the profile in Neuron Explorer, upload the complete output directory for a system or combined profile. For a device-only profile, upload the matching .neff and .ntff files individually. See Step 5: Upload your profile for the UI and command-line instructions.
Profile RMSNorm
For this first analysis, you profile RMSNorm, a small operation with a clear sequence of arithmetic and reduction operations.
You profile three implementations of the same operation:
- Naive RMSNorm with individual PyTorch operations: multiply, mean, add, reciprocal square root, and multiply.
- torch.rms_norm(), the PyTorch built-in lowered to a single fused Neuron kernel.
- Naive RMSNorm with torch.compile, using the same operations as the first implementation.
The scripts use the same inputs, profiler configuration, and schedule so the results are directly comparable.
All scripts to reproduce this work are available in the torch-neuronx examples directory.
Profile 1: Naive PyTorch Operations
Start by writing RMSNorm with individual PyTorch operations. Each line maps to one or more ATen operations and can create intermediate tensors.
device = torch.device('neuron')
x = torch.randn(4, 512, dtype=torch.bfloat16, device=device)
weight = torch.ones(512, dtype=torch.bfloat16, device=device)
def rms_norm_fn(t):
x_sq = t * t # square
mean_sq = x_sq.mean(dim=-1, keepdim=True) # mean of squares
rms = torch.rsqrt(mean_sq + 1e-6) # 1 / sqrt(mean_sq + eps)
return t * rms * weight # normalize + scale
Call the profiler with two arguments required for the Neuron backend:
- Include ProfilerActivity.PrivateUse1 to collect activity from the Neuron device.
- Pass NeuronConfig as experimental_config to enable the Neuron-specific Kineto features.
The schedule uses wait=2, warmup=1, and active=1 across four prof.step() calls. This records one warm iteration without including initial setup overhead.
Call torch.neuron.synchronize() inside the profiler loop. Without this synchronization, the profiler can export before the device finishes and leave the Neuron time column empty.
with torch.profiler.profile(
activities=[ProfilerActivity.CPU, ProfilerActivity.PrivateUse1],
experimental_config=nc,
record_shapes=False,
on_trace_ready=exporter.export_trace,
schedule=torch.profiler.schedule(wait=2, warmup=1, active=1),
) as prof:
for step in range(4):
with torch.no_grad():
_ = rms_norm_fn(x)
torch.neuron.synchronize()
prof.step()
After profiling finishes, inspect the key_averages() table. Sort by self_privateuse1_time_total to place the operations with the highest device cost first. The table shown here is filtered to the RMSNorm operations that report Neuron device time. The full key_averages() output also includes CPU-side events such as allocations and host-to-device transfers.
print(prof.key_averages().table(sort_by="self_privateuse1_time_total", row_limit=4))
The key averages table shows that three aten::mul calls account for the largest share of Neuron device time.
The table shows which ATen operations ran, how often they ran, and how much time they spent on the CPU and Neuron device. aten::mul has the highest total device time because the function performs three multiplications.
Self NEURON total: 2.355 ms
Each row maps directly to the function. aten::mul runs for x * x, x * rms, and the final multiplication by weight. aten::mean reduces the last dimension. aten::add adds the epsilon value and triggers a host-to-device copy because the scalar starts on the CPU. aten::rsqrt computes the reciprocal square root.
Neuron Explorer follows an aten::mul operation from framework dispatch through runtime scheduling to execution on the Neuron hardware lanes.
Read the Trace in Explorer
In Explorer, time flows from left to right and the event lifecycle flows from top to bottom. Each ATen operation follows the same pattern: creation in the framework lane, scheduling in the runtime lane as nrta_execute_schedule, and device execution in the hardware lane as nc_exec_running.
The system trace view organizes events into lanes. In the Default grouping mode, labels follow the pattern instance/PID/subsystem/identifier:
In eager mode, each operation runs independently. Use the dependency chain viewer to follow aten::mul.Tensor from the framework layer through runtime scheduling to device execution.
The dependency chain viewer highlights the lifecycle of one operation across the framework, runtime, and hardware lanes.
Profile 2: torch.rms_norm
Replace the individual operations with PyTorch's built-in rms_norm function. Neuron lowers the operation to one fused kernel:
def rms_norm_fn(t):
return torch.rms_norm(t, [512], weight)
The fused RMSNorm appears as one device operation in the key averages table.
Self NEURON total: 403.3 μs
The fused RMSNorm follows one framework-to-hardware path instead of dispatching each arithmetic operation separately.
Compare the Lowering Path
In the naive implementation, each operation crosses the dispatch boundary independently and follows its own framework-to-runtime-to-hardware path. In the fused implementation, one Python call covers the complete computation. The trace contains one dispatch and one hardware event, with no gaps between component operations.
Profile 3: Naive PyTorch with torch.compile
For the final example, compile the original naive RMSNorm function with torch.compile.
With @torch.compile(backend="neuron"), TorchDynamo captures the Python bytecode as an FX graph and AOTAutograd lowers it to ATen operations. The Neuron backend receives the complete graph instead of dispatching each operation separately. It fuses the operations, optimizes memory access, and compiles the graph into a single NEFF artifact. The artifact runs as one hardware dispatch.
@torch.compile(backend='neuron', fullgraph=True)
def naive_compiled(t):
x_sq = t * t
mean_sq = x_sq.mean(dim=-1, keepdim=True)
rms = torch.rsqrt(mean_sq + 1e-6)
return t * rms * weight
The compiled function appears as one Torch-Compiled Region in the key averages table.
Self NEURON total: 375.5 μs
The compiled graph follows one path through runtime scheduling and hardware execution.
The table shows one device execution, as in Profile 2, but the code takes a different path. Instead of calling a purpose-built fused operation, the compiler discovers the fusion from the original Python function. In this run, the Neuron time is comparable to torch.rms_norm, at 375 μs compared with 403 μs.
Side-by-Side Summary
Both torch.rms_norm and torch.compile fuse RMSNorm into one kernel dispatch. In these runs, either approach takes about one-sixth of the naive implementation's device time. Fusion removes the gaps between hardware executions and avoids materializing intermediate values such as x_sq and mean_sq between separately dispatched operations.
Conclusion
key_averages() shows which operations consume device time, while its full output also reveals CPU-side setup and transfers. Explorer provides the corresponding visual trace. You can follow one operation from framework dispatch through runtime scheduling to hardware execution and see the gaps that fusion removes.
You can use the same workflow on larger models: find a hotspot in the table, verify it in the trace, and then determine whether a fused operation or torch.compile can reduce its cost.
In Part 2, we will move into the device-level profile. You open the NEFF, inspect the TensorEngine, VectorEngine, and direct memory access (DMA) breakdown, and place the kernel on the roofline to determine whether it is compute-bound or memory-bound.
Special thanks to Aritra Roy Gosthipaty for the inspiration to begin this series, and to my peers Pinak Panigrahi, Samir Araujo, Akhil Azhikodan, Esha Lakhotia and the rest for reviewing!
Auto-fetched and translated by AI Atlas. All rights belong to the original author.