奇怪的 OpenCL 行为

Question

我遇到了非常奇怪的 OpenCL 行为。我已经链接了一个最小的代码示例。

如果我事先添加一个额外的操作 (g_idata[ai] = g_idata[ai-1])，则从某个随机索引（通常为 32 整除）开始的值不会写入数组。还值得注意的是，如果：

，我将得到正确的结果

只需读取值，然后写入文字（参见 SHOW_BUG）。
在开头添加if (ai >= n) g_idata[0]+=0;。查看注释行

已在 Intel 和 nvidia 上测试。


import numpy as np
import pyopencl as cl


ctx = cl.create_some_context()

prg = cl.Program(ctx, """
__kernel void prescan(__global float *g_idata, const int n) {
    int thid = get_global_id(0);
    
    int ai = thid*2+1;
    
    // if uncomment strings bellow the bug dissappears
    //if (ai >= n){
    //    g_idata[0]+=0;
    //}
    
    bool SHOW_BUG=1;
    // make a dummy operation
    if (SHOW_BUG)
        g_idata[ai] = g_idata[ai-1];
    else {
        g_idata[ai-1]; //dummy read
        g_idata[ai] = 3.14f; //constant write
    }
    barrier(CLK_GLOBAL_MEM_FENCE);
   
    //set 0,1,2,3... as result
    g_idata[thid] = thid;
}

""").build()

prescan_kernel = prg.prescan
prescan_kernel.set_scalar_arg_dtypes([None, np.int32])


def main():
    N = 512
    a_np = (np.random.random((N,))).astype(np.float32)
    queue = cl.CommandQueue(ctx)

    mf = cl.mem_flags
    a_g = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=a_np)

    global_size = (512,)
    local_size = None
    prescan_kernel(queue, global_size, local_size, a_g, N)
    cl.enqueue_copy(queue, a_np, a_g)

    corect = np.array(range(N))
    #assert np.allclose(a_np, 3.14), np.where(3.14 != a_np)
    assert np.allclose(a_np, corect), np.where(corect != a_np)


if __name__ == '__main__':
    for i in range(25):
        main()

Answer 1

根据 OpenCL 规范，您的代码中有几项会产生未定义的行为。

其中包括：

正在访问超出范围的内存。 N 个工作项的数组大小预计为 N*2+1。
多个工作项（线程）访问数组的相同索引（读取或写入）。

此外，障碍仅在工作组中同步 work-items/threads，因此它对您的代码没有影响。在讨论未定义的行为时，它可能在不同的平台上表现不同，有时会导致驱动程序崩溃，有时会导致 OS 崩溃。请解决这些问题，然后描述您的问题。

奇怪的 OpenCL 行为

strange OpenCL behavior

gpgpu

opencl

pyopencl