c - OpenCL、C - Pi 的莱布尼兹公式

Question

我正在尝试获得一些OpenCL的经验，环境已经设置好，我可以创建和执行内核。我目前正在尝试使用莱布尼茨公式并行计算 pi，但收到了一些奇怪的结果。

内核如下:

__kernel void leibniz_cl(__global float *space, __global float *result, int chunk_size) 
{
    __local float pi[THREADS_PER_WORKGROUP];
    pi[get_local_id(0)] = 0.;

    for (int i = 0; i < chunk_size; i += THREADS_PER_WORKGROUP) {
        // `idx` is the work item's `i` in the grander scheme
        int idx = (get_group_id(0) * chunk_size) + get_local_id(0) + i;
        float idx_f = 1 / ((2 * (float) idx) + 1);

        // Make the fraction negative if needed
        if(idx & 1)
            idx_f = -idx_f;

        pi[get_local_id(0)] += idx_f;
    }

    // Reduction within workgroups (in `pi[]`)
    for(int groupsize = THREADS_PER_WORKGROUP / 2; groupsize > 0; groupsize >>= 1) { 
        if (get_local_id(0) < groupsize) 
            pi[get_local_id(0)] += pi[get_local_id(0) + groupsize];

        barrier(CLK_LOCAL_MEM_FENCE);
    }

如果我在此处结束该函数，并将 result 设置为 !get_global_id(0) 的 pi[get_local_id(0)] (如第一组的减少)，打印 result 打印 -nan。

内核的其余部分:

    // Reduction amongst workgroups (into `space[]`)
    if(!get_local_id(0)) {
        space[get_group_id(0)] = pi[get_local_id(0)];

        for(int groupsize = get_num_groups(0) / 2; groupsize > 0; groupsize >>= 1) { 
            if(get_group_id(0) < groupsize)
                space[get_group_id(0)] += space[get_group_id(0) + groupsize];

            barrier(CLK_LOCAL_MEM_FENCE);
        }
    }

    barrier(CLK_LOCAL_MEM_FENCE);

    if(get_global_id(0) == 0)
        *result = space[get_group_id(0)] * 4;
}

返回 space[get_group_id(0)] * 4 返回 -nan 或一个非常大的数字，显然不是 pi 的近似值。

我无法确定这是我缺少的 OpenCL 概念还是一般的并行执行概念。如有任何帮助，我们将不胜感激。

链接

缩减模板:OpenCL float sum reduction

莱布尼茨公式:https://www.wikiwand.com/en/Leibniz_formula_for_%CF%80

Answer 1

也许这些不是代码最关键的问题，但它们可能是问题的根源:

1. 在局部还原之前，您绝对应该使用barrier(CLK_LOCAL_MEM_FENCE);。如果您知道工作组大小等于或小于并行运行相同指令的波前线程数，则可以避免这种情况 - AMD GPU 为 64，NVidia GPU 为 32。
2. 全局缩减必须在多次启动内核时完成，因为 barrier() 仅适用于同一工作组的工作项。将屏障插入内核的清晰且 100% 有效的方法是将其在需要全局屏障的地方分成两部分。