OpenCL 在运行时从主机函数创建内核

Question

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

举个例子，考虑这样一个简单的 Monte Carlo 集成：

/* 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（我几乎一无所知）更容易解决？

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

Answer 1

generating the kernel source with f defined at runtime and then compiling it

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

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

Maybe this kind of problem is easier to solve not using OpenCL but using SYCL

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

Answer 2

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

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

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

其他更复杂的方法是生成与各种 'f' 函数一样多的内核，并将 'dispatch' 逻辑移动到主机端，当您选择要排队的内核时调用某些 'f'.

OpenCL 在运行时从主机函数创建内核

OpenCL creating kernel from Host function at runtime

opencl

opencl-c

sycl