OpenCL Basic

title: Dot Product
class: center, middle

## OpenCL Basic ##

## Dot Product Demo ##

### OpenCL 向量內積 示範 ###

R04922067 楊翔雲, Morris

NTU CSIE

---

## 如何編譯 OpenCL? ##

```bash
gcc <...> -lOpenCL
```

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

```bash
echo $LD_LIBRARY_PATH
```

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

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

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

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

---

## 如何使用 OpenCL 標頭檔 ##

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

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

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

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

```bash
gcc -xc -E -v -
```

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

---

## 如果你使用 vim ##

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

下載 [opencl.vim](http://www.vim.org/scripts/script.php?script_id=3157) 後，移動到 `/usr/share/vim/vimXX/syntax/`

### .vimrc ###

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

```vim
au BufReadPost *.cl set syntax=opencl
```

---

## 裝置資訊 ##

### PCIe ###

```bash
lspci | grep -i "Nvidia"
```

### OpenCL ###

```bash
clinfo
```

### Nvidia ###

```bash
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
```

---

## 實驗上的重要資訊 ##

```bash
$ 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
```bash
cat /usr/include/CL/cl.h | grep "\-1"
```

---

## 進入正題 ##

---

## 易犯錯的部分 ##

### 讀取 kernel.cl 檔案 ###

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

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

```c
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`](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clCreateContext.html) 時只使用一個 GPU 即可，每多一個慢幾百毫秒！且還多用數個 thread 控制！

```c
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)
```

```c
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](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clCreateBuffer.html) 即可。

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

搭配 [clReleaseMemObject](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clReleaseMemObject.html) 使用。若驅動程式沒寫好，即使砍了主程式，GPU Memory 仍然會有問題，之前舊版曾遇過到 Memory Leak .. 等諸多問題。

```c
cl_int clReleaseMemObject(cl_mem memobj);
```

---

## 運行效率問題 - Buffer ##

題目要求的 `$N \le 16777216$`

```c
// -- 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`

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

`CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR` 並不會自動寫回 Host

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

請自己下 `clEnqueueReadBuffer` 讀回來

```c
clEnqueueReadBuffer(...)
```

---

## 如何設定常數變數 - Kernel ##

### main.c ###

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

### kernel.cl ###

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

---

## 如何動態設定 Local Memory ##

固定大小的寫法如下：

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

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

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

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

---

## 如何設置 local size ##

* 請參考 [CUDA Warps and Occupancy](http://on-demand.gputechconf.com/gtc-express/2011/presentations/cuda_webinars_WarpsAndOccupancy.pdf)

* 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](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clEnqueueNDRangeKernel.html) 自動調整，則設置 `local_work_size = NULL`

```c
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](http://docs.nvidia.com/gameworks/content/developertools/desktop/analysis/report/cudaexperiments/sourcelevel/memorytransactions.htm)

* Bank conflict

* Memory coalescing

* Register file capacity per SM

* Data transfer between CPU and GPU

* [PTX ISA instruction optimization](http://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-ld)

* ...

---

![](http://i.imgur.com/bJhmVvr.jpg)

---

![](http://i.imgur.com/s8NlOr9.jpg)

---

![](http://i.imgur.com/EaemMtn.jpg)

---

## 如果你寫過 CUDA ##

| CUDA    | OpenCL    |
|:---------|:-----------|
| Thread  | Work item |
| Block   | Work group|
| Grid    | NDRange (index space) |

---

## 如果你寫過 CUDA ##

| CUDA    | OpenCL    |
|:---------|:-----------|
| `threadIdx.x`  | `get_local_id(0)` |
| `blockIdx.x*blockDim.x + threadIdx.x`   | `get_global_id(0)`|
| `gridDim.x` | `get_num_groups(0)` |
| `blockIdx.x` | `get_group_id(0)` |
| `blockDim.x` | `get_local_size(0)` |
| `gridDim.x*blockDIm.x` | `get_global_size(0)` |

---

![](http://i.imgur.com/aEvsQLm.jpg)