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

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

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 (процессор) до 32 (графический процессор). ).
HYD_DIM_TOTAL известно во время компиляции и варьируется от 4 до 32 (обычно, но не обязательно, это степень числа 2).

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

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