Mercurial > hg > Papers > 2016 > kkb-master

changeset 11:12d1c2f53258

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
revision
author Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
date Sun, 14 Feb 2016 07:02:11 +0900
parents e4df300f45b6
children 2d6755608f67
files abstract.tex abstract_eng.tex cbc.tex cerium.tex gearsos.tex images/bitonic.bb images/bitonic.pdf images/bitonic_sort_01.bb images/bitonic_sort_01.pdf images/bitonic_sort_02.bb images/bitonic_sort_02.pdf images/bitonic_sort_03.bb images/bitonic_sort_03.pdf images/codesegment2.bb images/codesegment2.pdf images/images.graffle images/stream.bb images/stream.pdf images/wordcount.bb images/wordcount.pdf master_paper.pdf master_paper.sty master_paper.tex result/bitonic_sort/box result/bitonic_sort/cpu_1 result/bitonic_sort/cpu_12 result/bitonic_sort/cpu_2 result/bitonic_sort/cpu_4 result/bitonic_sort/cpu_8 result/bitonic_sort/gpu src/init_twice_cerium.cc src/sample.c src/sample_transform.c src/sync_dequeue.c src/sync_enqueue.c src/twice_cerium.cc src/twice_cuda.cu
diffstat 37 files changed, 747 insertions(+), 108 deletions(-) [+]
line wrap: on
line diff
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/abstract.tex	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,7 @@
+\begin{abstract}
+  本研究では Cerium を開発して得られた知見から Code Segment と Data Segment を用いた並列フレームワークの開発を行なっている。
+  Code Segment と Data Segment は処理とデータの単位である。
+  今回設計した Gears OS ではプログラムを Code Segment と Data Segment で記述する。
+  Code Segment と Data Segment で記述することでプログラム全体の並列度を高めて効率的に並列処理することを可能にする。
+  本論文では Gears OS の基本的な機能を設計し、CbC(Continuation based C) を用いて実装する。
+\end{abstract}
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/abstract_eng.tex	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,2 @@
+\begin{abstract_eng}
+\end{abstract_eng}
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/cbc.tex	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,47 @@
+\chapter{CbC}
+CbC は C から for 文、while 文といったループ制御構文や関数呼び出しを取り除き、Code Segment と goto による軽量継続を導入している。
+図:\ref{fig:cs} は goto による Code Segment の遷移を表したものである。
+
+Gears OS の実装には LLVM/Clang 上に実装した CbC を用いる。
+
+\begin{figure}[!h]
+  \begin{center}
+    \includegraphics[scale=0.7]{./images/codesegment2.pdf}
+  \end{center}
+  \caption{goto による Code Segment 間の継続}
+  \label{fig:cs}
+\end{figure}
+
+\section{Code Segment}
+CbC では処理の単位として Code Segment を用いる。
+Code Segment は CbC における最も基本的な処理単位であり、C の関数と異なり戻り値を持たない。
+Code Segment の宣言は C の関数の構文と同じように行い、型に \_\_code を用いる。
+前述した通り、Code Segment は戻り値を持たないので \_\_code はそれが関数ではなく Code Segment であることを示すフラグのようなものである。
+Code Segment の処理内容の定義も C の関数同様に行うが、CbC にはループ制御構文が存在しないのでループ処理は自分自身への再帰的な継続を行うことで実現する。
+
+現在の Code Segment から次の Code Segment への処理の移動は goto の後に Code Segment 名と引数を並べて記述するという構文を用いて行う。
+この goto による処理の遷移を継続と呼ぶ。
+C において関数呼び出しを繰り返し行う場合、呼び出された関数の引数の数だけスタックに値が積まれていくが、戻り値を持たない Code Segment ではスタックに値を積んでいく必要が無くスタックは変更されない。
+このようなスタックに値を積まない継続を軽量継続と呼ぶ。
+この軽量継続により並列化、ループ制御、関数コールとスタックの操作を意識した最適化がソースコードレベルで行えるようになる。
+
+\section{プロトタイプ宣言の自動化}
+Code Segment の処理単位は小さく、目的の計算を実現するためには多くの Code Segment を書く必要がある。
+Code Segment と同じ数だけプロトタイプ宣言を書く必要があり、好ましくない。
+また、tail call elimination を強制するためにはプロトタイプの宣言を正確に記述することを要求するためプログラマに対する負担が大きい。
+つまり、プロトタイプ宣言を自動的に行うようにすることで tail call elimnation の条件を安定して満たすことができ、プログラマの負担も減らすことができる。
+
+プロトタイプ宣言の自動化は、パーサーが Code Segment への継続の解析を行なった際にプロトタイプ宣言の有無を確認し、存在しない場合に接続先の Code Segment のプロトタイプ宣言を生成するというようにして行う。
+
+\section{Gear OS の構文サポート}
+Gears OS では Context から必要なデータを取り出して処理を行う。
+しかし、Context を直接扱うのはセキュリティ的に好ましくない。
+そこで Context から必要なデータを取り出して Code Segment に接続する stub を定義する。
+stub は接続される Code Segment から推論することが可能である。
+また、Code Segment の遷移には Meta Code Segment を挟む。
+Meta Code Segment への接続も省略して記述できるようにする。
+省略形のソースコード:\ref{sample}から実際にコンパイルされるソースコード:\ref{sample_trans}へ変換される。
+
+\lstinputlisting[label=sample, caption=省略形]{src/sample.c}
+\newpage
+\lstinputlisting[label=sample_trans, caption=変換後]{src/sample_transform.c}
--- a/cerium.tex	Tue Feb 09 19:16:59 2016 +0900
+++ b/cerium.tex	Sun Feb 14 07:02:11 2016 +0900
@@ -35,9 +35,10 @@
 TaskManger によって依存関係が解決されると ActiveTaskList に移され、実行可能な状態となる。
 実行可能な状態となった Task は set\_cpu で指定された Device に対応した Scheduler に転送し実行される。
 図:\ref{fig:createTask}は Cerium が Task を生成/実行する場合のクラスの構成である。
