今天终于将困扰我好久,令我程序优化进程受阻的问题找到了,其实仔细排查不难发现这个问题,但是当从整体来看时,问题很难准确抓到,而且由于基础不深,在这上面耽搁了很久,一边耽搁也在一边学习,黄天不负苦心人,终于让我找到问题所在!今天来谈一谈最近的收获。
CUDA并行代码优化时,我习惯将串行代码原生态在并行方式上实现,由于这是第一次用CUDA,所以并没有加入很多优化方法,在这个最基本的进程中,我觉得有四点需要格外注意。
第一点就是CUDA的内存地址空间大小分配问题以及数组越界问题,一定要保证好数组大小分配足够,检查好数组是否越界,在kernel函数中加入if()语句判断,保证线程不会访问超过数组实际大小的空间。
第二点就是内存分配与回调一定要注意,该存的时候要存好,该读的时候要读好,因为我们做的是异构运算,所以串行并行相互交织,这个问题就显得尤为重要,并且,记得在memcopy时加入设备同步,保证存取完之后再进行下一步操作。
第三点也是让我找到问题关键的点就是数据依赖性问题,这也是最重要的一个问题,并行优化十分要求互相之间没有依赖性,这不仅仅体现在迭代上,更体现在每一个计算上,比如下面的代码:
if (k < m){
dev_rowTotal[k] = (double)((dev_syndrome[k] == 1) ? -1 : 1);
}
//这里涉及一个数据依赖性问题!
//因为dev_ir[k]的取值是在0-m之间,所以会出现重复的情况,也就是说这个运算是一个有顺序的过程,后面的运算结果可能需要前面的结果,其数据存在依赖性!
//所以在这里,我们让其这样并行执行时不正确的,举例即,R数组中一共有m个值,这m个值都会进行不定次的乘等运算,只有最后一次的乘等运算才是其正确值,我们让这一语句并行化,就会造成最后取得的值只是其乘等中的一个值或一部分值,不能保证正确性。
//要解决这个问题,必须解决其数据依赖性!
__syncthreads();
dev_rowTotal[dev_ir[k]] *= dev_check_LLR[k];
这就是一个数据依赖性的典型例子,所以我们在并行优化的时候一定要谨慎仔细,确保每一语句每一逻辑,其是否存在依赖性。
第四点是同步问题,由于kernel中有许多前后调用的关系,所以保证栅栏同步很重要!