for-l oop прерывание замедляет ядро ​​OpenCL

Question

Изучение реализующей OpenCL программы, которая находит 8 наиболее похожих пикселей в радиусе 10 пикселей для каждого пикселя:

//-DBEST=8

__kernel void best_pixel_matches(
    __global const uint * image,
    const ushort width,
    const ushort height,
    // [ width ] [ height ] [ BEST ] [ 2 ]
    __global short * best,
    __global int * errors
)
{
    const int x = get_global_id(0);
    const int y = get_global_id(1);

    const uint pixel = image[ y * width + x ];

    const short SURROUND = 10;

    //local best pixels coordinates:
    short s[ BEST ][ 2 ];
    //... corresponding local best pixels computed deltas (storage instead of recompute speeds program about twice): 
    uint d[ BEST ];
    //... init with none:
    for ( short i = 0; i < BEST; ++ i ) d[ i ] = -1;

    for ( short sw = - SURROUND; sw < SURROUND; ++ sw ) {
    for ( short sh = - SURROUND; sh < SURROUND; ++ sh ) {
        //avoid inexistent pixels:
        const short tw = x + sw; if ( tw < 0 || tw >= width  ) continue;
        const short th = y + sh; if ( th < 0 || th >= height ) continue;
        //not with itself:
        if ( tw == x && th == y ) continue;
        const uint diff = (uint) abs( (long)pixel - (long)image[ th * width + tw ] );

        for ( short si = 0; si < BEST; ++ si ) {
            if ( d[si] == -1 || diff < d[si] ) {
                d[si] = diff;
                s[si][0] = tw;
                s[si][1] = th;
                break;//<-- this line causes 4 (!) times slow down
            }
        }
    } }

    //copy results from private memory to global:
    const long p = ( x * height + y ) * BEST;
    for ( short b = 0; b < BEST; ++ b ) {
        const long pb = ( p + b ) * 2;
        best[ pb     ] = s[b][0];
        best[ pb + 1 ] = s[b][1];
    }
}

Проблема заключалась в том, что потребовалось GPU 3593ms для изображения 2560 *1440* 1005 *, которое это почти то же самое время, что и простой C++ код процессора, который я имел раньше (~ 8500 мс). Затем я попытался настроить его здесь и там и иногда удалял break; строку и время выполнения -> 900 мс !

В чем причина такого удивительного ускорения? Эта строка break; просто сообщает программе, что нет необходимости проверять другие пиксели, поэтому предполагается уменьшить время выполнения, а не замедлить его? Может быть, есть и другие способы ускорить эту программу? :)

Answer 1

Вы не говорите, на каком устройстве вы это запускаете, но:

• В ЦП оптимизация включает в себя автовекторизацию, которая обычно очень чувствительна к ветвям. Так что коэффициент 4 не удивителен; Ваши числа итераций l oop известны во время сборки, поэтому они могут быть полностью развернуты и векторизованы, если нет раннего выхода. С ранним выходом код, вероятно, должен быть сериализован, поэтому вы теряете ускорение в 4, 8 или 16 раз из-за векторизации (в зависимости от типа процессора) и возвращаете только часть этого из раннего.
• Рабочие элементы GPU обычно выполняются в потоках с блокировкой, но изменение фактора 4 кажется маловероятным, если это происходит здесь; как правило, более вероятно, что break; не принесет никаких улучшений, если не будет высокой вероятности того, что весь волновой фронт воспользуется ярлыком. Однако некоторые модели графических процессоров предпочитают код SIMD, поэтому вы можете столкнуться с ситуацией, аналогичной описанной для CPU.

Обратите внимание, что удаление break, вероятно, заполнит все элементы BEST в d & s с тем же значением (если я правильно читаю ваш код), поэтому он не выдаст такой же вывод, как с break;.

Answer 2

Старайтесь поддерживать порядок доступа к памяти, непосредственно перед глобальной записью

 // here
 barrier(CLK_GLOBAL_MEM_FENCE);

 //copy results from private memory to global:
 const long p = ( x * height + y ) * BEST;

, если ваша проблема производительности связана с шаблоном доступа к памяти, это должно облегчить ее.

Если проблема возникает только из-за переходя между потоками, вы также можете попытаться сохранить порядок во втором l oop после самого внутреннего l oop:

for ( short si = 0; si < BEST; ++ si ) {
            if ( d[si] == -1 || diff < d[si] ) {
                d[si] = diff;
                s[si][0] = tw;
                s[si][1] = th;
                break;
            }
        }
barrier(CLK_LOCAL_MEM_FENCE); // local => should be faster for many devices

, чтобы на каждой итерации все соседние потоки объединялись и продолжали выполнять одни и те же инструкции пока новый разрыв не нарушит порядок.

Наконец, алгоритм выглядит как сортировка различий значений пикселей между ОБЪЕМНЫМИ пикселями путем уменьшения значений. Таким образом, самый внутренний l oop может быть фактически перемещен к самому внешнему, и все ветвления должны быть минимизированы, сразу после двух самых внутренних новых петель (которые были самыми внешними раньше). Но это увеличивает image[ th * width + tw ] чтение в BEST раз, поэтому оно может быть медленнее (возможно, не тогда, когда данные перемещаются в локальную память до этого). Но теперь d не обязательно должен быть массивом, поэтому ЛУЧШЕЕ число частных регистров сохраняется и может повысить производительность за счет уменьшения давления в регистре.