+
 \begin{figure}[!ht]
   \begin{center}
-    \includegraphics[scale=0.6]{./images/createTask.pdf}
+    \includegraphics[scale=0.6]{images/createTask.pdf}
   \end{center}
   \caption{TaskManager}
   \label{fig:createTask}
@@ -65,24 +66,39 @@
   依存関係が解決された Task は実行可能な状態となる。
 \end{itemize}
 
+ソースコード:\ref{inittwice_cerium} に Task を生成する例題を示す。
+
+input data として int 型の配列を受け取り、各要素を2倍にして output data に格納する twice という例題である。
+CPU を用いてデータ並列で実行する Task を生成している。
+set\_cpu で GPU を指定することで GPU を用いて実行される。
+
+\lstinputlisting[label=inittwice_cerium, caption=Task の生成]{src/init_twice_cerium.cc}
+
+CPU 上で実行される Task, GPU 上で実行される kernel はソースコード:\ref{twice_task_cerium}, ソースコード\ref{twice_task_cuda} の通りになる。
+
+Task には実行時に必要なデータが格納されている SchedTask, 設定した Input/Output Data が格納されている Buffer が渡される。
+
+\lstinputlisting[label=twice_task_cerium, caption=実行される Task]{src/twice_cerium.cc}
+\lstinputlisting[label=twice_task_cuda, caption=実行される kernel]{src/twice_cuda.cu}
+
 \section{Task のパイプライン実行}
 \begin{figure}[htpd]
   \begin{minipage}[t]{0.5\hsize}
     \begin{center}
-      \includegraphics[scale=0.4]{./images/cell_arch.pdf}
+      \includegraphics[scale=0.4]{images/cell_arch.pdf}
     \end{center}
     \caption{Cell Architecture}
     \label{fig:cellarch}
   \end{minipage}
   \begin{minipage}[t]{0.5\hsize}
     \begin{center}
-      \includegraphics[scale=0.4]{./images/gpu_arch.pdf}
+      \includegraphics[scale=0.4]{images/gpu_arch.pdf}
     \end{center}
     \caption{GPU Architecture}
     \label{fig:gpuarch}
   \end{minipage}
 \end{figure}
-Cell(図:\ref{fig:cellarch}) や GPU(図:\ref{fig:gpuarch}) のように異なるメモリ空間を持つ Device を計算資源として利用するにはデータの転送が必要になる。
+Cell(図:\ref{fig:cellarch})や GPU(図:\ref{fig:gpuarch})のように異なるメモリ空間を持つ Device を計算資源として利用するにはデータの転送が必要になる。
 このデータ転送がボトルネックとなり、並列度が低下してしまう。
 転送処理をオーバーラップし、並列度を維持するために Cerium では Task のパイプライン実行をサポートしている。
 
@@ -91,10 +107,268 @@
 TaskList でまとめられている Task は依存関係が解決されているので自由にパイプラインを組むことが可能である。
 実行完了は TaskList 毎ではなく、Task 毎に通知される。
 図:\ref{fig:scheduler}は TaskList を受け取り、Task をパイプラインで処理していく様子である。
+
 \begin{figure}[ht]
   \begin{center}
-    \includegraphics[scale=0.6]{./images/scheduler.pdf}
+    \includegraphics[scale=0.6]{images/scheduler.pdf}
   \end{center}
   \caption{Scheduler}
   \label{fig:scheduler}
 \end{figure}
