Moore Perf Compute 最佳实践

场景一：首次使用 - 性能分析初体验

问题： "我刚写了一个 MUSA kernel，想知道性能表现"

解决步骤：

# 1. 安装确认
which mcu
mcu --version

# 2. 简单分析
mcu ./my_program

# 3. 查看结果
mcu-ui
# File > Open File > report1.mcu-rep

预期结果： 生成 report1.mcu-rep，可在 GUI 中查看

场景二：只想分析特定的 Kernel

问题： "程序有上百个 kernel，只想分析其中一个"

解决步骤：

# 精确匹配
mcu -k matmul ./my_program

# 正则匹配
mcu -k regex:"matmul|relu" ./my_program

预期结果： 只分析指定 kernel，时间缩短，报告更小

场景三：多 GPU 环境分析

问题： "4 卡服务器，只想分析第 2 和第 4 张卡"

解决步骤：

mcu --devices 2,4 ./my_program

场景四：MPI 多进程应用分析

问题： "MPI 程序，多进程多节点，能分析吗？"

解决步骤：

# 单节点多 rank
mcu mpirun -np 4 ./my_mpi_program

# 多节点
mpirun -np 16 mcu ./my_mpi_program

场景五：分析结果太慢/开销太大

问题： "运行 mcu 后程序变慢，怎么办？"

解决步骤：

# 1. 减少分析的 kernel 数量
mcu -k my_important_kernel ./my_program

# 2. 降低采样频率
mcu --sampling-interval 200 ./my_program

# 3. 使用 normal 模式
mcu --dump-mode normal ./my_program

场景六：Roofline 分析

问题： "Roofline 图看不懂，该怎么分析？"

判断方法：

区域	诊断	优化方向
蓝色区域 (内存受限)	性能受内存带宽限制	优化内存访问
绿色区域 (计算受限)	性能受计算能力限制	优化计算逻辑

优化策略：

内存受限： 合并访问、共享内存缓存、减少传输
计算受限： 使用低精度、Tensor Core、操作融合

场景七：分析 Green Context 应用

问题： "使用了 Green Contexts，能分析吗？"

回答： 完全支持

识别方法：

标题栏 Attributes 按钮变绿色
Session 页面显示 Green Context Resources

场景八：报告打不开

问题： "报告在其他机器上打不开"

解决步骤：

使用英文路径（避免中文）
检查文件完整性

场景九：性能数据不可重现

问题： "每次运行数据都不一样，正常吗？"

原因： 系统负载变化、GPU 温度、后台服务

解决方案：

# 1. 关闭其他 GPU 程序
# 注意：将 other_gpu_app 替换为实际的进程名
pkill -f other_gpu_app

# 2. 多次运行取平均值
for i in {1..5}; do
    mcu -o run_$i.mcu-rep ./my_app
done

性能优化工作流

系统化优化流程

┌─────────────────────────────────────────────────────────┐
│ 第 1 步：建立基线                                        │
│ mcu -o baseline.mcu-rep ./app                           │
└─────────────────────────────────────────────────────────┘
                          ↓
┌─────────────────────────────────────────────────────────┐
│ 第 2 步：识别瓶颈                                        │
│ 查看 Roofline 位置、关键指标                             │
└─────────────────────────────────────────────────────────┘
                          ↓
┌─────────────────────────────────────────────────────────┐
│ 第 3 步：实施优化                                        │
│ 根据瓶颈类型选择优化策略                                 │
└─────────────────────────────────────────────────────────┘
                          ↓
┌─────────────────────────────────────────────────────────┐
│ 第 4 步：验证改进                                        │
│ mcu -o optimized.mcu-rep ./app                          │
│ 对比基线数据                                             │
└─────────────────────────────────────────────────────────┘

80/20 原则

聚焦关键内核：

分析整个应用：mcu -o full.mcu-rep ./app
识别占用 80% 执行时间的前 3-5 个 kernel
优先优化这些关键 kernel
获得最大性能提升

优化模式

模式：内存合并访问

适用场景： 内存受限 kernel

优化前：

