cuda_learning_02

发表于 更新于 分类于

cuda_learning_02

一、GPU 与 CUDA 结构
二、CUDA编程的要素
三、实践：PyTorch自定义CUDA算子

  内容主要来自知乎文章，CUDA（二）：GPU的内存体系及其优化指南，本文是学习笔记。

一、GPU的内存体系

各级内存及其特点

image-20241203173846140

CUDA 内存模型的层次结构

全局内存(global memory)

  最大，延迟最高、最长使用的内存，常说的“显存”中的大部分都是全局内存。可以用 cudaMemcpy函数将主机的数据复制到全局内存，或者反过来。

常量内存(constant memory)

  常量内存是指存储在片下存储的设备内存，但是通过特殊的常量内存缓存（constant cache）进行缓存读取，常量内存为只读内存。 常量内存数量有限，一共仅有 64 KB，由于有缓存，常量内存的访问速度比全局内存高，但得到高访问速度的前提是一个线程束中的线程(一个线程块中相邻的 32 个线程)要读取相同的常量内存数据。

  一个使用常量内存的方法是在核函数外面用 __constant__ 定义变量，并用 API 函数 cudaMemcpyToSymbol 将数据从主机端复制到设备的常量内存后 供核函数使用。

纹理内存(texture memory)和表面内存(surface memory)

  纹理内存和表面内存类似于常量内存，也是一 种具有缓存的全局内存，有相同的可见范围和生命周期，而且一般仅可读(表面内存也可 写)。不同的是，纹理内存和表面内存容量更大，而且使用方式和常量内存也不一样。

寄存器（register）

  寄存器是线程能独立访问的资源，它所在的位置与局部内存不一样，是在片上（on chip）的存储，用来存储一些线程的暂存数据。寄存器的速度是访问中最快的，但是它的容量较小。

  在核函数中定义的不加任何限定符的变量一般来说就存放于寄存器(register)中。 各种内建变量，如 gridDim、blockDim、blockIdx、 threadIdx 及 warpSize 都保存在特殊的寄存器中，以便高效访问。

  寄存器变量仅仅被一个线程可见，寄存器的生命周期也与所属线程的生命周期 一致，从定义它开始，到线程消失时结束。

局部内存(local memory)

  局部内存和寄存器几乎一 样，核函数中定义的不加任何限定符的变量有可能在寄存器中，也有可能在局部内存中。寄存器中放不下的变量，以及索引值不能在编译时就确定的数组，都有可能放在局部内存中。

  虽然称之为“局部内存”，但是其本质是设备全局内存（但不同于全局内存）中为每个线程单独分配的一块内存。所以，局部内存的延迟也很高，每个线程最多能使用高达512 KB的局部内存，但使用过多会降低程序的性能。

共享内存(shared memory)

  共享内存和寄存器类似，存在于芯片上，具有仅次于寄存器的读写速度，数量也有限。 一个使用共享内存的变量可以 __shared__ 修饰符来定义。

  共享内存对整个线程块可见，其生命周期也与整个线程块一致。共享内存的主要作用是减少对全局内存的访问，或者改善对全局内存的访问模式。

image-20241211170806477

各种内存特征表

L1和L2 缓存

  每个 SM 都有一个 L1 缓存，所有 SM 共享一个 L2 缓存。L1 和 L2 缓存都被用来存储局部内存和全局内存中的数据，也包括寄存器中溢出的部分，以减少延时。

  从物理结构上来说，在最新的GPU架构中，L1 缓存、纹理缓存及共享内存三者是统一的。但从编程的角度来看，共享内存是可编程的缓存(共享内存的使用完全由用户操控)，而L1 和 L2 缓存是不可编程的缓存(用户最多能引导编译器做一些选择)。

SM 构成及典型GPU的对比

一个 GPU 是由多个 SM 构成的。一个 SM 包含如下资源:

  • 一定数量的寄存器。

  • 一定数量的共享内存。

  • 常量内存的缓存。

  • 纹理和表面内存的缓存。

  • L1缓存。

  • 线程束调度器(warp scheduler) 。

  • 执行核心，包括:

    • 若干整型数运算的核心(INT32) 。
    • 若干单精度浮点数运算的核心(FP32) 。
    • 若干双精度浮点数运算的核心(FP64) 。
    • 若干单精度浮点数超越函数(transcendental functions)的特殊函数单元(Special Function Units，SFUs)。
    • 若干混合精度的张量核心(tensor cores)

  单精度浮点数超越函数（transcendental functions） 的 特殊函数单元（Special Function Units，SFUs） 是指一种硬件单元，它专门用于执行一些数学上超越（超出普通代数运算）函数的计算，如三角函数（sin, cos），指数函数（exp），对数函数（log），平方根（sqrt）等。

  张量核心（Tensor Cores）：主要用于加速深度学习中的矩阵运算，尤其是低精度浮点数运算（如 FP16）。

