GPU显存的理解-天翼云开发者社区

一、GPU显存的基本概念

GPU显存是显卡上的内存，用于存储图像、视频和其他图形数据。与CPU内存不同，GPU显存是专门为图形处理和计算任务设计的。GPU显存通常是GDDR（Graphics Double Data Rate）类型的内存，具有高带宽和低延迟的特点。

GPU显存的大小和速度对于图形处理和计算任务的性能至关重要。较大的显存可以容纳更多的图像和视频数据，从而提高图形处理和计算任务的效率。较快的显存可以更快地读写数据，从而提高图形处理和计算任务的响应速度。

二、GPU显存的作用

GPU显存在图形处理和计算任务中发挥着重要作用。以下是GPU显存的主要作用：

存储图像和视频数据：GPU显存可以存储图像和视频数据，以便显卡进行图形处理和计算任务。较大的显存可以容纳更多的图像和视频数据，从而提高图形处理和计算任务的效率。
加速图形处理：GPU显存可以加速图形处理，如纹理映射、光照计算和几何变换等。较快的显存可以更快地读写数据，从而提高图形处理的响应速度。
加速计算任务：GPU显存可以加速计算任务，如深度学习、科学计算和密码学等。GPU显存可以存储大规模的数据集，从而提高计算任务的效率。

三、系统内存与设备内存

3.1 内存的架构

了解内存，需要先了解内存的几个关键要素：位置、大小、速度与传输通道。位置是指存储的硬件在什么地方，通过位置可知道的信息很多，如硬件单元是在芯片内部还是芯片外部，能不能插拔（意味着可扩展）以及对应的上下游存储是什么；大小与速度是指内存数据能存多大，读写数据的速度有多快；传输通道：存储与存储之间能够通过什么传输协议/通道进行数据交换。下图列出一个简化的系统内存与设备内存架构示意图：

系统存储：

L1/L2/L3：多级缓存，其位置一般在CPU芯片内部；
System DRAM：片外内存，内存条；
Disk/Buffer：外部存储，如磁盘或者固态硬盘。

GPU设备存储：

L1/L2 cache：多级缓存，其位置在GPU芯片内部；
GPU DRAM：通常所指的显存；

设备存储还包含许多片上存储单元，后面进行详细介绍。

传输通道：

PCIE BUS：PCIE标准的数据通道，数据就是通过该通道从显卡到达主机；
BUS：总线。计算机内部各个存储之间交互数据的通道；
PCIE-to-PCIE：显卡之间通过PCIE直接传输数据；
NVLINK：显卡之间的一种专用的数据传输通道，由NVIDIA公司推出。

1.2 传输通道的速度

对于传输通道的速度需要一个基本的认识，比如常见的PCIE以及专用的NVLINK通道速度的了解能给我们优化算法提供思路。下面对PCIE和NVLINK的速度进行介绍（参照新一代DDR5内存条支持的速度51.2 GB/s）：

PCIE的速度

