Mercurial > hg > Members > innparusu > slides

comparison 2017/2017_09_26/slide.md @ 27:d005b4f353d3

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