--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/2017/2017_10_17/slide.md	Tue Oct 31 18:16:16 2017 +0900
@@ -0,0 +1,146 @@
+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 の並列処理機構の実装を行う。また、並列処理の検証をメタ計算として記述することで、 並列処理の信頼性を保証する。
+
+## 今週
+- cuda.c の抽象化
+  - 今までは一つ一つの例題で cuda.c を書き換えてたので、、
+
+## CudaBuffer
+- cuda 用の引数, DataGear の値を持った構造体を定義
+
+``` c
+union Data {
+    .....
+    struct CudaBuffer {
+        void** kernelParams;
+        int inputLen;
+        int outputLen;
+        union Data** inputData;
+        union Data** outputData;
+    } CudaBuffer;
+}
+```
+
+## 呼び出し
+- cudaの実行は stubで呼ばれる
+- cudaで変換する Gode Gear を指定する
+
+``` c
+__code twice_stub(struct Context* context) {
+#ifdef USE_CUDAWorker
+    if (context->gpu) {
+        struct Array* array = &context->data[context->idg]->Array;
+        struct CudaBuffer* buffer = &ALLOCATE(context, CudaBuffer)->CudaBuffer;
+        buffer->inputData = (union Data**)ALLOCATE_PTR_ARRAY(context, Array, 2);
+        buffer->inputData[0] = (union Data*)array->array;
+        buffer->inputData[1] = (union Data*)array;
+        buffer->outputData = NULL;
+        buffer->inputLen = 2;
+        buffer->outputLen = 0;
+        cudaExec(context, buffer, "c/examples/twice/CUDAtwice.ptx", "twice");
+        //continuationにそってGPUworkerに戻る
+        goto meta(context, context->next);
+    }
+#endif
+    goto twice(context,
+               &context->data[context->idg]->Array,
+               &context->data[context->idg+1]->MultiDim,
+               context->next,
+               Gearef(context, LoopCounter));
+}
+```
+
+## cudaRead
+- buffer に入れた DataGear を cuda に allocate, 送信
+
+``` c
+void cudaRead(struct CudaBuffer* buffer) {
+    buffer->kernelParams = (void **)calloc(buffer->inputLen + buffer->outputLen, sizeof(void *));
+    int paramCount = 0;
+    for (int i = 0; i < buffer->inputLen; i++) {
+        CUdeviceptr* deviceptr = (CUdeviceptr *)calloc(1, sizeof(CUdeviceptr));
+        // memory allocate
+        checkCudaErrors(cuMemAlloc(deviceptr, GET_SIZE(buffer->inputData[i])));
+        // Synchronous data transfer(host to device)
+        checkCudaErrors(cuMemcpyHtoD(*deviceptr, buffer->inputData[i], GET_SIZE(buffer->inputData[i])));
+        buffer->kernelParams[paramCount++] = deviceptr;
+    }
+
+    for (int i = 0; i < buffer->outputLen; i++) {
+        CUdeviceptr* deviceptr = (CUdeviceptr *)calloc(1, sizeof(CUdeviceptr));
+        // memory allocate
+        checkCudaErrors(cuMemAlloc(deviceptr, GET_SIZE(buffer->outputData[i])));
+        // Synchronous data transfer(host to device)
+        checkCudaErrors(cuMemcpyHtoD(*deviceptr, buffer->outputData[i], GET_SIZE(buffer->outputData[i])));
+        buffer->kernelParams[paramCount++] = deviceptr;
+    }
+}
+```
+
+## cudaExec2
+- *cudaLaunch* とか *cudaLaunchKernel* にしたかった. すでにcuda側で使われているらしい(この辺の名前は後で要修正)
+- cudaRead で設定した kernelParams(CUdeviceptr) を launchKernel に引き渡す
+
+``` c
+void cudaExec2(struct Context* context, struct CudaBuffer* buffer) {
+    // Asynchronous launch kernel
+    context->num_exec = 1;
+    if (context->iterate) {
+        struct MultiDimIterator* iterator = &context->iterator->iterator->MultiDimIterator;
+        checkCudaErrors(cuLaunchKernel(context->function,
+                    iterator->x, iterator->y, iterator->z,
+                    1, 1, 1,
+                    0, NULL, buffer->kernelParams, NULL));
+
+    } else {
+        checkCudaErrors(cuLaunchKernel(context->function,
+                    1, 1, 1,
+                    1, 1, 1,
+                    0, NULL, buffer->kernelParams, NULL));
+    }
+}
+```
+
+## cudaWrite
+- 実行した結果を指定した DataGear に書き出す
+  - input にそのまま結果を書き出す場合もある. flip 的なものも実装しないと
+- 後使用した領域の free
+
+``` c
+void cudaWrite(struct CudaBuffer* buffer) {
+    //結果を取ってくるコマンドを入力する
+    //コマンドの終了待ちを行う   
+    int paramCount = 0;
+    for (int i = 0; i < buffer->inputLen; i++) {
+        CUdeviceptr* deviceptr =  buffer->kernelParams[paramCount++];
+        checkCudaErrors(cuMemcpyDtoH(buffer->inputData[i], *deviceptr, GET_SIZE(buffer->inputData[i])));
+        cuMemFree(*deviceptr);
+        free(deviceptr);
+    }
+
+    for (int i = 0; i < buffer->outputLen; i++) {
+        CUdeviceptr* deviceptr =  buffer->kernelParams[paramCount++];
+        checkCudaErrors(cuMemcpyDtoH(buffer->outputData[i], *deviceptr, GET_SIZE(buffer->outputData[i])));
+        cuMemFree(*deviceptr);
+        free(deviceptr);
+    }
+    free(buffer->kernelParams);
+    // wait for stream
+    checkCudaErrors(cuCtxSynchronize());
+}
+```
+
+## 次は
+- RBTree の deletion
+- Perl script をどうにかする
+- 並列処理の検証をどうするか