+
+\section{マルチコアへの対応}
+Cell には MailBox という機能がある。
+MailBox を用いることで双方向のデータの受け渡しが可能になる。
+FIFO キュー構造を持つ MailBox に対応させる形で Synchronized Queue 用いて Multi Core CPU 用の TaskManager に MailBox を移植した。
+Synchronized Queue は Queue を操作しているスレッドが常に1つになるようにバイナリセマフォを用いて制御する。
+
+Cell では MailBox 以外に DMA 転送を使用してデータの受け渡しすることができる。
+DMA 転送は CPU を介さずに周辺装置とメモリ間でデータ転送を行う方式である。
+Cerium では DMA 転送を用いて Cell で実行することが可能である。
+Multi Core CPU 上で実行する場合、メモリ空間を共有しているので DMA 転送を行なっている部分をポインタ渡しを行うように修正し、直接アクセスさせることでデータ転送の速度の向上が見込める。
+
+\section{データ並列による実行}
+並列処理の方法としてタスク並列とデータ並列の2つがある。
+
+タスク並列は Task 毎にデータを準備し、管理スレッドが個別に生成した Task を CPU に割り当てることで並列処理する方法である。
+異なる処理を同時に実行することができるというメリットがあるが、データ群の各要素に対して同じ処理をしたいときタスク並列では要素毎に同じ処理をする Task を生成する必要があり、ほとんど同一な大量の Task によってメモリを圧迫する場合がある。
+また、大量な Task の生成自体が大きなオーバーヘッドになる。
+
+データ並列はあるデータ群を大量な Task で共有し、Task 実行時に処理範囲を計算し、その範囲にのみ処理を行うことで並列処理する方法である。
+実行スレッドで Task の生成・実行が行われるので、メモリの圧迫や Task 生成によるオーバーヘッドを抑えられる。
+並列化部分が全て同じ処理である場合、データ並列による実行のほうがタスク並列より有効である。
+
+いままで Cerium における並列処理はタスク並列だったが、データ並列のよる実行もサポートした。
+
+データ並列による実行では処理範囲を決定するための情報として index が必要になる。
+CPU による実行では SchedTask を参照(ソースコード:\ref{twice_task_cerium} 23行目)、GPU による実行では組み込み変数を参照(ソースコード:\ref{twice_task_cuda} 11行目)することで index を取得することができる。
+
+データの長さが10、CPU の数が4でデータ並列による実行をした場合の index の割当は表\ref{table:dataparallel_index} の通りになる。
+
+\begin{table}[htpb]
+  \begin{center}
+    \small
+    \begin{tabular}[htpb]{|c||c|c|c|c|}
+      \hline
+      stage & CPU0 & CPU1 & CPU2 & CPU3 \\
+      \hline
+      \hline
+      1 & 0 & 1 & 2 & 3 \\
+      \hline
+      2 & 4 & 5 & 6 & 7 \\
+      \hline
+      3 & 8 & 9 &   &   \\
+      \hline
+    \end{tabular}
+    \caption{index の割り当て}
+    \label{table:dataparallel_index}
+  \end{center}
+\end{table}
+
+\newpage
+
+\section{GPGPU への対応}
+GPU の演算資源を Cerium から利用するために OpenCL, CUDA を用いた GpuScheduler, CudaScheduler を実装した。
+OpenCL, CUDA 単体を用いて GPGPU を行う場合、依存関係を記述する必要がある
+しかし、Cerium には依存関係を解決する TaskManager があるので GpuScheduler, CudaScheduler は受け取った TaskList を元に GPU を制御して GPGPU を行えばよい。
+
+GPU はメモリ空間が異なる(図\ref{fig:gpuarch})のでデータ転送が大きなオーバーヘッドになる。
+なので、kernel 実行中にデータ転送を行うなどしてデータ転送をオーバーラップする必要がある。
+CUDA で GPU を制御するには同期命令を使う方法と非同期命令を使う方法があるが、同期命令ではデータ転送をオーバーラップすることが出来ないので非同期命令を利用して GPU を制御する。
+非同期命令は Stream に発行することで利用することができる。
+Stream に発行された命令は発行された順序で実行される。
+非同期命令と Stream を利用してデータ転送をオーバラップするには複数の Stream を準備して、Host から Device への転送・kernel の実行・Device から Host への転送を1セットとして各 Stream に発行することで実現できる。
+同期命令を使う場合と非同期命令を使う場合の実行の様子は図:\ref{fig:stream}の通りである。
+
+\begin{figure}[ht]
+  \begin{center}
+    \includegraphics[scale=0.4]{images/stream.pdf}
+  \end{center}
+  \caption{Overlap Data Transfer}
+  \label{fig:stream}
+\end{figure}
+
+\newpage
+
+\section{Cerium の評価}
+Bitonic Sort, Word Count, Fast Fourier Transform(FFT) の3つの例題を用いて Cerium を評価する。
+
+測定環境は表:\ref{table:firefly}、測定に用いる GPU は表\ref{table:k5000}の通りである。
+
+\begin{table}[htpb]
+  \begin{center}
+    \small
+    \begin{tabular}[htpb]{|c||c|}
+      \hline
+      Model & MacPro Mid 2010 \\
+      \hline
+      OS & Mac OS X 10.10.\\
+      \hline
+      Memory & 16GB \\
+      \hline
+      CPU & 2 x 6-Core Intel Xeon 2.66GHz \\
+      \hline
+      GPU & NVIDIA Quadro K5000 \\
+      \hline
+    \end{tabular}
+    \caption{測定環境}
+    \label{table:firefly}
+  \end{center}
+\end{table}
+
+\begin{table}[htpb]
+  \begin{center}
+    \small
+    \begin{tabular}[htpb]{|c||c|}
+      \hline
+      Cores & 1536 \\
+      \hline
+      Clock Speed & 706MHz \\
+      \hline
+      Memory Size & 4GB GDDR5 \\
+      \hline
+      Memory Bandwidth & 173 GB/s \\
+      \hline
+    \end{tabular}
+    \caption{Quadro K5000}
+    \label{table:k5000}
+  \end{center}
+\end{table}
+
+\subsection{Bitonic Sort}
+Bitonic Sort は並列処理に向いたソートアルゴリズムである。
+代表的なソートアルゴリズムである Quick Sort も並列処理することが、Quick Sort はソートの過程で並列度が変動するので自明な台数効果が出づらい。
+一方、Bitonic Sort は最初から最後まで並列度が変わらずに並列処理による恩恵を得やすい。
+図:\ref{fig:bitonic}は要素数8のデータに対する Bitonic Sort のソーティングネットワークである。
+
+\newpage
+
+\begin{figure}[!h]
+  \begin{center}
+    \includegraphics[scale=0.5]{images/bitonic.pdf}
+  \end{center}
+  \caption{Sorting Network : bitonic sort}
+  \label{fig:bitonic}
+\end{figure}
+
+Bitonic Sort の並列処理に用いられる Task は2点間のの比較・交換を行うだけの小さい処理なので、1コア当たりのクロック数よりもコアの数が結果に与える影響が大きいと考えられる。
+よって、通信時間を考慮しなければ CPU よりコア数が多い GPU が有利となる。
+
+Cerium を用いて Bitonic Sort を実装し、要素数$2^{20}$のデータに対してコア数・プロセッサの種類を変更して測定を行なった結果は表\ref{table:bitonic}、図\ref{fig:bitonic_box}の通りである。
+
+\begin{table}[!h]
+  \begin{center}
+    \small
+    \begin{tabular}[htpb]{|c||c|}
+      \hline
+      Processor & Time(ms) \\
+      \hline
+      \hline
+      1 CPU & 6143 \\
+      \hline
+      2 CPUs & 4633 \\
+      \hline
+      4 CPUs & 2557 \\
+      \hline
+      8 CPUs & 1630 \\
+      \hline
+      12 CPUs & 1318 \\
+      \hline
+      GPU &  155 \\
+      \hline
+    \end{tabular}
+    \caption{要素数$2^{20}$に対するソート}
+    \label{table:bitonic}
+  \end{center}
+\end{table}
+
+\newpage
+
+\begin{figure}[!h]
+  \begin{center}
+    \includegraphics[scale=1.0]{images/bitonic_sort_03.pdf}
+  \end{center}
+  \caption{要素数$2^{20}$に対するソート}
+  \label{fig:bitonic_box}
+\end{figure}
+
+1 CPU と 12 CPU では約4.6倍の速度向上が見られた。
+これは Task の粒度が小さいため1コア当たりのクロック数の高さが活かしづらく、並列化によるオーバーヘッドが結果に影響を与えたと考えられる。
+CPU を用いた並列化には Task の粒度をある程度大きくし1コア当たりの仕事量を増やして CPU のクロック数の高さを活かすことが重要であることがわかる。
+
+12 CPU と GPU では約8.5倍の速度向上が見られた。
+GPU の特徴であるコア数の多さによって CPU より高い並列度を発揮した結果だと考えられる。
+GPU の場合はその超並列性を活かすため Task を細かく分割することが重要であることがわかる。
+
+測定結果から CPU と GPU で並列化の方法を変更する必要があることがわかった。
+Cerium を用いてヘテロジニアス環境で並列実行する場合、混在しているプロセッサの特徴に合わせたスケジューリングを行い並列実行するように Scheduler を改良する必要がある。
+
+次に要素数も変更して測定を行なった。
+結果は図:\ref{fig:bitonic_result_2}、図:\ref{fig:bitonic_result_1}の通りである。
+
+\newpage
+
+\begin{figure}[!h]
+  \begin{center}
+    \includegraphics[scale=1.0]{images/bitonic_sort_02.pdf}
+  \end{center}
+  \caption{Bitonic Sort(from $2^{14}$ to $2^{17}$)}
+  \label{fig:bitonic_result_2}
+\end{figure}
+
+\begin{figure}[!h]
+  \begin{center}
+    \includegraphics[scale=1.0]{images/bitonic_sort_01.pdf}
+  \end{center}
+  \caption{Bitonic Sort(from $2^{14}$ to $2^{20}$)}
+  \label{fig:bitonic_result_1}
+\end{figure}
+
+GPGPU では通信時間を考慮する必要がある。
+図:\ref{fig:bitonic_result_2}を見ると要素数$2^{14}$のソートでは GPU が一番遅い。
+これはソート処理の時間より通信時間が大きいことが原因であると考えられる。
+通信時間を含めた処理時間が GPU が CPU を上回るのは要素数$2^{17}$を超えてからである。
+
+\subsection{Word Count}
+並列処理を行う際に Task を大量に生成する場合がある。
+一度に大量の Task を生成してしまうと Task がメモリを圧迫して処理速度が著しく低下する。
+改善策としては Task の生成と実行を平行して行えばよい。
+Cerium では Task を生成する Task を記述することが可能なので Task の生成と実行を平行して行うことができる。
+
+Word Count を並列処理する場合、与えられたテキストを分割して、分割されたデータごとに並列処理を行う。
+分割したデータの数だけ Task が必要なのでテキストサイズによっては一度に Task を生成するとメモリを圧迫する可能性がある。
+よって、Task を生成する Task が必要になる。
+Word Count の処理の流れは図\ref{fig:wordcount}の通りである。
+
+\begin{figure}[!h]
+  \begin{center}
+    \includegraphics[scale=0.7]{images/wordcount.pdf}
+  \end{center}
+  \caption{Word Count の流れ}
+  \label{fig:wordcount}
+\end{figure}
+
+Cerium が複雑な並列処理を記述可能でその上、高い並列度を保てること示すため Cerium 上に Word Count を実装し、測定を行なった。
+結果は図\ref{}
+
+\subsection{FFT}
+FFT は信号処理や画像処理、大規模シミュレーションに至るまで幅広い分野で活用されている計算である。
+バタフライ演算などの計算の性質上、大量の演算資源を持つ GPU と相性が良い。
+Cerium に実装した GPU 実行機構の評価を行うために適切な例題であると考えられる。
+
+\section{Cerium の問題点}
+Cerium では Task 間の依存関係を記述することで並列処理を実現する。
+しかし、本来 Task はデータが揃えば実行可能になるものである。
+Task 間の依存関係だけでは待っている Task が不正な処理を行いデータがおかしくなっても Task の終了は通知され、そのまま処理が続行されてしまう。
+その場合、どこでデータがおかしくなったのか特定するのは難しくデバッグに多くの時間が取られてしまう。
+また、Cerium の Task は汎用ポインタでデータを受け取るので型の情報がない。
+型の情報がないので Task を実行するまで正しい型かどうか判断することが出来ない。
+不正な型でも強制的に型変換され実行されるのでデータの構造を破壊する可能性がある。
+型システムによってプログラムの正しさを保証することも出来ず、バグが入り込む原因になる。
+
+Cerium の Allocator は Thread 間で共有されている。
+共有されているので、ある Thread がメモリを確保しようとすると他の Thread は終了を待つ必要がある
+その間メモリを確保することができないので処理が止まり、なにもしない時間が生まれてしまう。
+これが並列度の低下に繋がり、処理速度が落ちる原因になる。
+
+今回設計した Gears OS はこれらの問題を解決することを目的としている。
--- a/gearsos.tex	Tue Feb 09 19:16:59 2016 +0900
+++ b/gearsos.tex	Sun Feb 14 07:02:11 2016 +0900
@@ -119,6 +119,33 @@
 \lstinputlisting[label=enqueue, caption=Enqueue]{src/enqueue.c}
 \lstinputlisting[label=dequeue, caption=Dequeue]{src/dequeue.c}
 
