CUDA程序优化时中关于并行性杂谈。

  今天终于将困扰我好久，令我程序优化进程受阻的问题找到了，其实仔细排查不难发现这个问题，但是当从整体来看时，问题很难准确抓到，而且由于基础不深，在这上面耽搁了很久，一边耽搁也在一边学习，黄天不负苦心人，终于让我找到问题所在！今天来谈一谈最近的收获。

  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中有许多前后调用的关系，所以保证栅栏同步很重要！