AMD OpenCL大學課程是非常好的入門級OpenCL教程，通過看教程中的PPT，我們能夠很快的了解OpenCL機制以及編程方法。下載地址：http://developer.amd.com/zones/OpenCLZone/universities/Pages/default.aspx教程中的英文很簡單，我相信學OpenCL的人都能看得懂，而且看原汁原味的英文表述，更有利於我們了解各種術語的來龍去脈。
一、並行計算概述
在計算機術語中，並行性是指：把一個復雜問題,分解成多個能同時處理的子問題的能力。要實現並行計算，首先我們要有物理上能夠實現並行計算的硬件設備，比如多核CPU，每個核能同時實現算術或邏輯運算。
通常，我們通過GPU實現兩類並行計算：
任務並行：把一個問題分解為能夠同時執行的多個任務。
數據並行：同一個任務內，它的各個部分同時執行。
下面我們通過一個農場主雇佣工人摘蘋果的例子來描述不同種類的並行計算。
-
摘蘋果的工人就是硬件上的並行處理單元（process elements)。
樹就是要執行的任務。
蘋果就是要處理的數據。
串行的任務處理就如下圖所示，一個工人背着梯子摘完所有樹上的蘋果（一個處理單元處理完所有任務的數據）。
數據並行就好比農場主雇佣了好多工人來摘完一個樹上的蘋果（多個處理單元並行完成一個任務中的數據），這樣就能很快摘完一顆樹上的蘋果。
農場主也可以為每棵樹安排一個工人，這就好比任務並行。在每個任務內，由於只有一個工人，所以是串行執行的，但任務之間是並行的。
對一個復雜問題，影響並行計算的因素很多。通常，我們都是通過分解問題的方式來實施並算法行。
這又包括兩方面內容：
- 任務分解：把算法分解成很多的小任務，就像前面的例子中，把果園按蘋果樹進行划分，這時我們並不關注數據，也就是說不關注每個樹上到底有多少個蘋果。
- 數據分解：就是把很多數據，分成不同的、離散的小塊，這些數據塊能夠被並行執行，就好比前面例子中的蘋果。
通常我們按照算法之間的依賴關系來分解任務，這樣就形成了一個任務關系圖。一個任務只有沒有依賴任務的時候，才能夠被執行。
這有點類似於數據結構中的有向無環圖，兩個沒有連通路徑的任務之間可以並行執行。下面再給一個烤面包的例子，如果所示，預熱烤箱和購買面粉糖兩個任務之間可以並行執行。
對大多數科學計算和工程應用來說，數據分解一般都是基於輸出數據,例如：
- 在一副圖像中，對一個滑動窗口（例如：3*3像素)內的像素實施濾波操作，可以得到一個輸出像素的卷積。
- 第一個輸入矩陣的第i行乘以第二個輸入矩陣的第j列，得到的向量和即為輸出矩陣第i行，第j列的元素。
這種方法對於輸入和輸出數據是一對一，或者多對一的對應關系比較有效。
也有的數據分解算法是基於輸入數據的，這時，輸入數據和輸出數據一般是一對多的關系，比如求圖像的直方圖，我們要把每個像素放到對應的槽中(bins，對於灰度圖，bin數量通常是256）。一個搜索函數，輸入可能是多個數據，輸出卻只有一個值。對於這類應用，我們一般用每個線程計算輸出的一部分，然后通過同步以及原子操作得到最終的值，OpenCL中求最小值的kernel函數就是典型代表[可以看下ATI Stream Computing OpenCL programming guide第二章中求最小值的kernel例子]。
通常來說，怎樣分解問題和具體算法有關，而且還要考慮自己使用的硬件和軟件，比如AMD GPU平台和Nvdia GPU平台的優化就有很多不同。
二、常用基於硬件和軟件的並行
在上個實際90年代，並行計算主要研究如何在cpu上實施指自動的指令級並行。
- 同時發射多條指令（之間沒有依賴關系），並行執行這些指令。
- 在本教程中，我么不講述自動的硬件級並行，感興趣的話，可以看看計算機體系結構的教程。
高層的並行，比如線程級別的並行，一般很難自動化，需要程序員告訴計算機，該做什么，不該做什么。這時，程序員還要考慮硬件的具體指標，通常特定硬件都是適應於某一類並行編程，比如多核CPU就適合基於任務的並行編程，而GPU更適應於數據並行編程。
| Hardware type |
Examples |
Parallelism |
| Multi-core superscalar processors |
Phenom II CPU |
Task |
| Vector or SIMD processors |
SSE units (x86 CPUs) |
Data |
| Multi-core SIMD processors |
Radeon 5870 GPU |
Data |
現代的GPU有很多獨立的運算核（processor）組成，在AMD GPU上就是stream core，這些core能夠執行SIMD操作（單指令，多數據），所以特別適合數據並行操作。通常GPU上執行一個任務，都是把任務中的數據分配到各個獨立的core中執行。
在GPU上，我們一般通過循環展開，Loop strip mining 技術，來把串行代碼改成並行執行的。比如在CPU上，如果我們實現一個向量加法，代碼通常如下：
1: for(i = 0; i < n; i++)
2: {
3: C[i] = A[i] + B[i];
4: }
在GPU上，我們可以設置n個線程，每個線程執行一個加法，這樣大大提高了向量加法的並行性。
1: __kernel void VectorAdd(__global const float* a, __global const float* b, __global float* c, int n)
2: {
3: int i = get_global_id(0);
4: c[i] = a[i] + b[i];
5: }
上面這個圖展示了向量加法的SPMD（單指令多線程）實現，從圖中可以看出如何實施Loop strip mining 操作的。
GPU的程序一般稱作Kernel程序，它是一種SPMD的編程模型（the Single Program Multiple Data ）。SPMD執行同一段代碼的多個實例，每個實例對數據的不同部分進行操作。
在數據並行應用中，用loop strip mining來實現SPMD是最常用的方法：
- 在分布式系統中，我們用Message Passing Interface (MPI)來實現SPMD。
- 在共享內存並行系統中，我們用POSIX線程來實現SPMD。
- 在GPU中，我們就是用Kernel來顯現SPMD。
在現代的CPU上，創建一個線程的開銷還是很大的，如果要在CPU上實現SPMD，每個線程處理的數據塊就要盡量大，做更多的事情，以便減少平均線程開銷。但在GPU上，都是輕量級的線程，創建、調度線程的開銷比較小，所以我們可以做到把循環完全展開，一個線程處理一個數據。
GPU上並行編程的硬件一般稱作SIMD。通常，發射一條指令后，它要在多個ALU單元中執行（ALU的數量即使SIMD的寬度），這種設計減少了控制流單元以及ALU相關的其他硬件數量。
SIMD的硬件如下圖所示：
在向量加法中，寬度為4的SIMD單元，可以把整個循環分為四個部分同時執行。在工人摘蘋果的例子中，工人的雙手類似於SIMD的寬度為2。另外，我們要知道，現在的GPU硬件上都是基於SIMD設計，GPU硬件隱式的把SPMD線程映射到SIMD core上。對開發有人員來說，我們並不需要關注硬件執行結果是否正確，我們只需要關注它的性能就OK了。
CPU一般都支持並行級的原子操作，這些操作保證不同的線程讀寫數據，相互之間不會干擾。有些GPU支持系統范圍的並行操作，但會有很大開銷，比如Global Memory的同步。