在块中找到最大值的 OpenCL 障碍

Posted 2023-03-17

【中文标题】在块中找到最大值的 OpenCL 障碍

【英文标题】：OpenCL barrier of finding max in a block

【发布时间】：2015-08-14 10:37:01

【问题描述】：

我在 Nvidia 的开发者网站上找到了一段 OpenCL 内核示例代码 目的函数maxOneBlock是找出数组maxValue的最大值并存入maxValue[0]。

我完全了解循环部分，但对unroll 部分感到困惑：为什么展开部分不需要在每个步骤完成后同步线程？

e.g：当一个线程完成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];
    

【参考方案1】：

为什么展开部分不需要在每一步完成后同步线程？

样本不正确，每一步之后确实需要一个屏障。

看起来样本是用 warp-synchronous 风格编写的，这是一种利用 NVIDIA 硬件上 warp 的底层执行机制的方式，但如果底层执行机制发生变化或存在，不正确的同步会导致它中断编译器优化。

【讨论】：

即使是warp同步方式写的，unroll部分也需要在每一步之后限制线程。即第 1 步限制 32 个线程，第 2 个 16 个线程……等等。但事实并非如此，所有 32 个线程都运行了整个 unroll 代码。

是的，但没有使用他们的结果：他们只是免费做额外的工作。作者没有在每次迭代时禁用半线程，而是选择让它们运行。它使代码更简单，不会影响性能或最终结果。这是一种比较常见的技术。但这并没有使障碍成为可选的，在这方面，示例是错误的。