用户
 找回密码
 立即注册
NVIDIA 官方 Developer 论坛 门户 查看主题

GPU编程之进击的优化系列二--GPU矩阵乘法优化技术

发布者: nvadmin | 发布时间: 2016-6-24 11:14| 查看数: 40030| 评论数: 1|帖子模式

第二章 GPU矩阵乘法的高效实现
2.0 前言
2.1 指令级并行和数据预取
2.2 双缓冲区
2.3 宽数据内存事务
2.4 二级数据预取

2.5 细节调优

第二章 GPU矩阵乘法优化技术
前言
本章通过介绍开发GPU上的高效矩阵乘法的各种优化方法来展示一些GPU深度优化技术,其中所用到的大多数优化技术不仅在支持CUDA的设备上有效,对于GCN设备,MIC设备和CPU上同样有效。如非特别说明,在所有版本的实现中,数据都是列序存储;下面以双精度矩阵乘法为例详细介绍一些非常有用的优化技术。
2.0 指令级并行和数据预取
对于类似矩阵乘法这样的高计算密度类型的应用,努力提高每个线程中的ILP(instructionlevel parallism,指令级并行)比单纯使TLP(thread level parallism,线程级并行)更加有效,而且是达到接近峰值性能唯一的途径。ILP的真髓在于通过充分发掘单一线程内的指令流水线中的并行和并发性(如果看过第一章就会明白,单一线程内的并行存在于被双发的指令在不同功能单元上的操作,比如在LD/ST上的访存操作和计算单元上的FMA操作可以同时进行;而并发则存在于各个功能单元上的分时指令序列)以及配合寄存器的最低延迟和超高带宽的优势(在kepler和GCN设备上,寄存器的带宽大约是共享内存的6倍)。循环展开通常是提高ILP的有效方法,原因在于通过循环展开编译器可以在更长的无数据依赖的指令序列中对指令进行排序,同时减少循环索引和部分存储器地址的计算。
dgemm, loop unrolling&prefetch
__device__ __forceinline__ voidd_rank8x8( double* C, const double* A, const double* B )
{
double a[8],b;
            a[0]=A[0*16];
            a[1]=A[1*16];
            a[2]=A[2*16];
            a[3]=A[3*16];
            a[4]=A[4*16];
            a[5]=A[5*16];
            a[6]=A[6*16];
            a[7]=A[7*16];
#pragma unroll
            for( int i=0;i<8; ++i ){
                        b=B[i*16];
                        C[i*8+0]+=a[0]*b;
                        C[i*8+1]+=a[1]*b;
                        C[i*8+2]+=a[2]*b;
                        C[i*8+3]+=a[3]*b;
                        C[i*8+4]+=a[4]*b;
                        C[i*8+5]+=a[5]*b;
                        C[i*8+6]+=a[6]*b;
                        C[i*8+7]+=a[7]*b;
            }
}
__global__ voidcuk_dgemm_unroll( double* d_C, const double* d_A, const double* __restrict__ d_B, int n, int lda, int ldb, int ldc )
{
            __shared__ doublesmem[2048];
            double p0, p1,p2, p3, q0, q1, q2, q3, c[64]={0.f};
            int bx, by,i, x, y, u, v, k;
    bx=blockIdx.x;
            by=blockIdx.y;
            i=threadIdx.x;
            x=i&31;
            y=i>>5;
            u=i&15;
            v=i>>4;
            d_C+=((by<<7)+(i>>7))*ldc+(i&127);
            d_A+=y*lda+(bx<<7)+x;
            d_B+=((by<<7)+u)*ldb+v;
            double*St=&smem[(y<<7)+x];
            double*At=&smem[u];
            double*Bt=&smem[1024+v];
            p0=d_A[0*32];
            p1=d_A[1*32];
            p2=d_A[2*32];
            p3=d_A[3*32];
            q0=d_B[0*32*ldb];
q1=d_B[1*32*ldb];
            q2=d_B[2*32*ldb];
q3=d_B[3*32*ldb];
            for( k=n-8;k>=0; k-=8 )
            {
                        *(St+0*32)=p0;
                        *(St+1*32)=p1;
                        *(St+2*32)=p2;
                        *(St+3*32)=p3;
                        *(St+4*32)=q0;
                        *(St+5*32)=q1;
                        *(St+6*32)=q2;
                        *(St+7*32)=q3;                        
__syncthreads();                    
                        if(y<k){
                                    d_A+=lda<<3;d_B+=8;
                                    p0=d_A[0*32];
                                    p1=d_A[1*32];
                                    p2=d_A[2*32];
                                    p3=d_A[3*32];
                                    q0=d_B[0*32*ldb];
q1=d_B[1*32*ldb];
                                    q2=d_B[2*32*ldb];
q3=d_B[3*32*ldb];
                        }
            #pragma unroll
                        for( int i=0;i<8; ++i ){
                                    d_rank8x8(c, &At[i*128], &Bt[i*128] );
                        }__syncthreads();
            }
#pragma unroll
            for( int i=0;i<64; i+=8 ){
                        d_C[0*16]=c[i+0];
                        d_C[1*16]=c[i+1];
                        d_C[2*16]=c[i+2];
                        d_C[3*16]=c[i+3];
                        d_C[4*16]=c[i+4];
                        d_C[5*16]=c[i+5];
                        d_C[6*16]=c[i+6];
                        d_C[7*16]=c[i+7];
                        d_C+=ldc<<4;
            }
}
    第一个版本中我们主要的优化方法是通过循环展开获得更多的ILP以及通过数据预取和计算的重叠隐藏访存延迟。(注:循环展开并非在任何情况下都能提升指令级并行,待展开的计算中还需要尽可能的无数据依赖关系,待展开的循环体中的计算也应尽可能的简单;而且过度的循环展开还可能会造成寄存器溢出或是设备资源占用率过低从而导致性能下降)
    注意内核参数const double* __restrict__ d_B”,对于计算能力3.5+的设备编译器会自动使用LDG指令通过纹理管线加载指令加载数据从而可以利用纹理缓存减少非合并访问的影响。对于计算能力2.x或3.x的设备,显式的使用纹理通常比普通的缓存效果更好,但也受限于L2缓存有限的带宽和纹理操作的长延迟。纹理管线独立于普通的存储路径,可以提升可用的缓存带宽(纹理缓存的带宽+L1缓存的带宽),从而更加平衡负载,并减小不同数据之间的L1缓存污染(通过纹理缓存加载的数据不经过L1缓存)。对于计算能力5.x的设备,由于统一了L1缓存和纹理缓存,因此在不使用纹理滤波的情况下使用线性纹理和普通L1缓存已经没有本质的区别了。     此外所选择的的网格布局也会直接影响性能和代码的可读性。根据不同的架构选择每线程计算的元素数量,太少的话不能充分利用寄存器文件高带宽低延迟的优势;太多的话会造成寄存器溢出。对于计算能力2.x和3.0的设备,每个线程计算16个元素(1x16或者4x4),线程块的大小选择为256(计算能力2.x,block布局可设置为256,64x4或32x8,每个block计算64x64的子矩阵)或1024(计算能力3.0,block布局可设置为1024,128x8或32x32,每个block计算128x128的子矩阵)。对于计算能力3.5+,5.x的设备,比较好的选择是每线程计算8x8个元素(对于比较小的规模可以每个线程计算4x8个元素,每个block计算128x64大小的子矩阵),线程块的大小设置为256(block布局可设置为16x16,32x8,128x2或者256;每个block计算128x128的子矩阵,一个SM最多可部署一个活跃block(对于计算能力3.7的设备,每个SMX上的活跃block队列最多可部署2个)。
     我们采用256的block布局(但是采用32x8的布局进行内存访问),每个线程计算64个元素,每个block计算128x128的子矩阵。
                                       
2.1 双缓冲区技
    我们知道GPU网格调度器通过block内的warp(在AMD的GPU中称之为wavefront)切换隐藏内存访问和计算延迟;而通过block之间的切换隐藏块内同步延迟,但这里由于寄存器和共享内存资源的限制,一个SMX上最多只能部署一个活跃block(对于计算能力3.7的设备,由于具有双倍的共享内存和寄存器文件,所以每个SMX可以支持两个block并发),因而无法再通过增加每个SMX上可以并发的block数量来进一步提升效率。类似单个线程内的并行和并发,这时我们就需要考虑如何在单个block内寻找更多的并行和并发机制。容易看出主循环体内的并行和并发受限于块内同步,因此如果消除下一个循环步预取的数据和当前计算中所用数据的依赖关系,则第二个同步原语是可以消除的,为此我们通过双倍的共享内存作为双缓冲来消除这种影响:
dgemm,loop unrolling+double buffering
__global__ voidcuk_dgemm_unroll_db( double* d_C, const double* d_A, const double* __restrict__ d_B, int n, int lda, int ldb, int ldc )
{
            __shared__ doublesmem[4096];
            
            for( k=n-8;k>=0; k-=8 )
            {           
__syncthreads();                    
                        if(threadIdx.y<k){
                                    d_A+=lda<<3;d_B+=8;
                                    sidx^=0x800;
                                    smem[0*32+sidx]=d_A[0*32];
                                    smem[1*32+sidx]=d_A[1*32];
                                    smem[2*32+sidx]=d_A[2*32];
                                    smem[3*32+sidx]=d_A[3*32];
                                    smem[1024+0*32+sidx]=d_B[0*32*ldb];
smem[1024+1*32+sidx]=d_B[1*32*ldb];
                                    smem[1024+2*32+sidx]=d_B[2*32*ldb];
smem[1024+3*32+sidx]=d_B[3*32*ldb];
                        }
            #pragma unroll
                        for( int i=0;i<8; ++i ){
                                    d_rank8x8(c, &smem[i*128+u], &smem[i*128+v] );
                        }
                        u^=0x800; v^=0x800;
            }
            ……
}  
    使用双缓冲技术将当前计算的读取缓冲区和负责装载下一步数据的缓冲区相分离,从而加载下一步数据到共享内存的操作可以和当前的计算重叠执行,减少了同步开销从而减少流水线的中断。但是使用双缓冲时需要注意当前的设备上具体的效果,此外不同的计算规模效果也会不同。对于计算能力2.x的设备,由于双精度指令和存储指令不能双发(因此LD/ST指令和双精度浮点计算指令是顺序发射的,但在指令完成操作的这段延迟期间,其执行仍是重叠的),在其上使用双缓冲区的效果相比kepler和maxwell设备会小些;另外,如果计算规模比较小,那么使用双缓冲会由于更复杂的地址计算而降低效率。
2.2 宽内存事务
    我们还可以做得更好,对内层循环(d_rank8x8)进行分析可知,每个循环步共有512条FMA指令和128条共享内存加载指令,浮点计算指令所占比例(这里我们不考虑数据预取和缓冲区地址计算的相关指令,这样并不影响分析的有效性)是0.8。如果能够进一步提升这个比例,就有可能进一步提升性能。因为FMA操作的数量是固定的,因此可以通过使用128bit存储操作来缩减每个线程的内存事务的数量。首先需要对d_rank8x8进行改写:
#define mRank8x8(i0,i1,i2,i3){            \
            c[0*8+0]+=a[i0].x*b[i0].x;         \
            c[0*8+1]+=a[i0].y*b[i0].x;         \
            c[0*8+2]+=a[i1].x*b[i0].x;         \
            c[0*8+3]+=a[i1].y*b[i0].x;         \
            c[0*8+4]+=a[i2].x*b[i0].x;         \
            c[0*8+5]+=a[i2].y*b[i0].x;         \
            c[0*8+6]+=a[i3].x*b[i0].x;         \
            c[0*8+7]+=a[i3].y*b[i0].x;         \
            c[1*8+0]+=a[i0].x*b[i0].y;\
            c[1*8+1]+=a[i0].y*b[i0].y;\
            c[1*8+2]+=a[i1].x*b[i0].y;\
            c[1*8+3]+=a[i1].y*b[i0].y;\
            c[1*8+4]+=a[i2].x*b[i0].y;\
            c[1*8+5]+=a[i2].y*b[i0].y;\
            c[1*8+6]+=a[i3].x*b[i0].y;\
            c[1*8+7]+=a[i3].y*b[i0].y;\
            c[2*8+0]+=a[i0].x*b[i1].x;\
            c[2*8+1]+=a[i0].y*b[i1].x;\
            c[2*8+2]+=a[i1].x*b[i1].x;\
            c[2*8+3]+=a[i1].y*b[i1].x;\
            c[2*8+4]+=a[i2].x*b[i1].x;\
            c[2*8+5]+=a[i2].y*b[i1].x;\
            c[2*8+6]+=a[i3].x*b[i1].x;\
            c[2*8+7]+=a[i3].y*b[i1].x;\
            c[3*8+0]+=a[i0].x*b[i1].y;\
            c[3*8+1]+=a[i0].y*b[i1].y;\
            c[3*8+2]+=a[i1].x*b[i1].y;\
            c[3*8+3]+=a[i1].y*b[i1].y;\
            c[3*8+4]+=a[i2].x*b[i1].y;\
            c[3*8+5]+=a[i2].y*b[i1].y;\
            c[3*8+6]+=a[i3].x*b[i1].y;\
            c[3*8+7]+=a[i3].y*b[i1].y;\
            c[4*8+0]+=a[i0].x*b[i2].x;\
            c[4*8+1]+=a[i0].y*b[i2].x;\
            c[4*8+2]+=a[i1].x*b[i2].x;\
            c[4*8+3]+=a[i1].y*b[i2].x;\
            c[4*8+4]+=a[i2].x*b[i2].x;\
            c[4*8+5]+=a[i2].y*b[i2].x;\
            c[4*8+6]+=a[i3].x*b[i2].x;\
            c[4*8+7]+=a[i3].y*b[i2].x;\
            c[5*8+0]+=a[i0].x*b[i2].y;\
            c[5*8+1]+=a[i0].y*b[i2].y;\
            c[5*8+2]+=a[i1].x*b[i2].y;\
            c[5*8+3]+=a[i1].y*b[i2].y;\
            c[5*8+4]+=a[i2].x*b[i2].y;\
            c[5*8+5]+=a[i2].y*b[i2].y;\
            c[5*8+6]+=a[i3].x*b[i2].y;\
            c[5*8+7]+=a[i3].y*b[i2].y;\
            c[6*8+0]+=a[i0].x*b[i3].x;\
            c[6*8+1]+=a[i0].y*b[i3].x;\
            c[6*8+2]+=a[i1].x*b[i3].x;\
            c[6*8+3]+=a[i1].y*b[i3].x;\
            c[6*8+4]+=a[i2].x*b[i3].x;\
            c[6*8+5]+=a[i2].y*b[i3].x;\
            c[6*8+6]+=a[i3].x*b[i3].x;\
            c[6*8+7]+=a[i3].y*b[i3].x;\
            c[7*8+0]+=a[i0].x*b[i3].y;\
            c[7*8+1]+=a[i0].y*b[i3].y;\
            c[7*8+2]+=a[i1].x*b[i3].y;\
            c[7*8+3]+=a[i1].y*b[i3].y;\
            c[7*8+4]+=a[i2].x*b[i3].y;\
            c[7*8+5]+=a[i2].y*b[i3].y;\
            c[7*8+6]+=a[i3].x*b[i3].y;\
            c[7*8+7]+=a[i3].y*b[i3].y;\
}
#define mFetchSmem(i0,i1,i2,i3,k){    \
            a[i0]=smem[    k*64+ 0+u];      \
            a[i1]=smem[    k*64+16+u];     \
            a[i2]=smem[    k*64+32+u];     \
            a[i3]=smem[    k*64+48+u];     \
            b[i0]=smem[512+k*64+0+v];  \
            b[i1]=smem[512+k*64+16+v]; \
            b[i2]=smem[512+k*64+32+v]; \
            b[i3]=smem[512+k*64+48+v]; \
}
    通过使用128bit存储器操作,我们将共享内存加载指令数目减少了一半,现在只有64条共享内存加载指令,浮点计算指令所占比例约是0.89。虽然128bit的共享内存操作以及其引发的bank conflicts会带来更高的延迟,但是通过合理的安排计算足以将这些延迟抵消掉。
