CUDA与OpenCL架构

CPU 和 GPU之间浮点运算能力之因此存在这样的差别，缘由就在于CPU具备复杂的控制逻辑和大容量的缓存，适合进行控制转移，处理分支繁杂的任务，而GPU专为计算密集型、高度并行化的计算而设计。于是GPU具备更多ALU(算术运算单元)和高显存带宽的设计能使更多晶体管用于数据处理，而非数据缓存和流控制，以下图所示。

图 4 GPU中的更多晶体管用于数据处理

更具体地说，GPU专用于解决可表示为数据并行计算的问题——在许多数据元素上并行执行的程序，具备极高的计算密度（数学运算与存储器运算的比率）。因为全部数据元素都执行相同的程序，所以对精密流控制的要求不高；因为在许多数据元素上运行，且具备较高的计算密度，于是可经过计算隐藏存储器访问延迟，而没必要使用较大的数据缓存。

2. CUDA架构

CUDA是一种新的操做GPU计算的硬件和软件架构，它将GPU视做一个数据并行计算设备，并且无需把这些计算映射到图形API。

2.1 硬件架构

2.1.1 GPU困境

虽然GPU经过图形应用程序的算法存在以下几个特征：算法密集、高度并行、控制简单、分多个阶段执行以及前馈(Feed Forward)流水线等，可以在高度密集型的并行计算上得到较高的性能和速度，但在2007年之前GPU要实现这样的应用仍是存在许多困难的：

GPU 只能经过一个图形的API来编程，这不只加剧了学习负担更形成那些非图像应用程序处理这些 API 的额外开销。
因为DRAM内存带宽，一些程序会遇到瓶颈。
没法在 DRAM 上进行通用写操做。

因此NVIDIA于2006年11月在G80系列中引入的Tesla统一图形和计算架构扩展了GPU，使其超越了图形领域。经过扩展处理器和存储器分区的数量，其强大的多线程处理器阵列已经成为高效的统一计算平台，同时适用于图形和通用并行计算应用程序。从G80系列开始NVIDIA加入了对CUDA的支持。

2.1.2 芯片结构

具备Tesla架构的GPU是具备芯片共享存储器的一组SIMT（单指令多线程）多处理器。它以一个可伸缩的多线程流处理器（Streaming Multiprocessors，SMs）阵列为中心实现了MIMD（多指令多数据）的异步并行机制，其中每一个多处理器包含多个标量处理器（Scalar Processor，SP），为了管理运行各类不一样程序的数百个线程，SIMT架构的多处理器会将各线程映射到一个标量处理器核心，各标量线程使用本身的指令地址和寄存器状态独立执行。

图 5 GPU的共享存储器的SIMT多处理器模型

如上图所示，每一个多处理器(Multiprocessor)都有一个属于如下四种类型之一的芯片存储器：

每一个处理器上有一组本地 32 位寄存器（Registers）；
并行数据缓存或共享存储器(Shared Memory)，由全部标量处理器核心共享，共享存储器空间就位于此处；
只读固定缓存(Constant Cache)，由全部标量处理器核心共享，可加速从固定存储器空间进行的读取操做（这是设备存储器的一个只读区域）；
一个只读纹理缓存(Texture Cache)，由全部标量处理器核心共享，加速从纹理存储器空间进行的读取操做（这是设备存储器的一个只读区域），每一个多处理器都会经过实现不一样寻址模型和数据过滤的纹理单元访问纹理缓存。

多处理器 SIMT 单元以32个并行线程为一组来建立、管理、调度和执行线程，这样的线程组称为 warp 块(束)，即以线程束为调度单位，但只有全部32个线程都在诸如内存读取这样的操做时，它们就会被挂起，如图 7所示的状态变化。当主机CPU上的CUDA程序调用内核网格时，网格的块将被枚举并分发到具备可用执行容量的多处理器；SIMT 单元会选择一个已准备好执行的 warp 块，并将下一条指令发送到该 warp 块的活动线程。一个线程块的线程在一个多处理器上并发执行，在线程块终止时，将在空闲多处理器上启动新块。

图 6 CPU五种状态的转换

图 7 线程束调度变化

2.2 软件架构

