title: Gears OS
author: Tatsuki IHA
profile:
lang: Japanese
code-engine: coderay

## 研究目的
- 当研究室では  処理の単位を Code Gear、 データの単位を Data Gear を用いて 信頼性が高い並列処理を行う Gears OS を開発している
- Gears OS では Task を Code Gear と実行するときに必要な Input Data Gear と出力するための Output Data Gear の組で表現される。 Input Data Gear/Output Data Gear によって依存関係が決定し、それにそって並列実行を行う.
- 信頼性の確保はモデルチェック、検証等を使用して行う。この信頼性のための計算は通常の計算とは別の階層のメタ計算として記述する。
- また、 メタ計算は信頼性の他に CPU, GPU などの実行環境の切り替え, データ拡張等の柔軟性を提供する。
- 本研究では、 Gears OS の並列処理機構の実装を行う。また、並列処理の検証をメタ計算として記述することで、 並列処理の信頼性を保証する。

## 今週
- Gears での GPU の実行(結構無理やり)
  - bitonicSort

## GPUWorker
- cudaInit() はもともと, createGPUWorker() で行っている
- しかし, 実際 CUDA の API を呼び出すのは, createGPUWorker を呼んだスレッドではなく, createGPUWorker で作られた Thread
- そうすると code=201(CUDA_ERROR_INVALID_CONTEXT) がでる
  - cuda の API は cuInit() を呼んだ thread でしか呼び出せない
- thread を入れ替える技の痕跡があるが, thread 作った先の startCUDAWorker で cudaInit することに

``` c
static void startCUDAWorker(Worker* worker) {
    struct CUDAWorker* cudaWorker = &worker->worker->CUDAWorker;
    cudaInit(cudaWorker, 0);
    cudaWorker->context = NEW(struct Context);
    initContext(cudaWorker->context);
    Gearef(cudaWorker->context, Worker)->worker = (union Data*)worker;
    goto meta(cudaWorker->context, worker->taskReceive);
}
```

## MultiDimIterator
- 取り敢えず Iterator は全て CUDAWorker(作られていれば) にで実行するように
  - 将来的には meta部分で切り替えられるようにする
- CUDA に対応させるために, 前作った 1次元の iterator を 3次元に対応

``` c
Iterator* createMultiDimIterator(struct Context* context, int x, int y, int z) {
    struct Iterator* iterator = new Iterator();
    struct MultiDimIterator* multiDimIterator = new MultiDimIterator();
    iterator->iterator = (union Data*)multiDimIterator;
    iterator->exec = C_execMultiDimIterator;
    iterator->barrier = C_barrierMultiDimIterator;
    multiDimIterator->x = x;
    multiDimIterator->y = y;
    multiDimIterator->z = z;
    multiDimIterator->count = x * y * z;
    multiDimIterator->counterX = 0;
    multiDimIterator->counterY = 0;
    multiDimIterator->counterZ = 0;
    return iterator;
}
```

- exec時は, GPU があれば, task に gpu を設定して spawn する

``` c
__code execMultiDimIterator(struct MultiDimIterator* iterator, struct TaskManager* taskManager, struct Context* task, int numGPU, __code next(...)) {
    // No GPU device
    if (numGPU == 0) {
        goto meta(context, C_execMultiDimIterator1);
    }
    task->iterate = 1;
    task->gpu = 1;
    taskManager->taskManager = (union Data*)task->taskManager;
    taskManager->context = task;
    taskManager->next = next;
    goto meta(context, task->taskManager->spawn);
}
```

## TaskManager
- Iterator でgpu がセットされていれば,  CUDAWorker に投げる

```
if (task->gpu) {
    task->workerId = taskManagerImpl->sendGPUWorkerIndex;   
    if(++taskManagerImpl->sendGPUWorkerIndex >= taskManager->cpu) {
        taskManagerImpl->sendGPUWorkerIndex = taskManager->gpu;
    }
} else {
    task->workerId = taskManagerImpl->sendCPUWorkerIndex;
    if(++taskManagerImpl->sendCPUWorkerIndex >= taskManager->maxCPU) {
        taskManagerImpl->sendCPUWorkerIndex = taskManager->cpu;
    }
}
```

## CUDAExec()
- Task(Context) から Input&output DataGear をGPU に送信するために変換 & Task の実行をしている部分
- まだここは実行するタスクの内容で書き換えているのでCeriumのように一般化してデータを送信する必要あり
    - meta 部分で全部押し込んじゃう?
    - 必要なのはGPUに送るデータのサイズなので, meta部分でサイズを持たせて, それを使うようにする
- block の数は単純に iterator の値を適応(この辺は Cerium と一緒), 次にBenchmarkを載せるが, 遅い
- http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#thread-hierarchy
- http://horacio9573.no-ip.org/cuda/group__CUDA__EXEC_gb8f3dc3031b40da29d5f9a7139e52e15.html

``` c
checkCudaErrors(cuLaunchKernel(context->function,
            iterator->x, iterator->y, iterator->z,
            1, 1, 1,
            0, NULL, args, NULL));
```

## 一応Benchmark的なもの
- BitonicSort を GPUで動かしてみました
- 環境 firefly
- GPU Quadro k5000
- サイズは 2^23
- swap のタスク数は 2^22 にしたかったが, CUDA\_ERROR\_INVALID\_VALUE が出てしまった, 2^15=32768 だと Error なしだったのでそれで動かしてみました
- cpu は 2^15 で分割するとその分Contextを作ってしまい, とても遅いので32個のTaskに分割

