OpenCL 设置内核参数|极客笔记

为了执行一个具体内核，必须能够向内核函数传递参数。OpenCL使用如下函数来设置内核参数：

cl_int clSetKernelArg(cl_kernel kernel,
   cl_uint arg_index,
   size_t arg_size,
   const void ＊arg_value)

参数kernel为内核对象。参数arg_index为内核函数参数的索引，以从左到右的次序，第一个（即最左边的）参数的索引为0，第二个参数的索引为1，依此类推。

参数arg_size为参数大小。

如果参数为内存对象，则arg_size为内存对象类型的大小（例如，sizeof（cl_mem））；
如果参数用local修饰符声明，则arg_size为用来存储参数的缓冲区所需的字节数；
如果参数为sampler_t类型，则arg_size大小必须等于sizeof（cl_sampler）；如果参数为queue_t类型，则arg_size大小必须等于sizeof（cl_command_queue）；
如果参数为其他类型，arg_size大小为该类型参数大小（例如，对于一个cl_int参数，大小为sizeof（cl_int））。

参数arg_value为传入内核参数的一个指针。

如果参数为内存对象，则指针指向合适的缓冲、管道、图像或图像数组对象；
如果参数用local修饰符声明，arg_value必须为NULL；
如果参数为sampler_t，arg_value必须指向采样器对象；
如果参数为queue_t类型，arg_value必须指向设备队列对象。

对于不同参数类型，参数arg_size和arg_value具体使用，在后面的章节会有比较详细的例子。目前以前面例子hello_kernel1内核函数为例：

err = clSetKernelArg(kernel[0], 0, sizseof(cl_mem), &dev_a);
err = clSetKernelArg(kernel[0], 1, sizseof(cl_mem), &dev_b);
err = clSetKernelArg(kernel[0], 2, sizseof(cl_mem), &dev_result);

在hello_kernel1内核函数中有a、b和result三个参数，分别对应的参数索引为0、1、2。三个参数都为_global修饰符声明的内存对象，故参数arg_size为sizeof（cl_mem）。

对于共享虚拟内存（SVM）可以使用如下函数设置内核参数：

cl_int clSetKernelArgSVMPointer(cl_kernel kernel,
                                      cl_uint arg_index,
                                      const void ＊arg_value)

clSetKernelArg（）中参数意义一样。参数arg_value为指向SVM的指针。SVM指针只能用于指向global或者constant内存的指针。

对于SVM来说，仅仅是clSetKernelArgSVMPointer来设置内核函数参数值，这可能还不够，可以使用如下函数来向内核参数传递额外的信息：

cl_int clSetKernelExecInfo(cl_kernel kernel,
                                cl_kernel_exec_info param_name,
                                size_t param_value_size,
                                const void ＊param_value)

参数param_name指定向内存传递的信息。在表3-11中列出了支持的信息名称以及相应的描述。
参数param_value为指向内存的指针，该内存中的值由参数param_name确定。参数param_value_size为参数param_value指向的内存空间的字节大小。

内核中使用的对于粗粒度和细粒度缓冲SVM指针，如果不作为内核参数传递给内核，则必须使用clSetKernelExecInfo（）函数，param_name为CL_KERNEL_EXEC_INFO_SVM_PTRS来指定。例如，如果SVM内存A包含一个指向SVM内存B的指针，内核间接引用这个指针，则指向SVM内存B的指针必须作为参数传递或者使用clSetKernelExecInfo（）。下面的例子中，inputElements包含了inputFloats的指针，inputFloats不是作为参数传入内核，而是使用clSetKernelExecInfo（）函数：

typedef struct _Element
{
    float ＊internal;
    float ＊external;
    float value;
} Element;
……
//分配细粒度SVM空间
Element ＊inputElements = (Element ＊)clSVMAlloc(context,
                              CL_MEM_READ_ONLY | CL_MEM_SVM_FINE_GRAIN_BUFFER,
                              sizeof(Element) ＊ size, 0);
float ＊inputFloats = (float ＊)clSVMAlloc(context,
                        CL_MEM_READ_ONLY | CL_MEM_SVM_FINE_GRAIN_BUFFER,
                        sizeof(float) ＊ size, 0);
float ＊output = (float ＊)clSVMAlloc(context,
                        CL_MEM_READ_ONLY | CL_MEM_SVM_FINE_GRAIN_BUFFER,
                        sizeof(float) ＊ size, 0);
for(int i = 0; i 〈 size; i++)
{
    inputElements[i].value = i;
    inputFloats[i] = i + size;
}
for(int i = 0; i 〈 size; i++)
{
 int index = rand_index(i); //根据i产生随机数
 //inputElements元素值包含了inputFloats的地址空间值
 inputElements[i].internal = &(inputElements[index].value);
 inputElements[i].external = &(inputFloats[index]);
}
clSetKernelArgSVMPointer(kernel, 0, inputElements);
clSetKernelArgSVMPointer(kernel, 1, output);
//inputFloats通过clSetKernelExecInfo传入内核
clSetKernelExecInfo(kernel, CL_KERNEL_EXEC_INFO_SVM_PTRS,
                  sizeof(inputFloats), &inputFloats);
……
clEnqueueNDRangeKernel(cmdqueue, kernel, 1, 0, &globalsize, NULL,
                      NULL, NULL, NULL);

如上例子，分配了inputElements和inputFloats的SVM空间。inputElements中internal和external指向了inputFloats的地址空间位置。而inputFloats不是作为内核函数参数传入内核，此时就应该应用clSetKernelExecInfo（）函数，把inputFloats传入内核。