\section{Parallel Computing Platform}
\subsection{OpenCL}
OpenCL とは、Multi Core CPU と GPU のようなヘテロジニアスな環境を利用した並列計算を支援するフレームワークである。
演算用のプロセッサ(本研究では GPU)上で動作するプログラム OpenCL C という言語で記述する。
OpenCL C で記述したプログラムを GPU 上で実行させるために OpenCL Runtime API を利用する。
OpenCL ではオペレーティングシステムなどが処理されるメイン CPU のことを host、GPGPU が可能なグラフィックボードなどのことを device と定義している。
OpenCL Application は host 側のプログラムと device 側のプログラムが一体となって動作する。
この device 上で動作するプログラムを特別に kernel と呼ぶ。

\subsubsection{CommandQueue}
OpenCL では、device の操作に CommandQueue を使用する。
CommandQueue は device に Operation を送るための仕組みである。
kernel の実行、input buffer の読み込み、output buffer への書き込みなどが Operation となる。

CommandQueue に投入された Operation は投入された順序で実行される。
CommandQueue を生成するときプロパティを指定することで Operation を投入された順序を無視して(out of order)実行することが可能になる。
また複数の CommandQueue を生成し、device に投入することでも our of order で実行することが可能である。

out of order で実行する場合、データの依存関係を設定する必要がある。各 Operation を発行する関数には event\_wait\_list と event を指定することができ、これらを利用してデータの依存関係を設定することができる。
out of order 実行を可能にするプロパティをサポートしている device が少ないため、今回は複数の CommandQueue を用いる方法で実装を行なった。

\subsubsection{OpenCL におけるデータ並列}
3D グラフィックのような多次元のデータを処理する場合に高い並列度を保つには、データを分割して並列に実行する機能が必要である。
これを OpenCL ではデータ並列と呼んでいる。
OpenCL では次元数に対応する index があり、OpenCL は1つの記述から index の異なる複数の kernel を自動生成する。
その添字を global\_id と呼ぶ。
このとき入力されたデータは WorkItem という処理単位に分割される。

OpenCL は WorkItem に対して、それぞれを識別する ID(global\_id)を割り当てる。
kernel は get\_global\_id という API によって ID を取得し、取得した ID に対応するデータに対して処理を行うことでデータ並列を実現する。

データ並列による kernel 実行の場合、clEnqueueNDRangeKernel を使用する。
この関数の引数として WorkItem の数と次元数を指定することでデータ並列で実行することができる。

\subsection{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 と同様に、制御を行う CPU 側を host、GPU 側を device と定義している。
また、device 上で動作するプログラムも OpenCL と同様に kernel と呼ぶ。

\subsubsection{Stream}
CUDA には OpenCL の CommandQueue と似たような仕組みとして Stream がある。
Stream は host 側で発行された Operation を一連の動作として device で実行する。
Stream に発行された Operation は発行された順序で実行されることが保証されている。
異なる Stream に発行された Operation に依存関係が存在しない場合、Operation を並列に実行することができる。

Stream は cuStreamCreate という Driver API で生成される。
引数に Stream を指定しない API はすべて host 側をブロックする同期的な処理となる。
複数の Stream を同時に走らせ Operation を並列に実行するためには非同期な処理を行う API を利用する必要がある。

\subsubsection{CUDA におけるデータ並列}
CUDA では OpenCL の WorkItem に相当する単位を thread と定義している。
この thread をまとめたものを block と呼ぶ。
CUDA でデータ並列による kernel 実行をする場合、cuLaunchKernel API を使用する。
この関数は引数として各座標の block 数と各座標の block 1つ当たりの thread 数を指定することでデータ並列で実行できる。

cuLaunckKernel で 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:calculate_index}のようになる。

\begin{figure}[!h]
  \begin{center}
    \includegraphics[scale=0.4]{./images/culculate_index.pdf}
  \end{center}
  \caption{Calculate Index}
  \label{fig:calculate_index}
\end{figure}