现在的位置: 首页 > 综合 > 正文

CUDA编程指南阅读笔记（六）

2018年12月11日 ⁄ 综合 ⁄ 共 8371字 ⁄ 字号 小 中 大 ⁄ 评论关闭

4. CUDA C语言编程接口

        接上文。

4.3 CUDA C Runtime

4.3.3 共享内存（Shared Memory）

        共享内存是CUDA设备中非常重要的一个存储区域，有效地使用共享内存可以充分利用CUDA设备的潜能，极大提升程序性能。那么，共享内存有哪些特点呢？
        1、共享内存（shared Memory）是集成在GPU处理器芯片上的（on-chip），因此相比于存在于显存颗粒中的全局内存（global Memory）和本地内存（local Memory），它具有更高的传输带宽，一般情况下，共享内存的带宽大约是全局内存带宽的7-10倍。
        2、共享内存的容量很小。根据NVIDIA官方文档的说法，在计算能力1.x的设备中，每一个流多处理器（Streaming Multiprocessor）上的共享内存容量为16KB。对于计算能力2.x、3.0及3.5的设备该参数为48KB。因此共享内存是稀有资源。
        3、共享内存在物理上被划分为很多块，每一块被称为一个存储体（bank）。在同一时刻，CUDA设备可以同时访问多个存储体。因此，如果一次针对共享内存的访存操作需要读取n个地址，而这n个地址恰好分布在n个不同的存储体（bank）中，那么只需要一个存取周期就可以完成n个地址的访存任务了。对于计算能力1.x的设备，共享内存被平均划分为16个存储体。而对于计算能力2.x、3.0及3.5的设备此参数为32。在共享内存中，相邻两块32bit的数据分别属于相邻的两个存储体。存储体每两个时钟周期可以传输32位数据。
        4、共享内存既可以静态分配，也可以动态分配。
        从共享内存的这些特点中我们可以看出，它实际上相当于一个程序员可以操控的缓存（cache），下面，我们使用矩阵乘法的例子来说明如何有效使用共享内存。
        首先，我们使用最直观的方法来完成矩阵乘法C = A x B：读取A的每一行和B的每一列，顺次完成计算任务。矩阵乘法的示意图如下所示：


下面是矩阵乘法的CUDA C主要实现代码：

[cpp] view
plaincopy

  1. // Matrices are stored in row-major order:  
  2. // M(row, col) = *(M.elements + row * M.width + col)  
  3. typedef struct {  
  4.     int width;  
  5.     int height;  
  6.     float *elements;  
  7. } Matrix;  
  8.   
  9. // Thread block size  
  10. #define BLOCK_SIZE 16  
  11.   
  12. // Forward declaration of the matrix multiplication kernel  
  13. __global__ void MatMulKernel(const Matrix, const Matrix, Matrix);  
  14.   
  15. // Matrix multiplication - Host code  
  16. // Matrix dimensions are assumed to be multiples of BLOCK_SIZE  
  17. void MatMul(const Matrix A, const Matrix B, Matrix C) {  
  18.     // Load A and B to device memory  
  19.     Matrix d_A;  
  20.     d_A.width = A.width; d_A.height = A.height;  
  21.     size_t size = A.width * A.height * sizeof(float);  
  22.     cudaMalloc(&d_A.elements, size);  
  23.     cudaMemcpy(d_A.elements, A.elements, size, cudaMemcpyHostToDevice);  
  24.     Matrix d_B;  
  25.     d_B.width = B.width; d_B.height = B.height;  
  26.     size = B.width * B.height * sizeof(float);  
  27.     cudaMalloc(&d_B.elements, size);  
  28.     cudaMemcpy(d_B.elements, B.elements, size, cudaMemcpyHostToDevice);  
  29.   
  30.     // Allocate C in device memory  
  31.     Matrix d_C;  
  32.     d_C.width = C.width; d_C.height = C.height;  
  33.     size = C.width * C.height * sizeof(float);  
  34.     cudaMalloc(&d_C.elements, size);  
  35.   
  36.     // Invoke kernel  
  37.     dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);  
  38.     dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);  
  39.     MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);  
  40.   
  41.     // Read C from device memory  
  42.     cudaMemcpy(C.elements, d_c.elements, size, cudaMemcpyDeviceToHost);  
  43.   
  44.     // Free device memory  
  45.     cudaFree(d_A.elements);  
  46.     cudaFree(d_B.elements);  
  47.     cudaFree(d_C.elements);  
  48. }  
  49.   
  50. // Matrix multiplication kernel called by MatMul()  
  51. __global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) {  
  52.     // Each thread computes one element of C  
  53.     // by accumulating results into Cvalue  
  54.     float Cvalue = 0;  
  55.     int row  = blockIdx.y * blockDim.y + threadIdx.y;  
  56.     int col = blockIdx.x * blockDim.x + threadIdx.xl  
  57.     for (int e = 0; e < A.width; ++e)  
  58.         Cvalue += A.elements[row * A.width + e] * B.elements[e * B.width + col];  
  59.     C.elements[row * C.width + col] = Cvalue;  
  60. }  

