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

Question

我正在研究一种简单的并行归约算法来查找数组中的最小值，并且在我的算法中遇到了一些有趣的未定义行为。我在 Ubuntu 16.04 上运行 Intel 的 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;
    }
}

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

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

提前致谢！！

Answer 1

问题在于每个工作组使用一个线程启动内核，并且屏障仅在工作组内有效。请参阅对类似问题的回复：Open CL no synchronization despite barrier