dgemm,loop unrolling+double buffering+128bit LD/ST
__global__ voidcuk_dgemm_unroll_db_128b( double* d_C, const double2* d_A, const double2* __restrict__ d_B, int n, int lda, int ldb, int ldc )
{
            __shared__ double2smem[2048];
            double2 a[4],b[4];
            ……
            for( k=n-8;k>=0; k-=8 )
            {
                        __syncthreads();
                        if(y<k){
                                    d_A+=lda<<3;d_B+=8;
                                    sidx^=0x800;
                                    ((double*)smem)[0*32+sidx]=d_A[0*32];
                                    ((double*)smem)[1*32+sidx]=d_A[1*32];
                                    ((double*)smem)[2*32+sidx]=d_A[2*32];
                                    ((double*)smem)[3*32+sidx]=d_A[3*32];
                                    ((double*)smem)[1024+0*32+sidx]=d_B[0*32*ldb];
((double*)smem)[1024+1*32+sidx]=d_B[1*32*ldb];
                                    ((double*)smem)[1024+2*32+sidx]=d_B[2*32*ldb];
((double*)smem)[1024+3*32+sidx]=d_B[3*32*ldb];
                        }
            #pragma unroll
                        for( int i=0;i<8; ++i ){
                                    mFetchSmem(0,1,2,3,i)
                                    mRank8x8(0,1,2,3)
                        }
                        u^=0x400;v^=0x400;
            }
            ……
            
}
     合理使用128bit存储操作可以减少单个线程内的访存事务的数量(也会因此提升计算指令的整体比例)从而提升性能。对于双字宽数据,提升效果相对较小;因为相对于单字宽数据,双字宽内存事务的数量减少相对较低,因此计算指令所占比例的提升也较小,在单精度的情况下会得到更好的效果。
