# HG changeset patch # User Yuhi TOMARI # Date 1423998753 -32400 # Node ID 7956856211c59998ff92b96b7b92121f6125b92f # Parent 9b181351ca4db4a34d110daa27d47d3cb4825e48 move data multicore-parallel from GPGPU-chapter to Multicore-chapter diff -r 9b181351ca4d -r 7956856211c5 paper/chapter4.tex --- a/paper/chapter4.tex Sun Feb 15 12:04:36 2015 +0900 +++ b/paper/chapter4.tex Sun Feb 15 20:12:33 2015 +0900 @@ -25,3 +25,78 @@ 更に、DMA には prefetch 機能がある。 DMA の転送効率を向上させるため、 DMA で転送する送信データを予め取り込むことができる。 Cerium では DMA によるprefetch、DMA を用いないポインタ渡し、更に明示的なコピーによるデータ転送をサポートする。 + +\section{データ並列} +\label{sec:multicore_dataparallel} +並列プログラミングを行う際、並列化の方式としてタスク並列とデータ並列の2つがある。 +Cerium における並列処理は、タスク並列により実現されている。 + +タスク並列は1つのデータに対して異なる処理方法を適用し、それぞれ独立して実行させるものである。 +一方で1つの Task に対して多くのデータを与え、データごとに独立した処理を行わせる事をデータ並列という。 +プログラムを並列化する際はどちらの並列化手法が適しているか考える必要がある。 +独立した Task が充分にある場合はタスク並列が有効となる。 +処理対象となるデータが充分な数のサブデータへ分割することが可能で、 +サブデータへの処理を行うことが元のデータへの処理と同じになる場合、データ並列が有効となる。 + +本研究で Cerium でデータ並列による実行を可能にした。 +Cerium でデータ並列実行を行う場合、Task を spwan API でなく iterate API で生成すればよい。 +iterate API は複数の length を引数とし、 +length の値がデータ分割後に各 Task が担当するサイズ、length の個数がデータの次元数となる。 +これを元にScheduler が各 Task が担当する index を計算し、Task に set\_param する。 + +index の割り当ての例を表:\ref{table:dataparallel_index}に示す。 +データ数10個の入力を持つ Task に対して CPU 数4、 +一次元における分割でデータ並列実行した場合の index の割り当ては表:\ref{table:dataparallel_index}になる。 + +この例だと各 CPU に対する index の割り当ては CPU0 は index 0、4、8、 CPU1 は index 1、5、9、 +CPU2 は index 2、6、CPU3 は index 3、7となる。 + +\begin{tiny} + \begin{table}[htpb] + \begin{center} + \small + \begin{tabular}[htpb]{c||c|c|c|c} + \hline + stage & CPU0 & CPU1 & CPU2 & CPU3 \\ + \hline + \hline + 1 & 0 & 1 & 2 & 3 \\ + \hline + 2 & 4 & 5 & 6 & 7 \\ + \hline + 3 & 8 & 9 & & \\ + \hline + \end{tabular} + \caption{データ並列実行時の index の割り当て} + \label{table:dataparallel_index} + \end{center} + \end{table} +\end{tiny} + +並列プログラミングだと、並列化部分が全て同一の Task であるということは少なくない。 +その際、 Task 生成部分をループで回すことなく、iterate API により簡単な Syntax で記述することができる。 + +マルチコア CPU 上でデータ並列実行する場合、ソースコード:\ref{src:multiply_cpu}のように Task を記述する。 +なお、2つの input データの積を output データに格納する例題、 multiply を用いた。 + +\begin{lstlisting}[frame=lrbt,label=src:multiply_cpu, caption=Multiply(CPU),numbers=left] +static int +run(SchedTask *s, void *rbuf, void *wbuf) { + float *indata1, *indata2, *outdata; + + indata1 = (float*)s->get_input(rbuf, 0); + indata2 = (float*)s->get_input(rbuf, 0); + outdata = (float*)s->get_output(wbuf, 0); + + long id = (long)s->get_param(0); + outdata[id] = indata1[id] * indata2[id]; + return 0; +} +\end{lstlisting} + +Task 間で共有する Input/Output データと自分が計算を行う index は Scheduler により送られてきている。 +get\_input 、get\_output API を用いて Input/Output データを取得し、get\_param API で担当する index を取得する。 +後は各自担当範囲に対して計算を行うだけでよい。 + +データ並列で実行する場合、1つの Input と Output を各 Task 間で共有し、 +各 Task は自分が担当する index に対してのみ計算を行うため、少ないコピーに抑えられる。 diff -r 9b181351ca4d -r 7956856211c5 paper/chapter5.tex --- a/paper/chapter5.tex Sun Feb 15 12:04:36 2015 +0900 +++ b/paper/chapter5.tex Sun Feb 15 20:12:33 2015 +0900 @@ -40,12 +40,20 @@ \section{データ並列} 並列プログラミングにおいて、明示的な並列化部分はループ部分である。 - GPU は数百個のコアを有しており、ループ部分に対してデータ並列で処理を行うことで CPU より高速に演算を行う事ができる。 +GPU は数百個のコアを有しており、ループ部分に対してデータ並列で処理を行うことで CPU より高速に演算を行う事ができる。 プログラムの大部分がループであれば、データ並列による実行だけでプログラムの性能は向上する。 + OpenCL 、 CUDA ともにデータ並列をサポートしている。 -Task を実行する際にデータをどう分割するか指定し、kernel をデータ並列実行用に書き換えることで実現する。 -データ並列実行用の kernel は以下のように記述する。 -2つの input データの積を output データに格納する例題、 multiply を用いる。 +OpenCL と CUDA はTask を実行する際にデータをどう分割するか指定し、 +kernel にデータ並列用の処理を加えることで可能となる。 +\ref{sec:multicore_dataparallel}節で Cerium でマルチコア CPU におけるデータ並列を可能にした。 +GPGPU においてもデータ並列実行をサポートする。 +GPU 上でのデータ並列実行もマルチコア CPU と変わらず、iterate API によりデータ並列用の Task を生成することができる。 +iterate で Task を生成することで Scheduler が OpenCL 及び CUDA の API に適切なパラメタを渡している。 +Task の生成部分は マルチコア CPU と GPU で完全に同じ形式で記述できる。 + +データ並列実行の際、Task は以下のように記述する。 +なお、例題は multiply を用いている。 \begin{lstlisting}[frame=lrbt,label=src:multiply_opencl,caption=Multiply(OpenCL),numbers=left] __kernel void @@ -73,84 +81,21 @@ } \end{lstlisting} -このような kernel を分割数分生成する。 -分割数は kernel の生成時にそれぞれのフレームワークが 用意している API を用いて指定する。 -いずれの kernel も +このような Task を分割数分生成する。 +分割数は Task それぞれのフレームワークが用意している API を用いて指定する。 + \begin{itemize} \item 自分の計算する範囲を取得(ソースコード\ref{src:multiply_opencl}、ソースコード\ref{src:multiply_cuda}の7行目) \item 取得した範囲を計算(ソースコード\ref{src:multiply_opencl}、ソースコード\ref{src:multiply_cuda}の9行目) \end{itemize} -という手順で処理する。 -計算する範囲については OpenCL では取得用の API を用い、 CUDA では kernel の持つ組み込み変数から算出する。 - -Cerium でも データ並列実行をサポートする。 - GPU におけるデータ並列実行だけでなくマルチコア CPU 上でのデータ並列実行にも対応する。 -なお、マルチコア CPU 上で実行する場合も GPU 実行時の kernel -(ソースコード\ref{src:multiply_opencl}、ソースコード\ref{src:multiply_cuda}) となるべく近い形式で記述できるようにする。 -マルチコア CPU 上でデータ並列実行する場合、 kernel は以下のように記述する。 - -\begin{lstlisting}[frame=lrbt,label=src:multiply_cpu, caption=Multiply(CPU),numbers=left] -static int -run(SchedTask *s, void *rbuf, void *wbuf) { - float *indata1, *indata2, *outdata; - - indata1 = (float*)s->get_input(rbuf, 0); - indata2 = (float*)s->get_input(rbuf, 0); - outdata = (float*)s->get_output(wbuf, 0); - - long id = (long)s->get_param(0); - outdata[id] = indata1[id] * indata2[id]; - return 0; -} -\end{lstlisting} - -OpenCL 、CUDA と違い値を引数として直接渡すのではなく、メモリバッファから Load し、計算を行う。 -値渡しや修飾子等若干の違いはあるが、ほぼ同じ形式で kernel を記述することができる。 +いずれの Task も上記の手順で処理を行っている。 +計算する範囲について、 OpenCL では取得用の API を用い、 CUDA では kernel の持つ組み込み変数から算出する。 +マルチコア CPU では引数としてデータを直接渡していたが、OpenCL 、CUDA では上記の方法でメモリバッファから Load し、計算を行う。 +値渡しや修飾子等若干の違いはあるが、OpenCL 、CUDA ともにマルチコア CPU(ソースコード:\ref{src:multicore_cpu}) とほぼ同じ形式で kernel を記述することができる。 CPU、 OpenCL、 CUDA いずれか1つの記述から残りのコードも生成できるようにする事が望ましい。 -Cerium でデータ並列実行を行う場合、Task を spwan API でなく iterate API で生成すればよい。 -iterate API は複数の length を引数とし、 -length の値がデータ分割後に各 Task が担当するサイズ、length の個数がデータの次元数となる。 -これを元にScheduler が各 Task が担当する index を計算し、Task に set\_param する。 - -Task は実行時に get\_param することで set\_param した値を取得し、担当範囲をデータ並列を実行する。 -この get\_param が OpenCL における get\_global\_id API に相当する。 - -index の割り当ての例を示す。 -データ数10個の入力を持つ Task に対して CPU 数4、 -一次元における分割でデータ並列実行した場合の index の割り当ては表:\ref{table:dataparallel_index}になる。 - -この例だと各 CPU に対する index の割り当ては CPU0 は index 0、4、8、 CPU1 は index 1、5、9、 -CPU2 は index 2、6、CPU3 は index 3、7となる。 - -\begin{tiny} - \begin{table}[htpb] - \begin{center} - \small - \begin{tabular}[htpb]{c||c|c|c|c} - \hline - stage & CPU0 & CPU1 & CPU2 & CPU3 \\ - \hline - \hline - 1 & 0 & 1 & 2 & 3 \\ - \hline - 2 & 4 & 5 & 6 & 7 \\ - \hline - 3 & 8 & 9 & & \\ - \hline - \end{tabular} - \caption{データ並列実行時の index の割り当て} - \label{table:dataparallel_index} - \end{center} - \end{table} -\end{tiny} - -並列プログラミングだと、並列化部分が全て同一の Task であるということは少なくない。 -その際、 Task 生成部分をループで回すことなく、簡単な Syntax で記述することができる。 - データ並列で実行する場合、 Input と Output を各 Task 間で共有するため、少ないコピーに抑えられる。 -CPU ではメモリ領域を節約する事ができるが、 Task と Manager でメモリ領域が同じ(\ref{sec:shared_memory}節)なため、 +CPU ではメモリ領域を節約する事はできるが、 Task と Manager でメモリ領域が同じ(\ref{sec:shared_memory}節)なため、 コピーによるオーバーヘッドは少ない。 -しかし GPU は SharedMemory ではなく、データの転送がオーバーヘッドとなる。 -コピーを減らす事で並列度の向上が見込める。 +しかし GPU は SharedMemory ではなく、データの転送がオーバーヘッドとなるため、コピーを減らす事で並列度の向上が見込める。 diff -r 9b181351ca4d -r 7956856211c5 paper/chapter6.tex --- a/paper/chapter6.tex Sun Feb 15 12:04:36 2015 +0900 +++ b/paper/chapter6.tex Sun Feb 15 20:12:33 2015 +0900 @@ -9,7 +9,7 @@ 読み込んでいる間は他の CPU が動作せず、並列度が落ちてしまう。 そこで、 I/O 部分も並列に動作するよう実装した。 -Read を並列に行うには、File Open ではなく mmap を使う方法がある。 +Read を並列に行うには、File Open ではなく mmap を使う方法がある。Cerium でも mmap を使用していた。 mmap はすぐにファイルを読みに行くのではなく、まず仮想メモリ空間にファイルの中身を対応させる。 メモリ空間にアクセスが行われると、OS が対応したファイルを読み込む。 @@ -33,7 +33,8 @@ 読み込みが並列に実行されない場合、 Task が読み込み待ちを起こしてしまう。 読み込みが OS 依存となるため、環境によって左右されやすく、汎用性を損なってしまう。 -そこで、mmap を使わず、read を独立したスレッドで行い、読み込まれた部分に倒して並列に Task を起動する。 +そこで、mmap を使わず、read を独立したスレッドで行い、ファイルを一度に全て読み込むのではなく +ある程度の大きさ (Block) 分読み込み、読み込まれた部分に倒して並列に Task を起動する。 これを Blocked Read と呼ぶ。Blocked Read によるプログラミングは複雑になるが、高速化が期待できる。 \section{Blocked Read による I/O の並列化} @@ -58,6 +59,8 @@ その問題を解決するため、依存関係を設定する。 BlockedReadによる読み込みが終わってから TaskBlock が起動されるよう、 Cerium の API である wait\_for により依存関係を設定する。 +しかし、Task が BlockedRead を追い越すことによるロックはオーバーヘッドとなるため、起こさないようにしたい。 +つまり、 BlockedRead は連続で走っている必要がある。 以上を踏まえ、BlockedRead の実装を行った。 BlockedRead Task の生成はソースコード:\ref{src:blockedread_create}のように行う。 @@ -107,6 +110,7 @@ 受け取ったパラメタをそれぞれ pread 関数に渡すことで Blocked Read を実現している。 \newpage \section{I/O 専用 Thread の実装} +\label{sec:spe_problem} Cerium Task Manager では、各種 Task にデバイスを設定することができる。 SPE\_ANY 設定を使用すると、 Task Manager で CPU の割り振りを自動的に行う。 BlockedRead は連続で読み込まれなければならないが、 @@ -137,4 +141,5 @@ IO\_0 で実行される Task は BlockedRead のみなので、 IO\_0 のpriority を高く設定することで Blocked Read は連続で実行される。 -また、以上の事から I/O を含む処理では、 I/O を行う Thread の priority を高くする必要があるという知見を得られた。 +また、以上の事から I/O を含む並列処理において読み込みの Task と計算を行う Task を並列で走らせたい場合、 +I/O を行う Thread の priority を高くする必要があるという知見を得られた。 diff -r 9b181351ca4d -r 7956856211c5 paper/figures/GPU/fft/bar_plot.sh --- a/paper/figures/GPU/fft/bar_plot.sh Sun Feb 15 12:04:36 2015 +0900 +++ b/paper/figures/GPU/fft/bar_plot.sh Sun Feb 15 20:12:33 2015 +0900 @@ -22,7 +22,7 @@ set title "${fstitle} benchmark" # x,yの範囲をどこからどこまでにするのか - set xrange [ 0 : 4 ] + set xrange [ 0 : 5 ] set yrange [ 0 : 350 ] # set yrange [ 0 : 0.25] diff -r 9b181351ca4d -r 7956856211c5 paper/figures/GPU/fft/fft_dragonfly.txt --- a/paper/figures/GPU/fft/fft_dragonfly.txt Sun Feb 15 12:04:36 2015 +0900 +++ b/paper/figures/GPU/fft/fft_dragonfly.txt Sun Feb 15 20:12:33 2015 +0900 @@ -1,3 +1,4 @@ 1 1CPU 328.2477 2 6CPU 87.2767 3 OpenCL 54.6618 +4 OpenCL-original 54.0335 diff -r 9b181351ca4d -r 7956856211c5 paper/figures/GPU/fft/fft_firefly.txt --- a/paper/figures/GPU/fft/fft_firefly.txt Sun Feb 15 12:04:36 2015 +0900 +++ b/paper/figures/GPU/fft/fft_firefly.txt Sun Feb 15 20:12:33 2015 +0900 @@ -1,4 +1,5 @@ -1 1CPU 484.6164 -2 6CPU 153.434 -3 CUDA 138.4115 -4 OpenCL 176.6674 +1 1CPU 484.6164 +2 6CPU 153.434 +3 CUDA 138.4115 +4 OpenCL 176.6674 +5 OpenCL-original 133.999 diff -r 9b181351ca4d -r 7956856211c5 paper/figures/GPU/fft_dragonfly.pdf Binary file paper/figures/GPU/fft_dragonfly.pdf has changed diff -r 9b181351ca4d -r 7956856211c5 paper/figures/GPU/fft_firefly.pdf Binary file paper/figures/GPU/fft_firefly.pdf has changed diff -r 9b181351ca4d -r 7956856211c5 paper/master_paper.aux --- a/paper/master_paper.aux Sun Feb 15 12:04:36 2015 +0900 +++ b/paper/master_paper.aux Sun Feb 15 20:12:33 2015 +0900 @@ -53,88 +53,95 @@ \@writefile{lot}{\addvspace {10\p@ }} \@writefile{toc}{\contentsline {section}{\numberline {5.1}マルチコア上での実行の機構}{17}} \@writefile{toc}{\contentsline {section}{\numberline {5.2}DMA}{17}} -\@writefile{toc}{\contentsline {chapter}{\numberline {第6章}GPGPU への対応}{18}} +\@writefile{toc}{\contentsline {section}{\numberline {5.3}データ並列}{17}} +\newlabel{sec:multicore_dataparallel}{{5.3}{17}} +\@writefile{lot}{\contentsline {table}{\numberline {5.1}{\ignorespaces データ並列実行時の index の割り当て}}{18}} +\newlabel{table:dataparallel_index}{{5.1}{18}} +\newlabel{src:multiply_cpu}{{5.1}{18}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {5.1}Multiply(CPU)}{18}} +\@writefile{toc}{\contentsline {chapter}{\numberline {第6章}GPGPU への対応}{20}} \@writefile{lof}{\addvspace {10\p@ }} \@writefile{lot}{\addvspace {10\p@ }} -\@writefile{toc}{\contentsline {section}{\numberline {6.1}OpenCL および CUDA による実装}{18}} -\@writefile{toc}{\contentsline {section}{\numberline {6.2}データ並列}{19}} -\newlabel{src:multiply_opencl}{{6.1}{19}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.1}Multiply(OpenCL)}{19}} -\newlabel{src:multiply_cuda}{{6.2}{19}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.2}Multiply(CUDA)}{19}} -\newlabel{src:multiply_cpu}{{6.3}{20}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.3}Multiply(CPU)}{20}} -\@writefile{lot}{\contentsline {table}{\numberline {6.1}{\ignorespaces データ並列実行時の index の割り当て}}{20}} -\newlabel{table:dataparallel_index}{{6.1}{20}} -\@writefile{toc}{\contentsline {chapter}{\numberline {第7章}並列処理向けI/O}{22}} +\@writefile{toc}{\contentsline {section}{\numberline {6.1}OpenCL および CUDA による実装}{20}} +\@writefile{toc}{\contentsline {section}{\numberline {6.2}データ並列}{21}} +\newlabel{src:multiply_opencl}{{6.1}{21}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.1}Multiply(OpenCL)}{21}} +\newlabel{src:multiply_cuda}{{6.2}{21}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.2}Multiply(CUDA)}{21}} +\@writefile{toc}{\contentsline {chapter}{\numberline {第7章}並列処理向けI/O}{23}} \@writefile{lof}{\addvspace {10\p@ }} \@writefile{lot}{\addvspace {10\p@ }} -\@writefile{toc}{\contentsline {section}{\numberline {7.1}mmap}{22}} -\@writefile{lof}{\contentsline {figure}{\numberline {7.1}{\ignorespaces mmap の Model}}{22}} -\newlabel{fig:mmap}{{7.1}{22}} -\@writefile{toc}{\contentsline {section}{\numberline {7.2}Blocked Read による I/O の並列化}{23}} -\@writefile{lof}{\contentsline {figure}{\numberline {7.2}{\ignorespaces BlockedRead による WordCount}}{23}} -\newlabel{fig:blockedread}{{7.2}{23}} -\newlabel{src:blockedread_create}{{7.1}{24}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {7.1}BlockedRead を行う Task の生成}{24}} -\newlabel{src:blockedread_task}{{7.2}{24}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {7.2}BlockedRead Task}{24}} -\@writefile{toc}{\contentsline {section}{\numberline {7.3}I/O 専用 Thread の実装}{25}} -\@writefile{lof}{\contentsline {figure}{\numberline {7.3}{\ignorespaces BlockedRead と Task を同じ thread で動かした場合}}{25}} -\newlabel{fig:spe_any_blockedread}{{7.3}{25}} -\@writefile{lof}{\contentsline {figure}{\numberline {7.4}{\ignorespaces IO Thread による BlockedRead}}{25}} -\newlabel{fig:iothread__blockedread}{{7.4}{25}} -\@writefile{toc}{\contentsline {chapter}{\numberline {第8章}ベンチマーク}{26}} +\@writefile{toc}{\contentsline {section}{\numberline {7.1}mmap}{23}} +\@writefile{lof}{\contentsline {figure}{\numberline {7.1}{\ignorespaces mmap の Model}}{23}} +\newlabel{fig:mmap}{{7.1}{23}} +\@writefile{toc}{\contentsline {section}{\numberline {7.2}Blocked Read による I/O の並列化}{24}} +\@writefile{lof}{\contentsline {figure}{\numberline {7.2}{\ignorespaces BlockedRead による WordCount}}{24}} +\newlabel{fig:blockedread}{{7.2}{24}} +\newlabel{src:blockedread_create}{{7.1}{25}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {7.1}BlockedRead を行う Task の生成}{25}} +\newlabel{src:blockedread_task}{{7.2}{25}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {7.2}BlockedRead Task}{25}} +\@writefile{toc}{\contentsline {section}{\numberline {7.3}I/O 専用 Thread の実装}{26}} +\newlabel{sec:spe_problem}{{7.3}{26}} +\@writefile{lof}{\contentsline {figure}{\numberline {7.3}{\ignorespaces BlockedRead と Task を同じ thread で動かした場合}}{26}} +\newlabel{fig:spe_any_blockedread}{{7.3}{26}} +\@writefile{lof}{\contentsline {figure}{\numberline {7.4}{\ignorespaces IO Thread による BlockedRead}}{26}} +\newlabel{fig:iothread__blockedread}{{7.4}{26}} +\@writefile{toc}{\contentsline {chapter}{\numberline {第8章}ベンチマーク}{27}} \@writefile{lof}{\addvspace {10\p@ }} \@writefile{lot}{\addvspace {10\p@ }} -\@writefile{toc}{\contentsline {section}{\numberline {8.1}実験環境}{26}} -\@writefile{lot}{\contentsline {table}{\numberline {8.1}{\ignorespaces Ceriumを実行する実験環境1}}{26}} -\newlabel{tab:firefly_spec}{{8.1}{26}} -\@writefile{lot}{\contentsline {table}{\numberline {8.2}{\ignorespaces Ceriumを実行する実験環境2}}{26}} -\newlabel{tab:dragonfly_spec}{{8.2}{26}} -\@writefile{toc}{\contentsline {section}{\numberline {8.2}マルチコア}{27}} -\@writefile{lof}{\contentsline {figure}{\numberline {8.1}{\ignorespaces マルチコア CPU における Sort}}{27}} -\newlabel{fig:sort_on_multicore}{{8.1}{27}} -\@writefile{lof}{\contentsline {figure}{\numberline {8.2}{\ignorespaces マルチコア CPU における WordCount}}{28}} -\newlabel{fig:wordcount_on_multicore}{{8.2}{28}} -\@writefile{lof}{\contentsline {figure}{\numberline {8.3}{\ignorespaces Word Count による prefetch機能のベンチマーク}}{28}} -\newlabel{fig:prefetch_bench}{{8.3}{28}} -\@writefile{toc}{\contentsline {section}{\numberline {8.3}GPGPU}{29}} -\@writefile{lof}{\contentsline {figure}{\numberline {8.4}{\ignorespaces Word Count によるデータ並列実行のベンチマーク}}{29}} -\newlabel{fig:dataparallel}{{8.4}{29}} -\@writefile{lof}{\contentsline {figure}{\numberline {8.5}{\ignorespaces マルチコア CPU、OpenCL、CUDA における FFT}}{30}} -\newlabel{fig:fft_bench}{{8.5}{30}} -\@writefile{lof}{\contentsline {figure}{\numberline {8.6}{\ignorespaces MacPro 2013 における FFT}}{31}} -\newlabel{fig:fft_bench_dragonfly}{{8.6}{31}} -\@writefile{toc}{\contentsline {section}{\numberline {8.4}並列 I/O}{31}} -\@writefile{toc}{\contentsline {chapter}{\numberline {第9章}既存のプログラミングフレームワークとの比較}{32}} +\@writefile{toc}{\contentsline {section}{\numberline {8.1}実験環境}{27}} +\@writefile{lot}{\contentsline {table}{\numberline {8.1}{\ignorespaces Ceriumを実行する実験環境1}}{27}} +\newlabel{tab:firefly_spec}{{8.1}{27}} +\@writefile{lot}{\contentsline {table}{\numberline {8.2}{\ignorespaces Ceriumを実行する実験環境2}}{27}} +\newlabel{tab:dragonfly_spec}{{8.2}{27}} +\@writefile{toc}{\contentsline {section}{\numberline {8.2}マルチコア}{28}} +\@writefile{lof}{\contentsline {figure}{\numberline {8.1}{\ignorespaces マルチコア CPU における Sort}}{28}} +\newlabel{fig:sort_on_multicore}{{8.1}{28}} +\@writefile{lof}{\contentsline {figure}{\numberline {8.2}{\ignorespaces マルチコア CPU における WordCount}}{29}} +\newlabel{fig:wordcount_on_multicore}{{8.2}{29}} +\@writefile{lof}{\contentsline {figure}{\numberline {8.3}{\ignorespaces Word Count による prefetch機能のベンチマーク}}{29}} +\newlabel{fig:prefetch_bench}{{8.3}{29}} +\@writefile{toc}{\contentsline {section}{\numberline {8.3}GPGPU}{30}} +\@writefile{lof}{\contentsline {figure}{\numberline {8.4}{\ignorespaces Word Count によるデータ並列実行のベンチマーク}}{30}} +\newlabel{fig:dataparallel}{{8.4}{30}} +\@writefile{lof}{\contentsline {figure}{\numberline {8.5}{\ignorespaces マルチコア CPU、OpenCL、CUDA における FFT}}{31}} +\newlabel{fig:fft_bench}{{8.5}{31}} +\@writefile{lof}{\contentsline {figure}{\numberline {8.6}{\ignorespaces MacPro 2013 における FFT}}{32}} +\newlabel{fig:fft_bench_dragonfly}{{8.6}{32}} +\@writefile{toc}{\contentsline {section}{\numberline {8.4}並列 I/O}{32}} +\@writefile{lof}{\contentsline {figure}{\numberline {8.7}{\ignorespaces WordCount によるファイル読み込み方式のベンチマーク(MacPro2010)}}{33}} +\newlabel{fig:io_bench_firefly}{{8.7}{33}} +\@writefile{lof}{\contentsline {figure}{\numberline {8.8}{\ignorespaces WordCount によるファイル読み込み方式のベンチマーク(MacPro2013)}}{33}} +\newlabel{fig:io_bench_dragonfly}{{8.8}{33}} +\@writefile{toc}{\contentsline {chapter}{\numberline {第9章}既存のプログラミングフレームワークとの比較}{35}} \@writefile{lof}{\addvspace {10\p@ }} \@writefile{lot}{\addvspace {10\p@ }} -\@writefile{toc}{\contentsline {section}{\numberline {9.1}OpenCL}{32}} -\@writefile{lof}{\contentsline {figure}{\numberline {9.1}{\ignorespaces WorkItem ID}}{33}} -\newlabel{fig:workitem_id}{{9.1}{33}} -\@writefile{lot}{\contentsline {table}{\numberline {9.1}{\ignorespaces kernel で使用する ID 取得の API}}{33}} -\newlabel{table:kernel_id_api}{{9.1}{33}} -\@writefile{toc}{\contentsline {section}{\numberline {9.2}CUDA}{33}} -\@writefile{lof}{\contentsline {figure}{\numberline {9.2}{\ignorespaces Calculate Index example}}{34}} -\newlabel{fig:calculateIndex}{{9.2}{34}} -\@writefile{toc}{\contentsline {section}{\numberline {9.3}StarPU}{34}} -\@writefile{lof}{\contentsline {figure}{\numberline {9.3}{\ignorespaces StarPUにおけるデータ分割}}{35}} -\newlabel{fig:data_partition}{{9.3}{35}} -\@writefile{toc}{\contentsline {chapter}{\numberline {第10章}結論}{36}} +\@writefile{toc}{\contentsline {section}{\numberline {9.1}OpenCL}{35}} +\@writefile{lof}{\contentsline {figure}{\numberline {9.1}{\ignorespaces WorkItem ID}}{36}} +\newlabel{fig:workitem_id}{{9.1}{36}} +\@writefile{lot}{\contentsline {table}{\numberline {9.1}{\ignorespaces kernel で使用する ID 取得の API}}{36}} +\newlabel{table:kernel_id_api}{{9.1}{36}} +\@writefile{toc}{\contentsline {section}{\numberline {9.2}CUDA}{36}} +\@writefile{lof}{\contentsline {figure}{\numberline {9.2}{\ignorespaces Calculate Index example}}{37}} +\newlabel{fig:calculateIndex}{{9.2}{37}} +\@writefile{toc}{\contentsline {section}{\numberline {9.3}StarPU}{37}} +\@writefile{lof}{\contentsline {figure}{\numberline {9.3}{\ignorespaces StarPUにおけるデータ分割}}{38}} +\newlabel{fig:data_partition}{{9.3}{38}} +\@writefile{toc}{\contentsline {chapter}{\numberline {第10章}結論}{39}} \@writefile{lof}{\addvspace {10\p@ }} \@writefile{lot}{\addvspace {10\p@ }} -\newlabel{chapter:conclusion}{{10}{36}} -\@writefile{toc}{\contentsline {section}{\numberline {10.1}まとめ}{36}} -\@writefile{toc}{\contentsline {section}{\numberline {10.2}今後の課題}{36}} +\newlabel{chapter:conclusion}{{10}{39}} +\@writefile{toc}{\contentsline {section}{\numberline {10.1}まとめ}{39}} +\@writefile{toc}{\contentsline {section}{\numberline {10.2}今後の課題}{39}} \citation{*} \bibstyle{junsrt} \bibdata{master_paper} -\@writefile{toc}{\contentsline {chapter}{謝辞}{37}} +\@writefile{toc}{\contentsline {chapter}{謝辞}{40}} \bibcite{nobuyasu:2013a}{1} \bibcite{shoshi:2011a}{2} \bibcite{shoshi:2011b}{3} \bibcite{cassandra}{4} \bibcite{bigtable}{5} -\@writefile{toc}{\contentsline {chapter}{参考文献}{38}} -\@writefile{toc}{\contentsline {chapter}{発表文献}{39}} +\@writefile{toc}{\contentsline {chapter}{参考文献}{41}} +\@writefile{toc}{\contentsline {chapter}{発表文献}{42}} diff -r 9b181351ca4d -r 7956856211c5 paper/master_paper.dvi Binary file paper/master_paper.dvi has changed diff -r 9b181351ca4d -r 7956856211c5 paper/master_paper.lof --- a/paper/master_paper.lof Sun Feb 15 12:04:36 2015 +0900 +++ b/paper/master_paper.lof Sun Feb 15 20:12:33 2015 +0900 @@ -12,19 +12,21 @@ \addvspace {10\p@ } \addvspace {10\p@ } \addvspace {10\p@ } -\contentsline {figure}{\numberline {7.1}{\ignorespaces mmap の Model}}{22} -\contentsline {figure}{\numberline {7.2}{\ignorespaces BlockedRead による WordCount}}{23} -\contentsline {figure}{\numberline {7.3}{\ignorespaces BlockedRead と Task を同じ thread で動かした場合}}{25} -\contentsline {figure}{\numberline {7.4}{\ignorespaces IO Thread による BlockedRead}}{25} +\contentsline {figure}{\numberline {7.1}{\ignorespaces mmap の Model}}{23} +\contentsline {figure}{\numberline {7.2}{\ignorespaces BlockedRead による WordCount}}{24} +\contentsline {figure}{\numberline {7.3}{\ignorespaces BlockedRead と Task を同じ thread で動かした場合}}{26} +\contentsline {figure}{\numberline {7.4}{\ignorespaces IO Thread による BlockedRead}}{26} \addvspace {10\p@ } -\contentsline {figure}{\numberline {8.1}{\ignorespaces マルチコア CPU における Sort}}{27} -\contentsline {figure}{\numberline {8.2}{\ignorespaces マルチコア CPU における WordCount}}{28} -\contentsline {figure}{\numberline {8.3}{\ignorespaces Word Count による prefetch機能のベンチマーク}}{28} -\contentsline {figure}{\numberline {8.4}{\ignorespaces Word Count によるデータ並列実行のベンチマーク}}{29} -\contentsline {figure}{\numberline {8.5}{\ignorespaces マルチコア CPU、OpenCL、CUDA における FFT}}{30} -\contentsline {figure}{\numberline {8.6}{\ignorespaces MacPro 2013 における FFT}}{31} +\contentsline {figure}{\numberline {8.1}{\ignorespaces マルチコア CPU における Sort}}{28} +\contentsline {figure}{\numberline {8.2}{\ignorespaces マルチコア CPU における WordCount}}{29} +\contentsline {figure}{\numberline {8.3}{\ignorespaces Word Count による prefetch機能のベンチマーク}}{29} +\contentsline {figure}{\numberline {8.4}{\ignorespaces Word Count によるデータ並列実行のベンチマーク}}{30} +\contentsline {figure}{\numberline {8.5}{\ignorespaces マルチコア CPU、OpenCL、CUDA における FFT}}{31} +\contentsline {figure}{\numberline {8.6}{\ignorespaces MacPro 2013 における FFT}}{32} +\contentsline {figure}{\numberline {8.7}{\ignorespaces WordCount によるファイル読み込み方式のベンチマーク(MacPro2010)}}{33} +\contentsline {figure}{\numberline {8.8}{\ignorespaces WordCount によるファイル読み込み方式のベンチマーク(MacPro2013)}}{33} \addvspace {10\p@ } -\contentsline {figure}{\numberline {9.1}{\ignorespaces WorkItem ID}}{33} -\contentsline {figure}{\numberline {9.2}{\ignorespaces Calculate Index example}}{34} -\contentsline {figure}{\numberline {9.3}{\ignorespaces StarPUにおけるデータ分割}}{35} +\contentsline {figure}{\numberline {9.1}{\ignorespaces WorkItem ID}}{36} +\contentsline {figure}{\numberline {9.2}{\ignorespaces Calculate Index example}}{37} +\contentsline {figure}{\numberline {9.3}{\ignorespaces StarPUにおけるデータ分割}}{38} \addvspace {10\p@ } diff -r 9b181351ca4d -r 7956856211c5 paper/master_paper.log --- a/paper/master_paper.log Sun Feb 15 12:04:36 2015 +0900 +++ b/paper/master_paper.log Sun Feb 15 20:12:33 2015 +0900 @@ -1,4 +1,4 @@ -This is e-pTeX, Version 3.1415926-p3.4-110825-2.6 (utf8.euc) (TeX Live 2013) (format=platex 2013.5.30) 15 FEB 2015 12:03 +This is e-pTeX, Version 3.1415926-p3.4-110825-2.6 (utf8.euc) (TeX Live 2013) (format=platex 2013.5.30) 15 FEB 2015 20:10 entering extended mode \write18 enabled. %&-line parsing enabled. @@ -221,15 +221,15 @@ LaTeX Font Info: External font `cmex10' loaded for size (Font) <6> on input line 3. LaTeX Font Info: Font shape `JT1/mc/bx/n' in size <10> not available -(Font) Font shape `JT1/gt/m/n' tried instead on input line 29. +(Font) Font shape `JT1/gt/m/n' tried instead on input line 28. LaTeX Font Info: Font shape `JY1/mc/bx/n' in size <10> not available -(Font) Font shape `JY1/gt/m/n' tried instead on input line 29. +(Font) Font shape `JY1/gt/m/n' tried instead on input line 28. LaTeX Font Info: Font shape `JT1/mc/bx/n' in size <10.95> not available -(Font) Font shape `JT1/gt/m/n' tried instead on input line 29. +(Font) Font shape `JT1/gt/m/n' tried instead on input line 28. LaTeX Font Info: Font shape `JY1/mc/bx/n' in size <10.95> not available -(Font) Font shape `JY1/gt/m/n' tried instead on input line 29. +(Font) Font shape `JY1/gt/m/n' tried instead on input line 28. LaTeX Font Info: External font `cmex10' loaded for size -(Font) <10.95> on input line 29. +(Font) <10.95> on input line 28. [1 ]) @@ -357,54 +357,61 @@ ] 第 5 章(17ページ) -) (./chapter5.tex [17 - -] -第 6 章(18ページ) -[18 +[17 ] File: images/emblem-bitmap.eps Graphic file (type eps) - [19]) (./chapter6.tex + [18]) (./chapter5.tex File: images/emblem-bitmap.eps Graphic file (type eps) - [20] + [19] +第 6 章(20ページ) +[20 + +] File: images/emblem-bitmap.eps Graphic file (type eps) [21] -第 7 章(22ページ) + +LaTeX Warning: Reference `src:multicore_cpu' on page 22 undefined on input line + 94. + +) (./chapter6.tex +File: images/emblem-bitmap.eps Graphic file (type eps) + [22] +第 7 章(23ページ) File: ./images/mmap.pdf Graphic file (type pdf) -<./images/mmap.pdf> [22 +<./images/mmap.pdf> [23 ] File: ./images/blockedread.pdf Graphic file (type pdf) <./images/blockedread.pdf> File: images/emblem-bitmap.eps Graphic file (type eps) -[23] +[24] File: images/emblem-bitmap.eps Graphic file (type eps) - [24] + [25] File: ./images/speblockedread.pdf Graphic file (type pdf) <./images/speblockedread.pdf> -Overfull \hbox (2.16278pt too wide) in paragraph at lines 118--119 +Overfull \hbox (2.16278pt too wide) in paragraph at lines 122--123 [] [] -LaTeX Warning: Reference `fig:iothread_blockedread' on page 25 undefined on inp -ut line 128. +LaTeX Warning: Reference `fig:iothread_blockedread' on page 26 undefined on inp +ut line 132. File: ./images/iothread.pdf Graphic file (type pdf) <./images/iothread.pdf> -Overfull \hbox (47.1306pt too wide) in paragraph at lines 132--133 +Overfull \hbox (47.1306pt too wide) in paragraph at lines 136--137 [] [] ) (./chapter8.tex File: images/emblem-bitmap.eps Graphic file (type eps) - [25] -第 8 章(26ページ) + [26] +第 8 章(27ページ) File: ./figures/multicore/sort.pdf Graphic file (type pdf) -<./figures/multicore/sort.pdf> [26 +<./figures/multicore/sort.pdf> [27 ] File: ./figures/multicore/word_count.pdf Graphic file (type pdf) @@ -413,32 +420,42 @@ <./figures/dma/dmabench.pdf> File: images/emblem-bitmap.eps Graphic file (type eps) - [27] + [28] File: images/emblem-bitmap.eps Graphic file (type eps) - [28] + [29] File: ./figures/GPU/wordcount_dataparallel.pdf Graphic file (type pdf) <./figures/GPU/wordcount_dataparallel.pdf> File: images/emblem-bitmap.eps Graphic file (type eps) - [29] + [30] File: ./figures/GPU/fft_firefly.pdf Graphic file (type pdf) <./figures/GPU/fft_firefly.pdf> File: ./figures/GPU/fft_dragonfly.pdf Graphic file (type pdf) <./figures/GPU/fft_dragonfly.pdf> File: images/emblem-bitmap.eps Graphic file (type eps) - [30]) + [31] +File: ./figures/io/io_thread_firefly.pdf Graphic file (type pdf) + +<./figures/io/io_thread_firefly.pdf> +File: ./figures/io/io_thread_dragonfly.pdf Graphic file (type pdf) + <./figures/io/io_thread_dragonfly.pdf>) (./chapter9.tex File: images/emblem-bitmap.eps Graphic file (type eps) - [31] -第 9 章(32ページ) + [32] +File: images/emblem-bitmap.eps Graphic file (type eps) + +[33] +File: images/emblem-bitmap.eps Graphic file (type eps) + [34] +第 9 章(35ページ) File: ./images/workitem.pdf Graphic file (type pdf) -<./images/workitem.pdf> [32 +<./images/workitem.pdf> [35 ] File: images/emblem-bitmap.eps Graphic file (type eps) - [33] + [36] File: ./images/calculateIndex.pdf Graphic file (type pdf) <./images/calculateIndex.pdf> @@ -446,20 +463,20 @@ <./images/starpu_data_parallel.pdf>) (./conclusion.tex File: images/emblem-bitmap.eps Graphic file (type eps) - [34] + [37] File: images/emblem-bitmap.eps Graphic file (type eps) -[35] -第 10 章(36ページ) -) (./thanx.tex [36 +[38] +第 10 章(39ページ) +) (./thanx.tex [39 -]) (./master_paper.bbl [37 +]) (./master_paper.bbl [40 -]) (./appendix.tex [38 +]) (./appendix.tex [41 ]) No file master_paper.ind. -[39 +[42 ] (./master_paper.aux) @@ -470,12 +487,12 @@ ) Here is how much of TeX's memory you used: - 2536 strings out of 494008 - 35926 string characters out of 6154472 + 2546 strings out of 494008 + 36163 string characters out of 6154472 189780 words of memory out of 5000000 - 5902 multiletter control sequences out of 15000+600000 + 5909 multiletter control sequences out of 15000+600000 18223 words of font info for 71 fonts, out of 8000000 for 9000 745 hyphenation exceptions out of 8191 - 33i,12n,40p,207b,1954s stack positions out of 5000i,500n,10000p,200000b,80000s + 30i,12n,40p,213b,1954s stack positions out of 5000i,500n,10000p,200000b,80000s -Output written on master_paper.dvi (46 pages, 132264 bytes). +Output written on master_paper.dvi (49 pages, 140856 bytes). diff -r 9b181351ca4d -r 7956856211c5 paper/master_paper.lot --- a/paper/master_paper.lot Sun Feb 15 12:04:36 2015 +0900 +++ b/paper/master_paper.lot Sun Feb 15 20:12:33 2015 +0900 @@ -5,12 +5,12 @@ \contentsline {table}{\numberline {3.2}{\ignorespaces Task 側で使用する API}}{9} \addvspace {10\p@ } \addvspace {10\p@ } +\contentsline {table}{\numberline {5.1}{\ignorespaces データ並列実行時の index の割り当て}}{18} \addvspace {10\p@ } -\contentsline {table}{\numberline {6.1}{\ignorespaces データ並列実行時の index の割り当て}}{20} \addvspace {10\p@ } \addvspace {10\p@ } -\contentsline {table}{\numberline {8.1}{\ignorespaces Ceriumを実行する実験環境1}}{26} -\contentsline {table}{\numberline {8.2}{\ignorespaces Ceriumを実行する実験環境2}}{26} +\contentsline {table}{\numberline {8.1}{\ignorespaces Ceriumを実行する実験環境1}}{27} +\contentsline {table}{\numberline {8.2}{\ignorespaces Ceriumを実行する実験環境2}}{27} \addvspace {10\p@ } -\contentsline {table}{\numberline {9.1}{\ignorespaces kernel で使用する ID 取得の API}}{33} +\contentsline {table}{\numberline {9.1}{\ignorespaces kernel で使用する ID 取得の API}}{36} \addvspace {10\p@ } diff -r 9b181351ca4d -r 7956856211c5 paper/master_paper.pdf Binary file paper/master_paper.pdf has changed diff -r 9b181351ca4d -r 7956856211c5 paper/master_paper.toc --- a/paper/master_paper.toc Sun Feb 15 12:04:36 2015 +0900 +++ b/paper/master_paper.toc Sun Feb 15 20:12:33 2015 +0900 @@ -17,25 +17,26 @@ \contentsline {chapter}{\numberline {第5章}マルチコアへの対応}{17} \contentsline {section}{\numberline {5.1}マルチコア上での実行の機構}{17} \contentsline {section}{\numberline {5.2}DMA}{17} -\contentsline {chapter}{\numberline {第6章}GPGPU への対応}{18} -\contentsline {section}{\numberline {6.1}OpenCL および CUDA による実装}{18} -\contentsline {section}{\numberline {6.2}データ並列}{19} -\contentsline {chapter}{\numberline {第7章}並列処理向けI/O}{22} -\contentsline {section}{\numberline {7.1}mmap}{22} -\contentsline {section}{\numberline {7.2}Blocked Read による I/O の並列化}{23} -\contentsline {section}{\numberline {7.3}I/O 専用 Thread の実装}{25} -\contentsline {chapter}{\numberline {第8章}ベンチマーク}{26} -\contentsline {section}{\numberline {8.1}実験環境}{26} -\contentsline {section}{\numberline {8.2}マルチコア}{27} -\contentsline {section}{\numberline {8.3}GPGPU}{29} -\contentsline {section}{\numberline {8.4}並列 I/O}{31} -\contentsline {chapter}{\numberline {第9章}既存のプログラミングフレームワークとの比較}{32} -\contentsline {section}{\numberline {9.1}OpenCL}{32} -\contentsline {section}{\numberline {9.2}CUDA}{33} -\contentsline {section}{\numberline {9.3}StarPU}{34} -\contentsline {chapter}{\numberline {第10章}結論}{36} -\contentsline {section}{\numberline {10.1}まとめ}{36} -\contentsline {section}{\numberline {10.2}今後の課題}{36} -\contentsline {chapter}{謝辞}{37} -\contentsline {chapter}{参考文献}{38} -\contentsline {chapter}{発表文献}{39} +\contentsline {section}{\numberline {5.3}データ並列}{17} +\contentsline {chapter}{\numberline {第6章}GPGPU への対応}{20} +\contentsline {section}{\numberline {6.1}OpenCL および CUDA による実装}{20} +\contentsline {section}{\numberline {6.2}データ並列}{21} +\contentsline {chapter}{\numberline {第7章}並列処理向けI/O}{23} +\contentsline {section}{\numberline {7.1}mmap}{23} +\contentsline {section}{\numberline {7.2}Blocked Read による I/O の並列化}{24} +\contentsline {section}{\numberline {7.3}I/O 専用 Thread の実装}{26} +\contentsline {chapter}{\numberline {第8章}ベンチマーク}{27} +\contentsline {section}{\numberline {8.1}実験環境}{27} +\contentsline {section}{\numberline {8.2}マルチコア}{28} +\contentsline {section}{\numberline {8.3}GPGPU}{30} +\contentsline {section}{\numberline {8.4}並列 I/O}{32} +\contentsline {chapter}{\numberline {第9章}既存のプログラミングフレームワークとの比較}{35} +\contentsline {section}{\numberline {9.1}OpenCL}{35} +\contentsline {section}{\numberline {9.2}CUDA}{36} +\contentsline {section}{\numberline {9.3}StarPU}{37} +\contentsline {chapter}{\numberline {第10章}結論}{39} +\contentsline {section}{\numberline {10.1}まとめ}{39} +\contentsline {section}{\numberline {10.2}今後の課題}{39} +\contentsline {chapter}{謝辞}{40} +\contentsline {chapter}{参考文献}{41} +\contentsline {chapter}{発表文献}{42}