CUDA C++ 高效入门第五章 -- CUDA 内存组织之共享内存的用法

CUDA C 高效入门第五章 -- CUDA 内存组织之共享内存的用法1 资料2 正文2.1 使用共享内存进一步优化矩阵转置的效率2.2 使用共享内存实现数组归约3 总结1 资料在本章第一篇介绍的五种设备内存中共享内存是一种可直接被程序员操作的缓存。主要作用有两个一是缓存全局内存数据减少核函数访问全局内存的次数实现高效的线程块内部通信。二是提高全局内存访问的合并度。本文将以矩阵转置和数组归约为例讲解共享内存的合理使用。其中数组归约是 CUDA 编程真正意义上的 Hello World理解了数组归约就基本掌握了 CUDA 加速的核心思想面试必考。本文参考资料如下1樊哲勇CUDA编程 基础与实践 第六七八十二章2cuda-programming-guide 1.2 编程模型3cuda-programming-guide 2.1 CUDA C 入门4cuda-programming-guide 2.2 编写 CUDA SIMT 内核5cuda-programming-guide 2.4 统一内存与系统内存本系列博客汇总链接CUDA C 高效入门系列。2 正文2.1 使用共享内存进一步优化矩阵转置的效率1在上一篇文章中 CUDA C 高效入门第五章 – CUDA 内存组织之高效使用全局内存我们利用矩阵转置讲解了合并度对全局内存访问效率的影响。在文中的样例中两种转置核函数要么对矩阵的读是合并的要么对矩阵的写是合并的而无法实现读写同时合并。本文我们进一步优化矩阵转置的效率借助共享内存可以实现对矩阵的读写都是合并的效率最高。2使用共享内存优化矩阵转置思路其实很简单就是将全局内存的数据以合并读的方式缓存到共享内存然后再将共享内存里数据转置后以合并写的方式写回全局内存。不过在使用共享内存优化矩阵转置之前需要讲解共享内存的 bank 概念避免出现 bank 冲突。3共享内存 bank 及 bank 冲突共享内存不是铁板一块而是分为 32 个同等宽度4/8 字节的能被同时访问的内存 bank。bank 的标号从 0 31, 而每个 bank 有很多层每一层有 4/8 个字节绝大多数是 4 字节。一个 32×32的共享内存矩阵x 方向就是 bank 的标号y 方向就是 bank 的层级。一个线程束的 32 个线程如果同时访问不同的 32 个 bank则只触发一次内存事务memory transaction即数据传输data transfer效率最高。如果一个线程束的 32 个线程同时访问同一个 bank 的不同层的数据访问多少层就会出发多少次内存事务即出现 bank 冲突效率较低应当避免。如下图所示。4创建 05_cuda_memory/bank.cu请关注代码中的注释部分。#includecstdio#includecuda_error.cuhtypedeffloatreal;constintNUM_REPEATS10;constintTILE_DIM32;// 利用共享内存实现对全局内存的读写都是合并的在我的机器上平均运行时间是 0.0150368 ms。// 优于 memory6gobal.cu 的 transpose1 的 0.0213312 ms但比 memory6gobal.cu 的 transpose2 的 0.010992 ms 要慢说明还有优化空间__global__voidtranspose1(constreal*A,real*B,constintN){// 一片tile是 32×32每个线程块处理一片线程块的大小也是 32×32// 这里为每个线程块申请同样 32×32 大小的共享内存__shared__ real S[TILE_DIM][TILE_DIM];intixblockIdx.x*TILE_DIMthreadIdx.x;intiyblockIdx.y*TILE_DIMthreadIdx.y;// 这里是将每个线程块负责的一片子矩阵数据从全局数据总拷贝到当前线程块的共享内存中// 此时对全局内存 A 的读是合并的。if(ixNiyN){S[threadIdx.y][threadIdx.x]A[iy*Nix];}// 调用 __syncthreads 确保拷贝过程不被干扰// 一般来说在利用共享内存中的数据之前都要进行线程块内的同步操作以确保共享内存数组中的所有元素都已经更新完毕。__syncthreads();// 借助共享内存这里对全局内存 B 的写也可以使用合并的方式if(ixNiyN){B[iy*Nix]S[threadIdx.x][threadIdx.y];}}// 这个例子利用共享内存实现对全局内存的读写都是合并的且避免了 bank 冲突在我的机器上平均运行时间是 0.0088352 ms性能最好。__global__voidtranspose2(constreal*A,real*B,constintN){// transpose1 中共享内存矩阵 S 为 32×32// 写的时候线程束的每个线程访问不同的 bank没有 bank 冲突// 但读的时候线程束的每个线程访问同一个 bank 的不同层数据导致严重的 bank 冲突使得加速效果不明显// 这里将共享内存矩阵 S 改为 32×33于是矩阵的每一行的第一个元素在 bank 中的位置是错开的从而避免了共享内存的读取时的冲突__shared__ real S[TILE_DIM][TILE_DIM1];intixblockIdx.x*TILE_DIMthreadIdx.x;intiyblockIdx.y*TILE_DIMthreadIdx.y;if(ixNiyN){S[threadIdx.y][threadIdx.x]A[iy*Nix];}__syncthreads();if(ixNiyN){B[iy*Nix]S[threadIdx.x][threadIdx.y];}}voidtiming(constreal*d_A,real*d_B,constintN,constinttask){constintgrid_size_x(NTILE_DIM-1)/TILE_DIM;constintgrid_size_ygrid_size_x;constdim3block_size(TILE_DIM,TILE_DIM);constdim3grid_size(grid_size_x,grid_size_y);floatt_sum0;floatt2_sum0;for(intrepeat0;repeatNUM_REPEATS;repeat){cudaEvent_t start,stop;CHECK_CUDA_CALL(cudaEventCreate(start));CHECK_CUDA_CALL(cudaEventCreate(stop));CHECK_CUDA_CALL(cudaEventRecord(start));cudaEventQuery(start);switch(task){case0:transpose1grid_size,block_size(d_A,d_B,N);break;case1:transpose2grid_size,block_size(d_A,d_B,N);break;default:printf(Error: wrong task\n);exit(1);break;}CHECK_CUDA_CALL(cudaEventRecord(stop));CHECK_CUDA_CALL(cudaEventSynchronize(stop));floatelapsed_time;CHECK_CUDA_CALL(cudaEventElapsedTime(elapsed_time,start,stop));printf(Time %g ms\n,elapsed_time);if(repeat0){t_sumelapsed_time;t2_sumelapsed_time*elapsed_time;}CHECK_CUDA_CALL(cudaEventDestroy(start));CHECK_CUDA_CALL(cudaEventDestroy(stop));}constfloatt_avet_sum/NUM_REPEATS;constfloatt_errsqrt(t2_sum/NUM_REPEATS-t_ave*t_ave);printf(Time %g - %g ms\n,t_ave,t_err);}voidprint_matrix(constintN,constreal*A){for(intiy0;iyN;iy){for(intix0;ixN;ix){printf(%g\t,A[iy*Nix]);}printf(\n);}}// 利用共享内存改善全局内存的访问模式使得对全局内存的读和写都是合并的intmain(void){constintN128;constintN2N*N;constintMsizeof(real)*N2;real*h_A(real*)malloc(M);real*h_B(real*)malloc(M);for(inti0;iN2;i){h_A[i]i;}real*d_A,*d_B;CHECK_CUDA_CALL(cudaMalloc(d_A,M));CHECK_CUDA_CALL(cudaMalloc(d_B,M));CHECK_CUDA_CALL(cudaMemcpy(d_A,h_A,M,cudaMemcpyHostToDevice));printf(\ntranspose with shared memory bank conflict: \n);timing(d_A,d_B,N,0);printf(\ntranspose without shared memory bank conflict: \n);timing(d_A,d_B,N,1);CHECK_CUDA_CALL(cudaMemcpy(h_B,d_B,M,cudaMemcpyDeviceToHost));if(N8){printf(\nA \n);print_matrix(N,h_A);printf(\nB \n);print_matrix(N,h_B);}free(h_A);free(h_B);CHECK_CUDA_CALL(cudaFree(d_A));CHECK_CUDA_CALL(cudaFree(d_B));return0;}编译运行ycaoThinkpad-T14:~/cuda_junior$ ./run_demo.sh 05_cuda_memory/bank.cu -------------------- nvcc -archsm_75 -o /tmp/tmp.a02blXBWbf/run_demo_exec 05_cuda_memory/bank.cu -------------------- transpose with shared memory bank conflict: Time 0.01744 ms Time 0.011712 ms Time 0.02 ms Time 0.011616 ms Time 0.012288 ms Time 0.012096 ms Time 0.012288 ms Time 0.024064 ms Time 0.021888 ms Time 0.024544 ms Time 0.011648 ms Time 0.0162144 - 0.00536267 ms transpose without shared memory bank conflict: Time 0.00768 ms Time 0.007648 ms Time 0.007776 ms Time 0.008128 ms Time 0.00704 ms Time 0.007584 ms Time 0.007552 ms Time 0.007584 ms Time 0.008224 ms Time 0.007648 ms Time 0.020032 ms Time 0.0089216 - 0.00371629 ms2.2 使用共享内存实现数组归约1数组归约把数组中的所有元素通过某个操作如加法、求最大值合并成一个值。本文讲解数组求和。2两种归约方式第一使用 CPU 编程的做法是循环遍历;第二使用 GPU 编程的归约算法称之为折半归约binary reduction具体做法是将数组分为 BLOCK_SIZE 部分每部分交由一个线程块处理。对于每个 BLOCK_SIZE 子数组将后半部分的各个元素与前半部分对应的数组元素相加重复此过程最后得到的第一个数组元素就是最初的数组中各个元素的和。这个过程要求对同一个线程块的线程进行同步避免数据竞争。由于不同线程块负责不同的子数组因此线程块之间不需要进行线程同步。归约是 CUDA 编程的 Hello World会同时用到共享内存CUDA 线程同步折半归约算法。理解了归约就基本掌握了 CUDA 加速的核心思想面试必考。3共享内存的两种申请方式静态共享内存由编译时指定共享内存的大小。样例sharedreal s_y[BLOCK_SIZE];动态共享内存当调用核函数时指定共享内存大小。申请动态共享内存需要加 extern 关键字使用 []里面不填共享内存数组的长度。调用这个核函数时由 grid_size, block_size, sizeof(real) * block_size 的第三个参数指定共享内存大小样例externsharedreal s_y[];注意动态共享内存和静态共享内存没有性能的区别4创建 05_cuda_memory/reduce2gpu.cu展示折半归约的实现方式关注文中的注释部分。#includecstdio#includecstdint#includecuda_error.cuh// typedef float real;typedefdoublereal;constintNUM_REPEATS20;constintN1e8;constintMsizeof(real)*N;constintBLOCK_SIZE128;// 这个核函数只使用全局内存// 要求数组长度 N 必须被 BLOCK_SIZE 整除且 BLOCK_SIZE 为 2 的整数次方比如这里使用的 BLOCK_SIZE 为 128。__global__voidreduce_global(real*d_x,real*d_y){constinttidthreadIdx.x;// 这句等价于real *x d_x[blockDim.x * blockIdx.x];// 他的作用是从整个数组中取出 BLOCK_SIZE 个元素的首地址交由一个线程块处理real*xd_xblockDim.x*blockIdx.x;// blockDim.x 1 等价于 blockDim.x / 2; offset 1 等价于 offset / 2;// 使用位运算是因为效率比较高。for(intoffsetblockDim.x1;offset0;offset1){// 这个判断确保随着归约的进行越来越多的线程空闲下来if(tidoffset){x[tid]x[tidoffset];}// 用于同步单个线程块里的线程只能用于核函数也包括被核函数调用的设备函数// 他的作用是保证一个线程块中的所有线程在执行该语句后面的语句之前完全执行了该语句前面的语句,// 针对这个例子该函数确保每次归约涉及的读和写都能顺利走完不会被线程块的其他线程干扰避免数据竞争data race。__syncthreads();}// 这个函数归约的结果是将 N 长的数组变为 N/BLOCK_SIZE 长的数组而不是直接归约为 1请注意这一点// 后续的逻辑会将 d_y 拷贝到主机端使用循环完成最后的归约。if(tid0){d_y[blockIdx.x]x[0];}}// 这个核函数使用共享内存优化全局内存访问并优化 N 必须被 BLOCK_SIZE 整除的限制// 他不要求数组长度 N 必须被 BLOCK_SIZE 整除但 BLOCK_SIZE 必须为 2 的整数次方比如这里使用的 BLOCK_SIZE 为 128。// 提示在我的机器上reduce_global 和 reduce_shared 性能上并没有明显的区别// 因为使用共享内存本身也有别的开销比如多一次拷贝多一些 sync 操作// 一般来说在核函数中对共享内存访问的次数越多则使用共享内存带来的加速效果越明显。__global__voidreduce_shared(real*d_x,real*d_y){constinttidthreadIdx.x;constintnblockIdx.x*blockDim.xtid;// 共享内存推荐使用 s_* 前缀// 下面定义的共享内存数组会在每一个线程块中保留一份副本// 虽然数组变量名一致但每个线程块的副本都不一样每个线程块操作自己的副本互不干涉。__shared__ real s_y[BLOCK_SIZE];// 这里是将每个线程块负责的子数组数据从全局数据总拷贝到当前线程块的共享内存中从而减少对全局内存的访问// 由于这里有 n N 的限制超出则用 0.0 占位所以对于这个函数不要求数组长度 N 必须被 BLOCK_SIZE 整除// 调用 __syncthreads确保线程块内的所有线程在操作共享内存之前数据准备就绪。s_y[tid](nN)?d_x[n]:0.0;__syncthreads();// 下面的逻辑同 reduce_global()唯一的区别是对共享内存里的子数组进行归约// 归约完成后拷贝到 d_y 中后续的逻辑会将 d_y 拷贝到主机端使用循环完成最后的归约。for(intoffsetblockDim.x1;offset0;offset1){if(tidoffset){s_y[tid]s_y[tidoffset];}__syncthreads();}if(tid0){d_y[blockIdx.x]s_y[0];}}// 这里的核函数也不要求数组长度 N 必须被 BLOCK_SIZE 整除但 BLOCK_SIZE 必须为 2 的整数次方比如这里使用的 BLOCK_SIZE 为 128。__global__voidreduce_dynamic(real*d_x,real*d_y){constinttidthreadIdx.x;constintnblockIdx.x*blockDim.xtid;extern__shared__ real s_y[];s_y[tid](nN)?d_x[n]:0.0;__syncthreads();for(intoffsetblockDim.x1;offset0;offset1){if(tidoffset){s_y[tid]s_y[tidoffset];}__syncthreads();}if(tid0){d_y[blockIdx.x]s_y[0];}}realreduce(real*d_x,constintmethod){// 等价于 (N % BLOCK_SIZE 0) ? (N / BLOCK_SIZE) : (N / BLOCK_SIZE 1);// 也等价于(N - 1) / BLOCK_SIZE 1;intgrid_size(NBLOCK_SIZE-1)/BLOCK_SIZE;constintymemsizeof(real)*grid_size;constintsmemsizeof(real)*BLOCK_SIZE;real*d_y;CHECK_CUDA_CALL(cudaMalloc(d_y,ymem));real*h_y(real*)malloc(ymem);switch(method){case0:reduce_globalgrid_size,BLOCK_SIZE(d_x,d_y);break;case1:reduce_sharedgrid_size,BLOCK_SIZE(d_x,d_y);break;case2:reduce_dynamicgrid_size,BLOCK_SIZE,smem(d_x,d_y);break;default:printf(Error: wrong method\n);exit(1);break;}CHECK_CUDA_CALL(cudaMemcpy(h_y,d_y,ymem,cudaMemcpyDeviceToHost));// 在主机端完成最后的归约real result0.0;for(inti0;igrid_size;i){resulth_y[i];}free(h_y);CHECK_CUDA_CALL(cudaFree(d_y));returnresult;}voidtiming(real*h_x,real*d_x,constintmethod){real sum0;for(intrepeat0;repeatNUM_REPEATS;repeat){CHECK_CUDA_CALL(cudaMemcpy(d_x,h_x,M,cudaMemcpyHostToDevice));cudaEvent_t start,stop;CHECK_CUDA_CALL(cudaEventCreate(start));CHECK_CUDA_CALL(cudaEventCreate(stop));CHECK_CUDA_CALL(cudaEventRecord(start));cudaEventQuery(start);sumreduce(d_x,method);CHECK_CUDA_CALL(cudaEventRecord(stop));CHECK_CUDA_CALL(cudaEventSynchronize(stop));floatelapsed_time;CHECK_CUDA_CALL(cudaEventElapsedTime(elapsed_time,start,stop));printf(Time %g ms\n,elapsed_time);CHECK_CUDA_CALL(cudaEventDestroy(start));CHECK_CUDA_CALL(cudaEventDestroy(stop));}printf(sum %f\n,sum);}intmain(void){real*h_x(real*)malloc(M);for(inti0;iN;i){h_x[i]1.23;}real*d_x;CHECK_CUDA_CALL(cudaMalloc(d_x,M));printf(Using global memory only: \n);timing(h_x,d_x,0);printf(\nUsing static shared memory: \n);timing(h_x,d_x,1);printf(\nUsing dynamic shared memory: \n);timing(h_x,d_x,2);free(h_x);CHECK_CUDA_CALL(cudaFree(d_x));return0;}编译运行ycaoThinkpad-T14:~/cuda_junior$ ./run_demo.sh 05_cuda_memory/reduce2gpu.cu -------------------- nvcc -archsm_75 -o /tmp/tmp.42dHAAAWHz/run_demo_exec 05_cuda_memory/reduce2gpu.cu -------------------- Using global memory only: Time 25.918 ms Time 20.7884 ms Time 20.2744 ms Time 20.1758 ms Time 20.2002 ms Time 20.2006 ms Time 20.1852 ms Time 20.1754 ms Time 20.1991 ms Time 20.1588 ms Time 20.1999 ms Time 20.1513 ms Time 20.2177 ms Time 20.1412 ms Time 20.1749 ms Time 20.158 ms Time 20.1556 ms Time 20.2068 ms Time 20.2148 ms Time 20.144 ms Time 20.1867 ms sum 122999999.998770 Using static shared memory: Time 22.9675 ms Time 23.0166 ms Time 22.952 ms Time 22.9704 ms Time 22.9463 ms Time 22.9967 ms Time 22.9821 ms Time 22.9679 ms Time 22.9468 ms Time 22.9857 ms Time 22.9437 ms Time 22.9533 ms Time 22.9897 ms Time 22.983 ms Time 22.9381 ms Time 22.9978 ms Time 23.0136 ms Time 22.9995 ms Time 22.9832 ms Time 23.0116 ms Time 23.0163 ms sum 122999999.998770 Using dynamic shared memory: Time 22.9842 ms Time 22.9839 ms Time 22.9707 ms Time 22.9874 ms Time 22.9704 ms Time 22.9795 ms Time 22.9688 ms Time 22.9999 ms Time 22.9396 ms Time 22.985 ms Time 22.9873 ms Time 23.0034 ms Time 22.9637 ms Time 22.9838 ms Time 22.967 ms Time 23.0415 ms Time 23.0172 ms Time 22.9965 ms Time 23.1517 ms Time 22.9868 ms Time 22.9645 ms sum 122999999.9987703 总结本文所有代码都托管在本人的 github 上cuda_junior。