OpenCL Advanced

title: Multiple Devices
class: center, middle

## OpenCL Advanced ##

## Multiple Devices ##

### 注意事項 ###

R04922067 楊翔雲, Morris

NTU CSIE

---

## NVIDIA GPU ##

---

## 那麼一點特別 ##

* `clEnqueue{Read|Write}*`
  * 官方有提供 blocking / non-blocking 控制參數

* 沒意外地，AMD/Intel 按照規格完成

* 在 NVIDIA 上，**大部分** 屬於 blocking 操作

紅字部分在未來也許會更動

---

## 那麼一點特別 ##

* `clEnqueueNDRangeKernel` 
  * 沒有明確規範 blocking / non-blocking

* 在 AMD/Intel 上，大部分屬於 non-blocking 操作

* 在 NVIDIA 上，**大部分** 屬於 blocking 操作

紅字部分在未來也許會更動

---

## 多少有點疑惑 ##

### 你們在實作同一個 OpenCL？ ###

### ？？ ###

---

## OpenCL Execution Model ##

### 尋找遺失的實作細節 ###

* 它是我們這堂課可能還未理解的部分

* 在不同廠牌上的實作撲朔迷離

* 如果你還沒看過 [Advanced OpenCL Event Model Usage](http://sa09.idav.ucdavis.edu/docs/SA09-opencl-dg-events-stream.pdf)

* 現在讓你看看

---

## 讀後感 ##

* 便利性相當強大

* Separate/Cooperative Multi-Device Usage Model  
似乎就是我們所要需要的部分

* Combined Memory Pool  
在 Cooperative 部分有一個定義不明確的部分，請**小心**服用

---

## 除錯 Debug ##

---

## 錯誤案例 Race Condition ##

當使用多個 GPU 運作時，為避開 blocking 操作而使用 OpenCL

```c
cl_kernel clKrn;
...

#pragma omp parallel for firstprivate(clKrn) // what's firstprivate ?
for (int i = 0; i < Q; i++) {
    ...
    clStat = clSetKernelArg(clKrn, 0, sizeof(cl_uint), (void *) &N);
    ...//                   ^race condition
}
```

請建立多個 kernel object，如下

```c
cl_kernel clKrn[DEVICE_NUM];
...

#pragma omp parallel for
for (int i = 0; i < Q; i++) {
    ...
    clStat = clSetKernelArg(clKrn[DEVICE_ID], 0, sizeof(cl_uint), (void *) &N);
    ...//                   ^safe
}
```

---

## 錯誤案例 Error Argument Type ##

* 別忘了物件和位址的差別

* 請多下 `gcc -Wall` 幫助你檢查這些強制轉型

* 查好 API，小心行事

```c
clCreateCommandQueue(cl_context, cl_device_id, cl_command_queue_properties, cl_int*);
                              // ^object
clBuildProgram(cl_program, cl_uint num_devices, const cl_device_id *device_list,
                                // ^number                         ^begin address, not object
    const char*, void (*)(cl_program, void *user_data), void *)
```

```c
cl_device_id    clGPUID[MAXGPU];
...
clBuildProgram(clPrg[i], 1,   clGPUID+i, NULL, NULL, NULL);  // ok
clBuildProgram(clPrg[i], 1, &clGPUID[i], NULL, NULL, NULL);  // ok
clBuildProgram(clPrg[i], 1,  clGPUID[i], NULL, NULL, NULL);  // no, get far away from me.
```

---

### 錯誤案例 Parallel Kernel Execution ###

* 如果你要在 OpenCL 模擬出 CUDA 的 cudaStream，當心！

* 手動模擬存在許多無法預期的結果

```c
cl_kernel           clKrn;
cl_command_queue    clQue;

#pragma omp parallel for
for (int i = 0; i < Q; i++) {
    ...
    clEnqueueNDRangeKernel(clQue, clKrn, 1, globalOffset,
                globalSize, localSize, 0, NULL, NULL);
    clEnqueueReadBuffer(...);
    // I don't know what's the status happened
    // horrible !!!!!!
    ...
}
```

---

## 助教，為什麼程序沒有快？ ##

### 如何看待**你**的觀察 ###

---

## 如何偷窺 ##

+-----------------------------------------------------------------------------+
| Processes:                                                       GPU Memory |
|  GPU       PID  Type  Process name                               Usage      |
|=============================================================================|
|    0     19246    C   ./main3                                        101MiB |
|    1     19246    C   ./main3                                        113MiB |
|    2     19246    C   ./main3                                         64MiB |
+-----------------------------------------------------------------------------+
```

---

### 如何仔細偷窺 NVIDIA 在 OpenCL 的行為 ###

* 獻給使用舊版 CUDA 的使用者

* 請下載 [假 nvprof - Morris 改](https://gist.github.com/morris821028/6bb5709ac454c999e879b67cd7bb201e)

```bash
$ COMPUTE_PROFILE=1 COMPUTE_PROFILE_CONFIG=${CONFIG_FILE} <execute_command>
```

* 早期提供環境變數，運行後便產生 .log 檔案於目錄下

* 若您使用 CUDA 8.0 的最新版本，  
將沒辦法借助硬體得到 OpenCL 運行情況 (功能已拔除)

* 用 NVVP 也無法運行 OpenCL 的執行檔，僅限 CUDA

---

## Overhead ##

* Speedup/Efficiency 計算時，請說明有哪些因素，  
哪一種才是**合適**的？

* `clCreateContext();`  
在一個 GPU 和多個 GPU 上建立 Context 的消耗時間不同。

* 資料傳輸的帶寬 [PCI Express](https://en.wikipedia.org/wiki/PCI_Express)
  * 是否會是瓶頸？

* 雙通道技術 (Full Duplex)？

* ...

---

## 助教，程序 Debug 行為好怪 ##

### ... ###

---

## 在 kernel 運行 `printf` ##

* 在新的版本都支援在 kernel function 中使用 `printf`

* 印出資訊過多時，一部分輸出訊息可能會看不到

* 還是錯誤！

```c
__kernel void vecdot(...) {
  // sometimes it work, sometimes it not work
  printf("local id: %d, global id: %d\n", get_local_id(0), get_global_id(0));
}
```

Why? 請注意型態

```c
__kernel void vecdot(...) {
  // sometimes it work, sometimes it not work
  printf("local id: %d, global id: %d\n", (int) get_local_id(0), (int) get_global_id(0));
}
```

---

## 助教，為什麼程序會快？ ##

### ... ###

---

## 可怕的參數調校 ##

### Fine Grain or Coarse Grain ###

* 同時運作的數量有多少？剩下的 **那一輪** 的 **執行比例** 是多少？

* Coarse Grain 造成大測資達到最快，小測資效能不合乎推論

* Fine Grain 造成大小測資按理論復雜度成長，卻不是最快

```bash
$ clinfo
  Device Type:                                   CL_DEVICE_TYPE_GPU
  Device ID:                                     4318
  Max compute units:                             20                 // important!
```