快速优化检查清单

了解 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 --offload-arch=mp_31 -O2 app.mu -lmusart -L/usr/local/musa/lib -o app
参数说明：

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%

快速优化检查清单

性能瓶颈分析：屋顶模型（Roofline Model）

简单判断方法

占用率优化

查询最优配置

影响占用率的因素

减少寄存器使用

内存访问优化

合并访问

使用共享内存

避免存储体冲突（Bank Conflict）

数据结构优化

结构数组 vs 数组结构

内核优化技巧

循环展开

使用内置函数

性能分析工具

Moore Perf Compute 性能分析

占用率查询

优化检查清单

内核配置

内存访问

代码结构

性能对比示例

相关文档

这篇文章对您有帮助吗？

性能瓶颈分析：屋顶模型（Roofline Model）​

简单判断方法​

占用率优化​

查询最优配置​

影响占用率的因素​

减少寄存器使用​

内存访问优化​

合并访问​

使用共享内存​

避免存储体冲突（Bank Conflict）​

数据结构优化​

结构数组 vs 数组结构​

内核优化技巧​

循环展开​

使用内置函数​

性能分析工具​

Moore Perf Compute 性能分析​

占用率查询​

优化检查清单​

内核配置​

内存访问​

代码结构​

性能对比示例​

相关文档​

这篇文章对您有帮助吗？

性能瓶颈分析：屋顶模型（Roofline Model）

简单判断方法

占用率优化

查询最优配置

影响占用率的因素

减少寄存器使用

内存访问优化

合并访问

使用共享内存

避免存储体冲突（Bank Conflict）

数据结构优化

结构数组 vs 数组结构

内核优化技巧

循环展开

使用内置函数

性能分析工具

Moore Perf Compute 性能分析

占用率查询

优化检查清单

内核配置

内存访问

代码结构

性能对比示例

相关文档