CUDA基本优化方法

一、基于编程模型和执行模型的优化方法

1.选取合适的gridDim和blockDim

blockDim最好为32的整数倍：因为执行指令的基本单位为线程束，线程束内的所有线程统一执行广播下来的命令，而线程束的线程数量基本为32。当block被分到SM中去，其会被划分为多个线程束，若blockDim!=线程束内线程数整数倍，则会造成线程的浪费。

2.减少存在分支的if

因为线程束中的所有线程执行同一条的指令，若出现存在出现分支的if，即线程束内的部分线程符合if条件，部分符合else条件，因cuda编译不具备分支预测能力，则会在一个指令执行if，下一个指令执行else，两个指令中都有部分不符合条件的线程空闲，从而造成浪费，如下述代码所示：

__global__ void mathKernel1(float *c)
{int tid = blockIdx.x* blockDim.x + threadIdx.x;float a = 0.0;float b = 0.0;if (tid % 2 == 0){a = 100.0f;}else{b = 200.0f;}c[tid] = a + b;
}

线程束内奇数线程（threadIdx.x为奇数）会执行else，偶数线程执行if，分化很严重。对此，我们可以更改if条件，使整个线程束内都符合if或else即可：

__global__ void mathKernel2(float *c)
{int tid = blockIdx.x* blockDim.x + threadIdx.x;float a = 0.0;float b = 0.0;if ((tid/warpSize) % 2 == 0){a = 100.0f;}else{b = 200.0f;}c[tid] = a + b;
}

可以得到相同但是错乱的结果C，这个顺序是无所谓的，可以后期调整。

3.循环展开

减少循环判断的次数，本质上也是减少分支

for (int i=0;i<100;i++)
{a[i]=b[i]+c[i];
}

for (int i=0;i<100;i+=2)
{a[i]=b[i]+c[i];a[i+1]=b[i+1]+c[i+1];
}

最新的cuda应该不需要手动展开了，在for前加上#pragma unroll 循环次数命令即可。

二、基于内存的优化方法

1.合并访问和对齐访问

核函数运行时需要从全局内存（DRAM）中读取数据，只有两种粒度：128字节、32字节。粒度可以理解为最小单位，也就是核函数运行时每次读内存，哪怕是读一个字节的变量，也要读128字节，或者32字节。
合并访问是指所有线程访问连续的对齐的内存块，对于L1 cache，内存块大小支持32字节、64字节以及128字节为单位读取数据。前提是，访问必须连续，并且访问的地址是以32字节对齐。
我们把一次内存请求——也就是从内核函数发起请求，到硬件响应返回数据这个过程称为一个内存事务。当一个内存事务的首个访问地址是缓存粒度（32或128字节）的偶数倍的时候：比如二级缓存32字节的偶数倍64，128字节的偶数倍256的时候，这个时候被称为对齐内存访问。

规约中可以通过设计累加的位置来实现合并对齐访问，提高效率

__global__ void ReductionSum(float *d_a, float *d_partial_sum)
{//申请共享内存，存在于每个block中 __shared__ float partialSum[threadsPerBlock];//确定索引int i = threadIdx.x + blockIdx.x * blockDim.x;int tid = threadIdx.x;//传global memory数据到shared memorypartialSum[tid]=d_a[i];//传输同步__syncthreads();//在共享存储器中进行规约for(int stride = blockDim.x/2; stride > 0; stride/=2){if(tid<stride) partialSum[tid]+=partialSum[tid+stride];__syncthreads();}//将当前block的计算结果写回输出数组if(tid==0)  d_partial_sum[blockIdx.x] = partialSum[0];
}

2.利用缓存来进行交叉访问（非合并的访问）

在某些情况下，无法避免进行交叉的读或写，如求矩阵的转置。这种情况下可以利用从缓存中读取的高效性，进行交叉内存访问（虽然按照列读是不合并的，但是使用一级缓存加载过来的数据在后面会被使用，虽然一级缓存一次读取128字节的数据，其中只有一个单位是有用的，但是剩下的并不会被马上覆盖，粒度是128字节，但是一级缓存的大小有几k或是更大，这些数据很有可能不会被替换，所以，我们按列读取数据，虽然第一行只用了一个，但是下一列的时候，理想情况是所有需要读取的元素都在一级缓存中，这时候，数据直接从缓存里面读取），从而让写结果到全局内存的过程是合并访问（写的过程不走缓存）。

__global__ void transformNaiveCol(float * MatA,float * MatB,int nx,int ny)
{int ix=threadIdx.x+blockDim.x*blockIdx.x;int iy=threadIdx.y+blockDim.y*blockIdx.y;int idx_row=ix+iy*nx;int idx_col=ix*ny+iy;if (ix<nx && iy<ny){MatB[idx_row]=MatA[idx_col];//MatB[idx_col]=MatA[idx_row];   slower}
}

3.基于共享内存的优化

共享内存是在他所属的线程块被执行时建立，线程块执行完毕后共享内存释放，线程块和他的共享内存有相同的生命周期。
共享内存读写速度较快，且可编程，故可用它来存储中间变量，或者先将内存中的数据读到共享内存中来，再做后续的操作。

__global__ void setRowReadColIpad(int * out)
{__shared__ int tile[BDIMY][BDIMX+IPAD];unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;tile[threadIdx.y][threadIdx.x]=idx;__syncthreads();out[idx]=tile[threadIdx.x][threadIdx.y];
}

共享内存按列读取不会导致交叉访问那么严重的延迟，代码中申请了共享内存来进行交叉访问，避免了交叉的全局内存访问，还对申请的动态内存做了填充，防止读取冲突（有疑问）。