CUDA编程方法比较：矩阵乘法

本示例演示并比较了使用CUDA进行GPU编程的不同方法，重点关注矩阵乘法问题。通过使用各种技术实现相同的算法，我们可以了解编程复杂性、性能和代码可维护性之间的权衡。

您可以在 https://github.com/eunomia-bpf/basic-cuda-tutorial 找到代码

概述

矩阵乘法是一个非常适合并行化的经典问题。本示例使用七种不同的方法实现了简单的矩阵乘法C = A × B：

1. 标准CUDA C/C++
2. 使用内联PTX汇编的CUDA
3. CUDA统一内存
4. CUDA共享内存
5. Thrust（高级C++抽象）
6. CUDA流
7. CUDA动态并行

对于每种实现，我们测量并比较执行时间，并与CPU实现的结果进行验证，确保正确性。

编程方法解释

1. 标准CUDA C/C++

标准CUDA方法使用显式内存管理和内核函数：

__global__ void matrix_multiply_cuda(float *A, float *B, float *C, 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;
    }
}

这种实现需要手动进行： - 主机和设备上的内存分配 - 主机和设备之间的数据传输 - 内核启动配置 - 同步 - 内存释放

优势： - 对内存管理有直接控制 - 良好的性能 - 熟悉的编程模型

劣势： - 需要显式内存管理 - 代码更冗长 - 内存传输可能成为瓶颈

实现细节： - 每个线程计算输出矩阵的一个元素 - 线程组织在2D网格中，直接映射到输出矩阵维度 - 内核对每个线程具有O(n)的计算复杂度 - 所有矩阵元素都使用全局内存

2. 使用内联PTX汇编的CUDA

PTX（并行线程执行）是NVIDIA的低级汇编语言。我们可以直接在CUDA C++代码中嵌入PTX：

__device__ float multiply_add_ptx(float a, float b, float c) {
    float result;
    asm("fma.rn.f32 %0, %1, %2, %3;" : "=f"(result) : "f"(a), "f"(b), "f"(c));
    return result;
}

这种实现使用融合乘加（FMA）指令，在单个操作中计算a * b + c，可能会提高性能和数值精度。

优势： - 对特定操作的细粒度控制 - 可以利用特定架构的指令 - 对关键代码段可能具有更好的性能 - 直接访问CUDA C/C++中未公开的硬件特性

劣势： - 在不同架构之间的可移植性最差 - 最复杂的编写和维护 - 需要深入了解GPU架构 - 特定架构的优化可能会随着新硬件的出现而过时

实现细节： - 使用fma.rn.f32 PTX指令进行融合乘加 - .rn后缀指定舍入到最近偶数的舍入模式 - 该指令在兼容硬件上在单个时钟周期内执行 - PTX汇编允许精确控制指令选择和调度

3. CUDA统一内存

统一内存提供了一个可由CPU和GPU同时访问的单一内存空间：

cudaMallocManaged(&u_A, matrix_size);
cudaMallocManaged(&u_B, matrix_size);
cudaMallocManaged(&u_C, matrix_size);

内核代码与标准CUDA版本相同，但内存管理得到了简化。

优势： - 简化的内存管理 - 无需显式数据传输 - 更容易的编程模型 - CPU和GPU之间自动页面迁移 - 支持大于GPU内存的数据集

劣势： - 由于自动页面迁移，可能性能较低 - 对数据移动的控制较少 - 性能很大程度上取决于访问模式 - 首次访问开销和潜在的页面错误

实现细节： - 使用cudaMallocManaged()代替单独的malloc()和cudaMalloc() - CUDA运行时自动处理数据传输 - 内存页根据需求在CPU和GPU之间迁移 - 相同的指针可以在主机和设备上使用 - 可以使用预取提示来优化数据移动（本示例中未显示）

4. CUDA共享内存

共享内存是一种快速的片上内存，可由块中的所有线程访问：

