海光 DCU 性能分析实战：用 hipprof 从 API、Kernel 到 PMC 定位 PyTorch 瓶颈

本文围绕海光 DCU 的真实性能分析流程展开。我们先用一个可控的 vector-add 程序区分 HIP API、设备 kernel、调用栈、PMC 和资源泄漏检测的统计口径，再把同一套方法用到一个保持 69 通道完整输出的气象 Transformer。所有命令均在BW / gfx936、DTK 25.04.2 环境实测；文章中的优化收益来自严格 A/B，不使用profiler 插桩后的时间充当正常运行性能。

模型在 DCU 上“能跑”以后，下一步最容易犯的错误是只盯着设备利用率。利用率不满，并不能直接说明矩阵乘法没有吃满；它也可能是大量短 kernel、数据搬运、内存分配、同步、CPU 发射或不连续张量物化共同造成的。

我在气象 Transformer 上真正遇到过的情况正是如此。直觉上，Transformer 的主角应该是 GEMM；实际 trace 中，单类 direct-copy kernel 就占了约三分之一。直到把kernel 统计、调用栈和真实 shape 放到一起，优化方向才从“继续砍矩阵计算”转向“减少布局转换、中间张量和碎片化 launch”。

这篇文章重点回答五个问题：

一、先建立正确的性能分析分层

在海光 DCU 上分析 PyTorch 程序，可以把证据分成五层：

这五层不能互相替代。

例如，hipLaunchKernel 的 API 时间主要是主机提交开销；kernel CSV 中的时间是设备执行时间。两张表的百分比分母不同，不能把 hipLaunchKernel 0.3% 和vector_add kernel 100% 放在一起相加。

同样，PMC 是带硬件计数器插桩的诊断模式。它能解释一个 kernel 为什么慢，却可能大幅改变 kernel 自己的执行时间。最终性能必须退出 profiler 后重新测量。

二、环境与本文实测边界

本文使用：

OS Ubuntu 22.04.5 LTS
DTK path /opt/dtk-25.04.2
hipprof DTK 25.04.2 bundled tool
PyTorch 2.5.1
HIP runtime 6.3.25405
device BW
architecture gfx936

进入容器后先检查工具：

set +u
source /opt/dtk-25.04.2/env.sh
set -u

command -v hy-smi
command -v hipprof
hipprof –help | head -n 40

DTK 手册中几个容易混淆的参数是：

• –stats：只打印并导出统计数据，不生成完整可视化 JSON；
• –kernel-stack：额外导出 kernel 的主机端调用栈 CSV；
• –pmc：采集硬件计数器；
• –pmc-type 3：把 PMC 输出为适合批量解析的 CSV；
• –kernel-name：只对指定 kernel 做 PMC；
• –memory-usage：启动资源泄漏检查，但相比 –leak-check 减少函数栈提取；
• –mask-trace：只收集配置文件中指定的 HIP API，降低大 trace 的数据量。

本文不会使用浏览器可视化页面。hipprof 生成的 CSV 和数据库已经足够完成主要统计、归类和 A/B 决策，而且更适合归档和脚本复现。

三、先准备一个可控的最小程序

最小实验使用第一篇文章中的 vector_add.hip。它包含：

• 三次 hipMalloc；
• 两次 H2D 和一次 D2H hipMemcpy；
• 20 次 warmup；
• 100 次正式 kernel；
• HIP event 同步计时；
• 最大绝对误差检查；
• 资源销毁与显存释放。

编译时先确认当前 DCU 架构：

rocminfo | grep -E 'Name:.*gfx'

本文实机是 gfx936：

hipcc -O3 -std=c++17 \\
–offload-arch=gfx936 \\
vector_add.hip -o vector_add

正常运行：

device=BW arch=gfx936:sramecc+:xnack-
elements=1048576 mean_kernel_ms=0.0185488 max_abs_error=0

这里的 0.0185488 ms 是退出 profiler 后，由程序自己的 HIP event 得到的基线。后面所有插桩模式都要和它分开记录。

四、把不同 hipprof 模式拆成独立进程

不要一次叠上 –hip-trace –kernel-stack –pmc –memory-usage。不同模式的目标和扰动程度不同，叠加后既难解释，也可能制造新的瓶颈。

下面脚本把正常运行、stats、kernel-stack 和 PMC 放到四个目录。完整文件名为run_vector_add_profile.sh：

#!/usr/bin/env bash

# Run separate hipprof passes so trace, stack and PMC instrumentation do not
# contaminate one another. Each pass writes into its own output directory.
set -euo pipefail

