bbb1206
一般般會員
- 已加入
- 6/1/10
- 訊息
- 175
- 互動分數
- 0
- 點數
- 0
OpenCL 簡介
OpenCL 是由 Khronos Group 針對異質性計算裝置（heterogeneous device）進行平行化運算所設計的標準 API 以及程式語言。所謂的「異質性計算裝置」，是指在同一個電腦系統中，有兩種以上架構差異很大的計算裝置，例如一般的 CPU 以及顯示晶片，或是像 CELL 的 PPE 以及 SPE。目前，最為常見的就是所謂的 GPGPU 應用，也就是利用一般的顯示晶片（即 GPU）進行 3D 繪圖以外的計算工作。
過去 GPGPU 的應用，有各種不同的使用方式。最早的 GPGPU，多半是直接透過 3D 繪圖的 API 進行，例如 OpenGL 或 D3D 的 HLSL（High Level Shading Language）。但是，這樣做有很多缺點，主要是即使想要進行的運算和 3D 繪圖無關，仍需要處理很多 3D 繪圖方面的動作（例如建立 texture，安排 render-to-texture 動作等等）。這讓 GPGPU 變得十分複雜。後來開始有些嘗試把這些 3D 繪圖部份隱藏起來的想法，例如由 Stanford 大學設計的 BrookGPU，可以透過不同的 backend 將 Brook 程式轉換成由 CPU、Direct3D、或 OpenGL 來執行。另外，也有各家顯示卡廠商自行開發的系統，包括 ATI 針對其產品設計的 Close to Metal（以及後來的 AMD Stream），以及 NVIDIA 的 CUDA。Microsoft 也在 DirectX 11 中加入了特別為 GPGPU 設計的 DirectCompute。
由於各家廠商的 GPGPU 方案都是互不相容的（例如 AMD Stream 的程式無法在 NVIDIA 的顯示晶片上執行，而 CUDA 的程式也不能在 AMD 的顯示晶片上執行），這對 GPGPU 的發展是不利的，因為程式開發者必須為不同廠商的顯示晶片分別撰寫程式，或是選擇只支援某個顯示晶片廠商。由於顯示晶片的發展愈來愈彈性化，GPGPU 的應用範圍也增加，因此 Apple 決定提出一個統一的 GPGPU 方案。這個方案得到包括 AMD、IBM、Intel、NVIDIA 等相關廠商的支持，並很快就交由 Khronos Group 進行標準化。整個計畫只花了五個月的時間，並在 2008 年十二月時正式公開。第一個正式支援 OpenCL 的作業系統是 Apple 的 MacOS X 10.6 "Snow Leopard"。AMD 和 NVIDIA 也隨後推出了在 Windows 及 Linux 上的 OpenCL 實作。IBM 也推出了支援 CELL 的 OpenCL 實作。
OpenCL 的主要設計目的，是要提供一個容易使用、且適用於各種不同裝置的平行化計算平台。因此，它提供了兩種平行化的模式，包括 task parallel 以及 data parallel。目前 GPGPU 的應用，主要是以 data parallel 為主，這裡也是以這個部份為主要重點。所謂的 data parallel，指的是有大量的資料，都進行同樣的處理。這種形式的平行化，在很多工作上都可以見到。例如，影像處理的程式，經常要對一個影像的每個 pixel 進行同樣的動作（例如 Gaussian blur）。因此，這類工作很適合 data parallel 的模式。
OpenCL 的架構
OpenCL 包括一組 API 和一個程式語言。基本上，程式透過 OpenCL API 取得 OpenCL 裝置（例如顯示晶片）的相關資料，並將要在裝置上執行的程式（使用 OpenCL 程式語言撰寫）編繹成適當的格式，在裝置上執行。OpenCL API 也提供許多裝置控制方面的動作，例如在 OpenCL 裝置上取得一塊記憶體、把資料從主記憶體複製到 OpenCL 裝置上（或從 OpenCL 裝置上複製到主記憶體中）、取得裝置動作的資訊（例如上一個程式執行所花費的時間）等等。
例如，我們先考慮一個簡單的工作：把一群數字相加。在一般的 C 程式中，可能是如下：
float a[DATA_SIZE];
float b[DATA_SIZE];
float result[DATA_SIZE];
// ...
for(int i = 0; i < DATA_SIZE; i++) {
result = a + b;
}
在 OpenCL 中，則大致的流程是：
把 OpenCL 裝置初始化。
在 OpenCL 裝置上配置三塊記憶體，以存放 a、b、c 三個陣列的資料。
把 a 陣列和 b 陣列的內容，複製到 OpenCL 裝置上。
編譯要執行的 OpenCL 程式（稱為 kernel）。
執行編譯好的 kernel。
把計算結果從 OpenCL 裝置上，複製到 result 陣列中。
透過 data parallel 的模式，這裡的 OpenCL 程式非常簡單，如下所示：
__kernel void adder(__global const float* a, __global const float* b, __global float* result)
{
int idx = get_global_id(0);
result[idx] = a[idx] + b[idx];
}
在一般的版本中，是透過一個迴圈，執行 DATA_SIZE 次數的加法動作。而在 OpenCL 中，則是建立 DATA_SIZE 個 work item，每個 work item 都執行上面所示的 kernel。可以看到，OpenCL 程式語言和一般的 C 語言非常類似。__kernel 表示這個函式是在 OpenCL 裝置上執行的。__global 則表示這個指標是在 global memory 中（即 OpenCL 裝置上的主要記憶體）。而 get_global_id(0) 會傳回 work item 的編號，例如，如果有 1024 個 work item，則編號會分別是 0 ~ 1023（實際上編號可以是二維或三維，但在這裡先只考慮一維的情形）。
要如何讓上面這個簡單的 OpenCL kernel 實際在 OpenCL 裝置上執行呢？這就需要透過 OpenCL API 的幫助了。以下會一步一步說明使用 OpenCL API 的方法。
OpenCL 是由 Khronos Group 針對異質性計算裝置（heterogeneous device）進行平行化運算所設計的標準 API 以及程式語言。所謂的「異質性計算裝置」，是指在同一個電腦系統中，有兩種以上架構差異很大的計算裝置，例如一般的 CPU 以及顯示晶片，或是像 CELL 的 PPE 以及 SPE。目前，最為常見的就是所謂的 GPGPU 應用，也就是利用一般的顯示晶片（即 GPU）進行 3D 繪圖以外的計算工作。
過去 GPGPU 的應用，有各種不同的使用方式。最早的 GPGPU，多半是直接透過 3D 繪圖的 API 進行，例如 OpenGL 或 D3D 的 HLSL（High Level Shading Language）。但是，這樣做有很多缺點，主要是即使想要進行的運算和 3D 繪圖無關，仍需要處理很多 3D 繪圖方面的動作（例如建立 texture，安排 render-to-texture 動作等等）。這讓 GPGPU 變得十分複雜。後來開始有些嘗試把這些 3D 繪圖部份隱藏起來的想法，例如由 Stanford 大學設計的 BrookGPU，可以透過不同的 backend 將 Brook 程式轉換成由 CPU、Direct3D、或 OpenGL 來執行。另外，也有各家顯示卡廠商自行開發的系統，包括 ATI 針對其產品設計的 Close to Metal（以及後來的 AMD Stream），以及 NVIDIA 的 CUDA。Microsoft 也在 DirectX 11 中加入了特別為 GPGPU 設計的 DirectCompute。
由於各家廠商的 GPGPU 方案都是互不相容的（例如 AMD Stream 的程式無法在 NVIDIA 的顯示晶片上執行，而 CUDA 的程式也不能在 AMD 的顯示晶片上執行），這對 GPGPU 的發展是不利的，因為程式開發者必須為不同廠商的顯示晶片分別撰寫程式，或是選擇只支援某個顯示晶片廠商。由於顯示晶片的發展愈來愈彈性化，GPGPU 的應用範圍也增加，因此 Apple 決定提出一個統一的 GPGPU 方案。這個方案得到包括 AMD、IBM、Intel、NVIDIA 等相關廠商的支持，並很快就交由 Khronos Group 進行標準化。整個計畫只花了五個月的時間，並在 2008 年十二月時正式公開。第一個正式支援 OpenCL 的作業系統是 Apple 的 MacOS X 10.6 "Snow Leopard"。AMD 和 NVIDIA 也隨後推出了在 Windows 及 Linux 上的 OpenCL 實作。IBM 也推出了支援 CELL 的 OpenCL 實作。
OpenCL 的主要設計目的，是要提供一個容易使用、且適用於各種不同裝置的平行化計算平台。因此，它提供了兩種平行化的模式，包括 task parallel 以及 data parallel。目前 GPGPU 的應用，主要是以 data parallel 為主，這裡也是以這個部份為主要重點。所謂的 data parallel，指的是有大量的資料，都進行同樣的處理。這種形式的平行化，在很多工作上都可以見到。例如，影像處理的程式，經常要對一個影像的每個 pixel 進行同樣的動作（例如 Gaussian blur）。因此，這類工作很適合 data parallel 的模式。
OpenCL 的架構
OpenCL 包括一組 API 和一個程式語言。基本上，程式透過 OpenCL API 取得 OpenCL 裝置（例如顯示晶片）的相關資料，並將要在裝置上執行的程式（使用 OpenCL 程式語言撰寫）編繹成適當的格式，在裝置上執行。OpenCL API 也提供許多裝置控制方面的動作，例如在 OpenCL 裝置上取得一塊記憶體、把資料從主記憶體複製到 OpenCL 裝置上（或從 OpenCL 裝置上複製到主記憶體中）、取得裝置動作的資訊（例如上一個程式執行所花費的時間）等等。
例如，我們先考慮一個簡單的工作：把一群數字相加。在一般的 C 程式中，可能是如下：
float a[DATA_SIZE];
float b[DATA_SIZE];
float result[DATA_SIZE];
// ...
for(int i = 0; i < DATA_SIZE; i++) {
result = a + b;
}
在 OpenCL 中，則大致的流程是：
把 OpenCL 裝置初始化。
在 OpenCL 裝置上配置三塊記憶體，以存放 a、b、c 三個陣列的資料。
把 a 陣列和 b 陣列的內容，複製到 OpenCL 裝置上。
編譯要執行的 OpenCL 程式（稱為 kernel）。
執行編譯好的 kernel。
把計算結果從 OpenCL 裝置上，複製到 result 陣列中。
透過 data parallel 的模式，這裡的 OpenCL 程式非常簡單，如下所示：
__kernel void adder(__global const float* a, __global const float* b, __global float* result)
{
int idx = get_global_id(0);
result[idx] = a[idx] + b[idx];
}
在一般的版本中，是透過一個迴圈，執行 DATA_SIZE 次數的加法動作。而在 OpenCL 中，則是建立 DATA_SIZE 個 work item，每個 work item 都執行上面所示的 kernel。可以看到，OpenCL 程式語言和一般的 C 語言非常類似。__kernel 表示這個函式是在 OpenCL 裝置上執行的。__global 則表示這個指標是在 global memory 中（即 OpenCL 裝置上的主要記憶體）。而 get_global_id(0) 會傳回 work item 的編號，例如，如果有 1024 個 work item，則編號會分別是 0 ~ 1023（實際上編號可以是二維或三維，但在這裡先只考慮一維的情形）。
要如何讓上面這個簡單的 OpenCL kernel 實際在 OpenCL 裝置上執行呢？這就需要透過 OpenCL API 的幫助了。以下會一步一步說明使用 OpenCL API 的方法。
