Online/Offline 编译

最新推荐文章于 2021-07-07 11:04:15 发布

乱战狂歌

最新推荐文章于 2021-07-07 11:04:15 发布

阅读量1.8k

点赞数

分类专栏： OpenCL 文章标签： OpenCL 并行翻译

OpenCL 专栏收录该内容

4 篇文章 2 订阅

订阅专栏

很久之前翻译的一章OpenCL书籍，应该是本人的第一次翻译，水平够烂：

【原文】The OpenCL Programming Book 4.2 Online/Offline Compilation

在OpenCL中一个kernel可以通过在线和下线两种方式来进行编译

这两种方式的主要区别如下：

offline：二进制kernel由主机代码读取
online：kernel的代码文件由主机进行读取
PS:用在线和下线翻译很蛋疼，就直接用on/off了

offline编译时，kernel由一个OpenCL内核编译器进行预编译，并且生成可以用OpenCL API加载的二进制码。因为二进制kernel已经编译过了，所以主机代码开始执行到到内核代码执行之间的时间可以忽略。使用这种方法的问题是：为了在不同的平台上来执行程序，不同的的二进制kernel必须被包括在内，从而增加了可执行文件的大小。

online编译时，kernel是运行时使用OpenCL运行库创建的。这种方法被一般认为是JIL（Just in time，即时的）编译。这种方法的优势在于主机端的二进制可以以一种不依赖于设备的形式分布，并且改写可以改写kernel。这也使得在开发中更容易测试内核，因为它摆脱了每次都要编译内核的需要。然而这种方式不适合在需要实时处理的嵌入式系统。而且，由于代码是一种可读的形式，这种方法可能不适合商业应用。

OpenCL的运行库包含了一系列的关于上述操作的API。在某种程度上，因为OpenCL是一个用于各种环境的编程框架，支持online操作是正常的，实际上，一个单独的OpenCL编译器不可能同时支持nvidia，AMD和苹果的OpenCL环境。因此，为了在这些环境下产生二进制kernel，在主程序运行时，编译过的kernel必须在这过程写入一个文件。FOXC在另一方面包括一个独立的OpenCL内核编译器，使得生成二进制kernel的过程更直观（这句只是宣传这个公司的工具FOXC的）。

Online Compilation version：

#include <stdio.h>   
#include <stdlib.h>  
     
#ifdef __APPLE__   
#include <OpenCL/opencl.h>   
#else  
#include <CL/cl.h>   
#endif 
     
#define MEM_SIZE (128) 
#define MAX_SOURCE_SIZE (0x100000) 
     
int main() 
{  
    cl_platform_id platform_id = NULL;
    cl_device_id device_id = NULL;
    cl_context context = NULL;
    cl_command_queue command_queue = NULL;
    cl_mem memobj = NULL;
    cl_program program = NULL;
    cl_kernel kernel = NULL;
    cl_uint ret_num_devices;
    cl_uint ret_num_platforms;
    cl_int ret;
     
    float mem[MEM_SIZE];
     
    FILE *fp;
    const char fileName[] = "./kernel.cl";
    size_t source_size;
    char *source_str;
    cl_int i;
     
    /* Load kernel source code */
    fp = fopen(fileName, "r");
    if (!fp) {
     
     
    }
    source_str = (char *)malloc(MAX_SOURCE_SIZE);
    source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
    fclose(fp);
     
    /*Initialize Data */
    for (i = 0; i < MEM_SIZE; i++) {
     
    }
     
    /* Get platform/device information */
    ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
    ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);
     
    /* Create OpenCL Context */
    context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
     
    /* Create Command Queue */
    command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
     
    /* Create memory buffer*/
    memobj = clCreateBuffer(context, CL_MEM_READ_WRITE, MEM_SIZE * sizeof(float), NULL, &ret);
     
    /* Transfer data to memory buffer */
    ret = clEnqueueWriteBuffer(command_queue, memobj, CL_TRUE, 0, MEM_SIZE * sizeof(float), mem, 0, NULL, NULL);
     
    /* Create Kernel program from the read in source */
    program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);
     
    /* Build Kernel Program */
    ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
     
    /* Create OpenCL Kernel */
    kernel = clCreateKernel(program, "vecAdd", &ret);
     
    /* Set OpenCL kernel argument */
    ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj);
     
    size_t global_work_size[3] = {MEM_SIZE, 0, 0};
    size_t local_work_size[3] = {MEM_SIZE, 0, 0};
     
    /* Execute OpenCL kernel */
    ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
     
    /* Transfer result from the memory buffer */
    ret = clEnqueueReadBuffer(command_queue, memobj, CL_TRUE, 0, MEM_SIZE * sizeof(float), mem, 0, NULL, NULL);
     
    /* Display result */
    for (i=0; i < MEM_SIZE; i++) {
     
    }
     
    /* Finalization */
    ret = clFlush(command_queue);
    ret = clFinish(command_queue);
    ret = clReleaseKernel(kernel);
    ret = clReleaseProgram(program);
    ret = clReleaseMemObject(memobj);
    ret = clReleaseCommandQueue(command_queue);
    ret = clReleaseContext(context);
     
    free(source_str);
     
    return 0;
    }

Offline compilation version：

#include <stdio.h>   
#include <stdlib.h>  
     
#ifdef __APPLE__   
#include <OpenCL/opencl.h>   
#else  
#include <CL/cl.h>   
#endif 
     
#define MEM_SIZE (128) 
#define MAX_BINARY_SIZE (0x100000) 
     
