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

代碼下載：

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