https://www.cnblogs.com/wangshide/archive/2012/01/07/2315830.html

本篇结构：

背景

OpenCL是什么

框架组成

基本概念

编写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和CPU的区别，可以参考我之前的博文GPU编程–CPU和GPU的设计区别。

二、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中…

三、框架组成

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

OpenCL平台API：平台API定义了宿主机程序发现OpenCL设备所用的函数以及这些函数的功能，另外还定义为OpenCL应用创建上下文(上下文表示的是程序运行时所拥有的所有软硬件资源+内存+处理器)的函数。这里的平台指的是宿主机、OpenCL设备和OpenCL框架的组合。

OpenCL运行时API：平台API主要用来创建上下文，运行时API则强调使用这个上下文满足应用需求的函数集，用来管理上下文来创建命令队列以及运行时发生的其它操作。例如，将命令提交到命令队列的函数。

OpenCL编程语言：用来编写内核代码的编程语言，基于ISO C99标准的一个扩展子集，通常称为OpenCL C编程语言。

四、基本概念

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

下面参考上图，简单汇总一些OpenCL中的基本概念，在后续的OpenCL四个模型中也会再提到：

OpenCL平台由两部分组成，宿主机和OpenCL设备:

宿主机Host：宿主机一般为CPU，扮演组织者的角色。其作用包括定义kernel、为kernel指定上下文、定义NDRange和队列等；队列控制着kernel如何执行以及何时执行等细节。

设备Device：通常称为计算设备，可以是CPU、GPU、DSP或硬件提供以及OpenCL开发商支持的任何其他处理器。

SIMT（Single Instruction Multi Thread）：单指令多线程，GPU并行运算的主要方式，很多个多线程同时执行相同的运算指令，当然可能每个线程的数据有所不同，但执行的操作一致。

核函数（Kernel）：是在设备程序上执行运算的入口函数，在主机上调用。

工作项（Work-item）：跟CUDA中的线程（Threads）是同一个概念，N多个工作项（线程）执行同样的核函数，每个Work-item都有一个唯一固定的ID号，一般通过这个ID号来区分需要处理的数据。

工作组（Work-group）：跟CUDA中的线程块（Block）是同一个概念，N多个工作项组成一个工作组，Work-group内的这些Work-item之间可以通信和协作。

ND-Range：跟CUDA中的网格是同一个概念，定义了Work-group的组织形式。

上下文（Context）：定义了整个OpenCL的运行环境，包括Kernel、Device、程序对象和内存对象：

设备：OpenCL程序调用的计算设备。

内核：在计算设备上执行的并行程序。

程序对象：内核程序的源代码(.cl文件)和可执行文件。

内存对象：计算设备执行OpenCL程序所需的变量。

上面只简单列了部分概念，不完全，解释也很简略，目的只是提供一个对OpenCL的简单认识，后续再慢慢加深加全对OpenCL涉及到的概念理解。

五、编写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**

六、参考博文

OpenCL:一种异构计算架构

1 摘要
2 为什么需要OpenCL？
3 OpenCL架构
- 3.1 介绍
- 3.2 平台模型
- 3.3 执行模型
  - 3.3.1 内核
  - 3.3.2 上下文
  - 3.3.3 命令队列
- 3.4 内存模型
- 3.5 编程模型
4 基于OpenCL的编程示例
- 4.1 流程
- 4.2 图像旋转
  - 4.2.1 图像旋转原理
- 4.3 实现流程
- 4.4 图像旋转
5 总结
6 参考文献

1 摘要

由于晶体管功耗、物理性能的限制，CPU的发展受到了很大约束。人们转而寻找其它方式来提高系统性能，如多核处理器，异构平台等。开放式计算语言（OpenCL）的出现为当前大量存在的异构系统的并行计算提供了一个标准。OpenCL通过一系列API的定义，提供硬件独立的编程语言，为程序员提供了灵活而又高效的编程环境。本文通过对OpenCL计算架构的深入讨论，指出了OpenCL编程的优势及不足。并进行了相关编程实践，通过对不同设备的并行编程测试，表明如果采用OpenCL并行编程架构，能显著提高程序的运行效率。