int main() 
{  
    cl_platform_id platform_id = NULL;
    cl_device_id device_id = NULL;
    cl_context context = NULL;
    cl_command_queue command_queue = NULL;
    cl_mem memobj = NULL;
    cl_program program = NULL;
    cl_kernel kernel = NULL;
    cl_uint ret_num_devices;
    cl_uint ret_num_platforms;
    cl_int ret;
     
    float mem[MEM_SIZE];
     
    FILE *fp;
    char fileName[] = "./kernel.clbin";
    size_t binary_size;
    char *binary_buf;
    cl_int binary_status;
    cl_int i;
     
    /* Load kernel binary */
    fp = fopen(fileName, "r");
    if (!fp) {
     
     
    }
    binary_buf = (char *)malloc(MAX_BINARY_SIZE);
    binary_size = fread(binary_buf, 1, MAX_BINARY_SIZE, fp);
    fclose(fp);
     
    /* Initialize input data */
    for (i = 0; i < MEM_SIZE; i++) {
     
    }
     
    /* Get platform/device information */
    ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
    ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);
     
    /* Create OpenCL context*/
    context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
     
    /* Create command queue */
    command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
     
    /* Create memory buffer */
    memobj = clCreateBuffer(context, CL_MEM_READ_WRITE, MEM_SIZE * sizeof(float), NULL, &ret);
     
    /* Transfer data over to the memory buffer */
    ret = clEnqueueWriteBuffer(command_queue, memobj, CL_TRUE, 0, MEM_SIZE * sizeof(float), mem, 0, NULL, NULL);
     
    /* Create kernel program from the kernel binary */
    program = clCreateProgramWithBinary(context, 1, &device_id, (const size_t *)&binary_size,
    (const unsigned char **)&binary_buf, &binary_status, &ret);
     
    /* Create OpenCL kernel */
    kernel = clCreateKernel(program, "vecAdd", &ret);
    printf("err:%d\n", ret);
     
    /* Set OpenCL kernel arguments */
    ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj);
     
    size_t global_work_size[3] = {MEM_SIZE, 0, 0};
    size_t local_work_size[3] = {MEM_SIZE, 0, 0};
     
    /* Execute OpenCL kernel */
    ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
     
    /* Copy result from the memory buffer */
    ret = clEnqueueReadBuffer(command_queue, memobj, CL_TRUE, 0, MEM_SIZE * sizeof(float), mem, 0, NULL, NULL);
     
    /* Display results */
    for (i=0; i < MEM_SIZE; i++) {
     
    }
     
    /* Finalization */
    ret = clFlush(command_queue);
    ret = clFinish(command_queue);
    ret = clReleaseKernel(kernel);
    ret = clReleaseProgram(program);
    ret = clReleaseMemObject(memobj);
    ret = clReleaseCommandQueue(command_queue);
    ret = clReleaseContext(context);
     
    free(binary_buf);
     
    return 0;
    }

Kernel program：

__kernel void vecAdd(__global float* a)
{
    int gid = get_global_id(0);
    a[gid] += a[gid];
}

上诉两段代码几乎一样，所以我们主要关注不同点。

最主要的不同点是下面这段代码使用的是online编译模式。

035: fp = fopen(fileName, "r");
036: if (!fp) {
037: fprintf(stderr, "Failed to load kernel.\n");
038: exit(1);
039: }
040: source_str = (char *)malloc(MAX_SOURCE_SIZE);
041: source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
042: fclose(fp);

sourc_str变量只是一个仅仅包含源文件的数组，为了了执行kernel中的代码，它不需要需要用运行时（runtime）编译器。这段代码如下：

065: /* Create kernel program from the source */
066: program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);
067:
068: /* Build kernel program */
069: ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);

这段代码先由源文件创建，然后再编译。

现在来看offline编译的版本：

Reading the kernel binary：
036: fp = fopen(fileName, "r");
037: if (!fp) {
038: fprintf(stderr, "Failed to load kernel.\n");
039: exit(1);
040: }
041: binary_buf = (char *)malloc(MAX_BINARY_SIZE);
042: binary_size = fread(binary_buf, 1, MAX_BINARY_SIZE, fp);
043: fclose(fp);

上面的代码看起来很像在线版本，因为数据也是被读入到char型的缓存中。不同带的是缓存数据是可以直接执行的。这意味着这个kernel代码之前已经用OpenCL编译器编译过了。

（下面的话无视就好，为了尊重原作者才加入的）
用FOXC，这个操作可以这样进行
> /path-to-foxc/bin/foxc -o kernel.clbin kernel.cl

online编译要经过两个步骤来建立内核程序。在offline编译过程中， clCreateProgramWithSource被clCreateProgramWithBinary所取代。

067: program = clCreateProgramWithBinary(context, 1, &device_id, (const size_t *)&binary_size,
068: (const unsigned char **)&binary_buf, &binary_status, &ret);

由于kernel已经编译过了，所以没有必要向online模式一样再进行下一步编译。

总的来说，为了改变从online到offline改变编译的方式，主要的步骤是：

1.读二进制kernel
2.把clCreateProgramWithSource() 改成 clCreateProgramWithBinary()
3.去掉clBuildProgram()

这就是是两种方式的主要差别。

可以通过查看手册开了解相关API的具体细节。

（完）

乱战狂歌

关注

0
点赞
踩
2

收藏

觉得还不错? 一键收藏
0
评论
Online/Offline 编译

很久之前翻译的一章OpenCL书籍，应该是本人的第一次翻译，水平够烂：【原文】The OpenCL Programming Book 4.2 Online/Offline Compilation在OpenCL中一个kernel可以通过在线和下线两种方式来进行编译这两种方式的主要区别如下：offline：二进制kernel由主机代码读取online：kernel的
复制链接

扫一扫