OpenCL 簡介


時間剛好是 Siggraph08,所以最近一堆電腦繪圖方面的東西都出來了。像前幾天的 nVidia 的 GeForce Force WithinOpenGL 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 KernelCompute ProgramCompute Kernel 基本上類似於 CUDA 定義的 kernel,是最基本的的計算單元;而 Compute Program 則是 Compute Kernel 和內建函式的集合,類似一個動態函式庫(看不懂什麼意思…)。而應用程式則是會把要執行的 compute kernel 排序,在執行時則可以依序或不依照順序來執行。

在平行計算方面,有 data-paralleltask-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 memorylocal memoryconstant memoryglobal memory 四種;架構上大概是像右邊的圖一樣。不過由於這部分也就只有一張投影片,所以有的部分,也不是很清楚就是了。


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

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

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

第一種就是指定記憶體使用的記憶體空間的四個 qualifier__private__local__constant__global(雖然投影片沒提及,但實際上應該還有指定函式的 qualifier?)。第二種,則是內建的資料型別;像是純量、向量和影像(image2d_timage3d_tsampler_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 比較好奇的。


OpenCL 相關文章目錄

對「OpenCL 簡介」的想法

發表迴響

在下方填入你的資料或按右方圖示以社群網站登入:

WordPress.com 標誌

您的留言將使用 WordPress.com 帳號。 登出 /  變更 )

Google photo

您的留言將使用 Google 帳號。 登出 /  變更 )

Twitter picture

您的留言將使用 Twitter 帳號。 登出 /  變更 )

Facebook照片

您的留言將使用 Facebook 帳號。 登出 /  變更 )

連結到 %s

This site uses Akismet to reduce spam. Learn how your comment data is processed.