这段代码为什么得不到正确的结果？求助

加速计算 CUDA
1 
__device__ int3 Get3AdjacentTri(int center,int level)
{
   int octaTriCount = 1 << (level << 1);
   int octa=center>>(2*level);//八分体编号
   int octaID=center-(octa<<(2*level));//该ID在八分体内的编号
   int row=sqrtf((float)(octaID+0.5f)); //在八分体内的行号
   int col=octaID-row*row; //在八分体内的列号
   int rowCount=1<<level; //八分体的总行数
   int colCount=2*row+1; //所在行的总列数
   int leftTri = -1, rightTri = -1, topTri = -1;
   int rightOcta = -1;
   if (octa == 0 || octa == 1 || octa == 2)
   rightOcta = octa + 1;
   if (octa == 5 || octa == 6 || octa == 7)
   rightOcta = octa - 1;
   else if (octa == 3)
   rightOcta = 0;
   else if (octa == 4)
   rightOcta = 7;

   int leftOcta = -1;
   if (octa == 1 || octa == 2 || octa == 3)
   leftOcta = octa - 1;
   if (octa == 4 || octa == 5 || octa == 6)
   leftOcta = octa + 1;
   else if (octa == 0)
   leftOcta = 3;
   else if (octa == 7)
   leftOcta = 4;

   int topOcta = -1;
   if (octa < 4)
   {
   topOcta = octa + 4;
   }
   else
   {
   topOcta = octa - 4;
   }

   if (row == 0)//三角形B
   {
   rightTri = rightOcta * octaTriCount;
   leftTri = leftOcta * octaTriCount;
   topTri = 2;
   }
   else if (row > 0 && row < rowCount - 1 && col == 0)//三角形E
   {
   rightTri = center + 1;
   topTri = center + colCount + 1;
   leftTri = leftOcta * octaTriCount + octaID + colCount - 1;
   }
   else if (row == rowCount - 1 && col == 0)//三角形C
   {
   rightTri = center + 1;
   leftTri = leftOcta * octaTriCount + octaID + colCount - 1;
   topTri = topOcta * octaTriCount + octaTriCount - 1;
   }
   else if (row == rowCount - 1 && col > 0&&col<colCount-1 && col % 2 == 0)//三角形G
   {
   leftTri = center - 1;
   rightTri = center + 1;
   topTri = topOcta * octaTriCount + octaTriCount - col - 1;
   }
   else if (row == rowCount - 1 && col == colCount - 1)//三角形D
   {
   leftTri = center - 1;
   rightTri = rightOcta * octaTriCount + octaTriCount - colCount;
   topTri = topOcta * octaTriCount + octaTriCount - colCount;
   }
   else if (row > 0 && row < rowCount - 1 && col == colCount - 1)//三角形F
   {
   leftTri = center - 1;
   topTri = center + colCount + 1;
   rightTri = rightOcta * octaTriCount + octaID - colCount + 1;
   }
   else
   {
   leftTri = center - 1;
   rightTri = center + 1;
   if (col % 2 == 0)
   {
   topTri = center + colCount + 1;
   }
   else
   {
   topTri = center - colCount + 1;
   }
   }
   int3 neibhbour;
   neibhbour.x=leftTri;
   neibhbour.y=rightTri;
   neibhbour.z=topOcta;
   return neibhbour;
}

这是一段求QTM格网邻近三角形的代码，我想用CUDA实现，但计算结果总是不对，这个过程需不需要用原子操作啊，请高手帮忙
注：同时有很多个线程执行这些代码

2

我想当农民赚钱，可是我种植了树后，春天施肥，夏天浇水，却没有挣到，怎么回事啊？这个过程能挣钱的的啊。

注：我很认真的工作。

3

你好，谢谢你的回答，不过我不太明白你的意思！
我想问的是，我这个程序里用了很多的if else，这样会不会对CUDA程序有影响，这里需不需要用原子操作

4

LZ你好：

您的问题无法回答，您这样既无算法实现意图，也无其他信息，上来就贴一段代码，而且还只是一个__device__函数，而且目测是从CPU代码里面没怎么改直接放进去的，然后上来就问，这代码为何不对？

天知道为何不对。

这里面有无数因素能影响您的结果，您让斑竹和网友们一一往出猜么？

至于您3#补充的内容，同样无法回答。
1：if-else会不会对CUDA程序有影响？——要说实现逻辑，您写对了就无影响，CUDA和C一样提供了同样用法的if-else判断。但是如果您不是按照多线程的角度去考虑，可能会出各种问题。以及要说运行效率，那么如果算法合适或者实现较好，效率是比较好的，否则会严重影响效率。