APP="$(realpath "${1:-./vector_add}")"
OUT="${2:-hipprof_vector_add_$(date +%Y%m%d_%H%M%S)}"
mkdir -p "$OUT"/{normal,stats,stack,pmc}

set +u
if [[ -f /opt/dtk-25.04.2/env.sh ]]; then
source /opt/dtk-25.04.2/env.sh
elif [[ -f /opt/dtk/env.sh ]]; then
source /opt/dtk/env.sh
fi
set -u

"$APP" | tee "$OUT/normal/program.txt"

(
cd "$OUT/stats"
hipprof –hip-trace –stats "$APP" | tee hipprof.txt
)

(
cd "$OUT/stack"
hipprof –kernel-stack "$APP" | tee hipprof.txt
)

(
cd "$OUT/pmc"
hipprof –pmc –pmc-type 3 –kernel-name vector_add "$APP" | tee hipprof.txt
)

printf 'profile output: %s\\n' "$(realpath "$OUT")"

运行：

chmod +x run_vector_add_profile.sh
./run_vector_add_profile.sh ./vector_add ./profile_run

输出目录大致为：

profile_run/
├── normal/program.txt
├── stats/result_<pid>.hiptrace.csv
├── stats/result_<pid>.kernel.csv
├── stack/result_<pid>.stack.csv
└── pmc/pmc_results_<pid>.csv

这种目录隔离有三个好处：

五、第一步：用 –hip-trace –stats 分开看 API 和 kernel

命令：

hipprof –hip-trace –stats ./vector_add

本次独立冷进程得到的 HIP API 前几项：

name calls total_ms average_us percent
hipMalloc 3 240.481148 80160.382 97.846
hipMemcpy 3 2.432756 810.918 0.990
hipEventSynchronize 1 1.496596 1496.596 0.609
hipLaunchKernel 120 0.830707 6.922 0.338
hipDeviceSynchronize 1 0.377584 377.584 0.154

设备 kernel 表则只有：

kernel calls total_ms average_us percent
vector_add 120 2.255999 18.799 100.000

这两张表揭示了三个事实。

1. 冷进程的 API 百分比不等于稳态热点

hipMalloc 占 API 记录时间的 97.846%，但程序只在启动时分配三次。如果业务进程启动一次、随后推理几千次，直接为这三次分配重写 allocator，未必能改善稳定期吞吐。

相反，如果业务每个样本都在循环中分配和释放大张量，分配才可能是稳定期问题。因此要同时看调用次数、代码位置和业务生命周期，不能只按百分比排序。

2. API launch 时间和设备执行时间不是同一个指标

本次 hipLaunchKernel 主机 API 平均约 6.922 us，设备 kernel 平均约18.799 us。前者是提交，后者是设备执行。异步流水中两者还可能重叠，不能简单相加成“每个 kernel 25.721 us”。

3. –stats 很适合第一轮筛选

它不生成完整 JSON，数据量小，能够先回答：

• API 是否被分配、拷贝或同步主导；
• kernel 名称、调用次数和设备时间排名；
• 是否值得继续做 stack 或 PMC。

如果第一轮就发现某个 kernel 总占比不到 0.2%，通常不值得立即写专用 kernel。

六、第二步：用 –kernel-stack 找到上层调用者

命令：

hipprof –kernel-stack ./vector_add

vector-add 的 120 次 launch 都能回到：

kernel launches user frame
vector_add(float const*, float const*, float*, int) 120 vector_add()

真实 PyTorch 模型会更复杂。调用栈常常只能回到：

python()
libtorch_hip.so
librocblas.so

这是正常现象，不代表 –kernel-stack 没用。它至少能区分自定义扩展、rocBLAS、PyTorch elementwise 和 Runtime 路径。但如果想精确定位到模型中的某个 block，建议：

长 kernel 名称本身往往只能告诉你“它是什么”，调用栈和 range 才能告诉你“是谁触发了它”。

七、第三步：只对少数热点运行 PMC

PMC 使用有限的硬件计数器资源，不可能一次拿到所有指标。DTK 手册把指标分组，–pmc-type 3 可以导出 CSV，适合脚本解析。

本文只过滤 vector_add：

hipprof –pmc –pmc-type 3 \\
–kernel-name vector_add \\
./vector_add

解析得到的稳定资源字段：

records 120
grid_size 1048576
workgroup_size 256
wave_size 64
arch_vgpr 8
accum_vgpr 0
sgpr 16
lds_bytes 0
scratch_bytes 0
lds_bank_conflict 0