CUDA是一种新的操做GPU计算的硬件和软件架构，它将GPU视做一个数据并行计算设备，并且无需把这些计算映射到图形API。操做系统的多任务机制能够同时管理CUDA访问GPU和图形程序的运行库，其计算特性支持利用CUDA直观地编写GPU核心程序。目前Tesla架构具备在笔记本电脑、台式机、工做站和服务器上的普遍可用性，配以C/C++语言的编程环境和CUDA软件，使这种架构得以成为最优秀的超级计算平台。

图 8 CUDA软件层次结构

CUDA在软件方面组成有：一个CUDA库、一个应用程序编程接口（API）及其运行库(Runtime)、两个较高级别的通用数学库，即CUFFT和CUBLAS。CUDA改进了DRAM的读写灵活性，使得GPU与CPU的机制相吻合。另外一方面，CUDA 提供了片上（on-chip）共享内存，使得线程之间能够共享数据。应用程序能够利用共享内存来减小DRAM的数据传送，更少的依赖DRAM的内存带宽。

2.3 编程模型

CUDA程序构架分为两部分：Host和Device。通常而言，Host指的是CPU，Device指的是GPU。在CUDA程序构架中，主程序仍是由 CPU 来执行，而当遇到数据并行处理的部分，CUDA 就会将程序编译成 GPU 能执行的程序，并传送到GPU。而这个程序在CUDA里称作核（kernel）。CUDA容许程序员定义称为核的C语言函数，从而扩展了 C 语言，在调用此类函数时，它将由N个不一样的CUDA线程并行执行N次，这与普通的C语言函数只执行一次的方式不一样。执行核的每一个线程都会被分配一个独特的线程ID，可经过内置的threadIdx变量在内核中访问此ID。

在 CUDA 程序中，主程序在调用任何 GPU 内核以前，必须对核进行执行配置，即肯定线程块数和每一个线程块中的线程数以及共享内存大小。

2.3.1 线程层次结构

在GPU中要执行的线程，根据最有效的数据共享来建立块(Block)，其类型有一维、二维或三维。在同一个块内的线程可彼此协做，经过一些共享存储器来共享数据，并同步其执行来协调存储器访问。一个块中的全部线程都必须位于同一个处理器核心中。于是，一个处理器核心的有限存储器资源制约了每一个块的线程数量。在早起的 NVIDIA 架构中，一个线程块最多能够包含 512 个线程，而在后期出现的一些设备中则最多可支持1024个线程。通常 GPGPU 程序线程数目是不少的，因此不能把全部的线程都塞到同一个块里。但一个内核可由多个大小相同的线程块同时执行，于是线程总数应等于每一个块的线程数乘以块的数量。这些一样维度和大小的块将组织为一个一维或二维线程块网格(Grid)。具体框架如图 9所示。

图 9 线程块网格

核函数只能在主机端调用，其调用形式为：Kernel<<<Dg,Db, Ns, S>>>(param list)

Dg：用于定义整个grid的维度和尺寸，即一个grid有多少个block。为dim3类型。Dim3 Dg(Dg.x, Dg.y, 1)表示grid中每行有Dg.x个block，每列有Dg.y个block，第三维恒为1(目前一个核函数只有一个grid)。整个grid中共有Dg.x*Dg.y个block，其中Dg.x和Dg.y最大值为65535。
Db：用于定义一个block的维度和尺寸，即一个block有多少个thread。为dim3类型。Dim3 Db(Db.x, Db.y, Db.z)表示整个block中每行有Db.x个thread，每列有Db.y个thread，高度为Db.z。Db.x和Db.y最大值为512，Db.z最大值为62。一个block中共有Db.x*Db.y*Db.z个thread。计算能力为1.0,1.1的硬件该乘积的最大值为768，计算能力为1.2,1.3的硬件支持的最大值为1024。
Ns：是一个可选参数，用于设置每一个block除了静态分配的shared Memory之外，最多能动态分配的shared memory大小，单位为byte。不须要动态分配时该值为0或省略不写。
S：是一个cudaStream_t类型的可选参数，初始值为零，表示该核函数处在哪一个流之中。

以下是一个CUDA简单的求和程序：

图 10 CUDA求和程序

2.3.2 存储器层次结构

CUDA 设备拥有多个独立的存储空间，其中包括：全局存储器、本地存储器、共享存储器、常量存储器、纹理存储器和寄存器，如图 11所示。

图 11 CUDA设备上的存储器

CUDA线程可在执行过程当中访问多个存储器空间的数据，如图 12所示其中：

