\chapter{既存のマルチプラットフォームフレームワーク}
マルチプラットフォームでプログラムを動作させる場合、そのアーキテクチャを意識する必要がある。
マルチプラットフォームにはマルチコア CPU 、 GPU や Cell といったヘテロジニアスマルチコアのような
様々な構成がある。

\section{Architecture}
本研究では、 CPU の他に GPU 上でのプログラミング (GPGPU) にも対応する。

GPU(Graphics Processing Unit) は PC の画像処理を担当するユニットで、
レンダリングに特化したプロセッサが多く集まった構造を持つ。
一つ一つのプロセッサの構造は単純で、その機能は CPU に比べて限定的ではあるが
大量のデータを複数のプロセッサで並列処理することに長けている。

GPGPU (General Purpose computing on Graphics Processing Units) とは、
 GPU の高い演算能力を画像処理ではなく汎用計算に使用することである。

\section{Shared Memory}
\label{sec:shared_memory}
計算機にはメモリ空間が別の計算機と、共有メモリ(Shared Memory)な計算機がある。
GPU のメモリ空間(図:\ref{fig:gpuarch})はマルチコア CPU (図:\ref{fig:cpuarch})と違い、
共有メモリ(shared mermoy)でないので Host と Device 間で Data の共有ができない。
そのためマルチプラットフォーム環境に対応したフレームワークには、
Device と Host 間でデータの転送を行う API 備わっている。
しかし、異なる Device 間でデータの転送を行うとネックになる。
そのためデータの入出力を行う回数を減らす、入出力の処理をパイプライン処理にするなどの工夫が必要になる。


\begin{figure}[htpb]
  \begin{center}
    \includegraphics[scale=0.4]{./images/gpu_arch.pdf}
  \end{center}
  \caption{GPU Architecture}
  \label{fig:gpuarch}
\end{figure}

\begin{figure}[htpb]
  \begin{center}
    \includegraphics[scale=0.8]{./images/cpu_arch.pdf}
  \end{center}
  \caption{CPU Architecture}
  \label{fig:cpuarch}
\end{figure}
\newpage
%--------
% OpenCL
%--------

\section{OpenCL}
OpenCL とは、 Khronos Group の提供するマルチコア CPU と GPU といった、
ヘテロジニアス環境を利用した並列計算を支援するフレームワークである。

OpenCL では演算用プロセッサ側を Device 、制御用デバイス側を Host として定義する。
また、 Device 上で動作するプログラムの事を kernel と呼ぶ。

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

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

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

\begin{itemize}
\item OpenCL C言語
\item OpenCL Runtime API
\end{itemize}

OpenCL C は演算用プロセッサ上で動作する、 C 言語を拡張したプログラミング言語である。
一方で OpenCL Runtime API は OpenCL C で記述した kernel を Queuing するために Host が利用する API である。

Host では主に Data を input/output するメモリ資源の確保を行う。
OpenCL は host 側で memory buffer を作成してメモリのコピーを行う。
これらの処理や Task は Command Queue に enqueue することで実行される。

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

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

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

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

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

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

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

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

なお、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}
      \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}
      \caption{kernel で使用する ID 取得の API}
      \label{table:kernel_id_api}
    \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 を取得する。


%------
% CUDA
%------
\section{CUDA}
CUDA とは、半導体メーカーNVIDIA社が提供するGPUコンピューティング向けの総合開発環境である。

CUDA も OpenCL と同様、演算用プロセッサ (GPU) を Device 、制御用デバイス側を Host として定義する。
また、 Device 上で動作するプログラムの事も kernel と呼ぶ。

OpenCL における Command と CommandQueue に対応するものとして、 CUDA には Operation と Stream がある。
Stream は Host 側で発行された Operation を一連の動作として Device で実行する。
Operation は発行された順序で実行されることが保証されている。
更に、異なる Stream に発行された Operation も依存関係が存在しない場合、Operationは並列に実行される。
更に依存関係が存在しない、異なる Stream に発行された Operation は並列に実行される。

CUDAには主に3つの仕様がある。

