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, 中國海洋大學.