CUDA 性能优化实战笔记 + C++示例代码（第七章重点）

一、核心笔记（整合前七章核心）

1. 性能优化基础认知

核心目标：最大化 GPU 资源利用率（计算吞吐量、内存带宽），减少延迟（内存访问、指令执行）。
优化前提：通过性能分析工具定位瓶颈，避免盲目优化（如内存瓶颈优先优化内存访问，计算瓶颈优化指令执行）。
关键原则：
- 算法优先于代码：选择适合 GPU 架构的算法（如分块矩阵乘法）；
- 编译器优化为基础：Release 配置比 Debug 配置性能提升显著；
- 硬件特性适配：利用共享内存、融合指令等 GPU 专属特性。

2. NVIDIA Nsight Compute 性能分析工具

核心功能

可视化分析 GPU 核函数的执行细节（计算吞吐量、内存访问、缓存命中率、占用率等）；
自动生成优化建议（如长记分板阻塞、未使用融合指令等）；
支持多版本对比，量化优化效果。

关键使用步骤

配置性能计数器访问：
- 创建 /etc/modprobe.d/nvidia.conf，添加内容：options nvidia NVreg_RestrictProfilingToAdminUsers=0；
- 重启系统生效。
启动工具：在容器或本地终端执行 ncu-ui，打开图形界面。
创建项目：
- 选择目标平台（如 Linux x86_64）、可执行文件路径、工作目录；
- 选择分析模式（System Trace 快速定位热点，Profile 详细分析核函数）。
关键分析指标：
- GPU Speed of Light Throughput：计算/内存吞吐量（低于 60%通常存在瓶颈）；
- Roofline Chart：展示算术强度与性能的关系，判断是内存受限还是计算受限；
- Memory Chart：内存访问分布（全局内存/共享内存/L1/L2 缓存）；
- Occupancy：SM 资源利用率（线程数、寄存器、共享内存对占用率的影响）。

3. 核心优化策略（矩阵乘法实战）

（1）编译配置优化（最基础且高效）

Debug vs Release：
- Debug：无优化，含调试信息，性能差（如矩阵乘法耗时 330ms）；
- Release：开启`-O3`优化（循环展开、寄存器复用等），性能提升显著（耗时降至 47ms）。

CMake 配置命令：

cmake -DCMAKE_BUILD_TYPE=Release ..  # 切换Release模式
make  # 重新编译

（2）共享内存优化（突破内存瓶颈）

核心原理：将全局内存中的高频访问数据（如矩阵分块 Tile）预加载到共享内存（速度是全局内存的 100+倍），减少全局内存访问次数。
关键步骤：
1. 定义分块大小（Tile Size，如 16×16，适配共享内存容量）；
2. 线程分工加载矩阵分块到共享内存；
3. 块内同步（__syncthreads()）确保数据加载完成；
4. 基于共享内存计算部分和，减少全局内存访问。

（3）循环展开优化

核心原理：通过展开循环减少分支指令和循环控制开销，提高指令级并行度。
实现方式：使用编译器指令 #pragma unroll，自动展开循环（无需手动修改循环结构）。
适用场景：内层循环（如矩阵乘法中的 K 维度循环），循环次数固定且较小。

（4）融合指令优化

核心原理：融合乘加指令（FMA，如`fmaf(a,b,c) = a*b + c`）将两个指令合并为一个，减少指令数、避免中间舍入误差，提升计算吞吐量。
GPU 支持：所有现代 NVIDIA GPU 均支持 FMA，需手动替换普通乘加运算为 FMA 函数。

优化效果对比（1024×1024 矩阵乘法）

优化版本	耗时（ms）	核心优化点
朴素版（Debug）	330.37	无优化，全局内存直接访问
朴素版（Release）	47.33	编译器`-O3`优化
共享内存版（Debug）	443.30	分块加载到共享内存
共享内存+循环展开（Debug）	443.31	共享内存+`#pragma unroll`
共享内存+循环展开+FMA（Debug）	495.03	共享内存+循环展开+`fmaf`
整合所有优化（Release）	33.37	编译器优化+共享内存+循环展开+FMA

