我正在利用OpenCL的enqueue_kernel()函数从GPU动态地enqueue内核，以减少不必要的主机交互。下面是一个简化的例子，说明我在内核方面的尝试。
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()会无法enqueue而停止程序。发生这种情况是因为设备队列已经满了，它无法将另一个任务放入其中。但是，即使我做了大量的研究，我也无法理解为什么会这样。
我以为一旦队列中的一个任务（例如内核）完成，就会释放队列中的那个位置。所以我的队列一次最多应该只能达到1到2个任务的样子。但是这个程序会把整个262,144字节大小的设备命令队列填满，然后停止运行。
如果有人有任何想法，我将非常感谢一些潜在的洞察力，为什么会发生这种情况。我被卡住了，在解决这个问题之前，我无法继续工作。
先谢谢你
顺便说一下，我在Radeon RX 590卡上运行，并使用AMD APP SDK 3.0与OpenCL 2.0一起使用。
我不知道到底是哪里出了问题，但我注意到你发布的代码中的一些东西，这些反馈在评论中读起来太过冗长，所以这里就说说吧--这不是一个明确的答案，但却是一个接近的尝试。
在 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.
这可能是你的代码缩减版中的一个排版错误。
这可能又是由于你删减代码的方式造成的，但我不太明白为什么
所以我的队列应该真的是一次最多只能达到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 met)产生1个任务，至少有256个以上的工作项目，255个以上的任务有1个工作项目运行 kernelA. kernelD 行为类似。
因此，通过几次迭代，你可以很容易地最终获得几千个任务，用于运行 kernelA 排队。我真的不知道你的代码是怎么做的，但似乎是个好主意，检查一下是否减少了这数百个 kernelA 任务改善情况，以及是否可以或许修改 kernelA 这样你只需用一个范围来enqueue它一次，而不是从每个工作项目中enqueue一个1的工作大小。(或者类似于这样的思路--如果这样做更合理的话，也许每个组都可以enqueue一次。基本上，减少 enqueue_kernel 被调用。)
enqueue_kernel() 返回值你有没有实际检查过 enqueue_kernel? 它能准确地告诉你失败的原因，所以即使我上面的建议不可行，也许你可以设置一些全局状态，让 kernelA 如果中断了更多的任务，在更多任务耗尽后重新开始计算？