从Device端传回Host端的时候cudaMemcpy函数不起作用

  • 13 replies
  • 1394 views
从Device端传回Host端的时候cudaMemcpy函数不起作用
« 于: 十一月 02, 2022, 10:41:28 am »
请问一下我在GPU上计算完结果之后,想要传回CPU端,写了如下代码
程序代码: [选择]
void CSPH::ParticlesDataReturn()
{
CHECK(cudaMemcpy(TotalParCoor, d_TotalCoor, sizeof(double3)* nTotalPar, cudaMemcpyDeviceToHost));
CHECK(cudaMemcpy( Press,d_Press, sizeof(float)*nTotalPar, cudaMemcpyDeviceToHost));
CHECK(cudaMemcpy(velrhop,d_Velrhop , sizeof(float4)*nTotalPar, cudaMemcpyDeviceToHost));
CHECK(cudaMemcpy(SortPart, d_SortPart, sizeof(int)*nTotalPar, cudaMemcpyDeviceToHost));
CHECK(cudaMemcpy(RepForce, d_RepForce, sizeof(float3)*nFluidPar, cudaMemcpyDeviceToHost));
cudaError_t error = cudaGetLastError();
printf("CUDA error: %s\n", cudaGetErrorString(error));

for(int i=0;i<nFluidPar;i++)
{
if(SortPart[i]==0)
printf("%d, %f, %f, %f aaaaa\n ",i,RepForce[SortPart[i]].x,RepForce[SortPart[i]].y,RepForce[SortPart[i]].z);

}
}
在传回的这些数据中心,只有RepForce[].x,RepForce[].y,RepForce[].z的值还跟程序刚开始在CPU端赋的初值相同,而其他数组的值输出来就是在GPU上计算完成后的值。我在核函数中输出了d_RepForce[]的值是对的,感觉是从GPU传回CPU的时候cudaMemcpy函数没起作用,并且用cudaGetLastError();也显示no error。请问这是怎么回事呀

EDIT: 加上了代码的BBCode标注。
« 最后编辑时间: 十一月 03, 2022, 11:35:57 am 作者 屠戮人神 »

Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 回复 #1 于: 十一月 02, 2022, 12:35:15 pm »
请问一下我在GPU上计算完结果之后,想要传回CPU端,写了如下代码
void CSPH::ParticlesDataReturn()
{
   CHECK(cudaMemcpy(TotalParCoor, d_TotalCoor, sizeof(double3)* nTotalPar, cudaMemcpyDeviceToHost));
   CHECK(cudaMemcpy( Press,d_Press, sizeof(float)*nTotalPar, cudaMemcpyDeviceToHost));
   CHECK(cudaMemcpy(velrhop,d_Velrhop , sizeof(float4)*nTotalPar, cudaMemcpyDeviceToHost));
   CHECK(cudaMemcpy(SortPart, d_SortPart, sizeof(int)*nTotalPar, cudaMemcpyDeviceToHost));
   CHECK(cudaMemcpy(RepForce, d_RepForce, sizeof(float3)*nFluidPar, cudaMemcpyDeviceToHost));
   cudaError_t error = cudaGetLastError();
printf("CUDA error: %s\n", cudaGetErrorString(error));
   
for(int i=0;i<nFluidPar;i++)
{
   if(SortPart==0)
printf("%d, %f, %f, %f aaaaa\n ",i,RepForce[SortPart].x,RepForce[SortPart].y,RepForce[SortPart].z);

}
}
在传回的这些数据中心,只有RepForce[].x,RepForce[].y,RepForce[].z的值还跟程序刚开始在CPU端赋的初值相同,而其他数组的值输出来就是在GPU上计算完成后的值。我在核函数中输出了d_RepForce[]的值是对的,感觉是从GPU传回CPU的时候cudaMemcpy函数没起作用,并且用cudaGetLastError();也显示no error。请问这是怎么回事呀

