% 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 に継続する。