一文說清 OpenCL 框架

背景

說明：

  1. 介紹 =====

    1）用於編寫運行在 OpenCL device 上的kernels的語言（基於 C99）；

    2）OpenCL API，至於 Runtime 的實現交由各個廠家，比如 Intel 發佈的opencl_runtime_16.1.2_x64_rh_6.4.0.37.tgz

 以人工智能場景爲例來理解一下，假如在某個 AI 芯片上跑人臉識別應用，CPU 擅長控制，AI processor 擅長計算，軟件的 flow 就可以進行拆分，用 CPU 來負責控制視頻流輸入輸出前後處理，AI processor 來完成深度學習模型運算完成識別，這就是一個典型的異構處理場景，如果該 AI 芯片的 SDK 支持 OpenCL，那麼上層的軟件就可以基於 OpenCL 進行開發了。

話不多說，看看 OpenCL 的架構吧。

  1. OpenCL 架構 ============

OpenCL 架構，可以從平臺模型、內存模型、執行模型、編程模型四個角度來展開。

2.1 Platform Model

平臺模型：硬件拓撲關係的抽象描述

2.2 Execution Model

執行模型：Host 如何利用 OpenCL Device 的計算資源完成高效的計算處理過程

Context

OpenCL 的 Execution Model 由兩個不同的執行單元定義：1）運行在 OpenCL 設備上的 kernel；2）運行在 Host 上的 Host program；其中，OpenCL 使用 Context 代表 kernel 的執行環境：

Context 包含以下資源：

NDrange

有兩種方式來找到 work-item：

  1. 通過 work-item 的全局索引；

  2. 先查找到所在 work-group 的索引號，再根據局部索引號確定；

以一維爲例：

以二維爲例：

三維的方式也類似，略去。

2.3 Memory Model

內存模型：Host 和 OpenCL Device 怎麼來看待數據

OpenCL 的內存模型中，包含以下幾類類型的內存：

2.4 Programming Model

  1. 編程流程 =======

下邊來一個實際的代碼測試跑跑，Talk is cheap, show me the code!

  1. 示例代碼 =======

4.1 Host 端程序

#include <iostream>
#include <fstream>
#include <sstream>

#include <CL/cl.h>

const int DATA_SIZE = 10;

int main(void)
{
    /* 1. get platform & device information */
    cl_uint num_platforms;
    cl_platform_id first_platform_id;
    clGetPlatformIDs(1, &first_platform_id, &num_platforms);


    /* 2. create context */
    cl_int err_num;
    cl_context context = nullptr;
    cl_context_properties context_prop[] = {
        CL_CONTEXT_PLATFORM,
        (cl_context_properties)first_platform_id,
        0
    };
    context = clCreateContextFromType(context_prop, CL_DEVICE_TYPE_CPU, nullptr, nullptr, &err_num);


    /* 3. create command queue */
    cl_command_queue command_queue;
    cl_device_id *devices;
    size_t device_buffer_size = -1;

    clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, nullptr, &device_buffer_size);
    devices = new cl_device_id[device_buffer_size / sizeof(cl_device_id)];
    clGetContextInfo(context, CL_CONTEXT_DEVICES, device_buffer_size, devices, nullptr);
    command_queue = clCreateCommandQueueWithProperties(context, devices[0], nullptr, nullptr);
    delete [] devices;


    /* 4. create program */
    std::ifstream kernel_file("vector_add.cl", std::ios::in);
    std::ostringstream oss;

    oss << kernel_file.rdbuf();
    std::string srcStdStr = oss.str();
    const char *srcStr = srcStdStr.c_str();
    cl_program program;
    program = clCreateProgramWithSource(context, 1, (const char **)&srcStr, nullptr, nullptr);


    /* 5. build program */
    clBuildProgram(program, 0, nullptr, nullptr, nullptr, nullptr);


    /* 6. create kernel */
    cl_kernel kernel;
    kernel = clCreateKernel(program, "vector_add", nullptr);


    /* 7. set input data && create memory object */
    float output[DATA_SIZE];
    float input_x[DATA_SIZE];
    float input_y[DATA_SIZE];
    for (int i = 0; i < DATA_SIZE; i++) {
        input_x[i] = (float)i;
        input_y[i] = (float)(2 * i);
    }

    cl_mem mem_object_x;
    cl_mem mem_object_y;
    cl_mem mem_object_output;
    mem_object_x = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * DATA_SIZE, input_x, nullptr);
    mem_object_y = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * DATA_SIZE, input_y, nullptr);
    mem_object_output = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * DATA_SIZE, nullptr, nullptr);


    /* 8. set kernel argument */
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &mem_object_x);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &mem_object_y);
    clSetKernelArg(kernel, 2, sizeof(cl_mem), &mem_object_output);


    /* 9. send kernel to execute */
    size_t globalWorkSize[1] = {DATA_SIZE};
    size_t localWorkSize[1] = {1};
    clEnqueueNDRangeKernel(command_queue, kernel, 1, nullptr, globalWorkSize, localWorkSize, 0, nullptr, nullptr);


    /* 10. read data from output */
    clEnqueueReadBuffer(command_queue, mem_object_output, CL_TRUE, 0, DATA_SIZE * sizeof(float), output, 0, nullptr, nullptr);
    for (int i = 0; i < DATA_SIZE; i++) {
        std::cout << output[i] << " ";
    }
    std::cout << std::endl;


    /* 11. clean up */
    clRetainMemObject(mem_object_x);
    clRetainMemObject(mem_object_y);
    clRetainMemObject(mem_object_output);
    clReleaseCommandQueue(command_queue);
    clReleaseKernel(kernel);
    clReleaseProgram(program);
    clReleaseContext(context);

    return 0;
}

4.2 OpenCL Kernel 函數

內容如下：

__kernel void vector_add(__global const float *input_x,
 __global const float *input_y,
 __global float *output)
{
 int gid = get_global_id(0);
 
 output[gid] = input_x[gid] + input_y[gid];
}

4.3 輸出

參考

The OpenCL Specification

本文由 Readfog 進行 AMP 轉碼，版權歸原作者所有。
來源：https://mp.weixin.qq.com/s/XE-A4jnNJdjH_RyQb0ttzQ

猜你喜歡