在OpenCL内核中创建本地数组动态

Question

我有一个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函数...

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

Answer 1

这可以通过添加__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);

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