学习OpenCL :一种异构计算架构

前言：
在过去利用GPU对图像渲染进行加速的技术非常成熟，因为GPU是典型的单指令多数据（SIMD）的体系结构，擅长大规模的并行计算；而CPU是多指令单数据流（MISD）的体系结构，更擅长逻辑控制。

在当今数据量计算越发庞大的情况下，为了提升计算效率，人们希望将GPU大规模的并行计算的能力扩展到更多领域，而不只局限与图像渲染。这样，CPU只负责逻辑控制，GPU更多负责计算，这种一个CPU（控制单元）+几个GPU（有时可能再加几个CPU）（计算单元）的架构就是所谓的异构编程。

OpenCL就是这种情况下出现的，它是一种异构计算的标准，可以用来针对GPU编程。其实在OpenCL出来之前，NVIDIA就推出了GPGPU计算CUDA架构。只不过CUDA只能使用自家的N卡，对其他显卡不支持，而OpenCL则是一个通用的标准，对A卡，N卡等都支持，还支持CPU计算。
————————————————————————————————————————————

一、GPU编程 : GPU和CPU的设计区别

1.1 CPU简介

CPU (Central Processing Unit) 即中央处理器，是机器的“大脑”，也是布局谋略、发号施令、控制行动的“总司令官”。

CPU的结构主要包括运算器（ALU, Arithmetic and Logic Unit）、控制单元（CU, Control Unit）、寄存器（Register）、高速缓存器（Cache）和它们之间通讯的数据、控制及状态的总线。

CPU是基于低延时的设计，简单来说包括：计算单元、控制单元和存储单元，架构可参考下图：

CPU的特点是：

（1）CPU有强大的ALU（算术运算单元），它可以在很少的时钟周期内完成算术计算

当今的CPU可以达到64bit 双精度。执行双精度浮点源算的加法和乘法只需要1～3个时钟周期（CPU的时钟周期的频率是非常高的，达到1.532～3gigahertz(千兆HZ, 10的9次方)）。

（2）大的缓存可以降低延时

保存很多的数据放在缓存里面，当需要访问的这些数据，只要在之前访问过的，如今直接在缓存里面取即可。

（3）复杂的逻辑控制单元

当程序含有多个分支的时候，它通过提供分支预测的能力来降低延时。
数据转发。当一些指令依赖前面的指令结果时，数据转发的逻辑控制单元决定这些指令在pipeline中的位置并且尽可能快的转发一个指令的结果给后续的指令。这些动作需要很多的对比电路单元和转发电路单元。

————————————————————————————————————————————

1.2 GPU简介

GPU全称为Graphics Processing Unit，中文为图形处理器，就如它的名字一样，GPU最初是用在个人电脑、工作站、游戏机和一些移动设备（如平板电脑、智能手机等）上运行绘图运算工作的微处理器。

GPU是基于大的吞吐量设计，GPU简单架构参考下图:

GPU的特点是：

（1）有很多的ALU和很少的cache

缓存的目的不是保存后面需要访问的数据的，这点和CPU不同，而是为thread提高服务的。如果有很多线程需要访问同一个相同的数据，缓存会合并这些访问，然后再去访问dram（因为需要访问的数据保存在dram中而不是cache里面），获取数据后cache会转发这个数据给对应的线程，这个时候是数据转发的角色。但是由于需要访问dram，自然会带来延时的问题。

（2）GPU的控制单元（左边黄色区域块）可以把多个的访问合并成少的访问。

（3）GPU的虽然有dram延时，却有非常多的ALU和非常多的thread

为了平衡内存延时的问题，GPU可以充分利用多的ALU的特性达到一个非常大的吞吐量的效果。尽可能多的分配Threads。
————————————————————————————————————————————

1.3 并行计算

在对比CPU/GPU前，先了解下并行并行计算。

并行计算(Parallel Computing)是指同时使用多种计算资源解决计算问题的过程，是提高计算机系统计算速度和处理能力的一种有效手段。它的基本思想是用多个处理器来共同求解同一问题，即将被求解的问题分解成若干个部分，各部分均由一个独立的处理机来并行计算。

