3

我正在利用 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 一起使用）

4

1 回答 1

2

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

代码并没有完全按照评论所说的那样做

在kernelD中，您有：

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

…</p> 

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。

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

潜在的任务爆炸

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

所以我的队列实际上应该一次最多只能达到 1 或 2 个任务。

可以是真的。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 将（如果other requirements满足）产生 1 个任务，其中至少还有 256 个工作项，以及 255 个以上的任务，其中 1 个工作项正在运行kernelA。kernelD行为相似。

因此，通过几次迭代，您可以轻松地完成数千个任务以运行kernelA队列。我真的不知道你的代码做了什么，但检查减少这数百个任务是否能改善情况似乎是个好主意kernelA，以及你是否可以修改kernelA，以便你只用一个范围将它排入队列而不是入队每个工作项的工作大小为 1。（或者类似的东西——如果这更有意义的话，可能每组排队一次。基本上，减少enqueue_kernel被调用的次数。）

enqueue_kernel()返回值

你真的检查过返回值enqueue_kernel吗？它会告诉你它失败的确切原因，所以即使我上面的建议是不可能的，也许你可以设置一些全局状态，kernelA一旦更多任务耗尽，如果它被中断，它允许重新启动计算？

于 2020-06-02T15:40:51.267 回答