优化版本

耗时（ms）

核心优化点

朴素版（Debug）

330.37

无优化，全局内存直接访问

朴素版（Release）

47.33

编译器`-O3`优化

共享内存版（Debug）

443.30

分块加载到共享内存

共享内存+循环展开（Debug）

443.31

共享内存+`#pragma unroll`

共享内存+循环展开+FMA（Debug）

495.03

共享内存+循环展开+`fmaf`

整合所有优化（Release）

33.37

编译器优化+共享内存+循环展开+FMA

4. 常见优化误区

过度依赖编译器优化：Release 配置虽有效，但无法替代算法级优化（如共享内存分块）；
忽视内存访问模式：列优先访问矩阵会导致非连续内存访问，缓存命中率低；
盲目增加线程数：超过 GPU 资源限制会导致占用率下降，性能反而降低；
Debug 模式下评估优化效果：Debug 模式会禁用部分优化，需在 Release 模式下验证最终性能。

二、C++（CUDA）示例代码

示例 1：朴素矩阵乘法（优化基准）

核心功能

无任何优化的矩阵乘法，作为后续优化的性能基准，呼应第七章“朴素版本”。

#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <chrono>

using namespace std;
using namespace chrono;

#define CHECK_CUDA_ERR(err) \
    if (err != cudaSuccess) { \
        cerr << "CUDA错误：" << cudaGetErrorString(err) << "（行号：" << __LINE__ << "）" << endl; \
        exit(1); \
    }

const int MATRIX_SIZE = 1024;  // 矩阵尺寸
const dim3 BLOCK_DIM(16, 16); // 16×16线程/块

// GPU核函数：朴素矩阵乘法（无优化）
__global__ void matrixMulNaiveGPU(const float* d_A, const float* d_B, float* d_C, int width) {
    // 2D全局索引
    int row = threadIdx.y + blockIdx.y * blockDim.y;
    int col = threadIdx.x + blockIdx.x * blockDim.x;

    if (row < width && col < width) {
        float sum = 0.0f;
        // 直接访问全局内存，无优化
        for (int k = 0; k < width; ++k) {
            sum += d_A[row * width + k] * d_B[k * width + col];
        }
        d_C[row * width + col] = sum;
    }
}

// CPU串行矩阵乘法（基准）
void matrixMulCPU(const float* h_A, const float* h_B, float* h_C, int width) {
    for (int i = 0; i < width; ++i) {
        for (int j = 0; j < width; ++j) {
            float sum = 0.0f;
            for (int k = 0; k < width; ++k) {
                sum += h_A[i * width + k] * h_B[k * width + j];
            }
            h_C[i * width + j] = sum;
        }
    }
}