就目前的情况来看，异构系统有很高的性价比。相信在不久的将来，OpenCL将会成为计算机并行、异构计算的重要组成部分。

关键字：OpenCL，异构计算，CPU/GPU计算，并行计算

2 为什么需要OpenCL？

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

为了解决这一问题通常有两种方式，第一种是通过增加处理器的核心数目来为多任务，多线程等提供支持，从整体上提升系统的性能。第二种方式是通过异构的方式，例如可利用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将在异构并行编程中得到广泛的应用。

3 OpenCL架构

3.1 介绍

OpenCL为异构平台提供了一个编写程序，尤其是并行程序的开放的框架标准。 OpenCL所支持的异构平台可由多核CPU、GPU或其他类型的处理器组成。 OpenCL由两部分组成，一是用于编写内核程序(在OpenCL设备上运行的代码) 的语言，二是定义并控制平台的API。OpenCL提供了基于任务和基于数据两种并行计算机制，它极大地扩展了GPU 的应用范围，使之不再局限于图形领域。

OpenCL由Khronos Group维护。Khronos Group是一个非盈利性技术组织，维护着多个开放的工业标准，例如OpenGL和OpenAL。这两个标准分别用于三维图形和计算机音频方面。

OpenCL源程序既可以在多核CPU上也可以在GPU上编译执行，这大大提高了代码的性能和可移植性。OpenCL标准由相应的标准委员会制订，委员会的成员来自业界各个重要厂商（主要有：AMD，Intel，IBM和NVIDIA）。作为用户和程序员期待已久的东西，OpenCL带来两个重要变化：一个跨厂商的非专有软件解决方案；一个跨平台的异构框架以同时发挥系统中所有计算单元的能力。

OpenCL支持广泛的应用，将开发应用的过程一般化比较困难，但是，通常来说，一个基于异构平台的应用主要包含下面的步骤[ 3 ]：

找出组成异构平台的所有组件。
考察组件的特征，这样就能使得软件根据不同的硬件特征来实现。
创建在平台上运行的一组内核。
设置与计算相关的存储对象。
在合适的组件上以正确的顺序执行内核。
收集结果。

这些步骤通过在OpenCL内部的一系列API和内核编程环境来实现。这一实现采用“分治”策略。可将问题分解为下面的模型[ 1 ] 平台模型执行模型存储模型编程模型

这些概念是OpenCL整体架构的核心。这四个模型将贯穿在整个OpenCL的编程过程中。

下面就简要介绍这四个模型的相关内容。

3.2 平台模型

平台模型（如图1）指定有一个处理器（主机Host）来协调程序的执行，一个或多个处理器（设备Devices）来执行OpenCL C代码。在这里其实仅仅是一种抽象的硬件模型，这样就能方便程序员编写OpenCL C函数（称之为内核）并在不同的设备上执行。

OpenCL:简介 https://www.cnblogs.com/wangshide/archive/2012/01/07/2315830.htmlOpenCL:一种异构计算架构目录1 摘要2 为什么需要OpenCL？3 OpenCL架构4 基于OpenCL的编程示例

图中的设备可以被看成是CPU/GPU，而设备中的计算单元可以被看成是 CPU/GPU的核，计算单元的所有处理节点作为SIMD单元或SPMD单元（每个处理节点维护自己的程序计数器）执行单个指令流。抽象的平台模型更与当前的GPU的架构接近。

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

3.3 执行模型

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

3.3.1 内核

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

OpenCL:简介 https://www.cnblogs.com/wangshide/archive/2012/01/07/2315830.htmlOpenCL:一种异构计算架构目录1 摘要2 为什么需要OpenCL？3 OpenCL架构4 基于OpenCL的编程示例

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

__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次。合理设置节点数目，工作组数目能提高程序的并行度。

3.3.2 上下文

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

3.3.3 命令队列

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

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

3.4 内存模型

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

OpenCL:简介 https://www.cnblogs.com/wangshide/archive/2012/01/07/2315830.htmlOpenCL:一种异构计算架构目录1 摘要2 为什么需要OpenCL？3 OpenCL架构4 基于OpenCL的编程示例

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

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

3.5 编程模型

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

4 基于OpenCL的编程示例

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