image-20241211173038616

H100 的 SM 结构图

  1. L1 Instruction Cache

    用于缓存线程束（warp）的指令，提高指令读取的效率。

  2. Warp Scheduler 和 Dispatch Unit

    • Warp Scheduler（线程束调度器）:

      每个 SM 中有多个 Warp Scheduler，每个调度器可以每个时钟周期调度 32 个线程（一个 Warp）。调度线程束执行任务，包括加载指令、分配执行单元等。

    • Dispatch Unit（指令派发单元）：

      Warp Scheduler 将指令分配给不同的执行单元（如 FP32、FP64、INT32 核心或 SFUs 等），由 Dispatch Unit 具体派发。

  3. Register File（寄存器文件）

  4. 执行核心

  5. L0 Instruction Cache 和数据缓存

    • L0 Instruction Cache：

      每个 Warp Scheduler 附带的更小的指令缓存，用于加速最近使用的指令。

    • L1 Data Cache / Shared Memory：

      每个 SM 配备 256 KB 的共享内存或 L1 数据缓存。

  6. Load/Store 单元

    负责从全局内存中加载数据或将计算结果存储到全局内存中。

  7. 纹理单元（Tex）

    纹理单元专门用于处理纹理数据加载，通常在图形渲染中使用。

  8. Tensor Memory Accelerator

    针对张量核心操作的特殊加速器，用于处理张量内存的加载和存储。

    image-20241211174124976

主流GPU对比图

GPU 之外：近存计算与存算一体

  在GPU的层次结构之外，为了降低访存成本，获得更高的性能，近存计算与存算一体逐渐成为热门的方向。

  近存计算: Graphcore IPU

  存算一体: 后摩智能 H30

  存算一体或者存内计算的核心思想是，通过对存储器单元本身进行算法嵌入，使得计算可以在存储器单元内完成。

二、通过规约(Reduction)操作理解GPU内存体系

  关于reduce优化，CUDA（二）：GPU的内存体系及其优化指南(blog1)和深入浅出GPU优化系列：reduce优化(blog2)都做出了很详细的讲解，两篇文章都写得非常详细，优化的角度也是大同小异的，下面的内容主要是整理两篇文章中提到优化方向，并对所有提到优化角度做一个小小的总结。

  首先，算法reduce即求解$x=x_0 \bigotimes x1 \bigotimes x_2 \bigotimes x_3 \bigotimes … \bigotimes x_n$。其中$\bigotimes$可表示为求sum，min，max，avg等操作，最后获得的输出相比于输入一般维度上会递减。在GPU中，reduce采用了一种树形的计算方式，并且由于GPU没有针对global数据的同步操作，只能针对block的数据进行同步。所以，reduce一般分为两个阶段。

image-20241219223345835

blog1

1. cpu 版本

1
2
3
4
5
6
7
8
9
real reduce(const real *x, const int N)
{
    real sum = 0.0;
    for (int n = 0; n < N; ++n)
    {
        sum += x[n];
    }
    return sum;
}
1
2
3
4
5
sum = 33554432.000000.
mx_time = 935.548340.
mi_time = 530.885864.
tot_time = 11747.090820.
avg_time = 587.354553.

2. 仅使用全局内存

  每个线程负责其唯一id对应的那个位置的值的计算，N=1e8，每个线程块有128个线程。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
void __global__ reduce_global(real *d_x, real *d_y)
{
    const int tid = threadIdx.x;
    real *x = d_x + blockDim.x * blockIdx.x;

    for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1)
    {
        if (tid < offset)
        {
            x[tid] += x[tid + offset];
        }
        __syncthreads();
    }

    if (tid == 0)
    {
        d_y[blockIdx.x] = x[0];
    }
}

3.

笔记

技巧1：解决warp divergence

  确实能够解决warp divergence，只有当工作线程小于32的时候才会出现warp divergence。需要注意的是这份代码里每个线程并不是负责计算自己的唯一id对应的那个位置的数字每次迭代后的值。

技巧2：解决bank冲突

  确实会存在bank冲突，也确实是第一次迭代是2路冲突，第二次是4路，再是8路，16路，这是因为bank冲突是针对一个wrap里的线程说的。

  解决方法其实就是从大到小枚举offset。

技巧3：解决idle线程