int main() {
    const size_t data_size = MATRIX_SIZE * MATRIX_SIZE * sizeof(float);

    // 主机内存初始化（A、B全1，C全0）
    vector<float> h_A(MATRIX_SIZE * MATRIX_SIZE, 1.0f);
    vector<float> h_B(MATRIX_SIZE * MATRIX_SIZE, 1.0f);
    vector<float> h_C_CPU(MATRIX_SIZE * MATRIX_SIZE, 0.0f);
    vector<float> h_C_GPU(MATRIX_SIZE * MATRIX_SIZE, 0.0f);

    // CPU基准测试
    auto cpu_start = high_resolution_clock::now();
    matrixMulCPU(h_A.data(), h_B.data(), h_C_CPU.data(), MATRIX_SIZE);
    auto cpu_time = duration_cast<seconds>(high_resolution_clock::now() - cpu_start).count();
    cout << "CPU串行矩阵乘法耗时：" << cpu_time << " s" << endl;

    // GPU设备内存分配
    float *d_A, *d_B, *d_C;
    CHECK_CUDA_ERR(cudaMalloc((void**)&d_A, data_size));
    CHECK_CUDA_ERR(cudaMalloc((void**)&d_B, data_size));
    CHECK_CUDA_ERR(cudaMalloc((void**)&d_C, data_size));

    // 2D网格大小
    const dim3 GRID_DIM((MATRIX_SIZE + BLOCK_DIM.x - 1) / BLOCK_DIM.x,
                        (MATRIX_SIZE + BLOCK_DIM.y - 1) / BLOCK_DIM.y);

    // GPU计时
    cudaEvent_t gpu_start, gpu_stop;
    CHECK_CUDA_ERR(cudaEventCreate(&gpu_start));
    CHECK_CUDA_ERR(cudaEventCreate(&gpu_stop));
    CHECK_CUDA_ERR(cudaEventRecord(gpu_start, 0));

    // 数据拷贝+核函数执行
    CHECK_CUDA_ERR(cudaMemcpy(d_A, h_A.data(), data_size, cudaMemcpyHostToDevice));
    CHECK_CUDA_ERR(cudaMemcpy(d_B, h_B.data(), data_size, cudaMemcpyHostToDevice));
    matrixMulNaiveGPU<<<GRID_DIM, BLOCK_DIM>>>(d_A, d_B, d_C, MATRIX_SIZE);
    CHECK_CUDA_ERR(cudaGetLastError());
    CHECK_CUDA_ERR(cudaMemcpy(h_C_GPU.data(), d_C, data_size, cudaMemcpyDeviceToHost));

    // 计时结束
    CHECK_CUDA_ERR(cudaEventRecord(gpu_stop, 0));
    CHECK_CUDA_ERR(cudaEventSynchronize(gpu_stop));
    float gpu_time_ms;
    CHECK_CUDA_ERR(cudaEventElapsedTime(&gpu_time_ms, gpu_start, gpu_stop));

    // 结果验证与输出
    bool valid = true;
    for (int i = 0; i < 10; ++i) {
        if (abs(h_C_CPU[i] - MATRIX_SIZE) > 1e-5 || abs(h_C_GPU[i] - MATRIX_SIZE) > 1e-5) valid = false;
    }
    cout << "GPU朴素矩阵乘法（Debug模式）耗时：" << gpu_time_ms << " ms" << endl;
    cout << "结果验证：" << (valid ? "正确" : "错误") << endl;
    cout << "GPU加速比（相对CPU）：" << (double)cpu_time * 1000 / gpu_time_ms << "x" << endl;

    // 释放资源
    CHECK_CUDA_ERR(cudaFree(d_A));
    CHECK_CUDA_ERR(cudaFree(d_B));
    CHECK_CUDA_ERR(cudaFree(d_C));
    CHECK_CUDA_ERR(cudaEventDestroy(gpu_start));
    CHECK_CUDA_ERR(cudaEventDestroy(gpu_stop));

    return 0;
}

编译与运行

Debug 模式编译：cmake -DCMAKE_BUILD_TYPE=Debug .. && make；
运行命令：./matrix_mul_naive；

预期输出：

CPU串行矩阵乘法耗时：129 s
GPU朴素矩阵乘法（Debug模式）耗时：330.5 ms
结果验证：正确
GPU加速比（相对CPU）：390.3x

示例 2：整合所有优化的矩阵乘法（最终版本）

核心功能

集成“共享内存分块+循环展开+融合指令”，配合 Release 编译优化，达到最优性能，呼应第七章“整合优化”。

#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <chrono>

using namespace std;
using namespace chrono;

#define CHECK_CUDA_ERR(err) \
    if (err != cudaSuccess) { \
        cerr << "CUDA错误：" << cudaGetErrorString(err) << "（行号：" << __LINE__ << "）" << endl; \
        exit(1); \
    }

const int MATRIX_SIZE = 1024;  // 矩阵尺寸
const int TILE_SIZE = 16;      // 分块大小（与BLOCK_DIM一致）
const dim3 BLOCK_DIM(TILE_SIZE, TILE_SIZE);  // 16×16线程/块

