7

我有一个 OpenCL 内核，它需要将一个数组处理为多个数组，其中每个子数组和保存在本地缓存数组中。

例如，想象一下 fowling 数组：

[[1, 2, 3, 4], [10, 30, 1, 23]]
  • 每个工作组都有一个数组（在示例中我们有 2 个工作组）；

  • 每个工作项处理两个数组索引（例如将值索引乘以 local_id），其中工作项结果保存在工作组共享数组中。

    __kernel void test(__global int **values, __global int *result, const int array_size){
    __local int cache[array_size];

    // initialise
    if (get_local_id(0) == 0){
        for (int i = 0; i < array_size; i++)
            cache[i] = 0;
    }

    barrier (CLK_LOCAL_MEM_FENCE);

    if(get_global_id(0) < 4){
        for (int i = 0; i<2; i++)
            cache[get_local_id(0)] += values[get_group_id(0)][i] * 
                                                         get_local_id(0);
    }

    barrier (CLK_LOCAL_MEM_FENCE);

    if(get_local_id(0) == 0){
        for (int i = 0; i<array_size; i++)
            result[get_group_id(0)] += cache[i];
    }
}

问题是我无法通过使用内核参数来定义缓存数组大小，但我需要为了拥有一个动态内核。

如何动态创建它？就像 c 中的 malloc 函数...

或者唯一可用的解决方案是将临时数组发送到我的内核函数？

4

1 回答 1

18

这可以通过添加__local数组作为内核参数来实现：

__kernel void test(__global int **values, __global int *result, 
    const int array_size, __local int * cache)

并提供所需的内核参数大小：

clSetKernelArg(kernel, 3, array_size*sizeof(int), NULL);

本地内存将在内核调用时分配。请注意，可能需要进行额外检查以确保所需的本地内存大小不超过设备限制。

于 2013-07-10T16:40:02.730 回答