cuda的内存结构以及编程基础

lay_ls

于 2024-10-27 22:31:16 发布

阅读量684

点赞数 9

文章标签：硬件工程

本文链接：https://blog.youkuaiyun.com/lay523/article/details/143276191

版权

GPU的内存体系结构

GPU与CPU的简单对比

如下图所示，左图表示 CPU 结构，右图表示 GPU 结构。可以清楚的看到CPU 和 GPU 在结构上的巨大差异，与CPU相⽐，GPU的体系架构中的⼀个最⼤特点就是增加了⼤量的运算单元ALU。这种天然的结构差异使的CPU主要是来处理复杂的逻辑控制，而GPU则可以做大量的并行计算。

与CPU相⽐，每个核的结构简单了很多，通常不⽀持⼀些CPU中使⽤的较为复杂的调度机制。GPU早期主要是为了做图形处理和渲染⽽被提出的，由于图形渲染需要⼤量的并⾏计算能⼒，⽽GPU的并⾏度，也恰恰好给像矩阵计算这样的任务提供了较好的并⾏计算能⼒，因此，以英伟达GPU为例，很快推出了通⽤计算能⼒（如通过CUDA编程模型来实现⾼效的并⾏算法），深度学习的⾼效计算也是利⽤了GPU的通⽤计算能⼒。

但是，我们一般说的 GPU 计算并不是纯粹由 GPU 完成的，而是由 CPU 调度、GPU 计算共同完成的异构计算任务。在异构计算任务中，CPU 被称为主机（Host），GPU 被称为设备（device）。

v2-df49a98a67c5b8ce55f1a9afcf21d982_1440w

CUDA编程模型

上文介绍到，英伟达适配GPU的架构提出了一套针对并行计算的编程语言CUDA(Compute Unified Device Architecture) 。CUDA 程序一般使用 .cu 后缀，编译 CUDA 程序则使用 nvcc 编译器。一般而言，典型的CUDA程序的执行流程如下：

分配host内存，并进行数据初始化；
分配device内存，并从host将数据拷贝到device上；
调用CUDA的核函数在device上完成指定的运算；
将device上的运算结果拷贝到host上；
释放device和host上分配的内存。

int main() {
    主机代码;
    核函数调用;
    主机代码;
    核函数调用;
    ......
    return 0;  
}

__global__ void 核函数1(parameters) {
    ......
}

__global__ void 核函数2(parameters) {
    ......
}

通过伪代码可以看到，CUDA编程模型是一个异构模型，需要CPU和GPU协同工作。在CUDA中，**host(主机端)和device(设备端也就是GPU端)**是两个重要的概念，我们用host指代CPU及其内存，而用device指代GPU及其内存。CUDA程序中既包含host程序，又包含device程序，它们分别在CPU和GPU上运行。同时，host与device之间可以进行通信，这样它们之间可以进行数据拷贝。

v2-2b30613b9925ea9e1fc68a837f60f036_1440w

在 main 函数中，我们穿插地写主机代码和核函数调用代码。主机代码主要负责 CPU 和 GPU 的内存管理、计算任务的分派，而核函数调用主要负责完成主要的计算工作。而前缀 __global__ 用来定义一个核函数，在 CUDA 中，核函数只能返回 void 类型，这意味着当我们需要写计算结果时，应该在参数列表中传入一个用来存放计算结果的指针，然后将计算结果写回到这个指针指向的存储空间中。

除了 __global__ 以外，CUDA 程序中的函数还可以使用 __device__ 和 __host__ 来修饰函数

__divice__ 修饰的函数称为设备函数，只能被核函数或是其它设备函数调用，只能在设备中执行。
__host__ 修饰的函数称为主机函数，它就是主机端的普通 C++ 函数，在主机（CPU）中调用和执行，可以忽略。

v2-71d9417bb5549057a851ca4594e681d4_1440w

CUDA线程

在介绍cuda中的线程时，来简单介绍一下CUDA的内存模型。现代计算机中的内存往往存在一种组织结构（hierarchy）。在这种结构中，含有多种类型的内存，每种内存分别具有不同的容量和延迟（latency，可以理解为处理器等待内存数据的时间）。一般来说，延迟低（速度高）的内存容量小，延迟高（速度低）的内存容量大。

下图则展示了CUDA 内存模型的层次结构，每一种都有不同的作用域、生命周期以及缓存行为。

