Согласно спецификации, OpenCL попытается объединить доступ к памяти из соседних потоков, но я не могу найти правильного определения соседних потоков. Например, если у меня есть 2D-рабочая группа размером 32x32, я знаю, что рабочий элемент из той же строки (0, 0) и (1, 0) являются соседними потоками, но как насчет рабочего элемента из того же столбца (0, 0) и (0, 1)? Я спрашиваю об этом, потому что не могу объяснить свои результаты тестирования простого ядра транспонирования матрицы:
__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];
}
work group size execution time
x y
4 64 24ms
1 256 34ms
, когда размер рабочей группы равен 1x256, доступ на запись в память из соседних потоков можно объединить, но Я не ожидаю, что это произойдет, когда размер 4x64, однако он имеет лучшую производительность, чем 1x256. В чем причина этого?