第8章性能优化 - 榨干 GPU 的每一滴性能（Performance Optimization）

一、导读

1.1 本章学习目标

通过本章的学习，你将能够：

掌握 CUDA 性能分析的基本工具和方法

理解 Roofline 模型及其在性能分析中的应用

学会优化内存访问模式以提高带宽利用率

掌握指令级优化的技巧和方法

理解 occupancy 与性能的关系及优化策略

学会使用共享内存减少全局内存访问

掌握循环展开和寄存器优化的技术

能够识别和消除性能瓶颈

理解多 GPU 扩展的性能考虑

掌握系统化的性能调优方法论

1.2 为什么本章如此重要

编写正确的 CUDA 程序只是第一步，编写高效的 CUDA 程序才是真正挑战。即使算法正确，糟糕的性能也可能使程序在实际应用中不可用。性能优化是 CUDA 编程中最具挑战性也最有价值的技能。

GPU 提供了巨大的理论性能，但实际达到的性能往往远低于峰值。原因包括：

内存访问模式不佳导致带宽浪费
线程发散造成执行单元闲置
资源分配不当限制并行度
算法设计未充分利用硬件特性

本章将教你系统化的性能优化方法，从测量到分析再到优化，形成完整的调优闭环。这些技能不仅适用于 CUDA 编程，也是高性能计算领域的通用能力。

性能优化的经济价值：

在实际项目中，性能优化可以带来显著的经济效益：

降低成本：更高的性能意味着可以使用更少的硬件资源完成相同的工作

提高用户体验：实时应用（如游戏、视频处理）需要低延迟和高帧率

扩大应用范围：某些应用只有在达到一定性能阈值后才具有实用价值

节能减排：高效的代码消耗更少的能源，符合可持续发展理念

例如，在深度学习训练中，优化后的模型可以将训练时间从几周缩短到几天，大幅降低云计算成本。在科学模拟中，优化可以使原本需要数月才能完成的模拟在几天内完成。

1.3 核心问题引导

在深入学习之前，请思考以下问题：

如何准确测量 CUDA 程序的性能？

什么是 Roofline 模型？它如何指导优化？

内存带宽和计算能力，哪个是主要瓶颈？

如何提高内存访问的合并程度？

共享内存在什么情况下能带来最大收益？

寄存器溢出会有什么后果？

如何平衡 occupancy 和资源使用？

性能优化的系统化流程是什么？

二、核心概念详解

8.1 性能分析工具

nvprof（传统工具）

nvprof 是 NVIDIA 早期的性能分析工具，提供基本的内核执行时间、内存传输统计等信息。虽然已被 Nsight 系列取代，但在某些场景下仍然有用。

bash

nvprof ./my_cuda_program

Nsight Compute（推荐）

Nsight Compute 是现代 CUDA 性能分析的主要工具，提供：

详细的内核执行分析
内存访问模式可视化
占用率分析
指令级剖析
源码级别的性能映射

bash

ncu --set full ./my_cuda_program

Nsight Compute 的关键优势：

全面性：收集数百个性能指标

可视化：直观的图表展示瓶颈

建议：自动提供优化建议

比较：支持多次运行的对比分析

Nsight Systems

用于系统级分析，查看 CPU-GPU 交互、流重叠等。

bash

nsys profile ./my_cuda_program

Nsight Systems 特别适合分析：

CPU-GPU 同步点
多个流的并发执行
API 调用开销
多 GPU 系统的负载平衡

CUDA Events

程序内精确计时：

cpp

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start, stream);
// ... 要测量的操作 ...
cudaEventRecord(stop, stream);

cudaEventSynchronize(stop);
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);

CUDA Events 的优势：

高精度（微秒级）
异步记录，最小化测量开销
可以在程序中嵌入多个测量点

8.2 Roofline 模型

基本概念

Roofline 模型由 Williams 等人在 2009 年提出，用于可视化应用程序的性能上限。它将性能表示为算术强度（Arithmetic Intensity）的函数。

算术强度

算术强度定义为每字节数据传输所执行的浮点运算次数：

AI = FLOPs / Bytes

高算术强度的程序更有效地利用计算资源，因为它们对每个字节的数据执行更多计算。

