企业官网建设流程全解析

1 资料

在本章第一篇介绍的五种设备内存中，共享内存是一种可直接被程序员操作的缓存。主要作用有两个，一是缓存全局内存数据，减少核函数访问全局内存的次数，实现高效的线程块内部通信。二是提高全局内存访问的合并度。
本文将以矩阵转置和数组归约为例讲解共享内存的合理使用。其中数组归约是 CUDA 编程真正意义上的 Hello World，理解了数组归约，就基本掌握了 CUDA 加速的核心思想，面试必考。
本文参考资料如下：
（1）樊哲勇：CUDA编程 基础与实践 第六，七，八，十二章
（2）cuda-programming-guide 1.2 编程模型
（3）cuda-programming-guide 2.1 CUDA C++ 入门
（4）cuda-programming-guide 2.2 编写 CUDA SIMT 内核
（5）cuda-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，请关注代码中的注释部分。

#include"cstdio"#include"cuda_error.cuh"typedeffloatreal;constintNUM_REPEATS=10;constintTILE_DIM=32;// 利用共享内存实现对全局内存的读写都是合并的，在我的机器上，平均运行时间是 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];intix=blockIdx.x*TILE_DIM+threadIdx.x;intiy=blockIdx.y*TILE_DIM+threadIdx.y;// 这里是将每个线程块负责的一片子矩阵数据，从全局数据总拷贝到当前线程块的共享内存中，// 此时对全局内存 A 的读是合并的。if(ix<N&&iy<N){S[threadIdx.y][threadIdx.x]=A[iy*N+ix];}// 调用 __syncthreads ，确保拷贝过程不被干扰// 一般来说，在利用共享内存中的数据之前，都要进行线程块内的同步操作，以确保共享内存数组中的所有元素都已经更新完毕。__syncthreads();// 借助共享内存，这里对全局内存 B 的写，也可以使用合并的方式if(ix<N&&iy<N){B[iy*N+ix]=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_DIM+1];intix=blockIdx.x*TILE_DIM+threadIdx.x;intiy=blockIdx.y*TILE_DIM+threadIdx.y;if(ix<N&&iy<N){S[threadIdx.y][threadIdx.x]=A[iy*N+ix];}__syncthreads();if(ix<N&&iy<N){B[iy*N+ix]=S[threadIdx.x][threadIdx.y];}}voidtiming(constreal*d_A,real*d_B,constintN,constinttask){constintgrid_size_x=(N+TILE_DIM-1)/TILE_DIM;constintgrid_size_y=grid_size_x;constdim3block_size(TILE_DIM,TILE_DIM);constdim3grid_size(grid_size_x,grid_size_y);floatt_sum=0;floatt2_sum=0;for(intrepeat=0;repeat<=NUM_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:transpose1<<<grid_size,block_size>>>(d_A,d_B,N);break;case1:transpose2<<<grid_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(repeat>0){t_sum+=elapsed_time;t2_sum+=elapsed_time*elapsed_time;}CHECK_CUDA_CALL(cudaEventDestroy(start));CHECK_CUDA_CALL(cudaEventDestroy(stop));}constfloatt_ave=t_sum/NUM_REPEATS;constfloatt_err=sqrt(t2_sum/NUM_REPEATS-t_ave*t_ave);printf("Time = %g +- %g ms\n",t_ave,t_err);}voidprint_matrix(constintN,constreal*A){for(intiy=0;iy<N;iy++){for(intix=0;ix<N;ix++){printf("%g\t",A[iy*N+ix]);}printf("\n");}}// 利用共享内存改善全局内存的访问模式，使得对全局内存的读和写都是合并的intmain(void){constintN=128;constintN2=N*N;constintM=sizeof(real)*N2;real*h_A=(real*)malloc(M);real*h_B=(real*)malloc(M);for(inti=0;i<N2;++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(N<=8){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;}

编译运行

ycao@Thinkpad-T14:~/cuda_junior$ ./run_demo.sh 05_cuda_memory/bank.cu -------------------- nvcc -arch=sm_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 ms

2.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，展示折半归约的实现方式，关注文中的注释部分。

