用于嵌套 for 循环的 CUDA 网格步幅循环

Question

我正在使用CUDA网格跨步循环，结果看起来不错，但我不是100％理解为什么只需要对外循环进行跨步操作，而内循环没有任何变化。

__global__
void bodyForce(Body *p, float dt, int n) {
  
  int index = threadIdx.x + blockIdx.x * blockDim.x;
  int stride = blockDim.x * gridDim.x;


  for (int i = index; i < n; i += stride)
  {
    float Fx = 0.0f; float Fy = 0.0f; float Fz = 0.0f;

    for (int j = 0; j < n; j ++)
    {
      float dx = p[j].x - p[i].x;
      float dy = p[j].y - p[i].y;
      float dz = p[j].z - p[i].z;
      float distSqr = dx*dx + dy*dy + dz*dz + SOFTENING;
      float invDist = rsqrtf(distSqr);
      float invDist3 = invDist * invDist * invDist;

      Fx += dx * invDist3; Fy += dy * invDist3; Fz += dz * invDist3;
    }

    p[i].vx += dt*Fx; p[i].vy += dt*Fy; p[i].vz += dt*Fz;
  }
}

Answer 1

此代码似乎直接来自 NVIDIA DLI 介绍性 CUDA C++ course 的评估练习。

如果您正在学习本课程，那么网格步长循环就是您在将原始纯 CPU 代码转换为 GPU 加速代码的过程中编写的东西。所以代码是你选择写的任何东西。

外部 for 循环上的 grid-stride 循环有意义的原因是因为外部 for 循环迭代完成的工作是独立的。从一个迭代到下一个迭代的计算顺序无关紧要，因此这可以轻松/简单地跨 CUDA 线程并行化，the grid-stride loop 是分配外部 for 循环工作的一种方法（在原始 CPU-仅代码）跨 CUDA 线程。

内部 for 循环表示在迭代之间并非平凡独立的工作，因为一次迭代正在添加到前一次迭代的结果：

  Fx += dx * invDist3; Fy += dy * invDist3; Fz += dz * invDist3;

因此，跨 CUDA 线程分配这项工作将不会采用与用于并行化外部 for 循环相同的琐碎/机械方法。

当然可以，但这需要教授跨线程并行归约的概念，该课程当时尚未涉及。因此，明智的做法是不理会内部 for 循环，由单个线程处理，每个 Body。

这种认识仍然可以实现在该课程的评估练习中设定的绩效目标。