Mercurial > hg > Members > innparusu > slides

annotate 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
Ignore whitespace changes - Everywhere: Within whitespace: At end of lines:
rev   line source
28
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
1 title: Gears OS
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
2 author: Tatsuki IHA
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
3 profile:
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
4 lang: Japanese
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
5 code-engine: coderay
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
6
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
7 ## 研究目的
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
8 - 当研究室では 処理の単位を Code Gear、 データの単位を Data Gear を用いて 信頼性が高い並列処理を行う Gears OS を開発している
382cd93f2a60 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 によって依存関係が決定し、それにそって並列実行を行う.
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
10 - 信頼性の確保はモデルチェック、検証等を使用して行う。この信頼性のための計算は通常の計算とは別の階層のメタ計算として記述する。
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
11 - また、 メタ計算は信頼性の他に CPU, GPU などの実行環境の切り替え, データ拡張等の柔軟性を提供する。
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
12 - 本研究では、 Gears OS の並列処理機構の実装を行う。また、並列処理の検証をメタ計算として記述することで、 並列処理の信頼性を保証する。
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
13
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
14 ## 今週
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
15 - cuda.c の抽象化
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
16 - 今までは一つ一つの例題で cuda.c を書き換えてたので、、
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
17
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
18 ## CudaBuffer
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
19 - cuda 用の引数, DataGear の値を持った構造体を定義
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
20
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
21 ``` c
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
22 union Data {
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
23 .....
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
24 struct CudaBuffer {
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
25 void** kernelParams;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
26 int inputLen;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
27 int outputLen;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
28 union Data** inputData;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
29 union Data** outputData;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
30 } CudaBuffer;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
31 }
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
32 ```
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
33
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
34 ## 呼び出し
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
35 - cudaの実行は stubで呼ばれる
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
36 - cudaで変換する Gode Gear を指定する
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
37
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
38 ``` c
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
39 __code twice_stub(struct Context* context) {
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
40 #ifdef USE_CUDAWorker
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
41 if (context->gpu) {
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
42 struct Array* array = &context->data[context->idg]->Array;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
43 struct CudaBuffer* buffer = &ALLOCATE(context, CudaBuffer)->CudaBuffer;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
44 buffer->inputData = (union Data**)ALLOCATE_PTR_ARRAY(context, Array, 2);
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
45 buffer->inputData[0] = (union Data*)array->array;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
46 buffer->inputData[1] = (union Data*)array;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
47 buffer->outputData = NULL;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
48 buffer->inputLen = 2;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
49 buffer->outputLen = 0;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
50 cudaExec(context, buffer, "c/examples/twice/CUDAtwice.ptx", "twice");
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
51 //continuationにそってGPUworkerに戻る
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
52 goto meta(context, context->next);
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
53 }
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
54 #endif
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
55 goto twice(context,
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
56 &context->data[context->idg]->Array,
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
57 &context->data[context->idg+1]->MultiDim,
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
58 context->next,
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
59 Gearef(context, LoopCounter));
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
60 }
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
61 ```
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
62
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
63 ## cudaRead
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
64 - buffer に入れた DataGear を cuda に allocate, 送信
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
65
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
66 ``` c
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
67 void cudaRead(struct CudaBuffer* buffer) {
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
68 buffer->kernelParams = (void **)calloc(buffer->inputLen + buffer->outputLen, sizeof(void *));
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
69 int paramCount = 0;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
70 for (int i = 0; i < buffer->inputLen; i++) {
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
71 CUdeviceptr* deviceptr = (CUdeviceptr *)calloc(1, sizeof(CUdeviceptr));
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
72 // memory allocate
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
73 checkCudaErrors(cuMemAlloc(deviceptr, GET_SIZE(buffer->inputData[i])));
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
74 // Synchronous data transfer(host to device)
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
75 checkCudaErrors(cuMemcpyHtoD(*deviceptr, buffer->inputData[i], GET_SIZE(buffer->inputData[i])));
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
76 buffer->kernelParams[paramCount++] = deviceptr;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
77 }
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
78
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
79 for (int i = 0; i < buffer->outputLen; i++) {
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
80 CUdeviceptr* deviceptr = (CUdeviceptr *)calloc(1, sizeof(CUdeviceptr));
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
81 // memory allocate
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
82 checkCudaErrors(cuMemAlloc(deviceptr, GET_SIZE(buffer->outputData[i])));
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
83 // Synchronous data transfer(host to device)
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
84 checkCudaErrors(cuMemcpyHtoD(*deviceptr, buffer->outputData[i], GET_SIZE(buffer->outputData[i])));
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
85 buffer->kernelParams[paramCount++] = deviceptr;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
86 }
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
87 }
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
88 ```
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
89
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
90 ## cudaExec2
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
91 - *cudaLaunch* とか *cudaLaunchKernel* にしたかった. すでにcuda側で使われているらしい(この辺の名前は後で要修正)
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
92 - cudaRead で設定した kernelParams(CUdeviceptr) を launchKernel に引き渡す
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
93
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
94 ``` c
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
95 void cudaExec2(struct Context* context, struct CudaBuffer* buffer) {
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
96 // Asynchronous launch kernel
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
97 context->num_exec = 1;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
98 if (context->iterate) {
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
99 struct MultiDimIterator* iterator = &context->iterator->iterator->MultiDimIterator;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
100 checkCudaErrors(cuLaunchKernel(context->function,
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
101 iterator->x, iterator->y, iterator->z,
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
102 1, 1, 1,
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
103 0, NULL, buffer->kernelParams, NULL));
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
104
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
105 } else {
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
106 checkCudaErrors(cuLaunchKernel(context->function,
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
107 1, 1, 1,
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
108 1, 1, 1,
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
109 0, NULL, buffer->kernelParams, NULL));
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
110 }
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
111 }
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
112 ```
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
113
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
114 ## cudaWrite
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
115 - 実行した結果を指定した DataGear に書き出す
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
116 - input にそのまま結果を書き出す場合もある. flip 的なものも実装しないと
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
117 - 後使用した領域の free
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
118
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
119 ``` c
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
120 void cudaWrite(struct CudaBuffer* buffer) {
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
121 //結果を取ってくるコマンドを入力する
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
122 //コマンドの終了待ちを行う
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
123 int paramCount = 0;
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
124 for (int i = 0; i < buffer->inputLen; i++) {
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
125 CUdeviceptr* deviceptr = buffer->kernelParams[paramCount++];
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
126 checkCudaErrors(cuMemcpyDtoH(buffer->inputData[i], *deviceptr, GET_SIZE(buffer->inputData[i])));
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
127 cuMemFree(*deviceptr);
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
128 free(deviceptr);
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
129 }
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
130
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
131 for (int i = 0; i < buffer->outputLen; i++) {
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
132 CUdeviceptr* deviceptr = buffer->kernelParams[paramCount++];
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
133 checkCudaErrors(cuMemcpyDtoH(buffer->outputData[i], *deviceptr, GET_SIZE(buffer->outputData[i])));
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
134 cuMemFree(*deviceptr);
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
135 free(deviceptr);
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
136 }
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
137 free(buffer->kernelParams);
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
138 // wait for stream
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
139 checkCudaErrors(cuCtxSynchronize());
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
140 }
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
141 ```
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
142
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
143 ## 次は
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
144 - RBTree の deletion
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
145 - Perl script をどうにかする
382cd93f2a60 Update slide
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
146 - 並列処理の検証をどうするか