// GPU核函数：整合所有优化（共享内存+循环展开+FMA）
__global__ void matrixMulOptimizedGPU(const float* d_A, const float* d_B, float* d_C, int width) {
    // 共享内存：存储A和B的分块
    __shared__ float shared_A[TILE_SIZE][TILE_SIZE];
    __shared__ float shared_B[TILE_SIZE][TILE_SIZE];

    // 线程在块内的ID
    int tx = threadIdx.x;
    int ty = threadIdx.y;

    // 2D全局索引（映射到C矩阵的行和列）
    int row = ty + blockIdx.y * blockDim.y;
    int col = tx + blockIdx.x * blockDim.x;

    float sum = 0.0f;

    // 分块遍历K维度，加载分块到共享内存
    for (int tile_k = 0; tile_k < (width + TILE_SIZE - 1) / TILE_SIZE; ++tile_k) {
        // 加载A的当前分块（避免越界）
        if (row < width && (tile_k * TILE_SIZE + tx) < width) {
            shared_A[ty][tx] = d_A[row * width + tile_k * TILE_SIZE + tx];
        } else {
            shared_A[ty][tx] = 0.0f;
        }

        // 加载B的当前分块（避免越界）
        if (col < width && (tile_k * TILE_SIZE + ty) < width) {
            shared_B[ty][tx] = d_B[(tile_k * TILE_SIZE + ty) * width + col];
        } else {
            shared_B[ty][tx] = 0.0f;
        }

        __syncthreads();  // 等待所有线程加载完成

        // 循环展开+融合指令：计算当前分块的部分和
        #pragma unroll  // 编译器自动展开循环
        for (int k = 0; k < TILE_SIZE; ++k) {
            // FMA指令：sum = shared_A[ty][k] * shared_B[k][tx] + sum
            sum = fmaf(shared_A[ty][k], shared_B[k][tx], sum);
        }

        __syncthreads();  // 等待所有线程计算完成，避免覆盖共享内存
    }

    // 存储结果到C矩阵
    if (row < width && col < width) {
        d_C[row * width + col] = sum;
    }
}

int main() {
    const size_t data_size = MATRIX_SIZE * MATRIX_SIZE * sizeof(float);

    // 主机内存初始化（A、B全1）
    vector<float> h_A(MATRIX_SIZE * MATRIX_SIZE, 1.0f);
    vector<float> h_B(MATRIX_SIZE * MATRIX_SIZE, 1.0f);
    vector<float> h_C(MATRIX_SIZE * MATRIX_SIZE, 0.0f);

    // GPU设备内存分配
    float *d_A, *d_B, *d_C;
    CHECK_CUDA_ERR(cudaMalloc((void**)&d_A, data_size));
    CHECK_CUDA_ERR(cudaMalloc((void**)&d_B, data_size));
    CHECK_CUDA_ERR(cudaMalloc((void**)&d_C, data_size));

    // 2D网格大小
    const dim3 GRID_DIM((MATRIX_SIZE + BLOCK_DIM.x - 1) / BLOCK_DIM.x,
                        (MATRIX_SIZE + BLOCK_DIM.y - 1) / BLOCK_DIM.y);

    // GPU计时（含数据传输）
    cudaEvent_t gpu_start, gpu_stop;
    CHECK_CUDA_ERR(cudaEventCreate(&gpu_start));
    CHECK_CUDA_ERR(cudaEventCreate(&gpu_stop));
    CHECK_CUDA_ERR(cudaEventRecord(gpu_start, 0));

    // 数据拷贝+核函数执行
    CHECK_CUDA_ERR(cudaMemcpy(d_A, h_A.data(), data_size, cudaMemcpyHostToDevice));
    CHECK_CUDA_ERR(cudaMemcpy(d_B, h_B.data(), data_size, cudaMemcpyHostToDevice));
    matrixMulOptimizedGPU<<<GRID_DIM, BLOCK_DIM>>>(d_A, d_B, d_C, MATRIX_SIZE);
    CHECK_CUDA_ERR(cudaGetLastError());
    CHECK_CUDA_ERR(cudaMemcpy(h_C.data(), d_C, data_size, cudaMemcpyDeviceToHost));

    // 计时结束
    CHECK_CUDA_ERR(cudaEventRecord(gpu_stop, 0));
    CHECK_CUDA_ERR(cudaEventSynchronize(gpu_stop));
    float gpu_time_ms;
    CHECK_CUDA_ERR(cudaEventElapsedTime(&gpu_time_ms, gpu_start, gpu_stop));

    // 结果验证（前10个元素预期值为MATRIX_SIZE）
    bool valid = true;
    for (int i = 0; i < 10; ++i) {
        if (abs(h_C[i] - MATRIX_SIZE) > 1e-5) valid = false;
    }

    // 输出结果（Release模式下）
    cout << "GPU优化矩阵乘法（Release模式）耗时：" << gpu_time_ms << " ms" << endl;
    cout << "结果验证：" << (valid ? "正确" : "错误") << endl;

    // 释放资源
    CHECK_CUDA_ERR(cudaFree(d_A));
    CHECK_CUDA_ERR(cudaFree(d_B));
    CHECK_CUDA_ERR(cudaFree(d_C));
    CHECK_CUDA_ERR(cudaEventDestroy(gpu_start));
    CHECK_CUDA_ERR(cudaEventDestroy(gpu_stop));

    return 0;
}