这些字段可以这样理解：

• workgroup_size=256：每个 workgroup 256 个线程；
• wave_size=64：该 kernel 的本次记录显示以 wave64 执行；
• arch_vgpr=8、sgpr=16：这个简单 kernel 的寄存器压力较低；
• lds=0：没有使用片上共享存储；
• scratch=0：没有观察到寄存器溢出到 scratch；
• LDS bank conflict=0：没有 LDS，自然也没有 LDS bank 冲突。

不能仅凭这些字段宣称 kernel 已经“100% 吃满硬件”。要判断计算、L1/L2、访存停顿或 LDS 行为，还需要选择相应 PMC 分组，并结合业务目标分析。

八、最重要的警告：PMC 会改变被测程序

同一个可执行程序在四种模式下打印：

mode program mean kernel
normal 0.0185488 ms
hip-trace –stats 0.0187840 ms
kernel-stack 0.0189324 ms
PMC type 3 1.4807800 ms

相对正常运行：

• stats 约慢 1.27%；
• kernel-stack 约慢 2.07%；
• PMC 约慢 79.83x。

PMC 采集可能通过重复或重放 kernel 读取不同计数器，程序自己的 HIP event 也会感受到这类扰动。因此：

trace、stack 和 PMC 用于归因；退出 profiler 后的同步墙钟或 HIP event 才用于最终性能结论。

如果把 1.48078 ms 写成 vector-add 的“真实性能”，误差会接近两个数量级。

九、–memory-usage 不是峰值显存监控

命令：

hipprof –memory-usage ./vector_add

本次进程结束时打印：

host leakcheck memory size 170.06 MB
device0 leakcheck memory size 92.00 KB

vector-add 已显式释放三个设备指针，但 Runtime、profiler 和进程级资源仍可能出现在记录中。–memory-usage 的定位是资源泄漏检查，而不是业务稳定期的显存平均值，也不是进程生命周期峰值。

如果业务关注模型显存，应先明确口径：

• PyTorch 活动张量：memory_allocated()；
• caching allocator 保留：memory_reserved()；
• 进程峰值：外部监控或 profiler 峰值；
• 稳定期平均占用：固定周期外部采样；
• 资源未释放：–leak-check 或 –memory-usage。

不要看到 92 KB 就直接宣布用户代码存在 92 KB 显存泄漏。需要结合分配/释放记录、调用栈以及新进程复测再判断。

十、用脚本解析 CSV，而不是手工数几百个 kernel

下面的 summarize_hipprof.py 可以同时解析 API、kernel、stack 和 PMC CSV，并按名称模式聚合 PyTorch kernel。它不会依赖浏览器，也不会把私有绝对路径打印到摘要。

"""Summarize hipprof CSV exports without loading the browser UI."""

from __future__ import annotations

import argparse
import csv
from collections import defaultdict
from pathlib import Path
from typing import Iterable

CATEGORY_RULES = (
("vector_add", ("vector_add",)),
("direct_copy", ("direct_copy_kernel",)),
("softmax", ("softmax",)),
("add", ("cudafunctor_add", "addfunctor", "binary_internal::add")),
("layernorm", ("layer_norm", "layernorm")),
("gelu", ("gelu",)),
("roll", ("roll_cuda",)),
("gemm", ("cijk_", "gemm", "addmm")),
("convolution", ("conv3d", "convolution", "col_2_image")),
("cat_or_index", ("catarray", "index_elementwise", "index_kernel")),
)

def read_csv(path: Path) –> list[dict[str, str]]:
with path.open("r", encoding="utf-8-sig", newline="") as handle:
return list(csv.DictReader(handle))

def as_int(row: dict[str, str], key: str) –> int:
value = row.get(key, "").strip()
return int(float(value)) if value else 0

def as_float(row: dict[str, str], key: str) –> float:
value = row.get(key, "").strip()
return float(value) if value else 0.0

def classify_kernel(name: str) –> str:
lowered = name.lower()
for category, patterns in CATEGORY_RULES:
if any(pattern in lowered for pattern in patterns):
return category
return "other"

def print_api_summary(path: Path, top: int) –> None:
rows = [row for row in read_csv(path) if row.get("Name") != "Total"]
rows.sort(key=lambda row: as_int(row, "TotalDurationNs"), reverse=True)
print("\\n[HIP API summary]")
print("name\\tcalls\\ttotal_ms\\taverage_us\\tpercent")
for row in rows[:top]:
print(
f"{row['Name']}\\t{as_int(row, 'Calls')}\\t"
f"{as_int(row, 'TotalDurationNs') / 1e6:.6f}\\t"
f"{as_int(row, 'AverageNs') / 1e3:.3f}\\t"
f"{as_float(row, 'Percentage'):.3f}"
)

