CUDA 上的块间屏障

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

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

#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)
            ;
    }

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

不幸的是，您想要实现的目标（块间通信/同步）在 CUDA 中并不完全可能。 CUDA 编程指南指出“线程块需要独立执行：必须能够以任何顺序、并行或串行执行它们”。此限制的原因是为了允许线程块调度程序具有灵活性，并允许代码随内核数量进行不可知的扩展。唯一支持的块间同步方法是启动另一个内核：内核启动（在同一流内）是隐式同步点。

您的代码违反了块独立性规则，因为它隐式假设内核的线程块同时执行（参见并行）。但不能保证他们会这样做。为了了解为什么这对您的代码很重要，让我们考虑一个假设的只有一个核心的 GPU。我们还假设您只想启动两个线程块。在这种情况下，你的 spinloop 内核实际上会死锁。如果线程块 0 首先在核心上调度，那么当它到达屏障时，它将永远循环，因为线程块 1 永远没有机会更新计数器。因为线程块零永远不会被换出（线程块执行到完成），所以它在旋转时会使核心之一的线程块挨饿。

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

对于您的应用程序，尝试找到一个不依赖于块间同步的解决方案，因为（除非对 CUDA 编程模型进行含义更改）这是不可能的。