OpenCL编程步骤（四）：创建内核对象和设置内核参数

内核就是程序中声明的一个函数。对于程序中的任一函数，都可以通过加上限定符__kernel将其标识为内核。内核对象中封装了程序中的某个__kernel函数以及执行此函数时所需的参数。

1、创建内核

cl_kernel clCreateKernel (cl_program program,
                          const char *kernel_name,
                          cl_int *errcode_ret)

此函数用来创建内核对象。

• program，是一个程序对象，带有成功构建的执行体。
• kernel_name，是程序中一个声明时带有限定符__kernel的函数名。
• errcode，返回相应的错误码

成功创建内核对象，会将其返回，并将errcode_ret置为CL_SUCCESS。否则返回NULL，并返回下列错误码之一：

• CL_INVALID_PROGRAM，program无效
• CL_INVALID_PROGRAM_EXECUTABLE，没有为program成功构建执行体。
• CL_INVALID_KERNEL_NAME，program中找不到kernel_name。
• CL_INVALID_KERNEL_DEFINITION，kernel_name的函数定义在已经具有对应执行体的那些设备上有所区别。
• CL_INVALID_VALUE，kernel_name是NULL。
• CL_OUT_OF_RESOURCES，为设备上的OpenCL操作分配内存失败。
  
• CL_OUT_OF_HOST_MEMORY，为主机上的OpenCL操作分配内存失败。
  

clCreateKernel一次只为一个内核函数创建内核对象。要是有多个内核函数时，需要一个一个的进行创建吗？

OpenCL提供了另外一个函数来解决这个问题。

cl_int clCreateKernelInProgram (cl_program program,
                       cl_uint num_kernels,
                       cl_kernel *kernels,
                       cl_uint *num_kernls_ret)

这个函数是为program中所有的内核函数创建对应的内核对象。如果某个__kernel函数的定义在已经具有执行体的那些设备上不完全一样，则不会为其创建内核对象。

• program，一个程序对象，具有成功构建的执行体。
• num_kernels，程序对象中内核函数的数目。
• kernels，用来存储多返回的内核对象。如果为NULL，则忽略；否则，num_kernels的值必须大于或等于program中内核的数目。
• num_kernels_ret，program中的内核的数目，如果为NULL，则忽略。

这个函数一般调用两次，第一次确定program中的内核数；第二次具体创建内核对象。下面的代码块展示了这个函数的使用方法：

cl_uint numKernels;
errNum = clCreateKernelInProgram(program, NULL, NULL, &numKernels);
cl_kernel *kernels = new cl_kernel[numKernels];
errNum = clCreateKernelInProgram(program, numKernels, kernels, NULL);

2、设置内核参数

想要执行内核，就必须设置内核参数。

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

这个函数用来设置内核的某个参数。

• kernel，是内核对象。
• arg_index是参数索引。内核参数的索引值从最左边的0开始，一直到n-1，其中n是参数的总数。
• arg_size，参数的大小。这个大小由内核函数中参数的声明方式如何声明来确定。
• arg_value，传入内核函数的参数的一个指针。

3、查询和管理内核对象

cl_int clGetKernelInfo (cl_kernel kernel,
                        cl_kernel_info param_name,
                        size_t param_value_size,
                        void *param_value,
                        size_t *param_value_size_ret)

这个函数用来查询内核对象的相关信息。

• kernel，索要查询的内核对象。
• param_name，指定所要查询的信息。
• param_value，执行的内存用来存储查询结果。如果为NULL，则忽略。
• pram_value_size，param_value所指向内存块的大小。其值必须大于param_name对应的类型的大小。
• param_value_siez_ret，返回查询结果的实际大小。

这个函数也需要调用两次。

clGetKernelInfo所支持的param_name

cl_kernel_info返回类型CL_KERNEL_FUNCTION_NAMEchar []返回内核函数的名字。
CL_KERNEL_NUM_ARGScl_uint返回内核参数的个数。
CL_KERNLE_REFERENCE_COUNTcl_uint返回kernel的引用计数。
CL_KERNEL_CONTEXTcl_context返回kernel所在的上下文。
CL_KERNEL_PROGRAMcl_program返回kernel所关联的程序对象。
CL_KERNEL_ATTRIBUTESchar []返回程序源码中声明内核函数时所有通过__attributes__指定的特性。

另外一个查询内核信息的函数：

clGetKernelWorkGroupInfo (cl_kernel kernel,
                    cl_device_id device,
                    cl_kernel_work_group_info param_name,
                    size_t param_value_size,
                    void *param_value,
                    size_t *param_value_size_ret)

