Mercurial > hg > Papers > 2018 > parusu-master

changeset 30:9fa82713bd5c

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
Add files
author Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
date Fri, 02 Feb 2018 03:42:23 +0900
parents 39fe904d9a34
children 8793903e4a0d
files paper/fig/cudaArchitecture.graffle paper/fig/cudaArchitecture.pdf paper/fig/cudaArchitecture.xbb paper/fig/cudaDataArchitecture.graffle paper/fig/cudaDataArchitecture.pdf paper/fig/cudaDataArchitecture.xbb paper/fig/iterateTaskExec.graffle paper/fig/iterateTaskExec.pdf paper/fig/iterateTaskExec.xbb paper/gearsOS.tex paper/gpu.tex paper/interface.tex paper/master_paper.pdf paper/reference.bib paper/src/cuLaunchKernel.cbc paper/src/executorInterface.h paper/src/iterateCall.cbc paper/src/iteratePargoto.cbc paper/src/iteratorInterface.h
diffstat 19 files changed, 203 insertions(+), 15 deletions(-) [+]
line wrap: on
line diff
Binary file paper/fig/cudaArchitecture.graffle has changed
Binary file paper/fig/cudaArchitecture.pdf has changed
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/fig/cudaArchitecture.xbb	Fri Feb 02 03:42:23 2018 +0900
@@ -0,0 +1,8 @@
+%%Title: fig/cudaArchitecture.pdf
+%%Creator: extractbb 20170318
+%%BoundingBox: 0 0 446 488
+%%HiResBoundingBox: 0.000000 0.000000 446.000000 488.000000
+%%PDFVersion: 1.3
+%%Pages: 1
+%%CreationDate: Thu Feb  1 07:03:30 2018
+
Binary file paper/fig/cudaDataArchitecture.graffle has changed
Binary file paper/fig/cudaDataArchitecture.pdf has changed
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/fig/cudaDataArchitecture.xbb	Fri Feb 02 03:42:23 2018 +0900
@@ -0,0 +1,8 @@
+%%Title: fig/cudaDataArchitecture.pdf
+%%Creator: extractbb 20170318
+%%BoundingBox: 0 0 448 598
+%%HiResBoundingBox: 0.000000 0.000000 448.000000 598.000000
+%%PDFVersion: 1.3
+%%Pages: 1
+%%CreationDate: Fri Feb  2 00:45:11 2018
+
Binary file paper/fig/iterateTaskExec.graffle has changed
Binary file paper/fig/iterateTaskExec.pdf has changed
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/fig/iterateTaskExec.xbb	Fri Feb 02 03:42:23 2018 +0900
@@ -0,0 +1,8 @@
+%%Title: fig/iterateTaskExec.pdf
+%%Creator: extractbb 20170318
+%%BoundingBox: 0 0 749 310
+%%HiResBoundingBox: 0.000000 0.000000 749.000000 310.000000
+%%PDFVersion: 1.3
+%%Pages: 1
+%%CreationDate: Wed Jan 31 20:11:02 2018
+
--- a/paper/gearsOS.tex	Fri Feb 02 02:35:00 2018 +0900
+++ b/paper/gearsOS.tex	Fri Feb 02 03:42:23 2018 +0900
@@ -1,5 +1,5 @@
 \chapter{Gears OS の概念}
-Gears OS は信頼性をノーマルレベルの計算に対して保証し、拡張性をメタレベルの計算で実現することを目標に開発している OSである。
+Gears OS\cite{kkb-master} は信頼性をノーマルレベルの計算に対して保証し、拡張性をメタレベルの計算で実現することを目標に開発している OSである。
 
 Gears OS は処理の単位を Code Gear、データの単位を Data Gear と呼ばれる単位でプログラムを構成する。
 信頼性や拡張性はメタ計算として、通常の計算とは区別して記述する。
@@ -62,7 +62,7 @@
 プログラムの記述する際は、ノーマルレベルの計算の他に、メモリ管理、スレッド管理、CPUがGPUの資源管理等を記述しなければならない処理が存在する。
 これらの計算はノーマルレベルの計算と区別してメタ計算と呼ぶ。
 
-メタ計算は関数型言語では Monad\cite{moggi-monad} を用いて表現される。
+メタ計算は関数型言語では Monad\cite{moggi-monad} を用いて表現される\cite{kkb-sigos}。
 Monad は Haskell では実行時の環境を記述する構文として使われる。
 
 従来の OS では、メタ計算はシステムコールやライブラリーコールの単位で行われる。
--- a/paper/gpu.tex	Fri Feb 02 02:35:00 2018 +0900
+++ b/paper/gpu.tex	Fri Feb 02 03:42:23 2018 +0900
@@ -4,8 +4,8 @@
   % block とか grid とかの関係
 
 \chapter{CUDA への対応} 
