--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/chapter8.tex	Mon Mar 02 11:21:40 2015 +0900
@@ -0,0 +1,194 @@
+\chapter{ベンチマーク}
+\section{実験環境}
+今回使用する実験環境を表:\ref{tab:firefly_spec}、表:\ref{tab:dragonfly_spec}に示す.
+
+\begin{table}[!htbp]
+  \begin{center}
+    \begin{tabular}{|c||c|} \hline
+      名前 & 概要 \\ \hline \hline
+      Model & MacPro Mid 2010 \\ \hline
+      CPU & 6-Core Intel Xeon @2.66GHz \\ \hline
+      Serial-ATA Device & HDD ST4000VN000-1H4168\\ \hline
+      Memory & 16GB \\ \hline
+      OS & MacOSX 10.10.1 \\ \hline
+      Graphics & NVIDIA Quadro K5000 1536Core(4096MB) \\ \hline
+    \end{tabular}
+  \end{center}
+  \caption{Ceriumを実行する実験環境1}
+  \label{tab:firefly_spec}
+\end{table}
+
+\begin{table}[!htbp]
+  \begin{center}
+    \begin{tabular}{|c||c|} \hline
+      名前 & 概要 \\ \hline \hline
+      Model &  MacPro Late 2013 \\ \hline
+      CPU & 6-Core Intel Xeon E5@3.5GHz \\ \hline
+      Serial-ATA Device & Apple SSD SM0256 \\ \hline
+      Memory & 16GB \\ \hline
+      OS & MacOSX 10.10.1 \\ \hline
+      Graphics & AMD FirePro D700 2048Core(6144MB) \\ \hline
+    \end{tabular}
+  \end{center}
+  \caption{Ceriumを実行する実験環境2}
+  \label{tab:dragonfly_spec}
+\end{table}
+
+なお、表:\ref{tab:firefly_spec}と表:\ref{tab:dragonfly_spec}は
+CPU クロック数の他にも Strage や GPU の性能にも違いがある。
+実験環境1(表:\ref{tab:firefly_spec})は実験環境2(表:\ref{tab:dragonfly_spec})に比べてクロック数が低く、
+ Strage は HDD を使用している。GPU は NVIDIA を使用している。
+実験環境2(表:\ref{tab:dragonfly_spec})はクロック数が高く、 Strage に SSD を使用している。
+GPU は AMD を使用している。
+
+以上の環境で今回新たに実装したマルチコア、GPGPU 、並列 I/O のベンチマークを行う。
+
+\section{マルチコア}
+マルチコア CPU における並列実行について、
+Sort(図:\ref{fig:sort_on_multicore}) と
+WordCount(図:\ref{fig:wordcount_on_multicore}) 
+\if0、
+FFT(図:\ref{fig:fft_on_multicore})
+\fi
+によるベンチマークを行った。
+
+\begin{figure}[htpb]
+  \begin{center}
+    \includegraphics[scale=1.0]{./figures/multicore/sort.pdf}
+  \end{center}
+  \caption{マルチコア CPU における Sort}
+  \label{fig:sort_on_multicore}
+\end{figure}
+
+\begin{figure}[htpb]
+  \begin{center}
+    \includegraphics[scale=1.0]{./figures/multicore/word_count.pdf}
+  \end{center}
+  \caption{マルチコア CPU における WordCount}
+  \label{fig:wordcount_on_multicore}
+\end{figure}
+
+MacPro 2013 において、6CPU を使用した場合、1CPU を利用した場合と比較して、
+Sort は 5.2倍、WordCount は 4.9倍の速度向上が見られる。
+MacPro 2010 においても Sort は5.65倍、WordCount は5.0倍の速度向上が見られた。
+MacPro のコア数以上の Thread 数になると並列度の低下が見られる。
+
+Cerium は実行時に -pre オプションをつけることで使い分ける事ができる。
+DMA の prefetch 機能によるデータ転送のベンチマークを行った(図:\ref{fig:prefetch_bench})。
+
+\begin{figure}[htpb]
+  \begin{center}
+    \includegraphics[scale=1.0]{./figures/dma/dmabench.pdf}
+  \end{center}
+  \caption{Word Count による prefetch機能のベンチマーク}
+  \label{fig:prefetch_bench}
+\end{figure}
+
+測定の結果、CPU数が1の場合は prefetch オプションを入れると1.17\%、CPU数が6の場合は1.63\%の性能向上が見られた。
+\newpage
+\section{GPGPU}
+GPGPU を行う際はデータ並列による処理を行った時、充分に性能を発揮することができる(\ref{sec:shared_memory}節)。
+WordCount による OpenCL、CUDA、マルチコア CPU 上における
+データ並列実行の性能評価を行った(図:\ref{fig:dataparallel})。
+なお、MacPro 2013 (表:\ref{tab:dragonfly_spec}) は GPU が NVIDIA 製でないため、仕様上 CUDA による測定ができない。
+MacPro 2010 で測定を行った。
+
+\begin{figure}[htpb]
+  \begin{center}
+    \includegraphics[scale=1.0]{./figures/GPU/wordcount_dataparallel.pdf}
+  \end{center}
+  \caption{Word Count によるデータ並列実行のベンチマーク}
+  \label{fig:dataparallel}
+\end{figure}
+
+データ並列実行することにより、
+マルチコア CPU では 1.06倍の性能向上が見られた。
+GPU に関しては特に性能向上が大きく、OpenCL においては115倍、CUDA においては14倍の性能向上が見られた。
+この結果から、GPGPU を行う際はデータ並列による実行が必須であることがわかる。
+更に、マルチコア CPU においても性能向上が見られた。
+データ並列を行わない場合、OpenCL の実行時間が CUDA の約6倍かかっている。
+これはMacPro2010 の GPU が NVIDIA 製であることが要因として考えられる。
+CUDA は NVIDIA 製の GPU でのみ動作するため、 NVIDIA に特化した最適化が行われていると考えられる。
+
+データ並列によって性能向上は実現できたが、
+マルチコア CPU と比較すると実行時間が OpenCL では4.2倍、CUDA では 5.3倍かかっている。
+WordCount の例題は特定の文字を区切り文字とし、区切り文字による分岐でカウントアップしている。
+区切り文字による分岐がマルチコア CPU と比較して性能を落としていると考えられる。
+
+GPU は分岐命令を苦手としており、GPU で並列度を維持するには分岐を最小限にする必要がある。
+
+そこで、FFT(図:\ref{fig:fft_bench})による測定を行う。
+このベンチマークにより、GPU の制約に当てはまる Task であれば並列度を維持できることを示す。
+OpenCL、CUDA、マルチコアCPU の性能比較を行う。
+更に、Cerium を用いないで OpenCL を使用した場合(OpenCL-original)についても測定を行う。
+
+\begin{figure}[htpb]
+  \begin{center}
+    \includegraphics[scale=1.0]{./figures/GPU/fft_firefly.pdf}
+  \end{center}
+  \caption{マルチコア CPU、OpenCL、CUDA における FFT}
+  \label{fig:fft_bench}
+\end{figure}
+
+CUDA は 1CPU に比べて3.5倍、6CPU に比べて1.1倍の性能向上が見られる。
+OpenCL は 1CPU に比べて2.75倍の性能向上が見られたが、
+6CPU と比べると0.87倍、OpenCL-original と比べると0.76倍の性能低下が見られた。
+高性能の GPU を使用することで OpenCL でも並列度が向上が期待できる。
+また、Cerium では Task のパイプライン実行により、 GPU よりも上のレイヤでの並列化を行っている。
+CPU の性能を上げることで Scheduling 部分が高速になり、OpenCL-original と比較した場合の性能向上も見込める。
+
+そこで、MacPro 2013 にて測定を行った(図:\ref{fig:fft_bench_dragonfly})。
+
+\begin{figure}[htpb]
+  \begin{center}
+    \includegraphics[scale=1.0]{./figures/GPU/fft_dragonfly.pdf}
+  \end{center}
+  \caption{MacPro 2013 における FFT}
+  \label{fig:fft_bench_dragonfly}
+\end{figure}
+
+より高性能な GPU を搭載した計算機(表:\ref{tab:dragonfly_spec})で測定したところ、
+OpenCL が 1CPU に比べて6倍、6CPU に比べて1.6倍の性能向上が見られた。
+マルチコア CPU での実行速度も向上しているため、
+GPU の性能だけでなく、 CPU のクロック数やStrage に SSD を使用している事も性能向上の要因と考えられる。
+SSD はランダムアクセスでのデータ読み込み性能が高く、ディスク読み書きに関するオーバーヘッドの改善が見込める。
+
+Cerium による実行は OpenCL-original による実行とほぼ同じ性能で、約1\% OpenCL-original の方が速い。
+Cerium のパイプライン構造を利用しており、
+本来 OpenCL-original よりも高い並列度が期待できる、まだチューニングを行う余地がある事がわかる。
+
+\section{並列 I/O}
+Cerium の従来のファイル読み込みである mmap、一般的な File Open であるread、
+今回実装したBlocked Read を比較した測定を行った。
+なお、Blocked Read については IO Threads を使用した場合としてない場合(SPE\_ANY)両方の測定を行う。
+例題として Word Count を使用した測定を行った。
+図:\ref{fig:io_bench_firefly}がMacPro 2010における測定で、
+図:\ref{fig:io_bench_dragonfly}がMacPro 2013における測定となる。
+
+\begin{figure}[htpb]
+  \begin{center}
+    \includegraphics[scale=1.0]{./figures/io/io_thread_firefly.pdf}
+  \end{center}
+  \caption{WordCount によるファイル読み込み方式のベンチマーク(MacPro2010)}
+  \label{fig:io_bench_firefly}
+\end{figure}
+
+\begin{figure}[htpb]
+  \begin{center}
+    \includegraphics[scale=1.0]{./figures/io/io_thread_dragonfly.pdf}
+  \end{center}
+  \caption{WordCount によるファイル読み込み方式のベンチマーク(MacPro2013)}
+  \label{fig:io_bench_dragonfly}
+\end{figure}
+
+6CPU においてBlockedRead\_IO を使用した場合、mmap に比べて1.1倍、read に比べて1.58倍、
+BlockedRead\_SPE\_ANY に比べて1.34倍の性能向上が見られた。
+しかし、実験環境のコア数である6CPU 以上になると並列度は低下していき、
+8CPU からは BlockedRead による並列実行に比べて mmap によるファイルの先読みが有効に働いている。
+
+また、SPE\_ANY による BlockedRead は他の読み込み形式と違い、
+10CPU の時、CPU 数を増やしたにも関わらず極端に処理が遅くなっている。
+これは\ref{sec:spe_problem}節で述べた、ReadTask に対する 実行 Task の割り込みにより
+ロックがかかる問題が起きていると考えられる。
+IO Thread を用いた BlockedRead では極端な速度低下は起きていない。
+測定の結果から、IO Thread を用いることでロックの問題が解決できていることがわかる。

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/chapter9.tex	Mon Mar 02 11:21:40 2015 +0900
@@ -0,0 +1,128 @@
+\chapter{既存のプログラミングフレームワークとの比較}
+\section{OpenCL}
+多次元のデータ構造を扱う計算において高い並列度を保つには、
+多次元データを分割して並列に実行する機能が必要である。
+これをデータ並列実行という。 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 を取得する。
+
+\section{CUDA}
+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 のデータプールに登録されている必要がある。
+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 に割り当てることができる。