CUDA 上的块间障碍

Posted 2023-04-15

【中文标题】CUDA 上的块间障碍

【英文标题】：Inter-block barrier on CUDA

【发布时间】：2011-12-03 22:39:27

【问题描述】：

我想在 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)
            ;
    

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

【问题讨论】：

我怀疑您的内核实际上是由于取消引用 Barrier::wait 中的错误指针而崩溃。使用cudaGetLastError 在内核期间检查错误。

【参考方案1】：

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

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

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

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

【讨论】：

最新 CUDA SDK 中的 threadFenceReduction 示例怎么样？他们不在那里进行屏障同步，但通过使用全局内存栅栏实现与主题启动器想要的类似结果（实际上，代码几乎相同，但他们只是检查当前块是否是自旋锁而不是自旋锁最后完成它的执行）。

也许可以用内存栅栏实现求和，但 OP 的问题是关于块间同步。在任何情况下，在不依赖原子的情况下，以两阶段方法更好地实现 OP 中示例规模的缩减。一个更好的主意是直接调用thrust::reduce。

【参考方案2】：

块到块同步是可能的。看到这个paper。 论文没有详细说明它是如何工作的，但它依赖于 __syncthreads(); 的操作。为当前块创建暂停屏障，...同时等待其他块到达同步点。

论文中没有提到的一点是，只有当块的数量足够少或 SM 的数量对于手头的任务足够大时，同步才是可能的。即如果你有 4 个 SM 并尝试同步 5 个块，.. 内核将死锁。

通过他们的方法，我已经能够在许多块之间分散一个长的串行任务，与单个块方法相比，轻松节省了 30% 的时间。即块同步对我有用。

【讨论】：

但是和前面的答案有矛盾吗？

【参考方案3】：

看起来像编译器优化问题。我不擅长阅读 PTX 代码，但看起来编译器完全省略了 while-loop（即使使用 -O0 编译）：

.loc    3   41  0
cvt.u64.u32     %rd7, %ctaid.x; // Save blockIdx.x to rd7
ld.param.u64    %rd8, [__cudaparm__Z3sumPiS_S_7Barrier_cache];
mov.s32     %r8, %ctaid.x; // Now calculate ouput address
mul.wide.u32    %rd9, %r8, 4;
add.u64     %rd10, %rd8, %rd9;
st.global.s32   [%rd10+0], %r5; // Store result to cache[blockIdx.x]
.loc    17  128 0
ld.param.u64    %rd11, [__cudaparm__Z3sumPiS_S_7Barrier_barrier+0]; // Get *count to rd11
mov.s32     %r9, -1; // put -1 to r9
atom.global.add.s32     %r10, [%rd11], %r9; // Do AtomicSub, storing the result to r10 (will be unused)
cvt.u32.u64     %r11, %rd7; // Put blockIdx.x saved in rd7 to r11
mov.u32     %r12, 0; // Put 0 to r12
setp.ne.u32     %p3, %r11, %r12; // if(blockIdx.x == 0)
@%p3 bra    $Lt_0_5122;
ld.param.u64    %rd12, [__cudaparm__Z3sumPiS_S_7Barrier_sum];
ld.global.s32   %r13, [%rd12+0];
mov.s64     %rd13, %rd8;
mov.s32     %r14, 0;

在 CPU 代码的情况下，通过使用 volatile 前缀声明变量来防止此类行为。但是即使我们将count 声明为int __device__ count（并适当更改代码），添加volatile 说明符也只会破坏编译（错误loke argument of type "volatile int *" is incompatible with parameter of type "void *"）

我建议查看 CUDA SDK 中的 threadFenceReduction 示例。他们所做的和你做的几乎一样，但是在运行时选择进行最终求和的块，而不是预定义的，并且消除了while-loop，因为全局变量上的自旋锁应该是 非常慢。

【讨论】：

threadFenceReduction 在一个关键点上有所不同：不是最后执行的块将继续执行并终止。这意味着将有最后一个要执行的块。在 OP 的方案中，他希望所有线程等到最后一个块到达屏障，但这可能导致死锁。

@Tom 我并不是说这样做完全相同相同，但是栅栏允许实现类似的结果（不是在指令流方面，而是在内容方面输出数组）

没说你这样做 ;-) 这就是我的意思，OP 正在尝试建立一个全局障碍，这是一个坏主意（请参阅 Jared 的回答），但看看他的代码，他可以达到预期的效果与 threadFenceReduction 示例相同。 @anyoneelse 阅读此内容：threadfence 不 与屏障相同！查看编程指南或在线搜索“内存围栏”以获取更多信息。