-Gears OS では GPU での実行もサポートする。
-また、 CPU、 GPU の実行環境の切り替えは Meta Code Gear、 つまり stub Code Gear で切り替えを行う。
+Gears OS では GPU での実行もサポートする\cite{ikkun-sigos}。
+また、CPU、GPU の実行環境の切り替えは Meta Code Gear、つまり stub Code Gear で切り替えを行う。
 
 本章では、 Gears OS での CUDA 実行のサポートについて説明する。
 
@@ -47,13 +47,16 @@
 受け取ったデータ並列用の Task は Code Gear のメタレベルで kernel の実行を行う。
 
 \section{CUDAExectuor}
-CUDAExectuor は Host から Device へのデータの送信、 kernel の実行、 Device から Host への データの書き出しを行う。
+CUDAExectuor は \coderef{executorInterface} を実装しており、 Host から Device へのデータの送信、 kernel の実行、 Device から Host への データの書き出しを行う。
+
+\lstinputlisting[caption=executor Inteface, label=code:executorInterface]{./src/executorInterface.h}
 
 Gears OS では データは Data Gear で表現される。
 つまり、Host、Device 間でデータのやり取りを行うということは Data Gear を GPU のデータ領域に沿った形に適用する必要がある。
 Host から Device へデータを送信する際、 CUDA では cuMemAlloc 関数を使用して Device 側のデータ領域を確保する。
 cuMemAlloc 関数には確保するデータ領域のサイズを指定する必要がある。
-全ての Data Gear には Meta Data Gear として Data Gear のサイズを持っており、基本的にはこのサイズでデータ領域を取ればよいのだが、Data Gear によっては内部に更にポインタで Data Gear を持っている場合がある。
+全ての Data Gear には Meta Data Gear として Data Gear のサイズを持っており、基本的にはこのサイズでデータ領域を取ればよい。
+しかし、Data Gear によっては内部に更にポインタで Data Gear を持っている場合がある。
 このような Data Gear は Data Gear の実際のサイズではなく、ポインタのサイズで計算されてしまうため、そのままでは Device 用のデータ領域を確保することができない。
 
 この問題を解決するために、CUDABuffer という CUDA データ送信用の Data Gear を用意した.
@@ -62,9 +65,6 @@
 CUDABuffer に格納されている Data Gear のサイズを参照し、 cuMemAlloc 関数で Device のデータ領域を確保する。
 
 Host、Device、CUDABuffer 間の関係を\figref{cudaDataArchitecture} に示す。
-Host から Device にデータをコピーするには cuMemcpyHtoD 関数を使用して行う。
-この際に Host で指定するデータは CUDABuffer に格納されている Data Gear となる。
-kernel の実行後、結果を Device から Host にコピーする際は cuMemcpyDtoH 関数で行われる。
 
 \begin{figure}[htbp]
     \begin{center}
@@ -74,7 +74,21 @@
     \label{fig:cudaDataArchitecture}
 \end{figure}
 
+Host から Device にデータをコピーするには cuMemcpyHtoD 関数を使用して行う。
+この際に Host で指定するデータは CUDABuffer に格納されている Data Gear となる。
+kernel の実行後、結果を Device から Host にコピーする際は cuMemcpyDtoH 関数で行われる。
+
 kernel の実行はcuLaunchKernel 関数で行われる。
+cuLaunchKernel 関数には引数として各次元のblockサイズ、thread サイズ、kernel への引数等を渡す。
+
 実際にcuLaunchKernel 関数を使用している部分を \coderef{cuLaunchKernel} に示す。
 
+\lstinputlisting[caption=kernel に起動, label=code:cuLaunchKernel]{./src/cuLaunchKernel.cbc}
+
+Gears OS ではデータ並列 Task の際は Iterator Interface を持っており、 そこで指定した長さ、次元数に応じて cuLaunchKernel の引数を決定する(\coderef{cuLaunchKernel} 11-18行目)。
+
 \section{stub Code Gear による kernel の実行}
+Gears OS では stub Code Gear で CUDA による実行の切り替える。
+
+stub Code Gear での切り替えの際は CUDABuffer への Data の格納、 実行される kernel の読み込みを行う。
+実際にGPU で実行される
--- a/paper/interface.tex	Fri Feb 02 02:35:00 2018 +0900
+++ b/paper/interface.tex	Fri Feb 02 03:42:23 2018 +0900
@@ -12,6 +12,7 @@
  % interface 内部の Code Gearは 自由に 引数の Data Gear, 実装のData Gear にアクセス出来る。
  % inteface の実装の際は Code Gear に代入して書く
  % C++ で言うとCode Gearは virtual 