每一个线程都有一个私有的本地存储器。
每一个线程块都有一个共享存储器，该存储器对于块内的全部线程都是可见的，而且与块具备相同的生命周期。
全部线程均可访问相同的全局存储器。
此外还有两个只读的存储器空间，可由全部线程访问，这两个空间是常量存储器空间和纹理存储器空间。全局、固定和纹理存储器空间通过优化，适于不一样的存储器用途。纹理存储器也为某些特殊的数据格式提供了不一样的寻址模式以及数据过滤，方便 Host对流数据的快速存取。

图 12 存储器的应用层次

2.3.3 主机（Host）和设备（Device）

如图 13所示，CUDA 假设线程可在物理上独立的设备上执行，此类设备做为运行C语言程序的主机的协处理器操做。内核在GPU上执行，而C语言程序的其余部分在CPU上执行（即串行代码在主机上执行，而并行代码在设备上执行）。此外，CUDA还假设主机和设备均维护本身的DRAM，分别称为主机存储器和设备存储器。于是，一个程序经过调用CUDA运行库来管理对内核可见的全局、固定和纹理存储器空间。这种管理包括设备存储器的分配和取消分配，还包括主机和设备存储器之间的数据传输。

图 13 CUDA异构编程模型

2.4 CUDA软硬件

2.4.1 CUDA术语

因为CUDA中存在许多概念和术语，诸如SM、block、SP等多个概念不容易理解，将其与CPU的一些概念进行比较，以下表所示。

CPU	GPU	层次
算术逻辑和控制单元	流处理器(SM)	硬件
算术单元	批量处理器(SP)	硬件
进程	Block	软件
线程	thread	软件
调度单位	Warp	软件

图 14 NVIDIA 630显卡CUDA信息

2.4.2 硬件利用率

当为一个GPU分配一个内核函数，咱们关心的是如何才能充分利用GPU的计算能力，但因为不一样的硬件有不一样的计算能力，SM一次最多能容纳的线程数也不尽相同，SM一次最多能容纳的线程数量主要与底层硬件的计算能力有关，以下表显示了在不一样的计算能力的设备上，每一个线程块上开启不一样数量的线程时设备的利用率。

计算能力每一个线程块的线程数	1.0	1.1	1.2	1.3	2.0	2.1	3.0
64	67	67	50	50	33	33	50
96	100	100	75	75	50	50	75
128	100	100	100	100	67	67	100
192	100	100	94	94	100	100	94
256	100	100	100	100	100	100	100
……	……

3 OpenCL架构

3.1 简介

OpenCL（Open Computing Language），即开放运算语言，是一个统一的开放式的开发平台。OpenCL是首个提出的并行开发的开放式的、兼容的、免费的标准，它的目的是为异构系统通用提供统一开发平台。OpenCL最初是由苹果公司设想和开发，并在与AMD，IBM，英特尔和NVIDIA技术团队的合做之下初步完善。随后，苹果将这一草案提交至Khronos Group。

图 15 OpenCL历史版本

3.2 框架组成

OpenCL的框架组成能够划分为三个部分，分别为OpenCL平台API、OpenCL运行时API，以及OpenCL内核编程语言。

3.2.1 平台API

平台（Platform）这个词在OpenCL中拥有很是特定的含义，它表示的是宿主机、OpenCL设备和OpenCL框架的组合。多个OpenCL平台能够共存于一台异构计算机。举个例子，CPU开发人员和GPU开发人员能够在同一个系统上分别定义本身的OpenCL框架。这时就须要一种方法来查询系统中可用的OpenCL 框架，哪些OpenCL设备是可用的，以及这些OpenCL设备的特性。至关于CUDA的主机和设备之间的关系。

此外，为了造成一个给定的OpenCL应用平台，还须要对这些框架和设备所属的子集进行控制。这些功能都是由OpenCL平台API中的函数来解决的。此外，平台API还提供了为OpenCL建立上下文的函数。OpenCL的上下文规定了OpenCL应用程序的打开方式(至关是CUDA中核函数的调用)，这能够在宿主机程序代码中获得验证。

3.2.2 运行时API

