OpenCL嵌套循环产生意外结果

Question

我对OpenCL / C99有点陌生，但是不明白为什么下面的两个内核给出不同的结果。 X已被初始化为零，但在每个外循环步骤都需要“重新调零”，否则会得到错误的结果（请参见图）。注意，这里我没有调用任何并行性。据我所知，这完全是连续的：

kernels.cl的内容：

#pragma OPENCL EXTENSION cl_khr_fp64: enable
#define __CL_ENABLE_EXCEPTIONS

__kernel void nested_sum_succeeds(
    int L,
    __global __read_only double* X,
    __global double* Y
)
{
    for (int k=0; k<=L; k++) {
        Y[k] = 0;
        for (int j=0; j<=k; j++) {
            Y[k] += X[j];
        }
    }
}

__kernel void nested_sum_fails(
    int L,
    __global __read_only double* X,
    __global double* Y
)
{
    for (int k=0; k<=L; k++) {
        // Y[k] = 0;
        for (int j=0; j<=k; j++) {
            Y[k] += X[j];
        }
    }
}

script.py的内容：

import numpy as np
import pyopencl as cl
import pyopencl.array as cl_array
import matplotlib.pyplot as plt

with open("./kernels.cl") as fp:
    prog_str = fp.read()
ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx)
prog = cl.Program(ctx, prog_str).build()

L = 1000
X = np.linspace(0, 10, L)
X_dev = cl_array.to_device(queue, X)
Y_succeeds_dev = cl_array.to_device(queue, np.zeros(shape=X.shape, dtype=np.float64))
Y_fails_dev = cl_array.to_device(queue, np.zeros(shape=X.shape, dtype=np.float64))


nested_sum_succeeds = prog.nested_sum_succeeds
nested_sum_succeeds.set_scalar_arg_dtypes([
    np.int64,
    None,
    None,
])

nested_sum_succeeds(
    queue,
    (len(X),),
    None,
    L,
    X_dev.data,
    Y_succeeds_dev.data,
)


nested_sum_fails = prog.nested_sum_fails
nested_sum_fails.set_scalar_arg_dtypes([
    np.int64,
    None,
    None,
])

nested_sum_fails(
    queue,
    (len(X),),
    None,
    L,
    X_dev.data,
    Y_fails_dev.data,
)

np.allclose(Y_succeeds_dev.get(), Y_fails_dev.get()) #False

plt.ion()
plt.plot(Y_succeeds_dev.get())
plt.plot(Y_fails_dev.get())

结果：

Answer 1

注意，我这里没有调用任何并行性；

是，内核确实可以并行运行-已计划在第一维中运行len(x)工作项。一旦将其更改为使用1个工作项处理，就可以了：

nested_sum_succeeds(
    queue,
    (1,),
    None,
    L,
    X_dev.data,
    Y_succeeds_dev.data,
)

nested_sum_fails(
    queue,
    (1,),
    None,
    L,
    X_dev.data,
    Y_fails_dev.data,
)

然后np.allclose(Y_succeeds_dev.get(), Y_fails_dev.get())返回True。您也可以从Y[k] = 0;内核中删除该置零nested_sum_succeeds，因为它不需要。

此外，如果您想在其他设备上运行此内核，则需要进行一些较小的修复，因为并非所有编译器都会接受第一个内核参数的类型在内核int中，而在python中将其指定为[C0 ]，它必须与内核中的内容匹配，因此：

np.int64

还有适用于在其他设备上使用的另一件事，我将删​​除

nested_sum_succeeds.set_scalar_arg_dtypes([
    np.int32,
    None,
    None,
])

nested_sum_fails.set_scalar_arg_dtypes([
    np.int32,
    None,
    None,
])

访问限定符，它也不会在所有设备上编译。