关于核函数的流并发 || 核函数的嵌套并发问题

  • 3 replies
  • 248 views
关于核函数的流并发 || 核函数的嵌套并发问题
« 于: 八月 13, 2019, 05:55:43 pm »
目前依然是cuda小白,以下有错请指正,谢谢!
目前使用显卡GeForce RTX 2080;
SM 46
也就是说同一时刻可并行的线程为SM*1024 = 47104;

面临的问题:
当前已实现一个核函数,功能为对4096*3072图片的bayer格式转换为24位彩色图片;
dim3 block(32,32);
dim3 grid((ImgW*3 + block.x - 1) / block.x, (ImgH + block.y - 1) / block.y);
BayerKernel<<<grid, block>>>(........);

上面这个核函数单次运行仅需1.8ms;
接下里我有100张图片均需要转换,所以总耗时在190ms左右;
于是想到核函数并行,即标题所述;

但最终结果通过性能分析工具发现并没有并行,全是串行,顶多在两个核函数之前会有细微的并行;

通过网络查询学习最终得到以下结论:
想实现核函数的并行需满足以下条件:
    1.每个核函数的局部变量不能多。(有个上限,应该和硬件参数有关)
    2.同一时刻,多个核函数的总并发线程不能大于硬件的活跃线程数(1024*SM)(47104);
       也就是说,假如我开了10个流,同一时刻,10个流的总活跃线程数不能大于47104;

回到我遇到的问题,也就是说,我的核函数开的线程饱和了,用流并发或者嵌套并不能提高我的性能,
请问:
1.我的理解是否有误。
2.是否有方法提高性能。

Re: 关于核函数的流并发 || 核函数的嵌套并发问题
« 回复 #1 于: 八月 13, 2019, 07:57:00 pm »
你的问题用流可能确实没有太多性能提升的空间。正如在群里说过的,我认为可以尝试增大核函数的计算量。

比如:

void __global__ kernel_1(double *x)
{
    int n = blockDim.x * blockIdx.x + threadIdx.x;
    x[n]++;
}

kernel_1<<<grid_size, block_size>>>(x);
kernel_1<<<grid_size, block_size>>>(y);
kernel_1<<<grid_size, block_size>>>(z);

没有

void __global__ kernel_2(double *x, double *y, double *z)
{
    int n = blockDim.x * blockIdx.x + threadIdx.x;
    x[n]++;
    y[n]++;
    z[n]++;
}
kernel_2<<<grid_size, block_size>>>(x, y, z);

高效。



Re: 关于核函数的流并发 || 核函数的嵌套并发问题
« 回复 #2 于: 八月 14, 2019, 01:05:18 pm »
楼上两位说的不错。我稍微补充一点。

楼主的启动线程规模太大,此时并行多个grid无可能了。你的理解也是正确的,如果这种是可能的话,那么最大驻留线程数(也就是某一时刻真正在卡上并行的线程数量)这个指标,将无意义。所以这是不可能的。

此外如同bruce说的,你的确考虑可以将1个线程的工作量加大的(注意这并不总是会提高性能)。

Re: 关于核函数的流并发 || 核函数的嵌套并发问题
« 回复 #3 于: 八月 14, 2019, 02:35:38 pm »
感谢brucefan博士和屠神的回复,我会尝试把单线程工作量增大,尝试是否能提高性能