OpenCL 2.0 设备命令队列不断填​​满并停止执行

Question

我正在利用 OpenCL 的 enqueue_kernel() 函数从 GPU 中动态地对内核进行排队，以减少不必要的主机交互。这是我在内核中尝试做的事情的简化示例:

kernel void kernelA(args)
{
    //This kernel is the one that is enqueued from the host, with only one work item. This kernel
    //could be considered the "master" kernel that controls the logic of when to enqueue tasks
    //First, it checks if a condition is met, then it enqueues kernelB

    if (some condition)
    {
        enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(some amount, 256), ^{kernelB(args);});
    }
    else
    {
        //do other things
    }
}

kernel void kernelB(args)
{
    //Do some stuff

    //Only enqueue the next kernel with the first work item. I do this because the things
    //occurring in kernelC rely on the things that kernelB does, so it must take place after kernelB is completed,
    //hence, the CLK_ENQUEUE_FLAGS_WAIT_KERNEL
    if (get_global_id(0) == 0)
    {
        enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(some amount, 256), ^{kernelC(args);});
    }
}

kernel void kernelC(args)
{
    //Do some stuff. This one in particular is one step in a sorting algorithm

    //This kernel will enqueue kernelD if a condition is met, otherwise it will
    //return to kernelA 
    if (get_global_id(0) == 0 && other requirements)
    {
        enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(1, 1), ^{kernelD(args);});
    }
    else if (get_global_id(0) == 0)
    {
        enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(1, 1), ^{kernelA(args);});
    }
}

kernel void kernelD(args)
{
    //Do some stuff

    //Finally, if some condition is met, enqueue kernelC again. What this will do is it will
    //bounce back and forth between kernelC and kernelD until the condition is
    //no longer met. If it isn't met, go back to kernelA
    if (some condition)
    {
        enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(some amount, 256), ^{kernelC(args);});
    }
    else
    {
        enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(1, 1), ^{kernelA(args);});
    }
}

这就是程序的一般流程，除了一个问题外，它工作得很好，完全符合我的预期，按照我预期的顺序执行。在某些情况下，当工作负载非常高时，随机的一个 enqueue_kernel() 将无法入队并停止程序。发生这种情况是因为设备队列已满，无法将另一个任务放入其中。但即使经过广泛的研究，我终究无法弄清楚这是为什么。

我认为一旦队列中的任务(例如内核)完成，它就会释放队列中的那个位置。所以我的队列一次最多只能处理 1 到 2 个任务。但是这个程序会完全填满设备命令队列的整个 262,144 字节大小，然后停止运行。

如果有人有任何想法，我将不胜感激一些关于为什么会发生这种情况的潜在见解。我有点卡住了，在解决这个问题之前无法继续。

提前致谢!

(顺便说一句，我在 Radeon RX 590 卡上运行，并使用 AMD APP SDK 3.0 与 OpenCL 2.0 一起使用)

Answer 1

我不知道到底出了什么问题，但我注意到您发布的代码中有一些问题，并且此反馈太长/难以在评论中阅读，所以这里是 - 不是一个明确的答案，但是尝试靠近一点:

代码并不完全按照评论所说的去做

在 kernelD 中，您有:

//Finally, if some condition is met, enqueue kernelC again.

…

if (get_global_id(0) == 0)
{
    enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(some amount, 256), ^{kernelD(args);});
}

这实际上使 kernelD 本身再次入队，而不是像评论所暗示的那样 kernelC 。另一个条件分支将 kernelA 入队。

这可能是您代码的简化版本中的拼写错误。

潜在的任务爆炸

这可能再次归结为您删节代码的方式，但我不太明白如何

So my queue should really only reach a max of like 1 or 2 tasks at a time.

可以是真的。根据我的阅读，kernelC 和 kernelD 的所有工作项都会产生新任务；由于在每种情况下似乎都有不止 1 个工作项，这似乎很容易产生大量任务:

例如，在kernelC中:

if (get_global_id(0) == 0 && other requirements)
{
    enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(some amount, 256), ^{kernelD(args);});
}
else
{
    enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(1, 1), ^{kernelA(args);});
}

kernelB 将创建至少 256 个运行 kernelC 的工作项。在这里，工作项 0 将(如果 其他要求 满足)产生 1 个任务和至少 256 个以上的工作项，以及 255 个以上的任务和 1 个运行 kernelA 的工作项。 kernelD 的行为类似。

因此，通过几次迭代，您可以很容易地得到几千个用于运行 kernelA 的任务排队。我真的不知道你的代码做了什么，但检查一下减少这数百个 kernelA 任务是否可以改善情况，以及你是否可以修改 kernelA< 似乎是个好主意 以便您只需使用一个范围将其入队一次，而不是从每个工作项中入队 1 的工作大小。 (或者类似的东西——如果这更有意义的话，也许每组排队一次。基本上，减少 enqueue_kernel 被调用的次数。)

enqueue_kernel()返回值

你真的检查过 enqueue_kernel 的返回值了吗？它准确地告诉您失败的原因，所以即使我上面的建议不可行，也许您可​​以设置一些全局状态，这将允许 kernelA 在更多任务耗尽后重新启动计算，如果它是打扰了？