# HG changeset patch # User Tatsuki IHA # Date 1517658995 -32400 # Node ID ebcf093795f32e559b2ba135562a495c2f905314 # Parent 8793903e4a0de4450781b21e303819fe7304718b Add twice examples diff -r 8793903e4a0d -r ebcf093795f3 paper/evaluation.tex --- a/paper/evaluation.tex Fri Feb 02 04:14:20 2018 +0900 +++ b/paper/evaluation.tex Sat Feb 03 20:56:35 2018 +0900 @@ -1,3 +1,99 @@ \chapter{Gears OS の評価} + +\section{実験環境} +今回 Twice、 BitonicSort をそれぞれ CPU、GPU環境で Gears OS の測定を行う。 + +使用する実験環境を\tabref{powerEdge}、 GPU 環境を\tabref{gtx1070} に示す。 + +\begin{table}[htbp] + \begin{center} + \begin{tabular}{|l|l|} \hline + Model & Dell PowerEdgeR630 \\ \hline + OS & CentOS 7.4.1708 \\ \hline + Memory & 768GB \\ \hline + CPU & 2 x 18-Core Intel Xeon 2.30GHz \\ \hline + \end{tabular} + \caption{実行環境} + \label{tab:powerEdge} + \end{center} +\end{table} + +\begin{table}[htbp] + \begin{center} + \begin{tabular}{|l||l|} \hline + GPU & GeForce GTX 1070 \\ \hline + Cores & 1920 \\ \hline + Clock Speed & 1683MHz \\ \hline + Memory Size & 8GB GDDR5 \\ \hline + Memory Bandwidth & 256GB/s \\ \hline + \end{tabular} + \caption{GPU 環境} + \label{tab:gtx1070} + \end{center} +\end{table} + \section{Twice} +Twice は与えられた整数配列のすべての要素を2倍にする例題である。 + +Twice のTask生成の方針として、CPU の場合は配列ある程度の範囲に分割してTaskを生成する。 +これは要素毎に Task を生成するとその分の Context を生成するために時間を取ってしまうからである。 + +GPU での実行は データ並列を使用して行う。 +GPU でデータ並列を実行する場合は Context とコピーなどは発生しないため、1要素に1スレッドを割り振って実行を行う。 + +Twice は並列実行の依存関係もなく、データ並列での実行に適した課題である。 +そのため、 通信時間を考慮しなければ CPU よりコア数が多い GPU が有利となる。 + +要素数$2^{23}$ のデータに対する Twice の実行結果を \tabref{twice}、\figref{twice}に示す。 +CPU 実行の際は $2^{23}$ のデータを 64 個のTask に分割して並列実行を行っている。 +ここでの ``GPU`` は CPU、 GPU 間のデータの通信時間を含めた時間、 ``GPU(kernel only)`` は kernel のみの実行時間である。 + +\begin{table}[htbp] + \begin{center} + \begin{tabular}{|l||l|} \hline + Processor & Time(ms) \\ \hline + 1 CPU & 147.946 \\ \hline + 2 CPUs & 80.773\\ \hline + 4 CPUs & 40.527\\ \hline + 8 CPUs & 20.267\\ \hline + 16 CPUs & 10.936\\ \hline + 32 CPUs & 5.878\\ \hline + GPU & 542.816\\ \hline + GPU(kernel only)& 0.755\\ \hline + \end{tabular} + \caption{$2^{23}$ のデータに対する Twice} + \label{tab:twice} + \end{center} +\end{table} + +\begin{figure}[htbp] + \begin{center} + \includegraphics[scale=0.6]{./fig/twice.pdf} + \end{center} + \caption{$2^{23}$ のデータに対する twice} + \label{fig:twice} +\end{figure} + +1 CPU と 32 CPU では 約 25.1 倍の速度向上が見られた。 +ある程度の台数効果があると考えられる。 + +GPU での実行は kernel のみの実行時間は 32CPU に比べて 約 7.8 倍の実行向上が見られた。 +しかし、通信時間を含めると 1 CPU より著しく遅い結果となってしまった。 +CPU、GPU の通信時間かボトルネックになっている事がわかる。 + \section{BitonicSort} +BitonicSort は並列処理向けのソートアルゴリズムである。 +代表的なソートアルゴリズムである Quick Sort も並列処理 を行うことが可能であるが、 QuickSort では ソートの過程で並列度が変動するため、台数効果が出づらい。 +一方でBitonic Sort は最初から最後まで並列度が変わらずに並列処理を行う。 +\figref{bitonicNetwork} は要素数8のデータに対する BitonicSort のソートネットワークである。 + +\begin{figure}[htbp] + \begin{center} + \includegraphics[scale=0.6]{./fig/bitonicNetwork.pdf} + \end{center} + \caption{要素数8の BtionicNetwork} + \label{fig:bitonicNetwork} +\end{figure} + +BitonicSort はステージ毎に決まった2点間の要素の入れ替えを並列に実行することによってソートを行う。 + diff -r 8793903e4a0d -r ebcf093795f3 paper/fig/bitonicNetwork.pdf Binary file paper/fig/bitonicNetwork.pdf has changed diff -r 8793903e4a0d -r ebcf093795f3 paper/fig/bitonicNetwork.xbb --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/paper/fig/bitonicNetwork.xbb Sat Feb 03 20:56:35 2018 +0900 @@ -0,0 +1,8 @@ +%%Title: fig/bitonicNetwork.pdf +%%Creator: extractbb 20170318 +%%BoundingBox: 0 0 591 422 +%%HiResBoundingBox: 0.000000 0.000000 591.000000 422.000000 +%%PDFVersion: 1.3 +%%Pages: 1 +%%CreationDate: Sat Feb 3 19:57:43 2018 + diff -r 8793903e4a0d -r ebcf093795f3 paper/fig/twice.pdf Binary file paper/fig/twice.pdf has changed diff -r 8793903e4a0d -r ebcf093795f3 paper/fig/twice.svg --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/paper/fig/twice.svg Sat Feb 03 20:56:35 2018 +0900 @@ -0,0 +1,194 @@ + + + +Gnuplot +Produced by GNUPLOT 5.2 patchlevel 2 + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + 0 + + + + + 100 + + + + + 200 + + + + + 300 + + + + + 400 + + + + + 500 + + + + + 600 + + + + + 1 cpu + + + + + 2 cpus + + + + + 4 cpus + + + + + 8 cpus + + + + + 16 cpus + + + + + 32 cpus + + + + + gpu + + + + + gpu(only) + + + + + + + + + time(ms) + + + + + CPUs + + + + + twice benchmark + + + gnuplot_plot_1 + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff -r 8793903e4a0d -r ebcf093795f3 paper/fig/twice.xbb --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/paper/fig/twice.xbb Sat Feb 03 20:56:35 2018 +0900 @@ -0,0 +1,8 @@ +%%Title: fig/twice.pdf +%%Creator: extractbb 20170318 +%%BoundingBox: 0 0 600 480 +%%HiResBoundingBox: 0.000000 0.000000 600.000000 480.000000 +%%PDFVersion: 1.3 +%%Pages: 1 +%%CreationDate: Sat Feb 3 18:15:19 2018 + diff -r 8793903e4a0d -r ebcf093795f3 paper/gpu.tex --- a/paper/gpu.tex Fri Feb 02 04:14:20 2018 +0900 +++ b/paper/gpu.tex Sat Feb 03 20:56:35 2018 +0900 @@ -37,6 +37,8 @@ また、block 内の thread 数は blockDim という組み込み変数で取得でき、これも3次元のベクター型になっている。 CUDA では これらの組み込み変数から thread が対応するデータを割り出し、データ並列の処理を行う。 +\newpage + \section{CUDAWorker} CUDAWorker は TaskManager から送信される CUDA用の Task を取得し、実行を行う。 @@ -85,6 +87,8 @@ 実際にcuLaunchKernel 関数を使用している部分を \coderef{cuLaunchKernel} に示す。 +\newpage + \lstinputlisting[caption=kernel に起動, label=code:cuLaunchKernel]{./src/cuLaunchKernel.cbc} Gears OS ではデータ並列 Task の際は Iterator Interface を持っており、 そこで指定した長さ、次元数に応じて cuLaunchKernel の引数を決定する(\coderef{cuLaunchKernel} 11-18行目)。 diff -r 8793903e4a0d -r ebcf093795f3 paper/master_paper.pdf Binary file paper/master_paper.pdf has changed diff -r 8793903e4a0d -r ebcf093795f3 paper/src/cudaTwice.cu --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/paper/src/cudaTwice.cu Sat Feb 03 20:56:35 2018 +0900 @@ -0,0 +1,5 @@ +extern "C" { + __global__ void twice(int* array) { + array[i+(blockIdx.x*blockDim.x+threadIdx.x)*prefix] = array[i+(blockIdx.x*blockDim.x+threadIdx.x)*prefix]*2; + } +}