并行计算可分为时间上的并行和空间上的并行。

时间上的并行是指流水线技术，比如说工厂生产食品的时候分为四步：清洗-消毒-切割-包装。

如果不采用流水线，一个食品完成上述四个步骤后，下一个食品才进行处理，耗时且影响效率。但是采用流水线技术，就可以同时处理四个食品。这就是并行算法中的时间并行，在同一时间启动两个或两个以上的操作，大大提高计算性能。

空间上的并行是指多个处理机并发的执行计算，即通过网络将两个以上的处理机连接起来，达到同时计算同一个任务的不同部分，或者单个处理机无法解决的大型问题。

比如小李准备在植树节种三棵树，如果小李1个人需要6个小时才能完成任务，植树节当天他叫来了好朋友小红、小王，三个人同时开始挖坑植树，2个小时后每个人都完成了一颗植树任务，这就是并行算法中的空间并行，将一个大任务分割成多个相同的子任务，来加快问题解决速度。
————————————————————————————————————————————

1.4 CPU/GPU对比

CPU：

CPU的架构中需要大量的空间去放置存储单元（橙色部分）和控制单元（黄色部分），相比之下计算单元（绿色部分）只占据了很小的一部分，所以它在大规模并行计算能力上极受限制，而更擅长于逻辑控制。

CPU遵循的是冯诺依曼架构，其核心就是：存储程序，顺序执行。这使得CPU就像是个一板一眼的管家，人们吩咐的事情它总是一步一步来做。但是随着人们对更大规模与更快处理速度的需求的增加，这位管家渐渐变得有些力不从心。

GPU:

GPU的构成相对简单，有数量众多的计算单元和超长的流水线，特别适合处理大量的类型统一的数据。

GPU的工作大部分都计算量大，但没什么技术含量，而且要重复很多很多次。GPU就是用很多简单的计算单元去完成大量的计算任务，纯粹的人海战术。这种策略基于一个前提，就是并行计算的线程之间没有什么依赖性，是互相独立的。

GPU在处理能力和存储器带宽上相对于CPU有明显优势，在成本和功耗上也不需要付出太大代价。由于图形渲染的高度并行性，使得GPU可以通过增加并行处理单元和存储器控制单元的方式提高处理能力和存储器带宽。GPU设计者将更多的晶体管用作执行单元，而不是像CPU那样用作复杂的控制单元和缓存并以此来提高少量执行单元的执行效率。

但GPU无法单独工作，必须由CPU进行控制调用才能工作。CPU可单独作用，处理复杂的逻辑运算和不同的数据类型，但当需要大量的处理类型统一的数据时，则可调用GPU进行并行计算。

另外：CPU的整数计算、分支、逻辑判断和浮点运算分别由不同的运算单元执行，此外还有一个浮点加速器。因此，CPU面对不同类型的计算任务会有不同的性能表现。而GPU是由同一个运算单元执行整数和浮点计算，因此，GPU的整型计算能力与其浮点能力相似。
————————————————————————————————————————————

1.5 适于GPU计算的场景

尽管GPU计算已经开始崭露头角，但GPU并不能完全替代X86解决方案。很多操作系统、软件以及部分代码现在还不能运行在GPU上，所谓的GPU+CPU异构超级计算机也并不是完全基于GPU进行计算。一般而言适合GPU运算的应用有如下特征：

• 运算密集
• 高度并行
• 控制简单
• 分多个阶段执行

GPU计算的优势是大量内核的并行计算，瓶颈往往是I/O带宽，因此适用于计算密集型的计算任务。
————————————————————————————————————————————

1.6 GPU开发环境

CG(C for Graphics) 是为GPU编程设计的高级绘制语言，由NVIDIA和微软联合开发，微软版本叫HLSL，CG是NVIDIA版本。Cg极力保留C语言的大部分语义，并让开发者从硬件细节中解脱出来，Cg同时也有一个高级语言的其他好处，如代码的易重用性，可读性得到提高，编译器代码优化。