平台API提供函数建立好上下文以后，运行时API主要提供使用上下文提供的功能知足各类应用需求的函数。这是一个规模庞大且内容十分复杂的函数集。运行时API的第一个任务是建立一个命令队列。命令队列与设备相关联，并且一个上下文中能够同时存在多个活动的命令队列。有了命令队列，就能够经过调用运行时API提供的函数来进行内存对象的定义以及管理内存中的对象所依赖的全部其余对象。以上是内存对象的持有操做，另外还有释放操做，也是由运行时API提供的。

此外，运行时API还提供了建立动态库所须要的程序对象的函数，正是这些动态库实现了Kernel的定义。最后，运行时层的函数会发出与命令队列交互的命令。此外，管理数据共享和对内核的执行加以限制同步点也是由运行时API处理的。

3.2.3 内核编程语言

内核编程语言是用于编写OpenCL内核代码的。除了宿主机程序以外，内核程序也十分重要，它负责完成OpenCL中的实际工做。在部分OpenCL实现中用户能够跟其余语言编写的原生内核实现交互，但多数状况下内核是须要用户使用内核编程语言编写实现的。OpenCL C编程语言就是OpenCL中的内核编程语言，该编程语言是"ISO C99 标准"的一个扩展子集，也就是说它是由 ISO C99语言派生而来的。如今的OpenCL2.1还支持C++，是基于eISO/IEC JTC1 SC22 WG21 N 3690(C++14)。

3.2.4 适合平台

AMD

根据AMD官网所提供的内容，OpenCL在AMD显卡中只能适用X86核心的CPU架构，而对其余PowerPC和ARM架构则不适用；而且也不是全部的AMD显卡都能运行OpenCL，按其官网介绍只能是AMD Radeon、AMD FirePro和AMD Firestream三种类型的显卡；但对于操做系统则能够是Linux或Windows的系统，如表 1所示。

表 1 AMD OpenCL

CPU架构	显卡类型	操做系统	系统位数
X86	AMD Radeon	Linux/ Windows	32/64
	AMD FirePro	Linux/ Windows	32/64
	AMD Firestream	Linux/ Windows	32/64

NVIDIA

NVIDIA OpenCL是一种运行于具备CUDA能力GPU上的一种底层API，即OpenCL是运行于CUDA之上的一种API，从而若适用CUDA的平台，也一样适用OpenCL。根据NVIDIA官网最新版本的CUDA 7.5适合的平台如表 2所示。

表 2 NVIDIA OpenCL

操做系统	CPU架构	Distribution
Windows	X86_64	十、8.一、七、Server 2012 R二、Server 2008 R2
Linux	X86_64	Fedora、OpenSUSE、RHEL、CentOS、SLES、steamOS、Ubuntu.
Linux	ppc64le	Ubuntu
Mac OSX	x86_64	10.十一、10.十、10.9

3.3 计算架构

OpenCL 的设计目标是为开发人员提供一套移植性强且高效运行的解决方案。为了更好的描述OpenCL设计的核心理念，Khronos Group官方将OpenCL的计算架构分解成四个模型，分别平台模型（Platform Model）、内存模型（Memory Model）、执行模型（Execution Model）以及编程模型（Programming Model）。

3.3.1 平台模型（Platform Model）

从总体上来看，主机（host）端是负责掌管整个运算的全部计算资源，所以OpenCL 应用程序首先是由主机端开始，而后由程序将各个计算命令从主机端发送给每一个 GPU 设备处理单元，运行完毕以后最后由主机端结束。

图 16 OpenCL架构的平台模型

平台模型如图 16所示。从图中能够直观的看到，最基本处理单位是Processing Element，简称PE(处理单元)，而一个或多个PE组成了Compute Unit，简称CU(计算单元)，进而一个或多个CU就组成了Compute Device，即OpenCL设备。最后，一个或多个OpenCL设备链接到主机，并等待着处理主机发送的计算指令，因为PE是最基本处理单位，所以每条计算指令最终都归PE进行处理，而PE是在CU中的。

3.3.2 内存模型（Memory Model）

OpenCL将内核程序中用到的内存分为图 17所示的四种不一样的类型。

图 17 OpenCL内存模型

其中它们的读写特性分别为：