2.3 二级数据预取
    除了对全局内存的数据进行预取外,我们可以进一步加入共享内存的预取流水,从而通过将计算和共享内存操作重叠来进一步隐藏访问共享内存的延迟,同时减少计算中等待数据的流水线停顿。另外,对于双精度矩阵乘法在某些设备上使用双缓冲反而会降低性能,尤其是规模比较小的矩阵不建议使用;而加入对共享内存数据的预取在多数设备和规模上均能性能提升(作者在GTX780,GTX970和TeslaK20, Teslak40上测试中都获得了提升)。
dgemm,loop unrolling + double buffering + 128bit LD/ST + smem_prefetch
__global__ voidcuk_dgemm_unroll_db_128b_prefsmem( double* d_C, const double* d_A, const double* __restrict__ d_B, int n, int lda, int ldb, int ldc )
{
            ……
            __syncthreads();
            mFetchSmem(0,2,4,6,0)
            for( k=n-8;k>=0; k-=8 )
            {
                        if(y<k)
                        {
                                    d_A+=lda<<3;d_B+=8;
                                    sidx^=0x800;
                                    ((double*)smem)[0*32+sidx]=d_A[0*32];
                                    ((double*)smem)[1*32+sidx]=d_A[1*32];
                                    ((double*)smem)[2*32+sidx]=d_A[2*32];
                                    ((double*)smem)[3*32+sidx]=d_A[3*32];
                                    ((double*)smem)[1024+0*32+sidx]=d_B[0*32*ldb];
((double*)smem)[1024+1*32+sidx]=d_B[1*32*ldb];
                                    ((double*)smem)[1024+2*32+sidx]=d_B[2*32*ldb];
((double*)smem)[1024+3*32+sidx]=d_B[3*32*ldb];
                        }
                        mFetchSmem(1,3,5,7,1)
                        mRank8x8(0,2,4,6)
                        mFetchSmem(0,2,4,6,2)
                        mRank8x8(1,3,5,7)
                        mFetchSmem(1,3,5,7,3)
                        mRank8x8(0,2,4,6)
                        mFetchSmem(0,2,4,6,4)
                        mRank8x8(1,3,5,7)
                        mFetchSmem(1,3,5,7,5)
                        mRank8x8(0,2,4,6)
                        mFetchSmem(0,2,4,6,6)
                        mRank8x8(1,3,5,7)
                        mFetchSmem(1,3,5,7,7)
                        mRank8x8(0,2,4,6)
                        __syncthreads();
u^=0x400,v^=0x400;
                        mFetchSmem(0,2,4,6,0)
                        mRank8x8(1,3,5,7)
            }
            ……
}
2.4 细节调优
    以上代码都假设A的行数,B的列数以及C的行列数是128的倍数。所以如果矩阵大小不匹配block的计算尺寸,那么更好的实现是添加一个边界处理的内核专门处理边界。如果边界大小和block的计算尺寸相差不大(因此不会浪费太多显存空间),则可以通过填充零将矩阵补齐到匹配block的计算尺寸,这样就不再需要单独的边界处理内核,效率也会更高些。还有一点需要说明:将数据存取操作和浮点计算操作混合排列可以获得更高的效率(更高的IPC),这里为了更清晰的表达所以没有这样做。上述代码的性能仍有较大的提升空间,如果想获得接近峰值的效率,需要手工使用PTX或SASS对细节进行调优,否则仅仅使用nvcc是无法获得接近于峰值性能的。如果开发人员想通过使用PTX显式的控制诸如寄存器分配,指令排序等问题,可能得不到想要的结果。原因是PTX并非真正的面向本地机器指令的汇编语言,而是一种为了兼容不同计算能力的设备而设计的一种类似寄存器传输语言(RTL)的中间层伪汇编语言;从源代码到PTX的编译过程只会做一些初级的优化,多数深度优化,比如寄存器分配,指令重排,基本块融合等都是在从PTX代码到SASS原生汇编代码的过程中进行的(PTX和SASS的关系类似AMDGPU编程中的IL和GCN原生ISA的关系,这里所说的多数情况在IL和GCN ISA之间也同样存在)。因此若要显式的控制指令的执行顺序和寄存器分配等就需要直接使用SASS,不过NVIDIA现在并未给出直接使用汇编编写GPU内核程序的开发工具,也未开放任何关于指令集编码信息的文档,因此开发者需要自己想办法绕过这个限制。至少对于矩阵乘法,ptxas做的并不是太好(使用C/C++只能得到最高约75%峰值的效率,使用PTX可以超过80%的峰值,但是离最好的效果还有差距),主要有两个原因:
