OpenCL барьер нахождения максимума в блоке

Я нашел фрагмент примера кода ядра OpenCL на сайте разработчиков Nvidia. Функция maxOneBlock предназначена для нахождения максимального значения массива maxValue и сохранения его в maxValue[0].

Я полностью понял часть зацикливания, но запутался в части unroll: почему в части развертывания не нужно синхронизировать поток после выполнения каждого шага?

например: когда один поток выполняет сравнение localId и localId+32, как он гарантирует, что другой поток сохранил свой результат в localId+16?

Код ядра:

void maxOneBlock(__local float maxValue[],
                 __local int   maxInd[])
{
    uint localId   = get_local_id(0);
    uint localSize = get_local_size(0);
    int idx;
    float m1, m2, m3;

    for (uint s = localSize/2; s > 32; s >>= 1)
    {
        if (localId < s) 
        {
            m1 = maxValue[localId];
            m2 = maxValue[localId+s];
            m3 = (m1 >= m2) ? m1 : m2;
            idx = (m1 >= m2) ? localId : localId + s;
            maxValue[localId] = m3;
            maxInd[localId] = maxInd[idx];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    // unroll the final warp to reduce loop and sync overheads
    if (localId < 32)
    {
        m1 = maxValue[localId];
        m2 = maxValue[localId+32];
        m3 = (m1 > m2) ? m1 : m2;
        idx = (m1 > m2) ? localId : localId + 32;
        maxValue[localId] = m3;
        maxInd[localId] = maxInd[idx];


        m1 = maxValue[localId];
        m2 = maxValue[localId+16];
        m3 = (m1 > m2) ? m1 : m2;
        idx = (m1 > m2) ? localId : localId + 16;
        maxValue[localId] = m3;
        maxInd[localId] = maxInd[idx];

        m1 = maxValue[localId];
        m2 = maxValue[localId+8];
        m3 = (m1 > m2) ? m1 : m2;
        idx = (m1 > m2) ? localId : localId + 8;
        maxValue[localId] = m3;
        maxInd[localId] = maxInd[idx];

        m1 = maxValue[localId];
        m2 = maxValue[localId+4];
        m3 = (m1 > m2) ? m1 : m2;
        idx = (m1 > m2) ? localId : localId + 4;
        maxValue[localId] = m3;
        maxInd[localId] = maxInd[idx];

        m1 = maxValue[localId];
        m2 = maxValue[localId+2];
        m3 = (m1 > m2) ? m1 : m2;
        idx = (m1 > m2) ? localId : localId + 2;
        maxValue[localId] = m3;
        maxInd[localId] = maxInd[idx];

        m1 = maxValue[localId];
        m2 = maxValue[localId+1];
        m3 = (m1 > m2) ? m1 : m2;
        idx = (m1 > m2) ? localId : localId + 1;
        maxValue[localId] = m3;
        maxInd[localId] = maxInd[idx];
    }
}

parallel-processing max opencl barrier gpgpu
person melode11    schedule 01.06.2015    source источник

Ответы (1)


arrow_upward
0
arrow_downward

Почему часть развертывания не требует синхронизации потока после выполнения каждого шага?

Образец неверный, действительно требуется барьер после каждого шага.

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

person user703016    schedule 01.06.2015
comment
Даже если он написан в стиле варп-синхронности, часть unroll должна ограничивать потоки после каждого шага. т. е. 1-й шаг ограничивает 32 потока, 2-й 16 потоков... и т. д. Но это не так, все 32 потока выполняли весь код unroll. - person melode11; 01.06.2015
comment
Да, но их результат не используется: они просто делают лишнюю работу бесплатно. Вместо того, чтобы отключать половину потоков на каждой итерации, автор решил разрешить им работать. Это упрощает код, не влияет на производительность и конечный результат. Это относительно распространенная техника. Но это не делает барьеры необязательными, и в этом отношении выборка неверна. - person user703016; 01.06.2015