OpenACC 当前子句更新数据

Question

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

call to cuStreamSynchronize returned error 700: Illegal address during kernel execution
call to cuMemFreeHost returned error 700: Illegal address during kernel execution
srun: error: jrc0017: task 0: Exited with exit code 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;

我的代码能够运行没有问题。但是如果我 un-comment 它们，我会遇到内存问题。似乎 sim.part.vx 正在尝试更新数据，但编译器不知道是什么导致了内存问题。

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

Answer 1

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

OpenACC 当前子句更新数据

OpenACC present clause update data

c++

gpu

c++11

openacc