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	10、8.1、7、Server 2012 R2、Server 2008 R2
Linux	X86_64	Fedora、OpenSUSE、RHEL、CentOS、SLES、steamOS、Ubuntu.
Linux	ppc64le	Ubuntu
Mac OSX	x86_64	10.11、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, 中國海洋大學.

免責聲明！

本站轉載的文章為個人學習借鑒使用，本站對版權不負任何法律責任。如果侵犯了您的隱私權益，請聯系本站郵箱yoyou2525@163.com刪除。

猜您在找 從CUDA開始讀OpenCL OpenCL與Cuda技術 OpenCV GPU CUDA OpenCL 配置 GPU架構和CUDA架構 OpenCL:一種異構計算架構從零開始學習OpenCL開發（一）架構 GPU 的硬件基本概念，Cuda和Opencl名詞關系對應 GPGPU OpenCL/CUDA 高性能編程的10大注意事項 OpenCL CUDA 架構與編程概述