并行缩减算法中的 OpenCL 未定义行为

Question

我正在研究一个简单的并行缩减算法来查找数组中的最小值，并且在我的算法中遇到了一些有趣的未定义行为。我运行在 Ubuntu 16.04.

上使用英特尔的 OpenCL 1.2

以下内核是我正在尝试的运行目前给我错误的答案：

__kernel void Find_Min(int arraySize, __global double* scratch_arr, __global double* value_arr, __global double* min_arr){

    const int index = get_global_id(0);
    int length = (int)sqrt((double)arraySize);
    int start = index*length;
    double min_val = INFINITY;
    for(int i=start; i<start+length && i < arraySize; i++){
        if(value_arr[i] < min_val)
            min_val = value_arr[i];
    }
    scratch_arr[index] = min_val;

    barrier(CLK_GLOBAL_MEM_FENCE);
    if(index == 0){
        double totalMin = min_val;
        for(int i=1; i<length; i++){
            if(scratch_arr[i] < totalMin)
                totalMin = scratch_arr[i];
        }
        min_arr[0] = totalMin;
    }
}

当输入 {0,-1,-2,-3,-4,-5,-6,-7,-8} 数组时，它最终返回 -2。

这是出现未定义行为的地方。当我运行以下内核在屏障之前带有 printf 语句时，我得到了正确的答案 (-8)：

__kernel void Find_Min(int arraySize, __global double* scratch_arr, __global double* value_arr, __global double* min_arr){

    const int index = get_global_id(0);
    int length = (int)sqrt((double)arraySize);
    int start = index*length;
    double min_val = INFINITY;
    for(int i=start; i<start+length && i < arraySize; i++){
        if(value_arr[i] < min_val)
            min_val = value_arr[i];
    }
    scratch_arr[index] = min_val;
    printf("setting scratch[%i] to %f\n", index, min_val);

    barrier(CLK_GLOBAL_MEM_FENCE);
    if(index == 0){
        double totalMin = min_val;
        for(int i=1; i<length; i++){
            if(scratch_arr[i] < totalMin)
                totalMin = scratch_arr[i];
        }
        min_arr[0] = totalMin;
    }
}

我能想到的唯一可能发生的事情是我错误地使用了 barrier 命令，而 printf 正在做的所有事情导致内核延迟，它以某种方式同步调用，因此它们都在最后的还原步骤。但是如果没有 printf，内核 0 会在其他内核完成之前执行最终归约。

还有其他人对如何调试此问题有任何建议或提示吗？

提前致谢！！

Answer 1

问题是内核启动时每个工作组只有一个线程，而障碍仅在工作组内起作用。请参阅对类似问题的回复：Open CL no synchronization despite barrier

并行缩减算法中的 OpenCL 未定义行为

OpenCL undefined behavior in parallel reduction algorithm

opencl