OpenCL 簡介

時間剛好是 Siggraph08，所以最近一堆電腦繪圖方面的東西都出來了。像前幾天的 nVidia 的 GeForce Force Within、OpenGL 3.0 的 ForceWare 177.89、ATI 的 4850×2 和 4870×2，以及 Intel 的 Larrabee，都算最近 Siggraph 的東西吧∼

而雖然不見得和圖學有關，但是 GPGPU 最近也是 Siggraph 的一個熱門話題；像是 ATi 的 CTM、nVidia 的 CUDA，就是這兩家處理器大廠各自推出的 GPGPU 開發方案。不過，日前也有提過，已經有在規畫一個稱做 OpenCL（Open Computing Language，開放計算語言）的標準了∼由於 nVidia、AMD、Intel 等相關公司都在其列，或許以後有可能 OpenCL 會像是 GLSL 或 HLSL 一樣，成為多核心程式的一個共通標準？而日前也有消息指出，AMD 似乎有意放棄 CTM 而以 OpenCL 為主（參考《AMD ditches Close-To-Metal, focuses on DX11 and OpenCL》）。（話說，微軟沒有參予 OpenCL。不知道是不是打算以 DirectX 11 的 Compute Shader 來吃這塊市場呢？）

但是，OpenCL 的標準似乎也都還沒有什麼細節可以參考，也不知道到底會是怎樣？不過前幾天，在 Siggraph 08 的 Class 裡「Beyond Programmable Shading: Fundamentals」，到是出現了 OpenCL 的簡單的範例程式，可以讓大家一窺究竟了∼

這個 class 的投影片，可以在 http://s08.idav.ucdavis.edu/ 找到。其中除了 OpenCL 外，也包含了其他 CUDA、AMD Stream SDK、Larrabee、DirectX 11 Compute Shader 等等的各種 GPU 計算的資料。不過，還是先根據 OpenCL 的投影片，來看看他的程式架構吧∼

投影片的標題當然是 OpenCL，副標題則是「Parallel Computing on the GPU and CPU」，代表了他不只是 GPGPU 的程式語言，也還可以用在 CPU 上！而 OpenCL 設計的目的，除了是鍵立一個標準外，主要是希望可以：

• 用平易近人的語言，使用不同的計算支援
• 能夠在單一或多個處理器上平行進行
  • GPU、CPU、GPU CPU 或多 GPU
• 桌上型和手持裝置
• 設計成可以和圖型 API（如 OpenGL）一起運作

而 OpenCL 也是基於 C 的一個程式語言，在軟體上，分為 Platform Layer、Runtime、Compiler 三個部分：

• Platform Layer 用來管理計算裝置，提供啟始化裝置的界面，並用來建立 compute contexts 和 work-queues。
• Runtime 用來管理資源，並執行程式的 kernel
• 編譯器則是 ISO C99 的子集合，並加上了 OpenCL 特殊的語法。

在 OpenCL 的執行模型中，有所謂的 Compute Kernel 和 Compute Program。Compute Kernel 基本上類似於 CUDA 定義的 kernel，是最基本的的計算單元；而 Compute Program 則是 Compute Kernel 和內建函式的集合，類似一個動態函式庫（看不懂什麼意思…）。而應用程式則是會把要執行的 compute kernel 排序，在執行時則可以依序或不依照順序來執行。

在平行計算方面，有 data-parallel 和 task-parallel 兩種不同的計算模型。其中，data-parallel 的模式，就是 CUDA 的平行化的形式，這也是所有 OpenCL 計算裝置都必須要支援的！相較之下，task-parallel 模式，則比較接近一般在寫 CPU 程式的 threading 了∼這個模式，應該也不是所有裝置都會支援。

Data-Parallel Execution Model

在 data-parallel 的模型，定義了「N-Dimensional computation domain」（後面縮寫成「N-D domain」）。在這個 N-D domain 中，每一個獨立的項目，被稱做一個「work-item」（相當於 CUDA 的 thread），全部的 work-item 數量，則被定義為「global work size」。

每一個 work-item 會被平行地執行，同時也可以把幾個 work-item 作群組，成為「work-group」（相當於 CUDA 的 thread block）；而在 work-group 中的 work-item 可以互相溝通，也可以進行同步（synchronize）。而在 OpenCL 中，也是多個 work-gorup 同時被執行的。

Task-Parallel Execution Model

在文章中是說，某些計算裝置，像是 CPU，也可以執行 task-parallel 的 compute kernell。而在 OpenCL 的這個模式下，就是執行單一個 work-item 了∼

記憶體模型

而在 OpenCL 的記憶體模型方面，和 CUDA 一樣，分為很多種記憶體空間；不過，OpenCL 應該是沒有 CUDA 的 shared memory 和 texture memory 了。

在 OpenCL 裡的記憶體空間有 private memory、local memory、constant memory 和 global memory 四種；架構上大概是像右邊的圖一樣。不過由於這部分也就只有一張投影片，所以有的部分，也不是很清楚就是了。

上面都算是架構的東西，接下來，就是 OpenCL 語言的部分了∼

OpenCL 基本上是源自 ISO C99 的標準，和 CUDA 一樣，主要是針對 kernel 程式的部分而開發的。在這部分，一樣有加上一些限制，比如說 recursion、function pointers 都不能在 kernel 裡使用。

