GPGPU OpenCL 获取kernel函数编译信息

　　使用OpenCL编程时，kernel写成一个单独的文件或者将文件内容保存在一个string中。可以使用clBuildProgram对kernel进行编译链接(compiles & links)，如果失败，可以使用clGetProgramBuildInfo获取OpenCL编译器对kernel的编译信息。

1.clBuildProgram 

　cl_int clBuildProgram (

　　　　cl_program program,　　//program
　　　　cl_uint num_devices,　　//the number of device
　　　　const cl_device_id *device_list, 　　//devices id
　　　　const char *options,　　//the option of compiler
　　　　void (CL_CALLBACK *pfn_notify)(cl_program program, void *user_data),　　//the callback function
　　　　void *user_data)　　//the data of callback function
　　)

2.clGetProgramBuildInfo

　　cl_int clGetProgramBuildInfo (

　　　　cl_program program, 　　//program
　　　　cl_device_id device,　　//the id of device
　　　　cl_program_build_info param_name,
　　　　size_t param_value_size,
　　　　void *param_value,
　　　　size_t *param_value_size_ret
　　)

3.代码实例(获取编译器对kernel的编译信息)

3.1 kernel(build_info_kernel.cl)

__kernel void good(__global float *a,
                   __global float *b,
                   __global float *c) {
   
    *c = *a + *b;
}

__kernel void good(__global float *a,
                   __global float *b,
                   __global float *c) {
    __local int var=3;   
    int size=get_local_sze(0);
    *c = *a + *b;
}

 3.2 tool.h

#ifndef TOOLH
#define TOOLH
#include <CL/cl.h>
#include <string.h>
#include <stdio.h>
#include <stdlib.h>
#include <iostream>
#include <string>
#include <fstream>
using namespace std;

/** convert the kernel file into a string */
int convertToString(const char *filename, std::string& s);

/**Getting platforms and choose an available one.*/
int getPlatform(cl_platform_id &platform);

/**Step 2:Query the platform and choose the first GPU device if has one.*/
cl_device_id *getCl_device_id(cl_platform_id &platform);

/**获取编译program出错时，编译器的出错信息*/
int getProgramBuildInfo(cl_program program,cl_device_id device);
#endif

　   tool.cpp

#include <CL/cl.h>
#include <string.h>
#include <stdio.h>
#include <stdlib.h>
#include <iostream>
#include <string>
#include <fstream>
#include "tool.h"
using namespace std;

/** convert the kernel file into a string */
int convertToString(const char *filename, std::string& s)
{
    size_t size;
    char*  str;
    std::fstream f(filename, (std::fstream::in | std::fstream::binary));

    if(f.is_open())
    {
        size_t fileSize;
        f.seekg(0, std::fstream::end);
        size = fileSize = (size_t)f.tellg();
        f.seekg(0, std::fstream::beg);
        str = new char[size+1];
        if(!str)
        {
            f.close();
            return 0;
        }

        f.read(str, fileSize);
        f.close();
        str[size] = '';
        s = str;
        delete[] str;
        return 0;
    }
    cout<<"Error: failed to open file
:"<<filename<<endl;
    return -1;
}

/**Getting platforms and choose an available one.*/
int getPlatform(cl_platform_id &platform)
{
    platform = NULL;//the chosen platform

    cl_uint numPlatforms;//the NO. of platforms
    cl_int    status = clGetPlatformIDs(0, NULL, &numPlatforms);
    if (status != CL_SUCCESS)
    {
        cout<<"Error: Getting platforms!"<<endl;
        return -1;
    }

    /**For clarity, choose the first available platform. */
    if(numPlatforms > 0)
    {
        cl_platform_id* platforms =
            (cl_platform_id* )malloc(numPlatforms* sizeof(cl_platform_id));
        status = clGetPlatformIDs(numPlatforms, platforms, NULL);
        platform = platforms[0];
        free(platforms);
    }
    else
        return -1;
}

/**Step 2:Query the platform and choose the GPU device*/
cl_device_id *getCl_device_id(cl_platform_id &platform)
{
    cl_uint numDevices = 0;
    cl_device_id *devices=NULL;
    cl_int    status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &numDevices);
    if (numDevices > 0) //GPU available.
    {
        devices = (cl_device_id*)malloc(numDevices * sizeof(cl_device_id));
        status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, numDevices, devices, NULL);
    }
    return devices;
}

/**获取编译program出错时，编译器的出错信息*/
int getProgramBuildInfo(cl_program program,cl_device_id device)
{
    size_t log_size;
    char *program_log;
    /* Find size of log and print to std output */
    clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 
            0, NULL, &log_size);
    program_log = (char*) malloc(log_size+1);
    program_log[log_size] = '';
    clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 
            log_size+1, program_log, NULL);
    printf("%s
", program_log);
    free(program_log);
    return 0;
}

3.3 buildInfo.cpp

#include "tool.h"
#include <string.h>
#include <stdio.h>
#include <stdlib.h>
#include <iostream>
#include <string>
#include <fstream>
using namespace std;

void CL_CALLBACK checkData(cl_program platform, void* data){
    printf("%s
",(char*)data);
}

int main(int argc, char* argv[])
{
    cl_int    status;
    /** Getting platforms and choose an available one(first).*/
    cl_platform_id platform;
    getPlatform(platform);

    /**Query the platform and choose the GPU device.*/
    cl_device_id *devices=getCl_device_id(platform);

    /**Create context use the frist device.*/
    cl_context context = clCreateContext(NULL,1, devices,NULL,NULL,NULL);

    /**Create program object */
    const char *filename = "build_info_kernel.cl";
    string sourceStr;
    status = convertToString(filename, sourceStr);
    const char *source = sourceStr.c_str();
    size_t sourceSize[] = {strlen(source)};
    cl_program program = clCreateProgramWithSource(context, 1, &source, sourceSize, NULL);

    /**Build program. */
    //status=clBuildProgram(program, 1,devices,NULL,checkData,"sdf");
    status=clBuildProgram(program, 1,devices,NULL,NULL,NULL);
    if(status < 0)    //get the build info
        getProgramBuildInfo(program ,devices[0]);
    else
        printf("Build Success
");
    
    status = clReleaseProgram(program);    //Release the program object.
    status = clReleaseContext(context);//Release context.
    free(devices);

    getchar();
    return 0;
}

对kernel的编译结果：