CUDA上的块间障碍

Question

我想在CUDA上实现Inter-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

不幸的是,在CUDA中,您想要实现的目标(块间通信/同步)并不是很严格.CUDA编程指南指出"线程块需要独立执行:必须能够以任何顺序,并行或串行执行它们." 这种限制的原因是允许线程块调度程序具有灵活性,并允许代码在核心数量上进行扩展.唯一支持的块间同步方法是启动另一个内核:内核启动(在同一个流中)是隐式同步点.

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

有些人已经尝试过像你这样的方案并且已经看到了成功,因为调度程序碰巧偶然地安排了块,以便假设成功.例如,有一段时间启动尽可能多的线程块,因为GPU具有SM意味着这些块是真正同时执行的.但是当对驱动程序或CUDA运行时或GPU的更改使该假设无效时,他们感到很失望,从而破坏了他们的代码.

对于您的应用程序,尝试找到一个不依赖于块间同步的解决方案,因为(除非对CUDA编程模型有意义的改变),这是不可能的.

Answer 2

可以阻止块同步.见本文.
本文没有详细介绍它是如何工作的,但它依赖于__syncthreads()的操作; 为当前块创建暂停屏障,...等待其他块到达同步点.

本文未提及的一个项目是,只有当块数足够小或者SM的数量足够大以完成现有任务时才能进行同步.即如果你有4个SM并且正在尝试同步5个块,那么内核将会死锁.

通过他们的方法,我已经能够在许多块中传播一个长串行任务,在单块方法上轻松节省30%的时间.即块同步对我有用.