第三代的PCIE x16的理论速度是16GB/s，所以受限于PCIE的速度，CPU到GPU速度小于16GB/s，(PCIe x16 第六代 2021：128GB/s）

NVLINK的速度

第二代的速度300GB/s，第三代速度：600GB/s。当前(2022年)发布的H100用了第四代NVLink，全通道速度可达900GB/s.

四、设备内部的存储

前面提到GPU的内部存储分为片上存储和片下存储，指的硬件所在位置，为了满足GPU的应用场景，对存储功能进行了细分，包括：局部内存(local memory)、全局内存（global memory）、常量内存（constant memory）、图像/纹理（texture memory）、共享内存（shared memory）、寄存器（register）、L1/L2缓存、常量内存/纹理缓存（constant/texture cache），下面逐个介绍一下。

其中涉及到一些名词，可以参考CUDA手册/NVIDA芯片手册理解，这里先通俗地解释一下：

SM（Streaming Multiprocessors）：理解为一个GPU内数据处理的大单元，好比多核的CPU芯片里面的一个核，CPU的一个核一般是运行一个线程，而SM能够运行多个轻量线程；
nvcc：GPU程序的编译器，其实就是针对CUDA特殊化的gcc编译器；
block： thread线程的集合单位。比如让GPU完成一个矩阵数据的运算，然后我们给参与运算的thread编个队，队名叫做block，对多个block编队就成了grid单位。
warp： SM里面的运算执行单位，理解为运算时一个warp抓一把thread 扔进了计算core里面进行计算。

4.1 全局内存

全局内存（global memory）是数据常用的内存，它能被设备内的所有线程访问、全局共享，为片下（off chip）内存，前面提到的硬件HBM中的大部分都是用作全局内存。跟CPU架构一样，运算单元不能直接的使用全局内存的数据，需要经过缓存，其过程如下图所示：

4.2 L1/L2缓存

L1/L2缓存（Cache）数据缓存，这个存储跟CPU架构的类似。L2为所有SM都能访问到，速度比全局内存块，所以为了提高速度有些小的数据可以缓存到L2上面；L1用于存储SM内的数据，SM内的运算单元能够共享，但跨SM之间的L1不能相互访问。

对于开发者来说，需要注意L2缓存能够提速运算，比如CUDA11 A100 上面L2缓存能够设置至多40MB的持续化数据(persistent data)，L2上面的持续化数据能够拉升算子kernel的带宽和性能，设置持续化数据的举例如下（摘取自CUDA 官网的exmaple）：

cudaGetDeviceProperties( &prop, device_id);
// Set aside 50% of L2 cache for persisting accesses size_t size = min( int(prop.l2CacheSize * 0.50) , prop.persistingL2CacheMaxSize );
cudaDeviceSetLimit( cudaLimitPersistingL2CacheSize, size); 
// Stream level attributes data structure cudaStreamAttrValue attr ;

attr.accessPolicyWindow.base_ptr = /* beginning of range in global memory */ ;

attr.accessPolicyWindow.num_bytes = /* number of bytes in range */ ;

// hitRatio causes the hardware to select the memory window to designate as persistent in the area set-aside in L2 attr.accessPolicyWindow.hitRatio = /* Hint for cache hit ratio */

// Type of access property on cache hit attr.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting;
// Type of access property on cache miss attr.accessPolicyWindow.missProp = cudaAccessPropertyStreaming;

cudaStreamSetAttribute(stream,cudaStreamAttributeAccessPolicyWindow,&attr);

通过persistent data可以使得运算提速1.5倍，具体参看：CUDA Toolkit Documentation

4.3 局部内存(local memory)

局部内存(local memory) 是线程独享的内存资源，线程之间不可以相互访问，硬件位置是off chip状态，所以访问速度跟全局内存一样。局部内存主要是用来解决当寄存器不足时的场景，即在线程申请的变量超过可用的寄存器大小时，nvcc会自动将一部数据放置到片下内存里面。

注意，局部内存设置的过程是在编译阶段就会确定。

4.4 寄存器（register）

寄存器（register）是线程能独立访问的资源，它所在的位置与局部内存不一样，是在片上（on chip）的存储，用来存储一些线程的暂存数据。寄存器的速度是访问中最快的，但是它的容量较小。以目前较新的Ampere架构的GA102为例，每个SM上的寄存器总量256KB，使用时被均分为了4块，且该寄存器块的64KB空间需要被warp中线程平均分配，所以在线程多的情况下，每个线程拿到的寄存器空间相当小。寄存器的分配对SM的占用率（occupancy）存在影响，可以通过CUDA Occupancy Calculator 计算比较，举例：当registers从32增加到128时，occupancy从100%降低到了33.33%。

4.5 共享内存（shared memory)

共享内存（shared memory) 是一种在block内能访问的内存，存储硬件位于芯片上（on chip），访问速度较快，共享内存主要是缓存一些需要反复读写的数据。可以通过一个矩阵运算的例子说明shared memory的作用，比如完成矩阵运算C = A X B， Ai_row表示A的第i行数据， Bj_col表示B的第j列数据，cij表示第i行第j例的数值，有：

假设要得到C矩阵的第i行Ci_row的数据，上述运算需要进行N次，N为：B矩阵列宽大小。

对于该计算而言，运算中的Ai_row保持不变，Bj_col进行迭代更新。Ai_row假设使用global memory，则每次运算都需要重新加载，数据重复加载了N次。然而Ai_row数据是可以复用的，所以将Ai_row放入共享内存中，这样相同的数据避免反复加载（Ai_row数据加载是要1次），从而提高运算效率。相比只用全局内存，共享内存在上述矩阵运算上可以提升20~50GB/s的速度。

参考示例：