这种有可能的,最常见的情况是kernel启动配置失败(没有执行),而你没有使用经典的2次检查的方式来捕获错误,也就是是根据在Kernel启动<<<>>>后面的立刻一次CHECK(cudaGetLastError()); 和在后续的最近一次地方的同步调用(例如对cudaXXXSynchronize()或者其他同步操作的CHECK())。这两个分别代表检测"kernel是否能启动", 和"kernel启动后,后续执行是否成功(例如没有因访存而挂掉).

常见的是漏掉前者,而前者的错误可能会被后续的其他操作的某个地方的成功给覆盖掉。而只有后者的错误才是持续能返回的。所以往往常见认为kernel已经执行好了,结果没动弹,往往漏在这里。

当然,不能排除其他方面的原因,例如是否你的kernel真的啥都没干,不过前面说的是比较常见的情况,可以先排查下。

Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 回复 #2 于: 十一月 02, 2022, 12:38:17 pm »
这种有可能的,最常见的情况是kernel启动配置失败(没有执行),而你没有使用经典的2次检查的方式来捕获错误,也就是是根据在Kernel启动<<<>>>后面的立刻一次CHECK(cudaGetLastError()); 和在后续的最近一次地方的同步调用(例如对cudaXXXSynchronize()或者其他同步操作的CHECK())。这两个分别代表检测"kernel是否能启动", 和"kernel启动后,后续执行是否成功(例如没有因访存而挂掉).

常见的是漏掉前者,而前者的错误可能会被后续的其他操作的某个地方的成功给覆盖掉。而只有后者的错误才是持续能返回的。所以往往常见认为kernel已经执行好了,结果没动弹,往往漏在这里。

当然,不能排除其他方面的原因,例如是否你的kernel真的啥都没干,不过前面说的是比较常见的情况,可以先排查下。

以及,不妨就在"这个kernel"中使用printf,看看能否发现值已经被修改了。注意不是在另外一个kernel中,也不是观察"原始输入值", 这样可以很大程度的避免一些幻觉。

Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 回复 #3 于: 十一月 02, 2022, 08:05:20 pm »
您好,我在这个核函数里printf了确实是计算了,以及在调用这个核函数之后用了cudaGetLastError();仍然显示no error,您看这是怎么回事呢
void CalFluidRepForce(int n, double3* TotalCoorg, int3* CCg, int* nCubeg, int* SortPartg, float* OuterReigong, int2* BeginEndCubeg, float3* RepForceg)
   {
      unsigned gridesize = unsigned(n + SPHSIZE - 1) / SPHSIZE;
      KerCalFluidRepForce << <gridesize, SPHSIZE >> >(n, TotalCoorg, CCg, nCubeg, SortPartg, OuterReigong,BeginEndCubeg, RepForceg);      
      
      cudaError_t error12 = cudaGetLastError();
printf("CUDA error12: %s\n", cudaGetErrorString(error12));
      CHECK(cudaDeviceSynchronize());
      cudaError_t error26 = cudaGetLastError();
printf("CUDA error26: %s\n", cudaGetErrorString(error26));
   }

Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 回复 #4 于: 十一月 03, 2022, 09:54:18 am »
您好,我在这个核函数里printf了确实是计算了,以及在调用这个核函数之后用了cudaGetLastError();仍然显示no error,您看这是怎么回事呢
void CalFluidRepForce(int n, double3* TotalCoorg, int3* CCg, int* nCubeg, int* SortPartg, float* OuterReigong, int2* BeginEndCubeg, float3* RepForceg)
   {
      unsigned gridesize = unsigned(n + SPHSIZE - 1) / SPHSIZE;
      KerCalFluidRepForce << <gridesize, SPHSIZE >> >(n, TotalCoorg, CCg, nCubeg, SortPartg, OuterReigong,BeginEndCubeg, RepForceg);      
      
      cudaError_t error12 = cudaGetLastError();
printf("CUDA error12: %s\n", cudaGetErrorString(error12));
      CHECK(cudaDeviceSynchronize());
      cudaError_t error26 = cudaGetLastError();
printf("CUDA error26: %s\n", cudaGetErrorString(error26));
   }

