CUDA独立线程调度饥饿

Question

ThisNVIDIA 官方博客中的“无饥饿算法”部分声称 Volta 中的 ITS 添加了对无饥饿算法的支持，并提供了一些细节。 This 官方文档指出，Turing 具有与 Volta 相同的 ITS。为什么下面的代码会遇到致命饥饿？

编辑：我修改了代码以仅测试扭曲内饥饿。在 T4、2080 Ti 和 RTX 3070 上进行了尝试，CUDA 版本包括 11.5 和 12.1，使用适当的架构/代码参数。两种锁实现，

libcudacxx

和

legacy

，都不允许线程 1 获取锁，除了 RTX 3070 上的

legacy

，尽管事实上锁一次被释放了一整秒。

#include <cuda.h>
#include <cstdio>
#include <cuda/semaphore>
#include <cuda/atomic>

__device__ uint32_t something_very_slow(uint32_t x) {
  for (uint32_t i = 0; i / 1e7 < 1; ++i) {
    x *= 13;
    x += 1;
    x %= 123456789;
  }
  return x;
}

__device__ cuda::binary_semaphore<cuda::thread_scope_block> lock{1};
__device__ cuda::atomic<uint32_t, cuda::thread_scope_block> mask{0};
__device__ cuda::atomic<uint32_t, cuda::thread_scope_block> clobber{0};

__global__ void starvation_libcudacxx() {
  lock.acquire();
  printf("start thread %d\n", threadIdx.x);
  bool cont = false;
  do {
    printf("step thread %d\n", threadIdx.x);
    lock.release();
    clobber.fetch_add(something_very_slow(clobber.load()) + threadIdx.x);
    cont = mask.fetch_add(threadIdx.x) == 0;
    lock.acquire();
  } while (cont);
  printf("done: %d\n", clobber.load());
  lock.release();
}

__global__ void starvation_legacy() {
  __shared__ uint32_t lock, mask, clobber;
  if (threadIdx.x == 0) {
    lock = mask = clobber = 0;
  }
  __syncthreads();
  while (atomicCAS(&lock, 0, 1) == 1) {
  }
  printf("start thread %d\n", threadIdx.x);
  bool cont = false;
  do {
    printf("step thread %d\n", threadIdx.x);
    atomicExch(&lock, 0);
    atomicAdd(&clobber, something_very_slow(atomicAdd(&clobber, 0)) + threadIdx.x);
    cont = atomicAdd(&mask, threadIdx.x) == 0;
    while (atomicCAS(&lock, 0, 1) == 1) {
    }
  } while (cont);
  printf("done: %d\n", atomicAdd(&clobber, 0));
  atomicExch(&lock, 0);
}

int main() {
  starvation_libcudacxx<<<1, 2>>>();
  starvation_legacy<<<1, 2>>>();
  cudaDeviceSynchronize();
}

Answer 1

首先感谢你的问题，因为我最近也在阅读这个博客，你的代码帮助我更好地理解这个博客。

---

简短的回答：使用与博客相同的结构来使代码正常工作。而且在kernel中使用看似无限的循环并不是一个好习惯，这似乎与GPU内核设计相反，GPU内核设计只是一个small线程来运行。

然后线程网格之间的饥饿就会消失（我使用GTX 1650，这也是图灵架构）：

$ alias nvcc
nvcc='nvcc -arch=sm_75'
$ nvcc starvation.cu
$ ./a.out 
start thread 0
...
start thread 1
...
start thread 0
...

因此，至少避免了“无饥饿”的更坏情况，即多个网格中“只有一个线程”获得锁。但显然这似乎不是独立线程调度在更细粒度的线程级别应该做的事情。 附注由于我没有其他 GPU，因此我不知道

Volta 之前的 GPU

（例如 Pascal）上的行为是什么。

这里我展示了如何更改代码以使其部分工作。但为什么在硬件、驱动程序或其他（可能超出用户知识范围）中实现的

---

独立线程调度

无法使原始代码工作，这超出了我的能力范围。因此，要知道为什么它不使用更细粒度的线程级别来使您的原始代码工作可能并不容易。 希望像

Nvidia员工或CUDA专家

这样的人可以帮助改进这个答案。预先感谢。

详细解答：

---

将您的代码与原始博客代码进行比较后（我发布在这里是为了更好地查看）：

__device__ void insert_after(Node *a, Node *b)
{
    Node *c;
    lock(a); lock(a->next);
    c = a->next;

    a->next = b;
    b->prev = a;

    b->next = c;
    c->prev = b;

    unlock(c); unlock(a);
}

博客内核将

unlock

作为最终指令。因此，将您的代码更改为与博客相同的结构：

$ diff starvation_orig.cu starvation.cu     
--- starvation_orig.cu  2023-08-02 15:38:42.592480115 +0800
+++ starvation.cu       2023-08-02 15:39:14.219905308 +0800
@@ -1,3 +1,4 @@
+// https://stackoverflow.com/q/76497234/21294350
 #include <cuda.h>
 #include <cstdio>
 #include <cuda/semaphore>
@@ -20,15 +21,12 @@
   lock.acquire();
   printf("start thread %d\n", threadIdx.x);
   bool cont = false;
-  do {
-    printf("step thread %d\n", threadIdx.x);
-    lock.release();
-    clobber.fetch_add(something_very_slow(clobber.load()) + threadIdx.x);
-    cont = mask.fetch_add(threadIdx.x) == 0;
-    lock.acquire();
-  } while (cont);
-  printf("done: %d\n", clobber.load());
+  printf("step thread %d\n", threadIdx.x);
   lock.release();
+  printf("release lock\n");
+  clobber.fetch_add(something_very_slow(clobber.load()) + threadIdx.x,cuda::memory_order_seq_cst);
+  cont = mask.fetch_add(threadIdx.x) == 0;
+  printf("done: %d\n", clobber.load());
 }
 
 __global__ void starvation_legacy() {
@@ -39,10 +37,10 @@
   __syncthreads();
   while (atomicCAS(&lock, 0, 1) == 1) {
   }
-  printf("start thread %d\n", threadIdx.x);
+  printf("starvation_legacy start thread %d\n", threadIdx.x);
   bool cont = false;
   do {
-    printf("step thread %d\n", threadIdx.x);
+    printf("starvation_legacy step thread %d\n", threadIdx.x);
     atomicExch(&lock, 0);
     atomicAdd(&clobber, something_very_slow(atomicAdd(&clobber, 0)) + threadIdx.x);
     cont = atomicAdd(&mask, threadIdx.x) == 0;
@@ -54,7 +52,7 @@
 }
 
 int main() {
-  starvation_libcudacxx<<<1, 2>>>();
-  starvation_legacy<<<1, 2>>>();
-  cudaDeviceSynchronize();
+  while (1) {
+    starvation_libcudacxx<<<1, 2>>>();
+  }
 }