Сокращение суммы OpenCL в разных рабочих группах дает неверный результат - PullRequest
Новая
       22

Сокращение суммы OpenCL в разных рабочих группах дает неверный результат

0 голосов
/ 09 ноября 2018

Итак, в настоящее время я пытаюсь написать ядро ​​в OpenCL с целью суммирования каждой строки матрицы (g_idata) в массив (g_odata). Указанная матрица представлена ​​массивом с плавающей запятой с длиной column_count * row_count, а результирующий массив имеет длину row_count. В качестве такового я реализовал следующее ядро: 

#define T float
#define Operation(X, Y) ((X) + (Y))

__kernel void marrow_kernel( __global T *g_odata,__global T *g_idata, 
            const unsigned long column_count, const unsigned long row_count, __local volatile T* sdata) {

size_t tid = get_local_id(0);
size_t gid = get_global_id(0);

size_t row = gid / column_count;
size_t column = gid % column_count;


if(row < row_count && column < column_count)
{
    sdata[tid] = g_idata[gid];
}
barrier(CLK_LOCAL_MEM_FENCE);

if(row < row_count && column < column_count)
{
   size_t step = column_count / 2;
   size_t limit = column_count;

   while(step > 0)
   {
        if(column + step < limit) {
            if(tid + step < get_local_size(0))
            {
                sdata[tid] = Operation(sdata[tid], sdata[tid + step]);
            }
            else if (gid + step < column_count * row_count)
            {
                sdata[tid] = Operation(sdata[tid], g_idata[gid + step]);
            }
        }

        barrier(CLK_LOCAL_MEM_FENCE);

        step /= 2;
        limit /= 2;
   }
}


barrier(CLK_LOCAL_MEM_FENCE);

if(row < row_count && column == 0)
{
    g_odata[row] = column_count % 2 == 0 ? sdata[tid] : sdata[tid] + g_idata[gid + (column_count - 1)];
}

}

В настоящее время указанное ядро ​​создается с рабочей группой из 128 рабочих блоков. В настоящее время я не контролирую размер рабочей группы.

Теперь вот проблема: если, скажем, у меня есть строка, которая разделена между двумя различными рабочими группами, он вернет неправильный результат, так как он получит значение в g_idata, так как невозможно получить доступ к результат следующей рабочей группы локальной памяти. После первой итерации это неправильное значение, и оно повлияет на конечный результат операции.

Кто-нибудь может подсказать, как решить эту проблему?

...