而除了標準 C 的語法外，OpenCL 也另外加上了一些其他的東西。

第一種就是指定記憶體使用的記憶體空間的四個 qualifier（__private、__local、__constant、__global）（雖然投影片沒提及，但實際上應該還有指定函式的 qualifier？）。第二種，則是內建的資料型別；像是純量、向量和影像（image2d_t、image3d_t、sampler_t）的相關型別，以及基本的型別轉換函式。

第三種，則是其他內建的函數，在投影片中，OpenCL 把這些函數分成必要的（Required）和選擇性的（Optional）兩種。必要的函式包括了：

• work-item functions
• math.h
• read and write image
• relational
• geometric functions
• synchronization functions

而選擇性的，則有：

• double precision
• atomics to global and local memory
• selection of rounding mode
• writes to image3d_t surface

不過呢，在投影片也就只供這些資料了。究竟有哪些東西呢？細節部分，大概要等到詳細的文件出來吧？不過，目前 Heresy 完全找不到其他的相關資料就是了…（連軟體都沒有，也沒的試了）

而在接下來的投影片，則是提供了一份用 OpenCL 寫的 FFT 的範例程式；和 CUDA 一樣，分為 Host code 和 Compute kernel 兩部分。

下面就是 Host Code 的部分：

// create a compute context with GPU device 
context = clCreateContextFromType(CL_DEVICE_TYPE_GPU); 
// create a work-queue queue = clCreateWorkQueue(context, NULL, NULL, 0);
  
// allocate the buffer memory objects 
memobjs[0] = clCreateBuffer( context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                             sizeof(float)*2*num_entries, srcA); 
memobjs[1] = clCreateBuffer( context, CL_MEM_READ_WRITE,
                             sizeof(float)*2*num_entries, NULL);

 
// create the compute program 
program = clCreateProgramFromSource(context, 1, &fft1D_1024_kernel_src, NULL); 
// build the compute program executable 
clBuildProgramExecutable(program, false, NULL, NULL); 
// create the compute kernel 
kernel = clCreateKernel(program, “fft1D_1024”);
  
// create N-D range object with work-item dimensions 
global_work_size[0] = n; 
local_work_size[0] = 64; 
range = clCreateNDRangeContainer(context, 0, 1, global_work_size, local_work_size);
  
// set the args values 
clSetKernelArg(kernel, 0, (void *)&memobjs[0], sizeof(cl_mem), NULL); 
clSetKernelArg(kernel, 1, (void *)&memobjs[1], sizeof(cl_mem), NULL); 
clSetKernelArg(kernel, 2, NULL, sizeof(float)*(local_work_size[0] 1)*16, NULL); 
clSetKernelArg(kernel, 3, NULL, sizeof(float)*(local_work_size[0] 1)*16, NULL);
  
// execute kernel 
clExecuteKernel(queue, kernel, NULL, range, NULL, 0, NULL);  

看了上面的程式，Heresy 個人的感覺是比 CUDA 來得更為繁瑣…而且光就投影片來說，Heresy 很多地方還是沒能清楚；像是「fft1D_1024_kernel_src」就沒有在投影片中出現他的相關資訊。

而 Compute Kernel 的部分，則是：

// This kernel computes FFT of length 1024. The 1024 length FFT is decomposed into 
// calls to a radix 16 function, another radix 16 function and then a radix 4 function 
__kernel void fft1D_1024( __global float2 *in, __global float2 *out,
                                     __local float *sMemx, __local float *sMemy) 
{
     int tid = get_local_id(0);
     int blockIdx = get_group_id(0) * 1024   tid;
     float2 data[16];
      // starting index of data to/from global memory
     in = in   blockIdx; out = out   blockIdx;
     globalLoads(data, in, 64); // coalesced global reads
     fftRadix16Pass(data); // in-place radix-16 pass
     twiddleFactorMul(data, tid, 1024, 0);
      // local shuffle using local memory
     localShuffle(data, sMemx, sMemy, tid, (((tid & 15) * 65)   (tid >> 4)));
     fftRadix16Pass(data); // in-place radix-16 pass
     twiddleFactorMul(data, tid, 64, 4); // twiddle factor multiplication
     localShuffle(data, sMemx, sMemy, tid, (((tid >> 4) * 64)   (tid & 15)));
      // four radix-4 function calls
     fftRadix4Pass(data);
     fftRadix4Pass(data   4);
     fftRadix4Pass(data   8);
     fftRadix4Pass(data   12);
      // coalesced global writes
     globalStores(data, out, 64); 
}

基本上，這邊就更慘了…在沒有文件的情況下，Heresy 也無法分辨哪些函式是 OpenCL 內建的（還是其實都是？）…所以，看看就好吧∼

到這，主要就是這份 OpenCL 投影片的資料了∼說實話，資料時在不夠，只能大概做些參考。真的要下去玩，應該要等到有軟體、有完整的程式可以下載安裝了。而之前有提過，OpenCL 可能是 CUDA 的子集合，不過在 Heresy 看來，在概念上基本上可以這樣說，可是在程式開發上，應該還是會有相當的差異。至於如果 OpenCL 這個標準化的東西起來後，對於 CUDA 會有多大的影響呢？這點應該會是 Heresy 比較好奇的。

原始發表：http://heresy.spaces.live.com/blog/cns!E0070FB8ECF9015F!4487.entry