__global__ void matrix_kernel_1(float* _C,const float* _A,const float *_B,int _wa,int _wb) //_wa是A矩阵的宽度，_wb是矩阵B的宽度
{
    int bx = blockIdx.x; //Block X的当前位置
    int by = blockIdx.y;　//Block　y的当前位置
    int tx = threadIdx.x;
    int ty = threadIdx.y;

    //该block要处理的A ,A的取值方向是X轴方向， B的取值方向是Y轴方向
    int aBegin = _wa*(by*BLOCK_SIZE);//A(0,by) //在矩阵A上每个block的首地址
    int aEnd = aBegin + _wa - 1; //
    int aStep = BLOCK_SIZE;//offsetA  //因为A是横向取值，所以step是blocksize

    int bBegin = BLOCK_SIZE*bx;//B(bx,0) //矩阵B的首地址
    int bStep = BLOCK_SIZE*_wb;//offsetB //因为B是纵向取值，所以step是blocksize*_wb.
    
    float cSub = 0;//每一个线程计算一个像素点，分成wa/block 次来计算，每次计算一段A(sub) * B(sub)，最后累加得到C的结果．//假设矩阵都是n*n的，那么旧的basicMatrix每个线程都需要执行2n次globalMemory的访问，这里用到sharedMemory只需要执行2n／blocksize，每个线程可以提高blocksize倍，／／每个block里面的thread都是通过读取sharedMemory来执行计算的，速度会非常快．
    for (int a = aBegin,b = bBegin; a <= aEnd; a += aStep,b += bStep)
    {
        __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
        __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
        //每个线程负责一个元素拷贝,我们以block为单位来分析。假设blocksize=16, 一个block里面有 16*16个线程。        //每个block 可以填满需要用到的 As, 和Bs大小的矩阵。这里就是矩阵A里面的16*16的数据可以填满，保存在sharedMemory中。同样B矩阵也是。
        As[ty][tx] = _A[a + _wa*ty + tx];   
        Bs[ty][tx] = _B[b + _wb*ty + tx];

        __syncthreads();　//同步使得矩阵Ａ，和矩阵Ｂ的第一个tile*tile的数据保存在As和Bs里，供下面的计算使用．
        
        //每个线程负责计算一个子块i 和 子块j的子乘积宽度是block_size,执行到wa/block_size次，累加可得到Ｃ的值
        for (int k = 0; k < BLOCK_SIZE; ++k)
        {
            cSub += As[ty][k]*Bs[k][tx];　
        }

        __syncthreads();
    }

    //全局地址，向全局寄存器写回去
    //一个线程负责一个元素，一个block负责一个子块
    int cIndex = (by*BLOCK_SIZE + ty)*_wb + (bx*BLOCK_SIZE + tx);
    _C[cIndex] = cSub;
}

4.6 常量内存(constant memory)

常量内存(constant memory) 是指存储在片下存储的设备内存上，但是通过特殊的常量内存缓存（constant cache）进行缓存读取，常量内存为只读内存。为什么需要设立单独的常量内存？直接用global memory或者shared memory不行吗？

它主要是解决一个warp内多线程的访问相同数据的速度太慢的问题。所有运算的thread都需要访问一个constant_A的常量，在存储介质上面constant_A的数据只保存了一份，而内存的物理读取方式决定了这么多thread不能在同一时刻读取到该变量，所以会出现先后访问的问题，这样使得并行计算的thread出现了运算时差。常量内存正是解决这样的问题而设置的，它有对应的cache位置产生多个副本，让thread访问时不存在冲突，从而提高并行度。

参考示例：

__constant__ int c1 = 10;  // 声明__constant__ 即可。
__global__ void kernel1(int *d_dst) {
   int tId = threadIdx.x + blockIdx.x * blockDim.x;
   d_dst[tId] += c1;
}

4.7 图像/纹理内存（texture memory）

图像/纹理（texture memory）是一种针对图形化数据的专用内存，其中texture直接翻译是纹理的意思，但根据实际的使用来看texture应该是指通常理解的1D/2D/3D结构数据，相邻数据之间存在一定关系，或者相邻数据之间需要进行相同的运算。 texture内存的构成包含 global + cache + 处理单元，texture为只读内存。texture的优势：

texture memory 进行图像类数据加载时， warp内的thread访问的数据地址相邻，从而减少带宽的浪费。
texture 在运算之前能进行一些处理（或者说它本身就是运算），比如聚合、映射等。

一、GPU显存的基本概念

二、GPU显存的作用

GPU显存在图形处理和计算任务中发挥着重要作用。以下是GPU显存的主要作用：

存储图像和视频数据：GPU显存可以存储图像和视频数据，以便显卡进行图形处理和计算任务。较大的显存可以容纳更多的图像和视频数据，从而提高图形处理和计算任务的效率。
加速图形处理：GPU显存可以加速图形处理，如纹理映射、光照计算和几何变换等。较快的显存可以更快地读写数据，从而提高图形处理的响应速度。
加速计算任务：GPU显存可以加速计算任务，如深度学习、科学计算和密码学等。GPU显存可以存储大规模的数据集，从而提高计算任务的效率。

