0

我正在尝试一些 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];
}

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

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

4

2 回答 2

1

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

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

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

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

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

于 2020-08-03T19:56:04.713 回答
0

您可以使用带有传入选项“-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”时。

于 2021-06-10T08:53:05.177 回答