奇怪的OpenCL行为

Question

我面临着OpenCL非常奇怪的行为。我链接了一个最小的代码样本。

从一些随机索引(通常为32可除)开始，如果预先添加一个额外的操作(g_idata[ai] = g_idata[ai-1])，则不会将值写入数组。同样值得注意的是，如果：

1. 只是读取值，并编写一个文字(参见SHOW_BUG)。
2. 在开头添加if (ai >= n) g_idata[0]+=0;。见注释行

在英特尔和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规范，代码中有几样东西会产生未定义的行为。

其中包括：

1. 访问超出范围的内存.数组大小预期为N个work-items.
2. Multiple工作项(线程)的N*2+1，它们访问数组的相同索引(读或写)。

此外，屏障只同步工作组中的工作项/线程，因此它在代码中没有任何影响。在讨论未定义的行为时，它可能在不同的平台上表现不同，有时会导致驱动程序崩溃，有时甚至会破坏操作系统。请解决这些问题，然后描述你的问题。