Почему размер рабочей группы OpenCL оказывает огромное влияние на производительность графического процессора?

Question

Я тестирую простое ядро ​​транспонирования матрицы на Qualcomm Adreno 630 GPU, и я пытаюсь увидеть влияние различного размера рабочей группы, но, к удивлению, я получаю интересный результат, который я не могу объяснить. Вот мой код ядра:

__kernel void transpose(__global float *input, __global float *output, const int width, const int height)
    int i = get_global_id(0);
    int j = get_global_id(1);
    output[i*height + j] = input[j*width + i];
}

и ширина и высота равны 6400, результаты эксперимента (время выполнения - это разница между событиями END и START):

work group size      execution time
x     y
4    64              24ms
64   4               169ms
256  1               654ms
1    256             34ms
8    32              27ms
1    1024            375ms
1024 1               657ms
32   32              26ms

после этого я провел еще один эксперимент, в котором я изменил ширину и высоту с 6400 на 6401 (а также глобальный рабочий размер в вызове NDRangeKernel), и результат оказался еще более интересным:

work group size      execution time
x     y
4    64              28ms
64   4               105ms
256  1               359ms
1    256             31ms
8    32              32ms
1    1024            99ms
1024 1               358ms
32   32              32ms

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

Answer 1

Объединение памяти происходит, когда последовательные потоки обращаются к данным по последовательным адресам глобальной памяти в пределах 128-байтового выровненного сегмента. Затем обращения к памяти объединяются в один, что значительно снижает общую задержку.

В 2D-диапазоне объединение происходит только в направлении get_global_id(1) или j в вашем случае. В строке output[i*height + j] = input[j*width + i];, input[j*width + i]; - это неверно выровненное (не объединенное) чтение, а output[i*height + j] - объединенная запись. Доступ к объединенной памяти обычно намного быстрее, чем доступ с несогласованным доступом, но снижение производительности для объединенного / несогласованного чтения может сильно отличаться от объединенного / несогласованного чтения. На большинстве настольных архитектур графических процессоров комбинация несовмещенного чтения и объединенной записи выполняется быстрее, чем наоборот, см. Диаграмму ниже. Таким образом, ваша реализация уже должна быть более быстрым вариантом.

Since coalesced access is only possible along the j index, if you have a range of (x=256,y=1) (i along x-direction, j along y-direction), you do not get any coalescing. For (x=8,y=32), j is coalesced in groups of 32 8 times per thread block, so memory bandwidth is fairly saturated and performance is good.

If you want maximum possible performance, I'd suggest you go with 1D indexing. This way you have full control about coalescing and coalescing happens over the entire thread block. Your matrix transpose kernel then would look like this:

#define width 6400
__kernel void transpose(__global float *input, __global float *output) {
    const int n = get_global_id(0);
    int i = n/width;
    int j = n%width;
    output[i*height + j] = input[j*width + i];
}

Вы можете записать width в Ccode OpenCL во время выполнения C ++ и до времени компиляции OpenCL через конкатенация строк.