OpenCL简介

    开放计算语言（Open Computing Language, OpenCL）是非盈利技术联盟Khronos Group管理的异构编程框架。该框架充分利用了CPU、DSP、FPGA、GPU的计算能力。OpenCL支持多层次的并行，可以高效的映射到同构或异构的体系结构上。

OpenCL标准

    OpenCL API是按照C 的，由C和C++封装而成，并且有很多第三方语言的绑定。这些语言包括Java、Python以及.NET等。OpenCL是C99语言的子集，并适当地扩展了在众多异构设备上执行数据并行代码的能力。     OpenCL和许多CPU并发编程模型一样，语法类似于标准的C函数，主要区别在于它拥有额外的一些关键字和OpenCL Kernel实现的执行模型。在OpenCL编程中，编程人员应该考虑的是如何细粒度地表示程序中的并行性。一个典型的OpenCL Kernel应该是这样的： _kernel void vecadd(_global int *C, _global int *A, _global int *C) { int tid = get_global_id(0); //OpenCL intrisic function C[tid] = A[tid] + B[tid]; }

OpenCL规范

    OpenCL规范由四个模块组成：

1. 平台模型：定义了一个抽象的硬件模型，描述了宿主机和设备。供编程人员在上面编写在这些设备上执行的Kernel。
2. 执行模型：定义了如何在宿主机上配置OpenCL环境以及如何在设备上执行Kernel。
3. 内存模型：定义了Kernel中所使用的内存层次，无需考虑实际的底层架构。
4. 编程模型：定义了如何将并发模型映射到物理硬件上。

平台模型

    如下图描述，平台模型由一个宿主机以及一个或多个设备组成，一个设备可以分为多个CUs(Compute Units)，一个CU又可以进一步划分为多个PEs(Processing Elements)。设备上的计算具体映射到了每个PE之上。
    在平台模型中编程人员需要关注三种版本号：平台版本号、设备版本号、OpenCL支持版本号。     平台版本号表明了平台所支持的运行时功能，包括了OpenCL运行时所有与宿主机交互的APIs，例如上下文、内存对象、设备、命令队列相关的函数。     设备版本号表明了设备的能力，可以使用clGetDeviceInfo函数获取。设备信息通常包括资源限制和扩展功能等。     语言版本号表明了设备最高支持的语言版本。

执行模型

    OpenCL的执行模型分为 两部分，Kernel Program和Host Program，宿主程序定义了Kernel执行的上下文并控制其执行。执行模型中主要定义了kernel是怎样执行的。     这里首先讲一下执行模型中非常重要的两个概念——work-item和work-groups。在kernel由宿主机提交给设备执行的同时，定义了一个索引空间。每个work-item可以看做是kernel在设备上单次执行的实例，该实例在索引空间中拥有唯一的索引值。基于索引值的不同，每个work-item虽然执行同样的kernel程序，但却有自己的代码路径和数据路径。work-item被组织为work-groups。在同一work-groups中每个work-item拥有唯一的local ID，如果有多个work-groups，那么在不同的work-groups之间local ID不是唯一的。一个work-item索引必须采用global ID或local ID + work-groups ID才能唯一确定。          在执行模型中另一个非常重要的概念是NDRange，即N维索引空间，这里的N可以定义为1,2或3。它由一个长度为N的数组组成，每个数组元素表示该维度上工作节点的个数。工作节点的global ID、local ID以及work-groups ID都是N维的。上图是一个二维索引空间的实例，NDRange(Gx,Gy)，对应的工作组为NDRange(Sx,Sy)。每个维度的工作空间恰好分配为整数个工作组。这里要求Gx必须为Sx的整数倍，Gy必须为Sy的整数倍。

上下文与命令队列

    在OpenCL中，上下文（context）是一个存在于主机端的抽象容器。负责协调主机——设备之间的交互，管理设备上可用的内存对象，跟踪每个设备建立的kernel和程序。     宿主机负责定义kernel的上下文，上下文包括以下资源：

1. 设备：平台上可用的设备。
2. Kernels：运行在设备上的OpenCL C程序。
3. 程序对象：负责执行和实现Kernels的程序资源。
4. 内存对象：Kernel和宿主机可见的内存对象的集合，内存对象中包含kernel实例中用于计算的数据。

    新建上下文的API为clCreateContext()。 cl_context clCreateContext(const cl_context_properties *properties, cl_unit num_devices, void (CL_CALLBACK *pfn_notify)( const char *errinfo, const void *private_info, size_t cb, void *user_data), void *user_data, cl_int *errcode_ret)    API函数clCreateCommandQueue()用来建立命令队列并将其关联到某个设备。cl_command_queue clCreateCommandQueue(cl_context context cl_device_id device, cl_command_queue_properties properties, cl_int *errcode_ret)
    clCreateCommandQueue()中的属性参数properties是一个位域，用于开启程序剖析（profiling）命令（CL_QUEUE_PROFILING_ENABLE）、乱序执行命令（CL_QUEUE_OUT_OF_ORDER_EXEC_ENABLE）。

内存对象

    OpenCL应用程序用来处理大型数组或多维矩阵，在kernel开始执行之前需要将这些数据实际映射到设备上。在OpenCL中采用的方法是将这些数据封装为内存对象。OpenCL定义了两种内存对象：buffer和image。buffer类似于C语言中的数组，由malloc()分配，在内存中是连续存放的。而image被设计为不透明的对象，以利于进行数据填充和在特定的设备上进行性能优化。     创建buffer的内存对象的API为clCreateBuffer()，内存对象作为返回值返回。 cl_mem clCreateBuffer(cl_context context, cl_mem_flags flags, size_t size, void *host_ptr, cl_int *errcode_ret)    将宿主机内存传入或传出设备的API分别为clEnqueueWriteBuffer和clEnqueueReadBuffer。可通过设置参数选项blocking_write为CL_TRUE，使写函数工作在阻塞模式——数据完成到OpenCL buffer的传输后才返回。     image对象在后面用到的时候在详细介绍。

程序对象

    OpenCL中的程序（program）特指kernel函数的集合，是kernel被调度安排到设备上运行的单位。     OpenCL程序运行时通过调用一系列。API进行编译，编译系统对具体设备进行优化。OpenCL软件仅链接到公共运行层（称为ICD），所有平台特定的SDK通过一个动态链接接口委托给某个厂商的运行时。     新建kernel的步骤如下：

1. OpenCL源代码以字符串形式存储。如果代码存放在文件中，则必须读取到内存中，在内存中以字符串数组的形式存放。
2. 源代码通过调用clCreateProgramWithSource()转化为cl_program对象。
3. 调用clBuildProgram()在支持OpenCL的设备上编译程序对象。

向量加代码演示

#include #include #include const char *programSource = "_kernel void vecadd(_global int *A, _global int *B, _global int *C) " "{ " " int idx = get_global_id(0); " " C[idx] = A[idx] + B[idx]; " "} " int main() { int *A = NULL; int *B = NULL; int *C = NULL; const int elements = 2048; size_t datasize = sizeof(int) * elements; A = (int*)malloc(datasize); B = (int*)malloc(datasize); C = (int*)malloc(datasize); int i; for(i=0; i