OpenACC 存在子句更新数据

OpenACC present clause update data

本文关键字：数据 更新 子句 存在 OpenACC      更新时间：2023-10-16

我正在尝试为许多身体模拟进行openACC优化。目前，我面临一个导致以下内存问题的问题

调用 cuStreamSynchronize返回错误 700：内核执行
期间地址非法 调用 cuMemFreeHost返回错误 700：内核执行
期间地址非法 srun：错误：jrc0017：任务 0：退出，退出代码为 1

我正在使用pgc++编译器，我的编译器标志-acc -Minfo=accel -ta=tesla -fast -std=c++11，我不想使用-ta=tesla:managed，因为我想自己组织内存。

#pragma acc kernels present(sim.part.rx, sim.part.ry, sim.part.rz, sim.part.vx, sim.part.vy, sim.part.vz) 
{
  for(int idx = 0; idx < sim.num; ++idx) {     // Loop over target particle
    float
      prx = sim.part.rx[idx],                   // my position
      pry = sim.part.ry[idx],
      prz = sim.part.rz[idx];
    float Fx = 0.f, Fy = 0.f, Fz = 0.f;          // Force
    #pragma acc loop 
    for(int jdx = 0; jdx < sim.num; ++jdx) {   // Loop over interaction partners
      if(idx != jdx) {                          // No self-force
        const float dx = prx - sim.part.rx[jdx]; // Distance to partner
        const float dy = pry - sim.part.ry[jdx];
        const float dz = prz - sim.part.rz[jdx];
        const float h  = 1.f/sqrt(dx*dx + dy*dy + dz*dz + eps);
        const float h3 = h*h*h;
        Fx += dx*h3;                            // Sum up force
        Fy += dy*h3;
        Fz += dz*h3;
      }
    }
    sim.part.vx[idx] += sim.mass*dt*Fx;         // update velocity
    sim.part.vy[idx] += sim.mass*dt*Fy;
    sim.part.vz[idx] += sim.mass*dt*Fz;
  }
}

如果我删除下面的代码 

sim.part.vx[idx] += sim.mass*dt*Fx;         // update velocity
sim.part.vy[idx] += sim.mass*dt*Fy;
sim.part.vz[idx] += sim.mass*dt*Fz;

我的代码能够毫无问题地运行。但是如果我取消评论它们，我会遇到记忆问题。似乎sim.part.vx尝试更新数据，但编译器不知道哪个导致内存问题。

有谁知道如何解决这个问题？

我怀疑问题是sim和sim.part不在设备上（或者编译器没有意识到它们在设备上。作为一种解决方法，您可以尝试直接引入指向这些数组的指针吗？

float *rx = sim.part.rx, *ry = sim.part.ry, *rz = sim.part.rz, 
      *vx = sim.part.vx, *vy = sim.part.vy, *vz = sim.part.vz;
#pragma acc kernels present(rx, ry, rz, vx, vy, vz) 
{
  for(int idx = 0; idx < sim.num; ++idx) {     // Loop over target particle
    float
      prx = rx[idx],                   // my position
      pry = ry[idx],
      prz = rz[idx];
    float Fx = 0.f, Fy = 0.f, Fz = 0.f;          // Force
    #pragma acc loop 
    for(int jdx = 0; jdx < sim.num; ++jdx) {   // Loop over interaction partners
      if(idx != jdx) {                          // No self-force
        const float dx = prx - rx[jdx]; // Distance to partner
        const float dy = pry - ry[jdx];
        const float dz = prz - rz[jdx];
        const float h  = 1.f/sqrt(dx*dx + dy*dy + dz*dz + eps);
        const float h3 = h*h*h;
        Fx += dx*h3;                            // Sum up force
        Fy += dy*h3;
        Fz += dz*h3;
      }
    }
    vx[idx] += sim.mass*dt*Fx;         // update velocity
    vy[idx] += sim.mass*dt*Fy;
    vz[idx] += sim.mass*dt*Fz;
  }
}
sim

和 sim.part 是如何分配的？可以在构造函数和析构函数中使用非结构化数据指令，以确保 sim 和 sim.part 也在设备上。如果您已经这样做了，那么另一种可能的解决方案是将present(sim, sim.part)添加到现有的 present 子句中，以便编译器知道您也已经处理了这些数据结构。

相关文章：