两条限制线

内存带宽限制线：斜率为带宽的直线

```

Performance ≤ Bandwidth × AI

```

当程序受内存带宽限制时，性能随算术强度线性增长。

计算能力限制线：水平线，表示峰值 FLOPs

```

Performance ≤ Peak FLOPs

```

当程序受计算能力限制时，性能达到平台期，不再随算术强度增长。

拐点

两条线的交点称为拐点（Knee Point），对应的算术强度为：

AI_knee = Peak FLOPs / Bandwidth

拐点是区分内存受限和计算受限的临界点。

优化指导

AI < AI_knee：程序受内存带宽限制，优化内存访问（合并访问、共享内存缓存）
AI > AI_knee：程序受计算能力限制，优化指令效率（指令级并行、Tensor Core）

Performance
    ↑
    │        ┌────────────── ← 计算限制 (Peak FLOPs)
    │       /
    │      /
    │     / ← 内存限制 (Bandwidth × AI)
    │    /
    │   /
    └──┴──────────────────→ Arithmetic Intensity
       ↑
    AI_knee

实际应用示例：

假设某 GPU 的参数为：

峰值 FP32 性能：10 TFLOPS
峰值显存带宽：500 GB/s

则拐点算术强度为：

AI_knee = 10e12 / 500e9 = 20 FLOPs/byte

如果一个程序的算术强度为 5 FLOPs/byte，它是内存受限的；如果为 50 FLOPs/byte，则是计算受限的。

8.3 内存访问优化

合并访问（Coalesced Access）

当 warp 内的线程访问连续的内存地址时，硬件可以将多个访问合并为少数事务。

cpp

// 好的访问模式
__global__ void coalesced(float *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        data[idx] *= 2.0f; // 相邻线程访问相邻地址
    }
}

// 差的访问模式
__global__ void strided(float *data, int n, int stride) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx * stride < n) {
        data[idx * stride] *= 2.0f; // 步幅访问，无法合并
    }
}

合并访问的原理：

GPU 内存事务以 128 字节为单位
如果 warp 中的 32 个线程访问连续的 128 字节区域，只需一个事务
如果访问分散，可能需要多个事务，降低有效带宽

数据结构优化

cpp

// AoS (Array of Structures) - 可能导致非合并访问
struct Particle {
    float x, y, z;
    float vx, vy, vz;
};

// SoA (Structure of Arrays) - 更好的合并访问
struct Particles {
    float *x, *y, *z;
    float *vx, *vy, *vz;
};

SoA 布局的优势：

当只需要访问部分字段时，可以减少不必要的数据加载
更容易实现合并访问
更适合 SIMD/SIMT 架构

共享内存缓存

cpp

__global__ void sharedCacheOptimized(float *input, float *output, int n) {
    extern __shared__ float cache[];
    
    int tid = threadIdx.x;
    int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 加载到共享内存
    if (globalIdx < n) {
        cache[tid] = input[globalIdx];
    }
    __syncthreads();
    
    // 多次使用共享内存中的数据
    float result = 0.0f;
    for (int i = 0; i < windowSize; i++) {
        int neighborIdx = (tid + i) % blockDim.x;
        result += cache[neighborIdx];
    }
    
    if (globalIdx < n) {
        output[globalIdx] = result;
    }
}

共享内存的优势：

速度接近寄存器，远高于全局内存
可以用于线程间通信
可以作为软件管理的缓存

8.4 指令级优化

避免分支发散

cpp

// 有分支的版本
__global__ void withBranch(float *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        if (data[idx] > 0) {
            data[idx] = sqrt(data[idx]);
        } else {
            data[idx] = 0.0f;
        }
    }
}

// 无分支的版本
__global__ void branchless(float *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float val = data[idx];
        float result = (val > 0) ? sqrt(val) : 0.0f;
        data[idx] = result;
    }
}

分支发散的影响：

Warp 中的线程必须串行执行不同分支
性能下降与发散程度成正比
在某些情况下，可以使用谓词执行或查表法避免分支

使用内置函数

cpp

// 慢速精确版本
float result = sin(x);

// 快速近似版本
float result = __sinf(x);

