一个有关并行化的问题

  • 3 replies
  • 430 views
一个有关并行化的问题
« 于: 四月 16, 2019, 04:07:08 pm »
想要并行化的函数主要处理下面这个问题,
对于一个二维矩阵中的每一个位置点,需要进行一次计算得到一个结果,但是每个位置的结果要依赖于当前点位置上一列或上一行的结果,也就是说存在一定的依赖关系,所以并不能一个线程来计算一个位置点。现在有一个想法是这样的:
首先CPU中计算得到矩阵第一行的结果,然后设置一个线程负责一列的计算,当计算完列中的一个点之后循环向下计算下一个点,直到这一列计算结束。但现在的问题是单个位置的结果计算就比较复杂,如果一个线程来负责一列的计算的话计算量会过大,加速效果并不明显。卡是gtx1060,不知道有没有类似的更好一点的思想来处理这个问题来达到更高的并行度。

还一个问题是程序在Nsight下可以调试输出结果,但是主机端跑的时候貌似会跳过核函数未进行计算,此程序之前可以正常运行得到结果,但后来不知道什么原因就突然不能正常跑了。比较玄学
想知道有什么原因会导致上面这个问题,
麻烦[名词6]答疑,谢谢啦

Re: 一个有关并行化的问题
« 回复 #1 于: 四月 16, 2019, 05:04:38 pm »
想要并行化的函数主要处理下面这个问题,
对于一个二维矩阵中的每一个位置点,需要进行一次计算得到一个结果,但是每个位置的结果要依赖于当前点位置上一列或上一行的结果,也就是说存在一定的依赖关系,所以并不能一个线程来计算一个位置点。现在有一个想法是这样的:
首先CPU中计算得到矩阵第一行的结果,然后设置一个线程负责一列的计算,当计算完列中的一个点之后循环向下计算下一个点,直到这一列计算结束。但现在的问题是单个位置的结果计算就比较复杂,如果一个线程来负责一列的计算的话计算量会过大,加速效果并不明显。卡是gtx1060,不知道有没有类似的更好一点的思想来处理这个问题来达到更高的并行度。

还一个问题是程序在Nsight下可以调试输出结果,但是主机端跑的时候貌似会跳过核函数未进行计算,此程序之前可以正常运行得到结果,但后来不知道什么原因就突然不能正常跑了。比较玄学
想知道有什么原因会导致上面这个问题,
麻烦[名词6]答疑,谢谢啦

如果矩阵每一行算完了,才能继续下一行,正常情况下,我建议你使用CPU进行计算。真的。

但如果你同时有多种这样的矩阵,则可以考虑用GPU(等效于任务并行了)。

关于你的第二个问题,光这样看,看不出什么,我建议你执行如下步骤:
your_kernel<<<>>>(....);
cudaError_t r0 = cudaGetLastError();
cudaError_t r1 = cudaDeviceSynchronize();
分别检查r0和r1是否为cudaSuccess(0), 从而确定是否真的“kernel并未执行”,或者中途挂掉,而不是凭感觉。

Re: 一个有关并行化的问题
« 回复 #2 于: 四月 16, 2019, 07:58:09 pm »
多谢,

第一个问题确实是前一行依赖于后一行或者前一列依赖于后一列,您说的我大概了解了

关于第二个问题,已经试过在核函数后加
printf("%s\n",cudaGetErrorString(cudaGetLastError()));
打印结果为“no error”;

对核函数计时得到的时间为负数(有个细节是计时同步的时候卡顿了一下,然后就过去了);

另外,命令行中尝试了cuda-memcheck。结果首先显示no error,随后显示某线程越界访问(读写纹理内存上时);

还有一点就是用Nsight调试,能进入核函数内部,但会崩(因为有太多循环,所以还没想到好的方法去调);

您上面给出的两个debug代码中,cudaGetLastError()没问题,但是cudaDeviceSynchronize()的值为cudaErrorIllegalAddress(77);

Re: 一个有关并行化的问题
« 回复 #3 于: 四月 17, 2019, 03:55:11 pm »
多谢,

第一个问题确实是前一行依赖于后一行或者前一列依赖于后一列,您说的我大概了解了

关于第二个问题,已经试过在核函数后加
printf("%s\n",cudaGetErrorString(cudaGetLastError()));
打印结果为“no error”;

对核函数计时得到的时间为负数(有个细节是计时同步的时候卡顿了一下,然后就过去了);

另外,命令行中尝试了cuda-memcheck。结果首先显示no error,随后显示某线程越界访问(读写纹理内存上时);

还有一点就是用Nsight调试,能进入核函数内部,但会崩(因为有太多循环,所以还没想到好的方法去调);

您上面给出的两个debug代码中,cudaGetLastError()没问题,但是cudaDeviceSynchronize()的值为cudaErrorIllegalAddress(77);

(1)请严格按照我的两次严查要求来进行计算。不要认为你的那种写法的"no error"表示正常。不接受建议对你不利。

(2)kernel启动的常规检查需要两种,一种是对kernel启动本身,看看kernel能否正常启动和极快的速挂,这是用的跟随在<<<>>>后面的立刻cudaGetLastError得到的,另外始终则是运行中的,稍微往后出现的错误,此时可以在下一次同步调用的时候得到出错代码,例如之前建议的cudaDeviceSynchronize()时候的返回值。

(3)按照这两步要求得到了第一个代码成功,第二个代码无效访存,则应当看后者。此时表明你的kernel的确挂了。原因是访存问题。cuda-memcheck也提示了这点,表明是实情的确如此,你得去修改bug了。

(4)nsight请启用memory checker,然后在遇到第一个访存出错的地方,会自动停掉(你需要debug编译)

去改吧。