--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/Benchmark.tex	Thu Feb 20 02:57:54 2014 +0900
@@ -0,0 +1,5 @@
+\chapter{Benchmark}
+
+\section{WordCount}
+
+\section{FFT}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/Cerium.tex	Thu Feb 20 02:57:54 2014 +0900
@@ -0,0 +1,69 @@
+\chapter{Cerium}
+
+\section{Cerium TaskManager}
+TaskManager は、Task と呼ばれる分割されたプログラムを管理する。Task の単位はサブルーチンまたは関数として、Task 同士の依存関係を考慮しながら実行される。
+TaskManager の API を表:\ref{table:TaskManager_api}に示す。
+\begin{table}[htpb]
+  \begin{center}
+    \small
+    \begin{tabular}[htpb]{|c|l|} \hline
+      create\_task & Task を生成する \\ \hline
+      allocate     & 環境のアライメントに考慮した allocator \\ \hline
+      set\_inData  & Task への入力データのアドレスを追加 \\ \hline
+      set\_outData & Task からのデータ出力先アドレスを追加 \\ \hline
+      set\_param   & Task のパラメータ(32 bits) \\ \hline
+      wait\_for    & Task の依存関係の考慮 \\ \hline
+      set\_cpu     & Task を実行する Device の設定 \\ \hline
+      spawn        & Task を Queue に登録する \\ \hline
+      iterate      & データ並列で実行する Task として Queue に登録する \\ \hline
+    \end{tabular}
+    \caption{TaskManager API}
+    \label{table:TaskManager_api}
+  \end{center}
+\end{table}
+
+\section{Cerium における Task}
+Task は TaskManager を使って生成する。Task を生成すると際に、以下のような要素を設定することができる。
+
+\begin{itemize}
+  \item input data
+  \item output data
+  \item parameter
+  \item cpu type
+  \item dependency
+\end{itemize}
+
+input, output, parameter は関数でいうところの引数に相当する。
+cpu type は Task がどのような Device の組み合わせで実行されるかを示す。
+dependency は他の Task との依存関係を示している。Task の終了が通知され、その通知に従って依存関係を処理していく。
+例えば、Task B には Task A との依存関係がある場合、Task B は Wait Queue に投入され依存関係が解消されるのを待つ。
+Task A は Active Queue に投入され、cpu type によって示された Device に転送され実行される。
+Task A の処理が終了すると Mail で Task B へ通知される。Task B はその通知を受けると Active Queue に投入され、Task A と同様に指定された Device 上で実行される。
+
+\section{Task のスケジューリング}
+Device は、Task を一つずつ受け取るのではなく、ある程度まとめて受け取る。それを TaskList と呼んでいる。TaskList は各 Device 用に TaskManager 側で生成される。
+受け取った TaskList に沿って Task をパイプラインで実行していく。Task 毎に実行完了の Mail を送る。
+TaskList の Task をすべて実行すると、次の TaskList を要求する Mail を送る。実行の様子は図:\ref{fig:task_scheduler}のようになる。
+
+\begin{figure}[htpb]
+  \begin{center}
+    \includegraphics[scale=0.6]{./images/scheduler.pdf}
+  \end{center}
+  \caption{Task Scheduler}
+  \label{fig:task_scheduler}
+\end{figure}
+
+\newpage
+\section{Task の入出力}
+Task の入出力の API として、set\_inData, set\_outData がある。
+\begin{description}
+  
+\item[set\_inData(index, addr, size)] は、データを受け取る buffer の配列番号とデータのアドレス、そのデータのサイズを引数として入力する。
+  このデータは DMA 転送されるため、addr は 16 byte alignment が取れており、size は 16 byte の倍数である必要がある。
+  
+\item[set\_param(index, param)] は、データを受け取る buffer の配列番号と 32bit のデータを渡す。set\_inData で渡すには小さいデータを送るのに適している。
+  param はアドレスとしてではなく、値を Task オブジェクトが直接持っているので、 DMA 転送は行わない。
+  
+\item[set\_outData(index, addr, size)] は、Task のデータの出力先を指定する。使用方法は set\_inData と同じで、alignment, byte 数に気をつける必要がある。
+  
+\end{description}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/Conclusion.tex	Thu Feb 20 02:57:54 2014 +0900
@@ -0,0 +1,5 @@
+\chapter{結論}
+
+\section{まとめ}
+
+\section{今後の課題}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/GPGPU.tex	Thu Feb 20 02:57:54 2014 +0900
@@ -0,0 +1,10 @@
+\chapter{Cerium TaskManager の GPGPU への対応}
+本章では、まずはじめに GPU Programming の問題点について述べ、Cerium TaskManager への実装でそれをどう改善したのかについて説明する。
+
+\section{GPU Programming の問題点}
+
+\section{パイプライン実行}
+
+\section{OpenCL 実装}
+
+\section{CUDA 実装}