// 其他快速函数
__cosf(x), __tanf(x), __logf(x), __expf(x)
__fdividef(a, b)  // 快速除法
__rsqrtf(x)       // 快速平方根倒数

内置函数的特点：

牺牲一定精度换取速度
适合对精度要求不高的应用
可以通过编译选项 --use_fast_math 全局启用

循环展开

cpp

// 普通循环
__global__ void normalLoop(float *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    for (int i = 0; i < 4; i++) {
        data[idx * 4 + i] *= 2.0f;
    }
}

// 展开循环
__global__ void unrolledLoop(float *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    data[idx * 4 + 0] *= 2.0f;
    data[idx * 4 + 1] *= 2.0f;
    data[idx * 4 + 2] *= 2.0f;
    data[idx * 4 + 3] *= 2.0f;
}

循环展开的好处：

减少循环控制开销
增加指令级并行
允许编译器更好地调度指令

8.5 Occupancy 优化

影响 Occupancy 的因素

线程块大小：更大的块可以填充更多 warp 槽位

寄存器用量：每个线程使用的寄存器越多，并发线程越少

共享内存用量：每块使用的共享内存越多，并发块越少

查询最优配置

cpp

int minGridSize, blockSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, 
                                    myKernel, 0, 0);

限制寄存器使用

cpp

// 编译时限制
// nvcc --maxrregcount=32 kernel.cu

// 或使用启动边界
__launch_bounds__(256, 4)
__global__ void myKernel(...) {
    ...
}

Occupancy 与性能的权衡

高 occupancy 不一定意味着最佳性能：

有时降低 occupancy 可以增加每线程可用资源
更好的算法可能比更高的 occupancy 更重要
内存带宽受限的程序对 occupancy 不敏感

8.6 寄存器优化

寄存器压力

当内核使用过多寄存器时：

并发线程数减少

可能发生寄存器溢出到本地内存（慢速）

检查寄存器用量

bash

nvcc --ptxas-options=-v kernel.cu

输出类似：

Used 16 registers, 352 bytes cmem[0], 4 bytes cmem[2]

减少寄存器使用的技巧

减少局部变量数量

重用变量

避免大型数组

让编译器决定分配

8.7 多 GPU 扩展

数据并行

将数据分割到多个 GPU：

cpp

for (int gpu = 0; gpu < numGPUs; gpu++) {
    cudaSetDevice(gpu);
    // 在该 GPU 上处理数据子集
}

流水线并行

不同 GPU 处理流水线的不同阶段。

通信开销

使用 NVLink 或 PCIe 进行 GPU 间通信，注意带宽限制。

三、程序员的视角

8.1 性能优化流程

系统化调优步骤

基准测试：建立性能基线

剖析：使用工具识别瓶颈

假设：分析瓶颈原因

优化：实施改进

验证：确认改进效果

迭代：重复直到满足目标

常见瓶颈类型

瓶颈类型	特征	优化方向
内存带宽	低 AI，高内存事务	合并访问，共享内存
计算能力	高 AI，低 occupancy	指令优化，并行度
延迟	小问题规模	批量处理，异步
同步	高同步开销	减少同步，warp 原语

8.2 常见陷阱与误区

误区一：优化总是值得的

事实：过早优化可能浪费时间。先确保正确性，再针对热点优化。

误区二：最高 occupancy 等于最佳性能

事实：性能取决于多种因素，occupancy 只是其中之一。

陷阱：忽视数据传输开销

GPU 计算再快，如果数据传输慢，整体性能仍然受限。

陷阱：过度优化

追求极致性能可能导致代码难以维护。在性能和可维护性之间取得平衡。

8.3 与其他层级的联系

本章综合了前面所有章节的知识：

第3章：线程模型影响并行度
第4章：内存优化是性能关键
第5章：同步影响执行效率
第6章：流重叠提高吞吐量
第7章：库提供优化基线

四、代码实战

8.1 完整的性能优化案例：矩阵乘法

cpp

#include <stdio.h>
#include <cuda_runtime.h>
#include <cublas_v2.h>

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            printf("CUDA error: %s\n", cudaGetErrorString(err)); \
            exit(1); \
        } \
    } while(0)

