Я уже довольно давно пытаюсь заставить работать простое сканирование. Для небольших проблем вывод правильный, однако для больших выходов я получаю правильные результаты только иногда. Я проверил пример OpenCL от Apple и Я в основном делаю то же самое (за исключением конфликтов с банками, которые я игнорирую). Итак, вот код для первой фазы:
__kernel void
scan_init(__global int * input,
__global int * sums)
{
int gid = get_global_id(0);
int lid = get_local_id(0);
int chunk_size = get_local_size(0)*2;
int chunk = gid/chunk_size;
int offset = chunk*chunk_size;
reduction(input, offset);
// store sums
if(lid==0)
{
sums[chunk] = input[(chunk+1)*chunk_size-1];
}
downsweep(input, offset);
}
И сама функция редукции:
void reduction(__global int * input,
int offset)
{
int stride = 1;
int grp_size = get_local_size(0);
int lid = get_local_id(0);
for(int d = grp_size; d > 0; d>>=1)
{
barrier(CLK_GLOBAL_MEM_FENCE);
if(lid < d)
{
int ai = stride*(2*lid+1)-1+offset;
int bi = stride*(2*lid+2)-1+offset;
input[bi] += input[ai];
}
stride *= 2;
}
}
На втором этапе частичные суммы используются для построения суммы для каждого элемента:
void downsweep(__global int * input,
const unsigned int offset)
{
int grp_size = get_local_size(0);
int lid = get_local_id(0);
int stride = grp_size*2;
for(int d = 1; d <= grp_size; d *=2)
{
barrier(CLK_GLOBAL_MEM_FENCE);
stride >>=1;
if(lid+1 < d)
{
int src = 2*(lid + 1)*stride-1+offset;
int dest = src + stride;
input[dest]+=input[src];
}
}
}
Входные данные дополняются до размера, кратного локальному рабочему размеру. Каждая рабочая группа может сканировать фрагмент в два раза больше. Я сохраняю сумму каждого фрагмента в массиве sums, который я использую для проверки результата. Ниже приведен вывод для входного размера 4000 массива 1:
Chunk size: 1024
Chunks: 4
Scan global size: 4096
Local work size: 512
Sum size: 4
0:1024 1:1120 2:2904 3:928
Однако ожидаемый результат будет
0:1024 1:1024 2:1024 3:928
Если я снова запущу код, я получу:
0:1056 1:5376 2:1024 3:928
0:1024 1:1088 2:1280 3:992
0:5944 1:11156 2:3662 3:1900
0:7872 1:1056 2:2111 3:1248
Обращение к ядру выглядит следующим образом:
clEnqueueNDRangeKernel(cl_ctx->queue, scan_init, 1, NULL, &scan_global_size, &local_work_size, 0, NULL, NULL);
Где глобальный размер — 4096, а локальный — 512. Если я ограничу размер локальной рабочей группы до 64, результат будет выглядеть следующим образом:
0:128 1:128 2:128 3:288 4:128 5:128 6:192 7:192
8:192 9:254 10:128 11:256 12:128 13:360 14:128 15:128
16:128 17:128 18:128 19:288 20:128 21:128 22:128 23:128
24:192 25:128 26:128 27:192 28:128 29:128 30:128 31:32
И если я изменю размер ввода на 512 и любой размер кусков, все работает отлично!
Наконец, при использовании размера ввода 513 и размера группы 256 (то есть у меня есть два фрагмента, каждый из которых имеет 512 элементов, а второй имеет только первый элемент, установленный в 1), результат первой фазы:
0:1 1:2 2:1 3:6 4:1 5:2 6:1 7:14
8:1 9:2 10:1 11:6 12:1 13:2 14:1 15:28
16:1 17:2 18:1 19:6 20:1 21:2 22:1 23:14
24:1 25:2 26:1 27:6 28:1 29:2 30:1 31:56
32:1 33:2 34:1 35:6 36:1 37:2 38:1 39:14
40:1 41:2 42:1 43:6 44:1 45:2 46:1 47:28
48:1 49:2 50:1 51:6 52:1 53:2 54:1 55:14
56:1 57:2 58:1 59:6 60:1 61:2 62:1 63:148
Где это должно быть:
0:1 1:2 2:1 3:4 4:1 5:2 6:1 7:8
8:1 9:2 10:1 11:4 12:1 13:2 14:1 15:16
16:1 17:2 18:1 19:4 20:1 21:2 22:1 23:8
24:1 25:2 26:1 27:4 28:1 29:2 30:1 31:32
32:1 33:2 34:1 35:4 36:1 37:2 38:1 39:8
40:1 41:2 42:1 43:4 44:1 45:2 46:1 47:16
48:1 49:2 50:1 51:4 52:1 53:2 54:1 55:8
56:1 57:2 58:1 59:4 60:1 61:2 62:1 63:64
Я предполагаю, что это проблема с одновременным доступом к одним и тем же данным из разных потоков, однако этого не должно быть, поскольку каждая рабочая группа обрабатывает разные фрагменты входных данных. Любая помощь по этому вопросу будет принята с благодарностью!