#include<cstdio>#include<cstdint>#include"cuda_error.cuh"// typedef float real;typedefdoublereal;constintNUM_REPEATS=20;constintN=1e8;constintM=sizeof(real)*N;constintBLOCK_SIZE=128;// 这个核函数只使用全局内存，// 要求数组长度 N 必须被 BLOCK_SIZE 整除，且 BLOCK_SIZE 为 2 的整数次方，比如这里使用的 BLOCK_SIZE 为 128。__global__voidreduce_global(real*d_x,real*d_y){constinttid=threadIdx.x;// 这句等价于：real *x = &d_x[blockDim.x * blockIdx.x];// 他的作用是从整个数组中，取出 BLOCK_SIZE 个元素的首地址，交由一个线程块处理real*x=d_x+blockDim.x*blockIdx.x;// blockDim.x >> 1 等价于 blockDim.x / 2; offset >>= 1 等价于 offset /= 2;// 使用位运算，是因为效率比较高。for(intoffset=blockDim.x>>1;offset>0;offset>>=1){// 这个判断确保随着归约的进行，越来越多的线程空闲下来if(tid<offset){x[tid]+=x[tid+offset];}// 用于同步单个线程块里的线程，只能用于核函数（也包括被核函数调用的设备函数）// 他的作用是保证一个线程块中的所有线程，在执行该语句后面的语句之前，完全执行了该语句前面的语句,// 针对这个例子，该函数确保每次归约涉及的读和写都能顺利走完，不会被线程块的其他线程干扰，避免数据竞争（data race）。__syncthreads();}// 这个函数归约的结果，是将 N 长的数组，变为 N/BLOCK_SIZE 长的数组，而不是直接归约为 1，请注意这一点// 后续的逻辑会将 d_y 拷贝到主机端，使用循环完成最后的归约。if(tid==0){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){constinttid=threadIdx.x;constintn=blockIdx.x*blockDim.x+tid;// 共享内存推荐使用 s_* 前缀，// 下面定义的共享内存数组，会在每一个线程块中保留一份副本，// 虽然数组变量名一致，但每个线程块的副本都不一样，每个线程块操作自己的副本，互不干涉。__shared__ real s_y[BLOCK_SIZE];// 这里是将每个线程块负责的子数组数据，从全局数据总拷贝到当前线程块的共享内存中，从而减少对全局内存的访问// 由于这里有 n < N 的限制，超出则用 0.0 占位，所以对于这个函数，不要求数组长度 N 必须被 BLOCK_SIZE 整除// 调用 __syncthreads，确保线程块内的所有线程在操作共享内存之前，数据准备就绪。s_y[tid]=(n<N)?d_x[n]:0.0;__syncthreads();// 下面的逻辑同 reduce_global()，唯一的区别是对共享内存里的子数组进行归约，// 归约完成后，拷贝到 d_y 中，后续的逻辑会将 d_y 拷贝到主机端，使用循环完成最后的归约。for(intoffset=blockDim.x>>1;offset>0;offset>>=1){if(tid<offset){s_y[tid]+=s_y[tid+offset];}__syncthreads();}if(tid==0){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){constinttid=threadIdx.x;constintn=blockIdx.x*blockDim.x+tid;extern__shared__ real s_y[];s_y[tid]=(n<N)?d_x[n]:0.0;__syncthreads();for(intoffset=blockDim.x>>1;offset>0;offset>>=1){if(tid<offset){s_y[tid]+=s_y[tid+offset];}__syncthreads();}if(tid==0){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=(N+BLOCK_SIZE-1)/BLOCK_SIZE;constintymem=sizeof(real)*grid_size;constintsmem=sizeof(real)*BLOCK_SIZE;real*d_y;CHECK_CUDA_CALL(cudaMalloc(&d_y,ymem));real*h_y=(real*)malloc(ymem);switch(method){case0:reduce_global<<<grid_size,BLOCK_SIZE>>>(d_x,d_y);break;case1:reduce_shared<<<grid_size,BLOCK_SIZE>>>(d_x,d_y);break;case2:reduce_dynamic<<<grid_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 result=0.0;for(inti=0;i<grid_size;++i){result+=h_y[i];}free(h_y);CHECK_CUDA_CALL(cudaFree(d_y));returnresult;}voidtiming(real*h_x,real*d_x,constintmethod){real sum=0;for(intrepeat=0;repeat<=NUM_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);sum=reduce(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(inti=0;i<N;++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;}

编译运行

ycao@Thinkpad-T14:~/cuda_junior$ ./run_demo.sh 05_cuda_memory/reduce2gpu.cu -------------------- nvcc -arch=sm_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.998770

3 总结

本文所有代码都托管在本人的 github 上：cuda_junior。