#define TILE_SIZE 16

// 版本 1：朴素实现
__global__ void matmulNaive(float *C, const float *A, const float *B, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (row < N && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < N; k++) {
            sum += A[row * N + k] * B[k * N + col];
        }
        C[row * N + col] = sum;
    }
}

// 版本 2：共享内存分块
__global__ void matmulTiled(float *C, const float *A, const float *B, int N) {
    __shared__ float tileA[TILE_SIZE][TILE_SIZE];
    __shared__ float tileB[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    
    float sum = 0.0f;
    
    for (int t = 0; t < (N + TILE_SIZE - 1) / TILE_SIZE; t++) {
        // 加载 tile A
        if (row < N && t * TILE_SIZE + threadIdx.x < N)
            tileA[threadIdx.y][threadIdx.x] = A[row * N + t * TILE_SIZE + threadIdx.x];
        else
            tileA[threadIdx.y][threadIdx.x] = 0.0f;
        
        // 加载 tile B
        if (col < N && t * TILE_SIZE + threadIdx.y < N)
            tileB[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
        else
            tileB[threadIdx.y][threadIdx.x] = 0.0f;
        
        __syncthreads();
        
        // 计算
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += tileA[threadIdx.y][k] * tileB[k][threadIdx.x];
        }
        
        __syncthreads();
    }
    
    if (row < N && col < N) {
        C[row * N + col] = sum;
    }
}

// 版本 3：寄存器分块（进一步优化）
__global__ void matmulRegisterTiled(float *C, const float *A, const float *B, int N) {
    const int BM = 64, BN = 64, BK = 8;
    const int TM = 8, TN = 8;
    
    __shared__ float As[BM][BK];
    __shared__ float Bs[BK][BN];
    
    int baseRow = blockIdx.y * BM;
    int baseCol = blockIdx.x * BN;
    
    int threadRow = threadIdx.x / (BN / TN);
    int threadCol = threadIdx.x % (BN / TN);
    
    float regC[TM][TN] = {{0}};
    
    for (int t = 0; t < (N + BK - 1) / BK; t++) {
        // 协作加载共享内存
        for (int r = 0; r < BM; r += blockDim.x) {
            int row = r + threadIdx.x;
            if (baseRow + row < N && t * BK + threadCol < N)
                As[row][threadCol] = A[(baseRow + row) * N + t * BK + threadCol];
        }
        
        for (int c = 0; c < BN; c += blockDim.x) {
            int col = c + threadIdx.x;
            if (baseCol + col < N && t * BK + threadRow < N)
                Bs[threadRow][col] = B[(t * BK + threadRow) * N + baseCol + col];
        }
        
        __syncthreads();
        
        // 计算
        for (int k = 0; k < BK; k++) {
            for (int i = 0; i < TM; i++) {
                for (int j = 0; j < TN; j++) {
                    int row = threadRow * TM + i;
                    int col = threadCol * TN + j;
                    if (baseRow + row < N && baseCol + col < N) {
                        regC[i][j] += As[row][k] * Bs[k][col];
                    }
                }
            }
        }
        
        __syncthreads();
    }
    
    // 写回
    for (int i = 0; i < TM; i++) {
        for (int j = 0; j < TN; j++) {
            int row = baseRow + threadRow * TM + i;
            int col = baseCol + threadCol * TN + j;
            if (row < N && col < N) {
                C[row * N + col] = regC[i][j];
            }
        }
    }
}

void benchmarkKernel(void (*kernel)(float*, const float*, const float*, int), 
                     const char *name, float *d_C, const float *d_A, const float *d_B, int N) {
    dim3 block(TILE_SIZE, TILE_SIZE);
    dim3 grid((N + TILE_SIZE - 1) / TILE_SIZE, (N + TILE_SIZE - 1) / TILE_SIZE);
    
    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));
    
    // 预热
    kernel<<<grid, block>>>(d_C, d_A, d_B, N);
    CUDA_CHECK(cudaDeviceSynchronize());
    
    // 测量
    CUDA_CHECK(cudaEventRecord(start));
    for (int iter = 0; iter < 10; iter++) {
        kernel<<<grid, block>>>(d_C, d_A, d_B, N);
    }
    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));
    
    float ms = 0;
    CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
    ms /= 10; // 平均每次
    
    float gflops = 2.0f * N * N * N / (ms * 1e6);
    printf("%-20s: %.3f ms, %.1f GFLOPS\n", name, ms, gflops);
    
    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));
}

