您的位置: 首页  >  文章  >  CUDA编程指南阅读笔记（六）

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

分类: 文章 • 2024-02-04 14:41:10

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编程指南阅读笔记（六）

下面是矩阵乘法的CUDA C主要实现代码：
[cpp]view plaincopy
  1. //Matricesarestoredinrow-majororder:
  2. //M(row,col)=*(M.elements+row*M.width+col)
  3. typedefstruct{
  4. intwidth;
  5. intheight;
  6. float*elements;
  7. }Matrix;
  9. //Threadblocksize
  10. #defineBLOCK_SIZE16
  12. //Forwarddeclarationofthematrixmultiplicationkernel
  13. __global__voidMatMulKernel(constMatrix,constMatrix,Matrix);
  15. //Matrixmultiplication-Hostcode
  16. //MatrixdimensionsareassumedtobemultiplesofBLOCK_SIZE
  17. voidMatMul(constMatrixA,constMatrixB,MatrixC){
  18. //LoadAandBtodevicememory
  19. Matrixd_A;
  20. d_A.width=A.width;d_A.height=A.height;
  21. size_tsize=A.width*A.height*sizeof(float);
  22. cudaMalloc(&d_A.elements,size);
  23. cudaMemcpy(d_A.elements,A.elements,size,cudaMemcpyHostToDevice);
  24. Matrixd_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);
  30. //AllocateCindevicememory
  31. Matrixd_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);
  36. //Invokekernel
  37. dim3dimBlock(BLOCK_SIZE,BLOCK_SIZE);
  38. dim3dimGrid(B.width/dimBlock.x,A.height/dimBlock.y);
  39. MatMulKernel<<<dimGrid,dimBlock>>>(d_A,d_B,d_C);
  41. //ReadCfromdevicememory
  42. cudaMemcpy(C.elements,d_c.elements,size,cudaMemcpyDeviceToHost);
  44. //Freedevicememory
  45. cudaFree(d_A.elements);
  46. cudaFree(d_B.elements);
  47. cudaFree(d_C.elements);
  48. }
  50. //MatrixmultiplicationkernelcalledbyMatMul()
  51. __global__voidMatMulKernel(MatrixA,MatrixB,MatrixC){
  52. //EachthreadcomputesoneelementofC
  53. //byaccumulatingresultsintoCvalue
  54. floatCvalue=0;
  55. introw=blockIdx.y*blockDim.y+threadIdx.y;
  56. intcol=blockIdx.x*blockDim.x+threadIdx.xl
  57. for(inte=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，很明显，这为我们节省了非常多的全局内存带宽。优化后的矩阵计算示意图如下所示：
CUDA编程指南阅读笔记（六）
为了提升计算效率，我们为类型Matrix增加了一个成员变量stride。__device__函数用来获得和设置子矩阵的元素。下面是优化后的代码：
[cpp]view plaincopy
  1. //Matricesarestoredinrow-majororder;
  2. //M(row,col)=*(M.elements+row*M.stride+col)
  3. typedefstruct{
  4. intwidth;
  5. intheight;
  6. intstride;
  7. float*elements;
  8. }Matrix;
  10. //Getamatrixelement
  11. __device__floatGetElement(constMatrixA,introw,intcol){
  12. returnA.elements[row*A.stride+col];
  13. }
  15. //Setamatrixelement
  16. __device__voidSetElement(MatrixA,introw,intcol,floatvalue){
  17. A.elements[row*A.stride+col]=value;
  18. }
  20. //GettheBLOCK_SIZExBLOCK_SIZEsub-matrixAsubofAthatis
  21. //locatedcolsub-matricestotherightandrowsub-matricesdown
  22. //fromtheupper-leftcornerofA
  23. __device__MatrixGetSubMatrix(MatrixA,introw,intcol){
  24. MatrixAsub;
  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. returnAsub;
  30. }
  32. //Threadblocksize
  33. #defineBLOCK_SIZE16
  35. //Forwarddeclarationofthematrixmultiplicationkernel
  36. __global__voidMatMulKernel(constMatrix,constMatrix,Matrix);
  38. //Matrixmultiplication-Hostcode
  39. //MatrixdimensionsareassumedtobemultiplesofBLOCK_SIZE
  40. voidMatMul(constMatrixA,constMatrixB,MatrixC){
  41. //LoadAandBtodevicememory
  42. Matrixd_A;
  43. d_A.width=d_A.stride=A.width;
  44. d_A.height=A.height;
  45. size_tsize=A.width*A.height*sizeof(float);
  46. cudaMalloc(&d_A.elements,size);
  47. cudaMemcpy(d_A.elements,A.elements,size,cudaMemcpyHostToDevice);
  48. Matrixd_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);
  55. //AllocateCindevicememory
  56. Matrixd_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);
  62. //Invokekernel
  63. dim3dimBlock(BLOCK_SIZE,BLOCK_SIZE);
  64. dim3dimGrid(B.width/dimBlock.x,A.height/dimBlock.y);
  65. MatMulKernel<<<dimGrid,dimBlock>>>(d_A,d_B,d_C);
  67. //ReadCfromdevicememory
  68. cudaMemcpy(C.elements,d_C.elements,size,cudaMemcpyDeviceToHost);
  70. //Freedevicememory
  71. cudaFree(d_A.elements);
  72. cudaFree(d_B.elements);
  73. cudaFree(d_C.elements);
  74. }
  76. //MatrixmultiplicationkernelcalledbyMatMul()
  77. __global__voidMatMulKernel(MatrixA,MatrixB,MatrixC){
  78. //Blockrowandcolumn
  79. intblockRow=blockIdx.y;
  80. intblockCol=blockIdx.x;
  82. //Eachthreadblockcomputesonesub-matrixCsubofC
  83. MatrixCsub=GetSubMatrix(C,blockRow,blockCol);
  85. //EachthreadcomputesoneelementofCsub
  86. //byaccumulatingresultsintoCvalue
  87. floatCvalue=0;
  89. //ThreadrowandcolumnwithinCsub
  90. introw=threadIdx.y;
  91. intcol=threadIdx.x;
  93. //Lookoverallthesub-matricesofAandBthatarerequiredtocomputeCsub
  94. //Multiplyeachpairofsub-matricestogetherandaccumulatetheresults
  95. for(intm=0;m<(A.width/BLOCK_SIZE);++m){
  96. //Getsub-matrixAsubofA
  97. MatrixAsub=GetSubMatrix(A,blockRow,m);
  99. //Getsub-matrixBsubofB
  100. MatrixBsub=GetSubMatrix(B,m,blockCol);
  102. //SharedmemoryusedtostoreAsubandBsubrespectively
  103. __shared__floatAs[BLOCK_SIZE][BLOCK_SIZE];
  104. __shared__floatBs[BLOCK_SIZE][BLOCK_SIZE];
  106. //LoadAsubandBsubfromdevicememorytosharedmemory
  107. //Eachthreadloadsoneelementofeachsub-matrix
  108. As[row][col]=GetElement(Asub,row,col);
  109. Bs[row][col]=GetElement(Bsub,row,col);
  111. //Synchronizetomakesurethesub-matricesareloaded
  112. //beforestartingthecomputation
  113. __syncthreads();
  115. //MultiplyAsubandBsubtogether
  116. for(inte=0;e<BLOCK_SIZE;++e)
  117. Cvalue+=As[row][e]*Bs[e][col];
  119. //Synchronizetomakesurethattheprecedingcomputationisdonebefore
  120. //loadingtwonewsub-matricesofAandBinthenextiteration
  121. __syncthreads();
  122. }
  124. //WriteCsubtodevicememory
  125. //Eachthreadwritesoneelement
  126. SetElement(Csub,row,col,Cvalue);
  127. }

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

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

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编程指南阅读笔记（六）

