OpenCL сокращение от частного до локального, а затем глобального?

Question

Следующее ядро ​​вычисляет поле акустического давления, при этом каждая нить вычисляет собственный частный экземпляр вектора pressure, который затем следует суммировать в глобальную память. Я уверен, что код, который вычисляет вектор pressure, верен, но у меня все еще возникают проблемы с получением ожидаемого результата.

int gid  = get_global_id(0); 
int lid  = get_local_id(0); 
int nGroups = get_num_groups(0); 
int groupSize = get_local_size(0); 
int groupID = get_group_id(0); 

/* Each workitem gets private storage for the pressure field. 
* The private instances are then summed into local storage at the end.*/ 
private float2 pressure[HYD_DIM_TOTAL]; 
local float2 pressure_local[HYD_DIM_TOTAL]; 

/* Code which computes value of 'pressure' */ 

//wait for all workgroups to finish accessing any memory 
barrier(CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE); 

/// sum all results in a workgroup into local buffer: 
for(i=0; i<groupSize; i++){ 

    //each thread sums its own private instance into the local buffer 
    if (i == lid){ 
     for(iHyd=0; iHyd<HYD_DIM_TOTAL; iHyd++){ 
      pressure_local[iHyd] += pressure[iHyd]; 
     } 
    } 
    //make sure all threads in workgroup get updated values of the local buffer 
    barrier(CLK_LOCAL_MEM_FENCE); 
} 

/// copy all the results into global storage 
//1st thread in each workgroup writes the group's local buffer to global memory 
if(lid == 0){ 
    for(iHyd=0; iHyd<HYD_DIM_TOTAL; iHyd++){ 
     pressure_global[groupID +nGroups*iHyd] = pressure_local[iHyd]; 
    } 
} 

barrier(CLK_GLOBAL_MEM_FENCE); 

/// sum the various instances in global memory into a single one 
// 1st thread sums global instances 
if(gid == 0){ 

    for(iGroup=1; iGroup<nGroups; iGroup++){ 

     //we only need to sum the results from the 1st group onward 
     for(iHyd=0; iHyd<HYD_DIM_TOTAL; iHyd++){ 

      pressure_global[iHyd] += pressure_global[iGroup*HYD_DIM_TOTAL +iHyd]; 
      barrier(CLK_GLOBAL_MEM_FENCE); 
     } 
    } 
} 

Некоторые примечания по размерам данных: Общее количество потоков будет варьироваться в пределах от 100 до 2000, но могут иногда лежат вне этого интервала.
groupSize будет зависеть от аппаратного обеспечения, но в настоящее время я использую значения между 1 (cpu) и 32 (gpu).
HYD_DIM_TOTAL известен во время компиляции и варьируется от 4 до 32 (обычно, но необязательно, может быть мощностью 2).

Есть ли что-то откровенно неправильное с этим кодом сокращения?

PS: Я запускаю это на i7 3930k с AMD APP SDK 2.8 и на NVIDIA GTX580.

Answer 1

я замечаю две проблемы здесь, один большой, один меньше:

• Этот код предполагает, что у вас есть непонимание того, что делает барьер. Барьер никогда не синхронизируется между несколькими рабочими группами. Он синхронизируется только внутри рабочей группы. CLK_GLOBAL_MEM_FENCE делает его похожим на глобальную синхронизацию, но на самом деле это не так. Этот флаг просто заворачивает все обращения к текущему рабочему элементу в глобальную память. Таким образом, выдающиеся записи будут глобально наблюдаться после барьера с этим флагом. Но это не изменяет поведение синхронизации барьера, которое находится только в области рабочей группы. Глобальной синхронизации в OpenCL нет, кроме запуска другого NDRange или Task.
• Первый цикл for заставляет несколько рабочих элементов перезаписывать вычисления друг друга. Индексация давления_local с iHyd будет выполняться каждым рабочим элементом с тем же iHyd. Это приведет к неопределенным результатам.

Надеюсь, это поможет.