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

这是我写的一个 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也能发射成功运行;
求助!!求解!!

楼主的“发射”指得是启动/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;

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

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

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

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

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

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

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

额!

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

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

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

嘟嘟

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

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

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

:slight_smile:

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

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

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

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

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