for (int i = threadIdx.x; i < N; i += blockDim.x) {
    output[i] = input[i * stride];  // 未合并
}

优化后：

int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
    output[idx] = input[idx];  // 合并访问
}

模式：共享内存缓存

适用场景： 数据重用场景

优化前：

for (int k = 0; k < K; k++) {
    float a = A[row * K + k];
    float b = B[k * N + col];
    sum += a * b;
}

优化后：

__shared__ float As[TILE_SIZE][TILE_SIZE];
__shared__ float Bs[TILE_SIZE][TILE_SIZE];

// 加载到共享内存
// 从共享内存计算

模式：共享内存 Bank Conflict 优化

适用场景： kernel 使用共享内存分块、转置或重排数据，怀疑共享内存列访问导致性能下降

典型现象：

已经使用共享内存，但性能提升不明显
Memory Chart 中共享内存路径明显活跃
Memory Tables 的 Shared Memory 表中 Bank Conflicts 偏高

示例说明：

下面的矩阵转置样例中，优化前后算法逻辑保持不变，差别只在共享内存 tile 是否添加 padding：

优化前：

__shared__ float tile[32][32];

for (int j = 0; j < 32; j += 8) {
    tile[threadIdx.y + j][threadIdx.x] = in[(y + j) * width + x];
}

__syncthreads();

for (int j = 0; j < 32; j += 8) {
    out[(y + j) * height + x] = tile[threadIdx.x][threadIdx.y + j];
}

优化后：

__shared__ float tile[32][33];

for (int j = 0; j < 32; j += 8) {
    tile[threadIdx.y + j][threadIdx.x] = in[(y + j) * width + x];
}

__syncthreads();

for (int j = 0; j < 32; j += 8) {
    out[(y + j) * height + x] = tile[threadIdx.x][threadIdx.y + j];
}

为每行额外增加 1 个元素后，能够打散共享内存 bank 映射，从而减少按列读取时的冲突。

采集方式：

命令行：同时采集 MemoryWorkloadAnalysis、MemoryWorkloadAnalysis_Chart 和 MemoryWorkloadAnalysis_Tables

mcu --sections MemoryWorkloadAnalysis,MemoryWorkloadAnalysis_Chart,MemoryWorkloadAnalysis_Tables ./bank_conflict_transpose_musa

mcu --sections MemoryWorkloadAnalysis,MemoryWorkloadAnalysis_Chart,MemoryWorkloadAnalysis_Tables ./no_bank_conflict_transpose_musa

GUI：在 Connect to process 窗口的 Metrics 页面中，同时勾选 Memory Workload Analysis、Memory Workload Analysis Chart 和 Memory Workload Analysis Tables

在报告中如何定位：

打开报告的 Details 页面，进入 Memory Workload Analysis 分区。
先查看 Memory Chart，确认共享内存路径是否明显活跃。
再查看 Memory Tables，重点关注 Shared Memory 表中的 Bank Conflicts。
对比优化前后的两个报告，确认 Bank Conflicts 是否明显下降，再结合 Memory Throughput、执行时间等结果判断优化是否有效。

阅读建议

如果同时采集了 MemoryWorkloadAnalysis、MemoryWorkloadAnalysis_Chart 和 MemoryWorkloadAnalysis_Tables，可以在 Memory Workload Analysis 分区右上角的下拉框中切换 全部（All）、内存图表（Memory Chart） 和 内存表格（Memory Tables） 视图。

优化建议：

对 2D 共享内存 tile 的列访问，优先检查是否可以通过 padding 打散 bank 映射，例如 32 x 32 改为 32 x 33
如果 Bank Conflicts 下降但整体性能改善有限，还需要继续结合 L1 Hit Rate、L2 Hit Rate、设备内存 Throughput 等指标排查其他瓶颈
共享内存路径活跃并不一定单独说明存在 Bank Conflict，仍需结合 Shared Memory 表中的 Bank Conflicts 一起判断

相关文档：

性能分析：查看 Memory Workload Analysis、Memory Chart 和 Memory Tables 的结果解读方法
用户指南：查看 section 的启用方式，以及 Details 页面中的展示位置
MUSA SDK Programming Guide：查看共享内存 Bank Conflict 的原理，以及 padding、swizzle 等优化方法

