Cuda学习：Cuda加速 GPU缓存机制(后附demo) |
您所在的位置：网站首页 › gpu的主要作用是 › Cuda学习：Cuda加速 GPU缓存机制(后附demo) |
Cuda学习：Cuda加速 GPU缓存机制(后附demo)
|
Cuda是NVIDIA专为图形处理单元（GPU）上的同通用计算开发的并行计算平台和编程模型。借助CUDA，开发者能够利用GPU的强大性能显著加速计算应用。 矩阵转置是现性代数中的基本问题，转置意味着每一列与相应的一行互换，我们通过这个例子来深入GPU的内存结构和缓存机制。下图是一个简单的矩阵和它的转置  图1下面代码是在主机实现的单精度浮点值的错位转置法，矩阵在本职来说就是一维数组，假设矩阵存储在一位数组中，数组的长度是 2^{14}*2^{14}(n_{x}=2^{14},n_{y}=2^{14}) ，选择 2^{14} 作为矩阵的宽和高，一方面尽可能增加数据量达到对比明显的目的，尤其是现在显卡性能不断提高，数据量不够很难看出差别；一方面若再增加该尺寸cuda会报cudaErrorMemoryAllocation(2)，即设备内存申请失败。 void transposeHost(float*out,float*in ,const int nx,const int ny){ for(int iy=0;iy图2.矩阵转置在主机端运行耗时14.7s。我们尝试使用GPU的高性能实现加速效果，有以下两种转置核函数来加速运算 1、按行读取按列存储。 2、按列读取按行存储。 //按行读取按列存储 __global__ void transposeDiagonalRow(float* MatA, float* MatC, int nx, int ny) { unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x; unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y; if (ix < nx && iy < ny) { MatC[ix * ny + iy] = MatA[iy * nx + ix]; } } //按列读取按行存储 __global__ void transposeDiagonalCol(float* MatA, float* MatC, int nx, int ny) { unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x; unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y; if (ix < nx && iy < ny) { MatC[iy * nx + ix] = MatA[ix * ny + iy]; } }耗时对比：  表1. 测试显卡型号：NVIDIA GeForce RTX 2060表1 可知，按行读取按列存储的耗时比按列读取按行存储的耗时分别是20.76ms和14.88ms，前者必后者慢40%，差距明显，出现这么悬殊的结果，我们需要深入了解显卡的内存机构和缓存机制。 内存结构 一般来说，应用程序不会在某一时间点访问任意数据或运行任意代码，应用 程序往往遵循局部性原则，这表明它们可以在任意时间点访问相对较小的局部地址空间。有两种不同的局部性： 时间局部性空间局部性时间局部性是一个位置被引用，那么该数据在较短的时间周期内很可能会再次被引用，随着时间流逝，该数据被引用的可能性逐渐降低，空间局部性认为一个内存位置被引用，则附近的位置也可能被引用。现代计算机使用不断改进的低延迟低容量的内存层次结构来优化性能。一个内存层次结构由具有不同延迟、带宽和容量的多级内存组成。通常，随着从处理器到内存延迟的增加，内存的容量也在增加。一个典型的层次结构如图3所示。  图3.底部所示的存储类型通常有如下特点： 更低的每比特位的平均成本更高的容量更高的延迟更少的处理器访问频率CPU和GPU的主存都采用的是DRAM（动态随机存取存储器），而低延迟内存（如CPU一级缓存）使用的则是SRAM（静态随机存取存储器）。内存层次结构中最大且最慢的级别通常使用磁盘或闪存驱动来实现。在这种内存层次结构中，当数据被处理器频繁使用时，该数据保存在低延迟、低容量的存储器中；而当该数据被存储起来以备后用时，数据就存储在高延迟、大容量的存储器中。这种内存层次结构符合大内存低延迟的设想。 CADA内存模型提出了多种可编程内存的类型： 寄存器共享内存本地内存常量内存纹理内存全局内存 图4.图4是这些内存空间的层次结构，每种都有不用的生命域、生命周期和缓存行为。 寄存器（Registers）寄存器是GPU上运行速度最快的内存空间。核函数中声明的一个没有其他修饰的自变量，通常存储在寄存器中。下面的核函数中变量ix、iy、idx就存储在寄存器里。 __global__ void sumMatrixOnGPU2D(float* MatA, float* MatB, float* MatC, int nx, int ny) { unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x; unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y; unsigned int idx = iy * nx + ix; if (ix < nx && iy < ny) { MatC[idx] = MatA[idx] + MatB[idx]; } }本地内存(Local Memory)核函数中符合存储在寄存器中但不能进入被该核函数分配的寄存器空间中的变量将溢 出到本地内存中。编译器可能存在本地内存中的变量有： 1、在编译时使用未知索引引用的本地数组 2、可能会占用大量寄存器空间的较大本地结构体或数组 3、任何不满足核函数寄存器限定条件的变量 共享内存(Shared Memory)在核函数中使用如下修饰符修饰的变量存放在共享内存中 __shared__，因为共享内存是片上内存，所以与本地内存或全局内存相比，它具有更高的带宽和更低的延迟。下面的block求和的代码中 变量shared[32]就存储在共享内存里，共享内存对block内的所有线程都可见。 __inline__ __device__ int warpReduceSum(int val) { for (int offset = warpSize/2; offset > 0; offset /= 2) val += __shfl_down(val, offset); return val; } __inline__ __device__ int blockReduceSum(int val) { __shared__ int shared[32]; // Shared mem for 32 partial sums int lane = threadIdx.x % warpSize; int wid = threadIdx.x / warpSize; val = warpReduceSum(val); // Each warp performs partial reduction if (lane==0) shared[wid]=val; // Write reduced value to shared memory __syncthreads(); // Wait for all partial reductions //read from shared memory only if that warp existed val = (threadIdx.x < blockDim.x / warpSize) ? shared[lane] : 0; if (wid==0) val = warpReduceSum(val); //Final reduce within first warp return val; }常量内存（Constant Memory）常量内存存储在设备内存中，并在每个SM专用的常量缓存中缓存。常量变量用__constant__修饰，常量变量必须在全局空间内和所有核函数之外进行声明。常量内存是静态声明的，并对同一编译单元中所有的核函数可见。核函数只能从常量内存中读取数据。因此，常量内存必须在主机端使用下面的函数来初始化： cudaMemcpyToSymbol(const void*symbol, const void* src, size_t count);纹理内存（Texture Memory）纹理内存驻留在设备内存中，并在每个SM的只读缓存中缓存。纹理内存是一种通过指定的只读缓存访问的全局内存。 全局内存（Global Memory）全局内存是GPU中最大、延迟最高并且最常使用的内存。global指的是其作用域和生命周期。它的声明可以在任何SM设备上被访问到，并且贯穿应用程序的整个生命周期。一个全局内存变量可以被静态声明或动态声明。 cudaMalloc就是常见的在设备端的全局内存中申请内存。 //malloc device global memory cudaMalloc(void**devPtr, size_t size); //transfer data from host to device cudaMemcpy(void*dst, const void*src, size_t size, cudaMemcpyHostToDevice);下表总结了各类存储器的主要特征。  表2.GPU缓存 在GPU上有4中缓存： 一级缓存二级缓存只读常量缓存只读纹理缓存每个SM都有一个一级缓存，所有的SM共享一个二级缓存。一级和二级缓存都被用来在存储本地内存和全局内存中的数据，也包括寄存器溢出的部分。在CPU上，内存的加载和存储都可以被缓存。但是，在GPU上只有内存加载操作可以被缓存，内存存储操作不能被缓存。 如图5所示，全局内存通过缓存实现加载和存储。全局内存是一个逻辑内存空间，可以通过核函数访问它。核函数的内存请求通常是DRAM设备和芯片上内存之间以128或32字节交换的（该数据来自书籍《cuda-programming》仅供参考）。  图5.所有对全局内存的访问都会通过二级缓存，也有许多访问通过一级缓存，这取决于访问类型和GPU架构。全局内存加载请求首先尝试通过一级缓存，如果一级缓存缺失，该请求转向二级缓存。如果二级缓存缺失，则请求由DRAM完成。内存的存储操作相对简单。一级缓存不能用在Fermi或Kepler GPU上进行存储操作，在发送到设备内存之前存储操作只通过二级缓存。 热知识： CPU一级缓存优化了时间和空间局部性。GPU一级缓存是专为空间局部性而不是为时间局部性设计的。频繁访问一个一级缓存的内存位置不会增加数据留在缓存中的概率。  图6.按行读取按列存储 图7.按列读取按行存储图6所示为按行读取按列存储方法，图7所示为按列读取按行存储方法。按列读取操作将这次访问中没有使用到的数据留在一级缓存（图7左侧灰色横条纹），意味着下一个读操作可能会在缓存上执行而不在全局内存上执行。因为写操作不在一级缓存中缓存，所以对按列执行写操作的例子而言，任何缓存都没有意义。 结论：按列读取和按行存储可以获得更好的性能，这个性能优化的源头是一级缓存的使用。 代码地址： |
【本文地址】
今日新闻 |
推荐新闻 |