Заставить все потоки в рабочей группе выполнять одну и ту же ветвь if/else

Question

Я хотел бы использовать оптимизацию локальной/разделяемой памяти, чтобы уменьшить доступ к глобальной памяти, поэтому у меня в основном есть эта функция

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

Теперь сложность в том, что вызывается в ядре только в одной из двух возможных ветвей if/else. Если только некоторые потоки в рабочей группе выполняют ветвь else, все остальные потоки не должны выбирать ветвь if для оптимизации локальной памяти в test_optimizedработать. Поэтому я придумал обходной путь: условие для каждого потока в рабочей группе atomic_or-ed в целое число, а затем целое число, одинаковое для всех потоков, проверяется на ветвление. Это гарантирует, что если 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;
    }
}

Приведенный выше код работает безупречно и примерно на 25% быстрее, чем с test_unoptimizedфункция. Но атомарное вклинивание в одну и ту же локальную память из всех потоков в рабочей группе кажется мне немного хаком, и оно эффективно работает только для небольшого размера рабочей группы ( def_ws) 32, 64 или 128, но не 256 или выше.

Используется ли этот трюк в других кодах и есть ли у него название? Если нет: есть ли лучший способ сделать это?

Answer 1

С OpenCL 1.2 или старше я не думаю, что есть способ сделать это быстрее. (Мне неизвестны какие-либо соответствующие расширения поставщиков, но проверьте список своей реализации на наличие чего-либо многообещающего.)

С OpenCL 2.0+ вы можете использовать функции рабочей группы, в данном случае конкретноwork_group_any()для такого рода вещей.