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

我了解您的要求——您希望删除内核之间的映像写入/读取周期。使用您描述的内核，这将是无效的。在现有内核中，您“读取体素和一些周围的体素”——假设这意味着读取 7 个体素。如果您在内核 2 和 3 中执行相同的读取模式，则总共有 21 次读取（和 3 次写入）。如果您以某种方式将这三个内核链接到一个编写单个输出体素的内核中，则需要从更多源体素中读取才能获得相同的结果（因为每个读取步骤都在增加半径）。

内核写入/读取链接会有所帮助的场景是单进/单出内核，例如独立于其邻居修改颜色的图像处理。为此，您需要对内核进行更高级别的描述，以及可以根据您拥有的操作生成所需内核的内容。

Answer 2

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

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

Answer 3

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

OpenCL 命令队列默认为 IN ORDER，无需在内核调用之间传输数据。

只需将数据留在设备上（不要执行映射/取消映射和读/写操作），将两个内核排入队列并等待它们完成。这是一段代码 sn-p 的外观：

// 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。