CUDA（ComputeUnified DeviceArchitecture，统一计算架构） 是由NVIDIA所推出的一种集成技术，是该公司对于GPGPU的正式名称。通过这个技术，用户可利用NVIDIA的GeForce8以后的GPU和较新的QuadroGPU进行计算。亦是首次可以利用GPU作为C-编译器的开发环境。NVIDIA营销的时候，往往将编译器与架构混合推广，造成混乱。实际上，CUDA架构可以兼容OpenCL或者自家的C-编译器。无论是CUDAC-语言或是OpenCL，指令最终都会被驱动程序转换成PTX代码，交由显示核心计算。

ATIStream是AMD针对旗下图形处理器（GPU）所推出的通用并行计算技术。利用这种技术可以充分发挥AMDGPU的并行运算能力，用于对软件进行加速或进行大型的科学运算，同时用以对抗竞争对手的NVIDIA CUDA技术。与CUDA技术是基于自身的私有标准不同，ATIStream技术基于开放性的OpenCL标准。

OpenCL(Open Computing Language，开放计算语言) 是一个为异构平台编写程序的框架，此异构平台可由CPU，GPU或其他类型的处理器组成。OpenCL由一门用于编写kernels（在OpenCL设备上运行的函数）的语言（基于C99）和一组用于定义并控制平台的API组成。OpenCL提供了基于任务分区和数据分区的并行计算机制。

OpenCL类似于另外两个开放的工业标准OpenGL和OpenAL，这两个标准分别用于三维图形和计算机音频方面。OpenCL扩展了GPU用于图形生成之外的能力。OpenCL由非盈利性技术组织KhronosGroup掌管。

下面是对几种GPU开发环境的简单评价：

CG：优秀的图形学学开发环境，但不适于GPU通用计算开发。

ATIStream：硬件上已经有了基础，但只有低层次汇编才能使用所有的硬件资源。高层次的brook是基于上一代GPU的，缺乏良好的编程模型。

OpenCL：开放标准，抽象层次较低，较多对硬件的直接操作，代码需要根据不同硬件优化。

CUDA：仅能用于NVIDIA的产品，发展相对成熟，效率高，拥有丰富的文档资源。

————————————————————————————————————————————

二、OpenCL学习之路

2.1 OpenCL是什么，为什么需要OpenCL

OpenCL(全称为Open Computing Langugae，开放运算语言)是第一个面向异构系统(此系统中可由CPU，GPU或其它类型的处理器架构组成)的并行编程的开放式标准。它是跨平台的。

OpenCL由两部分组成，一是用于编写kernels（在OpenCL设备上运行的函数）的语言，二是用于定义并控制平台的API(函数)。OpenCL提供了基于任务和基于数据两种并行计算机制，它极大地扩展了GPU的应用范围，使之不再局限于图形领域。

OpenCL是一种标准，intel、Nvidia、ARM、AMD、QUALCOMM、Apple都有其对应的OpenCL实现。像NVDIA将OpenCL实现集成到它的CUDA SDK中，而AMD则将其实现后放在AMD APP （Accelerated Paral Processing）SDK中…

在过去的几十年里，计算机产业发生了巨大的变化。计算机性能 的不断提高为当前各种应用提供了有力的保障。对于计算机的速度 而言，正如摩尔定律描述的那样，是通过晶体管数目增加来提高频率 的方式实现的。但是到了二十一世纪初期以后，这种增长方式受到 了一些限制，晶体管尺寸变得已经很小，其物理特性决定了很难再通 过大规模地增加晶体管的数目来提升频率，且由于功耗也以非线性 的速度增加，因此这种方式受到很大的限制。在未来，这一趋势会继续 成为影响计算机系统最为重要的因素之一。

为了解决这一问题通常有两种方式，第一种是通过 增加处理器的核心数目来为多任务，多线程等提供支持，从整体 上提升系统的性能。第二种方式是通过异构的方式，例如可 利用CPU（Central Processing Unit）、GPU（Graphic Processing Unit）、甚至APU(Accelerated Processing Units， CPU与GPU的融合)等计算设备的计算能力从而 来既提高系统的速度。

