是否有方法能够阻止某些块直到满足某些条件？

Question

我想阻止一些块，直到一个变量被设置为特定值。所以我编写这段代码来测试一个简单的 do-while 循环是否有效。

__device__ int tag = 0;
__global__ void kernel() {
    if ( threadIdx.x == 0 ) {
        volatile int v;
        do {
            v = tag;
        }
        while ( v == 0 );
    }
    __syncthreads();
    return ;
}

然而，它不起作用（没有发生死循环，很奇怪）。

我想问是否有任何其他方法能够阻止某些块，直到满足某些条件，或者对代码进行一些更改是否有效。

Answer 1

目前在 CUDA 中没有可靠的方法来执行块间同步。

有一些黑客方法可以在总线程数适中的块之间实现某种方式的锁定或阻塞，但它们利用了执行模型中未定义的行为，这些行为不能保证在所有硬件上以相同的方式运行或在将来继续工作。确保块之间同步或阻塞的唯一可靠方法是我们单独启动内核。如果您的算法在没有块间同步的情况下无法工作，那么您要么需要新的算法，要么您的应用程序非常不适合 GPU 架构。

Answer 2

这是我尝试看看是否有效的一种黑客方法。

#include <stdio.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <cuda_runtime_api.h>

__global__ static
void kernel(int *count, float *data)
{
    count += threadIdx.x;
    data += gridDim.x * threadIdx.x;
    int i = blockIdx.x;
    if (i < gridDim.x - 1) {
        data[i] = i + 1;
        atomicAdd(count, 1);
        return;
    }

    while (atomicMin(count, i) != i);

    float tmp = i + 1;
    for (int j = 0; j < i; j++) tmp += data[j];

    data[i] = tmp;
}

int main(int argc, char **args)
{
        int num = 100;
    if (argc >= 2) num = atoi(args[1]);

    int bytes = num * sizeof(float) * 32;
    float *d_data; cudaMalloc((void **)&d_data, bytes);
    float *h_data = (float *)malloc(bytes);
    for (int i = 0; i < 32 * num; i++) h_data[i] = -1; // Being safe                                                                                                                           

    int h_count[32] = {1};
    int *d_count; cudaMalloc((void **)&d_count, 32 * sizeof(int));
    cudaMemcpy(d_count, &h_count, 32 * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_data, h_data, bytes, cudaMemcpyHostToDevice);
    kernel<<<num, 32>>>(d_count, d_data);
    cudaMemcpy(&h_count, d_count, 32 * sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(h_data, d_data, bytes, cudaMemcpyDeviceToHost);

    for (int i = 0; i < 32; i++) {
        printf("sum of first %d from thread %d is %d \n", num, i, (int)h_data[num -1]);
        h_data += num;
    }

    cudaFree(d_count);
    cudaFree(d_data);
    free(h_data - num * 32);
}

我不能保证这永远有效。但我的卡（320M）上的断点似乎是 num = 5796。也许每张卡都有某种不同的硬件限制？

编辑

这个问题的答案是 n * (n + 1) / 2 > 2^24 对于 n > 5795（这是单精度限制）。超出此点的整数值的准确性是未定义的。感谢 talonmies 指出这一点。

./a.out 5795
sum of first 5795 from thread 0 is 16793910 
sum of first 5795 from thread 1 is 16793910 
sum of first 5795 from thread 2 is 16793910 
sum of first 5795 from thread 3 is 16793910 
sum of first 5795 from thread 4 is 16793910 
sum of first 5795 from thread 5 is 16793910 
sum of first 5795 from thread 6 is 16793910 
sum of first 5795 from thread 7 is 16793910 
sum of first 5795 from thread 8 is 16793910 
sum of first 5795 from thread 9 is 16793910 
sum of first 5795 from thread 10 is 16793910 
sum of first 5795 from thread 11 is 16793910 
sum of first 5795 from thread 12 is 16793910 
sum of first 5795 from thread 13 is 16793910 
sum of first 5795 from thread 14 is 16793910 
sum of first 5795 from thread 15 is 16793910 
sum of first 5795 from thread 16 is 16793910 
sum of first 5795 from thread 17 is 16793910 
sum of first 5795 from thread 18 is 16793910 
sum of first 5795 from thread 19 is 16793910 
sum of first 5795 from thread 20 is 16793910 
sum of first 5795 from thread 21 is 16793910 
sum of first 5795 from thread 22 is 16793910 
sum of first 5795 from thread 23 is 16793910 
sum of first 5795 from thread 24 is 16793910 
sum of first 5795 from thread 25 is 16793910 
sum of first 5795 from thread 26 is 16793910 
sum of first 5795 from thread 27 is 16793910 
sum of first 5795 from thread 28 is 16793910 
sum of first 5795 from thread 29 is 16793910 
sum of first 5795 from thread 30 is 16793910 
sum of first 5795 from thread 31 is 16793910 

--

我编辑了以前的代码，该代码仅使用一个块。这更能代表现实世界的线程/块（内存访问很奇怪并且会非常慢，但这样做是为了快速移植我的旧测试代码以使用多个线程）。

看起来在某些情况下你可以跨块同步，但主要取决于你事先了解某些事情（对于这种特殊情况，我只同步 n - 1 个块，然后对最后一个块执行疯狂无用的计数）。

这只是一个概念证明，请勿认真对待代码