1

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。是否可以以某种方式初始化全局变量？我知道我可以把它写成代码，但是还有其他方法吗？

PS 我正在使用 AMD OpenCL

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

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

4

3 回答 3

4 
#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 上，我得到 {102, 103} 与 OCL 实现（AMD，英特尔）。所以我不能强调，我个人不会这样做，因为这只是时间问题。

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

于 2013-07-02T21:51:32.143 回答
1

前段时间，我努力在自己的程序中完成这项工作。我没有找到任何方法通过一些 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); 之类的东西编译你的缓冲区
  • 瞧！
于 2013-07-03T17:48:53.967 回答
-1

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

于 2013-07-02T21:59:04.400 回答