三、系统内存与设备内存

3.1 内存的架构

系统存储：

L1/L2/L3：多级缓存，其位置一般在CPU芯片内部；
System DRAM：片外内存，内存条；
Disk/Buffer：外部存储，如磁盘或者固态硬盘。

GPU设备存储：

L1/L2 cache：多级缓存，其位置在GPU芯片内部；
GPU DRAM：通常所指的显存；

设备存储还包含许多片上存储单元，后面进行详细介绍。

传输通道：

PCIE BUS：PCIE标准的数据通道，数据就是通过该通道从显卡到达主机；
BUS：总线。计算机内部各个存储之间交互数据的通道；
PCIE-to-PCIE：显卡之间通过PCIE直接传输数据；
NVLINK：显卡之间的一种专用的数据传输通道，由NVIDIA公司推出。

1.2 传输通道的速度

PCIE的速度

第三代的PCIE x16的理论速度是16GB/s，所以受限于PCIE的速度，CPU到GPU速度小于16GB/s，(PCIe x16 第六代 2021：128GB/s）

NVLINK的速度

第二代的速度300GB/s，第三代速度：600GB/s。当前(2022年)发布的H100用了第四代NVLink，全通道速度可达900GB/s.

四、设备内部的存储

其中涉及到一些名词，可以参考CUDA手册/NVIDA芯片手册理解，这里先通俗地解释一下：

SM（Streaming Multiprocessors）：理解为一个GPU内数据处理的大单元，好比多核的CPU芯片里面的一个核，CPU的一个核一般是运行一个线程，而SM能够运行多个轻量线程；
nvcc：GPU程序的编译器，其实就是针对CUDA特殊化的gcc编译器；
block： thread线程的集合单位。比如让GPU完成一个矩阵数据的运算，然后我们给参与运算的thread编个队，队名叫做block，对多个block编队就成了grid单位。
warp： SM里面的运算执行单位，理解为运算时一个warp抓一把thread 扔进了计算core里面进行计算。

4.1 全局内存

4.2 L1/L2缓存

cudaGetDeviceProperties( &prop, device_id);
// Set aside 50% of L2 cache for persisting accesses size_t size = min( int(prop.l2CacheSize * 0.50) , prop.persistingL2CacheMaxSize );
cudaDeviceSetLimit( cudaLimitPersistingL2CacheSize, size); 
// Stream level attributes data structure cudaStreamAttrValue attr ;

attr.accessPolicyWindow.base_ptr = /* beginning of range in global memory */ ;

attr.accessPolicyWindow.num_bytes = /* number of bytes in range */ ;

// hitRatio causes the hardware to select the memory window to designate as persistent in the area set-aside in L2 attr.accessPolicyWindow.hitRatio = /* Hint for cache hit ratio */

// Type of access property on cache hit attr.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting;
// Type of access property on cache miss attr.accessPolicyWindow.missProp = cudaAccessPropertyStreaming;

cudaStreamSetAttribute(stream,cudaStreamAttributeAccessPolicyWindow,&attr);

通过persistent data可以使得运算提速1.5倍，具体参看：CUDA Toolkit Documentation

4.3 局部内存(local memory)

注意，局部内存设置的过程是在编译阶段就会确定。

4.4 寄存器（register）

4.5 共享内存（shared memory)

假设要得到C矩阵的第i行Ci_row的数据，上述运算需要进行N次，N为：B矩阵列宽大小。

参考示例：