-ソースコード:\ref{enqueue} とソースコード:\ref{dequeue} はシングルスレッドでは正常に動作するが、並列実行すると期待した値にならない。
+ソースコード:\ref{enqueue} とソースコード:\ref{dequeue} はシングルスレッドでは正常に動作するが、マルチスレッドでは期待した動作を達成できない可能性がある。
+並列実行すると同じメモリ位置にアクセスされる可能性があり、データの一貫性が保証できないからである。
+データの一貫性を並列実行時でも保証するために Compare and Swap(CAS) を利用して Queue の操作を行うように変更する必要がある。
+CAS はデータの比較・置換をアトミックに行う命令である。
+メモリからのデータの読み出し、変更、メモリへのデータの書き出しという一連の処理を、CAS を利用することで処理の間に他のスレッドがメモリに変更を加えていないということを保証することができる。
+CAS に失敗した場合は置換は行わず、再びデータの読み出しから始める。
+
+ソースコード:\ref{enqueue} 44行目の putQueue3, 51行目の putQueue4, ソースコード:\ref{dequeue} 2行目の getQueue が実際に Queue を操作している Code Gear である。
+これらの Code Gear から CAS を利用したソースコード:\ref{sync_enqueue}, ソースコード:\ref{sync_dequeue} の Code Gear に接続を変更することでスレッドセーフな Queue として扱うことが可能になる。
+Code Gear は Gears OS における最小の処理単位となっており、接続を変更することでプログラムの振る舞いを柔軟に変更することができる。
+
+\lstinputlisting[label=sync_enqueue, caption=Enqueue using CAS]{src/sync_enqueue.c}
+\lstinputlisting[label=sync_dequeue, caption=Dequeue using CAS]{src/sync_dequeue.c}
+
+\section{Worker}
 
 \section{Red-Black Tree}