模式：占用率优化

适用场景： 低占用率 kernel

优化方法：

调整 block size（128/256/512）
减少寄存器使用
优化共享内存使用

模式：操作融合

适用场景： 计算受限 kernel

优化前：

float temp = a * b;
float result = temp + c;

优化后：

float result = fma(a, b, c);  // Fused Multiply-Add

模式：低精度计算

适用场景： 精度允许场景

优化方法：

FP32 → FP16
FP16 → INT8
使用 Tensor Core

案例研究

案例：GEMM Kernel 优化

初始状态：

计算利用率：32%
内存利用率：71%
Roofline: 内存受限区域
性能：1.8 TFLOPS (28% 峰值)

优化措施：

实现共享内存分块 (128×128)
改进全局加载的内存合并
减少寄存器压力 (96 → 64)

优化结果：

计算利用率：58% (+26%)
内存利用率：82% (+11%)
性能：3.4 TFLOPS (+89%)

案例：卷积 Kernel 优化

初始状态：

占用率：25%
Block 大小：100 (不是 32 的倍数!)
性能：0.4 TFLOPS (15% 峰值)

优化措施：

Block 大小改为 128 (4 个线程束)
减少寄存器到 48/线程
重构内存访问模式

优化结果：

占用率：62% (+37%)
性能：1.1 TFLOPS (+175%)

案例：多内核流水线

问题： 15 个内核，不清楚优化哪个

方法：

分析整个应用：mcu -o full.mcu-rep ./inference_engine
发现 3 个内核占 80% 时间
聚焦优化前 3 个内核

结果：

优化前：50 ms 总推理时间
优化后：32 ms 总推理时间
改进：36% 端到端

经验： 聚焦关键的少数，而非琐碎的多数

团队协作

基线管理

perf-reports/
├── baseline/              # 基线报告
├── optimization/          # 优化迭代
├── regression/            # 回归测试
└── production/            # 生产监控

性能审查清单

- [ ]  基线已建立
- [ ]  使用代表性工作负载
- [ ]  已测量性能改进
- [ ]  无性能退化
- [ ]  报告已分享

常见问题速查表

问题	快速答案
多 GPU 只想分析一个	`--devices 0`
想看所有 kernel 名	GUI 的 Result 下拉框
分析太慢	`-k <kernel>` 减少范围
报告打不开	检查版本一致性
数据不可重现	多次运行取平均值

场景一：首次使用 - 性能分析初体验​

场景二：只想分析特定的 Kernel​

场景三：多 GPU 环境分析​

场景四：MPI 多进程应用分析​

场景五：分析结果太慢/开销太大​

场景六：Roofline 分析​

场景七：分析 Green Context 应用​

场景八：报告打不开​

场景九：性能数据不可重现​

性能优化工作流​

系统化优化流程​

80/20 原则​

优化模式​

模式：内存合并访问​

模式：共享内存缓存​

模式：共享内存 Bank Conflict 优化​

模式：占用率优化​

模式：操作融合​

模式：低精度计算​

案例研究​

案例：GEMM Kernel 优化​

案例：卷积 Kernel 优化​

案例：多内核流水线​

团队协作​

基线管理​

性能审查清单​

常见问题速查表​

这篇文章对您有帮助吗？

场景一：首次使用 - 性能分析初体验

场景二：只想分析特定的 Kernel

场景三：多 GPU 环境分析

场景四：MPI 多进程应用分析

场景五：分析结果太慢/开销太大

场景六：Roofline 分析

场景七：分析 Green Context 应用

场景八：报告打不开

场景九：性能数据不可重现

性能优化工作流

系统化优化流程

80/20 原则

优化模式

模式：内存合并访问

模式：共享内存缓存

模式：共享内存 Bank Conflict 优化

模式：占用率优化

模式：操作融合

模式：低精度计算

案例研究

案例：GEMM Kernel 优化

案例：卷积 Kernel 优化

案例：多内核流水线

团队协作

基线管理

性能审查清单

常见问题速查表