__global__ void matrix_kernel_1(float* _C,const float* _A,const float *_B,int _wa,int _wb) //_wa是A矩阵的宽度，_wb是矩阵B的宽度
{
    int bx = blockIdx.x; //Block X的当前位置
    int by = blockIdx.y;　//Block　y的当前位置
    int tx = threadIdx.x;
    int ty = threadIdx.y;

    //该block要处理的A ,A的取值方向是X轴方向， B的取值方向是Y轴方向
    int aBegin = _wa*(by*BLOCK_SIZE);//A(0,by) //在矩阵A上每个block的首地址
    int aEnd = aBegin + _wa - 1; //
    int aStep = BLOCK_SIZE;//offsetA  //因为A是横向取值，所以step是blocksize

    int bBegin = BLOCK_SIZE*bx;//B(bx,0) //矩阵B的首地址
    int bStep = BLOCK_SIZE*_wb;//offsetB //因为B是纵向取值，所以step是blocksize*_wb.
    
    float cSub = 0;//每一个线程计算一个像素点，分成wa/block 次来计算，每次计算一段A(sub) * B(sub)，最后累加得到C的结果．//假设矩阵都是n*n的，那么旧的basicMatrix每个线程都需要执行2n次globalMemory的访问，这里用到sharedMemory只需要执行2n／blocksize，每个线程可以提高blocksize倍，／／每个block里面的thread都是通过读取sharedMemory来执行计算的，速度会非常快．
    for (int a = aBegin,b = bBegin; a <= aEnd; a += aStep,b += bStep)
    {
        __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
        __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
        //每个线程负责一个元素拷贝,我们以block为单位来分析。假设blocksize=16, 一个block里面有 16*16个线程。        //每个block 可以填满需要用到的 As, 和Bs大小的矩阵。这里就是矩阵A里面的16*16的数据可以填满，保存在sharedMemory中。同样B矩阵也是。
        As[ty][tx] = _A[a + _wa*ty + tx];   
        Bs[ty][tx] = _B[b + _wb*ty + tx];

        __syncthreads();　//同步使得矩阵Ａ，和矩阵Ｂ的第一个tile*tile的数据保存在As和Bs里，供下面的计算使用．
        
        //每个线程负责计算一个子块i 和 子块j的子乘积宽度是block_size,执行到wa/block_size次，累加可得到Ｃ的值
        for (int k = 0; k < BLOCK_SIZE; ++k)
        {
            cSub += As[ty][k]*Bs[k][tx];　
        }

        __syncthreads();
    }

    //全局地址，向全局寄存器写回去
    //一个线程负责一个元素，一个block负责一个子块
    int cIndex = (by*BLOCK_SIZE + ty)*_wb + (bx*BLOCK_SIZE + tx);
    _C[cIndex] = cSub;
}

4.6 常量内存(constant memory)

参考示例：

__constant__ int c1 = 10;  // 声明__constant__ 即可。
__global__ void kernel1(int *d_dst) {
   int tId = threadIdx.x + blockIdx.x * blockDim.x;
   d_dst[tId] += c1;
}

4.7 图像/纹理内存（texture memory）

texture memory 进行图像类数据加载时， warp内的thread访问的数据地址相邻，从而减少带宽的浪费。
texture 在运算之前能进行一些处理（或者说它本身就是运算），比如聚合、映射等。

应用商城

合作伙伴

开发者

支持与服务

了解天翼云

GPU显存的理解

一、GPU显存的基本概念

二、GPU显存的作用

三、系统内存与设备内存

3.1 内存的架构

1.2 传输通道的速度

四、设备内部的存储

4.1 全局内存

4.2 L1/L2缓存

4.3 局部内存(local memory)

4.4 寄存器（register）

4.5 共享内存（shared memory)

4.6 常量内存(constant memory)

4.7 图像/纹理内存（texture memory）

GPU显存的理解

一、GPU显存的基本概念

二、GPU显存的作用

三、系统内存与设备内存

3.1 内存的架构

1.2 传输通道的速度

四、设备内部的存储

4.1 全局内存

4.2 L1/L2缓存

4.3 局部内存(local memory)

4.4 寄存器（register）

4.5 共享内存（shared memory)

4.6 常量内存(constant memory)

4.7 图像/纹理内存（texture memory）

活动

应用商城

合作伙伴

开发者

支持与服务

了解天翼云

GPU显存的理解

一、GPU显存的基本概念

二、GPU显存的作用

三、 系统内存与设备内存

3.1 内存的架构

1.2 传输通道的速度

四、 设备内部的存储

4.1 全局内存

4.2 L1/L2缓存

4.3 局部内存(local memory)

4.4 寄存器（register）

4.5 共享内存（shared memory)

4.6 常量内存(constant memory)

4.7 图像/纹理内存（texture memory）

GPU显存的理解

一、GPU显存的基本概念

二、GPU显存的作用

三、 系统内存与设备内存

3.1 内存的架构

1.2 传输通道的速度

四、 设备内部的存储

4.1 全局内存

4.2 L1/L2缓存

4.3 局部内存(local memory)

4.4 寄存器（register）

4.5 共享内存（shared memory)

4.6 常量内存(constant memory)

4.7 图像/纹理内存（texture memory）

三、系统内存与设备内存

四、设备内部的存储

三、系统内存与设备内存

四、设备内部的存储