OpenCL 教程 第三章 程序模型

#第三章 程序模型
创建程序的过程实际上就是将程序读入主机的过程，并将程序的内容保存在cl_program_id类型的结构体中，跟其他语言一样，我们读取程序后需要对程序进行编译和链接来创建程序，编译程序的接口如下：

cl_int clBuildProgram (cl_program program,                     //需要编译的程序对象                                   
                       cl_uint num_devices,                    //程序有关的设备列表中的设备数量
                       const cl_device_id *device_list,        //程序有关的设备列表，如果为NULL表示为所有与程序有关的设备编译程序
                       const char *options,                    //编译选项
                       void (CL_CALLBACK *pfn_notify)          //错误信息回调函数
                       (cl_program program,void *user_data),
                       void *user_data)                        //当发生错误时回调函数会访问该参数

常用的编译选项：

同样可以将编译和链接的过程分开执行：

cl_int clCompileProgram (cl_program program,                    //需要编译的程序对象
                         cl_uint num_devices,                   //程序有关的设备列表中的设备数量
                         const cl_device_id *device_list,       //程序有关的设备列表，如果为NULL表示为所有与程序有                                                                   关的设备编译程序
                         const char *options,                   //编译选项
                         cl_uint num_input_headers,             //程序文件包含的头文件数目
                         const cl_program *input_headers,       //头文件程序列表
                         const char **header_include_names,     //头文件程序名称列表
                         void (CL_CALLBACK *pfn_notify)         //错误信息回调函数
                         (cl_program program,void *user_data),  
                         void *user_data)                       //当发生错误时回调函数会访问该参数

举例：

//程序文件内容
#include <foo.h>
#include <mydir/myinc.h>
__kernel void hello (...)
{
...
}
//主机端程序内容
cl_program foo_pg = clCreateProgramWithSource(context,1, &foo_header_src, NULL, &err);
cl_program myinc_pg = clCreateProgramWithSource(context,1, &myinc_header_src, NULL, &err);
// let’s assume the program source described above is given
// by program_A and is loaded via clCreateProgramWithSource
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_data

cl_program clLinkProgram (cl_context context,                  //与设备有关的上下文  
                          cl_uint num_devices,                 //程序有关的设备列表中的设备数量
                          const cl_device_id *device_list,     //程序有关的设备列表，如果为NULL表示为所有与程序有                                                                  关的设备编译程序
                          const char *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)                 //程序状态指示

当程序创建和编译完成后，我们可以通过调用接口来获取生成日志，编译信息。

cl_int clGetProgramBuildInfo (cl_program program,               //获取信息的程序对象
                              cl_device_id device,              //获取信息的设备对象
                              cl_program_build_info param_name, //获取的信息名称
                              size_t param_value_size,          //获取的信息内容大小
                              void *param_value,                //信息内容指针
                              size_t *param_value_size_ret)     //返回的信息内容大小

程序信息参数：

实验代码：

main.cpp:
#include"opencl.h"
#define filesname "hello_world.cl"
char *readKernelSourceFile(const char *filename, size_t *filelength){
    FILE *file = NULL;
    size_t source_length;
    char *source_string;
    int ret;
    file = fopen(filename, "rb");
    if (file == NULL){
        //printf("%s at %d :Can't open %s\n", __FILE__, __LINE__ - 2, filename);
        return NULL;
    }
    fseek(file, 0, SEEK_END);
    source_length = ftell(file);
    fseek(file, 0, SEEK_SET);
    source_string = (char *)malloc(source_length + 1);
    //source_string[source_length] = '\0';
    ret = fread(source_string, source_length, 1, file);
    if (ret == 0){
        printf("%s at %d :Can't read source %s \n", __FILE__, __LINE__ - 2, filename);
        return NULL;
    }
    fclose(file);
    if (source_length != 0){
        *filelength = source_length;
    }
    source_string[source_length] = '\0';
    return source_string;
}
int main(){
    cl_uint platform_num;
    clGetPlatformIDs(0, NULL, &platform_num);
    cl_platform_id *platforms = new cl_platform_id[platform_num];
    clGetPlatformIDs(platform_num, platforms, NULL);
    cl_uint device_num;
    clGetDeviceIDs(platforms[1], CL_DEVICE_TYPE_ALL, 0, NULL, &device_num);
    cl_device_id *devices = new cl_device_id[device_num];
    clGetDeviceIDs(platforms[1], CL_DEVICE_TYPE_ALL, device_num, devices, NULL);
    cl_int err;
    cl_context context = clCreateContext(NULL, device_num, devices, NULL, NULL, &err);
    checkError(err, "Can't create context");
    cl_command_queue *queues = new cl_command_queue[device_num];
    for (cl_uint i = 0; i < device_num; i++){
        queues[i] = clCreateCommandQueue(context, devices[i], CL_QUEUE_PROFILING_ENABLE, &err);
        checkError(err, "Can't create queue");
    }
    size_t file_length;
    char * files_string = readKernelSourceFile(filesname, &file_length);
    cl_program program = clCreateProgramWithSource(context, 1, (const char **)&files_string, &file_length, &err);
    checkError(err, "Failed to create program with source");
    err = clBuildProgram(program, device_num, devices, NULL, NULL, NULL);
    checkError(err, "Can't build the program");
    size_t value_size;
    clGetProgramBuildInfo(program, devices[0], CL_PROGRAM_BUILD_LOG, NULL, NULL, &value_size);
    char *bulid_value = new char[value_size];
    clGetProgramBuildInfo(program, devices[0], CL_PROGRAM_BUILD_LOG, value_size, bulid_value, NULL);
    cout << string(bulid_value) << endl;
    return 0;
}

hello_world.cl:

__kernel void hello_world(){
    int A,B;
    a+b;    //test build log
}

实验结果：