1 指令的排序效果不理想,也就是LD/ST指令,FMA,IADD,LOP,ISETP等指令之间的排列位置不够理想,最好的指令之间的排列距离要综合考虑指令管线的深度,dual issue,以及各种操作的发射延迟和计算延迟;对指令放置在不同位置进行测试选优也往往是必不可少的。对于这些情况人工通常可以比编译器做的更好。
2 寄存器bank冲突(bank-conflict)是最主要的原因。对于2-way,每条FMA指令发射到流水线会增加一个时钟的延迟,3-way则会增加2个时钟周期,这会严重影响效率(每个循环步增加的时钟周期=存在端口冲突的指令数x (n_conflict_ways-1),这里每个循环步的FMA指令数是512),ptxas为了避免寄存器bank冲突以及强制将寄存器对齐到指令的操作数端口而增加了很多冗余的MOV指令(循环体的追尾效应,使用cuobjdump查看SASS代码会看到在循环体的末尾处多出很多MOV操作),虽然MOV操作是直接在寄存器文件内部移动数据(意味着无需像其它操作那样先从数据端口读取寄存器中的数据,然后再将数据通过数据端口写入寄存器),但仍然不是免费的,哪怕其延迟很低。即使这样仍然不能保证完全消除寄存的端口冲突。虽然可以通过在循环体前端移除数据预取代码中的条件代码以及通过循环剥离移除循环体内的条件代码来消除或是减少格外的寄存器复制操作,但由于第一因素的存在,所以性能的提升仍然受限。
下面给出各个版本的测试结果(每个测例循环100次取平均,单位ms)
version                                 GTX970                                       
                        size           1024x1024   2048x2048  4096x4096