Global memory:工做区内的全部工做节点均可以自由的读写其中的任何数据。OpenCL C语言提供了全局缓存（Global buffer）的内建函数。
Constant memory: 工做区内的全部工做节点能够读取其中的任何数据但不能够对数据内容进行更改，在内核程序的执行过程当中保持不变。主机端负责分配和初始化常量缓存（Constant buffer）。
Local memory: 只有同一工做组中的工做节点才能够对该类内存进行读写操做。它既能够为 OpenCL 的执行分配一块私有内存空间，也能够直接将其映射到一块全局缓存（Global buffer）上。特色是运行速度快。
Private memory: 只有当前的工做节点能对该内存进行访问和读写操做。一个工做节点内部的私有缓存（Private buffer）对其余节点来讲是不可见的。

表 3 OpenCL各类存储器的分配方式和访问权限

存储器类型	主机		内核
存储器类型	分配方式	访问权限	分配方式	访问权限
Global	动态分配	可读、可写	不可分配	可读、可写
Constant	动态分配	可读、可写	静态分配	只读
Local	动态分配	不可访问	静态分配	可读、可写
Private	不可分配	不可访问	静态分配	可读、可写

3.3.3 执行模型（Execution Model）

OpenCL的执行模型是应用程序经过主机端对OpenCL设备端上的内核程序进行管理，该模型分为两个模块:一个是在主机端执行的管理程序，也称为Hostprogram，另外一个是主机端的Hostprogram所管理的在OpenCL上执行的程序，也被称做Kernels。在执行Kernels前，先要创建一个索引空间，来对设备里的每一个节点进行标识，每一个节点都将执行相同的kernel程序。在每一个工做组中，都有一个局部ID，每一个节点在全局里还有个全局 ID，OpenCL使用NDRange来定义这个索引空间。

图 18 OpenCL执行模型

如图 18所示的OpenCL执行模型，其过程能够细分为以下的步骤完成：

查询链接主机上的OpenCL设备；
建立一个关联到OpenCL设备的context；
在关联的设备上建立可执行程序；
从程序池中选择kernel程序；
从主机或设备上建立存储单元；
若是须要将主机的数据复制到OpenCL设备上的存储单元上；
执行kernel程序执行；
从OpenCL设备上复制结果到主机上。

3.3.4 编程模型（Programming Model）

OpenCL支持两种编程模型，分别为数据并行编程模型和任务并行编程模型，并支持上面由这两种编程模型混合的混合编程模型。

数据并行编程模型

OpenCL提供一个分层的数据并行编程模型，即典型的SIMD计算模型，其特色是每一个数据经由一样的指令序列处理，而处理数据的次序是不肯定的，而且每一个数据的处理是不相干的，即任一线程的计算不得依赖于其它线程的结果(包括中间结果)。

任务并行编程模型

任务并行模型中的每一个内核是在一个独立的索引空间中执行的，也就是说，执行内核的计算机单元内只有一个工做组，其中只有一个工做项。在这样的模型中，每一个线程均可以执行不一样的带啊，着至关于MIMD的计算模型，适合多核心CPU。

4. CUDA与OpenCL之间的差别

CUDA和OpenCL都是实现计算机异构并行计算架构，然而CUDA是针对NVIDIA公司的GPU，而OpenCL是一种通用的计算框架。二者基本的差异为：

表 4 CUDA与OpenCL基本差异

	CUDA	OpenCL
技术类型	控制	开源和VIP服务
出现时间	2006年	2008年
SDK企业	NVIDIA	具体根据企业
SDK是否免费	Yes	依赖企业
实现企业	仅NVIDIA	Apple、NVIDIA、AMD、IBM
支持系统	Windows, Linux, Mac OS X; 32 and 64‐bit	依赖具体企业
支持设备类型	仅NVIDIA GPU	多种类型
支持嵌入式设备	NO	Yes

4.1 硬件架构

4.1.1 芯片结构

CUDA和OpenCL的芯片结构相似，都是按等级划分的，并逐渐提升等级。然而OpenCL更具通用性并使用更加通常的技术，如OpenCL经过使用Processing Element代替CUDA的Processor，同时CUDA的模型只能在NVIDIA架构的GPU上运行。

图 19 OpenCL与CUDA芯片结构

4.1.2 存储结构

CUDA和OpenCL的存储模型如图 20所示，二者的模型类型，都是将设备和主机的存储单元独立分开，它们的都是按等级划分并须要程序员进行精确的控制，并都能经过API来查询设备的状态、容量等信息。而OpenCL模型更加抽象，并为不一样的平台提供更加灵活的实现，在CUDA模型的Local Memory在OpenCL没有相关的概念。对于CUDA和OpenCL模型的相似概念，经过表 5列出二者对存储单元命名的差别。