def print_kernel_summary(path: Path, top: int) –> None:
rows = [row for row in read_csv(path) if row.get("Name") != "Total"]
rows.sort(key=lambda row: as_int(row, "TotalDurationNs"), reverse=True)
total_ns = sum(as_int(row, "TotalDurationNs") for row in rows)

print("\\n[Kernel summary]")
if total_ns <= 0:
print("no non-zero duration rows")
return
print("name\\tcalls\\ttotal_ms\\taverage_us\\tpercent")
for row in rows[:top]:
display_name = row["Name"]
if len(display_name) > 100:
display_name = display_name[:97] + "…"
print(
f"{display_name}\\t{as_int(row, 'Calls')}\\t"
f"{as_int(row, 'TotalDurationNs') / 1e6:.6f}\\t"
f"{as_int(row, 'AverageNs') / 1e3:.3f}\\t"
f"{as_int(row, 'TotalDurationNs') / total_ns * 100:.3f}"
)

category_ns: dict[str, int] = defaultdict(int)
category_calls: dict[str, int] = defaultdict(int)
for row in rows:
category = classify_kernel(row["Name"])
category_ns[category] += as_int(row, "TotalDurationNs")
category_calls[category] += as_int(row, "Calls")

print("\\n[Pattern-based kernel categories]")
print("category\\tcalls\\ttotal_ms\\tpercent")
for category, duration_ns in sorted(
category_ns.items(), key=lambda item: item[1], reverse=True
):
print(
f"{category}\\t{category_calls[category]}\\t"
f"{duration_ns / 1e6:.6f}\\t{duration_ns / total_ns * 100:.3f}"
)

def sanitize_frame(frame: str) –> str:
frame = frame.strip()
if not frame:
return "unknown"
if " [" in frame:
frame = frame.split(" [", 1)[0]
if "/" in frame:
prefix, separator, suffix = frame.rpartition("/")
if separator:
frame = suffix
return frame

def is_runtime_frame(frame: str) –> bool:
lowered = frame.lower()
return lowered.startswith("python(") or any(
marker in lowered
for marker in (
"libgalaxy",
"libamdhip",
"libc.so",
"libtorch",
"libpython",
"python3",
)
)

def print_stack_summary(path: Path, top: int) –> None:
rows = read_csv(path)
grouped: dict[tuple[str, str], list[tuple[int, str]]] = defaultdict(list)
for row in rows:
key = (row.get("str_name", "unknown"), row.get("index", "unknown"))
grouped[key].append((as_int(row, "stack_sq"), row.get("addrname", "")))

kernel_launches: dict[str, int] = defaultdict(int)
kernel_origins: dict[str, dict[str, int]] = defaultdict(lambda: defaultdict(int))
for (kernel_name, _), frames in grouped.items():
kernel_launches[kernel_name] += 1
origin = "runtime-only"
for _, frame in sorted(frames):
cleaned = sanitize_frame(frame)
if cleaned and not is_runtime_frame(cleaned):
origin = cleaned
break
kernel_origins[kernel_name][origin] += 1

print("\\n[Kernel stack summary]")
print("kernel\\tlaunches\\tmost_common_user_frame")
for kernel_name, launches in sorted(
kernel_launches.items(), key=lambda item: item[1], reverse=True
)[:top]:
origins = kernel_origins[kernel_name]
origin = max(origins.items(), key=lambda item: item[1])[0]
display_name = kernel_name if len(kernel_name) <= 90 else kernel_name[:87] + "…"
print(f"{display_name}\\t{launches}\\t{origin}")

def unique_values(rows: Iterable[dict[str, str]], key: str) –> str:
values = sorted({row.get(key, "").strip() for row in rows if row.get(key, "").strip()})
return ",".join(values) if values else "n/a"

def print_pmc_summary(path: Path) –> None:
rows = read_csv(path)
if not rows:
print("\\n[PMC summary]\\nno rows")
return

