强制一个工作组中的所有线程执行同一个 if/else 分支

Question

我想用local/shared内存优化来减少全局内存访问，所以我基本都有这个功能

float __attribute__((always_inline)) test_unoptimized(const global float* data, ...) {
    // ...
    for(uint j=0; j<def_data_length; j++) {
        const float x = data[j];
        // do sime computation with x, like finding the minimum value ...
    }
    // ...
    return x_min;
}

并对其进行通常的 local/shared 内存优化：

float __attribute__((always_inline)) test_optimized(const global float* data, ...) {
    // ...
    const uint lid = get_local_id(0); // shared memory optimization (only works with first ray)
    local float cache_x[def_ws];
    for(uint j=0; j<def_data_length; j+=def_ws) {
        cache_x[lid] = data[j+lid];
        barrier(CLK_LOCAL_MEM_FENCE);
        #pragma unroll
        for(uint k=0; k<min(def_ws, def_data_length-j); k++) {
            const float x = cache_x[k];
            // do sime computation with x, like finding the minimum value ...
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }
    // ...
    return x_min;
}

现在的困难是 test_optimized 在内核中仅在两个可能的 if/else 分支之一中被调用。如果工作组中只有部分线程执行else-branch，则所有其他线程不得选择if-branch以使test_optimized中的本地内存优化起作用。所以我创建了一个解决方法：将工作组中每个线程的条件 atomic_or 编辑成一个整数，然后检查所有线程都相同的整数是否有分支。这确保了，如果线程块中的 1 个或多个线程选择了 else 分支，则所有其他线程也会这样做。

kernel void test_kernel(const global float* data, global float* result...) {
    const uint n = get_global_id(0);
    
    // ...
    const bool condition = ...; // here I get some condition based on the thread ID n and global data

    local uint condition_any; // make sure all threads within a workgroup are in the if/else part
    condition_any = 0u;
    barrier(CLK_LOCAL_MEM_FENCE);
    atomic_or(&condition_any, condition);
    barrier(CLK_LOCAL_MEM_FENCE);

    if(condition_any==0u) {
        // if-part is very short
        result = 0;
        return;
    } else {
        // else-part calls test_optimized function
        const float x_min = test_optimized(data, ...);
        result = condition ? x_min : 0;
    }
}

以上代码运行完美，比 test_unoptimized 函数快约 25%。但是从工作组中的所有线程原子地干扰到同一个本地内存中对我来说有点像黑客，它只对小工作组大小有效运行（def_ws）32、64 或 128，但不是 256 或更大。

这个技巧在其他代码中使用过吗？它有名字吗？如果没有：有更好的方法吗？

Answer 1

对于 OpenCL 1.2 或更早版本，我认为没有任何方法可以更快地完成此操作。（我不知道有任何相关的供应商扩展，但请检查您的实施列表中是否有任何有前途的内容。）

使用 OpenCL 2.0+，您可以使用工作组功能，在这种情况下特别是 work_group_any() 用于此类事情。

强制一个工作组中的所有线程执行同一个 if/else 分支

Force all threads in a work group to execute the same if/else branch

gpgpu

atomic

shared-memory

opencl

gpu-local-memory