不妨发一下你的kernel的代码,然后在kernel后面立刻取回(cudaXXXSynchronize()这里可以省略,如果你后面立刻同步取回的话),而不要在另外一个函数里,避免你的CPU上有代码中间干别的事情,而对你造成了错觉。然后我看下。

Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 回复 #5 于: 十一月 03, 2022, 10:07:33 am »
非常感谢,我的Kernel是:
__global__ void KerCalFluidRepForce(int n, double3* totalcoor, int3* cc, int* ncube, int* sortpart, float* outerReigon, int2* beginendcube, float3* Repforce)
   {
      const int p = blockIdx.x*blockDim.x + threadIdx.x; //-Number of particle.

      if (p < n)//n 指的是全体粒子数
      {

         const int p1 = p;
         double shift = 0.0;
      
      if(sortpart[p1] >= 0 && sortpart[p1] <CONSTANTS.nBound2Start)
{
Repforce[sortpart[p1]].x=0.0;
Repforce[sortpart[p1]].y=0.0;
Repforce[sortpart[p1]].z=0.0;
}


         //-Obtains basic data of particle p1.
         const double posp1x = __ldg(&totalcoor[p1].x);//获取当前p1粒子位置
         const double posp1y = __ldg(&totalcoor[p1].y);//获取当前p1粒子位置
         const double posp1z = __ldg(&totalcoor[p1].z);//获取当前p1粒子位置
   
                  
         //-Obtains neighborhood search limits.获取邻居搜索的极限
         int ini1, fin1, ini2, fin2, ini3, fin3;//三个方向上的极限,

         ini1 = cc[p1].x - 1;//cc.x为粒子所在格子的x方向坐标
         fin1 = cc[p1].x +1;

         ini2 = cc[p1].y - min(cc[p1].y, 1);//cc.y为粒子所在格子的y方向坐标
         fin2 = cc[p1].y + ((ncube[1] - cc[p1].y - 1)>0 ? 1 : 0);

         ini3 = cc[p1].z - min(cc[p1].z, 1);//cc.z为粒子所在格子的z方向坐标
         fin3 = cc[p1].z + ((ncube[2] - cc[p1].z - 1)>0 ? 1 : 0);

         //-Interaction with Particles.
         for (int m3 = ini3; m3 < fin3 + 1; m3++)
         {
            for (int m2 = ini2; m2 < fin2 + 1; m2++)
            {
               for (int m1 = ini1; m1 < fin1+1; m1++)
               {
                  //periodic condition on the x directions
                  int k1;
                  if (m1 == -1)
                  {
                     k1 = ncube[0] - 1;
                     shift = outerReigon[0];
                  }
                  else if (m1 == ncube[0])
                  {
                     k1 = 0;
                     shift = -outerReigon[0];
                  }
                  else
                  {
                     k1 = m1;
                     shift = 0.0;
                  }

                  int pini, pfin = 0;
                  cusph::ParticleRange(beginendcube, ncube, k1, m2, m3, pini, pfin);
                           
                                              for (int p2 = pini; p2 < pfin; p2++)
                  {
                     if (sortpart[p1] >= 0 && sortpart[p1] <CONSTANTS.nBound2Start && sortpart[p2] >= CONSTANTS.nBound1Start && sortpart[p2] < CONSTANTS.nTotalPar)//sortpart[p1]-FluidPar,sortpart[p2]-Bound1Par
                     {
                                       
                        double FRep;
                        double chi, eta, feta;
                        float h = 0.5;//initial distance between the boundary particle and fluid particle
                        float RepCoef = 0.01;

                        const double posp2x = __ldg(&totalcoor[p2].x);//获取当前p2粒子位置
                        const double posp2y = __ldg(&totalcoor[p2].y);//获取当前p2粒子位置
                        const double posp2z = __ldg(&totalcoor[p2].z);//获取当前p2粒子位置

                        double drx = posp1x - posp2x + shift;
                        double dry = posp1y - posp2y;
                        double drz = posp1z - posp2z;
                        double rr2 = drx*drx + dry*dry + drz*drz;
                        double rr = sqrt(rr2);

                        if (rr < 1.5*h)
                           chi = 1.0 - rr / (1.5*h);
                        else
                           chi = 0.0;
                        eta = rr / (0.75*h);
                        if (eta < 2.0 / 3.0)
                           feta = 2.0 / 3.0;
                        else if (eta < 1.0)
                           feta = 2.0 * eta - 1.5*eta*eta;
                        else if (eta < 2.0)
                           feta = 0.5*(2 - eta)*(2 - eta);
                        else
                           feta = 0.0;
                        if (chi>0 && feta>0)
                        {
                           FRep = RepCoef*CONSTANTS.SoundSpeed*CONSTANTS.SoundSpeed*chi * feta / (rr*rr) *CONSTANTS.Mass;


                           Repforce[sortpart[p1]].x += FRep*drx;
                           Repforce[sortpart[p1]].y += FRep*dry;
                           Repforce[sortpart[p1]].z += FRep*drz;
                              
                        }
                  
                  
                     }//sortpart[p1]-FluidPar,sortpart[p2]-Bound1Par
                  }//p2
               }//m1
            }//m2
         }//m3

if(p1==3203)
      printf("%d, %f,%f ,%f BBBB\n ",sortpart[p1],Repforce[sortpart[p1]].x,Repforce[sortpart[p1]].y,Repforce[sortpart[p1]].z);

      }//if (p < n)
   }