全局内存（global memory）：全局内存是GPU中最大、延迟最高、最长使用的内存，通常说的“显存”中的大部分都是全局内存。全局内存的声明可以上被核函数的所有线程访问到，并且贯穿应用程序的整个生命周期。全局内存由于没有存放在GPU的芯片上，因此具有较高的延迟和较低的访问速度。全局内存的主要作用是为核函数提供数据，并在主机与设备及设备与设备之间传递数据。
寄存器内存（register）：在核函数中定义的不加任何限定符的变量一般来说就存放于寄存器（register）中。核函数中定义的不加任何限定符的数组有可能存放于寄存器中，但也有可能存放于局部内存中。另外，以前提到过的各种内建变量，如 gridDim、blockDim、blockIdx、 threadIdx 及 warpSize 都保存在特殊的寄存器中。在核函数中访问这些内建变量是很高效的。
局部内存（local memory）：局部内存和寄存器几乎一样。核函数中定义的不加任何限定符的变量有可能在寄存器中，也有可能在局部内存中。寄存器中放不下的变量，以及索引值不能在编译时就确定的数组，都有可能放在局部内存中。
共享内存（shared memory）：共享内存和寄存器类似，存在于芯片上，具有仅次于寄存器的读写速度，数量也有限。不同于寄存器的是，共享内存对整个线程块可见，其生命周期也与整个线程块一致。也就是说，每个线程块拥有一个共享内存变量的副本。共享内存变量的值在不同的线程块中可以不同。一个线程块中的所有线程都可以访问该线程块的共享内存变量副本，但是不能访问其他线程块的共享内存变量副本。
常量内存（constant memory）：常量内存（constant memory）是有常量缓存的全局内存，数量有限，一共仅有 64 KB。它的可见范围和生命周期与全局内存一样。不同的是，常量内存仅可读、不可写。由于有缓存，常量内存的访问速度比全局内存高，但得到高访问速度的前提是一个线程束中的线程（一个线程块中相邻的 32 个线程）要读取相同的常量内存数据。
纹理内存（texture memory）：纹理内存（texture memory）和表面内存（surface memory）类似于常量内存，也是一种具有缓存的全局内存，有相同的可见范围和生命周期，而且一般仅可读（表面内存也可写）。不同的是，纹理内存和表面内存容量更大，而且使用方式和常量内存也不一样。

介绍完cuda的内存结构之后，我们来看一下cuda是如何组织线程结构的，CUDA中⾸先将⼀组线程（通常不超1024个）组成⼀个线程块（block），每个线程块中的线程⼜可以分成多个warp被调度到GPU核上执⾏，⼀个线程块可以在⼀个SM上运⾏。多个线程块⼜可以组成⼀个⽹格（grid）。线程块和⽹格分别通过⼀个3维整数类型描述其⼤⼩( blockDim 和 gridDim )，每个线程都可以通过 threadIdx 和 blockIdx 来确定其属于哪个线程块以及哪个线程。

核函数在device上执行时实际上是启动很多线程，一个kernel所启动的所有线程称为一个网格（grid），同一个网格上的线程共享相同的全局内存空间。

网格（grid）：由多个线程块组成，同一个网格上的线程共享相同的全局内存空间(global memory)
线程块(block)：由多个线程组成
线程：GPU并行的最小单元
线程束：线程束包含32个线程，这些线程同时执行相同的指令，但是每个线程都包含自己的指令地址计数器和寄存器状态，也有自己独立的执行路径。

v2-aa6aa453ff39aa7078dde59b59b512d8_1440w

v2-6456af75530956da6bc5bab7418ff9e5_1440w

上面讲到，线程⼜可以分成多个warp被调度到GPU核上执⾏，⼀个线程块可以在⼀个SM上运⾏。可以看见真正在GPU上体现的硬件架构就是一个SM也被称为流处理器，一个 GPU 是由多个 SM 构成的。一个 SM 包含如下资源:

一定数量的寄存器。
一定数量的共享内存。
常量内存的缓存。
纹理和表面内存的缓存。
L1缓存。
线程束调度器(warp scheduler) 。
执行核心，包括:
- 若干整型数运算的核心(INT32) 。
- 若干单精度浮点数运算的核心(FP32) 。
- 若干双精度浮点数运算的核心(FP64) 。
- 若干单精度浮点数超越函数(transcendental functions)的特殊函数单元(Special Function Units，SFUs)。
- 若干混合精度的张量核心(tensor cores)
从软件的角度来讲：
1. 线程处理器 (SP)(coda核心) 对应线程 (thread)。
2. 多核处理器 (SM) 对应线程块 (thread block)。
3. 设备端 (device) 对应线程块组合体 (grid)。