CUDA C 高效入门第五章 -- CUDA 内存组织之高效使用全局内存1 资料2 正文2.1 全局内存访问的合并度2.2 以矩阵转置为例展示合并访问与非合并访问的性能区别2.2.1 矩阵转置2.2.2 样例分析3 总结1 资料上一篇文章我们系统介绍了各类 GPU 设备内存的特性和简单使用方法。内存的合理使用直接影响程序的性能。而在所有的五种设备内存中全局内存容量最大访问速度最慢往往会成为一个 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内存事务memory transaction对全局内存的访问将触发内存事务也就是数据传输data transfer一次数据传输处理的数据量在默认情况下是 32 字节。当 CUDA 线程从全局内存请求数据时相关的线程束会根据每个线程访问的数据大小以及内存地址在线程间的分布情况将该线程束中所有线程的内存请求合并成若干次数据传输即若干次 32 字节传输。2内存事物对数据地址的要求在一次数据传输中要求从全局内存转移到 L2 缓存的一片内存的首地址一定是一个最小粒度这里是 32 字节的整数倍。而使用 CUDA 运行时 API 函数如我们常用的 cudaMalloc分配的内存的首地址至少是 256 字节的整数倍因此满足数据传输对地址的要求。3合并度degree of coalescing或者叫内存利用率它等于线程束请求的字节数除以由该请求导致的所有数据传输处理的字节数。内存利用率越高核函数中与全局内存访问有关的部分的性能就更好利用率低则意味着对显存带宽的浪费。4好的情况如果线程束中的连续线程请求内存中连续的 4 字节那么该线程束将总共请求 128 字节的内存并且这所需的 128 字节将通过四次 32 字节的内存事务获取。这实现了 100% 的内存利用率。也就是说100% 的内存流量都被该线程束利用了。下图展示了完美合并度的情况。5坏的情况连续的线程访问内存中彼此相距 32 字节或更远的数据元素。在这种情况下线程束将被迫为每个线程发起一次 32 字节的内存事务内存传输的总字节数将是 32 字节 * 32 线程/线程束 1024 字节。然而实际使用的内存量仅为 128 字节线程束中每个线程 4 字节因此内存利用率仅为 128 / 1024 12.5%。这是对内存系统非常低效的使用。下图展示了非合并访问的情况。2.2 以矩阵转置为例展示合并访问与非合并访问的性能区别2.2.1 矩阵转置1矩阵转置是线性代数里的一个基本操作是将矩阵所有元素的位置沿左上-右下轴翻转。样例如下2.2.2 样例分析1我们编程实现矩阵的转置即 05_cuda_memory/memory6global.cu代码中的注释部分介绍了不同合并度对程序性能的影响。#includecstdio#includecuda_error.cuhtypedeffloatreal;constintNUM_REPEATS10;constintTILE_DIM32;// 这里的矩阵拷贝读和写都实现了 100% 的合并度速度最快时间最短。在我的机器上单次拷贝平均耗时0.0068736 ms__global__voidcopy(constreal*A,real*B,constintN){// 核函数是可以直接访问普通全局变量的// 由于 blockDim.x 和 blockDim.y 等于 TILE_DIM因此也可以替换为 blockDim.x 和 blockDim.y。constintixblockIdx.x*TILE_DIMthreadIdx.x;constintiyblockIdx.y*TILE_DIMthreadIdx.y;if(ixNiyN){B[iy*Nix]A[iy*Nix];}}// 这里的矩阵转置读的时候实现了 100% 的合并度而写的时候没有实现合并。在我的机器上单次拷贝平均耗时0.0213312 ms__global__voidtranspose1(constreal*A,real*B,constintN){constintixblockIdx.x*blockDim.xthreadIdx.x;constintiyblockIdx.y*blockDim.ythreadIdx.y;if(ixNiyN){B[ix*Niy]A[iy*Nix];}}// 这里的矩阵转置读的时候没有实现合并而写的时候实现了 100% 的合并度在我的机器上单次拷贝平均耗时0.010992 ms// 这里的耗时比上面的转置更短是因为读虽然没有合并但全局内存的读默认情况下是有缓存机制的而写是没有的请读者注意。// 因此在不能同时满足读取和写入都是合并的情况下一般来说应当尽量做到合并的写入。__global__voidtranspose2(constreal*A,real*B,constintN){constintixblockIdx.x*blockDim.xthreadIdx.x;constintiyblockIdx.y*blockDim.ythreadIdx.y;if(ixNiyN){B[iy*Nix]A[ix*Niy];}}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:copygrid_size,block_size(d_A,d_B,N);break;case1:transpose1grid_size,block_size(d_A,d_B,N);break;case2: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(copy: \n);timing(d_A,d_B,N,0);printf(\ntranspose with coalesced read: \n);timing(d_A,d_B,N,1);printf(\ntranspose with coalesced write: \n);timing(d_A,d_B,N,2);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/memory6global.cu -------------------- nvcc -archsm_75 -o /tmp/tmp.zv3WIq2KZu/run_demo_exec 05_cuda_memory/memory6global.cu -------------------- copy: Time 0.009376 ms Time 0.006848 ms Time 0.006464 ms Time 0.007808 ms Time 0.008064 ms Time 0.006464 ms Time 0.00624 ms Time 0.006464 ms Time 0.006656 ms Time 0.006656 ms Time 0.006816 ms Time 0.006848 - 0.000573329 ms transpose with coalesced read: Time 0.020416 ms Time 0.01824 ms Time 0.02048 ms Time 0.020512 ms Time 0.019424 ms Time 0.02208 ms Time 0.019136 ms Time 0.020448 ms Time 0.019328 ms Time 0.019872 ms Time 0.02048 ms Time 0.02 - 0.000994693 ms transpose with coalesced write: Time 0.01152 ms Time 0.012288 ms Time 0.019296 ms Time 0.010304 ms Time 0.012224 ms Time 0.011712 ms Time 0.012288 ms Time 0.010752 ms Time 0.011904 ms Time 0.01232 ms Time 0.010272 ms Time 0.012336 - 0.00244812 ms3 总结本文所有代码都托管在本人的 github 上cuda_junior。