在循环中排队到设备端队列答案

【问题标题】：Enqueuing to device side queue in a loop在循环中排队到设备端队列
【发布时间】：2020-03-05 21:52:59
【问题描述】：

在我的代码中，我有 kernelA 和 kernelB。 kernelB 取决于 kernelA 结果。我对这个内核进行了无数次迭代，每次迭代都取决于上一次迭代的结果。

截取的主机端入队代码是这样的：

for(int x = 0; x < iterations; ++x)
{
    queue.enqueueNDRangeKernel(kernelA, cl::NullRange, cl::NDRange(3*256, 1), cl::NDRange(256, 1));
    queue.enqueueNDRangeKernel(kernelB, cl::NullRange, cl::NDRange(256, 1), cl::NDRange(256, 1));
}
queue.finish();

上面的代码运行良好。

现在我想移植上述代码以使用设备端排队，但我在 AMD GPU 上遇到了问题。内核代码：

__attribute__((reqd_work_group_size(256, 1, 1)))
__kernel void kernelA(...){}

__attribute__((reqd_work_group_size(256, 1, 1)))
__kernel void kernelB(...){}

__attribute__((reqd_work_group_size(1, 1, 1)))
__kernel void kernelLauncher(...)
{
    queue_t default_queue = get_default_queue();
    clk_event_t ev1, ev2;

    for (int x = 0; x < iterations; ++x)
    {
        void(^fnKernelA)(void) = ^{ kernelA(
        ... // kernel params come here
            ); };

        if (x == 0)
        {
            enqueue_kernel(default_queue,
                CLK_ENQUEUE_FLAGS_NO_WAIT,
                ndrange_1D(3 * 256, 256),
                0, NULL, &ev1,
                fnKernelA);
        }
        else
        {
            enqueue_kernel(default_queue,
                CLK_ENQUEUE_FLAGS_NO_WAIT,
                ndrange_1D(3 * 256, 256),
                1, &ev2, &ev1, // ev2 sets dependency on kernelB here
                fnKernelA);
        }

        void(^fnKernelB)(void) = ^{ kernelB(
        ... // kernel params come here
            ); };

        enqueue_kernel(default_queue,
            CLK_ENQUEUE_FLAGS_NO_WAIT,
            ndrange_1D(256, 256),
            1, &ev1, &ev2,  // ev1 sets dependency on kernelA here
            fnKernelB);
    }
}

主机代码：

queue.enqueueNDRangeKernel(kernelLauncher, cl::NullRange, cl::NDRange(1, 1), cl::NDRange(1, 1));

问题是在 AMD GPU 上运行时从内核返回的结果是错误的。有时内核也会挂起，这可能表明内核同步可能有问题。相同的代码在 Intel CPU 上运行良好，不确定是运气还是内核中的同步点有问题。

更新： enqueue_kernel 在第 1025 个入队命令上失败，错误为 -1。我试图获得更详细的错误（在构建过程中添加了-g）但无济于事。我将设备队列大小增加到最大值，但这并没有改变任何东西（在第 1025 个入队命令上仍然失败）。删除 kernelA 和 kernelB 的内容也没有改变任何东西。有什么想法吗？

【问题讨论】：

在您的第一个代码 sn-p（从主机排队）中，kernelA 的工作组大小设置为 3*256。在第二个片段中（从设备排队），kernelA 的工作组大小设置为256?这些应该是一样的吗？
@jprice kernelA 和 kernelB 两个版本完全相同。 kernelA 在这两种情况下都作为 256 个工作项的 3 个工作组启动。请注意，内核端入队设置为ndrange_1D(3 * 256, 256)。
好的，但是问题顶部的主机队列代码会为kernelA 启动一个3*256 工作项工作组，而不是3 工作组@ 987654338@ 工作项。
@jprice，是的，你说得对，这只是我在准备剪报时打错了字，对此感到抱歉。
很公平。看不到代码有任何其他问题。如果只运行一次迭代，它会产生正确的结果吗？您是否尝试过检查 enqueue_kernel 函数是否总是返回 CLK_SUCCESS？

标签： opencl

【解决方案1】：

回答一个老问题，希望能在未来节省一些时间。如果您在设备上查询CL_DEVICE_MAX_ON_DEVICE_EVENTS，它将返回 1024。这是您可以“在设备上”排队的最大事件数。这就是它在 1025 队列上失败的原因。如果您在不同的 GPU（如 Intel）上运行 OpenCL 代码，您可能会很幸运地得到一个真正的错误代码，即 CLK_DEVICE_QUEUE_FULL 或 -161。 AMD 忽略了-g 选项，并且似乎从来没有在设备上排队失败时返回-1。

【讨论】：