快速优化检查清单
了解 MT GPU 程序性能优化的基本概念和技术
性能瓶颈分析：屋顶模型（Roofline Model）
屋顶模型（Roofline Model）是可视化 GPU 性能上限的有效方式。它直观地显示内核性能与硬件极限之间的关系 ，帮助开发者快速定位性能瓶颈类型（计算受限或内存受限）。
核心指标是 算术强度（Arithmetic Intensity）：指每从内存中读取一个字节的数据，程序可以执行多少浮点运算，单位为 FLOP/Byte。
性能 (每秒浮点运算次数)
↑
│ ╭─────────────── 峰值计算性能
│ ╱│
│ ╱ │
│ ╱ │
│ ╱ │
│ ╱ │
│ ╱ │
│ ╱ 内存带宽极限
│ ╱ │
│ ╱ │
│╱ │
└───────────────────┴────────────→ 算术强度
0 脊点 (FLOP/Byte)
| 区域 | 诊断 | 优化策略 |
|---|---|---|
| 脊点左侧 | 内存瓶颈 | 增加算术强度（共享内存、数据重用、合并访问） |
| 脊点右侧 | 计算瓶颈 | 减少浮点运算 或使用更低精度（如 FP16、INT8） |
| 脊点位置 | 内存/计算平衡点 | 算法设计的理论目标 |
算术强度计算公式：
算术强度 = 每秒浮点运算数 / 每秒内存带宽总数 (Bytes)
简单判断方法
使用 muBLAS 数学库函数对比自定义内核：
#include <mublas.h>
// 计算密集型：矩阵乘法（调用 mublas）
// 特点：O(N³) 计算量，O(N²) 数据量，算术强度高
// C = alpha * A * B + beta * C, A 是 m×k, B 是 k×n, C 是 m×n
void computeBound(float* A, float* B, float* C, int m, int n, int k) {
mublasHandle_t handle;
mublasCreate(&handle);
const float alpha = 1.0f;
const float beta = 0.0f;
// mublas 矩阵乘法（列主序）：C = alpha * A * B + beta * C
// 对于 N×N 矩阵：计算量 = 2N³，数据量 = 3N²
// 算术强度 = 2N³ / 3N² = 2N/3 → N 越大越趋向计算瓶颈
mublasSgemm(handle, MUBLAS_OP_N, MUBLAS_OP_N,
m, n, k, &alpha, A, m, B, k, &beta, C, m);
mublasDestroy(handle);
}
// 内存密集型：简单向量运算
// 特点：O(N) 计算量，O(N) 数据量，算术强度低
__global__ void memoryBound(float* data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
data[idx] = data[idx] * 2.0f + 1.0f; // 2 次加载，1 次存储，2 次计算
}
}
💻 编译命令：使用 MT GPU 编译器编译上述代码
参数说明：
mcc：MUSA C/C++ 编译器--offload-arch=mp_31：指定目标架构（MTT S5000）-O2：标准优化（官方推荐）-lmusart：链接 MUSA 运行时库
判断技巧：
- 计算瓶颈：每个数据元素参与多次计算（如矩阵乘法、卷积）
- 内存瓶颈：每个数据元素只参与少量计算（如向量加法、数乘）
占用率优化
📖 术语说明：线程束是 GPU 调度的基本单位，MT GPU 与 CUDA 一致，每个线程束包含 32 个线程。流多处理器是 GPU 的核心计算单元。
占用率是指流多处理器上活跃线程束数量与最大线程束数量的比率。
查询最优配置
int minGridSize, blockSize;
// 参数说明：
// minGridSize：推荐使用的最小网格（Grid）大小
// blockSize：推荐使用的线程块（Block）大小
// myKernel：要分析的内核（Kernel）函数
// 0, 0：共享内存大小和寄存器数量（0 表示使用默认值）
musaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize,
myKernel, 0, 0);
// 使用推荐配置
myKernel<<<minGridSize, blockSize>>>(data);
影响占用率的因素
| 因素 | 影响 | 优化方法 |
|---|---|---|
| 寄存器使用 | 寄存器越多，占用率越低 | 减少局部变量、使用共享内存 |
| 共享内存 | 共享内存越多，占用率越低 | 合理分配共享内存 |
减少寄存器使用
// ❌ 寄存器使用过多
__global__ void highRegKernel(float* data) {
float temp[10]; // 大数组容易溢出到内存
for (int i = 0; i < 10; i++) {
temp[i] = data[i];
}
}
// ✅ 使用共享内存
__global__ void lowRegKernel(float* data) {
__shared__ float shared[256];
int tid = threadIdx.x;
shared[tid] = data[tid]; // 使用共享内存
}
内存访问优化
合并访问
// ✅ Good: 连续线程访问连续地址
__global__ void coalescedAccess(float* data) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
data[idx] = data[idx] * 2.0f;
}
// ❌ Bad: 跨步访问
__global__ void stridedAccess(float* data, int stride) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
data[idx * stride] = data[idx * stride] * 2.0f;
}
使用共享内存
// 使用共享内存减少全局内存访问
__global__ void sharedMemoryKernel(float* input, float* output, int n) {
__shared__ float cache[256];
int tid = threadIdx.x;
int idx = blockIdx.x * blockDim.x + threadIdx.x;
// 加载到共享内存
cache[tid] = input[idx];
__syncthreads();
// 在共享内存中计算
if (tid > 0 && tid < 255) {
output[idx] = (cache[tid-1] + cache[tid] + cache[tid+1]) / 3.0f;
}
}
避免存储体冲突（Bank Conflict）
// ❌ 32-way Bank Conflict
__shared__ float sharedA[32][32];
value = sharedA[threadIdx.x][col]; // 所有线程访问同一列
// ✅ 添加 padding
__shared__ float sharedB[32][33]; // 33 而不是 32
value = sharedB[threadIdx.x][col]; // 无冲突
数据结构优化
结构数组 vs 数组结构
结构数组（AoS）（不推荐）:
struct Point {
float x, y, z;
};
__global__ void processAoS(Point* points) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float x = points[idx].x * 2.0f; // 只访问 x，但加载了整个结构
}
数组结构（SoA）（推荐）:
struct Points {
float* x;
float* y;
float* z;
};
__global__ void processSoA(Points points, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float x = points.x[idx] * 2.0f; // 连续访 问，合并访问
}
内核优化技巧
循环展开
循环展开可以减少循环控制开销，提高指令级并行度。关键是要保证合并访问：
// ✅ 推荐：循环展开 + 合并访问
// 每个线程处理 4 个连续元素，相邻线程访问相邻地址
// 注意: 确保 gridDim.x * blockDim.x * 4 >= n，或使用尾循环处理剩余元素
__global__ void unrolledKernel(float* data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int stride = blockDim.x * gridDim.x;
int base = idx;
#pragma unroll 4
for (int i = 0; i < 4; i++) {
int offset = base + i * stride;
if (offset < n) {
data[offset] = data[offset] * 2.0f;
}
}
}
// ❌ 不推荐：非合并访问
// 相邻线程访问的地址不连续，带宽利用率低
__global__ void stridedUnrolledKernel(float* data, int n) {
int idx = blockIdx.x * blockDim.x * 4 + threadIdx.x;
int stride = blockDim.x * gridDim.x * 4;
#pragma unroll 4
for (int i = 0; i < stride; i += blockDim.x * gridDim.x * 4) {
data[idx + i] = data[idx + i] * 2.0f;
data[idx + i + 1] = data[idx + i + 1] * 2.0f;
data[idx + i + 2] = data[idx + i + 2] * 2.0f;
data[idx + i + 3] = data[idx + i + 3] * 2.0f;
}
}
使用内置函数
// ❌ 普通计算
float result = a * b + c;
// ✅ 使用融合乘加（融合乘加）
float result = fmaf(a, b, c); // 更快、更精确
性能分析工具
Moore Perf Compute 性能分析
📖 工具说明：
mcu（Moore Compute Utility）是 MT GPU 性能分析工具，用于分析内核性能、内存带宽、占用率等指标。
# 生成性能报告
mcu -o report ./application
# 查看 内核时间
mcu --metrics kernel_runtime -o report ./application
# 查看内存带宽
mcu --metrics gld_throughput,gst_throughput -o report ./application
# 常用指标说明：
# kernel_runtime: 内核执行时间
# gld_throughput: 全局内存加载吞吐量
# gst_throughput: 全局内存存储吞吐量
# achieved_occupancy: 实际占用率
详细用法参考: 性能分析工具
占用率查询
void checkOccupancy(void* kernel) {
int numBlocks;
musaOccupancyMaxActiveBlocksPerMultiprocessor(
&numBlocks, kernel, 256, 0);
printf("Max blocks per MP: %d\n", numBlocks);
}
优化检查清单
内核配置
| 原则 | 建议 |
|---|---|
| 线程块大小 | 使用线程束大小的倍数：128、256、512、1024 |
| 占用率 | 目标 > 50%，使用 musaOccupancyMaxPotentialBlockSize |
| 寄存器使用 | 减少局部变量，使用共享内存 |
- 线程块大小是否为 128/256/512/1024？
- 占用率是否超过 50%？
- 是否使用了
musaOccupancyMaxPotentialBlockSize？
内存访问
| 原则 | 建议 |
|---|---|
| 合并访问 | 相邻线程访问相邻地址 |
| 避免跨步 | 不使用 idx * stride 模式 |
| 共享内存 | 避免存储体冲突（Bank Conflict），添加 padding |
| 数据重用 | 频繁访问的数据放入共享内存 |
- 全局内存访问是否合并？
- 是否避免了跨步访问？
- 共享内存是否避免了存储体冲突（Bank Conflict）？
- 频繁访问的数据是否放入共享内存？
代码结构
- 是否减少了寄存器使用？
- 是否使用了循环展开（
#pragma unroll）？ - 是否使用了内置函数（如
fmaf）？ - 数据结构是否适合并行访问（SoA vs AoS）？
性能对比示例
优化前：
- 内核执行时间：10 ms
- 内存带宽：50 GB/s
- 流多处理器利用率：30%
优化后（合并访问 + 共享内存）：
- 内核执行时间：2 ms (5 倍提升)
- 内存带宽：250 GB/s
- 流多处理器利用率：75%
