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

Question

Я нашел образец кода ядра 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]; 
    } 
} 

Answer 1

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

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

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