--- a/paper/Parallel.tex	Mon Feb 17 17:40:20 2014 +0900
+++ b/paper/Parallel.tex	Thu Feb 20 02:57:54 2014 +0900
@@ -1,5 +1,4 @@
 \chapter{Parallel Computing Platform}
-\pagenumbering{arabic}
 
 \section{OpenCL}
 OpenCL とは、Multi Core CPU と GPU のようなヘテロジニアスな環境を利用した並列計算を支援するフレームワークである。
@@ -14,7 +13,7 @@
 
 オペレーティングシステムなどが処理される、メイン CPU などのことを host、GPGPU を搭載したグラフィックボードなどのことを device と定義している。
 OpenCL では device に CPU を割り当てることも可能である。OpenCL Application は host 側のプログラムと device 側のプログラムが一体となって動作する。
-この device 側で動作するプログラムを OpenCL では、特別に kernel と呼ぶ。
+この device 上で動作するプログラムを OpenCL では、特別に kernel と呼ぶ。
 
 \subsection{Command Queue}
 OpenCL では、デバイスの操作に Command Queue を使用する。Command Queue は device に OpenCL の Operation を送るために仕組みである。
@@ -30,12 +29,12 @@
 
 \subsection{Memory Access}
 host 側は主にデータを input/output する Memory の確保を行う。
-GPU の Memory 空間(図:\ref{fig:gpuarch})や Cell の Memory 空間(図:\ref{fig:cellarch})は Multi Core CPU(図:\ref{fig:cpuarch})とは異なり、共有 Memory ではないため host と kernel 間でデータの共有ができない。
+GPU の Memory 空間(図:\ref{fig:gpuarch})や Cell の Memory 空間(図:\ref{fig:cellarch})は Multi Core CPU(図:\ref{fig:cpuarch})とは異なり、Shared Memory ではないため host と device 間でデータの共有ができない。
 アクセスするには Memory 空間ごとコピーしなければならない。
 
 OpenCL では host 側で Memory Buffer を作成して Memory のコピーを行う。
 データの読み込みは clEnqueueReadBuffer、書き込みは clEnqueueWriteBuffer という API でそれぞれ行われる。
-前節で述べた通り、これらの Operation を Command Queue に Enqueue する。そして、event によってデータの依存解消が解消されると実行される。
+前節で述べた通り、これらの Operation を Command Queue に Enqueue する。そして、データの依存解消が解消されると実行される。
 
 \begin{figure}[htpd]
   \begin{center}
@@ -69,20 +68,122 @@
 多次元のデータ構造がある場合に高い並列度を保つには、それを分割して並列に実行する機能が必要である。
 これを OpenCL ではデータ並列と呼んでいる。
 OpenCL では次元数に対応する index があり、OpenCL は一つの記述から index の異なる複数の kernel を自動生成する。
-その添字を global\_id と呼ぶ。このとき入力されたデータはワークアイテムという処理単位に分割される。
+その添字を global\_id と呼ぶ。このとき入力されたデータは WorkItem という処理単位に分割される。
 
-OpenCL はワークアイテムに対して、それぞれを識別する ID(global\_id) を割り当てる。
+OpenCL は WorkItem に対して、それぞれを識別する ID(global\_id) を割り当てる。
 kernel は get\_global\_id という API によって ID を取得し、取得した ID に対応するデータに対して処理を行うことでデータ並列を実現する。
-この ID によって取得してきたワークアイテムをグローバルワークアイテムという。
-また、ワークアイテムは3次元までにデータを渡すことができる。
+また、WorkItem は3次元までにデータを渡すことができる。
 
-データ並列による kernel 実行の場合、clEnqueueNDRangeKernel API を使用するが、この関数の引数としてワークアイテムのサイズと次元数を指定することでデータ並列で実行できる。
+データ並列による kernel 実行の場合、clEnqueueNDRangeKernel API を使用する。この関数の引数として WorkItem の数と次元数を指定することでデータ並列で実行できる。
 
 \begin{figure}[htpd]
   \begin{center}
