OpenCL：如何正确链接内核

Question

好的，所以我有两个内核，它们都接受输入和输出图像并执行一些有意义的操作：

#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable

kernel void Kernel1(read_only image3d_t input, write_only output)
{
    //read voxel and some surrounding voxels
    //perform some operation
    //write voxel
}


#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable

kernel void Kernel2(read_only image3d_t input, write_only output)
{
    //read voxel and some surrounding voxels
    //perform some other operation
    //write voxel
}


#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable

kernel void KernelCombined(read_only image3d_t input, write_only output)
{
    //read voxel and some surrounding voxels

    //...
    //perform operation of both kernels (without read, write)
    //...
    //write voxel
}

现在我想在某些情况下链接内核，所以我可以做的是先调用内核 1，然后调用内核 2。但这意味着，我在两者之间进行了不必要的写入和读取。我也可以编写第三个内核来完成这两项工作，但维护复制粘贴代码似乎很烦人。据我所知，我无法真正将每个内核的内容放在一个单独的函数中，因为我无法传递 image3d_t 输入。

问题：有什么巧妙的方法可以链接这两个内核吗？也许 OpenCL 已经在做一些我不知道的聪明事了？

编辑： 添加了我想要实现的示例。

Answer 1

如果您使用的是支持 opencl 2.0 的设备，这是可能的。 enqueue_kernel 允许内核对另一个内核进行排队，就像主机上的 EnqueueNDRange。

如果您使用的是 opencl 1.2 —— 可能所有 1.x，您需要 return 到主机并调用下一个内核（或者下一个内核已经排队）。不过，您不需要在内核之间将缓冲区复制回主机，因此至少您无需为多次传输付费。

Answer 2

据我从你的描述中了解到，你不应该做任何特别的事情，即使使用 OpenCL 1.2 也能正常工作。

默认情况下，OpenCL 命令队列是有序的，不需要在内核调用之间传输数据。

只需将数据留在设备上（不要做 map/unmap 和 Read/Write），将两个内核加入队列并等待它们完成。这是它的外观的代码片段：

// Enqueue first kernel
clSetKernelArg(kernel1, 0, sizeof(cl_mem), in);
clSetKernelArg(kernel1, 1, sizeof(cl_mem), out);
clEnqueueNDRange(..., kernel1, ...);

// Enqueue second kernel
clSetKernelArg(kernel2, 0, sizeof(cl_mem), in);
clSetKernelArg(kernel2, 1, sizeof(cl_mem), out);
clEnqueueNDRange(..., kernel2, ...);

// Flush the queue and wait for the results
clFlush(...);  // Start the execution
clWait(...);   // Wait until all operations in the queue are done

当使用 OOO（乱序）队列时，可以使用事件（参见 clEnqueueNDRangeKernel 中的最后 3 个参数）来指定内核之间的依赖关系，并在管道的末尾执行 clWaitForEvents .

Answer 3

我明白您的要求 -- 您希望删除内核之间的图像写入/读取周期。使用您描述的内核，这不会有效。在现有的内核中，你 "read voxel and some surrounding voxels" —— 假设这意味着读取 7 个体素。如果您在内核 2 和内核 3 中执行相同的读取模式，则共有 21 次读取（和 3 次写入）。如果您以某种方式将这三个内核链接到一个写入单个输出体素的内核中，它将需要从更多的源体素中读取以获得相同的结果（因为每个读取步骤都会增加半径）。

内核 write/read 链接会有所帮助的场景将适用于 single-in/single-out 内核，例如图像处理，其中颜色独立于其邻居进行修改。为此，您需要对内核进行更高级别的描述，以及可以根据您的操作生成所需内核的内容。

OpenCL：如何正确链接内核

OpenCL: How to properly chain kernels

opencl