1

我正在尝试对许多身体模拟进行 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正在尝试更新数据，但编译器不知道导致内存问题的原因。

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

4

1 回答 1

2

我怀疑问题出在设备上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 子句中，以便编译器知道您也已经处理了这些数据结构。

于 2016-01-27T13:41:52.793 回答