<table  border="1" align='center' width='50%'>
<tbody>
<tr>
<td style="text-align: center;">Processors</td>
<td style="text-align: center;">Time(ms)</td>
</tr>
<tr>
<td style="text-align: center;">1 CPU</td>
<td style="text-align: right;">31161</td>
</tr>
<tr>
<td style="text-align: center;">2 CPUs</td>
<td style="text-align: right;">16277</td>
</tr>
<tr>
<td style="text-align: center;">4 CPUs</td>
<td style="text-align: right;">8466</td>
</tr>
<tr>
<td style="text-align: center;">8 CPUs</td>
<td style="text-align: right;">4457</td>
</tr>
<tr>
<td style="text-align: center;">12 CPUs</td>
<td style="text-align: right;">3299</td>
</tr>
<tr>
<td style="text-align: center;">GPU(CUDA)</td>
<td style="text-align: right;">12083</td>
</tr>
</tbody>
</table>

## ちょっと GPU が遅いので
- block 内の Threadも指定してみる
- 指定値は 取り敢えず最大値の 1024
- Task 数は 2^22

```
checkCudaErrors(cuLaunchKernel(context->function,
            iterator->x/1024, iterator->y, iterator->z,
            1024, 1, 1,
            0, NULL, args, NULL));
```

<table  border="1" align='center' width='50%'>
<tbody>
<tr>
<td style="text-align: center;">Processors</td>
<td style="text-align: center;">Time(ms)</td>
</tr>
<tr>
<td style="text-align: center;">4 CPUs</td>
<td style="text-align: right;">8466</td>
</tr>
<tr>
<td style="text-align: center;">8 CPUs</td>
<td style="text-align: right;">4457</td>
</tr>
<tr>
<td style="text-align: center;">12 CPUs</td>
<td style="text-align: right;">3299</td>
</tr>
<tr>
<td style="text-align: center;">GPU(CUDA)</td>
<td style="text-align: right;">12083</td>
</tr>
<tr>
<td style="text-align: center;">GPU(CUDA Block Threads)</td>
<td style="text-align: right;">5642</td>
</tr>
</tbody>
</table>

## いろいろtips
- cuda sample の ./deviceQuery で, GPU の性能を見れる

```
firefly@one$ /Developer/NVIDIA/CUDA-8.0/samples/1_Utilities/deviceQuery/deviceQuery
/Developer/NVIDIA/CUDA-8.0/samples/1_Utilities/deviceQuery/deviceQuery Starting...

 CUDA Device Query (Runtime API) version (CUDART static linking)

Detected 1 CUDA Capable device(s)

Device 0: "Quadro K5000"
  CUDA Driver Version / Runtime Version          8.0 / 8.0
  CUDA Capability Major/Minor version number:    3.0
  Total amount of global memory:                 4096 MBytes (4294508544 bytes)
  ( 8) Multiprocessors, (192) CUDA Cores/MP:     1536 CUDA Cores
  GPU Max Clock rate:                            706 MHz (0.71 GHz)
  Memory Clock rate:                             2700 Mhz
  Memory Bus Width:                              256-bit
  L2 Cache Size:                                 524288 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
  Maximum Layered 1D Texture Size, (num) layers  1D=(16384), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(16384, 16384), 2048 layers
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total number of registers available per block: 65536
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  2048
  Maximum number of threads per block:           1024
  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and kernel execution:          Yes with 2 copy engine(s)
  Run time limit on kernels:                     Yes
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support:                        Disabled
  Device supports Unified Addressing (UVA):      Yes
  Device PCI Domain ID / Bus ID / location ID:   0 / 5 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 8.0, CUDA Runtime Version = 8.0, NumDevs = 1, Device0 = Quadro K5000
```

## 自分のpc に cuda を入れたときにちょっとハマった場所
- cuda version
  - V8.0.61
- Apple clang version
  - Apple LLVM version 8.1.0 (clang-802.0.42)
- The version ('80100') of the host compiler ('Apple clang') is not supported と怒られる
- https://github.com/arrayfire/arrayfire/issues/1384#issuecomment-291471533 を見ると,  8.0.61 では xcode 8.3 の CLT は not support(8.1.0 は xocde8.3 用?になってる) なので 8.2用を入れることに
  - 8.2 を入れるとApple LLVM version 8.0.0 (clang-800.0.42.1)になりました
  - 複数入れられるかなぁ
    - xcode 経由で入れると pathが変わるらしい(dmg 経由で入れると /Library/Developer/CommandLineTools/usr/bin)
    - xcode 経由と言うより, xcode にくっついてる感じ preference -> location から変更できる
    - ``sudo xcode-select -s path/to`` で切り替え
- めでたくcompileできた
- 最近 CUDA の Versionが 9 になっており, Apple LLVM 8.1.0 でも動くようになっていた
- http://docs.nvidia.com/cuda/cuda-installation-guide-mac-os-x/index.html#system-requirements
- しかし, xcode の version も 9.0 にあがっており, Apple LLVM の version 9.0.0に上がっているため, 動くかどうか不明

## 次は
- RBTree の deletion
- cuda.c の抽象化
- Perl script をどうにかする
- 並列処理の検証をどうするか