4.1 流程

OpenCL:简介 https://www.cnblogs.com/wangshide/archive/2012/01/07/2315830.htmlOpenCL:一种异构计算架构目录1 摘要2 为什么需要OpenCL？3 OpenCL架构4 基于OpenCL的编程示例

4.2 图像旋转

4.2.1 图像旋转原理

图像旋转是指把定义的图像绕某一点以逆时针或顺时针方向旋转一定的角度，通常是指绕图像的中心以逆时针方向旋转。假设图像的左上角为(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度

OpenCL:简介 https://www.cnblogs.com/wangshide/archive/2012/01/07/2315830.htmlOpenCL:一种异构计算架构目录1 摘要2 为什么需要OpenCL？3 OpenCL架构4 基于OpenCL的编程示例

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

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

5 总结

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

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

6 参考文献

【1】 Aaftab Munshi. The OpenCL Specification Version1.1 Document Revision:44[M]. Khronos OpenCL Working Group. 2011.6.1.

【2】Aaftab Munshi. 倪庆亮译. OpenCL规范 Version1.0 Document Revision:48[M]. Khronos OpenCL Working Group. 2009.10.6.

【3】Aaftab Munshi, Benedict R. Gaster, Timothy G. Mattson, James Fung, Dan Ginsburg. OpenCL Programming Guide [M]. Addison-Wesley Professional. 2011.7.23.

【4】Benedict Gaster, Lee Howes, David R. Kaeli and Perhaad Mistry. Heterogeneous Computing with OpenCL[M]. Morgan Kaufmann, 1 edition. 2011.8.31.

【5】Slo-Li Chu, Chih-Chieh Hsiao. OpenCL: Make Ubiquitous Supercomputing Possible[J]. IEEE International Conference on High Performance Computing and Communications. 2010 12th 556-561.

【6】John E. Stone, David Gohara, Guochun Shi. OpenCL: A parallel programming standard for heterogeneous computing systems[J]. Copublished by the IEEE CS and the AIP. 2010.5/6 66-72.

【7】Kyle Spafford, Jeremy Meredith, Jeffrey Vetter. Maestro:Data Orchestration and Tuning for OpenCL Devices[J]. P. D'Ambra,M.Guarracino, and D.Talia (Eds.):Euro-Par 2010,Part II,LNCS6272, pp. 275–286, 2010. \copyright Springer-Verlag Berlin Heidelberg 2010.

Author: Let it be!

Date: 2011-11-13 00:12:07

HTML generated by org-mode 6.35i in emacs 24

OpenCL:简介 https://www.cnblogs.com/wangshide/archive/2012/01/07/2315830.htmlOpenCL:一种异构计算架构目录1 摘要2 为什么需要OpenCL？3 OpenCL架构4 基于OpenCL的编程示例

https://www.cnblogs.com/wangshide/archive/2012/01/07/2315830.html

OpenCL:一种异构计算架构

目录

1 摘要

2 为什么需要OpenCL？

3 OpenCL架构

3.1 介绍

3.2 平台模型

3.3 执行模型

3.4 内存模型

3.5 编程模型

4 基于OpenCL的编程示例

4.1 流程

4.2 图像旋转

5 总结

6 参考文献

继续阅读

CUDA4.0 + VS2010 + Win7 64 + VAssistX 安装经验

用Intel HAXM加速Android Emulator

服务器硬件配置深度学习训练/GPU服务器硬件配置

ios开发之离屏渲染前言

k8s调用gpu，nvidia-docker使用gpu

CentOS上Docker安装GPU支持Nvidia-docker

Window10 Tensorflow 2.1 GPU 安装和测试

Vulkan demo运行

Unity Shader 系列笔记一：内置变量

AM5728 Opencl 案例汇总：实现sobel算法，计算向量和，矩阵转置

AMD-SDK的学习[1]--AdvancedConvolution

TensorFlow错题集

Android 过度绘制简单的分析概念对比分析参考文章

pyTorch 使用多GPU训练

【cpu or gpu】【tensorflow】怎么查看用的是CPU还是GPU

一种解决思路： ImportError: libcublas.so.10.0: cannot open shared object file: No such file