如何在 OpenCL 的程序（全局）范围内使用数组

Question

AMD OpenCL 编程指南，第 6.3 节常量内存优化：

全局范围的常量数组。这些数组被初始化， 全局作用域，并且在常量地址空间中（如在 OpenCL 规范的第 6.5.3 节）。如果数组的大小是 低于 64 kB，它被放置在硬件常量缓冲区中；否则，它 使用全局内存。这方面的一个例子是数学查找表 功能。

我想使用这个“全局范围的常量数组”。我在纯 C 中有这样的代码

#define SIZE 101
int *reciprocal_table;

int reciprocal(int number){
  return reciprocal_table[number];
}

void kernel(int *output)
{
  for(int i=0; i < SIZE; i+)
    output[i] = reciprocal(i);
}

我想把它移植到 OpenCL 中

__kernel void kernel(__global int *output){
  int gid = get_global_id(0);

  output[gid] = reciprocal(gid);
}

int reciprocal(int number){
  return reciprocal_table[number];
}

我应该如何处理全局变量reciprocal_table？如果我尝试在其中添加 __global 或 __constant，则会收到错误消息：

global variable must be declared in addrSpace constant

我不想将 __constant int *reciprocal_table 从 kernel 传递到 reciprocal。是否可以以某种方式初始化全局变量？我知道我可以把它写成代码，但是还有其他方式吗？

附：我正在使用 AMD OpenCL

UPD 上面的代码只是一个例子。我有很多功能更复杂的代码。所以我想在程序范围内创建数组以在所有函数中使用它。

UPD2更改了示例代码并添加了来自编程指南的引用

Answer 1

#define SIZE 2
int constant array[SIZE] = {0, 1};

kernel void
foo (global int* input,
     global int* output)
{
    const uint id = get_global_id (0);
    output[id] = input[id] + array[id];
}

我可以使用 Intel 和 AMD 编译上述内容。它也可以在没有初始化数组的情况下工作，但是您将不知道数组中的内容，并且由于它位于常量地址空间中，因此您无法分配任何值。

程序全局变量必须在 __constant 地址空间中，如标准中第 6.5.3 节所述。

更新现在，我完全理解了这个问题：

对我有用的一件事是在常量空间中定义数组，然后通过传递覆盖数组的内核参数constant int* array 来覆盖它。 这仅在 GPU 设备上产生了正确的结果。 AMD CPU 设备和 Intel CPU 设备没有覆盖数组地址。它也可能不符合标准。

它的外观如下：

#define SIZE 2
int constant foo[SIZE] = {100, 100};

int
baz (int i)
{
  return foo[i];
}

kernel void
bar (global int* input,
     global int* output,
     constant int* foo)
{
    const uint id = get_global_id (0);
    output[id] = input[id] + baz (id);
}

对于 input = {2, 3} 和 foo = {0, 1}，这会在我的 HD 7850 设备（Ubuntu 12.10，Catalyst 9.0.2）上生成 {2, 4}。 但是在 CPU 上，我通过任一 OCL 实现（AMD、英特尔）得到 {102, 103}。所以我不能强调，我个人不会这样做，因为这只是时间问题。

另一种实现方式是在运行时使用主机计算 .h 文件，其中包含数组的定义（或预定义），并在编译时通过编译器将它们传递给内核选项。当然，这需要为每个不同的 LUT 重新编译 clProgram/clKernel。

Answer 2

前段时间，我努力在自己的程序中完成这项工作。 我没有找到任何方法通过一些 clEnqueueWriteBuffer 左右从主机初始化常量或全局范围数组。唯一的方法是在你的 .cl 源文件中明确地写出来。

所以在这里我从主机初始化它的技巧是利用您实际上是从主机编译源代码这一事实，这也意味着您可以在编译之前更改您的 src.cl 文件。

首先我的 src.cl 文件显示：

__constant double lookup[SIZE] = { LOOKUP };    // precomputed table (in constant memory).

double func(int idx) {
  return(lookup[idx])
}

__kernel void ker1(__global double *in, __global double *out)
{
   ... do something ...
   double t = func(i)
   ...
}

注意查找表是用 LOOKUP 初始化的。

然后，在宿主程序中，在编译您的 OpenCL 代码之前：

• 在 host_values[] 中计算我的查找表的值

• 在您的主机上，运行类似：

  char *buf = (char*) malloc( 10000 );
  int count = sprintf(buf, "#define LOOKUP ");    // actual source generation !
  for (int i=0;i<SIZE;i++) count += sprintf(buf+count, "%g, ",host_values[i]);
  count += sprintf(buf+count,"\n");
  

• 然后读取源文件 src.cl 的内容并将其放在 buf+count 处。

• 您现在有了一个源文件，其中包含您刚刚从主机计算的明确定义的查找表。
• 用类似 clCreateProgramWithSource(context, 1, (const char **) &buf, &src_sz, err); 之类的东西编译你的缓冲区
• 瞧！

Answer 3

看起来“数组”是一种查找表。您将需要 clCreateBuffer 和 clEnqueueWriteBuffer 以便 GPU 拥有它的副本以供使用。