+Gears OS では Persistent Data Gear の管理に木構造を用いる。
+
+\section{TaskManager}
+Gears OS の TaskManager は WaitTaskQueue に入っている Task の依存関係を解決する。
+Task には Input/Output Data Gear の情報が格納されている。
+Input Data Gear は Task に必要な Data Gear で揃ったら Task は実行可能な状態になる。
+Output Data Gear は Task が Persistent Data Tree に書き出す Data Gear である。
+この Input と Output の関係が依存関係となる。
+TaskManager は Persistent Data Tree を監視しており、WaitTaskQueue に入っている Task の Input Data Gear が揃っているのを確認したら実行可能な Task として AcitiveTaskQueue へ移動させる。
+
+TaskManager は Worker の管理も行う。
+メインとなる Context には Worker の情報が格納されており、TaskManager はこの Context を参照して Worker の起動・停止を行う。
+ソースコード\ref{init_worker}は Worker を起動する Code Gear である。
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/images/bitonic.bb	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,5 @@
+%%Title: ./images/bitonic.pdf
+%%Creator: extractbb 20140317
+%%BoundingBox: 0 0 591 422
+%%CreationDate: Fri Feb 12 02:40:15 2016
+
Binary file images/bitonic.pdf has changed
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/images/bitonic_sort_01.bb	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,5 @@
+%%Title: ./images/bitonic_sort_01.pdf
+%%Creator: extractbb 20140317
+%%BoundingBox: 0 0 360 216
+%%CreationDate: Sat Feb 13 22:47:36 2016
+
Binary file images/bitonic_sort_01.pdf has changed
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/images/bitonic_sort_02.bb	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,5 @@
+%%Title: ./images/bitonic_sort_02.pdf
+%%Creator: extractbb 20140317
+%%BoundingBox: 0 0 360 216
+%%CreationDate: Sun Feb 14 00:45:20 2016
+
Binary file images/bitonic_sort_02.pdf has changed
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/images/bitonic_sort_03.bb	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,5 @@
+%%Title: ./images/bitonic_sort_03.pdf
+%%Creator: extractbb 20140317
+%%BoundingBox: 0 0 360 216
+%%CreationDate: Sat Feb 13 23:41:22 2016
+
Binary file images/bitonic_sort_03.pdf has changed
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/images/codesegment2.bb	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,5 @@
+%%Title: ./images/codesegment2.pdf
+%%Creator: extractbb 20140317
+%%BoundingBox: 0 0 584 275
+%%CreationDate: Wed Feb 10 03:34:03 2016
+
Binary file images/codesegment2.pdf has changed
Binary file images/images.graffle has changed
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/images/stream.bb	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,5 @@
+%%Title: ./images/stream.pdf
+%%Creator: extractbb 20140317
+%%BoundingBox: 0 0 952 356
+%%CreationDate: Fri Feb 12 01:28:07 2016
+
Binary file images/stream.pdf has changed
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/images/wordcount.bb	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,5 @@
+%%Title: ./images/wordcount.pdf
+%%Creator: extractbb 20140317
+%%BoundingBox: 0 0 657 356
+%%CreationDate: Sun Feb 14 02:17:07 2016
+
Binary file images/wordcount.pdf has changed
Binary file master_paper.pdf has changed
--- a/master_paper.sty	Tue Feb 09 19:16:59 2016 +0900
+++ b/master_paper.sty	Sun Feb 14 07:02:11 2016 +0900
@@ -79,41 +79,49 @@
 \footskip \headsep
 
 
-%年度
+% 年度
 \def\@year{}
 \def\year#1{\gdef\@year{#1}}
 \def\@eyear{}
 \def\eyear#1{\gdef\@eyear{#1}}
-%和文タイトル
+
+% タイトル
 \def\jetitle{}
 \def\jtitle#1{\gdef\@title{#1}}
-%英文タイトル
 \def\@etitle{}
 \def\etitle#1{\gdef\@etitle{#1}}
-%所属
-%\def\@belongto{}
-%\def\belongto#1{\gdef\@belongto{#1}}
-\def\@affiliation{}
-\def\affiliation#1{\gdef\@affiliation{#1}}
-\let\belongto\affiliation
-%名前
+
+% 名前
 \def\@author{}
 \def\author#1{\gdef\@author{#1}}
 \def\@eauthor{}
 \def\eauthor#1{\gdef\@eauthor{#1}}
 
