Межблочный барьер на CUDA

Question

Я хочу реализовать межблочный барьер в 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)
            ;
    }

Программа завершается нормально. Но я ожидаю получить бесконечный цикл в этом случае.

Answer 1

К сожалению, то, чего вы хотите достичь (межблочная связь / синхронизация), не является строго возможным в CUDA. В руководстве по программированию CUDA говорится, что «блоки потоков должны выполняться независимо: должна быть возможность выполнять их в любом порядке, параллельно или последовательно». Причина этого ограничения состоит в том, чтобы обеспечить гибкость в планировщике потокового блока и позволить коду независимо масштабироваться в зависимости от количества ядер. Единственный поддерживаемый метод межблочной синхронизации - запуск другого ядра: запуски ядра (в том же потоке) являются точками неявной синхронизации.

Ваш код нарушает правило независимости блоков, потому что оно неявно предполагает, что блоки потоков вашего ядра выполняются одновременно (см. Параллельно). Но нет гарантии, что они делают. Чтобы понять, почему это важно для вашего кода, давайте рассмотрим гипотетический графический процессор с одним ядром. Мы также предполагаем, что вы хотите запустить только два потоковых блока. В этом случае ваше ядро ​​spinloop фактически заходит в тупик. Если нулевой блок потока запланирован на ядре первым, он будет зацикливаться вечно, когда доберется до барьера, потому что первый блок потока никогда не имеет возможности обновить счетчик. Поскольку нулевой блок потока никогда не заменяется (блоки потока выполняются до их завершения), он останавливает блок потока одного из ядер, пока он вращается.

Некоторые люди пробовали схемы, подобные вашей, и добились успеха, потому что планировщик случайно спланировал блоки таким образом, что предположения сработали. Например, было время, когда запускалось столько потоковых блоков, сколько у графического процессора имеется SM, означало, что блоки действительно выполнялись одновременно. Но они были разочарованы, когда изменение драйвера, среды выполнения CUDA или GPU опровергло это предположение, нарушив их код.

Для вашего приложения попытайтесь найти решение, которое не зависит от межблочной синхронизации, потому что (за исключением изменения значения в модели программирования CUDA) это просто невозможно.

Answer 2

Возможна блокировка синхронизации блока.См. Эту статью .
В статье не очень подробно рассказывается о том, как она работает, но она опирается на операцию __syncthreads ();создать барьер паузы для текущего блока, ... ожидая, пока другие блоки дойдут до точки синхронизации.

Один элемент, который не отмечен в статье, заключается в том, что синхронизация возможна только в том случае, есликоличество блоков достаточно мало или количество SM достаточно велико для выполнения поставленной задачи.т.е. если у вас есть 4 SM и вы пытаетесь синхронизировать 5 блоков, ... ядро ​​заходит в тупик.

Благодаря их подходу я смог распределить длинную последовательную задачу среди множества блоков, легко сэкономив 30%время на подходе одного блока.т.е. у меня сработала блочная синхронизация.

Answer 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;

В случае кода ЦП такое поведение предотвращается объявлением переменной с префиксом volatile. Но даже если мы объявим count как int __device__ count (и соответствующим образом изменим код), добавление спецификатора volatile просто нарушит компиляцию (с ошибками loke argument of type "volatile int *" is incompatible with parameter of type "void *")

Я предлагаю посмотреть на пример threadFenceReduction из CUDA SDK. Там они делают то же самое, что и вы, но блок для окончательного суммирования выбирается во время выполнения, а не предопределено, и исключается while -loop, поскольку спин-блокировка для глобальной переменной должна быть очень медленно.