图 20 CUDA与OpenCL存储模型

表 5 CUDA与OpenCL存储器对比

OpenCL	CUDA
Host memory	Host memory
Global memory	Global or Device memory
Global memory	Local memory
Constant memory	Constant memory
Global memory	Texture memory
Local memory	Shared memory
Private memory	Registers

4.2 软件架构

4.2.1 应用框架

一个典型的应用框架都包含有libraries、API、drivers/compilies和runtime system等来支持软件开发。CUDA和OpenCL也拥有类似的特性，都拥有runtime API和library API，但具体环境下的建立和复制API是不一样的，而且OpenCL能够经过平台层查询设备的信息；CUDA的kernel能够直接经过NVIDIA 驱动执行，而OpenCL的kernel必须经过OpenCL驱动，但这样可能影响到性能。由于OpenCL毕竟是一个开源的标准，为了适应不一样的CPU、GPU和设备都可以获得正常执行；而CUDA只针对NVIDIA的GPU产品。

图 21 CUDA与OpenCL应用框架

4.2.2 编程模型

开发模型

CUDA和OpenCL应用的开发模型基本一致，都是由Host和Device程序组成。程序首先开始执行Host程序，而后由Host程序激活Device程序kernel执行。其中二者也存在一些差异，如表 6所示。

表 6 CUDA与OpenCL开发模型比较

	CUDA	OpenCL
精确的host和device代码分离	Yes	Yes
定制的kernel编程语言	Yes	Yes
并行的kernel编程语言	Yes	仅有OpenCL C或具体的企业语言
支持数据并行kernels	Yes	Yes
支持任务并行kernels	No	Yes
多编程接口	Yes，包括OpenCL	仅支持标准C的API
host和device结合程度	Yes，效率很是高	No，分离的编译而且kernel和API调用是不相干的
Graphics支持	OpenGL and Direct3D	OpenGL

kernel 编程

kernel程序是指Device设备上执行的代码，它是直接在设备上执行，受具体设备的限制，具体二者的差异，如表 7所示。

表 7 kernel编程差别

	CUDA	OpenCL
基于开发语言版本	基本C和C++、C++14	C99
访问work-item方式	经过内置的变量	经过内置函数
内置vector类型	基本vector类型，没有操做和函数	vector、literals类型，并内置操做和函数
Voting函数	Yes (CC 1.2 or greater)	No
Atomic函数	Yes (CC 1.1 or greater)	Only as extension
异步内存空间复制和预取函数	No	Yes
支持C++语言功能	Yes，受限，但大部分功能都支持	No

Host 编程

Host端基本是串行的，CUDA和OpenCL的差异主要表如今调用Device的API的差别，因此表 8描述了二者之间API的差别。

表 8 Host端可用的API比较

C Runtime for CUDA	CUDA Driver API	OpenCL API
Setup
	Initialize driver Get device(s) (Choose device) Create context	Initialize plauorm Get devices Choose device Create context Create command queue
Device and host memory buffer setup
Allocate host memory Allocate device memory for input Copy host memory to device memory Allocate device memory for result	Allocate host memory Allocate device memory for input Copy host memory to device memory Allocate device memory for result	Allocate host memory Allocate device memory for input Copy host memory to device memory Allocate device memory for result
Initialize kernel
	Load kernel module (Build program) Get module function	Load kernel source Create program object Build program Create kernel object bound to kernel function
Execute the kernel
Setup execution configuration Invoke the kernel (directly with its parameters)	Setup kernel arguments Setup execution configuration Invoke the kernel	Setup kernel arguments Setup execution configuration Invoke the kernel
Copy results to host
Copy results from device memory	Copy results from device memory	Copy results from device memory
Cleanup
Cleanup all set up above	Cleanup all set up above	Cleanup all set up above

4.3 性能

本节根据学术上对CUDA和OpenCL的研究，比较二者的性能，其中本文简单以[1-3]研究成功比较CUDA和OpenCL之间的性能差别，若需详细了解CUDA和OpenCL之间的性能差别能够参考[4-15]。

4.3.1 AES实现

Wang^[1]提出一种在XTS模式的AES实现，并对OpenCL和CUDA性能进行比较。如图 22和图 23所示，整体性能CUDA要比OpenCL好10%~20%之间。

图 22 the performance for OpenCL and CUDA in NVIDIA GTX 285

