第四章 GPU编程优化技术总结 4.1.0 CUDA设备上的优化技术 4.1.1 访存优化 4.1.2 指令优化 4.1.3 内核调用优化 4.2.0 GCN设备上的优化技术 4.2.1 访存优化 4.2.2 指令优化 4.2.3 内核调用优化 4.3 构建性能可移植的程序 第四章 GPU编程优化技术总结 我们在两个章节分别讲述针对CUDA和GCN这两大目前主流的GPU并行计算的设备。但是诸如合并访问,如何避免共享内存的bank conflicts以及简单的指令优化等基本内容这里不再叙述,有需要的可以参考<<CUDA Programming Guide>>和<<AMD Accelerated Parallel ProcessingOpenCL Programming Guide>>,这里仅给出一些不常见的优化技巧。 4.1.0 CUDA设备的优化技术 4.1.1 访存优化 1 在计算能力为2.0或以上的CUDA设备上,当一个warp内的所有线程访问同一个地址时,可以使用统一加载操作将一个数据通过缓存广播到warp内的所有线程中,从而提升性能。虽然CUDAprogramming guide上提到当访问数据的地址和线程号无关且是只读数据时,编译器会自动使用LDU加载指令,但有时编译器并不能得到我们想要的结果。比如: __global__ void add( float* d_a, const float* d_b, int n) { intwarpid=(blockDim.x>>5)*blockIdx.x+(threadIdx.x>>5); if(warpid>=n)return; d_a[threadIdx.x]+=d_b[warpid]; } 查看PTX代码,我们发现编译器并未使用LDU,因此我们就需要显示的使用内联PTX汇编来达到我们的目的: #if defined(_WIN64)|| defined(__x86_64)||defined(_M_X64)|| defined(_M_IA64)||defined(_M_AMD64) #define PTX_PTR “l” #else #define PTX_PTR “r” #endif __device____forceinline__ float __ldu(const float* p ) { float val; asm volatile("ldu.global.f32 {%0}, [%1];" : "=f"(val) : PTX_PTR(p)); returnval; } 根据测试,在满足使用LDU的情况下的所有设备中均能获得性能提升,即使是在计算能力3.5+的设备上其效果也要略好于使用LDG(使用纹理缓存)。例如卷积神经网络的计算中,当每个通道中对应的是一个标量的偏置值,那么在卷积计算后对通道施加偏置的操作就可以通过LDU操作高效的完成(虽然也可以通过共享内存,但是使用LDU具有更简洁的实现,并具有轻微的性能优势)。 4.1 指令优化 1 对于可以完成相同计算的指令集合,应尽可能选择具有更低延迟以及更高混合比例的指令集合,比如在某些设备上双精度可以和内存加载存储指令双发,但是却无法和单精度以及整数指令双发。 2 同时对同一个数组进行多次等距寻址时尽量将不变的索引在开始处加到数组的基址上,这样可以减少地址的计算或是便于基址+常量寻址,从而减少指令数量。 3 对于存在大量计算的循环中如果某些指令,如数据存取指令无需复杂的寻址计算,那么考虑对每个存储操作使用断定,便于编译器将计算和存储指令混合排列从而利用指令的双发(dualissue)机制。 4 使用某些特定的常量,可以将数据融入指令码中,从而具有更小的代码体积。在kepler和maxwell设备上的32位浮点数和整数的双操作数(输入)指令支持全精度的常量,比如 c=a+128.f 对应的SASS指令为 FADD R2,R0,128, c=a+10007.f对应的SASS指令为 FADD32I R2,R0,10007, 这些立即数会被嵌入指令的编码序列中,但是对于三操作数指令(如FMA)则会将常数放入常量内存的第2个bank中(猜测原因是受限于指令编码的长度,因为多出的一个操作数需要额外的位数表示寄存器索引),因此当一个计算序列中使用FMA不能减少指令数量时(亦即和使用FMUL,FADD数量相同),如果涉及到立即数,则尽量不要使用FMA代替FMUL和FADD,因为操作数直接嵌入指令编码具有更小的代码体积,除非出于精度考虑。如将 temp.x=c*b.x+(-s)*b.y; temp.y=c*b.y+s*b.x; b.x=a.x-temp.x; b.y=a.y-temp.y; a.x+=temp.x; a.y+=temp.y; 转换为 temp.x=b.x; temp.y=b.y; b.x=a.x+(-c)*temp.x+s*temp.y; b.y=a.y+(-c)*temp.y+(-s)*temp.x; a.x+=c*temp.x+(-s)*temp.y; a.y+=c*temp.y+s*temp.x; 并不能减少指令数量,也不会带来性能提升,除非是出于精度考虑。对于双精度数据,有规律的常量也可被嵌入到指令码中,比如0.5,0.25,0.125,0.0625,0.03125,…,1.0,1.5, 1.25,…, 64.0, 128.0, 65536.0, …;但是无规则的常量会被放入常量内存的第2个bank中,比如 c=a+128.0 对应的SASS指令为 DADD R2, R7, 128 c=a+790045.7对应的SASS指令为 DADD R2, R7, c[0x2][0x0] 可以包含在双精度指令码中的常量的具体规则为: … +-512.0 +-256.0 +-128.0,+-128.5 +-64.0,+-64.5, 64.25 +-32.0,+-32.5, +-32.25, +-32.125 +-16.0,+-16.5, +-16.25, +-16.125, +-16.0625 … 4.1.2 分支优化 1 使用小的局部数组消除多分支或是简化复杂的条件代码计算。 2 巧妙的利用位操作和局部数组消除分支,例如第三章中通过局部数组简化了主分割面的选择,而通过巧妙的位操作减少了确定下个待遍历节点的分支。 3 分析算法看是否能将不同的路径分配到不同的warp中或block中,同时保证warp或block中的指令路径相同;或是将问题进行拆分成多个内核进行处理。 4 在多分支结构中将判据按照命中的概率从高到低进行排列。 5 某些情况下使用对函数指针列表的寻址代替switch逻辑已消除对大批量分支判断的遍历,同时可以生成更小的代码。 4.1.3 内核调用优化 通常不是限制性能的地方,但是当很多内核在一个循环中被比较长时间的跨距调用时(因此设备驱动的热身会被过长的间隔抹消或是被其它内核的调用覆盖掉内核参数缓存),尤其是当内核具有很多参数时,每次内核参数都需要从内存到设备上的内核参数缓存的复制过程,有时这也会给效率带来较大影响,这里根据作者经验总结了几个方法来优化内核的启动时间: 1 如果内核参数很多,对于指针类型的参数,考虑合并多个指针变量,并在内核内部解引用,这样做有时也会减轻寄存器压力(但也不要想当然,任何时候都应该试着查看编译后的寄存器使用情况)。例如,假设四个长度均为1024的数组(指针合并并不要求每个指针指向的数组大小一样): __global__ void …( …,const int * d_a, const int * d_b, const int * d_c, const int * d_d, … ) { … d_a+=tidx; d_b+=tidx; d_c+=tidx; d_d+=tidx; } 可以改成如下形式: __global__ void … ( …, const int * d_a, … ) { … d_a+=tidx; const int *d_b=d_a+1024; const int *d_c=d_b+1024; const int *d_d=d_c+1024; /* 或者通过对d_a的常量偏移分别访问各个数组: d_a, d_a+1024,d_a+2048, d_a+3072 */ … } 2 从PTX指令到本地SASS汇编指令并不是严格一一对应的,在这个翻译的过程中ptxas会进行实际的寄存器分配,指令的替换和重排等优化,因此很多时候你无法通过使用PTX达到控制指令执行顺序和寄存器分配的目的(ptxas做的并不够好,一个实际的例子就是对于矩阵乘法,如果想要达到接近峰值的效率,必须直接对SASS指令进行重拍以及对寄存器进行细致的分配以最小化指令计算延迟和寄存器bankconflicts引起的指令流水线停顿。但是很可惜,NVIDIA并未开放本地汇编的编程环境,甚至连SASS ISA的指令编码格式都未公开,因此需要程序员自己绕开种种限制开发自己的第三方GPU汇编器),但是PTX仍然能在一定程度上影响最终得到的SASS结果,这需要开发者耐心的对指令的顺序和逻辑进行调整并观察最终编译出来的SASS代码。 3 创建CUDA上下文时使用CU_CTX_LMEM_RESIZE_TO_MAX标志,以避免那些具有寄存器溢出的内核在下次启动时重新在设备内存上为寄存器溢出分配局部内存,这样会造成当前线程中的CUDA上下文中所包含的的所有流上的数据传输和内核计算操作中断(即使操作是异步的)。 4.2.0 GCN设备上的优化技术 CUDA设备上的分支优化和内核调用优化方法同样可以用在GCN设备上,所以本节不再做重复的叙述。 4.2.1 访存优化 1 虽然GCN设备上一个wavefront对应的连续256字节对齐数据具有最高的传输效率(每个线程4字节),但是当遇到计算密集型的问题时,如第一章中所讲的那样使用宽向量加载和存储操作可以具有更高的效率。 2 GCN设备上的缓存结构并不具备在wavefront线程间的广播机制,因此如果多个线程访问同一个或少数几个数据,更好的方式是通过局部内存,例如: #if(get_local_id(0)<4){ l_data=g_data[get_local_id(0)]; }barrier(CLK_LOCAL_MEM_FENCE); 而不是 data=g_data[get_local_id(0)&3] 3 将不同block内的全局数据访问尽量分散到不同的全局内存channel和bank中,如果多个同时进行全局内存数据访问的不同block访问的数据位于同一个channel或bank中,则内存操作会串行执行,对效率的影响很大,必要时显式的对block进行调度。 4 GCN设备上的共享内存可以不经过寄存器直接访问(有点类似fermi之前的CUDA设备),因此可以省去volatile关键字。 5.2.2 指令优化 1 由于GCN设备具有独立的标量计算单元,所以支持整数计算和浮点计算指令的双发,合理调度指令的顺序可以更好的隐藏指令的发射和计算延迟,比如将浮点计算指令和预取数据的地址计算指令交叉排列。 2 当指令中包含了一些特定的常量值时,编译器可以生成更小的代码,因为这些特殊的常量对应了指令的二进制编码中特定的几个比特位。这些值是 0,1~64,, -1~-16,+-0.5, +-1.0,+-2.0, +-4.0,1.0/(2*PI) 对1.0/(2*PI)内嵌常量的支持更多的是考虑到诸如FFT等图像计算方面的应用,但是只有矢量指令才支持在指令码中内嵌1.0/(2*PI)常量。同时自定义的PI值可能无法匹配指令支持的值,因此最好通过使用OpenCL中的内置的定义。 3 和比fermi更早期的CUDA设备类似,目前所有GCN(1.0~1.3)设备上对24位整数乘法提供原生支持,因此使用24位整数乘法具有更高的效率。 4 对于GCN1.1,GCN1.2的设备,在诸如归约和扫描的应用中尽量使用OpenCL内置的归约和扫面函数,这样可以帮助编译器生成DPP指令从而可以使用硬件上的数据并行引擎执行跨通道计算(无需通过LDS中转)。 4.3 构建性能可移植的程序 OpenCL是为跨平台的高性能并发程序设计而制定的开放式规范,虽然理论上使用OpenCL开发的程序可以在任何支持OpenCL的平台上运行,但是实际上受限于不同平台对OpenCL支持的力度以及不同硬件架构上得差异,使得同一个OpenCL程序在两个不同的设备上的性能表现可能差别很大(甚至这两个设备在理论上的技术指标很接近)。很多实际的应用,仅仅拥有代码的可移植性是不够的,因此本章主要讨论如何利用OpenCL的运行时编译系统构建性能可移植的程序。为了写出性能可移植的程序,不仅仅需要对同一厂商的不同架构的设备做针对性的优化,同时还要针对不同厂商的设备给出不同的优化实现;这一过程虽然增加了开发的时间和难度,但是从给与用户更佳体验的角度来说是完全值得的。下面我们以并行规约为例讲解如何开发性能可移植的OpenCL程序。 4.3.1 CUDA设备上的并行规约 #ifdef CUDA_DEVICE #if CUDA_SM<30 #define SMEM_SIZE264 inline voidwarp_reduce_add( double& s, __local volatile double* sptr, int lane ) { if(lane<16) { *sptr=s; s+=*(sptr+16); *sptr=s; s+=*(sptr+ 8); *sptr=s; s+=*(sptr+ 4); *sptr=s; s+=*(sptr+ 2); *sptr=s; s+=*(sptr+ 1); } } #else #define SMEM_SIZE 8 inline double__shfl( double val, int mask ) { double out; asm volatile ("{ \n\t" ".reg.b32 slo, shi,dlo, dhi ; \n\t" "mov.b64 { slo, shi}, %1 ; \n\t" "shfl.down.b32 dlo,slo, %2, 0x1f ; \n\t" "shfl.down.b32 dhi,shi, %2, 0x1f ; \n\t" "mov.b64 %0, { dlo,dhi } ; \n\t" "}" :"=d"(out) : "d"(val), "r"(mask) ); return out; } inline voidwarp_reduce_add( double& s ) { s+=__shfl(s,16); s+=__shfl(s, 8); s+=__shfl(s, 4); s+=__shfl(s, 2); s+=__shfl(s, 1); } #endif #else defined(GCN_DEVICE) #define SMEM_SIZE 260 #endif inline voidblock_reduce_add( double& s, __local double* smem, int lane, int warpid ) { #if CUDA_SM<30 __local volatile double*sptr=&smem[get_local_id(0)]; warp_reduce_add( s, sptr, lane ); if( lane==0 ){ smem[256+warpid]=s; } barrier(CLK_LOCAL_MEM_FENCE); sptr+=256; if(get_local_id(0)<4){ s=*sptr; s+=*(sptr+4); *sptr=s; s+=*(sptr+2); *sptr=s; s+=*(sptr+1); } #else warp_reduce_add( s ); if(lane==0){ smem[warpid]=s; } barrier(CLK_LOCAL_MEM_FENCE); if(get_local_id(0)<8) { s=smem[get_local_id(0)]; s+=__shfl(s,4); s+=__shfl(s,2); s+=__shfl(s,1); } #endif } __kernel voidkReduceAdd( __global double *g_mapped, __global double *g_temp, __globalunsigned int * g_mutex, __globalconst double * g_a, int n ) { __local unsigned int l_mutex; __local double l_temp[SMEM_SIZE]; double c=0; inti=(get_group_id(0)<<8)+get_local_id(0); unsigned intstride=get_num_groups(0)<<8; while(i<n){ c+=g_a; i+=stride;} const intlane=get_local_id(0)&31; const intwarpid=get_local_id(0)>>5; block_reduce_add( c, l_temp, lane,warpid ); if(get_local_id(0)==0){ __global double*g_out=(get_num_groups(0)>1)?&g_temp[get_group_id(0)]:g_mapped; *g_out=c; } if(get_num_groups(0)>1) { barrier(CLK_GLOBAL_MEM_FENCE); if(get_local_id(0)==0){ l_mutex=atom_add( &g_mutex, 1 ); } barrier(CLK_LOCAL_MEM_FENCE); if(l_mutex==(get_num_groups(0)-1)) { c=(get_local_id(0)<get_num_groups(0))?g_temp[get_local_id(0)]:0; block_reduce_add( c, l_temp, lane, warpid); if(get_local_id(0)==0){ g_mapped[0]=c; g_mutex=0; } } } } 首先看warp_reduce_add函数,我们使用指针操作,并把访问共享内存的次数降到了最小,如果在主函数内 改成 smem[threadIdx.x]=c然后将warp_reduce_add和block_reduce_add改成如下形式 inline void warp_reduce_add( __local volatiledouble* sptr ) { if(lane<16) { sptr[0]+=sptr[16]; sptr[0]+=sptr[8]; sptr[0]+=sptr[4]; sptr[0]+=sptr[2]; sptr[0]+=sptr[1]; } } inline double block_reduce_add(double* smem, int lane, int warpid ) { volatile double*sptr=&smem[threadIdx.x]; warp_reduce_add( s,sptr, lane ); if( lane==0 ){ smem[256+warpid]=s; }barrier(CLK_LOCAL_MEM_FENCE); sptr+=256; if(get_local_id(0)<4){ sptr[0]+=sptr[4]; sptr[0]+=sptr[2]; sptr[0]+=sptr[1]; } return smem[0]; } 那么会多出9次共享内存的访问操作。同时注意到在不支持warpshuffle操作的设备上每个block我们多分配了64字节(SMEM_SIZE=256+8)的共享内存,这是为了减少快内同步而做的优化,如果不加上这额外的共享内存,那么我们必须像下面这样多加一次同步: warp_reduce_add(s, sptr, lane ); barrier(CLK_LOCAL_MEM_FENCE); if( lane==0 ){ smem[warpid]=s; } barrier(CLK_LOCAL_MEM_FENCE); 如果warp_reduce_add之后不加同步的话,那么就无法保证来自其他warp的第一个线程对共享内存写入之前第一个warp的计算已经完成,就会有写冲突的问题。这一优化策略对于GCN设备上的实现同样适用,这其实也是对双缓冲技术的一个变相应用。 4.3.2 GCN设备上的并行规约 GCN设备上的实现和CUDA设备上的实现很相似,稍微的不同在CUDA设备上锁步计算的粒度是32,而GCN设备上的锁步粒度则是64(和CUDA设备上基于显式的warp并发一样,单个wavefront内的线程是以锁步的方式并行执行的,因此无需同步),因此为了在GCN架构上具有更好的亲和性,我们需要将锁步并发粒度改成64。 inline voidwavefront_reduce_add( double& s, __local double* sptr, int lane ) { if(lane<32) { *sptr=s; s+=*(sptr+32); *sptr=s; s+=*(sptr+16); *sptr=s; s+=*(sptr+ 8); *sptr=s; s+=*(sptr+ 4); *sptr=s; s+=*(sptr+ 2); *sptr=s; s+=*(sptr+ 1); } } inline voidblock_reduce_add( double& s, __local double* smem, int lane, intwavefront_id ) { __local volatile double*sptr=&smem[get_local_id(0)]; wavefront_reduce_add( s, sptr, lane); if( lane==0 ){ smem[256+wavefront_id]=s; } barrier(CLK_LOCAL_MEM_FENCE); sptr+=256; if(get_local_id(0)<2){ *sptr=s; s+=*(sptr+2); *sptr=s; s+=*(sptr+1); } } __kernel voidkReduceAdd( __global double * g_mapped, __global double *g_temp, __global unsigned int * g_mutex, __global const double * g_a, int n ) { __local unsigned int l_mutex; __local double l_temp[SMEM_SIZE]; double c=0; int i=(get_group_id(0)<<8)+get_local_id(0); unsigned intstride=get_num_groups(0)<<8; while(i<n){ c+=g_a; i+=stride;} const intlane=get_local_id(0)&63; const intwavefront_id=get_local_id(0)>>6; block_reduce_add( c, l_temp, lane,wavefront_id ); if(get_local_id(0)==0){ __globaldouble* g_out=(get_num_groups(0)>1)?&g_temp[get_group_id(0)]:g_mapped; *g_out=c; } if(get_num_groups(0)>1) { barrier(CLK_GLOBAL_MEM_FENCE); if(get_local_id(0)==0){ l_mutex=atom_add( &g_mutex, 1 ); } barrier(CLK_LOCAL_MEM_FENCE); if(l_mutex==(get_num_groups(0)-1)) { c=(get_local_id(0)<get_num_groups(0))?g_temp[get_local_id(0)]:0; block_reduce_add( c, l_temp, lane,wavefront_id ); if(get_local_id(0)==0){ g_mapped[0]=c; g_mutex=0; } } } } 现在,我们需要将所有的版本整合进同一个OpenCL内核程序中,这个可以简单的通过预处理指令来实现 #define CUDA_GPU 0 #define GCN_GPU 1 #pragmaOPENCL EXTENSION cl_khr_fp64:enable #if DEVICE == CUDA_DEVICE 包含CUDA设备版本的代码 #elif DEVICE == GCN_GPU 包含GCN设备版本的代码 #elif DEVICE == … 包含针对其它设备优化的版本 … #endif ‘DEVICE’以及’CUDA_SM’并没有在设备代码端定义,所以我们需要将它们作为命令行参数传递给OpenCL运行时编译系统,这样就可以让OpenCL驱动程序在运行时根据当前的设备选择合适的版本进行编译,从而在多个不同设备上都可以获得很高的性能 sprintf ( options, “-DDEVICE=dev_type –DCUDA_SM=cc”, …); clBuildProgram(prog, options, … ); 小结 GPU计算优化技术和方法种类繁多,每个人也可能习惯于自己的方法。但总体来说不会有太多不同,即使对于不同架构的设备很多优化技术也依然是通用的。开发性能可移植的程序从某种意义上来说不仅仅是一种挑战,也是对程序员技能的一种考验和磨练;从用户的角度考虑,他们是性能可移植性的程序的最终受益者,因此这样做也就更具有现实意义。有些优化技术在一些情况下可能适用,而另一些情况则可能适得其反,因此实际的验证必不可少。 参考资料 1《CUDAProgramming Guide》 2《AMD GCN Architecture Whitepaper》 3《AMDSouthern Island Series Instruction Set Architecture》 4《Graphics Core Next Architecture,Generation3》 5《AMD Parallel Processing OpenCL Programming Guide》
|