异构系统越来越普遍，对于支持这种环境的计算而言，也正受到越来越多 的关注。当前，不同厂商通常仅仅提供对于自己设备编程的实现。对于异 构系统一般很难用同种风格的编程语言来实现机构编程，而且将不同的设备 作为统一的计算单元来处理的难度也是非常大的。

开放式计算语言（Open Computing Language:OpenCL），旨在满足这一重要需求。 通过定义一套机制，来实现硬件独立的软件开发环境。利用OpenCL可以充分利 用设备的并行特性，支持不同级别的并行，并且能有效映射到由CPU，GPU， FPGA（Field－Programmable Gate Array）和将来出现的设备 所组成的同构或异构，单设备或多设备的系统。OpenCL定义了运行时， 允许用来管理资源，将不同类型的硬件结合在同种执行环境中，并且很有希望 在不久的将来，以更加自然的方式支持动态地平衡计算，功耗和其他资源。

我相信在不久的将来，OpenCL将在异构并行编程中得到广泛的应用。

————————————————————————————————————————————

2.2 OpenCL架构

这些概念可能会比较难理解，没关系，后续看了一些相关的例子应该就容易理解了。

• OpenCL平台API： 平台API定义了宿主机程序发现OpenCL设备所用的函数以及这些函数的功能，另外还定义为OpenCL应用创建上下文(上下文表示的是程序运行时所拥有的所有软硬件资源+内存+处理器)的函数。这里的平台指的是宿主机、OpenCL设备和OpenCL框架的组合。
• OpenCL运行时API： 平台API主要用来创建上下文，运行时API则强调使用这个上下文满足应用需求的函数集，用来管理上下文来创建命令队列以及运行时发生的其它操作。例如，将命令提交到命令队列的函数。
• OpenCL编程语言： 用来编写内核代码的编程语言，基于ISO C99标准的一个扩展子集，通常称为OpenCL C编程语言。

————————————————————————————————————————————

2.3 基本概念及模型

OpenCL程序同CUDA程序一样，也是分为两部分，一部分是在主机（以CPU为核心）上运行，一部分是在设备（以GPU为核心）上运行。设备有一个或多个计算单元，计算单元又包含一个或多个处理单元。在设备上运行的程序被称为核函数。但是对于核函数的编写，CUDA一般直接写在程序内，OpenCL是写在一个独立的文件中，并且文件后缀是.cl ，由主机代码读入后执行，这一点OpenCL跟OpenGL中的渲染程序很像。

————————————————————————————————————————————

2.3.1 平台模型

平台模型（如上图）指定有一个处理器（主机Host）来协调程序的执行， 一个或多个处理器（设备Devices）来执行OpenCL C代码。 在这里其实仅仅是一种抽象的硬件模型，这样就能方便程序员 编写OpenCL C函数（称之为内核）并在不同的设备上执行。
图中的设备可以被看成是CPU/GPU，而设备中的计算单元可以被看成是 CPU/GPU的核，计算单元的所有处理节点作为SIMD单元或SPMD单元（每个 处理节点维护自己的程序计数器）执行单个指令流。 抽象的平台模型更与当前的GPU的架构接近。

平台可被认为是不同厂商提供的OpenCL API的实现。如果一个平台选定之后一般只能 运行该平台所支持的设备。就当前的情况来看，如果选择了Intel的OpenCL SDK 就只能使用Intel的CPU来进行计算了，如果选择AMD的APP SDK则能进行AMD的CPU和AMD的 GPU来进行计算。一般而言，A公司的平台选定之后不能与B公司的平台进行通信。
————————————————————————————————————————————

2.3.2 执行模型

在执行模型中最重要的是内核，上下文和命令队列的概念。上下文管理多个设备， 每个设备有一个命令队列，主机程序将内核程序提交到不同的命令队列上执行。

2.3.2.1 内核