编译与运行

Release 模式编译：cmake -DCMAKE_BUILD_TYPE=Release .. && make；
运行命令：./matrix_mul_optimized；

预期输出（RTX 2060 为例）：

GPU优化矩阵乘法（Release模式）耗时：33.4 ms
结果验证：正确

性能对比：比朴素 Debug 版本快 9.9 倍，比朴素 Release 版本快 1.4 倍。

示例 3：Nsight Compute 性能分析脚本（辅助工具）

核心功能

提供一键式性能分析脚本，快速生成核函数的详细分析报告，呼应第七章“Profiling 工具使用”。

#!/bin/bash
# 脚本名称：profile_matrix_mul.sh
# 功能：使用Nsight Compute分析优化后的矩阵乘法核函数

# 编译Release版本
cmake -DCMAKE_BUILD_TYPE=Release ..
make -j4

# 运行Nsight Compute，生成分析报告（输出报告文件：matrix_mul_report%i.ncu-rep）
ncu -o matrix_mul_report --kernel-name matrixMulOptimizedGPU --set full ./matrix_mul_optimized

echo "性能分析完成！报告文件已生成（matrix_mul_report*.ncu-rep）"
echo "可通过命令查看报告：ncu --import matrix_mul_report_0.ncu-rep"

使用方法

将脚本保存为 profile_matrix_mul.sh；
赋予执行权限：chmod +x profile_matrix_mul.sh；
运行脚本：./profile_matrix_mul.sh；
查看报告：ncu --import matrix_mul_report_0.ncu-rep（或打开`ncu-ui`导入报告）。

三、关键说明

1. 代码与知识点关联

示例 1（朴素版本）：作为优化基准，展示无优化时的 GPU 性能，突出全局内存访问的瓶颈；
示例 2（整合优化版本）：核心展示第七章的四大优化策略，每个优化点均有对应代码实现（共享内存分块、#pragma unroll、`fmaf`函数）；
示例 3（分析脚本）：简化 Nsight Compute 的使用流程，帮助快速定位优化效果和潜在瓶颈。

2. 优化关键注意事项

分块大小（Tile Size）：建议设为 16 或 32（适配 GPU Warp 大小和共享内存容量），过大易导致共享内存不足，过小则降低数据复用率；
循环展开：仅适用于内层循环（循环次数固定），外层循环展开可能增加寄存器占用，降低占用率；
FMA 指令：需确保编译器未自动优化（Release 模式下可能已自动融合，手动替换可明确控制）；
性能验证：必须在 Release 模式下测试最终性能，Debug 模式的优化效果不具参考性。

3. 常见问题排查

共享内存越界：分块加载时需判断索引是否超出矩阵范围，超出部分填 0；
同步缺失：__syncthreads() 未正确使用会导致数据竞争，核函数结果错误；
占用率低：通过 Nsight Compute 的 Occupancy 面板调整线程数、寄存器使用量或共享内存大小；
分析工具无法访问 GPU：确保 nvidia.conf 配置正确，且已重启系统。