CUDA 程序挂起：过滤器锁定

Question

我对@9​​87654321@ 给出的代码应用了过滤器锁定。锁定算法参考Filter Lock Algorithm。

#include <stdio.h>

__device__
void releaseLock(int i, int* level) {
  level[i] = -1;
}

__device__
bool sameOrHigher(int i, int j, int *level, int n) {
  for(int k = 0; k < n; k++) {
    if (k!= i && level[k] >= j) return true;
  }
  return false;
}

__device__
void acquireLockWithNeighbours(int i, int *level, int *victim, int n)
{
  for (int j = 0; j < n; j++) {
    level [i] = j;
    victim [j] = i;
    // wait while conflicts exist
    while (sameOrHigher(i, j, level, n) && victim[j] == i);
  }
}

__global__
void saxpy(int n, float a, float *x, float *y, int *level, int *victim)
{
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  if(i >= n) return;
  acquireLockWithNeighbours(i, level, victim, n);
  if (i < n) y[i] = a*x[i] + y[i];
  releaseLock(i, level);
}

int main(void)
{
  int N = 1024;
  float *x, *y, *d_x, *d_y;
  x = (float*)malloc(N*sizeof(float));
  y = (float*)malloc(N*sizeof(float));

  cudaMalloc(&d_x, N*sizeof(float));
  cudaMalloc(&d_y, N*sizeof(float));

  int *l, *v, *d_l, *d_v;
  l = (int*)malloc(N*sizeof(int));
  v = (int*)malloc(N*sizeof(int));
  cudaMalloc(&d_l, N*sizeof(int));
  cudaMalloc(&d_v, N*sizeof(int));

  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
    l[i] = -1;
    v[i] = -1;
  }

  cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);
  cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice);
  cudaMemcpy(d_l, l, N*sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_v, v, N*sizeof(int), cudaMemcpyHostToDevice);

  saxpy<<<(N+255)/256, 256>>>(N, 2.0f, d_x, d_y, d_l, d_v);

  cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost);

  float maxError = 0.0f;
  cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost);

  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = max(maxError, abs(y[i]-4.0f));
  printf("Max error: %f\n", maxError);
}

问题是：1 给出的代码可以正常执行。但是上面的代码卡住了。

任何建议将不胜感激：

1. 代码中存在的任何错误，只要删除就会使其运行。

2. 还有其他方法来编码锁吗？我想模拟一个图，其中每个节点都有执行能力。我正在寻找一个锁，其中每个节点都与邻居互斥执行。每个节点都知道它的邻居；对于我可以使用的锁类型的任何建议将不胜感激。

Answer 1

代码中存在的任何错误，只要删除就会使其运行。

1. 我对“过滤器锁定”算法不太熟悉，但我只是在处理您提供的代码。

2. 您发布的代码无法编译，您有两个这样的实例：

   float maxError = 0.0f;
   

   我们可以删除第一个。

3. 最近的问题似乎是您正在使用全局内存在线程之间进行通信，但编译器对您不利。编译器假定允许将全局内存值“优化”到寄存器中，除非存在某些其他条件（线程/内存栅栏/屏障、原子、volatile 装饰）。正因为如此，即使在只有两个线程的情况下，无论是在同一个线程块中还是在不同的线程块中，您都可以见证死锁。每个线程都可以写入全局内存，但不能保证线程会看到那里写入的其他值。有些可能会被“看到”，有些则不会。这显然不能可靠地使用该算法。 “修复”此问题的一种可能方法是使用 volatile keyword 装饰用于线程间通信的全局变量，这会阻止我描述的“优化”效果。

4. 在 cmets 中已经提到的另一个问题是这种结构：

   while (sameOrHigher(i, j, level, n) && victim[j] == i);
   

   在 Volta 之前的 GPU 架构上通常不能很好地工作，其中线程间争用发生在同一 warp 中。由于扭曲以锁步方式执行，因此执行可能会“挂起”在等待锁但从未获得锁的线程上，从而阻止将“释放”锁的线程。 换句话说，while 循环的执行停留在等待的线程上，绝不允许可以前进的线程实际前进。在cuda标签上有各种各样的问题，here是一个例子，here是另一个带有分析的例子。 Volta 和其他 GPU 架构引入了一种新的线程调度模型，可以帮助缓解这种情况。在 Volta 之前的处理器上，我的一般建议是不要尝试使用涉及 Warp 内争用的锁。 （请参阅下面的链接了解我的建议。）

5. 随着线程数量的增加，这个锁系统碰巧会以指数方式变得更糟（就获取锁的工作而言）。稍后我们将看到一个这样的例子。我不是 Java 专家，但我怀疑这种过滤器锁定机制可能适用于 4 个线程或 8 个线程，但 Volta 处理器在任何给定时间都可以有超过 100,000 个线程在运行。在考虑这种机制时，这可能是一个真正的问题。

