这是我写的一个 kerner,kerner发射总是失败，求助

加速计算 CUDA
1

这是我写的一个 kerner,

pd_DEMCellBox, pd_DES_POS, pd_nebrPPTab, pd_nebrTabLen, d_vOff均为全局内存上的变量，确保已经cudaMalloc成功。

global void SearchOverCells(long *pd_DEMCellBox, int demcellsx,

                        int demcellsy, int demcellsz,

                        float *pd_DES_POS,const float rrNebr,

                        int particles, long *pd_nebrPPTab,

                        int *pd_nebrTabLen, int *d_vOff)

{

int i = blockDim.x * blockIdx.x + threadIdx.x;

int ci1, cj1, ci2, cj2;

int m2x, m2y, m2z, m1x, m1y, m1z;

long LL, JJ;

float RelPos[3], DIST;

cj1 = pd_DEMCellBox[i * 81];

if(cj1 == 0)

   return;

for(int j=0; j<14; j++)

{

   m1x = (i % (demcellsx * demcellsy)) % demcellsx;

   m2x = m1x + d_vOff[j*3];

   if(m1x == 0)

       m1y = (i % (demcellsx * demcellsy)) / demcellsx;

   else

       m1y = (i % (demcellsx * demcellsy)) / demcellsx + 1;

   m2y =  + d_vOff[j*3 + 1];

   if(i % (demcellsx * demcellsy) == 0)

       m1z = i / (demcellsx * demcellsy);

   else

       m1z = i / (demcellsx * demcellsy) + 1; 

   m2z = m1z + d_vOff[j*3 + 2];

   if(m2x < 0 || m2x >= demcellsx ||

      m2y < 0 || m2y >= demcellsy ||

      m2z < 0 || m2z >= demcellsz)

      continue;

   int index = demcellsx * demcellsy * m2z + m2x * m2y + m2x; 

   cj2 = pd_DEMCellBox[index * 81];

   if(cj2 == 0)

       continue;

   for(ci1=1; ci1<=cj1; ci1++)

   {

       for(ci2=1; ci2<=cj2; ci2 ++)

       {

          LL = pd_DEMCellBox[i*81 + ci1];

          JJ = pd_DEMCellBox[index*81 + ci2];

          if ( m1x!=m2x || m1y!=m2y || m1z!=m2z || LL > JJ)

          {

              RelPos[0] = pd_DES_POS[JJ] - pd_DES_POS[LL];

RelPos[1] = pd_DES_POS[particles + JJ] - pd_DES_POS[particles + LL];

              if(fabs(RelPos[0])>rrNebr || fabs(RelPos[1])>rrNebr) 

                 continue;

RelPos[2] = pd_DES_POS[2particles + JJ] - pd_DES_POS[2particles + LL];

              DIST = ABS_V3(RelPos);

              if(DIST <= rrNebr )

              {

                  pd_nebrPPTab[(*pd_nebrTabLen) * 2] = LL;

                  pd_nebrPPTab[(*pd_nebrTabLen) * 2 + 1] = JJ;

                  (*pd_nebrTabLen)++;

              }

          }

       }

   }

}

}

这个kerner老是发射不成，本人做过如下测试。

红色部分代码屏蔽掉后，kerner能发射成功；

将红色部分之外的代码屏蔽掉，kerner内只留红色体代码，kerner也能发射成功运行；
求助!!求解!!

2

楼主的“发射”指得是启动/launch/<<<>>>吗？

如果是，那么可能是你的kernel有BUG.

我仔细找了一下您的“红色”部分，发现这里可能不妥当：
多个线程竞争读写pd_nebrTabLen里的值。

我建议您尝试使用原子读写，以保证您的安全。

我举个例子，将您对pd_nebrPPTab头部的计数的读写，
pd_nebrPPTab[(*pd_nebrTabLen) * 2] = LL;
pd_nebrPPTab[(*pd_nebrTabLen) * 2 + 1] = JJ;
(*pd_nebrTabLen)++;
改为：
your_type(你没有在代码里给出) count = atomicAdd(pd_nebrTabLen, 1);
pd_nebrTab[count * 2] = LL;
pd_nebrTab[count * 2 + 1] = JJ;

可以参考上述修改避免多个线程的竞争读写。

至于其他问题，欢迎其他会员/版主继续回帖。
祝您调试愉快！

3

此外，您还缺少对您的pd_nebrPPTab缓冲区头部进行初始化的过程。

如果不初始化，可以会导致头部的计数是随机值。

如果您已经初始化，但没有贴出代码，请忽略此条。

4

另外，觉得这个kernel里面循环套循环，各种分支判断，似乎复杂了点，能否拆分改下算法什么的。

5

“pd_nebrPPTab缓冲区头部进行初始化”, 斑竹大侠的意思是不是pd_nebrPPTab指向的内存块要初始化。

6

额！

我看错了。我昨日没看到Len后缀。我以为这是*pd_nebrPPTab(即第一个元素）。
不好意思（变量名有点长。。看混2个名字了。。。）

既然你用了单独的变量作为长度。那么只初始化这个变量好了。不用初始化头部了。
原帖已经修改。请重新阅读。

请无视第二个回帖。
只看第一个即可。（原子操作的建议）。

嘟嘟

7

我觉得显卡的处理器应该还能处理这些循环吧，想简单点实现

8

请楼主先尝试修改正确，然后才考虑是否需要优化的才好。

现在已经一天过去了，楼主您修改的如何了？

:slight_smile:

9

建议在kernel调用后加上以下代码看看能否输出具体错误信息：

printf(“%s\n”, cudaGetErrorString(cudaGetLastError()));

更好的方法是建议用cuda-gdb跑一次你的程序，看看输出的具体错误信息是什么。

方法：
nvcc -g -G file.cu
cuda-gdb ./a.out

10

王总亲自上阵了~~~:victory: