第1章基本概念
要进行Cuda核函数的优化，需要先了解GPU硬件的特性，以及代码中Block,Grid和GPU硬件的对应关系。
1.1.Cuda并行模型
Cuda的并行模型为SIMT,单指令多线程。32个为一组的线程，执行的指令是一样的，因此如果线程有分支，且32个线程执行的分支不一样，则所有分支的指令都要执行，导致效率降低。CPU的指令集SSE,AVX2等属于SIMD。SIMD是严格的单指令多数据，无法处理逻辑分支。SIMT相比SIMD，有其灵活性，但效率相对低。
1.2.Cuda硬件特性
一个GPU由多个流多处理器(SM)组成，比如4080有76个SM，而3060有28个SM。每个SM由多个CudaCore组成，一般有128个。NV的GPU还有专用的TersorCore，用于深度学习推理和训练，可以实现比CudaCore高10倍的效率。
1.3.GPU硬件和Cuda概念对应关系
软件的一个线程就跑到一个Cuda核心上。一个Block上的所有线程会跑到同一个SM上，线程通常32个为一组同时执行，称为wrap。SM上有128个cuda核心，因此一个SM可能会同时处理一个block或者多个block上的wrap。
第2章核函数优化方法
2.1.多个核函数并行
对于计算规模较小的情况，开启的grid数量小于实际的硬件SM处理器，则不能充分发挥GPU的优势。这种情况下可以使用核函数并行的方法优化，或许将这样的计算任务放到CPU更合适。
多个核函数并行需要硬件支持，在cuda-z工具中可以查看。
同时4个流上的核函数并发执行：
for (int i = 0; i < n_streams; i++)
{
kernel_1 << <grid, block, 0, streams[i] >> > (d_A);
kernel_2 << <grid, block, 0, streams[i] >> > (d_A);
kernel_3 << <grid, block, 0, streams[i] >> > (d_A);
kernel_4 << <grid, block, 0, streams[i] >> > (d_A);
}
编程时，将没有依赖关系的核函数放到不同的流上，由GPU调度最优的并行策略。如果所有的核函数放到一个流上，GPU没有机会调度了。
2.2.核启动参数优化
一般blocksize使用32 * 16 or 32 * 8，gridsize根据图像大小，以及一个线程处理一个像素，计算得到。
根据Nppi对3*3的核的启动参数，nppi库对小核一个线程会处理32个像素。实际测试发现，一个线程处理2 * 2，4 * 4个像素，均会有一定的提升。具体如下：
1 * 1 2 * 2 3 * 3的核，少开grid，一个线程处理多个像素，相比一个线程处理一个像素，在4080上，可以提升1倍效率，3060上提升20%。
1 * 1的核函数有：Inspect，Sub，Add，GBR2Gray等
2 * 2的核函数有：2 * 2滤波，Remap等
3 * 3的核函数有：Sobel，Gauss，Morph等。
一个线程处理多个像素核函数优化方法：
__global__ void kernelGauss(MyT* datain, MyT* dataout, int nW, int nH)
{
int col = blockIdx.x * blockDim.x + threadIdx.x;;
int row = blockIdx.y * blockDim.y + threadIdx.y;
if (col >= nW - 1 || row >= nH - 1 || col < 1 || row < 1)
return;
int gx = 0;
gx += datain[(row - 1) * nW + col - 1];
gx += 2 * datain[row * nW + col - 1];
gx += datain[(row + 1) * nW + col - 1];
gx += 2 * datain[(row - 1) * nW + col];
gx += 4 * datain[row * nW + col];
gx += 2 * datain[(row + 1) * nW + col];
gx += datain[(row - 1) * nW + col + 1];
gx += 2 * datain[row * nW + col + 1];
gx += datain[(row + 1) * nW + col + 1];
dataout[row * nW + col] = gx >> 4;
}
__global__ void kernelGauss_4x4(MyT* datain, MyT* dataout, int nW, int nH)
{
int _col = blockIdx.x * blockDim.x + threadIdx.x;;
int _row = blockIdx.y * blockDim.y + threadIdx.y;
if (_col*4 >= nW - 5 || _row*4 >= nH - 5)
return;
int gx = 0;
#pragma unroll
for (int j = 0; j < 4; j++)
{
int row = 4 * _row + j;
#pragma unroll
for (int i = 0; i < 4; i++)
{
int nS4 = 0;
int col = 4 * _col + i;
gx += datain[(row - 1) * nW + col - 1];
gx += datain[(row + 1) * nW + col - 1];
gx += datain[(row - 1) * nW + col + 1];
gx += datain[(row + 1) * nW + col + 1];
nS4 += datain[row * nW + col - 1];
nS4 += datain[(row - 1) * nW + col];
nS4 += datain[(row + 1) * nW + col];
nS4 += datain[row * nW + col + 1];
gx += nS4 * 2;
gx += 4 * datain[row * nW + col];
dataout[row * nW + col] = gx >> 4;
}
}
}
5 * 5的核，使用一个线程处理多个像素的策略，提升不明显，在4080上大概有7%（8k*8k的图像做灰度膨胀算法，570us提升到530us。）。
2.3.3 * 3滤波核共享内存的方法
对3 * 3的滤波核，由于线程间重复读取全局内存，自然想到使用共享内存。实际编码后发现，在4080和3060显卡上均观察不到效率提升。
在没有加载边部像素的情况下，耗时才和未优化的版本持平。
__global__ void kernelSobel(MyT* datain, MyT* dataout, int nW, int nH, float fS)
{
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
if (col >= nW - 1 || row >= nH - 1 || col < 1 || row < 1)
return;
int gx = 0;
gx = max(gx, datain[(row - 1) * nW + col - 1]);
gx = max(gx, datain[row * nW + col - 1]);
gx = max(gx, datain[(row + 1) * nW + col - 1]);
gx = max(gx, datain[(row - 1) * nW + col]);
gx = max(gx, datain[row * nW + col]);
gx = max(gx, datain[(row + 1) * nW + col]);
gx = max(gx, datain[(row - 1) * nW + col + 1]);
gx = max(gx, datain[row * nW + col + 1]);
gx = max(gx, datain[(row + 1) * nW + col + 1]);
dataout[row * nW + col] = gx;
}
__global__ void kernelSobel_sharedMem(MyT* datain, MyT* dataout, int nW, int nH, float fS)
{
__shared__ MyT sharedMem[BLOCK_SIZE_Y + BLOCK_BORDER * 2][BLOCK_SIZE_X + BLOCK_BORDER * 2];
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
int localCol = threadIdx.x + 1;
int localRow = threadIdx.y + 1;
//加载数据
if (col < nW && row < nH)
{
sharedMem[localRow][localCol] = datain[row * nW + col];
边部加载 46->37(16 6)
//if (threadIdx.x == 0 && col > 0) //左边列
//{
// sharedMem[localRow][0] = datain[row * nW + col - 1];
//}
//if (threadIdx.x == blockDim.x - 1 && col < nW - 1) {
// sharedMem[localRow][BLOCK_SIZE_X + 1] = datain[row * nW + col + 1];
//}
//if (threadIdx.y == 0 && row > 0)
//{
// sharedMem[0][localCol] = datain[(row - 1) * nW + col];
//}
//if (threadIdx.y == blockDim.y - 1 && row < nH - 1)
//{
// sharedMem[BLOCK_SIZE_Y + 1][localCol] = datain[(row + 1) * nW + col];
//}
}
__syncthreads();
if (col >= nW - 1 || row >= nH - 1 || col < 1 || row < 1)
return;
int gx = 0;
gx = max(gx, sharedMem[localRow - 1][localCol - 1]);
gx = max(gx, sharedMem[localRow][localCol - 1]);
gx = max(gx, sharedMem[localRow + 1][localCol - 1]);
gx = max(gx, sharedMem[localRow - 1][localCol]);
gx = max(gx, sharedMem[localRow][localCol]);
gx = max(gx, sharedMem[localRow + 1][localCol]);
gx = max(gx, sharedMem[localRow - 1][localCol + 1]);
gx = max(gx, sharedMem[localRow][localCol + 1]);
gx = max(gx, sharedMem[localRow + 1][localCol + 1]);
dataout[row * nW + col] = gx;
}
观察Nppi的算子，在33滤波核的时候，从Nsight中没有发现使用共享内存。55及以上的滤波核，Nppi使用的共享内存优化。
2.4.5 * 5滤波核共享内存的方法
5 * 5的膨胀腐蚀核，可以采用共享内存的方式优化：
先进行3 * 3的形态学操作，将结果放到共享内存上，接着从共享内存取值，继续做33形态学，就完成了55的灰度膨胀腐蚀操作。在4080上大概有15%的提升（8k*8k的图像，570us提升到490us）。
分别为未优化(570us)，1个线程处理4个像素(530us)，使用共享内存优化的版本(490us)。
__global__ void kernel5x5(MyT* datain, MyT* dataout, int nW, int nH, float fS)
{
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
if (col >= nW - 2 || row >= nH - 2 || col < 2 || row < 2)
{
dataout[row * nW + col] = 0;
return;
}
int gx = 0;
gx = max(gx, datain[(row - 2) * nW + col - 2]);
gx = max(gx, datain[(row - 1) * nW + col - 2]);
gx = max(gx, datain[row * nW + col - 2]);
gx = max(gx, datain[(row + 1) * nW + col - 2<
最低0.47元/天 解锁文章
1191
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