print("\\n[PMC resource summary]")
print(f"records\\t{len(rows)}")
print(f"kernel\\t{unique_values(rows, 'KernelName')}")
for label, key in (
("grid_size", "grd"),
("workgroup_size", "wgr"),
("wave_size", "wave_size"),
("arch_vgpr", "arch_vgpr"),
("accum_vgpr", "accum_vgpr"),
("sgpr", "sgpr"),
("lds_bytes", "lds"),
("scratch_bytes", "scr"),
("lds_bank_conflict", "SQ_LDS_BANK_CONFLICT"),
):
print(f"{label}\\t{unique_values(rows, key)}")

def main() –> None:
parser = argparse.ArgumentParser()
parser.add_argument("–hiptrace", type=Path)
parser.add_argument("–kernel", type=Path)
parser.add_argument("–stack", type=Path)
parser.add_argument("–pmc", type=Path)
parser.add_argument("–top", type=int, default=10)
args = parser.parse_args()

if not any((args.hiptrace, args.kernel, args.stack, args.pmc)):
parser.error("provide at least one hipprof CSV path")
if args.hiptrace:
print_api_summary(args.hiptrace, args.top)
if args.kernel:
print_kernel_summary(args.kernel, args.top)
if args.stack:
print_stack_summary(args.stack, args.top)
if args.pmc:
print_pmc_summary(args.pmc)

if __name__ == "__main__":
main()

vector-add 的解析命令：

python summarize_hipprof.py \\
–hiptrace profile_run/stats/result_<pid>.hiptrace.csv \\
–kernel profile_run/stats/result_<pid>.kernel.csv \\
–stack profile_run/stack/result_<pid>.stack.csv \\
–pmc profile_run/pmc/pmc_results_<pid>.csv

注意 <pid> 是占位符，实际运行时应替换成当前目录中的文件名；不要把带尖括号的文本直接交给 shell。

模式归类是启发式的。PyTorch 或 DTK 版本变化后，kernel 名称可能改变，发布报告前应抽查每个分类的代表行，避免把名称相似但语义不同的 kernel 合并。

十一、真实案例：69 通道气象 Transformer 的热点是什么

我们对一个保持 69 通道、原空间分辨率和完整预测路径的气象 Transformer 分别运行了 stats 和 kernel-stack。设备 kernel CSV 共记录：

kernel calls 1869
kernel time 825.152715 ms

按名称模式聚合后：

category calls total_ms percent
direct_copy 843 266.933896 32.350
gemm 316 181.074114 21.944
add 192 103.964575 12.599
softmax 56 90.451776 10.962
layernorm 96 77.525224 9.395
gelu 44 35.815013 4.340
roll 120 31.863178 3.861
other 134 19.596949 2.375
cat_or_index 60 10.896634 1.321
convolution 8 7.031356 0.852

这里的 825.152715 ms 是捕获窗口内所有设备 kernel duration 的累计值，用于比较同一CSV 内各类别的相对占比；它不是无插桩环境下的单样本端到端墙钟时间，也不应和后文的整模稳态时间直接比较。

这里有两个看似矛盾、其实都成立的结论：

如果只优化 GEMM，即使把这部分无限加速，Amdahl 上限也受剩余约 78% 时间约束。因此后续重点转向：

• QKV 布局转换和切片物化；
• bias 与激活中间张量；
• LayerNorm 前后的非连续 rearrange；
• roll、copy、cat 和 index 等碎片化路径。

为什么不能只看 HIP API 表

同一全模型进程的 HIP API 表中，hipModuleLoadData、hipHostMalloc、hipMalloc和 hipModuleGetFunction 占比较高，因为 trace 包含模型加载和 Runtime 初始化。这些数字适合分析冷启动，不应直接替代稳态 kernel 分类。

十二、从热点归因到真正有效的融合

最终通过数值门和整模 A/B 的方案包含两部分。

1. QKV 后处理融合

原路径把 QKV projection 组织为连续的：

[B, P, T, 3, H, D]

其中 B 是 batch，P 是窗口组，T 是窗口内 token 数，H 是 head 数，D 是单个 head 的维度。

随后 permute 到：

[3, B, H, P, T, D]

Q、K、V 切片变成非连续视图；在这条真实路径中，后续计算触发了额外的布局物化与复制。融合 kernel 读取一次交错 QKV，直接写出 packed contiguous 布局，并融合Q scaling。

2. fc1 bias + exact GELU 融合

原路径：

no-bias GEMM -> BF16 bias add -> exact GELU -> output

融合路径把 bias 物化和 exact erf GELU 合并到一个随源码编译的 HIP kernel 中，避免一个完整中间张量。

最终外部 2 ms 采样的 A-B-B-A：

