# HG changeset patch # User Yuhi TOMARI # Date 1425262900 -32400 # Node ID 7990a2abbf052f409ec7960877d30e12970f069c # Parent 4bb76093b65a9feba8230997568b3be2100b99dc add file diff -r 4bb76093b65a -r 7990a2abbf05 paper/chapter8.tex --- /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 を用いることでロックの問題が解決できていることがわかる。 diff -r 4bb76093b65a -r 7990a2abbf05 paper/chapter9.tex --- /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 に割り当てることができる。