opnecl dot товар. Здесь я пытаюсь сохранить результат в локальной переменной, которая сбрасывается до нуля каждый раз, когда

Question

мой код ядра.

__kernel void OUT__1__1527__(__constant float *A,__constant float *B,__global float *res) 
{ 
    int i = get_global_id(0); 
    float C=0; 
    if (i <= 5 - 1) { 
    C += (A[i] * B[i]); 
    *res=C; 
} 
} 

& В имеют значение {1,2,3,4,5} обоих. для этого ядра я получаю результат 25, который равен 5 * 5, когда я хочу, чтобы результат был равен 55. (1 * 1 + 2 * 2 + 3 * 3 + 4 * 4 + 5 * 5)

какой код необходимо вставить для синхронизации и где код должен быть вставлен.

Answer 1

Для этих вещей нет никакого волшебного исправления. Это типичная проблема сокращения (где вы хотите объединить несколько результатов в одну переменную).

Вы можете использовать атоматику (но не с поплавковыми значениями), если это не является узким местом вашего алгоритма (т. Е. Вы выполняете другой более дорогостоящий процесс в другом месте ядра). Но если это Ядро ядра. Затем вы должны перейти, чтобы полностью изменить свой алгоритм.

Вы можете начать читать это: http://developer.amd.com/resources/documentation-articles/articles-whitepapers/opencl-optimization-case-study-simple-reductions/

Кроме того, ваш код не так, он фактически никогда не «сливать» s любые данные в res. он устанавливает только res в значение C. И C является приватным для каждого потока, поэтому он ничего не суммирует. Только последняя нить на самом деле выигрывает datarace, что приводит к ответу будучи 25.

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

inline void AtomicAdd(volatile __global float *source, const float operand) { 
    union { 
     unsigned int intVal; 
     float floatVal; 
    } newVal; 
    union { 
     unsigned int intVal; 
     float floatVal; 
    } prevVal; 
    do { 
     prevVal.floatVal = *source; 
     newVal.floatVal = prevVal.floatVal + operand; 
    } while (atomic_cmpxchg((volatile __global unsigned int *)source, prevVal.intVal, newVal.intVal) != prevVal.intVal); 
} 

__kernel void OUT__1__1527__(__constant float *A,__constant float *B,__global float *res) 
{ 
    int i = get_global_id(0); 
    if (i <= 5 - 1) { 
    AtomicAdd(res, (A[i] * B[i])); 
} 
}