可以看出，为了计算矩阵C的任何一个元素，程序都需要从全局内存（global memory）中获得矩阵A的一行和矩阵B的一列。因此，完成这一计算矩阵A被读取了B.width次，矩阵B被读取了A.height次。
        现在我们来使用共享内存（shared memory）实现矩阵乘法。假设矩阵C可以被划分为若干个较小的子方阵Csub，我们使用一个线程块（thread block）来负责某一子方阵的计算，线程块中的每一个线程（thread）正好负责子方阵Csub中一个元素的计算。这样划分后，任何一个结果子方阵Csub'（尺寸为block_size
* block_size）都是与该方阵具有相同行索引的尺寸为A.width * block_size的A的子矩阵Asub和与该方阵具有相同列索引的尺寸为block_size * B.height的B的子矩阵Bsub相乘所得到。

        为了匹配设备的计算资源，两个子矩阵Asub和Bsub被划分为尽可能多的分离的维度为block_size的子方阵，Csub的值便是这些子矩阵相乘后相加所得到的结果。子矩阵乘法的执行顺序都是首先将它们从全局内存（global
memory）拷贝到共享内存（shared memory）（线程块中的每一个线程正好负责方阵一个元素的拷贝），然后由线程自己完成相应元素的计算任务，利用寄存器存储局部结果，最后将寄存器的内容与新得到的计算结果依此累加起来得到最终运算结果并将其传输到全局内存（global memory）中。
        通过使用这种分治的计算策略，共享内存得到了很好的利用，采用这种方案计算完成时全局内存中矩阵A被访问的次数为B.width / block_size，矩阵B被访问的次数为A.height / block_size，很明显，这为我们节省了非常多的全局内存带宽。优化后的矩阵计算示意图如下所示：

        为了提升计算效率，我们为类型Matrix增加了一个成员变量stride。__device__函数用来获得和设置子矩阵的元素。下面是优化后的代码：

