\section{OpenCL}
OpenCL とは、マルチコア CPU と GPU のようなヘテロジニアスな環境を利用した並列計算を支援するフレームワークである。
このフレームワークを用いて Cerium を GPGPU に対応させる。

OpenCL には主に2つの仕様がある。

\begin{itemize}
\item OpenCL C言語
\item OpenCL ランタイム API
\end{itemize}
OpenCL C は演算用プロセッサ(本研究では GPU )上で動作する、 C 言語を拡張したプログラミング言語である。
一方で OpenCL ランタイム API は OpenCL C で記述したプログラムを GPU 上で実行させるため、
制御用のプロセッサ(本研究では CPU )が利用する API である。

OpenCL では GPU 側を kernel 、制御デバイス側を host として定義する。

\subsection{Command Queue}
OpenCL では、デバイスの操作に Command Queue を使用する。
 Command Queue は kernel に命令を送るための仕組みである。
 Command Queue は clCreateCommandQueue という OpenCL API で作成され、
 Command Queueが所属するコンテキストや実行対象となるデバイスを指定する。

kernel の実行、input data への書き込み、 output data の読み込みといった
メモリ操作はこの Command Queue を通して行われる。

\subsection{メモリアクセス}
host 側は主に data を input/output するメモリ資源の確保を行う。
GPU のメモリ空間(図:\ref{fig:gpuarch})や Cell のメモリ空間(図:\ref{fig:cellarch})
はマルチコア CPU (図:\ref{fig:cpuarch})と違い、共有メモリでないため host と kernel ( task )間で data の共有ができない。
アクセスするにはメモリ空間間でコピーしなければならない。

GPGPU では host 側で memory buffer を作成してメモリのコピーを行う。
これらの処理や Task は Command Queue に enqueue することで実行される。
\begin{figure}[htpb]
  \begin{center}
    \includegraphics[scale=0.3]{./images/gpu_arch.pdf}
  \end{center}
  \caption{Gpu Architecture}
  \label{fig:gpuarch}
\end{figure}

\begin{figure}[htpb]
  \begin{center}
    \includegraphics[scale=0.5]{./images/cell_arch.pdf}
  \end{center}
  \caption{Cell Architecture}
  \label{fig:cellarch}
\end{figure}

\begin{figure}[htpb]
  \begin{center}
    \includegraphics[scale=0.5]{./images/cpu_arch.pdf}
  \end{center}
  \caption{Cpu Architecture}
  \label{fig:cpuarch}
\end{figure}

\subsection{データ並列}
多次元のデータ構造がある場合に高い並列度を保つには、それを分割して並列に実行する機能が必要である。
これを OpenCL ではデータ並列と読んでいる。
OpenCL は次元数に対応する index があり、 OpenCL は一つの記述から異なる index を持つ複数の kernel を自動生成する。
その添字を global\_id と呼ぶ。この時入力されたデータはワークアイテムという処理単位に分割される。

OpenCL はワークアイテムに対してそれぞれを識別する ID ( global\_id )を割り当てる。
kernel は get\_global\_id API によって ID を取得し、取得した ID に対応するデータに対して処理を行い、
データ並列を実現する。
この ID によって取得してきたワークアイテムをグローバルワークアイテムという。
また、ワークアイテムは3次元までのデータを渡すことができる。

データ並列による kernel 実行の場合は clEnqueueNDRangeKernel API を使用するが、
この関数の引数としてワークアイテムのサイズと次元数を指定することでデータ並列で実行できる。

\begin{figure}[htpb]
  \begin{center}
    \includegraphics[scale=0.60]{./images/workitem.pdf}
  \end{center}
  \caption{WorkItem ID}
  \label{fig:workitem_id}
\end{figure}

\subsection{ワークグループ}
前節でワークアイテムという処理単位について述べたが、
さらに複数個のグローバルワークアイテムを work\_group という単位にまとめることができる。
work\_group 内では同期やローカルメモリの共有が可能となる。

グローバルワークアイテム(ワークアイテム全体)の個数と、
ローカルワークアイテム(グループ一つ辺りのアイテム)の個数を指定することでワークアイテムを分割する。
なお、このときグローバルワークアイテム数はローカルアイテム数の整数倍でなければ
clEnqueueNDRangeKernel API 呼び出しは失敗する。

ローカルアイテム数は0を指定することで、コンパイル時に最適化させることができる。
したがってローカルアイテムのサイズは0を指定するのが一般的である。

なお、 work\_group を設定した場合は global\_id の他に work\_group\_id 、local\_id が
それぞれの kernel に割り当てられる(図:\ref{fig:workitem_id})。

なお、work\_groupを設定した場合はglobal\_idの他にwork\_group\_id、local\_idが
それぞれのkernelに割り当てられる(図:\ref{fig:workitem_id})。

kernel 側からそれぞれ ID に対応した API を使用して、各 ID を取得する。
取得した ID から自分が担当する index を計算して導く。
表:\ref{table:kernel_id_api}は kernel 側で使用できる、 ID を取得するための API となる。
\begin{tiny}
  \begin{table}[htpb]
    \begin{center}
      \caption{kernel で使用する ID 取得の API}
      \label{table:kernel_id_api}
      \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}
    \end{center}
  \end{table}
\end{tiny}
なお、 local\_id 、global\_id を取得する API は引数に0、1、2の値を set することができる。
id は x, y, z 座標があり、それぞれが 0, 1, 2 に対応している。
例えば get\_global\_id(1) と呼び出した場合は y 座標の、
get\_global\_id(1) と呼び出した場合は z 座標の global\_id を取得する。