OpenCL 编译Program对象|极客笔记

OpenCL是一个跨平台工业标准，OpenCL C代码只有在选择设备对象后才能确定将会运行在何种计算设备上。

clBuildProgram

在运行前需要对程序对象进行构建，如以下代码所示：

cl_int clBuildProgram(cl_program program,
                            cl_uint num_deivces,
                            const cl_device_id ＊device_list,
                            const char ＊options,
                            void(CL_CALLBACK ＊pfn_notify)(
                                  cl_program program,
                                  void ＊user_data),
                            void ＊user_data)

clBuildProgram（）会为指定的所有设备构建程序对象（或者，如果没有指定设备列表，则为与上下文关联的所有设备构建程序对象）。这一步实际等价于在一个C程序上调用编译器与链接器。第4个参数设置编译器选项，这些编译选项与常用的编译器（如gcc）类似，但有很大一部分是专门针对OpenCL。编译器选项可以分为以下类型：

预处理选项；
数学选项；
优化选项；
其他选项。

这几个选项的详细情况，如下表所示。

编译选项

下面两行代码展示了如何使用上表中的这些选项。为一个设备（device）编译和链接program，使用的编译链接选项为：-cl-std、-cl-mad-enable和-Werror。

const char options[] = " -cl-std = CL1.1 -cl-mad-enable
                              -Werror";
clBuildProgram(program, 1, &device, option, NULL, NULL);

clCompileProgram

大家可能会发现，在上述讲述构件程序内容中，我们使用clBuildProgram（）把编译和链接过程合在一起处理了。但是从OpenCL 1.2开始，可以把编译和链接分开。分开编译和链接支持如下使用：

分开编译和链接步骤。编译程序源码，产生一个编译过的二进制对象。在一个单独的步骤中，这个二进制对象与其他编译过的程序对象链接成一个程序。
内嵌头文件。在编译选项中-I选项可以指定包含程序源码的头文件目录。从OpenCL 1.2开始扩展了该功能，运行头文件代码来源于其他程序对象。
库。使用链接器链接编译过的对象和库到一个程序中，或者把编译过的二进制生成一个库。

可以使用如下函数实现编译：

cl_int clCompileProgram(cl_program program,
                            cl_uint num_devices,
                            const cl_device_id ＊device_list,
const char ＊options,
const_cl_program ＊input_header,
const char ＊＊header_include_name,
void(CL_CALLBACK ＊pfn_notify)
                    (cl_program program, void ＊user_data),
void ＊user_data)

对于clCompileProgram（），参数program为编译出的程序对象。参数input_header，是使用clCreateProgramWithSource创建的内嵌头文件代码的程序数组。参数num_input_headers指定参数input_header中程序的数量。其他参数与clBuildProgram（）中参数类似。

#include 〈foo.h>
#include 〈mydir/myinc.h>
kernel void   image_filter (int n, int m,
                              constant float ＊filter_weights,
read_only image2d_t src_image,
                              write_only image2d_t dst_image)
{
    ...
}

对于如上内核代码，包含了两个头文件foo.h和mydir/myinc.h。下列代码展示了这两个头文件在程序对象中如何传递：

cl_program foo_pg = clCreateProgramWithSource(context, 1,
&foo_header_src, NULL, &err);
cl_program myinc_pg = clCreateProgramWithSource(context, 1,
&myinc_header_src, NULL, &err);
cl_program input_headers[2] = { foo_pg, myinc_pg };
char ＊input_header_names[2] = { "foo.h", "mydir / myinc.h" };
clCompileProgram(program_A,
                    0, NULL,        //num_devices & device_list
                    NULL,           //compile_options
                    2,              //num_input_headers
                    input_headers,
                    input_header_names,
                    NULL, NULL);    //pfn_notify & user_da

上述代码展示了内核代码与已编译过的头文件二进制对象如何一起编译。两个头文件先用clCreateProgramWithSource（）函数编译为二进制。内核代码包含了两个头文件，故num_input_headers设置为2，input_header_name数组分为两个头文件名称。

clLinkProgram

对于已编译好的对象，可使用如下函数链接成程序对象：

cl_program clLinkProgram(cl_context context,
                              cl_uint num_devices,
                              const cl_device_id ＊device_list,
                              const chat ＊options,
                              cl_uint num_input_programs,
const cl_program ＊input_programs,
void(CL_CALLBACK ＊pfn_notify)cl_program
                        ( program,
                        void ＊user_data),
      void ＊user_data,
      cl_int ＊errcode_ret)

参数input_programs是链接到程序对象中的已编译过的二进制或库的程序对象。参数num_input_programs指定参数input_programs中程序对象数目。其他参数与clBuildProgram（）中参数类似。

下面通过一个例子详细讲解分开编译和链接的使用。kernel.cl文件内容如下：

#include"add.h"
kernel void call_test(global float ＊A, const float b)
{
    int index = get_global_id(0);
    A[index] = add(A[index], b);
}

add.h文件内容如下：

float add(float a, float b)
{
    return (a + b) + 100.0;
}

对应的主机端代码如下所示，基于减少篇幅的原因，去掉了许多不重要的代码。

……
cl_program program_cl, program_head, program;
unsigned char ＊program_file;
size_t program_size       //读取内核源码
read_file(&program_file, &program_size, "kernel.cl");
//创建内核程序对象
program_cl = clCreateProgramWithSource(context, 1, &program_file,
                                              &program_size, &err);
//读取头文件源码
read_file(&program_file, &program_size, "add.h");
//创建头文件程序对象
program_ head = clCreateProgramWithSource(context, 1, &program_file,
                                                  &program_size, &err);
const char ＊input_head_names[1] = {"add.h"};
cl_program input_head[1] = {program_head};
//把头文件程序对象与内核程序对象一起编译
err = clCompileProgram(program_cl, 0, NULL, 0, 1, input_head,
                            input_head_names, NULL, NULL);
//链接为程序
program = clLinkProgram(context, num_device, devices, NULL, 1,
&program_cl, NULL, NULL, &err);
……

大家可能会发现上表编译选项中-I选项也可以完成如上功能，对于该实现方式，在此就不列出代码。