[cpp] view
plaincopy

  1. // Matrices are stored in row-major order;  
  2. // M(row, col) = *(M.elements + row * M.stride + col)  
  3. typedef struct {  
  4.     int width;  
  5.     int height;  
  6.     int stride;  
  7.     float* elements;  
  8. } Matrix;  
  9.   
  10. // Get a matrix element  
  11. __device__ float GetElement(const Matrix A, int row, int col) {  
  12.     return A.elements[row * A.stride + col];  
  13. }  
  14.   
  15. // Set a matrix element  
  16. __device__ void SetElement(Matrix A, int row, int col, float value) {  
  17.     A.elements[row * A.stride + col] = value;  
  18. }  
  19.   
  20. // Get the BLOCK_SIZExBLOCK_SIZE sub-matrix Asub of A that is  
  21. // located col sub-matrices to the right and row sub-matrices down  
  22. // from the upper-left corner of A  
  23. __device__ Matrix GetSubMatrix(Matrix A, int row, int col) {  
  24.     Matrix Asub;  
  25.     Asub.width = BLOCK_SIZE;  
  26.     Asub.height = BLOCK_SIZE;  
  27.     Asub.stride = A.stride;  
  28.     Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row + BLOCK_SIZE * col];  
  29.     return Asub;  
  30. }  
  31.   
  32. // Thread block size  
  33. #define BLOCK_SIZE 16  
  34.   
  35. // Forward declaration of the matrix multiplication kernel  
  36. __global__ void MatMulKernel(const Matrix, const Matrix, Matrix);  
  37.   
  38. // Matrix multiplication - Host code  
  39. // Matrix dimensions are assumed to be multiples of BLOCK_SIZE  
  40. void MatMul(const Matrix A, const Matrix B, Matrix C) {  
  41.     // Load A and B to device memory  
  42.     Matrix d_A;  
  43.     d_A.width = d_A.stride = A.width;  
  44.     d_A.height = A.height;  
  45.     size_t size = A.width * A.height * sizeof(float);  
  46.     cudaMalloc(&d_A.elements, size);  
  47.     cudaMemcpy(d_A.elements, A.elements, size, cudaMemcpyHostToDevice);  
  48.     Matrix d_B;  
  49.     d_B.width = d_B.stride = B.width;  
  50.     d_B.height = B.height;  
  51.     size = B.width * B.height * sizeof(float);  
  52.     cudaMalloc(&d_B.elements, size);  
  53.     cudaMemcpy(d_B.elements, B.elements, size, cudaMemcpyHostToDevice);  
  54.   
  55.     // Allocate C in device memory  
  56.     Matrix d_C;  
  57.     d_C.width = d_C.stride = C.width;  
  58.     d_C.height = C.height;  
  59.     size = C.width * C.height * sizeof(float);  
  60.     cudaMalloc(&d_C.elements, size);  
  61.   
  62.     // Invoke kernel  
  63.     dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);  
  64.     dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);  
  65.     MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);  
  66.   
  67.     // Read C from device memory  
  68.     cudaMemcpy(C.elements, d_C.elements, size, cudaMemcpyDeviceToHost);  
  69.   
  70.     // Free device memory  
  71.     cudaFree(d_A.elements);  
  72.     cudaFree(d_B.elements);  
  73.     cudaFree(d_C.elements);  
  74. }  
  75.   
  76. // Matrix multiplication kernel called by MatMul()  
  77. __global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) {  
  78.     // Block row and column  
  79.     int blockRow = blockIdx.y;  
  80.     int blockCol = blockIdx.x;  
  81.   
  82.     // Each thread block computes one sub-matrix Csub of C  
  83.     Matrix Csub = GetSubMatrix(C, blockRow, blockCol);  
  84.   
  85.     // Each thread computes one element of Csub  
  86.     // by accumulating results into Cvalue  
  87.     float Cvalue = 0;  
  88.   
  89.     // Thread row and column within Csub  
  90.     int row = threadIdx.y;  
  91.     int col = threadIdx.x;  
  92.   
  93.     // Look over all the sub-matrices of A and B that are required to compute Csub  
  94.     // Multiply each pair of sub-matrices together and accumulate the results  
  95.     for (int m = 0; m < (A.width / BLOCK_SIZE); ++m) {  
  96.         // Get sub-matrix Asub of A  
  97.         Matrix Asub = GetSubMatrix(A, blockRow, m);  
  98.           
  99.         // Get sub-matrix Bsub of B  
  100.         Matrix Bsub = GetSubMatrix(B, m, blockCol);  
  101.   
  102.         // Shared memory used to store Asub and Bsub respectively  
  103.         __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];  
  104.         __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];  
  105.   
  106.         // Load Asub and Bsub from device memory to shared memory  
  107.         // Each thread loads one element of each sub-matrix  
  108.         As[row][col] = GetElement(Asub, row, col);  
  109.         Bs[row][col] = GetElement(Bsub, row, col);  
  110.   
  111.         // Synchronize to make sure the sub-matrices are loaded  
  112.         // before starting the computation  
  113.         __syncthreads();  
  114.   
  115.         // Multiply Asub and Bsub together  
  116.         for (int e = 0; e < BLOCK_SIZE; ++e)  
  117.             Cvalue += As[row][e] * Bs[e][col];  
  118.   
  119.         // Synchronize to make sure that the preceding computation is done before  
  120.         // loading two new sub-matrices of A and B in the next iteration  
  121.         __syncthreads();  
  122.     }  
  123.   
  124.     // Write Csub to device memory  
  125.     // Each thread writes one element  
  126.     SetElement(Csub, row, col, Cvalue);  
  127. }

同一个block中的线程在95行的for循环中获取到的Asub，Bsub，Csub是一样的，每个线程就负责Csub内元素的计算

http://blog.csdn.net/csgxy123/article/details/10018531

返回
【上篇】
【下篇】

作者:

抱歉!评论已关闭.

返回首页

Copyright © 2013-2018 学步园  保留所有权利.
软文销售 QQ客服：2265327166 点击这里给我发消息（其他合作也可洽谈）
必威体育
必威电竞