核函数写global memory问题

  • 6 replies
  • 544 views
核函数写global memory问题
« 于: 十二月 11, 2019, 04:46:51 pm »
核函数中部分程序如下

__syncthreads();
if(blockIdx.x<48)
{
     offset = blockIdx.x * 64 + threadIdx.x;
     PQ_bus_cal[offset]=blockIdx.x;
}

PQ_bus_cal为float* 型指针,指向global memory
该核函数调用了64个BLOCK每个BLOCK 64线程,前面的运算需要用完所有block,但是这一步只需要48个。

函数执行以后调用cudaDeviceSynchronize(),然后用cudamemcpy复制该块内存至HOST,cout出来以后发现,0~44 block正常输出,45~47没有对global memory写数据(数据还是上一步运算的结果)

请问一下有地方可以查一下
« 最后编辑时间: 十二月 11, 2019, 04:48:00 pm 作者 jetlin1992 »

Re: 核函数写global memory问题
« 回复 #1 于: 十二月 12, 2019, 08:10:27 pm »
__syncthreads()并不能将你的64个blocks,切分成前面一部分干活,和后面一部分干活的。它只有block内部的同步效果,并不能将多个blocks的活动,同步然后切换成逻辑的两部分的。

建议的解决方案:
(1)看看能否通过两次kernel启动来完成,每个kernel只做前一部分和后一部分(推荐)。
(2)查看Cooperative Launch能否满足你的64个blocks同时在设备上驻留,然后全局同步(不推荐)。(和具体的卡有关,卡的SM数量越多,越容易成功。如果不能满足同时驻留,则会启动失败,同时无法全局同步)。


Re: 核函数写global memory问题
« 回复 #2 于: 十二月 13, 2019, 08:22:37 am »
__syncthreads()并不能将你的64个blocks,切分成前面一部分干活,和后面一部分干活的。它只有block内部的同步效果,并不能将多个blocks的活动,同步然后切换成逻辑的两部分的。

建议的解决方案:
(1)看看能否通过两次kernel启动来完成,每个kernel只做前一部分和后一部分(推荐)。
(2)查看Cooperative Launch能否满足你的64个blocks同时在设备上驻留,然后全局同步(不推荐)。(和具体的卡有关,卡的SM数量越多,越容易成功。如果不能满足同时驻留,则会启动失败,同时无法全局同步)。

全局同步一般要怎么做呢,同时驻留应该是满足的

Re: 核函数写global memory问题
« 回复 #3 于: 十二月 13, 2019, 12:56:50 pm »
看官方的《CUDA Programmig Guide》(CUDA 9+)的附录C.2.5: Uses of intra-block cooperative groups. 这个限制也是超多,我也不建议使用。

Re: 核函数写global memory问题
« 回复 #4 于: 十二月 13, 2019, 01:40:57 pm »
看官方的《CUDA Programmig Guide》(CUDA 9+)的附录C.2.5: Uses of intra-block cooperative groups. 这个限制也是超多,我也不建议使用。
谢谢

Re: 核函数写global memory问题
« 回复 #5 于: 十二月 13, 2019, 02:56:17 pm »
你最快的方案是重开kernel。

前一个kernel的整体死亡,和后一个kernel的整体重新开始,是最容易使用的“全局同步“方式。(然而你可能认为这个方法土,就无视了我)

Re: 核函数写global memory问题
« 回复 #6 于: 十二月 13, 2019, 06:32:32 pm »
你最快的方案是重开kernel。

前一个kernel的整体死亡,和后一个kernel的整体重新开始,是最容易使用的“全局同步“方式。(然而你可能认为这个方法土,就无视了我)

没有没有,最后我也是决定重开kernel,只是想看看有没有更高效的方法