内核是执行模型的核心，能在设备上执行。当一个内核执行之前，需要指定一个 N-维的范围（NDRange）。一个NDRange是一个一维、二维或三维的索引空间。 还需要指定全局工作节点的数目，工作组中节点的数目。如下图所示， 全局工作节点的范围为{12, 12}，工作组的节点范围为{4, 4}，总共有9个工作组。

例如一个向量相加的内核程序：

__kernel void VectorAdd(__global int *A, __global int *B, __global int *C){
    int id = get_global_id(0);
    C[id]  = A[id] + B[id];
}

如果定义向量为1024维，特别地，我们可以定义全局工作节点为1024， 工作组中节点为128，则总共有8个组。定义工作组主要是为有些仅需在 组内交换数据的程序提供方便。当然工作节点数目的多少要受到设备的限制。 如果一个设备有1024个处理节点，则1024维的向量，每个节点计算一次就能完成。 而如果一个设备仅有128个处理节点，那么每个节点需要计算8次。合理设置 节点数目，工作组数目能提高程序的并行度。

2.3.2.2 上下文

一个主机要使得内核运行在设备上，必须要有一个上下文来与设备进行交互。 一个上下文就是一个抽象的容器，管理在设备上的内存对象，跟踪在设备上 创建的程序和内核。

2.3.2.3 命令队列

主机程序使用命令队列向设备提交命令，一个设备有一个命令队列，且与上下文 相关。命令队列对在设备上执行的命令进行调度。这些命令在主机程序和设备上 异步执行。执行时，命令间的关系有两种模式：(1)顺序执行，(2)乱序执行。

内核的执行和提交给一个队列的内存命令会生成事件对象。 这用来控制命令的执行、协调宿主机和设备的运行。
————————————————————————————————————————————

2.3.3 内存模型

一般而言，不同的平台之间有不同的存储系统。例如，CPU有高速缓存而GPU就没有。 为了程序的可移植性，OpenCL定义了抽象的内存模型，程序实现的时候只需关注抽 象的内存模型，具体向硬件上的映射由驱动来完成。内存空间的定义及与硬件的映 射大致如图所示。

内存空间在程序内部可以用关键字的方式指定，不同的定义与数据存在的位置 相关，主要有如下几个基本概念：

• 全局内存： 所有工作组中的所有工作项都可以对其进行读写。工作项可以 读写此中内存对象的任意元素。对全局内存的读写可能会被缓存，这取决于设备的能力。
• 不变内存： 全局内存中的一块区域，在内核的执行过程中保持不变。 宿主机负责对此中内存对象的分配和初始化。
• 局部内存： 隶属于一个工作组的内存区域。它可以用来分配一些变量， 这些变量由此工作组中的所有工作项共享。在OpenCL设备上，可能会将其实现成一块专有的内存区域，也可能将其映射到全局内存中。
• 私有内存： 隶属于一个工作项的内存区域。一个工作项的私有内存中所定义的变量对另外一个工作项来说是不可见的。
  ————————————————————————————————————————————

2.3.4 编程模型

OpenCL支持数据并行，任务并行编程，同时支持两种模式的混合。对于同步 OpenCL支持同一工作组内工作项的同步和命令队列中处于同一个上下文中的 命令的同步。

————————————————————————————————————————————

2.4 基于OpenCL的编程示例

在本小节中以图像旋转的实例，具体介绍OpenCL编程的步骤。 首先给出实现流程，然后给出实现图像旋转的C循环实现和OpenCL C kernel实现。

2.4.1 编写OpenCL程序的基本步骤

1）获取平台–>clGetPlatformIDs
2）从平台中获取设备–>clGetDeviceIDs
3）创建上下文–>clCreateContext
4）创建命令队列–>clCreateCommandQueue
5）创建缓存->clCreateBuffer
6）读取程序文件，创建程序–>clCreateProgramWithSource
7）编译程序–>clBuildProgram
8）创建内核–>clCreateKernel
9）为内核设置参数–>clSetKernelArg
10）将内核发送给命令队列，执行内核–>clEnqueueNDRangeKernel
11）获取计算结果–>clEnqueueReadBuffer
12）释放资源–>clReleaseXX**