2：if-else判断和原子操作无任何直接联系。用if-else的不一定要用原子操作，用原子操作的不一定要用if-else判断。一个使用原子操作的算法实现也经常可以修改为不用原子操作的实现。这一切都取决于对算法实现的安排和权衡以及算法自身的一些情况。

然后您就这样劈头盖脸地来一句，“这里需不需要用原子操作”。
算法是您自己设计的，需不需要用原子操作您应该比我清楚，我最多只能告诉您原子操作自身是干什么用的，而绝对不具备某种超能力自动脑补您的一切详细信息并给出您心目中的答案。

5

感谢版主的耐心回答，是我对问题描述的不太清楚！！！现在已经找到出现错误的原因，跟if-else及原子操作都没有关系，是由于本人粗心大意，将代码第93行写错，正确的应为neighbour.z=topTri;实在是太不应该了！
我之所以问if-else是在GPU中执行时否会有影响，是因为我看过一些资料，说CUDA的分支预测、控制不如CPU那么好，所以我觉得用了这么多的分支控制语句，是否会对执行的结果产生影响！现在知道了原因不在这儿了，但使用过多的这种语句是不是会对运行效率有影响呢，如果在程序中确实要用有很多分支，应该怎么样改进呢？
至于原子操作，我只是稍有了解，不知道什么时候必须要使用，代码得不到正确的结果，我又想会不会是该使用的时候我没有使用。有点乱投医的感觉了！
再次感谢您的帮助！！！我以后再问问题尽量把问题说的清楚一点！

6

LZ您好：

1：关于GPU中分支的效率问题，简单地说，如果所有分支都是按warp对齐的，那么并无大碍。因为GPU运行中，单个warp总是同时执行的，如果warp内部有分支（即一个warp内部的线程部分走if，部分走else），那么可能会出现假执行的方式维护逻辑的正确性，即整体运行两次，各个线程丢弃不需要的结果；也可能出现部分线程等待的执行方式。
以及这样的话会带来一些性能的损失。所以如果分支的行为是确定的，那么尽量通过编程技巧使分支按照warp对齐；如果分支十分复杂或者是预先无法知道的随机性分支，那么程序效率可能会受到很大的影响，以及在随机分支的时候，也常常伴有各种无法合并的访存。

您可以参考论坛中以前的讨论帖：
http://cudazone.nvidia.cn/forum/forum.php?mod=viewthread&tid=5963

http://cudazone.nvidia.cn/forum/forum.php?mod=viewthread&tid=6391&extra=&page=1

论坛暂不提供搜索功能，我先只找到这两个帖子。

7

2：关于原子操作，简单地说您可以把握两点。

1）原子操作自身的意义，原子操作指的是不可分割的操作，他会把一系列操作打包成一个整体，这样中途不会被打断。我们考虑一个多线程情况下的累加。如果线程a先读取某变量cat，然后+1，回写；在a回写前，线程b也读取cat，并+1，回写。那么实际上变量cat只被累加了一次，有一个线程没有正确累加上去。如果使用原子操作，a先读取了cat变量，那么a回写完成之前，b是无法读取cat的，这就避免了多线程情况下之前的问题。

2）原子操作一般在程序中用于维护某种强制串行执行的功能。但是注意，只保证串行，而不保证顺序。比如前面的累加，就必须一个线程干完活，另外一个线程才能继续。

此外，原子操作通常是比较慢的，如果各个GPU线程（GPU线程数量通常较大）都直接使用原子操作，那么性能会受到很大影响。可以考虑在block内部利用block内部的通信和同步手段，先将局部的结果予以整理和组合，最后一次性累加到全局的变量中，这样效果较好。

还有一个需要注意的地方是，原子操作的写入结果并不与常规读取保持一致。
也就是如果您对变量cat进行原子累加，以及对cat进行普通读取，那么您读取到的cat的值可能不是cat的真实值。您需要使用原子操作的返回值来得到cat的真实值。

大致如上，原子操作的不同计算能力版本GPU的支持情况以及如何手工实现更多原子操作功能，请参阅programming guide。

祝您好运~

8

从您的回答中，我对在GPU中使用分支及原子操作有了更多的了解，再次感谢！
祝您一切顺利！！！

9

不客气的，欢迎您常来论坛~

祝您编码顺利！

10

感谢超版将本帖移回开发版，使之能继续为其他网友提供参考。

让我们努力共创和谐向上的论坛环境！