OpenCL 在运行时从 Host 函数创建内核答案

【问题标题】：OpenCL creating kernel from Host function at runtimeOpenCL 在运行时从 Host 函数创建内核
【发布时间】：2020-11-22 15:24:09
【问题描述】：

我正在尝试一些 OpenCL，想知道是否有办法将函数作为参数传递给内核，或者最接近的可用匹配项是什么（使用 OpenCL 1.2）。

例如，考虑这样一个简单的蒙特卡洛积分：

/* this is 1/(2^32) */
#define MULTI (2.3283064365386962890625e-10)

/* for more information see: https://arxiv.org/pdf/2004.06278v2.pdf*/
uint
squares(ulong ctr, ulong key)
{
  ulong x, y, z;
  y = x = ctr * key;
  z = y + key;
  x = x * x + y;
  x = (x >> 32) | (x << 32);                /* round 1 */
  x = x * x + z; x = (x >> 32) | (x << 32); /* round 2 */
  return (x * x + y) >> 32;                 /* round 3 */
}

void
kernel
reduce(ulong  key,
       float  low,
       float  high,
       global float* partialSums,
       local  float* localSums)
{
  uint lid = get_local_id(0);

  float rand = squares(get_global_id(0), key) * MULTI;
  localSums[lid] = f((rand * (high - low)) + low);

  for (uint stride =  get_local_size(0) / 2; stride > 0; stride /= 2) {
    barrier(CLK_LOCAL_MEM_FENCE);

    if (lid < stride)
      localSums[lid] += localSums[lid + stride];
  }

  if (lid == 0)
    partialSums[get_group_id(0)] = localSums[0];
}

我发现Passing a function as an argument in OpenCL 告诉我传递函数指针不起作用。所以我猜想在运行时生成带有 f 定义的内核源然后编译它会起作用（以前是否做过？如果是这样，我在哪里可以找到它？）。也许这种问题不使用 OpenCL 而是使用 SYCL（我几乎一无所知）更容易解决？

我对此比较陌生，所以如果这种问题以完全不同的方式解决，请告诉我。

【问题讨论】：

标签： opencl opencl-c sycl

【解决方案1】：

使用在运行时定义的 f 生成内核源代码，然后编译它

是的，它可以做到。您可以从头开始创建整个源代码，然后使用经典的 clCreateProgram + clBuildProgram。

另一种选择是将您的程序拆分为静态和动态生成的部分，然后在运行时通过 clCompileProgram（静态部分仅一次）分别编译它们，然后将它们与 clLinkProgram 链接。这可能会更快一些。

也许这种问题不使用 OpenCL 而使用 SYCL 更容易解决

实际上用 SYCL 可能更难解决；我不确定 SYCL 是否完全支持动态（运行时）编译。

【讨论】：

SYCL 会使这变得更难是不正确的。虽然 SYCL（注意 SYCL，而不是 SyCL——CL 不是来自 OpenCL）没有运行时编译，除非使用 OpenCL 互操作性，但它是 C++ 单源编程模型，因此您可以拥有模板化内核。例如，您可以使用 lambda 函数来实例化您的内核，这可能就是这里所需要的。所以我会鼓励在这样的用例中使用 SYCL。
@illuhad 感谢您提供的信息。模板很好，但它们不能替代运行时编译——它们在编译时被实例化。所以确实如你所说，这取决于这里到底需要什么。
当然，它们在一定程度上涵盖了不同的用例。但是 OP 本身并没有专门要求运行时编译，而是更普遍地要求一种“将函数作为参数传递”的方法，这在我看来像是在编译时可能可行的东西。 IMO，在绝大多数情况下，与运行时编译相比，使用模板和 lambda 函数自定义算法绝对是传递函数作为参数的首选方式，因为它具有简单性、类型安全性和鲁棒性。
我编辑了我的问题以使用 SYCL 而不是 SyCL :)。如果至少知道函数的类型（即多项式或其他），那么我会看到使用编译时机制来实现它是多么有意义。但是，如果我不知道该函数将具有哪种形状，那么除了运行时编译之外别无选择，不是吗？尽管如此，我还是会研究 SYCL (hipSYCL? ;)，因为它看起来很舒服。
@Malte 我不确定你对函数的“形状”是什么意思。您的功能需要一个固定的签名。使用 lambda 函数，您还可以为函数捕获其他参数。目前我不清楚为什么这对任意多项式不起作用。也许你可以更详细地描述你需要什么？如果您使用 OpenCL 进行运行时编译，您可能需要查看 boost.compute。 AFAIK 他们有一些方便的包装器可以在运行时注入代码。当然，我也很乐意帮助您解决 SYCL（hipSYCL 与否；））问题。

【解决方案2】：

您可以使用带有传入选项“-create-library”的 clCreateProgram + clLinkProgram 创建一个函数“f”的 OpenCL 库。

按照内核的这种方法，您应该传递额外的整数参数 f_idx，编码要调用的 'f' 的实际实例，并在内核主体中而不是实际的 'f' 调用中执行 f_dispatch(f_idx, f_params)。其中 f_dispatch 将是在内核附近定义的函数，并将 f_idx 值“表转换”为由 f_idx 编码的某些“f(f_params)”的实际调用。

这是处理任务的经典 C 方法，虽然 OpenCL C 是某种 C99，不允许使用函数指针，但它似乎是处理您的任务的合理方法。

其他更复杂的方法是生成尽可能多的内核，因为您有各种“f”函数，并将“调度”逻辑移动到主机端，当您选择要排队的内核以调用某些“f”时。

【讨论】：