它是计算粒子排斥力的一个核函数,我在最后输出的if(p1==3203)就是输出经过核函数计算之后的sortpart[p1]粒子的排斥力三个分量,这个输出可以得到正确结果

Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 回复 #6 于: 十一月 03, 2022, 10:11:10 am »
还有一个问题是,我在这个核函数刚开始写的赋值代码
 if(sortpart[p1] >= 0 && sortpart[p1] <CONSTANTS.nBound2Start)
{
Repforce[sortpart[p1]].x=0.0;
Repforce[sortpart[p1]].y=0.0;
Repforce[sortpart[p1]].z=0.0;
}
对最后的计算结果影响很大,加上他就是对的,不加就是错的,但是我明明在CPU端已经给他赋值为0了再传回到GPU端上,按说在这里是否赋值为0没有影响才对的吧

Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 回复 #7 于: 十一月 03, 2022, 11:41:32 am »
(1)

CPU端的检查代码:

程序代码: [选择]
for(int i=0;i<nFluidPar;i++)
{
    if(SortPart[i]==0)
        printf("%d, %f, %f, %f aaaaa\n ",i,RepForce[SortPart[i]].x,RepForce[SortPart[i]].y,RepForce[SortPart[i]].z);
}

GPU端的检查代码:
程序代码: [选择]
if(p1==3203)
      printf("%d, %f,%f ,%f BBBB\n ",sortpart[p1],Repforce[sortpart[p1]].x,Repforce[sortpart[p1]].y,Repforce[sortpart[p1]].z);

这两个似乎不太等价,可能会对你的printf的结果的查看造成干扰,前面的%d输出的是i的索引,后面的%d对应的则是sortpart[索引]变换后的。这里会不会正好对你造成误会?


(2)此外,将你的cudaMalloc改成cudaMallocManaged(), 然后取消掉最后一个cudaMemcpy回传,看看结果是否正确。Unified Memory的GPU写入结果,CPU可以直接看到(有一定的限制条件,这里先不管),不需要回传,这样可以排除或者确认是否是cudaMemcpy()本身的问题。



Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 回复 #8 于: 十一月 03, 2022, 11:52:14 am »
还有一个问题是,我在这个核函数刚开始写的赋值代码
 if(sortpart[p1] >= 0 && sortpart[p1] <CONSTANTS.nBound2Start)
{
Repforce[sortpart[p1]].x=0.0;
Repforce[sortpart[p1]].y=0.0;
Repforce[sortpart[p1]].z=0.0;
}
对最后的计算结果影响很大,加上他就是对的,不加就是错的,但是我明明在CPU端已经给他赋值为0了再传回到GPU端上,按说在这里是否赋值为0没有影响才对的吧