2.4.2 图像旋转原理

图像旋转是指把定义的图像绕某一点以逆时针或顺时针方向旋转一定的角度， 通常是指绕图像的中心以逆时针方向旋转。假设图像的左上角为(l, t), 右下角为(r, b)，则图像上任意点(x, y) 绕其中心(xcenter, ycenter)逆时针旋转θ角度后， 新的坐标位置(x’,y’)的计算公式为：

x′ = (x - xcenter) cosθ - (y － ycenter) sinθ + xcenter,
y′ = (x - xcenter) sinθ + (y － ycenter) cosθ + ycenter.

C代码：

void rotate(
      unsigned char* inbuf,
      unsigned char* outbuf,
      int w, int h,
      float sinTheta,
      float cosTheta)
{
   int i, j;
   int xc = w/2;
   int yc = h/2;
   for(i = 0; i < h; i++)
   {
     for(j=0; j< w; j++)
     {
       int xpos =  (j-xc)*cosTheta - (i - yc) * sinTheta + xc;
       int ypos =  (j-xc)*sinTheta + (i - yc) * cosTheta + yc;
       if(xpos>=0&&ypos>=0&&xpos<w&&ypos<h)
          outbuf[ypos*w + xpos] = inbuf[i*w+j];
     }
   }
}

OpenCL C kernel代码：

#pragma OPENCL EXTENSION cl_amd_printf : enable
__kernel  void image_rotate(
      __global uchar * src_data,
      __global uchar * dest_data,        //Data in global memory
      int W,    int H,                   //Image Dimensions
      float sinTheta, float cosTheta )   //Rotation Parameters
{
   const int ix = get_global_id(0);
   const int iy = get_global_id(1);
   int xc = W/2;
   int yc = H/2;
   int xpos =  ( ix-xc)*cosTheta - (iy-yc)*sinTheta+xc;
   int ypos =  (ix-xc)*sinTheta + ( iy-yc)*cosTheta+yc;
   if ((xpos>=0) && (xpos< W)   && (ypos>=0) && (ypos< H))
      dest_data[ypos*W+xpos]= src_data[iy*W+ix];
}

旋转45度：

正如上面代码中所给出的那样，在C代码中需要两重循环来计算横纵坐标上新的坐标位置。其实，在图像旋转的算法中每个点的计算可以独立进行，与其它点的坐标位置没有关系，所以并行处理较为方便。OpenCL C kernel代码中用了并行 处理。

上面的代码在Intel的OpenCL平台上进行了测试，处理器为双核处理器，图像大小 为4288*3216，如果用循环的方式运行时间稳定在0.256s左右，而如果用OpenCL C kernel并行的方式，运行时间稳定在0.132秒左右。GPU的测试在NVIDIA的GeForce G105M显卡上进行，运行时间稳定在0.0810s左右。从循环的方式，双核CPU并行以及GPU并行计算已经可以看出，OpenCL编程的确能大大提高执行效率。
————————————————————————————————————————————

2.5 总结

通过对OpenCL编程的分析和实验可以得出，用OpenCL编写的应用具有很好的移植性，能在不同的设备上运行。OpenCL C kernel一般用并行的方式处理，所以能极大地提高程序的运行效率。

异构并行计算变得越来越普遍，然而对于现今存在的OpenCL版本来说，的确还存在很多不足，例如编写内核，需要对问题的并行情况做较为深入的分析，对于内存的管理还是需要程序员来显式地申明、显式地在主存和设备的存储器之间进行移动， 还不能完全交给系统自动完成。从这些方面，OpenCL的确还需加强，要使得人们能高效而又灵活地开发应用，还有很多工作要完成。

参考博文：

https://blog.csdn.net/w1992wishes/article/details/80426476
https://www.cnblogs.com/wangshide/archive/2012/01/07/2315830.html
https://www.zybuluo.com/w1992wishes/note/1129621
https://www.iteye.com/blog/jubincn-1163960
https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#axzz4YFZTibzg