c - CUDA 上的 block 间屏障

Question

我想在 CUDA 上实现一个 block 间屏障，但是遇到了一个严重的问题。

我不明白为什么它不起作用。

#include <iostream>
#include <cstdlib>
#include <ctime>

#define SIZE 10000000
#define BLOCKS 100 

using namespace std;

struct Barrier {
    int *count;

    __device__ void wait() {
        atomicSub(count, 1);
        while(*count)
            ;
    }

    Barrier() {
        int blocks = BLOCKS;
        cudaMalloc((void**) &count, sizeof(int));
        cudaMemcpy(count, &blocks, sizeof(int), cudaMemcpyHostToDevice);
    }

    ~Barrier() {
        cudaFree(count);
    }
};


__global__ void sum(int* vec, int* cache, int *sum, Barrier barrier)
{
    int tid = blockIdx.x;

    int temp = 0;
    while(tid < SIZE) {
        temp += vec[tid];
        tid += gridDim.x;
    }

    cache[blockIdx.x] = temp;

    barrier.wait();

    if(blockIdx.x == 0) {
        for(int i = 0 ; i < BLOCKS; ++i)
            *sum += cache[i];
    }
}

int main()
{
    int* vec_host = (int *) malloc(SIZE * sizeof(int));    
    for(int i = 0; i < SIZE; ++i)
        vec_host[i] = 1;

    int *vec_dev;
    int *sum_dev;
    int *cache;
    int sum_gpu = 0;

    cudaMalloc((void**) &vec_dev, SIZE * sizeof(int));
    cudaMemcpy(vec_dev, vec_host, SIZE * sizeof(int), cudaMemcpyHostToDevice);
    cudaMalloc((void**) &sum_dev, sizeof(int));
    cudaMemcpy(sum_dev, &sum_gpu, sizeof(int), cudaMemcpyHostToDevice);
    cudaMalloc((void**) &cache, BLOCKS * sizeof(int));
    cudaMemset(cache, 0, BLOCKS * sizeof(int));

    Barrier barrier;
    sum<<<BLOCKS, 1>>>(vec_dev, cache, sum_dev, barrier);

    cudaMemcpy(&sum_gpu, sum_dev, sizeof(int), cudaMemcpyDeviceToHost);

    cudaFree(vec_dev);
    cudaFree(sum_dev);
    cudaFree(cache);
    free(vec_host);
    return 0;
}

事实上，即使我将 wait() 重写为如下

    __device__ void wait() {
        while(*count != 234124)
            ;
    }

程序正常退出。但我希望在这种情况下会出现无限循环。

Answer 1

不幸的是，您想要实现的( block 间通信/同步)在 CUDA 中并非严格可行。 CUDA 编程指南指出“线程 block 需要独立执行:必须能够以任何顺序(并行或串行)执行它们。”此限制的原因是允许线程 block 调度程序具有灵 active ，并允许代码随内核数量不可知地扩展。唯一受支持的 block 间同步方法是启动另一个内核:内核启动(在同一流内)是隐式同步点。

您的代码违反了 block 独立性规则，因为它隐含地假定您的内核线程 block 并发执行(参见并行)。但不能保证他们会这样做。要了解为什么这对您的代码很重要，让我们考虑一个只有一个内核的假想 GPU。我们还假设您只想启动两个线程 block 。在这种情况下，您的自旋循环内核实际上会死锁。如果第一个线程 block 0 被调度到核心上，当它到达屏障时它将永远循环，因为线程 block 1 永远没有机会更新计数器。因为线程 block 零永远不会被换出(线程 block 执行到它们完成)，它在旋转时使核心之一的线程 block 处于饥饿状态。

有些人已经尝试过像您这样的方案并取得了成功，因为调度器碰巧以假设成功的方式偶然地调度了 block 。例如，曾经有一段时间启动与 GPU 拥有的 SM 一样多的线程 block 意味着这些 block 是真正并发执行的。但是当对驱动程序或 CUDA 运行时或 GPU 的更改使该假设无效并破坏了他们的代码时，他们感到很失望。

对于您的应用程序，请尝试找到不依赖于 block 间同步的解决方案，因为(除非对 CUDA 编程模型进行重大更改)这是不可能的。