+% Impl の説明いれてない
 
 \chapter{Gears OS のモジュール化}
 Gears OS は stub Code Gear という Meta Code Gear で Context という全ての Code Gear と Data Gear を持った Meta Data Gear から値を取りだし、ノーマルレベルの Code Gear に値を渡す。
Binary file paper/master_paper.pdf has changed
--- a/paper/reference.bib	Fri Feb 02 02:35:00 2018 +0900
+++ b/paper/reference.bib	Fri Feb 02 03:42:23 2018 +0900
@@ -1,6 +1,5 @@
 @inproceedings{queue,
-    author = {Michael, Maged M. and Scott, Michael L.},
-    title = {Simple, Fast, and Practical Non-blocking and Blocking Concurrent Queue Algorithms},
+    author = {Michael, Maged M. and Scott, Michael L.}, title = {Simple, Fast, and Practical Non-blocking and Blocking Concurrent Queue Algorithms},
     booktitle = {Proceedings of the Fifteenth Annual ACM Symposium on Principles of Distributed Computing},
     series = {PODC '96},
     year = {1996},
@@ -36,7 +35,7 @@
 }
 
 @inproceedings{nobu-prosym,
-    author = "大城,信康 and 河野,真治",
+    author = "大城信康 and 河野真治",
     title = "Continuation based C の GCC4.6 上の実装について",
     booktitle = "第53回プログラミング・シンポジウム予稿集",
     year  = "2012",
@@ -46,7 +45,7 @@
 }
 
 @inproceedings{mitsuki-prosym,
-    author = "宮城,光希 and 河野,真治",
+    author = "宮城光希 and 河野真治",
     title = "Code Gear と Data Gear を持つ Gears OS の設計",
     booktitle = "第59回プログラミング・シンポジウム予稿集",
     year  = "2018",
@@ -65,7 +64,7 @@
 }
 
 @article{kkb-sigos,
-    author  = "小久保,翔平 and 伊波,立樹 and 河野,真治",
+    author  = "小久保翔平 and 伊波立樹 and 河野真治",
     title   = "Monad に基づくメタ計算を基本とする Gears OS の設計",
     journal = "第133回情報処理学会システムソフトウェアとオペレーティング・システム研究会(OS)",
     month   = "May",
@@ -73,13 +72,21 @@
 }
 
 @article{parusu-sigos,
-    author  = "伊波,立樹 and 東恩納, 琢偉 and 河野,真治",
+    author  = "伊波立樹 and 東恩納琢偉 and 河野真治",
     title   = "Code Gear、 Data Gear に基づく OS のプロトタイプ",
     journal = "第137回情報処理学会システムソフトウェアとオペレーティング・システム研究会(OS)",
     month   = "May",
     year    = 2016
 }
 
+@article{ikkun-sigos,
+    author  = "東恩納琢偉 and 伊波立樹 and 河野真治",
+    title   = "Gears OS における並列処理",
+    journal = "第140回 情報処理学会システムソフトウェアとオペレーティング・システム研究会(OS)",
+    month   = "May",
+    year    = 2017
+}
+
 @misc{opencl,
     title = {OpenCL | NVIDIA Developer},
     howpublished = {\url{https://developer.nvidia.com/opencl}},
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/src/cuLaunchKernel.cbc	Fri Feb 02 03:42:23 2018 +0900
@@ -0,0 +1,21 @@
+int computeblockDim(int count, int maxThreadPerBlock) {
+    return count < maxThreadPerBlock ? count : maxThreadPerBlock;
+}
+
+__code execCUDAExecutor(struct CUDAExecutor* executor, struct Context* task, __code next(...)) {
+    // check data parallelism task
+    if (task->iterate) {
+        struct MultiDimIterator* iterator = &task->iterator->iterator->MultiDimIterator;
+
+        // compute block thread size 
+        int blockDimX = computeblockDim(iterator->x, executor->maxThreadPerBlock);
+        int blockDimY = computeblockDim(iterator->y, executor->maxThreadPerBlock);
+        int blockDimZ = computeblockDim(iterator->z, executor->maxThreadPerBlock);
+
+        checkCudaErrors(cuLaunchKernel(task->function,
+                    iterator->x/blockDimX, iterator->y/blockDimY, iterator->z/blockDimZ,
+                    blockDimX, blockDimY, blockDimZ,
+                    0, NULL, (void**)executor->kernelParams, NULL));
+    }
+    ...
+}
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/src/executorInterface.h	Fri Feb 02 03:42:23 2018 +0900
@@ -0,0 +1,9 @@
+typedef struct Executor<Impl>{
+    union Data* Executor;
+    struct Context* task;
+    __code next(...);
+    // method
+    __code read(Impl* executor, struct Context* task, __code next(...));
+    __code exec(Impl* executor, struct Context* task, __code next(...));
+    __code write(Impl* executor, struct Context* task, __code next(...));
+}
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/src/iterateCall.cbc	Fri Feb 02 03:42:23 2018 +0900
@@ -0,0 +1,97 @@
+
+#include "../context.h"
+#interface "Iterator.h"
+#interface "TaskManager.h"
+#include <stdio.h>
+
+Iterator* createMultiDimIterator(struct Context* context, int x, int y, int z) {
+    struct Iterator* iterator = new Iterator();
+    struct MultiDimIterator* multiDimIterator = new MultiDimIterator();
+    iterator->iterator = (union Data*)multiDimIterator;
+    iterator->exec = C_execMultiDimIterator;
+    iterator->barrier = C_barrierMultiDimIterator;
+    multiDimIterator->x = x;
+    multiDimIterator->y = y;
+    multiDimIterator->z = z;
+    multiDimIterator->count = x * y * z;
+    multiDimIterator->counterX = 0;
+    multiDimIterator->counterY = 0;
+    multiDimIterator->counterZ = 0;
+    return iterator;
+}
+
+/**
+ * create iterateTask with index, that copy from task argument
+ * @return created iterateTask
+ * @param task      task of the copy source
+ * @x     index
+ */
+struct Context* createMultiDimIterateTask(struct Context* task, int x, int y, int z) {
+    struct Context* task1 = NEW(struct Context);
+    initContext(task1);
+    task1->taskManager = task->taskManager;
+    task1->next     = task->next;
+    task1->iterate  = 1;
+    task1->iterator = task->iterator;
+    task1->idgCount = task->idgCount;
+    task1->idg      = task->idg;
+    task1->maxIdg   = task->maxIdg;
+    for(int i = task1->idg; i < task1->maxIdg; i++) {
+        task1->data[i] = task->data[i];
+    }
+
+    // create index data gear and register input data to iterate task
+    struct MultiDim* multiDim = &ALLOCATE_DATA_GEAR(task1, MultiDim)->MultiDim;
+    multiDim->x = x;
+    multiDim->y = y;
+    multiDim->z = z;
+    task1->data[task1->maxIdg++] = (union Data*)multiDim;
+    task1->odg      = task->odg + 1;
+    task1->maxOdg   = task->maxOdg + 1;
+    for (int i = task1->odg; i < task1->maxOdg; i++) {
+        task1->data[i] = task->data[i-1];
+    }
+
+    return task1;
+}
+
+__code execMultiDimIterator(struct MultiDimIterator* iterator, struct Context* task, int numGPU, __code next(...)) {
+    // No GPU device
+    if (numGPU == 0) {
+        goto execMultiDimIterator1();
+    }
+    task->iterate = 1;
+    task->gpu = 1;
+    struct TaskManager* taskManager = task->taskManager;
+    goto taskManager->spawn(task, next(...));
+}
+
+__code execMultiDimIterator1(struct MultiDimIterator* iterator, struct Context* task, __code next(...)) {
+    int x = iterator->counterX;
+    int y = iterator->counterY;
+    int z = iterator->counterZ;
+    struct Context* iterateTask = createMultiDimIterateTask(task, x, y, z);
+    struct TaskManager* taskManager = task->taskManager;
+    goto taskManager->spawn(iterateTask, execMultiDimIterator2);
+}
+
+__code execMultiDimIterator2(struct MultiDimIterator* iterator, struct Context* task, __code next(...)) {
+    if (++iterator->counterX >= iterator->x) {
+        iterator->counterX = 0;
+        if (++iterator->counterY >= iterator->y) {
+            iterator->counterY = 0;
+            if (++iterator->counterZ >= iterator->z) {
+                iterator->counterZ = 0;
+                goto next(...);
+            }
+        }
+    }
+    goto execMultiDimIterator1();
+}
+
+__code barrierMultiDimIterator(struct MultiDimIterator* iterator, struct Context* task, __code next(...), __code whenWait(...)) {
+    if (task->gpu || __sync_fetch_and_sub(&iterator->count, 1) == 1) {
+        goto next(...);
+    }
+    goto whenWait(...);
+}
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/src/iteratePargoto.cbc	Fri Feb 02 03:42:23 2018 +0900
@@ -0,0 +1,4 @@
+__code code1() {
+    par goto printIterator(input, output, iterate(2), __exit);
+    goto code2();
+}
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/src/iteratorInterface.h	Fri Feb 02 03:42:23 2018 +0900
@@ -0,0 +1,11 @@
+typedef struct Iterator<Impl>{
+        union Data* iterator;
+        struct Context* task;
+        int numGPU;
+        __code next(...);
+        __code whenWait(...);
+
+        // method
+        __code exec(Impl* iterator, struct Context* task, int numGPU, __code next(...));
+        __code barrier(Impl* iterator, struct Context* task, __code next(...), __code whenWait(...));
+} Iterator;