2

我想使用本地/共享内存优化来减少全局内存访问，所以我基本上有这个功能

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

并对其进行通常的本地/共享内存优化：

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-ed 转换为一个整数，然后检查对于所有线程都相同的整数是否有分支。这确保了，如果线程块中的 1 个或多个线程选择 else-branch，所有其他线程也会这样做。

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%。但是对我来说，从工作组中的所有线程以原子方式将一点干扰到相同的本地内存中似乎有点像 hack，它只对小型工作组大小 ( def_ws) 32、64 或 128 有效运行，而不是 256 或更大。

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

4

1 回答 1

1

对于 OpenCL 1.2 或更早版本，我认为没有办法更快地做到这一点。（我不知道有任何相关的供应商扩展，但请检查您的实现列表是否有任何希望。）

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

于 2021-07-18T14:32:53.327 回答