Members/innparusu/slides: 2017/2017_09

author	Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
date	Tue, 26 Sep 2017 19:06:19 +0900
parents
children

rev	line source
27 d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	1 title: Gears OS
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	2 author: Tatsuki IHA
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	3 profile:
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	4 lang: Japanese
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	5 code-engine: coderay
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	6
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	7 ## 研究目的
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	8 - 当研究室では処理の単位を Code Gear、データの単位を Data Gear を用いて信頼性が高い並列処理を行う Gears OS を開発している
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	9 - Gears OS では Task を Code Gear と実行するときに必要な Input Data Gear と出力するための Output Data Gear の組で表現される。 Input Data Gear/Output Data Gear によって依存関係が決定し、それにそって並列実行を行う.
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	10 - 信頼性の確保はモデルチェック、検証等を使用して行う。この信頼性のための計算は通常の計算とは別の階層のメタ計算として記述する。
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	11 - また、メタ計算は信頼性の他に CPU, GPU などの実行環境の切り替え, データ拡張等の柔軟性を提供する。
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	12 - 本研究では、 Gears OS の並列処理機構の実装を行う。また、並列処理の検証をメタ計算として記述することで、並列処理の信頼性を保証する。
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	13
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	14 ## 今週
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	15 - Gears での GPU の実行(結構無理やり)
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	16 - bitonicSort
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	17
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	18 ## GPUWorker
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	19 - cudaInit() はもともと, createGPUWorker() で行っている
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	20 - しかし, 実際 CUDA の API を呼び出すのは, createGPUWorker を呼んだスレッドではなく, createGPUWorker で作られた Thread
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	21 - そうすると code=201(CUDA_ERROR_INVALID_CONTEXT) がでる
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	22 - cuda の API は cuInit() を呼んだ thread でしか呼び出せない
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	23 - thread を入れ替える技の痕跡があるが, thread 作った先の startCUDAWorker で cudaInit することに
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	24
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	25 ``` c
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	26 static void startCUDAWorker(Worker* worker) {
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	27 struct CUDAWorker* cudaWorker = &worker->worker->CUDAWorker;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	28 cudaInit(cudaWorker, 0);
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	29 cudaWorker->context = NEW(struct Context);
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	30 initContext(cudaWorker->context);
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	31 Gearef(cudaWorker->context, Worker)->worker = (union Data*)worker;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	32 goto meta(cudaWorker->context, worker->taskReceive);
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	33 }
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	34 ```
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	35
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	36 ## MultiDimIterator
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	37 - 取り敢えず Iterator は全て CUDAWorker(作られていれば) にで実行するように
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	38 - 将来的には meta部分で切り替えられるようにする
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	39 - CUDA に対応させるために, 前作った 1次元の iterator を 3次元に対応
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	40
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	41 ``` c
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	42 Iterator* createMultiDimIterator(struct Context* context, int x, int y, int z) {
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	43 struct Iterator* iterator = new Iterator();
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	44 struct MultiDimIterator* multiDimIterator = new MultiDimIterator();
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	45 iterator->iterator = (union Data*)multiDimIterator;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	46 iterator->exec = C_execMultiDimIterator;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	47 iterator->barrier = C_barrierMultiDimIterator;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	48 multiDimIterator->x = x;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	49 multiDimIterator->y = y;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	50 multiDimIterator->z = z;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	51 multiDimIterator->count = x * y * z;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	52 multiDimIterator->counterX = 0;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	53 multiDimIterator->counterY = 0;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	54 multiDimIterator->counterZ = 0;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	55 return iterator;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	56 }
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	57 ```
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	58
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	59 - exec時は, GPU があれば, task に gpu を設定して spawn する
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	60
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	61 ``` c
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	62 __code execMultiDimIterator(struct MultiDimIterator* iterator, struct TaskManager* taskManager, struct Context* task, int numGPU, __code next(...)) {
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	63 // No GPU device
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	64 if (numGPU == 0) {
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	65 goto meta(context, C_execMultiDimIterator1);
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	66 }
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	67 task->iterate = 1;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	68 task->gpu = 1;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	69 taskManager->taskManager = (union Data*)task->taskManager;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	70 taskManager->context = task;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	71 taskManager->next = next;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	72 goto meta(context, task->taskManager->spawn);
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	73 }
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	74 ```
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	75
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	76 ## TaskManager
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	77 - Iterator でgpu がセットされていれば, CUDAWorker に投げる
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	78
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	79 ```
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	80 if (task->gpu) {
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	81 task->workerId = taskManagerImpl->sendGPUWorkerIndex;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	82 if(++taskManagerImpl->sendGPUWorkerIndex >= taskManager->cpu) {
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	83 taskManagerImpl->sendGPUWorkerIndex = taskManager->gpu;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	84 }
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	85 } else {
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	86 task->workerId = taskManagerImpl->sendCPUWorkerIndex;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	87 if(++taskManagerImpl->sendCPUWorkerIndex >= taskManager->maxCPU) {
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	88 taskManagerImpl->sendCPUWorkerIndex = taskManager->cpu;
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	89 }
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	90 }
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	91 ```
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	92
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	93 ## CUDAExec()
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	94 - Task(Context) から Input&output DataGear をGPU に送信するために変換 & Task の実行をしている部分
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	95 - まだここは実行するタスクの内容で書き換えているのでCeriumのように一般化してデータを送信する必要あり
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	96 - meta 部分で全部押し込んじゃう?
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	97 - 必要なのはGPUに送るデータのサイズなので, meta部分でサイズを持たせて, それを使うようにする
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	98 - block の数は単純に iterator の値を適応(この辺は Cerium と一緒), 次にBenchmarkを載せるが, 遅い
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	99 - http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#thread-hierarchy
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	100 - http://horacio9573.no-ip.org/cuda/group__CUDA__EXEC_gb8f3dc3031b40da29d5f9a7139e52e15.html
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	101
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	102 ``` c
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	103 checkCudaErrors(cuLaunchKernel(context->function,
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	104 iterator->x, iterator->y, iterator->z,
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	105 1, 1, 1,
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	106 0, NULL, args, NULL));
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	107 ```
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	108
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	109 ## 一応Benchmark的なもの
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	110 - BitonicSort を GPUで動かしてみました
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	111 - 環境 firefly
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	112 - GPU Quadro k5000
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	113 - サイズは 2^23
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	114 - swap のタスク数は 2^22 にしたかったが, CUDA\_ERROR\_INVALID\_VALUE が出てしまった, 2^15=32768 だと Error なしだったのでそれで動かしてみました
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	115 - cpu は 2^15 で分割するとその分Contextを作ってしまい, とても遅いので32個のTaskに分割
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	116
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	117 <table border="1" align='center' width='50%'>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	118 <tbody>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	119 <tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	120 <td style="text-align: center;">Processors</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	121 <td style="text-align: center;">Time(ms)</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	122 </tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	123 <tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	124 <td style="text-align: center;">1 CPU</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	125 <td style="text-align: right;">31161</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	126 </tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	127 <tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	128 <td style="text-align: center;">2 CPUs</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	129 <td style="text-align: right;">16277</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	130 </tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	131 <tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	132 <td style="text-align: center;">4 CPUs</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	133 <td style="text-align: right;">8466</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	134 </tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	135 <tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	136 <td style="text-align: center;">8 CPUs</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	137 <td style="text-align: right;">4457</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	138 </tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	139 <tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	140 <td style="text-align: center;">12 CPUs</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	141 <td style="text-align: right;">3299</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	142 </tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	143 <tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	144 <td style="text-align: center;">GPU(CUDA)</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	145 <td style="text-align: right;">12083</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	146 </tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	147 </tbody>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	148 </table>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	149
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	150 ## ちょっと GPU が遅いので
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	151 - block 内の Threadも指定してみる
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	152 - 指定値は取り敢えず最大値の 1024
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	153 - Task 数は 2^22
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	154
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	155 ```
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	156 checkCudaErrors(cuLaunchKernel(context->function,
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	157 iterator->x/1024, iterator->y, iterator->z,
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	158 1024, 1, 1,
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	159 0, NULL, args, NULL));
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	160 ```
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	161
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	162 <table border="1" align='center' width='50%'>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	163 <tbody>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	164 <tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	165 <td style="text-align: center;">Processors</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	166 <td style="text-align: center;">Time(ms)</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	167 </tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	168 <tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	169 <td style="text-align: center;">4 CPUs</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	170 <td style="text-align: right;">8466</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	171 </tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	172 <tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	173 <td style="text-align: center;">8 CPUs</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	174 <td style="text-align: right;">4457</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	175 </tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	176 <tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	177 <td style="text-align: center;">12 CPUs</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	178 <td style="text-align: right;">3299</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	179 </tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	180 <tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	181 <td style="text-align: center;">GPU(CUDA)</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	182 <td style="text-align: right;">12083</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	183 </tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	184 <tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	185 <td style="text-align: center;">GPU(CUDA Block Threads)</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	186 <td style="text-align: right;">5642</td>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	187 </tr>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	188 </tbody>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	189 </table>
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	190
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	191 ## いろいろtips
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	192 - cuda sample の ./deviceQuery で, GPU の性能を見れる
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	193
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	194 ```
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	195 firefly@one$ /Developer/NVIDIA/CUDA-8.0/samples/1_Utilities/deviceQuery/deviceQuery
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	196 /Developer/NVIDIA/CUDA-8.0/samples/1_Utilities/deviceQuery/deviceQuery Starting...
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	197
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	198 CUDA Device Query (Runtime API) version (CUDART static linking)
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	199
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	200 Detected 1 CUDA Capable device(s)
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	201
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	202 Device 0: "Quadro K5000"
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	203 CUDA Driver Version / Runtime Version 8.0 / 8.0
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	204 CUDA Capability Major/Minor version number: 3.0
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	205 Total amount of global memory: 4096 MBytes (4294508544 bytes)
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	206 ( 8) Multiprocessors, (192) CUDA Cores/MP: 1536 CUDA Cores
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	207 GPU Max Clock rate: 706 MHz (0.71 GHz)
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	208 Memory Clock rate: 2700 Mhz
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	209 Memory Bus Width: 256-bit
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	210 L2 Cache Size: 524288 bytes
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	211 Maximum Texture Dimension Size (x,y,z) 1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	212 Maximum Layered 1D Texture Size, (num) layers 1D=(16384), 2048 layers
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	213 Maximum Layered 2D Texture Size, (num) layers 2D=(16384, 16384), 2048 layers
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	214 Total amount of constant memory: 65536 bytes
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	215 Total amount of shared memory per block: 49152 bytes
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	216 Total number of registers available per block: 65536
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	217 Warp size: 32
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	218 Maximum number of threads per multiprocessor: 2048
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	219 Maximum number of threads per block: 1024
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	220 Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	221 Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535)
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	222 Maximum memory pitch: 2147483647 bytes
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	223 Texture alignment: 512 bytes
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	224 Concurrent copy and kernel execution: Yes with 2 copy engine(s)
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	225 Run time limit on kernels: Yes
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	226 Integrated GPU sharing Host Memory: No
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	227 Support host page-locked memory mapping: Yes
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	228 Alignment requirement for Surfaces: Yes
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	229 Device has ECC support: Disabled
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	230 Device supports Unified Addressing (UVA): Yes
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	231 Device PCI Domain ID / Bus ID / location ID: 0 / 5 / 0
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	232 Compute Mode:
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	233 < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	234
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	235 deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 8.0, CUDA Runtime Version = 8.0, NumDevs = 1, Device0 = Quadro K5000
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	236 ```
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	237
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	238 ## 自分のpc に cuda を入れたときにちょっとハマった場所
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	239 - cuda version
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	240 - V8.0.61
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	241 - Apple clang version
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	242 - Apple LLVM version 8.1.0 (clang-802.0.42)
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	243 - The version ('80100') of the host compiler ('Apple clang') is not supported と怒られる
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	244 - 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用を入れることに
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	245 - 8.2 を入れるとApple LLVM version 8.0.0 (clang-800.0.42.1)になりました
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	246 - 複数入れられるかなぁ
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	247 - xcode 経由で入れると pathが変わるらしい(dmg 経由で入れると /Library/Developer/CommandLineTools/usr/bin)
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	248 - xcode 経由と言うより, xcode にくっついてる感じ preference -> location から変更できる
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	249 - ``sudo xcode-select -s path/to`` で切り替え
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	250 - めでたくcompileできた
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	251 - 最近 CUDA の Versionが 9 になっており, Apple LLVM 8.1.0 でも動くようになっていた
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	252 - http://docs.nvidia.com/cuda/cuda-installation-guide-mac-os-x/index.html#system-requirements
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	253 - しかし, xcode の version も 9.0 にあがっており, Apple LLVM の version 9.0.0に上がっているため, 動くかどうか不明
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	254
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	255 ## 次は
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	256 - RBTree の deletion
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	257 - cuda.c の抽象化
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	258 - Perl script をどうにかする
d005b4f353d3 Update slide Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> parents: diff changeset	259 - 並列処理の検証をどうするか

27

d005b4f353d3 Update slide