+% 学位
 \def\thesis{修士(工学)学位論文}
 \def\ethesis{Master's Thesis of Engineering}
 
+% 大学
 \def\university{琉球大学}
 \def\euniversity{University of the Ryukyus}
 
+% 所属
 \def\department{大学院理工学研究科}
 \def\edepartment{Graduate School of Engineering and Science}
 
+% 専攻
 \def\course{情報工学専攻}
 \def\ecourse{Infomation Engineering Course}
 
+% 署名
+\def\commission{論　文　審　査　会}
+
+% 指導教員
+\def\@chife{}
+\def\chife#1{\gdef\@chife{#1}}
+\def\@echife{}
+\def\echife#1{\gdef\@echife{#1}}
+
 %表紙
 %\renewcommand{\maketitle}{%
 %\newpage\null
@@ -135,33 +143,62 @@
 %%\par\vskip 1.5em
 %}
 \renewcommand{\maketitle}{%
-\newpage\null
-\thispagestyle{empty}
-\vskip 1cm%
-\begin{center}%
-  \let\footnote\thanks
-  {\Large\bf\mc\thesis\par}
-  {\Large\bf\ethesis \par}
-    \vskip 2.0 em
-  {\LARGE\bf\mc\@title \par}
-  {\LARGE\textbf{\@etitle} \par}
-    \vskip 2.0 em
-  {\large\mc\@year \par}
-  {\large\@eyear \par}
-    \vskip 1.0 em
-  {\large\bf\mc\@author \par}
-  {\large\textbf\@eauthor \par}
-    \vskip 2.0 em
-  {\large \@affiliation \par}
-    \vskip 2.0 em
-  {\large\bf\mc\university \par}
-  {\large\bf\mc\department \par}
-  {\large\bf\mc\course \par}
-    \vskip 1.0 em
-  {\large\textbf\ecourse \par}
-  {\large\textbf\edepartment \par}
-  {\large\textbf\euniversity \par}
-\end{center}}
+  \newpage\null
+  \thispagestyle{empty}
+  \vskip 1cm%
+  
+  \begin{center}%
+    \let\footnote\thanks
+        {\Large\bf\thesis\par}
+        {\Large\bf\ethesis\par\vskip 0.5em}
+        {\LARGE\bf\mc\@title\par}
+        {\LARGE\bf{\@etitle}\par\vskip 0.5 em}
+        {\large\mc\@year\par}
+        {\large\@eyear\par\vskip 0.2 em}
+        {\large\bf\mc\@author\par}
+        {\large\bf\@eauthor\par\vskip 1.0 em}
+        {\includegraphics[clip,keepaspectratio=true,scale=0.5]{images/u-ryukyu-Mark.eps}\vskip 2.0 em}
+        {\large\bf\mc\university\par}
+        {\large\bf\mc\department\par}
+        {\large\bf\mc\course\par\vskip 0.2 em}
+        {\large\textbf\ecourse \par}
+        {\large\textbf\edepartment \par}
+        {\large\textbf\euniversity \par\vskip 0.2 em}
+        {\large\bf\mc\@chife\par}
+        {\large\bf\@echife\par}
+  \end{center}
+}
+
+\newcommand{\makecommission} {
+  \thispagestyle{empty}
+  \null
+  \vskip 10 em
+  \begin{center}
+    本論文は、修士(工学)の学位論文として適切であると認める。
+  \end{center}
+  \vskip 10 em
+  \begin{minipage}{0.5\hsize}
+    \begin{center}
+    \end{center}
+  \end{minipage}
+  \begin{minipage}{0.5\hsize}
+    \begin{center}
+      \bf\commission
+      \vskip 3 em
+      \underline{　　　　　　　　　　　　　　　　　　印}\\
+      (主　査)　　　　和田　知久　　　　
+      \vskip 2 em
+      \underline{　　　　　　　　　　　　　　　　　　印}\\
+      (副　査)　　　　高良　富夫　　　　
+      \vskip 2 em
+      \underline{　　　　　　　　　　　　　　　　　　印}\\
+      (副　査)　　　　長田　智和　　　　
+      \vskip 2 em
+      \underline{　　　　　　　　　　　　　　　　　　印}\\
+      (副　査)　　　　河野　真治　　　　
+    \end{center}
+  \end{minipage}
+}
 
 %abstract
 \renewenvironment{abstract}{%
@@ -169,17 +206,17 @@
       \thispagestyle{empty}
       \null\vfil
       \@beginparpenalty\@lowpenalty
-      {\Huge \bfseries \abstractname}%
+       {\Huge \bfseries \abstractname}%
       \begin{center}%
         \@endparpenalty\@M
       \end{center}
 }%
 
 \newenvironment{abstract_eng}{%
-      \titlepage
-      \thispagestyle{empty}
-      \null\vfil
-      \@beginparpenalty\@lowpenalty
+  \titlepage
+  \thispagestyle{empty}
+  \null\vfil
+  \@beginparpenalty\@lowpenalty
       {\Huge \bfseries \abstractengname}%
       \begin{center}%
         \@endparpenalty\@M
@@ -197,12 +234,14 @@
   \if@restonecol\twocolumn\fi
 }
 
+% 目次用ページ番号
 \newcommand\frontmatter{%
   \cleardoublepage
   %\@mainmatterfalse
   \pagenumbering{roman}
 }
 
