Mercurial > hg > Members > innparusu > slides

comparison 2017/2017_10_17/slide.md @ 28:382cd93f2a60

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, 31 Oct 2017 18:16:16 +0900
parents
children
comparison
equal deleted inserted replaced
27:d005b4f353d3 28:382cd93f2a60
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 - cuda.c の抽象化
16 - 今までは一つ一つの例題で cuda.c を書き換えてたので、、
17
18 ## CudaBuffer
19 - cuda 用の引数, DataGear の値を持った構造体を定義
20
21 ``` c
22 union Data {
23 .....
24 struct CudaBuffer {
25 void** kernelParams;
26 int inputLen;
27 int outputLen;
28 union Data** inputData;
29 union Data** outputData;
30 } CudaBuffer;
31 }
32 ```
33
34 ## 呼び出し
35 - cudaの実行は stubで呼ばれる
36 - cudaで変換する Gode Gear を指定する
37
38 ``` c
39 __code twice_stub(struct Context* context) {
40 #ifdef USE_CUDAWorker
41 if (context->gpu) {
42 struct Array* array = &context->data[context->idg]->Array;
43 struct CudaBuffer* buffer = &ALLOCATE(context, CudaBuffer)->CudaBuffer;
44 buffer->inputData = (union Data**)ALLOCATE_PTR_ARRAY(context, Array, 2);
45 buffer->inputData[0] = (union Data*)array->array;
46 buffer->inputData[1] = (union Data*)array;
47 buffer->outputData = NULL;
48 buffer->inputLen = 2;
49 buffer->outputLen = 0;
50 cudaExec(context, buffer, "c/examples/twice/CUDAtwice.ptx", "twice");
51 //continuationにそってGPUworkerに戻る
52 goto meta(context, context->next);
53 }
54 #endif
55 goto twice(context,
56 &context->data[context->idg]->Array,
57 &context->data[context->idg+1]->MultiDim,
58 context->next,
59 Gearef(context, LoopCounter));
60 }
61 ```
62
63 ## cudaRead
64 - buffer に入れた DataGear を cuda に allocate, 送信
65
66 ``` c
67 void cudaRead(struct CudaBuffer* buffer) {
68 buffer->kernelParams = (void **)calloc(buffer->inputLen + buffer->outputLen, sizeof(void *));
69 int paramCount = 0;
70 for (int i = 0; i < buffer->inputLen; i++) {
71 CUdeviceptr* deviceptr = (CUdeviceptr *)calloc(1, sizeof(CUdeviceptr));
72 // memory allocate
73 checkCudaErrors(cuMemAlloc(deviceptr, GET_SIZE(buffer->inputData[i])));
74 // Synchronous data transfer(host to device)
75 checkCudaErrors(cuMemcpyHtoD(*deviceptr, buffer->inputData[i], GET_SIZE(buffer->inputData[i])));
76 buffer->kernelParams[paramCount++] = deviceptr;
77 }
78
79 for (int i = 0; i < buffer->outputLen; i++) {
80 CUdeviceptr* deviceptr = (CUdeviceptr *)calloc(1, sizeof(CUdeviceptr));
81 // memory allocate
82 checkCudaErrors(cuMemAlloc(deviceptr, GET_SIZE(buffer->outputData[i])));
83 // Synchronous data transfer(host to device)
84 checkCudaErrors(cuMemcpyHtoD(*deviceptr, buffer->outputData[i], GET_SIZE(buffer->outputData[i])));
85 buffer->kernelParams[paramCount++] = deviceptr;
86 }
87 }
88 ```
89
90 ## cudaExec2
91 - *cudaLaunch* とか *cudaLaunchKernel* にしたかった. すでにcuda側で使われているらしい(この辺の名前は後で要修正)
92 - cudaRead で設定した kernelParams(CUdeviceptr) を launchKernel に引き渡す
93
94 ``` c
95 void cudaExec2(struct Context* context, struct CudaBuffer* buffer) {
96 // Asynchronous launch kernel
97 context->num_exec = 1;
98 if (context->iterate) {
99 struct MultiDimIterator* iterator = &context->iterator->iterator->MultiDimIterator;
100 checkCudaErrors(cuLaunchKernel(context->function,
101 iterator->x, iterator->y, iterator->z,
102 1, 1, 1,
103 0, NULL, buffer->kernelParams, NULL));
104
105 } else {
106 checkCudaErrors(cuLaunchKernel(context->function,
107 1, 1, 1,
108 1, 1, 1,
109 0, NULL, buffer->kernelParams, NULL));
110 }
111 }
112 ```
113
114 ## cudaWrite
115 - 実行した結果を指定した DataGear に書き出す
116 - input にそのまま結果を書き出す場合もある. flip 的なものも実装しないと
117 - 後使用した領域の free
118
119 ``` c
120 void cudaWrite(struct CudaBuffer* buffer) {
121 //結果を取ってくるコマンドを入力する
122 //コマンドの終了待ちを行う
123 int paramCount = 0;
124 for (int i = 0; i < buffer->inputLen; i++) {
125 CUdeviceptr* deviceptr = buffer->kernelParams[paramCount++];
126 checkCudaErrors(cuMemcpyDtoH(buffer->inputData[i], *deviceptr, GET_SIZE(buffer->inputData[i])));
127 cuMemFree(*deviceptr);
128 free(deviceptr);
129 }
130
131 for (int i = 0; i < buffer->outputLen; i++) {
132 CUdeviceptr* deviceptr = buffer->kernelParams[paramCount++];
133 checkCudaErrors(cuMemcpyDtoH(buffer->outputData[i], *deviceptr, GET_SIZE(buffer->outputData[i])));
134 cuMemFree(*deviceptr);
135 free(deviceptr);
136 }
137 free(buffer->kernelParams);
138 // wait for stream
139 checkCudaErrors(cuCtxSynchronize());
140 }
141 ```
142
143 ## 次は
144 - RBTree の deletion
145 - Perl script をどうにかする
146 - 並列処理の検証をどうするか