图 23 the runtime for OpenCL and CUDA in NVIDIA GTX 285

4.3.2 三维可视化加速模型

上海理工大学^[3]提出合理设计内核函数实现改进的光线投射算法在GPU上并行和并发运行的三维可视化加速模型，该模型实现代码可不用修改在两大主流显卡平台NVIDIA和AMD上任意移植，经过实验证实比较OpenCL与CUDA之间的性能。

表 9 各类实现光线投射算法的三维可视化模型的运行效果对比表

4.3.3 MAGMA和DGEMM算法

做者^[2]已经在先前的版本中使用CUDA实现了MAGMA(Matrix Algebra on GPU and multicore architectures)和DGEMM算法，如今将其实现移植到OpenCL API，并对二者的性能进行比较。在NVIDIA处理器上进行测试，其结果是CUDA的性能要高于OpenCL。

图 24 DGEMM performance on Tesla C2050 under OpenCL and CUDA

4.4 总结

CUDA与OpenCL的功能和架构类似，只是CUDA只针对NVIDIA的产品，而OpenCL是一种通用性框架，可使用多种品牌的产品，因此CUDA的性能通常状况下要比OpenCL的性能要高10%~20%之间。

4.4.1 CUDA与OpenCL的类似点

关注数据并行计算模型；
将主机和设备的程序和存储分离；
提供定制和标准C语言对设备进行编程；
设备、执行和存储模型是现相似的；
OpenCL已经能够在CUDA之上进行实现了。

4.4.2 CUDA和OpenCL主要的差别点

CUDA是属于NVIDIA公司的技术框架，只有NVIDIA的设备才能执行；
OpenCL是一个开源的框架，其目标是定位不一样的设备；
CUDA拥有更多的API和帮助文档；
CUDA投入市场的时间更早，因此获得更多的支持，而且在研究、产品和应用都比OpenCL丰富；
CUDA有很是多的文档，但也更加模糊。

References

1.Wang, X., et al. AES finalists implementation for GPU and multi-core CPU based on OpenCL. in Anti-Counterfeiting, Security and Identification (ASID), 2011 IEEE International Conference on. 2011: IEEE.

2. Du, P., et al., From CUDA to OpenCL: Towards a performance-portable solution for multi-platform GPU programming. Parallel Computing, 2012. 38(8): p. 391-407.

袁健与高勃, 基于 OpenCL 的三维可视化加速模型. 小型微型计算机系统, 2015. 36(002): 第327-331页.

3. Karimi, K., N.G. Dickson and F. Hamze, A performance comparison of CUDA and OpenCL. arXiv preprint arXiv:1005.2581, 2010.

4. McConnell, S., et al. Scalability of Self-organizing Maps on a GPU cluster using OpenCL and CUDA. in Journal of Physics: Conference Series. 2012: IOP Publishing.

5. Fang, J., A.L. Varbanescu and H. Sips. A comprehensive performance comparison of CUDA and OpenCL. in Parallel Processing (ICPP), 2011 International Conference on. 2011: IEEE.

6. Oliveira, R.S., et al., Comparing CUDA, OpenCL and OpenGL implementations of the cardiac monodomain equations, in Parallel Processing and Applied Mathematics. 2012, Springer. p. 111-120.

7. Harvey, M.J. and G. De Fabritiis, Swan: A tool for porting CUDA programs to OpenCL. Computer Physics Communications, 2011. 182(4): p. 1093-1099.

8. 林乐森, 基于 OpenCL 的 AES 算法并行性分析及加速方案, 2012, 吉林大学.

9. 易卓霖, 基于 GPU 的并行支持向量机的设计与实现, 2011, 西南交通大学.

10. 蒋丽媛等, 基于 OpenCL 的连续数据无关访存密集型函数并行与优化研究. 计算机科学, 2013. 40(3): 第111-115页.

11. 詹云, 赵新灿与谭同德, 基于 OpenCL 的异构系统并行编程. 计算机工程与设计, 2012. 33(11): 第4191-4195页.

12. 王晗, 基于多核环境下的多线程并行程序设计方法研究, 2014, 中原工学院.

13. 黄文慧, 图像处理并行编程方法的研究与应用, 2012, 华南理工大学.

14. 刘寿生, 虚拟现实仿真平台异构并行计算关键技术研究, 2014, 中国海洋大学.