OpenCL - использование атомной редукции для двойного

Question

Я знаю, что атомные функции с OpenCL-1.x не рекомендуются, но я просто хочу понять атомный пример.

Следующий код ядра не работает хорошо, он производит случайные конечные значения для вычисления суммы всех значений массива (уменьшение суммы):

#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable 

void atom_add_double(volatile __local double *val, double delta) 
{ 
    union { 
    double f; 
    ulong i; 
    } old, new; 

    do 
    { 
    old.f = *val; 
    new.f = old.f + delta; 
    } 
    while (atom_cmpxchg((volatile __local ulong *)val, old.i, new.i) != old.i); 

} 

__kernel void sumGPU (__global const double *input, 
       __local double *localInput, 
       __global double *finalSum 
       ) 
{ 

    uint lid = get_local_id(0); 
    uint gid = get_global_id(0); 
    uint localSize = get_local_size(0); 
    uint groupid = get_group_id(0); 
    local double partialSum; 
    local double finalSumTemp; 

// Initialize sums 
    if (lid==0) 
    { 
    partialSum = 0.0; 
    finalSumTemp = 0.0; 
    } 
    barrier(CLK_LOCAL_MEM_FENCE); 

    // Set in local memory 
    int idx = groupid * localSize + lid; 
    localInput[lid] = input[idx]; 

    // Compute atom_add into each workGroup 
    barrier(CLK_LOCAL_MEM_FENCE); 
    atom_add_double(&partialSum, localInput[lid]); 
    // See and Check if barrier below is necessary 
    barrier(CLK_LOCAL_MEM_FENCE); 

    // Final sum of partialSums 
    if (lid==0) 
    { 
    atom_add_double(&finalSumTemp, partialSum); 
    *finalSum = finalSumTemp; 
    } 

}     

версии с global id стратегии работает хорошо, но версия выше , который проходит при использовании local memory (разделяемая память), не дает ожидаемых результатов (значение *finalSum является случайным для каждого исполнения).

Здесь Буфера и ядро ​​арг, что я положил в моем коде хоста:

// Write to buffers 
    ret = clEnqueueWriteBuffer(command_queue, inputBuffer, CL_TRUE, 0, 
     nWorkItems * sizeof(double), xInput, 0, NULL, NULL); 
    ret = clEnqueueWriteBuffer(command_queue, finalSumBuffer, CL_TRUE, 0, 
         sizeof(double), finalSumGPU, 0, NULL, NULL); 

// Set the arguments of the kernel 
    clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&inputBuffer); 
    clSetKernelArg(kernel, 1, local_item_size*sizeof(double), NULL); 
    clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&finalSumBuffer); 

и, наконец, я прочитал finalSumBuffer, чтобы получить значение суммы.

Я думаю, что моя проблема исходит скорее из кода ядра, но я не могу найти, где ошибка.

Если бы кто-нибудь мог понять, что случилось, это было бы неплохо рассказать мне.

Благодаря

UPDATE 1:

Я почти удалось выполнить это сокращение. После предложений, предложенных huseyın tuğrul buyukisik, я изменил код ядра, как это:

#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable 

void atom_add_double(volatile __local double *val, double delta) 
{ 
    union { 
    double d; 
    ulong i; 
    } old, new; 

    do 
    { 
    old.d = *val; 
    new.d = old.d + delta; 
    } 
    while (atom_cmpxchg((volatile __local ulong *)val, old.i, new.i) != old.i); 

} 

__kernel void sumGPU (__global const double *input, 
       __local double *localInput, 
       __local double *partialSum, 
       __global double *finalSum 
       ) 
{ 

    uint lid = get_local_id(0); 
    uint gid = get_global_id(0); 
    uint localSize = get_local_size(0); 
    uint groupid = get_group_id(0); 

    // Initialize partial sums 
    if (lid==0) 
    partialSum[groupid] = 0.0; 


    barrier(CLK_LOCAL_MEM_FENCE); 
    // Set in local memory 
    int idx = groupid * localSize + lid; 
    localInput[lid] = input[idx]; 

    // Compute atom_add into each workGroup 
    barrier(CLK_LOCAL_MEM_FENCE); 
    atom_add_double(&partialSum[groupid], localInput[lid]); 
    // See and Check if barrier below is necessary 
    barrier(CLK_LOCAL_MEM_FENCE); 

    // Compute final sum 
    if (lid==0) 
    *finalSum += partialSum[groupid]; 

}     

Как сказал Гусейн, мне не нужно использовать атомарные функции для окончательной суммы всех частичных сумм ,

Так что я сделал в конце:

// Compute final sum 
    if (lid==0) 
    *finalSum += partialSum[groupid]; 

Но, к сожалению, окончательная сумма не дает ожидаемого значения и значение является случайным (например, с nwork-items = 1024 и size-WorkGroup = 16, я получаю случайные значения в . порядок [1e+3 - 1e+4] вместо 5.248e+05 ожидается

Вот настройка аргументов в принимающем код:

// Set the arguments of the kernel 
    clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&inputBuffer); 
    clSetKernelArg(kernel, 1, local_item_size*sizeof(double), NULL); 
    clSetKernelArg(kernel, 2, nWorkGroups*sizeof(double), NULL); 
    clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *)&finalSumBuffer); 

Не могли бы вы увидеть, где моя ошибка в коде ядра?

Благодаря

Answer 1

не ошибка, но логический вопрос:

atom_add_double(&finalSumTemp, partialSum); 

работает только один раз в каждой группе (по нулевой локальной индексируется нить).

Таким образом, вы просто делаете

finalSumTemp = partialSum 

так атомарных здесь не требуется.

---

Существует раса условие

*finalSum = finalSumTemp; 

между рабочими группами, где каждый нулевой индекс локальный поток пишет в тот же адрес. Итак, это должно быть добавление атома (для учебных целей) или может быть написано на разных ячейках, которые должны быть добавлены на стороне хоста, такие как sum_group1 + sum_group2 + ... = общая сумма.

---

int idx = groupid * localSize + lid; 
localInput[lid] = input[idx]; 

здесь, используя GroupID подозрительна для нескольких устройств суммирования. Поскольку каждое устройство имеет свои собственные индексы индексов глобальных индексов и индексов рабочих групп, два устройства могут иметь одинаковые значения идентификатора группы для двух разных групп. Некоторое смещение, связанное с устройством, должно использоваться при использовании нескольких устройств. Такие, как:

idx= get_global_id(0) + deviceOffset[deviceId]; 

---

Кроме того, если атомная операция inavoidable, и если ровно N раз в действие, он может быть перемещен в один поток (например, 0-индексированные нити) и петельные для N раз (вероятно, быстрее) во втором ядре, если эта латентность атомной операции не может быть скрыта другими средствами.