boosting parallel reduction OpenCL

Question

I have an algorithm, performing two-staged parallel reduction on GPU to find the smallest elemnt in a string. I know that there is a hint on how to make it work faster, but I don't know what it is. Any ideas on how I can tune this kernel to speed my program up? It is not necessary to actually change algorithm, may be there are other tricks. All ideas are welcome.

Thank you!

__kernel
void reduce(__global float* buffer,
            __local float* scratch,
            __const int length,
            __global float* result) {    
    int global_index = get_global_id(0);
    float accumulator = INFINITY
        while (global_index < length) {
            float element = buffer[global_index];
            accumulator = (accumulator < element) ? accumulator : element;
            global_index += get_global_size(0);
    }
    int local_index = get_local_id(0);
    scratch[local_index] = accumulator;
    barrier(CLK_LOCAL_MEM_FENCE);
    for(int offset = get_local_size(0) / 2;
        offset > 0;
        offset = offset / 2) {
            if (local_index < offset) {
                float other = scratch[local_index + offset];
                float mine = scratch[local_index];
                scratch[local_index] = (mine < other) ? mine : other;
            }
        barrier(CLK_LOCAL_MEM_FENCE);
    }
    if (local_index == 0) {
        result[get_group_id(0)] = scratch[0];
    }
}

Evgeny Panasyuk Evgeny Panasyuk · Accepted Answer · 2013-04-28T12:14:42

accumulator = (accumulator < element) ? accumulator : element;

Use fmin function - it is exactly what you need, and it may result in faster code (call to built-in instruction, if available, instead of costly branching)

global_index += get_global_size(0);

What is your typical get_global_size(0)?

Though your access pattern is not very bad (it is coalesced, 128byte chunks for 32-warp) - it is better to access memory sequentially whenever possible. For instance, sequential access may aid memory prefetching (note, OpenCL code can be executed on any device, including CPU).

Consider following scheme: each thread would process range

[ get_global_id(0)*delta ,  (get_global_id(0)+1)*delta )

It will result in fully sequential access.

boosting parallel reduction OpenCL

1 Answers