int main() {
    int N = 1024;
    size_t size = N * N * sizeof(float);
    
    float *h_A = (float *)malloc(size);
    float *h_B = (float *)malloc(size);
    
    for (int i = 0; i < N * N; i++) {
        h_A[i] = (float)rand() / RAND_MAX;
        h_B[i] = (float)rand() / RAND_MAX;
    }
    
    float *d_A, *d_B, *d_C;
    CUDA_CHECK(cudaMalloc((void **)&d_A, size));
    CUDA_CHECK(cudaMalloc((void **)&d_B, size));
    CUDA_CHECK(cudaMalloc((void **)&d_C, size));
    
    CUDA_CHECK(cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice));
    
    printf("Matrix Multiplication Performance (N=%d):\n", N);
    printf("==========================================\n");
    
    benchmarkKernel(matmulNaive, "Naive", d_C, d_A, d_B, N);
    benchmarkKernel(matmulTiled, "Tiled (Shared Mem)", d_C, d_A, d_B, N);
    
    // cuBLAS 参考
    cublasHandle_t handle;
    cublasCreate(&handle);
    
    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));
    
    float alpha = 1.0f, beta = 0.0f;
    CUDA_CHECK(cudaEventRecord(start));
    for (int iter = 0; iter < 10; iter++) {
        cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, N, N, N, 
                    &alpha, d_A, N, d_B, N, &beta, d_C, N);
    }
    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));
    
    float ms = 0;
    CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
    ms /= 10;
    float gflops = 2.0f * N * N * N / (ms * 1e6);
    printf("%-20s: %.3f ms, %.1f GFLOPS\n", "cuBLAS", ms, gflops);
    
    cublasDestroy(handle);
    
    CUDA_CHECK(cudaFree(d_A));
    CUDA_CHECK(cudaFree(d_B));
    CUDA_CHECK(cudaFree(d_C));
    free(h_A);
    free(h_B);
    
    return 0;
}

8.2 内存带宽测试

cpp

#include <stdio.h>
#include <cuda_runtime.h>

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            printf("CUDA error: %s\n", cudaGetErrorString(err)); \
            exit(1); \
        } \
    } while(0)

// 带宽测试内核
__global__ void bandwidthTest(float *output, const float *input, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        output[idx] = input[idx] * 2.0f + 1.0f;
    }
}

int main() {
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);
    
    printf("Device: %s\n", prop.name);
    printf("Theoretical Memory Bandwidth: %.1f GB/s\n\n", 
           prop.memoryClockRate * 2.0 * prop.memoryBusWidth / 8.0 / 1e6);
    
    int sizes[] = {1 << 20, 1 << 22, 1 << 24, 1 << 26};
    int numSizes = 4;
    
    for (int s = 0; s < numSizes; s++) {
        int n = sizes[s];
        size_t bytes = n * sizeof(float);
        
        float *h_data = (float *)malloc(bytes);
        float *d_input, *d_output;
        
        CUDA_CHECK(cudaMalloc((void **)&d_input, bytes));
        CUDA_CHECK(cudaMalloc((void **)&d_output, bytes));
        
        int threadsPerBlock = 256;
        int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
        
        cudaEvent_t start, stop;
        CUDA_CHECK(cudaEventCreate(&start));
        CUDA_CHECK(cudaEventCreate(&stop));
        
        // 预热
        bandwidthTest<<<blocksPerGrid, threadsPerBlock>>>(d_output, d_input, n);
        CUDA_CHECK(cudaDeviceSynchronize());
        
        // 测量
        CUDA_CHECK(cudaEventRecord(start));
        for (int iter = 0; iter < 100; iter++) {
            bandwidthTest<<<blocksPerGrid, threadsPerBlock>>>(d_output, d_input, n);
        }
        CUDA_CHECK(cudaEventRecord(stop));
        CUDA_CHECK(cudaEventSynchronize(stop));
        
        float ms = 0;
        CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
        ms /= 100;
        
        // 计算有效带宽（读 + 写）
        float effectiveBw = 2.0 * bytes / (ms * 1e6);
        
        printf("Size: %6d KB -> Time: %.3f ms, Effective BW: %.1f GB/s\n", 
               bytes / 1024, ms, effectiveBw);
        
        CUDA_CHECK(cudaFree(d_input));
        CUDA_CHECK(cudaFree(d_output));
        CUDA_CHECK(cudaEventDestroy(start));
        CUDA_CHECK(cudaEventDestroy(stop));
        free(h_data);
    }
    
    return 0;
}