cublas                          0.019500    0.141330   1.109940
unroll                      0.017320    0.136030   1.081080
unroll+128b                 0.016690   0.132290    1.047700
unroll+128b+prefsmem           0.016670    0.131820   1.042550
unroll+db+128b           0.017000    0.135720   1.070310
unroll+db+128b+prefsmem     0.016850    0.134790   1.068920                 
本章小结
          通过合理组合使用循环展开,128bit存储操作以及对全局内存和共享内存的两级数据预取可以有效提升矩阵乘法的效率。在这些优化技术中,循环展开,数据预取,128bit内存通常是有效的,双缓冲技术则需视情况而定。
小结
    一开始我们就从一个性能很高的版本开始更进一步的优化,之所以没有从一个更加简单的版本(比如CUDAprogramming guide上的例子)开始逐步走完全程,并不是因为作者太懒(好吧,我承认是有这个原因)。其一是因为各种相关书籍都有介绍,除了占用篇幅,作者实在找不到把那些内容放在本书的理由。更主要的原因是作者更建议一开始就从尽可能高效的方法去实现,这个建议对任何项目都适用。假设你正在做一个比较大的项目,可能会想着先让项目运行起来,至于效率以后可以逐步改进。这种想法对有些项目可能适用。但是对于更多的项目这样做可能会大幅度拖延开发进度,因为当你准备优化时,可能会发现为了改进性能,不得不改变数据结构(因为有时候大幅度的性能提升必须依赖特定的数据结构和算法)和大范围的改写代码逻辑。有些情况下一个简单的数据转换接口即可在时间和效率上得到折中(通常意味着不是最好);但很多时候对数据结构或是算法的改变却会牵一发而动全身甚至从架构到代码都需要重新设计开发。所以强烈建议从一开始就用尽可能优化的方法作为起点,用尽可能优质的实现构建你的项目,这样前期可能会需要多些时间,但是之后会发现这样可以大大提高整体的开发进度,同时项目的质量也会得到保障。总之,对一个低质量的程序通过渐进的方法不断优化的工作效率远低于直接使用优化的方法;正所谓快刀斩乱麻,有时候屏蔽已有的东西,可以避免被潜意识中经验性或现有的代码规则所限制。
参考资料
       1《cuda programming guide》
       2《kepler tuning guide》
       3《maxwell tuning guide》
       4《PTX ISA version 4.2》
       5《nvidia-kepler-gk110-architecture-whitepaper》
6Performance Upper Bound Analysis and Optimization of SGEMM on Fermi and KeplerGPUs,by Junjie Lai, Andr_e Seznec。
7Anatomy of High-Performance Matrix Multiplication,KAZUSHIGEGOTO,The University of Texas at Austin and ROBERT A. VAN DE GEIJN TheUniversity of Texas at Austin。
8On Reducing TLB misses in Matrix Multiplication,kazushige goto, Robert van degeijn, november 1, 2002。


最新评论

maybe0524 发表于 2019-6-21 08:36:53
请问下楼主,有什么示例代码可以体验一下,优化前后的差距呢