如何将可变数量的不连续缓冲区传递给内核

How to pass a variable number of non-contiguous buffers to a kernel

本文关键字：缓冲区 不连续 内核      更新时间：2023-10-16

我在std::vector<cl::Buffer>中有一组相同大小的非连续 cl 缓冲区（非确定性长度）。我想在缓冲区中进行某种每像素的减少（在本例中为乘法）。我想出的唯一两种方法（似乎都有不必要的开销/完全不可读）

  1. 多次调用内核。
  2. 将所有缓冲区复制到连续内存。
  3. 以编程方式编写内核函数：

这样

int num_values = 4;
std::stringstream kernel;
kernel << 
"__global void pix_reduce_"<<num_values<<
"  (int out*,int* a1";
for (int i = 1; i < num_values; i++)
    kernel << ",int* a"<<i;
kernel << ",int numel)" <<
"{"<<
"  const int global_idx= get_global_id(0);"<<
"  if (global_idx<numel)"<<
"  {"<<
"    int out_val = 1;";
for (int i = 0; i < num_values; i++)
  kernel << "out_val*=a"<<i<<"[global_idx];";
kernel << "out[global_idx]=out_val;}}";

所有这些方法都有点臭。多次调用内核会一遍又一遍地增加内核调用开销。复制缓冲区是完全不必要的工作。编写内核是完全不可读的，并且需要在每个新计数上产生额外的开销。你如何解决这个问题？

如果您知道每个缓冲区的长度，并且它们都相等。那为什么不这样重写呢：

主机端：

//Use one struct to organize the data
struct MyStruct{
    cl_int values[10];
    // ...
};
//Create one only cl::Buffer of N elements of type MyStruct
cl::Buffer mybuff = cl::Buffer(context, CL_MEM_READ_ONLY, in_size * sizeof(MyStruct));
//Run the kernel

设备端：

//Use one struct to organize the data
typedef struct{
    int values[10];
    // ...
}MyStruct;
__global void pix_reduce
  (MyStruct out*, MyStruct* in)
{
   // ...
}

您甚至可以将长度"10"定义为 X，并将其放入代码中（主机和 CPU 端）。另一个不错的选择是一个连续内存的大数组，并将分离长度传递给内核。

注意：我知道很难停止使用指针，但这是要付出的代价，因为内存区域不同。相信我，我已经完成了非常复杂的动态长度结构，并且经过一些思考和重写，它就可以工作了，并且代码非常整洁。

编辑：另一种可以动态更改的解决方案

主机端：

//Data organized as jagged array
std::vector<std::vector<cl_int> > my_data;
//It is supposed to be filled my_Data
//....
//Create one only cl::Buffer of NxM elements of type int
cl::Buffer mybuff = cl::Buffer(context, CL_MEM_READ_ONLY, my_data.size() * my_data[0].size() * sizeof(cl_int));
cl::Buffer mybuff_out = cl::Buffer(context, CL_MEM_WRITE_ONLY, my_data.size() * my_data[0].size() * sizeof(cl_int));
// Copy the jagged array
for(int i=0; i<my_data.size(); i++)
queue.enqueueWriteBuffer(mybuff, CL_FALSE, i * my_data[0].size() * sizeof(cl_int), my_data[0].size() * sizeof(cl_int), &my_data[i][0]);
//Set kernel
kernel.SetArgs(0, mybuff);
kernel.SetArgs(1, mybuff_out);
kernel.SetArgs(2, (cl_int)my_data[0].size());
//Run the kernel
queue.EnqueueNDRangeKernel(...);

设备端：

__kernel void pix_reduce
  (_global const int * in, _global int * out, const int size)
{
     unsigned int id = get_global_id(0);
     for(int i=0; i<size; i++){
          out[id*size+i] = out[id*size+i]; //Or any other operation using size as the separation between chunks
     }
}
相关文章：