\begin{itemize}
\item CUDA C
\item CUDA Runtime API
\item CUDA Driver API
\end{itemize}
CUDA C は GPU 上で動作する、C 言語を拡張したプログラミング言語である。
CUDA Runtime API も CUDA Driver API も CUDA C で記述した Kernel を Queueing するために
 Host が利用するAPIである。
Driver API は Runtime APIに比べ、プログラマが管理しなければならないリソースが多くなる代わり、
より柔軟な処理を行う事ができる。

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

CUDA では OpenCL の WorkItemに相当する単位を thread として定義している。
この thread をまとめた単位として block がある。

CUDAでデータ並列による kernel 実行を行う場合、cuLaunchKernelAPIを使用する。
この関数は引数として各座標の block 数、
各座標の block 1つ辺りの thread 数を指定することによりデータ並列実行を行う。

cuLaunchKernel で kernel を実行すると各 thread に対して blockID と threadID が割り当てられる。
CUDA には OpenCLと異なり、IDを取得するAPIが存在しない。
それに代わり、 kernel に組み込み変数が準備されている。
その組み込み変数を参照し、対応するデータに対し処理を行うことでデータ並列実行を実現する。
組み込み変数は以下の3つである。

\begin{itemize}
\item uint3 blockDim
\item uint3 blockIdx
\item uint3 threadIdx
\end{itemize}

3つの組み込み変数はベクター型で、 blockDim.x とすると x 座標の thread 数を参照することができる。
同じように blockID 、 threadID の x 座標を参照することができる。
blockDim.x * blockIdx.x + threadIdx.x とする事で OpenCL における get\_global\_id(0) で
取得できる ID に相当する値を算出する事ができる。

例としてある kernel で get\_global\_id(0) の値が 8 の時、
 CUDA では 図\ref{fig:calculateIndex}のように算出する。


\begin{figure}[htpb]
  \begin{center}
    \includegraphics[scale=0.5]{./images/calculateIndex.pdf}
  \end{center}
  \caption{Calculate Index example}
  \label{fig:calculateIndex}
\end{figure}

\section{StarPU}
StarPUはフランス国立情報学自動制御研究所 (INRIA) の StarPU 開発チームの提供する、
ヘテロジニアス環境向けのフレームワークである。
 GPU の制御に OpenCL と CUDA を用いており、どちらかを選択することで GPU 上で実行することができる。

 OpenCL と CUDA における実行の単位は kernel だったが、 StarPU では実行の単位を Task と定義している。

 StarPU では Task を制御するためにcodeletと呼ばれる構造体を使う。
 codelet を Task 生成時にポインタ渡しすることで、
 演算を行うリソースや実行する関数等を指定することができる。
 CPU と GPU で並列に実行する例を\ref{src:codelet}に示す。

\begin{lstlisting}[frame=lrbt,label=src:codelet,caption=codeletの例,numbers=left]
starpu_codelet codelet = {
    .where     = STARPU_CPU|STARPU_CUDA,
    .cpu_func  = cpu_function,
    .cuda_func = cuda_function,
};
\end{lstlisting}

計算に必要なデータは、 StarPU のデータプールに登録されている必要がある。
StarPU ではデータを starpu\_data\_handle という型で登録する。
Task はこの handle を参照することで値を参照することができる。

\begin{figure}[htpb]
  \begin{center}
    \includegraphics[scale=0.5]{./images/starpu_data_parallel.pdf}
  \end{center}
  \caption{StarPUにおけるデータ分割}
  \label{fig:data_partition}
\end{figure}

図:\ref{fig:data_partition}に StarPU におけるデータ並列実行の流れを示す。
 StarPU では配列の初期化や代入を行った後、
 starpu\_data\_register 関数を使って StarPU のデータプールに登録する。

 データ並列で実行する場合、更にデータを分割する必要がある。
 starpu\_data\_partition 関数を用いる事で分割を行うことができる。
 分割数を指定することで、データプールに登録したデータを chunk と呼ばれる単位に分割する。
 starpu\_task\_submit 関数により chunk を CPU や GPU に割り当てることができる。