如何編譯 OpenCL?

gcc <...> -lOpenCL

• 請注意 -lOpenCL 在連結函式庫編譯上的順序，在這裡請放在最後面，

echo $LD_LIBRARY_PATH

• -L/<path> 可以幫你設定特別的路徑位置找到 Library

想偷懶不加 -L/，可以加以下的路徑到 ~/.bashrc

export LD_LIBRARY_PATH=/usr/local/cuda-8.0/lib64/:$LD_LIBRARY_PATH

miwa 工作站上，已經幫你解決了

如何使用 OpenCL 標頭檔

#ifdef _APPLE_
#include <OpenCL/OpenCL.h>
#else
#include <CL/cl.h>
#endif

• 我們在 Linux 下開發，直接使用 #include <CL/cl.h> 即可

• -I/<path> 可以幫你設定指定的標頭檔路徑

你可以藉由下述指令查看標頭檔引用的路徑

gcc -xc -E -v -

miwa 工作站上，已經幫你解決了

如果你使用 vim

kernel code 的語法高亮可能會對你有點痛苦

下載 opencl.vim 後，移動到 /usr/share/vim/vimXX/syntax/

.vimrc

有時候設定 syntex=c 也足以應付

au BufReadPost *.cl set syntax=opencl

裝置資訊

PCIe

lspci | grep -i "Nvidia"

OpenCL

clinfo

Nvidia

nvidia-smi                              // show information
sudo nvidia-smi -pm 1                   // on persistence mode
sudo nvidia-smi -pm 0                   // off persistence mode
watch -n 0.1 nvidia-smi                 // watch per 0.1 second

實驗上的重要資訊

$ clinfo
  Max work items dimensions:                     3
    Max work items[0]:                           1024
    Max work items[1]:                           1024
    Max work items[2]:                           64
...
  Address bits:                                  64
...
  Local memory size:                             49152
...
  Error correction support:                      0
  Unified memory for Host and Device:            0
...
  Queue properties:
    Out-of-Order:                                Yes

如果本身就沒支援 ... 就 ...

運行錯誤

完全沒有錯誤資訊？

• 請換個裝置再試試

• 請想好算法再試一次

有 Error Code

• 請到 CL/cl.h 下找相關資訊

• 如果要搜尋 Error Code = -1

  cat /usr/include/CL/cl.h | grep "\-1"

易犯錯的部分

讀取 kernel.cl 檔案

• 沒有察覺到檔案讀取不完整

• 別打錯檔案名字，批改娘上會替你重新命名檔案名稱

static char clSrc[32767];
FILE *codefin = fopen(clSrcFile, "r");
assert(codefin != NULL);
size_t clSrcLen = fread(clSrc, 1, 32767, codefin);
assert(feof(codefin) != 0); // check

運行效率問題 - Context

如果只使用一個 GPU 裝置，clCreateContext 時只使用一個 GPU 即可，每多一個慢幾百毫秒！且還多用數個 thread 控制！

cl_context clCreateContext(cl_context_properties *properties,
    cl_uint num_devices,
    const cl_device_id *devices,
    void *pfn_notify (const char *errinfo, const void *private_info, size_t cb, void *user_data),
    void *user_data,
    cl_int *errcode_ret)

clStat = clGetPlatformIDs(1, &clPlatID, &clPlatN);
CheckFailAndExit(clStat);
clStat = clGetDeviceIDs(clPlatID, CL_DEVICE_TYPE_GPU, MAXGPU, clGPUID, &clGPUN);
CheckFailAndExit(clStat);
*clCtx = clCreateContext(NULL, 1, clGPUID, NULL, NULL, &clStat);  // oneeeeeee!
CheckFailAndExit(clStat);

運行效率問題 - Buffer

我們有多筆測資，每次共用同一個足夠大的 buffer，運行 一次 clCreateBuffer 即可。

cl_mem clCreateBuffer(cl_context context,
    cl_mem_flags flags,
    size_t size,
    void *host_ptr,
    cl_int *errcode_ret)

搭配 clReleaseMemObject 使用。若驅動程式沒寫好，即使砍了主程式，GPU Memory 仍然會有問題，之前舊版曾遇過到 Memory Leak .. 等諸多問題。

cl_int clReleaseMemObject(cl_mem memobj);

運行效率問題 - Buffer

題目要求的 $N \le 16777216$

// -- create all buffers
cl_mem_flags clOutBuffFlag = CL_MEM_READ_WRITE;
*clMemOut = clCreateBuffer(*clCtx, clOutBuffFlag, sizeof(uint32_t)*MAXN,
        NULL, &clStat);
CheckFailAndExit(clStat);

• clCreateBuffer 很耗時間

• clEnqueueReadBuffer 很耗時間

• 傳輸和使用量越少越好

運行行為問題 - Buffer

如果裝置沒有支援 Unified memory for Host and Device

$ clinfo
...
  Unified memory for Host and Device:            0
...

CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR 並不會自動寫回 Host

cl_mem bufferC = 
  clCreateBuffer(context, 
     CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, ...);

請自己下 clEnqueueReadBuffer 讀回來

clEnqueueReadBuffer(...)

如何設定常數變數 - Kernel

main.c

int N = 128;
clStat = clSetKernelArg(*clKrn, 0, sizeof(cl_int), (void *) &N);
N = N/2; // N = 64
  ...

kernel.cl

__kernel void vecdot(int N, ...) {  // N = 128
  ...
}

如何動態設定 Local Memory

固定大小的寫法如下：

__kernel void vecdot(...) {
  __local unsigned int localArr[128];
  ...
}

如果不想固定大小，可以從 host 透過 clSetKernelArg 設定。

clStat = clSetKernelArg(*clKrn, 0, sizeof(unsigned int)*128, NULL);
                                   ^以 bytes 宣告整個陣列大小
                                                             ^用 NULL 表示 local memory 的參數

__kernel void vecdot( __local unsigned int localArr[]) {
  ...
}

如何設置 local size

• 請參考 CUDA Warps and Occupancy

• SM - Streaming multi-processors

• 每一個 SM 包含 32 個 processing cores

• 32 條 thread 組成一個 warp

• GTX 1080 有 20 個 SM (compute unit in OpenCL)

• 相當於最多 640 條 thread 同時運行

如何設置 local size

• 充分利用 warp

• 充分利用 warp scheduling

• 充分利用 SM

大致上

• LOCAL_SIZE >= WARP >= 32 效能就非常明顯地改善

• 讓 OpenCL 自動調整，則設置 local_work_size = NULL

cl_int clEnqueueNDRangeKernel (...,
   const size_t *local_work_size, ...)

內積計算

• 由於 $N \le 16777216 = 2^{24}$ 且傳輸量不可以太多

• 我們希望一個 thread 負責一個 chunk 的內積總和

• 隨著 chunk 增加所需要的 thread 就越少

• 例如 $\text{chunk} = 256 = 2^8$ 時，需要 $2^{16}$ 條 thread

• 一個 group 為充分使用 warp，安排 256 work items

其他效能細節

• Memory transactions

• Bank conflict

• Memory coalescing

• Register file capacity per SM

• Data transfer between CPU and GPU

• PTX ISA instruction optimization

• ...

如果你寫過 CUDA

如果你寫過 CUDA