--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/gpu.tex	Mon Jul 09 10:41:54 2018 +0900
@@ -0,0 +1,96 @@
+% TODO 
+% CUDABuffer はいらないかも
+% CUDA のアーキテクチャの図がいる?
+  % block とか grid とかの関係
+
+\chapter{CUDA への対応} 
+Gears OS では GPU での実行もサポートする\cite{ikkun-sigos}。
+また、CPU、GPU の実行環境の切り替えは Meta Code Gear、つまり stub Code Gear で切り替えを行う。
+
+本章では、 Gears OS での CUDA 実行のサポートについて説明する。
+
+\section{CUDA}
+CUDA\cite{cuda} とは NVIDA 社が提供する GPU コンピューティング向けの総合開発環境である。
+
+CUDA は演算用プロセッサ(GPU) を Device、制御用デバイス側(CPU) を Host として定義する。
+また、Device 上で実行するプログラムのことを kernel と呼ぶ。
+
+また、CUDA には CUDA Runtime API と CUDA Driver API の2種類存在する。
+Driver API は Runtime API に比べて管理すべきリソースが多いが、 Runtime API より柔軟な処理を行うことが出来る。
+Gears OS では Driver API を用いて GPU 実行の実装を行う。
+
+CUDA では処理の最小の単位を thread と定義しており、それをまとめたものを block と呼ぶ。
+block と thread は それぞれ3次元まで展開することが出来る。
+\figref{cudaArchitecture} に thread、block を2次元で展開した例を示す。
+
+\begin{figure}[htbp]
+    \begin{center}
+        \includegraphics[scale=0.7]{./fig/cudaArchitecture.pdf}
+    \end{center}
+    \caption{blockサイズ(3,3)、 threadサイズ(3,3) に展開}
+    \label{fig:cudaArchitecture}
+\end{figure}
+
+kernel を起動すると、各 thread に対して block ID と thread ID が付与される。
+この ID は blockIdx、threadIdx といった組み込み変数で取得できる。
+これらの変数は3次元のベクター型になっており、blockIdx.x とすると x座標の block ID が取得でき、 threadIdx.x とするとx座標の thread Id を取得できる。
+また、block 内の thread 数は blockDim という組み込み変数で取得でき、これも3次元のベクター型になっている。
+CUDA では これらの組み込み変数から thread が対応するデータを割り出し、データ並列の処理を行う。
+
+\section{CUDAWorker}
+CUDAWorker は TaskManager から送信される CUDA用の Task を取得し、実行を行う。
+
+CUDAWorker は CPUWorker と同じく初期化の際にスレッドが生成される。
+生成されたスレッドはCUDAライブラリ初期化や後述する CUDAExectuor の生成を行う。
+
+データ並列用の Task は CUDAWorker に送信する際は Task のコピーを行わず送信する。
+受け取ったデータ並列用の Task は Code Gear のメタレベルで kernel の実行を行う。
+
+\section{CUDAExectuor}
+CUDAExectuor は \coderef{executorInterface} に示す ExecutorInterfaceを実装しており、 Host から Device へのデータの送信(read)、 kernel の実行(exec)、 Device から Host への データの書き出しを行う(write)。
+
+\lstinputlisting[caption=executor Inteface, label=code:executorInterface]{./src/executorInterface.h}
+
+Gears OS では データは Data Gear で表現される。
+つまり、Host、Device 間でデータのやり取りを行うということは Data Gear を GPU のデータ領域に沿った形に適用する必要がある。
+Host から Device へデータを送信する際、 CUDA では cuMemAlloc 関数を使用してサイズを指定し、Device 側のデータ領域を確保する。
+全ての Data Gear には Meta Data Gear として Data Gear のサイズを持っており、基本的にはこのサイズでデータ領域を取ればよい。
+しかし、Data Gear によっては内部に更にポインタで Data Gear を持っている場合がある。
+このような Data Gear は Data Gear の実際のサイズではなく、ポインタのサイズで計算されてしまうため、そのままでは Device 用のデータ領域を確保することができない。
+
+この問題を解決するために、CUDABuffer という CUDA データ送信用の Data Gear を用意した.
+CUDABuffer には Data Gear の内部にポインタを持たない Data Gear まで展開した Input/Output Data Gear を格納される。
+Data Gear を CUDABuffer に格納する処理は CUDAExectuor では行わず、実行される Task の stub Code Gear で行われる。
+CUDABuffer に格納されている Data Gear のサイズを参照し、cuMemAlloc 関数で Device のデータ領域を確保する。
+
+Host、Device、CUDABuffer 間の関係を\figref{cudaDataArchitecture} に示す。
+
+\begin{figure}[htbp]
+    \begin{center}
+        \includegraphics[scale=0.7]{./fig/cudaDataArchitecture.pdf}
+    \end{center}
+    \caption{Host､ Device 間のデータの関係}
+    \label{fig:cudaDataArchitecture}
+\end{figure}
+
+Host から Device にデータをコピーするには cuMemcpyHtoD 関数を使用して行う。
+この際に Host で指定するデータは CUDABuffer に格納されている Data Gear となる。
+
+kernel の実行後、結果を Device から Host にコピーする際は cuMemcpyDtoH 関数で行われる。
+Host のコピーされたデータは Output Data Gear も含んでいるため、 コピー後は Output Data Gear への書き出す処理に継続する。
+
+kernel の実行はcuLaunchKernel 関数で行われる。
+cuLaunchKernel 関数には引数として各次元のblockサイズ、thread サイズ、kernel への引数等を渡す。
+Gears OS ではデータ並列 Task の際は Iterator Interface を持っており、 そこで指定した長さ、次元数に応じて cuLaunchKernel の引数を決定する。
+
+% 少ないけどコードはなるべく載せたくない(メタ部分 + 複雑)
+\section{stub Code Gear による kernel の実行}
+Gears OS では stub Code Gear で CUDA による実行を切り替える。
+
+stub Code Gear での切り替えの際は CUDABuffer への Data の格納、実行される kernel の読み込みを行う。
+実際に GPU で実行されるプログラムは \coderef{cudaTwice} のように記述する。
+
+\lstinputlisting[caption=配列の要素を二倍にする例題, label=code:cudaTwice]{./src/cudaTwice.cu}
+
+通常、stub Code Gear は対応した Code Gear に継続するが、CUDA で実行する際は CUDAExectuor の Code Gear に継続する。
+