changeset 27:d005b4f353d3

Update slide
author Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
date Tue, 26 Sep 2017 19:06:19 +0900
parents ed48cf95acab
children 382cd93f2a60
files 2017/2017_07_25/slide.md 2017/2017_08_08/slide.md 2017/2017_09_05/slide.md 2017/2017_09_19/slide.md 2017/2017_09_26/slide.md
diffstat 5 files changed, 678 insertions(+), 0 deletions(-) [+]
line wrap: on
line diff
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/2017/2017_07_25/slide.md	Tue Sep 26 19:06:19 2017 +0900
@@ -0,0 +1,119 @@
+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 の並列処理機構の実装を行う。また、並列処理の検証をメタ計算として記述することで、 並列処理の信頼性を保証する。
+
+## 今週
+- bitonicSort をGears上で動かす
+
+
+## index Iterator
+- Iterateなtaskにindexをもたせる
+- taskの情報をcopyする際に新しく Integerのdata Gear を生成し, それを input Data Gear にする
+
+``` c
+struct Context* createOneDimIterateTask(struct Context* task, int x) {
+    struct Context* task1 = NEW(struct Context);
+    initContext(task1);
+    task1->taskManager = task->taskManager;
+    task1->next     = task->next;
+    task1->iterate  = 1;
+    task1->iterator = task->iterator;
+    task1->idgCount = task->idgCount;
+    task1->idg      = task->idg;
+    task1->maxIdg   = task->maxIdg;
+    for(int i = task1->idg; i < task1->maxIdg; i++) {
+        task1->data[i] = task->data[i];
+    }
+
+    // create index data gear and register input data to iterate task
+    struct Integer* dataX = &ALLOCATE_DATA_GEAR(task1, Integer)->Integer;
+    dataX->value = x;
+    task1->data[task1->maxIdg++] = (union Data*)dataX;
+
+    task1->odg      = task->odg + 1;
+    task1->maxOdg   = task->maxOdg + 1;
+    for (int i = task1->odg; i < task1->maxOdg; i++) {
+        task1->data[i] = task->data[i-1];
+    }
+
+    return task1;
+}
+```
+
+## bitonicSort 
+- Cerium で使っていた GPU 向けの例題を実装
+
+<div style="text-align: center;">
+    <img src="./pictures/bitonic.svg" alt="message" width="500">
+</div>
+
+- 並列実行されるtask は swap
+
+``` c
+__code bitonicSwap(struct SortArray* inputArray, struct Integer* block, struct Integer* first, struct Integer* i, struct LoopCounter* loopCounter, __code next(struct SortArray* output, ...)) {
+    if (loopCounter->i < inputArray->prefix) {
+        int index = loopCounter->i + i->value * inputArray->prefix;
+        int position = index/block->value;
+        int index1 = index+block->value*position;
+        int index2 = (first->value == 1)? ((block->value<<1)*(position+1))-(index1%block->value)-1 : index1+block->value;
+        struct Integer** array = inputArray->array;
+        if (array[index2]->value < array[index1]->value) {
+            struct Integer *tmp = array[index1];
+            array[index1] = array[index2];
+            array[index2] = tmp;
+        }
+        loopCounter->i++;
+        goto meta(context, C_bitonicSwap);
+    }
+    loopCounter->i = 0;
+    output->array = inputArray->array;
+    goto meta(context, next);
+}
+```
+
+## bitonicSort 実測
+- 最初は1回の swap で 1 Task にしていた
+- しかし, swap は処理的には軽い Task なので, Workerの一連の処理(Get task から Data Gear の書き出しまで)のほうが長くなってしまい, CPU を増やすと逆に遅くなった
+- なので, ある程度の大きさの範囲のswap を 1Task にして実測した
+    - 実行環境: firefly
+    - core: 12
+    - length: 2^20
+    - Task数: 2^4
+
+<table  border="1" align='center' width='50%'>
+  <tbody>
+    <tr>
+      <td style="text-align: center;">Number of 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;">7986</td>
+    </tr>
+    <tr>
+      <td style="text-align: center;">2 CPUs</td>
+      <td style="text-align: right;">3993</td>
+    </tr>
+    <tr>
+      <td style="text-align: center;">4 CPUs</td>
+      <td style="text-align: right;">2142</td>
+    </tr>
+    <tr>
+      <td style="text-align: center;">8 CPUs</td>
+      <td style="text-align: right;">1204</td>
+    </tr>
+    <tr>
+      <td style="text-align: center;">12 CPUs</td>
+      <td style="text-align: right;">1078</td>
+    </tr>
+  </tbody>
+</table>
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/2017/2017_08_08/slide.md	Tue Sep 26 19:06:19 2017 +0900
@@ -0,0 +1,217 @@
+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 の並列処理機構の実装を行う。また、並列処理の検証をメタ計算として記述することで、 並列処理の信頼性を保証する。
+
+## 今週
+- par goto の変換script
+
+
+## par goto
+
+``` c
+__code code2(...) {
+...
+    par goto mult(integer1, integer2, integer3, __exit);
+    goto createTask1();
+}
+```
+
+- 実行するtaskは↓のような感じ
+``` c
+__code mult(struct Integer* input1, struct Integer* input2, __code next(struct Integer* output, ...)) {
+  ...
+}
+```
+
+- ↓ に変換できれば良い
+
+``` c
+__code code2(...) {
+...
+    struct Context** tasks = (struct Context**)ALLOC_ARRAY(context, Context, 1);
+    int taskCount = 0;
+
+    struct Context* task = NEW(struct Context);
+    initContext(task);
+    task->next = C_mult;
+    task->idgCount = 2;
+    task->idg = task->dataNum;
+    task->data[task->idg] = (union Data*)integer1;
+    task->data[task->idg+1] = (union Data*)integer2;
+    task->maxIdg = task->idg + 2;
+    task->odg = task->maxIdg;
+    task->data[task->odg] = (union Data*)integer3;
+    task->maxOdg = task->odg + 1;
+    tasks[taskCount] = task;
+    taskCount++;
+
+    taskManager->contexts = tasks;
+    // goto crateTask1();
+    taskManager->next1 = C_createTask1;
+    goto meta(context, taskManager->taskManager->TaskManager.spawnTasks);
+}
+```
+
+## par goto
+- 中途半端な感じですがこんな変換はできた
+
+``` c
+__code code2(...) {
+...
+    par goto mult(integer1, integer2, integer3, __exit);
+}
+```
+
+``` c
+__code code2(...) {
+...
+    struct Context** tasks = (struct Context**)ALLOC_ARRAY(context, Context, ?); // ? の部分は一つのcode gear block の par goto の数が入る
+    int taskCount = 0;
+
+    struct Context* task = NEW(struct Context);
+    initContext(task);
+    task->next = C_mult;
+    task->idgCount = 2;
+    task->idg = task->dataNum;
+    task->data[task->idg] = (union Data*)integer1;
+    task->data[task->idg+1] = (union Data*)integer2;
+    task->maxIdg = task->idg + 2;
+    task->odg = task->maxIdg;
+    task->data[task->odg] = (union Data*)integer3;
+    task->maxOdg = task->odg + 1;
+    tasks[taskCount] = task;
+    taskCount++;
+    // まだ par goto が残ってる
+    par goto mult(integer1, integer2, integer3, __exit);
+    
+    // まだ, TaskManager の spawn に飛ばす変換をしていない
+    goto crateTask1();
+}
+```
+
+## par goto script
+- generate\_stub を編集した形
+- par goto を見つけると, その cs名.cbcを見に行く
+
+``` perl
+sub getDataGear {
+    my ($filename) = @_;
+    my ($codeGearName, $name, $inTypedef);
+    open my $fd,"<",$filename or die("can't open $filename $!");
+    while (<$fd>) {
+            if (/^typedef struct (\w+)\s*<(.*)>/) {
+            ....
+            } elsif(/^(.*)par goto (\w+)\((.*)\)/) {
+                my $codeGearName = $2;
+                if ( -f "$codeGearName.cbc") {
+                    &getCodeGear("$codeGearName.cbc");
+                }
+            }
+```
+
+## par goto script
+- input と output の数を取りに行く
+``` perl
+sub getCodeGear {
+    my ($filename) = @_;
+    open my $fd,"<",$filename or die("can't open $filename $!");
+    my ($name,$impln);
+    while (<$fd>) {
+        if (/^(\w+)(\*)+ create(\w+)\(/) {
+            ....
+        }
+        if (defined $name) {
+            ....
+        } elsif (/^\_\_code (\w+)\((.*)\)(.*)/) {
+            my $codeGearName = $1;
+            my $args = $2;
+            my $inputCount = 0;
+            my $outputCount = 0;
+            my $inputIncFlag = 1;
+            while($args) {
+                if ($args =~ s/(^\s*,\s*)//) {
+                }
+                if ($args =~ s/^(\s)*\_\_code\s+(\w+)\(//) {
+                    $inputIncFlag = 0;
+                } elsif ($args =~ s/^(struct|union) (\w+)(\*)+\s(\w+)//) {
+                    if($inputIncFlag) {
+                        $inputCount++;
+                    }
+                    else {
+                        $outputCount++;
+                    }
+                } elsif ($args =~ s/(.*,)//) {
+                } else {
+                    last;
+                }
+            }
+            // ここで code gear 毎のinput, output の数を持っておく
+            $codeGear{$codeGearName}->{"input"} = $inputCount;
+            $codeGear{$codeGearName}->{"output"} = $outputCount;
+        }
+    }
+}
+
+```
+
+## par goto script
+- par gotoの変換
+
+```perl
+sub generateDataGear {
+    ...
+    while(<$fd>) {
+            if (...) {
+            } elsif(/^(.*)par goto (\w+)\((.*)\)/) {
+                # handling par goto statement
+                # convert it to the parallel 
+                my $prev = $1;
+                my $codeGearName = $2;
+                my @dataGears = split(/,\s*/, $3);
+                my $nextCodeGear = pop(@dataGears);
+                my $inputCount = $codeGear{$codeGearName}->{'input'};
+                my $outputCount = $codeGear{$codeGearName}->{'output'};
+                if (! $inParGoto) {
+                    $inParGoto = 1;
+                    my $initTasks = << "EOFEOF";
+${prev}struct Context** tasks = (struct Context**)ALLOC_ARRAY(context, Context, ?);
+${prev}int taskCount = 0;
+EOFEOF
+                    print $fd $initTasks;
+                }
+
+                my $initTask = << "EOFEOF";
+${prev}struct Context* task = NEW(struct Context);
+${prev}initContext(task);
+${prev}task->next = C_$codeGearName;
+${prev}task->idgCount = $inputCount;
+${prev}task->idg = task->dataNum;
+${prev}task->maxIdg = task->idg + $inputCount;
+${prev}task->odg = task->maxIdg;
+${prev}task->maxOdg = task->odg + $outputCount;
+EOFEOF
+                print $fd $initTask;
+                for my $i (0..$inputCount-1) {
+                    print $fd "${prev}task->data[task->idg+$i] = (union Data*)@dataGears[$i];\n";
+                }
+
+                for my $i (0..$outputCount-1) {
+                    print $fd "${prev}task->data[task->odg+$i] = (union Data*)@dataGears[$inputCount+$i];\n";
+                }
+
+                print $fd "${prev}tasks[taskCount] = task;\n";
+                print $fd "${prev}taskCount++;\n";
+            }
+```
+
+## 残っていること
+- par goto で実行される cs の stub 生成
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/2017/2017_09_05/slide.md	Tue Sep 26 19:06:19 2017 +0900
@@ -0,0 +1,18 @@
+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 の並列処理機構の実装を行う。また、並列処理の検証をメタ計算として記述することで、 並列処理の信頼性を保証する。
+
+## 今週
+- par goto の変換script
+
+
+## par goto
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/2017/2017_09_19/slide.md	Tue Sep 26 19:06:19 2017 +0900
@@ -0,0 +1,65 @@
+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
+- 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;
+}
+```
+
+## TaskManager
+- 取り敢えず Iterator は全て CUDAWorker(作られていれば) に sendするように変更
+
+## CUDAExec()
+- Task(Context) から Input&output DataGear をGPU に送信するために変換 & Task の実行をしている部分
+- まだここは実行するタスクの内容で書き換えているのでCeriumのように一般化してデータを送信する必要あり
+- block の数は単純に iterator の値を適応(この辺は Cerium と一緒), 次にBenchmarkを載せるが, 遅い
+
+## 一応Benchmark的なもの
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/2017/2017_09_26/slide.md	Tue Sep 26 19:06:19 2017 +0900
@@ -0,0 +1,259 @@
+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 をどうにかする
+- 並列処理の検証をどうするか