OpenCL 学习step by step (3) 存储kernel文件为二进制

     在教程二中，我们通过函数convertToString，把kernel源文件读到一个string串中，然后用函数clCreateProgramWithSource装入程序对象，再调用函数clBuildProgram编译程序对象。其实我们也可以直接调用二进制kernel文件，这样，当不想把kernel文件给别人看的时候，起到一定的保密作用。在本教程中，我们会把读入的源文件存储一个二进制文件中，并且还会建立一个计时器类，用来记录数组加法在cpu和gpu端分别执行的时间。

     首先我们建立工程文件gclTutorial2,在其中增加类gclFile,该类主要用来读取文本kernel文件，或者读写二进制kernel文件。

class gclFile
{
public:
    gclFile(void);
    ~gclFile(void);

    //打开opencl kernel源文件(文本模式)
    bool open(const char* fileName);

    //读写二进制kernel文件
    bool writeBinaryToFile(const char* fileName, const char* birary, size_t numBytes);
    bool readBinaryFromFile(const char* fileName);

…

}

gclFile中三个读写kernel文件的函数代码为：

bool gclFile::writeBinaryToFile(const char* fileName, const char* birary, size_t numBytes)
{
    FILE *output = NULL;
    output = fopen(fileName, "wb");
    if(output == NULL)
        return false;

    fwrite(birary, sizeof(char), numBytes, output);
    fclose(output);

    return true;
}


bool gclFile::readBinaryFromFile(const char* fileName)
{
    FILE * input = NULL;
    size_t size = 0;
    char* binary = NULL;

    input = fopen(fileName, "rb");
    if(input == NULL)
    {
        return false;
    }

    fseek(input, 0L, SEEK_END); 
    size = ftell(input);
    //指向文件起始位置
    rewind(input);
    binary = (char*)malloc(size);
    if(binary == NULL)
    {
        return false;
    }
    fread(binary, sizeof(char), size, input);
    fclose(input);
    source_.assign(binary, size);
    free(binary);

    return true;
}

bool gclFile::open(const char* fileName)   //!< file name
{
    size_t      size;
    char*       str;

    //以流方式打开文件
    std::fstream f(fileName, (std::fstream::in | std::fstream::binary));

    // 检查是否打开了文件流
    if (f.is_open())
    {
        size_t  sizeFile;
        // 得到文件size
        f.seekg(0, std::fstream::end);
        size = sizeFile = (size_t)f.tellg();
        f.seekg(0, std::fstream::beg);

        str = new char[size + 1];
        if (!str)
        {
            f.close();
            return  false;
        }

        // 读文件
        f.read(str, sizeFile);
        f.close();
        str[size] = '\0';

        source_  = str;

        delete[] str;

        return true;
    }

    return false;
}

现在，在main.cpp中，我们就可以用gclFile类的open函数来读入kernel源文件了：

//kernel文件为add.cl

gclFile kernelFile;
if(!kernelFile.open("add.cl"))
    {
    printf("Failed to load kernel file \n");
    exit(0);
    }
const char * source = kernelFile.source().c_str();
size_t sourceSize[] = {strlen(source)};
//创建程序对象
cl_program program = clCreateProgramWithSource(
    context,
    1,
    &source,
    sourceSize,
    NULL);

    编译好kernel后，我们可以通过下面的代码，把编译好的kernel存储在一个二进制文件addvec.bin中，在教程四中，我们将会直接装入这个二进制的kernel文件。

//存储编译好的kernel文件
char **binaries = (char **)malloc( sizeof(char *) * 1 ); //只有一个设备
size_t *binarySizes = (size_t*)malloc( sizeof(size_t) * 1 );

status = clGetProgramInfo(program, 
    CL_PROGRAM_BINARY_SIZES,
    sizeof(size_t) * 1, 
    binarySizes, NULL);
binaries[0] = (char *)malloc( sizeof(char) * binarySizes[0]);
status = clGetProgramInfo(program, 
    CL_PROGRAM_BINARIES,
    sizeof(char *) * 1, 
    binaries, 
    NULL);    
kernelFile.writeBinaryToFile("vecadd.bin", binaries[0],binarySizes[0]);

    我们还会建立一个计时器类gclTimer,用来统计时间，这个类主要用QueryPerformanceFrequency得到时钟频率，用QueryPerformanceCounter得到流逝的ticks数，最终得到流逝的时间。函数非常简单

class gclTimer
{
public:
    gclTimer(void);
    ~gclTimer(void);

private:

    double _freq;
    double _clocks;
    double _start;
public:
    void Start(void); // 启动计时器
    void Stop(void); //停止计时器
    void Reset(void); //复位计时器
    double GetElapsedTime(void); //计算流逝的时间
};

下面我们在cpu端执行数组加法时，增加计时器的代码：

gclTimer clTimer;
clTimer.Reset();
clTimer.Start();
//cpu计算buf1,buf2的和
for(i = 0; i < BUFSIZE; i++)
    buf[i] = buf1[i] + buf2[i];
clTimer.Stop();
printf("cpu costs time:%.6f ms \n ", clTimer.GetElapsedTime()*1000 );

同理在gpu执行kernel代码，以及copy gpu结果到cpu时候，增加计时器代码：

//执行kernel,Range用1维，work itmes size为BUFSIZE,
cl_event ev;
size_t global_work_size = BUFSIZE;

clTimer.Reset();
clTimer.Start();
clEnqueueNDRangeKernel( queue,
    kernel,
    1,
    NULL,
    &global_work_size,
    NULL, 0, NULL, &ev);
status = clFlush( queue );
waitForEventAndRelease(&ev);
 //clWaitForEvents(1, &ev);

clTimer.Stop();
printf("kernal total time:%.6f ms \n ", clTimer.GetElapsedTime()*1000 );

//数据拷回host内存
cl_float *ptr;
clTimer.Reset();
clTimer.Start();
cl_event mapevt;
ptr = (cl_float *) clEnqueueMapBuffer( queue,
    buffer,
    CL_TRUE,
    CL_MAP_READ,
    0,
    BUFSIZE * sizeof(cl_float),
    0, NULL, &mapevt, NULL );
status = clFlush( queue );
waitForEventAndRelease(&mapevt);
 //clWaitForEvents(1, &mapevt);

clTimer.Stop();
printf("copy from device to host:%.6f ms \n ", clTimer.GetElapsedTime()*1000 );

最终程序执行界面如下，在bufsize为262144时，在我的显卡上gpu还没有cpu快，在程序目录，我们可以看到也产生了vecadd.bin文件了。

完整的代码请参考：

工程文件gclTutorial2

代码下载：

https://files.cnblogs.com/mikewolf2002/gclTutorial.zip