五、深入思考

5.1 性能优化的哲学

性能优化既是科学也是艺术。科学的一面在于有系统的测量和分析方法；艺术的一面在于需要直觉和经验来判断优化的方向。

Knuth 的警告

"过早优化是万恶之源"

这句话提醒我们，在确保正确性和清晰性之前，不要陷入微优化的陷阱。

Amdahl 定律的启示

优化对整体性能影响最大的部分。花 10 小时优化占 1% 运行时间的代码，不如花 1 小时优化占 50% 运行时间的代码。

5.2 前沿方向

自动调优：使用机器学习自动寻找最优参数
编译器优化：更智能的代码生成
新硬件特性：Tensor Core、Sparse Core 等专用单元

5.3 延伸阅读

"Programming Massively Parallel Processors" Chapter 8-10
NVIDIA Nsight Compute 用户指南
CUDA Best Practices Guide

六、实验笔记

实验 1：完整性能分析流程

目标：使用 Nsight Compute 分析并优化一个内核。

步骤：

编写一个中等复杂度的内核

使用 Nsight Compute 进行详细分析

识别主要瓶颈

实施优化

比较优化前后的性能

实验 2：Roofline 分析

目标：绘制应用程序的 Roofline 图。

步骤：

测量内核的 FLOPs 和内存访问量

计算算术强度

查询设备的峰值性能参数

在 Roofline 图上标出当前位置

根据位置确定优化方向

七、本章小结

本章介绍了 CUDA 性能优化的系统方法。我们学习了性能分析工具、Roofline 模型、内存优化、指令优化等技术。

关键要点包括：

使用 Nsight 工具进行精确的性能分析
Roofline 模型帮助识别瓶颈类型
合并访问是内存优化的核心
共享内存可以减少全局内存访问
Occupancy 是影响性能的重要因素之一
系统化调优流程：测量→分析→优化→验证

性能优化是一个持续的过程，需要耐心和实践。希望本章提供的知识和方法能够帮助你编写出高效的 CUDA 程序。

八、习题

解释 Roofline 模型的两个限制线及其含义。

什么是算术强度？如何计算？

列举三种提高内存合并访问的方法。

为什么共享内存可以提高性能？有什么限制？

如何查询内核的寄存器用量？

解释 occupancy 与性能之间的关系。

使用 Nsight Compute 分析一个内核，报告主要瓶颈。

编写代码演示循环展开的效果。

什么情况下应该优先优化内存访问而不是计算？

设计一个性能优化实验，比较不同线程块大小的影响。

第8章 性能优化 - 榨干 GPU 的每一滴性能（Performance Optimization）

一、导读

1.1 本章学习目标

1.2 为什么本章如此重要

1.3 核心问题引导

二、核心概念详解

8.1 性能分析工具

8.2 Roofline 模型

8.3 内存访问优化

8.4 指令级优化

8.5 Occupancy 优化

8.6 寄存器优化

8.7 多 GPU 扩展

三、程序员的视角

8.1 性能优化流程

8.2 常见陷阱与误区

8.3 与其他层级的联系

四、代码实战

8.1 完整的性能优化案例：矩阵乘法

8.2 内存带宽测试

五、深入思考

5.1 性能优化的哲学

5.2 前沿方向

5.3 延伸阅读

六、实验笔记

实验 1：完整性能分析流程

实验 2：Roofline 分析

七、本章小结

八、习题

第8章性能优化 - 榨干 GPU 的每一滴性能（Performance Optimization）