__global__ void matrix_multiply_shared(float *A, float *B, float *C, int n) {
    __shared__ float shared_A[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float shared_B[BLOCK_SIZE][BLOCK_SIZE];

    // 将矩阵的瓦片加载到共享内存中
    // ...

    // 使用更快的共享内存进行计算
    // ...
}

这种实现将矩阵分成适合共享内存的瓦片，减少了全局内存访问。

优势： - 内存访问速度快得多（通常比全局内存快100倍） - 减少全局内存带宽需求 - 对内存受限的应用有更好的性能 - 支持线程块内的数据重用

劣势： - 共享内存大小有限（通常每个SM 48KB-64KB） - 更复杂的编程模型 - 需要仔细管理瓦片大小 - 如果设计不当，可能会导致存储体冲突

实现细节： - 矩阵被分成大小为BLOCK_SIZE × BLOCK_SIZE的瓦片 - 每个线程块从每个输入矩阵加载一个瓦片到共享内存 - 块内的线程使用__syncthreads()同步 - 每个线程计算输出瓦片的一个元素 - 该算法将全局内存访问减少了BLOCK_SIZE倍

5. Thrust高级实现

Thrust是CUDA的C++模板库，提供高级抽象：

void run_thrust_implementation(float *h_A, float *h_B, float *h_C, int n) {
    thrust::device_vector<float> d_A(h_A, h_A + n * n);
    thrust::device_vector<float> d_B(h_B, h_B + n * n);
    thrust::device_vector<float> d_C(n * n);

    // 创建2D索引空间并转换
    // ...
}

优势： - 代码最简洁 - 类似于STL的高级抽象 - 自动内存管理 - 高度可重用的组件 - 减少开发时间和减少错误

劣势： - 对实现细节的控制较少 - 对特定用例可能性能较低 - 可能更难调试 - 抽象可能带来的开销

实现细节： - 使用thrust::device_vector进行自动GPU内存管理 - 利用thrust::transform算法和自定义函子 - 使用thrust::make_zip_iterator和thrust::counting_iterator创建2D索引空间 - 函子为每个元素实现矩阵乘法 - 内存传输由Thrust容器自动处理

6. CUDA流

CUDA流使GPU上的并发操作成为可能：

void run_cuda_streams_implementation(float *h_A, float *h_B, float *h_C, int n) {
    // 创建多个CUDA流
    const int numStreams = 4;
    cudaStream_t streams[numStreams];

    // 在流之间分配工作
    for (int i = 0; i < numStreams; i++) {
        // 异步内存传输和内核启动
        cudaMemcpyAsync(..., streams[i]);
        matrix_multiply_cuda<<<grid, threads, 0, streams[i]>>>(...);
        cudaMemcpyAsync(..., streams[i]);
    }

    // 同步所有流
    for (int i = 0; i < numStreams; i++) {
        cudaStreamSynchronize(streams[i]);
    }
}

优势： - 计算与数据传输重叠 - 支持并发内核执行 - 更好地利用GPU资源 - 可以显著提高整体吞吐量 - 适合处理独立的数据块

劣势： - 更复杂的同步 - 更难调试和推理 - 如果设计不当，可能会出现竞态条件 - 效益取决于硬件能力

实现细节： - 将矩阵分成水平条带，每个流一个 - 每个流独立处理其条带 - 使用cudaMemcpyAsync()进行非阻塞内存传输 - 在不同的流中使用流参数启动内核 - 每个流有自己的独立执行的命令队列 - 最终同步确保所有操作完成

7. 动态并行

动态并行允许CUDA内核启动嵌套内核：

__global__ void matrix_multiply_dynamic_parent(float *A, float *B, float *C, int n) {
    // 根据线程索引计算子矩阵位置
    // ...

    // 为这个子矩阵启动子内核
    multiply_submatrix<<<dimGrid, dimBlock>>>(A, B, C, n, row_start, col_start, subsize);
}

优势： - 支持动态细化工作的自适应算法 - 允许递归问题分解 - 可以高效处理不规则工作负载 - 减少CPU-GPU协调开销 - 对分治算法更自然的表达

劣势： - 从设备启动内核的额外开销 - 更复杂的资源管理 - 需要较新的GPU架构（计算能力3.5+） - 可能更深的调用栈和更多的寄存器使用

实现细节： - 父内核将矩阵分成大的子矩阵 - 父网格中的每个线程启动一个子网格处理一个子矩阵 - 子内核处理更小、更易管理的块 - 需要-rdc=true编译标志来支持可重定位设备代码 - 父网格在内核结束时自动与所有子网格同步

内存层次结构和访问模式

CUDA中不同类型的内存具有截然不同的性能特征：

内存类型	访问速度	范围	生命周期	缓存
全局内存	最慢	主机和所有线程	应用程序	仅L2
共享内存	快	线程块	内核执行	片上
寄存器	最快	单个线程	线程生命周期	片上
常量内存	中等	所有线程（读）	应用程序	特殊缓存
纹理内存	中等	所有线程（读）	应用程序	特殊缓存
本地内存	慢	单个线程	线程生命周期	仅L2
统一内存	可变	主机和所有线程	应用程序	系统管理

我们的实现展示了利用这种内存层次结构的不同方式：

1. 标准CUDA：对所有数据使用全局内存
2. PTX汇编：与标准CUDA相同的内存模式，但使用优化的指令
3. 统一内存：使用自动管理的内存
4. 共享内存：显式地将数据缓存在片上共享内存中
5. Thrust：通过容器抽象内存管理
6. CUDA流：使用全局内存，但传输重叠
7. 动态并行：使用全局内存，但具有层次化的访问模式

高级优化技术

除了示例中显示的方法外，其他优化技术包括：

1. 线程束洗牌指令：在不使用共享内存的情况下交换线程束内的数据
2. Tensor核心：用于矩阵运算的专用硬件（在较新的GPU上）
3. 持久线程：保持线程驻留以处理多个工作项
4. 寄存器阻塞：在寄存器中存储部分结果以减少内存流量
5. 内存合并：确保对齐、连续的内存访问模式
6. 指令级并行：调度独立操作以隐藏延迟
7. 循环展开：减少循环开销并增加指令级并行性
8. 函数内联：消除函数调用开销
9. 占用率优化：平衡资源使用以最大化活动线程束

选择合适的方法

使用这个决策框架选择适当的实现：

• 当需要平衡控制和可读性时使用标准CUDA
• 仅在性能关键部分可以从特定指令中受益时使用PTX汇编
• 对于具有适度性能需求的更容易开发的情况使用统一内存
• 当内存访问是瓶颈且可以利用内存局部性时使用共享内存
• 对于快速开发使用Thrust，特别是在实现标准算法时
• 当需要重叠计算和数据传输时使用CUDA流
• 对于从递归分解或自适应细化中受益的问题使用动态并行

编译器标志和优化级别

不同的编译器标志可以显著影响性能：

# 基本编译
nvcc -o basic_version file.cu

# 优化编译
nvcc -O3 -arch=sm_70 -o optimized_version file.cu

# 使用动态并行
nvcc -rdc=true -O3 -arch=sm_70 -o dynamic_parallel_version file.cu

# 使用快速数学（可能降低精度）
nvcc -O3 -use_fast_math -arch=sm_70 -o fast_math_version file.cu

我们的示例使用-O3进行高度优化，使用-rdc=true支持动态并行。

构建和运行

编译示例：

make basic03

运行：

./basic03

程序将： 1. 生成随机矩阵 2. 在CPU上计算结果作为参考 3. 运行每个GPU实现 4. 验证每个实现的正确性 5. 打印时间信息和与CPU相比的加速比

分析和性能分析

要更深入地分析这些实现的性能，请使用NVIDIA的分析工具：

# 基本分析
nvprof ./basic03

# 详细时间线
nvprof --export-profile timeline.nvvp ./basic03

# 对于较新版本，使用Nsight Systems
nsys profile ./basic03

需要监控的关键指标： - 全局内存加载/存储吞吐量 - 共享内存加载/存储吞吐量 - 实现的占用率 - SM效率 - 线程束执行效率 - 内存带宽利用率

未来方向和高级主题

GPU计算领域不断发展。一些值得探索的高级主题：

1. 多GPU编程：在多个GPU之间分配工作
2. 异构计算：最佳地结合CPU和GPU计算
3. 混合精度：在适当的地方使用较低精度以获得更好的性能
4. Tensor核心编程：利用专用硬件进行矩阵运算
5. 基于图的执行：使用CUDA图进行优化的工作流执行
6. 协作组：更灵活的线程同步能力
7. CUDA感知MPI：分布式系统中的直接GPU到GPU通信

进一步阅读

• CUDA C++编程指南
• CUDA PTX ISA文档
• CUDA统一内存
• CUDA共享内存
• Thrust文档
• CUDA流和并发
• CUDA动态并行
• CUDA最佳实践指南

Share on Share on