这个函数是查询内核对象针对某个设备的相关信息。

clGetKernelWorkGroupInfo所支持的param_name

cl_kernel_work_group_info返回类型CL_KERNEL_GLOBAL_WORK_SIZEsize_t[3]利用此机制，应用可以查询用来在device上执行内核的全局索引空间的大小（即clEnqueueNDRangeKernel的参数global_work_size）。
这要求device是自定义设备或所执行的内核是内建的，否则clGetKernelWorkGroupInfo会返回CL_INVALID_VALUE.
CL_KERNEL_WORK_GROUP_SIZEsize_t应用可以查询用来在device上执行内核的工作组的大小。
CL_KERNEL_COMPILE_WORK_GROUP_SIZEsize_t [3]返回限定符__arrtibute__((reqd_work_group_size(X, Y, Z))) 所指定的工作组的大小。
CL_KERNEL_LOCAL_MEM_SIZEcl_ulong返回内核所使用的局部内存的大小。
对于任一带有限定符__local的指针参数，如果没有为其指定内存大小，则假定为0。
CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLEsize_t返回所期望的工作组大小的粒度。仅为性能的建议。
CL_KERNEL_PRIVATE_MEM_SIZEcl_ulong内核中每个工作项至少要使用多少私有内存。

4、查询内核参数信息

cl_int clGetKernelArgInfo (cl_kernel kernel,
                     cl_uint arg_index,
                     cl_kernel_arg_info param_name,
                     size_t param_value_size,
                     void *param_value,
                     size_t *param_value_size_ret) 

此函数会返回内核参数的相关信息。只有满足下列条件时，内核参数信息才可用：

• kernel所关联的程序对象是用clCreateProgramWithSource创建的。
• 用cl{Build | Compile}program构建程序执行体时，在参数options中指定了-cl-kernel-arg-info。

cl_kernel_arg_info的类型：

cl_kernel_arg_info返回类别CL_KERNEL_ARG_ADDRESS_QUALIFIERcl_kernel_arg_address_qualifier返回参数的地址限定符。所返回的值可以是下列之一：

• CL_KERNEL_ARG_ADDRESS_GLOBAL
• CL_KERNEL_ARG_ADDRESS_LOCAL
• CL_KERNEL_ARG_ADDRESS_CONSTANT 
• CL_KERNEL_ARG_ADDRESS_PRIVATE

如果没有指定地址限定符，则返回缺省的地址限定符CL_KERNEL_ADDRESS_PRIVATE。  
CL_KERNEL_ARG_ACCESS_QUALIFIERcl_kernel_arg_access_qualifier返回参数的访问限定符。可以是下列值之一：

• CL_KERNEL_ARG_ACCESS_READ_ONLY
• CL_KERNEL_ARG_ACCESS_WRITE_ONLY 
• CL_KERNEL_ARG_ACCESS_READ_WRITE 
• CL_KERNEL_ARG_ACCESS_NONE

如果参数类型不是图像，则返回CL_KERNEL_ARG_ACCESS_NONE。如果参数类别是图像，则会返回所制定的访问限定符或缺省的访问限定符。      
CL_KERNEL_ARG_ARG_TYPE_NAMEchar []返回arg_index所指定的参数的类型名。
CL_KERNEL_ARG_TYPE_QUALIFIERcl_kernel_arg_type_qualifier返回参数的类型限定符。可能是：

• CL_KERNEL_ARG_TYPE_CONST
• CL_KERNEL_ARG_TYPE_RESTRICT 
• CL_KERNEL_ARG_TYPE_VOLATILE
• 以上三种的组合；或者
• CL_KERNEL_ARG_TYPE_NONE   

CL_KERNEL_ARG_NAMEchar []返回参数的名字。

示例程序：

__kernel void hello_kernel(__global const float *a,
                           __global const float *b,
                           __global float *result)
{
    int gid = get_global_id(0);
    result[gid] = a[gid] + b[gid];
}

kernel = clCreateKernel(program, "hello_kernel", &errNum);
if (kernel == NULL || errNum != CL_SUCCESS)
{
    cerr << "Failed to create kernel." << endl;
    Cleanup(context, commandQueue, program, kernel, memObjects);
    return -1;
}
 
errNum = clSetKernelArg(kernel, 0, sizeof(cl_mem), &memObjects[0]);
errNum |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &memObjects[1]);
errNum |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &memObjects[2]);
 
if (errNum != CL_SUCCESS)
{
    cerr << "Error setting kernel arguments." << endl;
    Cleanup(context, commandQueue, program, kernel, memObjects);
    return -1;
}