下面是矩阵乘法的CUDA C主要实现代码：
[cpp]view plaincopy
  1. //Matricesarestoredinrow-majororder:
  2. //M(row,col)=*(M.elements+row*M.width+col)
  3. typedefstruct{
  4. intwidth;
  5. intheight;
  6. float*elements;
  7. }Matrix;
  9. //Threadblocksize
  10. #defineBLOCK_SIZE16
  12. //Forwarddeclarationofthematrixmultiplicationkernel
  13. __global__voidMatMulKernel(constMatrix,constMatrix,Matrix);
  15. //Matrixmultiplication-Hostcode
  16. //MatrixdimensionsareassumedtobemultiplesofBLOCK_SIZE
  17. voidMatMul(constMatrixA,constMatrixB,MatrixC){
  18. //LoadAandBtodevicememory
  19. Matrixd_A;
  20. d_A.width=A.width;d_A.height=A.height;
  21. size_tsize=A.width*A.height*sizeof(float);
  22. cudaMalloc(&d_A.elements,size);
  23. cudaMemcpy(d_A.elements,A.elements,size,cudaMemcpyHostToDevice);
  24. Matrixd_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);
  30. //AllocateCindevicememory
  31. Matrixd_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);
  36. //Invokekernel
  37. dim3dimBlock(BLOCK_SIZE,BLOCK_SIZE);
  38. dim3dimGrid(B.width/dimBlock.x,A.height/dimBlock.y);
  39. MatMulKernel<<<dimGrid,dimBlock>>>(d_A,d_B,d_C);
  41. //ReadCfromdevicememory
  42. cudaMemcpy(C.elements,d_c.elements,size,cudaMemcpyDeviceToHost);
  44. //Freedevicememory
  45. cudaFree(d_A.elements);
  46. cudaFree(d_B.elements);
  47. cudaFree(d_C.elements);
  48. }
  50. //MatrixmultiplicationkernelcalledbyMatMul()
  51. __global__voidMatMulKernel(MatrixA,MatrixB,MatrixC){
  52. //EachthreadcomputesoneelementofC
  53. //byaccumulatingresultsintoCvalue
  54. floatCvalue=0;
  55. introw=blockIdx.y*blockDim.y+threadIdx.y;
  56. intcol=blockIdx.x*blockDim.x+threadIdx.xl
  57. for(inte=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，很明显，这为我们节省了非常多的全局内存带宽。优化后的矩阵计算示意图如下所示：
CUDA编程指南阅读笔记（六）
为了提升计算效率，我们为类型Matrix增加了一个成员变量stride。__device__函数用来获得和设置子矩阵的元素。下面是优化后的代码：
[cpp]view plaincopy
  1. //Matricesarestoredinrow-majororder;
  2. //M(row,col)=*(M.elements+row*M.stride+col)
  3. typedefstruct{
  4. intwidth;
  5. intheight;
  6. intstride;
  7. float*elements;
  8. }Matrix;
  10. //Getamatrixelement
  11. __device__floatGetElement(constMatrixA,introw,intcol){
  12. returnA.elements[row*A.stride+col];
  13. }
  15. //Setamatrixelement
  16. __device__voidSetElement(MatrixA,introw,intcol,floatvalue){
  17. A.elements[row*A.stride+col]=value;
  18. }
  20. //GettheBLOCK_SIZExBLOCK_SIZEsub-matrixAsubofAthatis
  21. //locatedcolsub-matricestotherightandrowsub-matricesdown
  22. //fromtheupper-leftcornerofA
  23. __device__MatrixGetSubMatrix(MatrixA,introw,intcol){
  24. MatrixAsub;
  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. returnAsub;
  30. }
  32. //Threadblocksize
  33. #defineBLOCK_SIZE16
  35. //Forwarddeclarationofthematrixmultiplicationkernel
  36. __global__voidMatMulKernel(constMatrix,constMatrix,Matrix);
  38. //Matrixmultiplication-Hostcode
  39. //MatrixdimensionsareassumedtobemultiplesofBLOCK_SIZE
  40. voidMatMul(constMatrixA,constMatrixB,MatrixC){
  41. //LoadAandBtodevicememory
  42. Matrixd_A;
  43. d_A.width=d_A.stride=A.width;
  44. d_A.height=A.height;
  45. size_tsize=A.width*A.height*sizeof(float);
  46. cudaMalloc(&d_A.elements,size);
  47. cudaMemcpy(d_A.elements,A.elements,size,cudaMemcpyHostToDevice);
  48. Matrixd_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);
  55. //AllocateCindevicememory
  56. Matrixd_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);
  62. //Invokekernel
  63. dim3dimBlock(BLOCK_SIZE,BLOCK_SIZE);
  64. dim3dimGrid(B.width/dimBlock.x,A.height/dimBlock.y);
  65. MatMulKernel<<<dimGrid,dimBlock>>>(d_A,d_B,d_C);
  67. //ReadCfromdevicememory
  68. cudaMemcpy(C.elements,d_C.elements,size,cudaMemcpyDeviceToHost);
  70. //Freedevicememory
  71. cudaFree(d_A.elements);
  72. cudaFree(d_B.elements);
  73. cudaFree(d_C.elements);
  74. }
  76. //MatrixmultiplicationkernelcalledbyMatMul()
  77. __global__voidMatMulKernel(MatrixA,MatrixB,MatrixC){
  78. //Blockrowandcolumn
  79. intblockRow=blockIdx.y;
  80. intblockCol=blockIdx.x;
  82. //Eachthreadblockcomputesonesub-matrixCsubofC
  83. MatrixCsub=GetSubMatrix(C,blockRow,blockCol);
  85. //EachthreadcomputesoneelementofCsub
  86. //byaccumulatingresultsintoCvalue
  87. floatCvalue=0;
  89. //ThreadrowandcolumnwithinCsub
  90. introw=threadIdx.y;
  91. intcol=threadIdx.x;
  93. //Lookoverallthesub-matricesofAandBthatarerequiredtocomputeCsub
  94. //Multiplyeachpairofsub-matricestogetherandaccumulatetheresults
  95. for(intm=0;m<(A.width/BLOCK_SIZE);++m){
  96. //Getsub-matrixAsubofA
  97. MatrixAsub=GetSubMatrix(A,blockRow,m);
  99. //Getsub-matrixBsubofB
  100. MatrixBsub=GetSubMatrix(B,m,blockCol);
  102. //SharedmemoryusedtostoreAsubandBsubrespectively
  103. __shared__floatAs[BLOCK_SIZE][BLOCK_SIZE];
  104. __shared__floatBs[BLOCK_SIZE][BLOCK_SIZE];
  106. //LoadAsubandBsubfromdevicememorytosharedmemory
  107. //Eachthreadloadsoneelementofeachsub-matrix
  108. As[row][col]=GetElement(Asub,row,col);
  109. Bs[row][col]=GetElement(Bsub,row,col);
  111. //Synchronizetomakesurethesub-matricesareloaded
  112. //beforestartingthecomputation
  113. __syncthreads();
  115. //MultiplyAsubandBsubtogether
  116. for(inte=0;e<BLOCK_SIZE;++e)
  117. Cvalue+=As[row][e]*Bs[e][col];
  119. //Synchronizetomakesurethattheprecedingcomputationisdonebefore
  120. //loadingtwonewsub-matricesofAandBinthenextiteration
  121. __syncthreads();
  122. }
  124. //WriteCsubtodevicememory
  125. //Eachthreadwritesoneelement
  126. SetElement(Csub,row,col,Cvalue);
  127. }

相关推荐