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

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

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

\subsection{Data Parallel Execution}
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}