CUDA

title: CUDA
class: center, middle

## CUDA ##

## Something ##

### 一些瑣事 ###

R04922067 楊翔雲, Morris

NTU CSIE

---

## 如何編譯 CUDA? ##

CUDA 編譯程序分了數個步驟

* 若要搭配 OpenMP，透過 `-Xcompiler "-fopenmp"`

* 若要設定 [PTX ISA](http://docs.nvidia.com/cuda/parallel-thread-execution/#axzz4iktzhbwR)，透過 `-arch` 參數

* 指定特定版本 (`-code`)，將 PTX 特別優化包成二進制的版本，否則透過 PTX 在 real time 編譯優化

> Parallel Thread eXecution virtual machine: PTX

```bash
$ nvcc -Xcompiler "-fopenmp" \
       -arch sm_30 \
       -code sm_30 \
       -std=c++14 \
       main.cu \
       -o main
```

---

## 如何更方便地觀察運行結果 ##

* 使用 NVIDIA Visual Profiler

* 運行一次執行檔，便可觀察每一步運行的時間和呼叫細節

```bash
$ nvvp
```

### 如何運行遠端 GUI 程序 ###

```bash
$ ssh -X <username>@<remote-server>
```

* Windows OS: MobaXterm / Xming / ...

* Mac OS: iTerm / ...

---

## Thrust Library ##

Develop high-performance applications rapidly with Thrust! @[Github](https://github.com/thrust/thrust)

### 高尚的競爭是一切卓越才能的源泉 ###

---

## What's Thrust? ##

* 類似 C++ 的標準函式庫 STL

* 面向高效能平行算法

* GPUs / multicore CPUs 可使用

* 與現有的技術結合
  * CUDA - NVIDIA

* TBB - Intel Threading Building Blocks

* OpenMP

> 好像有個傢伙缺席

---

## 如何切換平行環境/一鍵換裝 ##

不用改動任何一行程序，編譯時改動參數

### OpenMP ###

```bash
$ nvcc \
        -Xcompiler -fopenmp \
        -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_OMP \
        ...
```

### CUDA ###

```bash
$ nvcc \
        -Xcompiler -fopenmp \
        -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CUDA \
        ...
```

---

## 工作站 ##

* Thrust Library 是內建函數，有需要請自行查閱

```bash
/usr/local/cuda-8.0/include/thrust/
├── detail
│   ├── allocator
│   ├── complex
│   ├── config
│   ├── dispatch
│   ├── functional
│   ├── mpl
│   ├── range
│   ├── type_traits
│   └── util
├── iterator
│   └── detail
├── random
│   └── detail
└── system
    ├── cpp
    ├── cuda
    ├── detail
    ├── omp
    └── tbb
```

---

## Workshop ##

### [20016. Labeling String](https://judgegirl.csie.org/problem/0/20016) ###

### 平行算法設計 ###

#### 專注平行算法三十年 ####

---

## 題目描述 ##

* 給定長度為 `$N$` 的字串 `$S$`，`$S$` 只由英文字母和空白字元組成

* 對於每個位置計算，往前有多少連續英文字母

* 約束連續英文字母長度不超過 `$K \le 500$`

* 例如：

```c
S = [  a  vy  lu  rah   yhy    fibcl ]
P = [00100120012001230001230000123450]
```

---

## 循序算法 ##

* 輸入型態為 `char` (1 bytes)

* 輸出型態為 `int` (4 bytes)

* 類似經典的 **最長平台問題**

* 時間複雜度 `$O(N)$`

```c
void serial_trans(const char S[], int P[], int N) {
    for (int i = 0, sum = 0; i < N; i++) {
        if (isalnum(S[i]))
            sum++;
        else
            sum = 0;
        P[i] = sum;
    }
}
```

---

## CUDA 常見錯誤 ##

**執行時間過慢**

```c
void labeling(const char *cuStr, int *cuPos, int strLen) {
    // call kernel function / Thrust library
}
```

* 由於只能在指定函數內撰寫，請注意呼叫次數將影響
* 根據給定的最大值限制，只宣告一次 global memory，避免重複 malloc

### C style ###

```c
#define MAXSTRLEN (1<<24)
void labeling(const char *cuStr, int *cuPos, int strLen) {
    static int *P0 = NULL;
    if (P0 == NULL)
        cudaMalloc((void **) &P0, MAXSTRLEN * sizeof(int));
    ...
}
```

---

## CUDA 常見錯誤 ##

### C++ style ###

```cpp
namespace MM {
   class MemoryPool {
    public:
        // ...
        MemoryPool() { ... }
        ~MemoryPool() { ... }
    } pool;
}
```

---

## 粗魯的平行算法 ##

* 對於每個位置，往前計算有多少個連續英文字母

* 時間複雜度 `$O(NK/p + K)$`

```c
void par_trans(const char S[], int P[], int N) {
    #pragma omp parallel for
    for (int i = 0; i < N; i++) {
        int sum = 0;
        for (int j = i; j >= 0 && isalnum(S[j]); j--)
            sum++;
        P[i] = sum;
    }
}
```

---

## 細緻些的平行算法 ##

* 處理器不多的情況下，先前的算法效率不高

* 複雜度不變，每個執行緒負責一段，在這一段中運行循序算法

```c
const int BLK = 8;  // set BLK as you wish
void par_2stage_trans(const char S[], int P[], int N) {
    int M = (N+BLK-1)/BLK;
    #pragma omp parallel for
    for (int i = 0; i < M; i++) {
        int l = i*BLK, r = min((i+1)*BLK, N);
        int sum = 0;
        for (int j = l-1; j >= 0 && isalnum(S[j]); j--)
            sum++;
        for (int j = l; j < r; j++) {
            if (isalnum(S[j]))
                sum++;
            else
                sum = 0;
            P[j] = sum;
        }
    }
}
```

---

## 接下來的算法 ##

### 僅供參考 ###

---

## 優雅的平行算法 ##

為排版方便，索引值以尾數表示之

1. 將空白部分填入索引值，反之填入 `-1`，得 `P0`  
```c
i = [01234567890123456789012345678901]
S = [  a  vy  lu  rah   yhy    fibcl ]
P0= [01-34--78--12---678---2345-----1]
```

2. 對初步得到的位置資訊 `P0` 運行 Prefix Maximum，得 `P1`  
```c
P0= [01-34--78--12---678---2345-----1]
P1= [01134447888122226788882345555551]
```

3. 對於 `P1`，將每個元素拿索引值減去，得 `P`  
```c
P1= [01134447888122226788882345555551]
P = [00100120012001230001230000123450]
```

---

## 分析－優雅的平行算法 ##

* 第一步，`$ (i, S_i) \mapsto (i+1)[S_i \in \Sigma] -1$`，轉換 `$O(N/p)$`

* 第二步，前綴算法 `$ O(N \log K / p + \log K) $`

* 第三步，`$ (i, P_i) \mapsto i - P_i$`，轉換 `$O(N/p)$`

* 整體而言，不論連續字母多寡 `$O(N \log N / p + \log N)$`

---

## 實作－優雅的平行算法 ##

* 運行 kernel call 越少越好
  * 手動撰寫時，**第一步** 和 **第二步** 合為一項操作

* 讓并發 (concurrency) 效果更好，減少 SM 的空窗期

* 減少 global memory 的存取次數
  * 設計的算法中，時常會需要 `$P_i$` 和 `$S_i$` 作為依據

* 存取次數從 `$2N$` 降到 `$N$`，效能提升 `$\sim 2 $`倍

---

## Thrust－優雅的平行算法 ##

* 如何帶入索引值進行轉換？

```c
template<class T> struct MM {
    const char *base;
    MM(const char *base): base(base) {}
    __host__ __device__ T operator()(const T& index) const { 
        return TRANS(index, base[index]); 
    };
};
thrust::tabulate(..., MM<int32_t>(cuStr));
```

* 如何運行前綴最大值 Prefix Maximum？

```c
thrust::inclusive_scan(thrust::device, ..., thrust::maximum<int>());
```

* 另一種算法：如何一鍵完成？

```c
thrust::inclusive_scan_by_key(thrust::device, ...);
```

---

## 加速－特別約束 ##

* 上述的算法能應付任何情況的 `$K$`

* 加速！加速！有效地利用 `$K \le 500$`

* 那麼全手動，無法倚靠 Thrust 給予我們的便利性

* 由於 `$K \le 500$`，我們仿造先前的算法
  * 每個 block 處理每段 512 個元素

* 用 local memory 完成 `$O(K \log K)$` 的後綴最大值

* 連續的一段字母 **最多** 橫跨 2 個 block

---

## 算法－特別約束 ##

### 範例說明 ###

#### block 設定 ####

* 每個 block 處理 5 個元素

* 每一個 block 中有 5 個 thread 合作

---

## 算法－特別約束 ##

### 步驟 ###

每一個 block 按照先前提到的 **步驟一**、**步驟二** 完成。在步驟二中，`R1` 由每個 block 獨立計算前綴最大值預設最小值 `-1`

```c
i = [01234|56789|01234|56789|01234|56789|01]
S = [  a  |vy  l|u  ra|h   y|hy   | fibc|l ]
R0= [01-34|--78-|-12--|-678-|--234|5----|-1]
R1= [01134|--788|-1222|-6788|--234|55555|-1]
```

接著，進入 **步驟三**

```c
i = [01234|56789|01234|56789|01234|56789|01]
S = [  a  |vy  l|u  ra|h   y|hy   | fibc|l ]
R1= [01134|--788|-1222|-6788|--234|55555|-1]
R2= [00100|12001|10012|10001|12000|01234|10]
```

---

## 算法－特別約束 ##

### 修正 ###

連續字母跨 block，**可能** 需要修正的 block 以星號 (*) 標注

```c
i = [01234|56789|01234|56789|01234|56789|01]
S = [  a  |vy  l|u  ra|h   y|hy   | fibc|l ]
R2= [00100|12001|10012|10001|12000|01234|10]
           *     *     *     *           *
```

星號部分往前繼找最後一個值 (非零) 進行線性修正

```c
R2= [00100|12001|10012|10001|12000|01234|10]
                 *     *     *           *
R3= [00100|12001|20012|30001|23000|01234|50]
```

大功告成！

---

## 獻給未來助教 ##

---

## 沙盒內抓取裝置 ##

在沙盒內完全抓不到裝置資訊，原因在於沒有掛載

```bash
$ ls /dev/nvidia*
```

將列出來的數個路徑，放入沙盒掛載設定檔中。

下面只是部分情況，未來可能模組名稱、數目可能會變動

```bash
dev/nvidia0
dev/nvidia1
dev/nvidia2
dev/nvidiactl
dev/nvidia-uvm
...
```

---

## 沙盒運作 ##

能抓取裝置資訊，卻無法運行 GPU kernel 程序，原因在於動態連結函式庫，如掛載 CUDA 8.0 環境時，通常缺少

* `libnvidia-fatbinaryloader.so.375.39`

* `libnvidia-ptxjitcompiler.so.375.39`

* `libnvidia-compiler.so.375.39`

以上三個檔案，主要支援動態的編譯的相關庫

* 隨著每次更新，它們的所在位置很容易變動

* 找出來，並且把它 `ln -s` 到執行檔會載入的路徑下，如 `/lib/x86_64-linux-gnu/`

---

## 更新事項 ##

* 編譯所需的 `nvcc` 在 CUDA 安裝時的路徑下

* 在多版本并存的時候特別小心，如 `which nvcc` 是否來自於同一版本，否則
  * 運行時，發生動態連結函式庫無法匹配 ... 等

* `nvvp`、`nvidia-smi` ... 等相關工具無法 **正常** 運行，  
  意即部分功能仍可能正常運作

---

## 工作站設定 ##

為讓其他使用者登入可以直接使用 `nvcc` 編譯器，以 bash 為例

```bash
sudo vim /etc/bash.bashrc
```

加入以下幾點 (請注意版本號)

```bash
# CUDA

export PATH=$PATH:/usr/local/cuda-8.0/bin
export LD_LIBRARY_PATH=/usr/lib/nvidia-375/:$LD_LIBRARY_PATH
export LD_LIBRARY_PATH=/usr/local/cuda-8.0/lib64/:$LD_LIBRARY_PATH
```