metric reference fused delta
sampled mean VRAM 2789.217304 MiB 2770.666960 MiB -18.550344 MiB
steady mean time 0.188963599 s 0.177391734 s -6.1239%
完整输出代理指标 33.900791087 33.894619270 -0.006171817

这组结果说明：

• 时间和采样显存同时下降，不是使用持久大 workspace 换速度；
• 完整输出代理指标只变化 -0.00617；
• 微基准收益能够落到整模，但整模提升远小于某个单 kernel 的加速倍数。

这里的完整输出代理指标是项目内部的同权重 A/B 数值门，只用于判断候选是否改变完整模型输出，不是跨模型或跨硬件的通用 benchmark。

并不是所有相似融合都安全。另一个 fc2 residual+bias kernel 只节省约 0.47 ms，却让完整输出代理指标从约 33.90 降到 33.44，因此直接停止。profiler 只能告诉你“哪里值得看”，不能替代数值与业务正确性门槛。

十三、一套可以复用的 DCU 性能优化流程

步骤 1：正常环境建立基线

记录：

• warmup；
• 重复次数；
• 同步位置；
• 输入 shape 和 dtype；
• 端到端时间；
• 业务数值指标；
• 显存统计口径。

步骤 2：stats 找 API 和 kernel 热点

先用：

hipprof –hip-trace –stats <command>

只选总占比足够高、调用次数和业务位置合理的候选。

步骤 3：stack 或 range 找到调用路径

hipprof –kernel-stack <command>

调用栈太浅时，缩小模型，或加入 ROCTX range 并使用 –hiptx-trace。

步骤 4：真实 shape 微基准

对 reference 和 candidate 同时做：

• 数值差异；
• 20–30 次 warmup；
• 100 次左右重复；
• 显式同步；
• 多轮交错 A/B。

步骤 5：只对少数 kernel 跑 PMC

hipprof –pmc –pmc-type 3 \\
–kernel-name <hot-kernel> \\
<microbenchmark>

PMC 用于解释资源行为，不用于最终时间。

步骤 6：回到完整模型做 A-B-B-A

保持相同权重、输入、进程条件和采样方式。时间、显存和完整输出必须一起通过。

步骤 7：冷启动和发布检查

source-only 扩展还要验证：

• 新容器第一次编译；
• 构建产物是否只写入临时目录；
• 缓存是否按源码、Runtime 和架构失效；
• 退出 profiler 后的发布路径是否仍然使用相同代码。

十四、常见误区对照表

十五、发布前验收清单

• 正常、stats、stack、PMC 分进程运行；
• API 与设备 kernel 的时间口径没有混用；
• 记录 kernel 调用次数，而不只看百分比；
• 对长 kernel 名称做抽样核对后再归类；
• stack 或 range 已经定位到上层路径；
• PMC 只过滤少数热点 kernel；
• PMC 插桩时间没有写成正式性能；
• 微基准使用真实 shape、dtype 和同步位置；
• candidate 先通过数值门，再测性能；
• 完整模型使用交错 A/B，时间、显存和输出一起验证；
• 显存报告明确峰值、活动、保留或采样平均口径；
• 报告中的绝对时间标注具体 DCU、DTK 和软件版本；
• CSV、脚本和原始结论都有可审计归档。

结语

hipprof 最有价值的地方不是生成一张很复杂的时间线，而是把“感觉模型慢”拆成可以验证的问题：主机在分配还是发射，设备在 GEMM 还是 copy，某个 kernel 是计算、缓存、LDS 还是寄存器受限，以及一个微优化最终能不能落到整模。

在本文的气象 Transformer 中，真正改变方向的是 direct-copy 占比，而不是某个漂亮的峰值算力数字。最终有效的方案也不是继续删除模型结构，而是沿着 profiler 证据，融合 QKV 布局和 fc1 bias/GELU 数据流，并用完整输出和采样显存守住边界。

可以复用的核心顺序只有一句话：

正常基线 -> stats 找热点 -> stack 找来源 -> 真实 shape 微基准
-> PMC 解释行为 -> 整模 A-B-B-A -> 冷启动发布验证

工具负责提供证据，数值和端到端结果负责做最终决定。

参考资料

• 光合开发者社区工具下载：https://developer.hpccube.com/tool/
• DTK 25.04.2《hipprof 使用手册》
• DTK 25.04.2《Rocprofiler API 开发手册》
• PyTorch HIP semantics：https://docs.pytorch.org/docs/stable/notes/hip.html
• PyTorch Profiler：https://docs.pytorch.org/docs/stable/profiler.html