+% 本文用ページ番号
 \newcommand\mainmatter{%
   \cleardoublepage
   % \@mainmattertrue
--- a/master_paper.tex	Tue Feb 09 19:16:59 2016 +0900
+++ b/master_paper.tex	Sun Feb 14 07:02:11 2016 +0900
@@ -5,26 +5,27 @@
 \usepackage{here}
 \usepackage{listings,jlisting}
 \usepackage{comment}
+\usepackage[deluxe, multi]{otf}
+
 %\input{dummy.tex} %% font
 
-\jtitle{Code Segment と Data Segment によって構成される Gears OS の設計}
-\etitle{Design of Gears OS with consist of Code and Data Segment}
+\jtitle{Code Segment と Data Segment を持つ Gears OS の設計}
+\etitle{Design of Gears OS with Code and Data Segment}
 \year{2016年 3月}
 \eyear{March 2016}
-\affiliation{\center%
-  \includegraphics[clip,keepaspectratio,width=.15\textwidth]
-                  {images/u-ryukyu-Mark.eps}
-}
-
 \author{小久保 翔平}
 \eauthor{Shohei KOKUBO}
+\chife{指導教員：教授　和田 知久}
+\echife{Supervisor: Prof. Tomohisa WADA}
+
 \marklefthead{% 左上に挿入
   \begin{minipage}[b]{.4\textwidth}
     琉球大学大学院学位論文(修士)
 \end{minipage}}
+
 \markleftfoot{% 左下に挿入
   \begin{minipage}{.8\textwidth}
-    Code Segment と Data Segment によって構成される Gears OS の設計
+    Code Segment と Data Segment を持つ Gears OS の設計
 \end{minipage}}
 
 \newcommand\figref[1]{図 \ref{fig:#1}}
@@ -64,9 +65,11 @@
 \maketitle
 \newpage
 
+\makecommission
+
 %要旨
-\begin{abstract}
-\end{abstract}
+\input{abstract.tex}
+\input{abstract_eng.tex}
 
 %目次
 \tableofcontents
@@ -83,7 +86,7 @@
 %chapters
 \chapter{並列分散環境下におけるプログラミング}
 \input{cerium.tex}
-\chapter{CbC}
+\input{cbc.tex}
 \input{gearsos.tex}
 \chapter{比較}
 \section{Cerium}
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/result/bitonic_sort/box	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,18 @@
+1 6.145623
+1 6.139483
+1 6.144892
+2 4.539460
+2 4.436463
+2 4.925628
+3 2.530406
+3 2.570707
+3 2.571939
+4 1.631456
+4 1.627970
+4 1.632292
+5 1.322997
+5 1.321029
+5 1.312275
+6 0.149487
+6 0.152230
+6 0.163515
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/result/bitonic_sort/cpu_1	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,21 @@
+16384 0.054535
+16384 0.053521
+16384 0.054224
+32768 0.119125
+32768 0.118421
+32768 0.118529
+65536 0.261226
+65536 0.258138
+65536 0.259855
+131072 0.575518
+131072 0.574977
+131072 0.573440
+262144 1.275700
+262144 1.280920
+262144 1.266228
+524288 2.803241
+524288 2.796131
+524288 2.793901
+1048576 6.145623
+1048576 6.139483
+1048576 6.144892
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/result/bitonic_sort/cpu_12	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,21 @@
+16384 0.028451
+16384 0.030065
+16384 0.030341
+32768 0.044924
+32768 0.045511
+32768 0.043224
+65536 0.076149
+65536 0.075513
+65536 0.076242
+131072 0.138685
+131072 0.138920
+131072 0.138807
+262144 0.285647
+262144 0.284760
+262144 0.289191
+524288 0.640933
+524288 0.642066
+524288 0.616291
+1048576 1.322997
+1048576 1.321029
+1048576 1.312275
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/result/bitonic_sort/cpu_2	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,21 @@
+16384 0.040895
+16384 0.039489
+16384 0.039098
+32768 0.085891
+32768 0.084144
+32768 0.086287
+65536 0.177633
+65536 0.167117
+65536 0.172507
+131072 0.372879
+131072 0.393514
+131072 0.393431
+262144 0.794246
+262144 0.934616
+262144 0.906198
+524288 2.164428
+524288 2.217274
+524288 2.032049
+1048576 4.539460
+1048576 4.436463
+1048576 4.925628
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/result/bitonic_sort/cpu_4	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,21 @@
+16384 0.028437
+16384 0.027519
+16384 0.027354
+32768 0.056601
+32768 0.056044
+32768 0.057884
+65536 0.117800
+65536 0.121380
+65536 0.121663
+131072 0.244752
+131072 0.246284
+131072 0.248601
+262144 0.549241
+262144 0.545955
+262144 0.517979
+524288 1.192022
+524288 1.123386
+524288 1.195495
+1048576 2.530406
+1048576 2.570707
+1048576 2.571939
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/result/bitonic_sort/cpu_8	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,21 @@
+16384 0.024300
+16384 0.024318
+16384 0.023870
+32768 0.041917
+32768 0.042845
+32768 0.042200
+65536 0.081457
+65536 0.080062
+65536 0.081057
+131072 0.169538
+131072 0.164242
+131072 0.166205
+262144 0.351782
+262144 0.355738
+262144 0.356107
+524288 0.771449
+524288 0.752787
+524288 0.755115
+1048576 1.631456
+1048576 1.627970
+1048576 1.632292
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/result/bitonic_sort/gpu	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,21 @@
+16384 0.083421
+16384 0.093149
+16384 0.078294
+32768 0.084040
+32768 0.095510
+32768 0.087000
+65536 0.100729
+65536 0.090225
+65536 0.085858
+131072 0.105753
+131072 0.103785
+131072 0.109697
+262144 0.115981
+262144 0.119174
+262144 0.108622
+524288 0.122353
+524288 0.127722
+524288 0.125356
+1048576 0.149487
+1048576 0.152230
+1048576 0.163515
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/init_twice_cerium.cc	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,33 @@
+void
+twice_init(TaskManager *manager, int* data, int length)
+{
+    /**
+     * Create Task
+     *   create_task(Task ID);
+     */
+    HTask* twice = manager->create_task(TWICE_TASK);
+    
+    /**
+     * Set of Device
+     *   set_cpu(CPU or GPU)
+     */
+    twice->set_cpu(SPE_ANY);
+
+    /**
+     * Set of Input Data
+     *   set_inData(index, address of input data, size of input data);
+     */
+    twice->set_inData(0, data, sizeof(int)*length);
+   
+    /**
+     * Set of OutPut area
+     *   set_outData(index, address of output area, size of output area);
+     */
+    twice->set_outData(0, data, sizeof(int)*length);
+
+    /**
+     * Enqueue Task
+     *   iterate(Number of Tasks)
+     */
+    twice->iterate(length);
+}
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/sample.c	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,11 @@
+// Code Gear
+__code code1(struct Allocate* allocate) {
+    allocate->size = sizeof(struct Data1);
+    
+    goto allocator(allocate, Code2);
+}
+
+// Code Gear
+__code code2(struct Data1* data1) {
+    // processing
+}
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/sample_transform.c	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,22 @@
+// Code Gear
+__code code1(struct Context* context, struct Allocate* allocate) {
+    allocate->size = sizeof(struct Data1);
+    context->next  = Code2;
+    
+    goto meta(context, Allocator);
+}
+
+// Meta Code Gear(stub)
+__code code1_stub(struct Context* context) {
+    goto code1(context, &context->data[Allocate]->allocate);
+}
+
+// Code Gear
+__code code2(struct Context* context, struct Data1* data1) {
+    // processing
+}
+
+// Meta Code Gear(stub)
+__code code2_stub(struct Context* context) {
+    goto code2(context, &context->data[context->dataNum]->data1);
+}
--- a/src/sync_dequeue.c	Tue Feb 09 19:16:59 2016 +0900
+++ b/src/sync_dequeue.c	Sun Feb 14 07:02:11 2016 +0900
@@ -5,7 +5,6 @@
 
     struct Element* first = queue->first;
     if (__sync_bool_compare_and_swap(&queue->first, first, first->next)) {
-        //    if (OSAtomicCompareAndSwapPtr(first, first->next, (void*)&queue->first)) {
         queue->count--;
 
         context->next = GetQueue;
@@ -19,10 +18,3 @@
         goto meta(context, GetQueue);
     }
 }
-
-// Meta Code Gear(stub)
-__code getQueue_stub(struct Context* context) {
-    goto getQueue(context,
-                  &context->data[ActiveQueue]->queue,
-                  &context->data[Node]->node);
-}
--- a/src/sync_enqueue.c	Tue Feb 09 19:16:59 2016 +0900
+++ b/src/sync_enqueue.c	Sun Feb 14 07:02:11 2016 +0900
@@ -1,34 +1,3 @@
-// allocate Element
-__code putQueue1(struct Context* context, struct Allocate* allocate) {
-    allocate->size = sizeof(struct Element);
-    allocator(context);
-
-    goto meta(context, PutQueue2);
-}
-
-// Meta Code Gear(stub)
-__code putQueue1_stub(struct Context* context) {
-    goto putQueue1(context, &context->data[Allocate]->allocate);
-}
-
-// write Element infomation
-__code putQueue2(struct Context* context, struct Element* new_element, struct Element* element, struct Queue* queue) {
-    new_element->task = element->task;
-
-    if (queue->first)
-        goto meta(context, PutQueue3);
-    else
-        goto meta(context, PutQueue4);
-}
-
-// Meta Code Gear(stub)
-__code putQueue2_stub(struct Context* context) {
-    goto putQueue2(context,
-                   &context->data[context->dataNum]->element,
-                   &context->data[Element]->element,
-                   &context->data[ActiveQueue]->queue);
-}
-
 // Enqueue(normal)
 __code putQueue3(struct Context* context, struct Queue* queue, struct Element* new_element) {
     struct Element* last = queue->last;
@@ -43,11 +12,6 @@
     }
 }
 
-// Meta Code Gear(stub)
-__code putQueue3_stub(struct Context* context) {
-    goto putQueue3(context, &context->data[ActiveQueue]->queue, &context->data[context->dataNum]->element);
-}
-
 // Enqueue(nothing element)
 __code putQueue4(struct Context* context, struct Queue* queue, struct Element* new_element) {
     if (__sync_bool_compare_and_swap(&queue->first, 0, new_element)) {
@@ -59,8 +23,3 @@
         goto meta(context, PutQueue3);
     }
 }
-
-// Meta Code Gear(stub)
-__code putQueue4_stub(struct Context* context) {
-    goto putQueue4(context, &context->data[ActiveQueue]->queue, &context->data[context->dataNum]->element);
-}
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/twice_cerium.cc	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,28 @@
+static int
+twice(SchedTask *s,void *rbuf, void *wbuf)
+{
+    /**
+     * Get Input Data
+     *   get_input(input data buffer, index)
+     */
+    int* input = (int*)s->get_input(rbuf, 0);
+
+    /**
+     * Get Output Data
+     *   get_output(output data buffer, index)
+     */
+    int* output = (int*)s->get_output(wbuf, 0);
+
+    /**
+     * Get index(x, y, z)
+     *   SchedTask member
+     *   x : SchedTask->x
+     *   y : SchedTask->y
+     *   z : SchedTask->z
+     */
+    long i = s->x;
+
+    output[i] = input[i]*2;
+    
+    return 0;
+}
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/twice_cuda.cu	Sun Feb 14 07:02:11 2016 +0900
@@ -0,0 +1,16 @@
+__global__ void
+twice(int* input, int* output)
+{
+    /**
+     * Get index(x, y, z)
+     *   kernel built-in variables
+     *   x : blockIdx.x * blockDim.x + threadIdx.x
+     *   y : blockIdx.y * blockDim.y + threadIdx.y
+     *   z : blockIdx.z * blockDim.z + threadIdx.z
+     */
+    long i = blockIdx.x * blockDim.x + threadIdx.x;
+
+    output[i] = input[i]*2;
+    
+    return 0;
+}