这是另外一个问题了,我们建议每次只考虑一个问题,并进行讨论。否则容易歪楼。

如果你能遇到这种情况,排除了没有正确初始化/传输的可能后,则较大的可能则是你有多个索引值p(或者p1), 能进行sortpart[p]这种查表映射后,映射到相同的位置,从而引发GPU上的竞态。我建议你先解决正事再讨论这个。这里只是简单说一说。你也可以自行排查一下(我看不到该表,不能确定或者排除此可能)。

Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 回复 #9 于: 十一月 03, 2022, 11:58:39 am »
非常感谢您的回复!那我们先解决上面的问题!

前面的i,也是3203,与后面的p1=3202是一样的值,两部分代码中的sortpart[p1]以及sortpart均为0,我觉得这里应该没有问题。

您说的第二个问题是这个意思吗:在分配d_RepForce[](用于GPU端的排斥力变量)的空间时换成cudaMallocManaged(),以及不用写cudaMemcpy回传代码,直接Printf就可以看到在GPU计算后改变的值对吧?

Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 回复 #10 于: 十一月 03, 2022, 12:02:20 pm »
非常感谢您的回复!那我们先解决上面的问题!

前面的i,也是3203,与后面的p1=3202是一样的值,两部分代码中的sortpart[p1]以及sortpart均为0,我觉得这里应该没有问题。

您说的第二个问题是这个意思吗:在分配d_RepForce[](用于GPU端的排斥力变量)的空间时换成cudaMallocManaged(),以及不用写cudaMemcpy回传代码,直接Printf就可以看到在GPU计算后改变的值对吧?

如果第一点你能确保就好。

第二点是你这样的理解的(如果你用Windows,请确保CPU在使用的时候,GPU已经所有工作都完成了,没有在进行中;Linux可以暂时无视这点)。注意将不再存在d_RepForce和RepForce, 只有1个缓冲区,CPU和GPU都用。

Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 回复 #11 于: 十一月 03, 2022, 08:11:43 pm »
您好,再次打扰了,我改成了统一内存,但是发现仍然和不使用统一内存得到的结果是相同的。CPU端输出与在核函数中输出不同,我是在Linux系统下运行的程序,那就说明应该不是之前cudaMemcpy函数的问题对吧

Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 回复 #12 于: 十一月 04, 2022, 09:41:59 am »
您好,使用统一内存的具体情况是这样的:我本来是有CPU端变量RepForce,以及GPU端变量d_RepForce,换成统一内存之后,我是将CPU端定义的变量删掉了,直接在GPU端这个变量上改的,因此现在得程序中d_RepForce是一个统一内存变量。
如附件一的图上,如果我在这个计算排斥力的函数 CalFluidRepForce() 中进行1161行的这种输出,那么会显示内存的段错误,但是我这里参数列表就有这个RepForceg呀,为什么会显示段错误呢?如果我按1162行这样输出,也就是直接使用统一内存变量名,这样会显示d_RepForce没有定义,但是我是包含着定义这个统一内存变量的头文件的,可能因为给统一内存分配空间是在另一个.cu文件导致的吗。

如附件二的图,如果我在调用计算排斥力的函数 CalFluidRepForce() 的函数CalRepForce()中按着722行形式输出,也就是也直接使用统一内存变量名,这样可以输出结果,但与核函数中输出结果不同,CalRepForce()所在.cu文件包含定义这个统一内存变量的头文件,并且统一内存的空间也是在这个.cu文件中分配的。

Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 回复 #13 于: 十一月 06, 2022, 11:08:00 am »
您好,非常感谢您的回复,问题找到了,如我给您发的第二张图,最后一个参数传递写错了,本应该是d_RepForce,但是我写成了d_SPHForce