Есть ли способ блокировать некоторые блоки, пока не будут выполнены некоторые условия?

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 (что является пределом одинарной точности),Точность целочисленных значений за пределами этой точки не определена.Спасибо талантливым за указание на это.

./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 блоков, прежде чем выполнить безумно бесполезный подсчетпоследний блок).

Это только подтверждение концепции, не принимайте код всерьез