-    \includegraphics[scale=0.4]{./images/workitem.pdf}
+    \includegraphics[scale=0.7]{./images/workitem.pdf}
   \end{center}
   \caption{WorkItem ID}
   \label{fig:workitem}
 \end{figure}
-  
+
+\newpage
+\subsection{Work Group}
+前節で WorkItem という処理単位について述べたが、さらに複数個の WorkItem を WorkGroup という単位にまとめることができる。
+WorkGroup 内では同期やローカルメモリの共有が可能になる。
+
+グローバルワークアイテム(ワークアイテム全体)の個数と、ローカルワークアイテム(グループ一つ辺りのアイテム)の個数を指定することでワークアイテムを分割する。
+なお、このときグローバルワークアイテム数はローカルアイテム数の整数倍でなければ clEnqueueNDRangeKernel API の呼び出しは失敗する。
+
+ローカルアイテムに0を指定することで、コンパイル時に最適化させることができる。
+したがって、ローカルアイテムのサイズは0を指定するのが一般的である。
+
+なお、WorkGroup を設定した場合、global\_id の他に work\_group\_id, local\_id がそれぞれの kernel に割り当てられる(図:\ref{fig:workitem})。
+
+kernel からそれぞれ ID に対応した API を使用して、各 ID を取得する。
+取得した ID から自分が担当する index を計算して導く。
+表:\ref{table:kernel_id_api}は kernel で使用できる ID を取得するための API となる。
+
+\begin{table}[htpb]
+  \begin{center}
+    \small
+    \begin{tabular}[htpb]{|c|l|} \hline
+      get\_group\_id  & work\_group\_id を取得  \\ \hline
+      get\_local\_id  & local\_id を取得 \\ \hline
+      get\_global\_id & global\_id を取得 \\ \hline
+    \end{tabular}
+    \caption{kernel で使用する ID 取得の API}
+    \label{table:kernel_id_api}
+  \end{center}
+\end{table}
+
+なお、local\_id, global\_id を取得する API は引数に0, 1, 2の値を指定することができる。
+id は x, y, z 座標があり、それぞれが0, 1, 2に対応している。
+例えば get\_global\_id(1) と呼び出した場合は y 座標の、get\_global\_id(2) と呼び出した場合は z 座標の global\_id を取得する。
+
+\section{CUDA}
+CUDA とは、半導体メーカー NVIDIA 社が提供する GPU コンピューティング向けの総合開発環境でプログラム記述、コンパイラ、デバッガなどから構成される。
+プログラム言語である CUDA C は C 言語ベースに拡張を加えたものである。
+
+CUDA には CUDA Runtime API と CUDA Driver API の2種類がある。
+Driver API は Runtime API と比べてプログラマが管理すべきリソースが多い。しかし、Runtime API より柔軟な処理を行うことができる。
+今回は Driver API を使用して実装した。
+
+CUDA も OpenCL と同様に、GPU 側を device、制御を行う CPU 側を host と定義している。また、device 上で動作するプログラムも OpenCL と同様に kernel と呼ぶ。
+
+\subsection{Stream}
+CUDA には OpenCL の Command Queue と似たような仕組みとして Stream がある。
+Stream は host 側の発行された Operation を一連の動作として device で実行する。Stream 内の Operation は発行された順序で実行されることが保証されている。
+異なる Stream での Operation でデータの依存関係などが解消され実行可能な場合、同時に実行することができ、Interleave させることができる。
+Stream は cuStreamCreate という Driver API で生成される。
+OpenCL と異なり、コンテキストと実行対象となる device を指定する必要はないが、コンテキストを作成した Thread と同一の Thread でないと Stream が生成できないという制約がある。
+引数に Stream を指定しない API はすべて host 側をブロックする同期的な処理となる。複数の Stream を同時に走らせ Operation を並列に実行するためには非同期処理を行う API を利用する必要がある。
+
+Stream 内の Operation を同期を行う方法はいくつかある。
+一つ目は cuStreamSynchronize API を利用した同期方法である。
+cuStraemSynchronize API の引数に Stream を指定すると、指定した Stream に発行されたすべての Operation が終了するまで host をブロックする。
+host をブロックすることなく、Stream に発行された Operation が終了したかどうかを調べるには cuStreamQuery API を利用する。
+二つ目は event を利用した同期方法である。
+各 Operation に event\_wait\_list や event を指定して同期を取る OpenCL とは異なり、CUDA では cuStreamWaitEvent で Stream と event を指定して同期を取る。
+cuStreamWaitEvent は指定した event が cuEventRecord されるまで指定された Stream をブロックする。
+event は cuEventCreate で作成され、Operation の直後で cuEventRecord に event を指定することで直前に行う Operation の終了を待つことが可能になる。
+この event は別の Stream で cuEventRecord されるものでも待つことができる。
+
+\subsection{Memory Access}
+CUDA も OpenCL と同様に Shared Memory ではないため host と device 間でデータの共有ができない。アクセスするには Memory 空間ごとコピーする必要がる。
+CUDA でのデータの読み込みは cuMemcpyHtoD、書き込みは cuMemcpyDtoH という API でそれぞれ行われる。しかし、これらの API は同期的に実行されてしまう。
+非同期処理にしたい場合、読み込みで cuMemcpyHtoDAsync、書き込みで cuMemcpyDtoHAsync という API をそれぞれ利用することで非同期に行うことができる。
+
+\subsection{Data Parallel Execution}
+CUDA では OpenCL の WorkItem に相当する単位を Thread と定義してる。この Thread をまとめたものを block と呼ぶ。
+CUDA でデータ並列による kernel 実行を行う場合、cuLaunchKernel API を使用する。
+この関数の引数として各座標の block 数と各座標の block 一つ当たりの thread 数を指定することでデータ並列で実行できる。
+cuLaunchKernel で kernel を起動すると各 Thread に対して block ID と Thread ID が付与される。
+CUDA には OpenCL とは異なり、ID を取得する API は存在しない。
+代わりとして、kernel には組み込み変数が準備されており、それを参照することで対応するデータに対し処理を行うことでデータ並列を実現する。
+組み込み変数は以下の通りである。
+\begin{itemize}
+  \item uint3 blockDim
+  \item uint3 blockIdx
+  \item uint3 threadIdx
+\end{itemize}
+各組み込み変数はベクター型で、blockDim.x とすると x 座標の Thread 数を参照することができる。
+blockIdx.x とすると x 座標の block ID が参照でき、threadIdx.x とすると x 座標の thread ID を参照することができる。
+blockDim.x * blockIdx.x + threadIdx.x を計算すると OpenCL の get\_global\_id(0) で取得できる ID に相当する ID を得ることができる。
+例えば、ある kernel で get\_global\_id(0) の返り値が13の場合、 CUDA では図:\ref{fig:culculate_index}のようにすることで ID を算出することができる。
+
+\begin{figure}[!h]
+  \begin{center}
+    \includegraphics[scale=0.4]{./images/culculate_index.pdf}
+  \end{center}
+  \caption{Culculate Index}
+  \label{fig:culculate_index}
+\end{figure}
+
+\newpage
+
+\section{Porting to OpenCL to CUDA}
+本項では OpenCL で記述された Application を CUDA に移植する方法について説明する。
+
+\subsection{Sequential Execution}
+
+\subsection{Parallel Execution}
+
+\section{StarPU}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/Thanks.tex	Thu Feb 20 02:57:54 2014 +0900
@@ -0,0 +1,1 @@
+\chapter{謝辞}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/images/culculate_index.bb	Thu Feb 20 02:57:54 2014 +0900
@@ -0,0 +1,5 @@
+%%Title: ./images/culculate_index.pdf
+%%Creator: extractbb 20090506
+%%BoundingBox: 0 0 562 226
+%%CreationDate: Tue Feb 18 19:17:52 2014
+

Binary file paper/images/culculate_index.pdf has changed

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/images/culculate_index.xbb	Thu Feb 20 02:57:54 2014 +0900
@@ -0,0 +1,8 @@
+%%Title: ./images/culculate_index.pdf
+%%Creator: extractbb 20090506
+%%BoundingBox: 0 0 562 226
+%%HiResBoundingBox: 0.000000 0.000000 562.000000 226.000000
+%%PDFVersion: 1.3
+%%Pages: 1
+%%CreationDate: Tue Feb 18 19:17:55 2014
+

--- a/paper/thesis-paper.tex	Mon Feb 17 17:40:20 2014 +0900
+++ b/paper/thesis-paper.tex	Thu Feb 20 02:57:54 2014 +0900
@@ -54,7 +54,20 @@
 %Parallel Computing Platform
 \input{Parallel.tex}
 
+%Cerium
+\input{Cerium.tex}
+
+%Cerium TaskManager の GPGPU への対応
+\input{GPGPU.tex}
+
+%Benchmark
+\input{Benchmark.tex}
+
+%結論
+\input{Conclusion.tex}
+
 %謝辞
+\input{Thanks.tex}
 
 %参考文献
 \nocite{*}