如果我们结合上面的想法，下面的代码提供了一个有用的测试用例进行实验。我们消除了额外的maxError 定义，对volatile 进行了一些适当的修饰，还添加了一个工具，以便我们可以从命令行测试不同大小的N。包括在 V100 处理器上运行的一些测试：

$ cat t1946.cu
#include <stdio.h>
#include <stdlib.h>

__device__
void releaseLock(int i, volatile int* level) {
  level[i] = -1;
  __threadfence();
}

__device__
bool sameOrHigher(int i, int j, volatile int *level, int n) {
  for(int k = 0; k < n; k++) {
    if (k!= i && level[k] >= j) return true;
  }
  return false;
}

__device__
void acquireLockWithNeighbours(int i, volatile int *level, volatile int *victim, int n)
{
  for (int j = 0; j < n; j++) {
    level [i] = j;
    victim [j] = i;
    __threadfence();
    // wait while conflicts exist
    while (sameOrHigher(i, j, level, n) && victim[j] == i);
  }
}

__global__
void saxpy(int n, float a, float *x, float *y, volatile int *level, volatile int *victim)
{
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  if(i >= n) return;
  acquireLockWithNeighbours(i, level, victim, n);
  if (i < n) y[i] = a*x[i] + y[i];
  releaseLock(i, level);
}

int main(int argc, char *argv[])
{
  int N = 2;
  if (argc > 1) N = atoi(argv[1]);
  float *x, *y, *d_x, *d_y;
  x = (float*)malloc(N*sizeof(float));
  y = (float*)malloc(N*sizeof(float));

  cudaMalloc(&d_x, N*sizeof(float));
  cudaMalloc(&d_y, N*sizeof(float));

  int *l, *v, *d_l, *d_v;
  l = (int*)malloc(N*sizeof(int));
  v = (int*)malloc(N*sizeof(int));
  cudaMalloc(&d_l, N*sizeof(int));
  cudaMalloc(&d_v, N*sizeof(int));

  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
    l[i] = -1;
    v[i] = -1;
  }

  cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);
  cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice);
  cudaMemcpy(d_l, l, N*sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_v, v, N*sizeof(int), cudaMemcpyHostToDevice);

  saxpy<<<(N+255)/256, 256>>>(N, 2.0f, d_x, d_y, d_l, d_v);

  cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost);

  cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost);

  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = max(maxError, abs(y[i]-4.0f));
  printf("Max error: %f\n", maxError);
}
$ nvcc -arch=sm_70 -o t1946 t1946.cu
$ time ./t1946 128
Max error: 0.000000

real    0m1.023s
user    0m0.467s
sys     0m0.552s
$ time ./t1946 256
Max error: 0.000000

real    0m4.694s
user    0m2.984s
sys     0m1.706s
$ time ./t1946 512
Max error: 0.000000

real    0m27.764s
user    0m18.215s
sys     0m9.543s
$ time ./t1946 1024
Max error: 0.000000

real    3m9.205s
user    2m6.902s
sys     1m2.288s
$

我们可以看到，虽然事情看起来很有效，但是当我们将 N 的大小翻倍时，我们将每一步的执行时间增加了大约 6 倍。但是我们知道，当我们到达 1024 的N 时，我们已经涵盖了经线内和经线间的情况，以及块间争用的情况。即使这个代码仍然只是挂在一个 pre- volta 处理器，即使N 为 2。我没有耐心等待代码在 2048 年的N 上花费多长时间，但我猜这个 V100 大约需要 20 分钟。

还有什么其他方式来编码锁吗？

是的，还有其他方法。如果您在cuda 标签上进行一些搜索，尤其是像lock 或critical section 或mutex 这样的关键字，您会找到示例。我熟悉的大多数使用原子来解决争用。我想我们可以观察到，随着线程数变大，这种当前形式的“过滤器锁定”算法基本上变得毫无用处。然而，这些其他原子方法（特别是如果我们将争用限制为线程块级别 I suggest）可以解决大规模问题，而无需我们在此处看到的那种开销。

我可以对这种“过滤器锁定”算法提出的另一个批评是，它似乎希望知道有多少线程在运行（或者至少上限是多少）。在一般情况下，原子锁机制不需要这些知识，并且可以设计为在没有这些知识的情况下正常工作（再次参见上面我建议的示例）。

（我在上面的代码中包含的__threadfence()instructions可能不是必需的，但它们可能会使代码整体执行得更快一些。）

关于这个：

我正在寻找一个锁，其中每个节点都与邻居互斥执行。每个节点都知道它的邻居；对于我可以使用的锁类型的任何建议将不胜感激。

我会说这种想法可能与如何从 GPU 获得最佳结果不一致。如果您的算法的性质是图形节点偶尔必须相互同步，但在很大程度上可以独立执行（即以任何顺序），那么您可能会没事。但是，如果您的算法受活动周期支配，仅允许单个节点执行任何操作，这通常可能与 GPU 的有效使用不一致。如果单个节点完成的处理足以“饱和”GPU，那么它可能没问题。否则你可能会得到令人失望的表现。这种观点或多或少与您使用哪种特定类型的锁来安排互斥无关。