从另一个OpenCL内核调用OpenCL内核

Question

我在这里看到一篇帖子，说我们可以从OpenCL内核中调用函数。但是在我的情况下，我需要将复杂的函数并行化（由所有可用线程运行），那么我是否也需要将该函数作为内核并像主内核函数一样直接调用它？或者有什么其他可能的解决方案？提前感谢。

- Akhtar Ali

Answer 1

你可以从内核中调用助手函数，它们将以与内核相同的方式并行化，可以将其想象为内联在内核代码中。因此，每个工作项将为其处理的工作集调用助手函数。

float4 helper_function(float4 input)
{
   return input.x + input.y + input.z + input.w;
}
__kernel kernel_function(const float4* arr, float4* out)
{
  id = get_global_id(0);
  out[id] = helper_function(arr[id]);
}

- sramij

---

补充sramij的回答，从内核本身调用另一个内核被称为动态并行性。为此，您需要支持OpenCL 2.0的设备。可以参考https://dev59.com/XmnWa4cB1Zd3GeqP2qj2。 - Meluha

Answer 2

OpenCL 2.0规范 新增了一项动态并行处理的新功能。

6.13.17 Enqueuing Kernels 
OpenCL 2.0 allows a kernel to independently enqueue to the same device, without host 
interaction. ...

在下面的示例中，my_func_B 在设备上将 my_func_A 排队：

kernel void
my_func_A(global int *a, global int *b, global int *c)
{
 ...
}

kernel void
my_func_B(global int *a, global int *b, global int *c)
{
 ndrange_t ndrange;
 // build ndrange information
 ...
 // example – enqueue a kernel as a block
 enqueue_kernel(get_default_queue(), ndrange, ^{my_func_A(a, b, c);});
 ...
}

- zr.

Answer 3

如果我正确理解您的问题，您希望从内核中独立地对缓冲区进行完整传递。我认为这在内核内部是不可能的，因此您需要将“内部”传递的代码创建为一个单独的内核，并且还要从主机代码分别调用该内核。该内核的输出不必读回到主机内存中，而可以在设备内存中在您的内核调用之间保留。

- Emil Styrke