Mercurial > hg > Papers > 2015 > yuhi-master
changeset 15:712576635154
gpgpu
| author | Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp> |
|---|---|
| date | Mon, 09 Feb 2015 11:32:28 +0900 |
| parents | 65a4b30a5a53 |
| children | d7cf4a51597f |
| files | paper/chapter1.tex paper/chapter2.tex paper/chapter3.tex paper/chapter4.tex paper/chapter5.tex paper/chapter6.tex paper/chapter7.tex paper/master_paper.aux paper/master_paper.dvi paper/master_paper.lof paper/master_paper.log paper/master_paper.lot paper/master_paper.pdf paper/master_paper.tex paper/master_paper.toc |
| diffstat | 15 files changed, 386 insertions(+), 241 deletions(-) [+] |
line wrap: on
line diff
--- a/paper/chapter1.tex Wed Feb 04 16:38:20 2015 +0900 +++ b/paper/chapter1.tex Mon Feb 09 11:32:28 2015 +0900 @@ -1,56 +1,83 @@ \chapter{既存のマルチプラットフォームフレームワーク} +マルチプラットフォームでプログラムを動作させる場合、そのアーキテクチャを意識する必要がある。 +マルチプラットフォームにはマルチコア CPU 、 GPU や Cell といったヘテロジニアスマルチコアのような +様々な構成がある。 + +\section{Architecutre} +本研究では、 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 で記述したプログラムを演算用プロセッサ上で実行させるため、 -制御用のプロセッサが利用する API である。 - -OpenCL では演算用プロセッサ側を device 、制御用デバイス側を host として定義する。 -また、 Device 上で動作するプログラムの事を kernel と呼ぶ。 - -OpenCL では、デバイスの操作に Command Queue を使用する。 Command Queue は Device に命令を送るための仕組みである。 - Command Queue は clCreateCommandQueue という OpenCL API で作成され、 - Command Queueが所属するコンテキストや実行対象となるデバイスを指定する。 - -kernel の実行、input data への書き込み、 output data の読み込みといった -メモリ操作はこの Command Queue を通して行われる。 - -host では主に data を input/output するメモリ資源の確保を行う。 -GPU のメモリ空間(図:\ref{fig:gpuarch})はマルチコア CPU (図:\ref{fig:cpuarch})と違い、 -共有メモリでないため host と kernel(Task)間で data の共有ができない。 -\begin{figure}[htpb] - \begin{center} - \includegraphics[scale=0.4]{./images/gpu_arch.pdf} - \end{center} - \caption{Gpu Architecture} - \label{fig:gpuarch} -\end{figure} +OpenCL C は演算用プロセッサ上で動作する、 C 言語を拡張したプログラミング言語である。 +一方で OpenCL Runtime API は OpenCL C で記述した kernel を Queuing するために Host が利用する API である。 -\begin{figure}[htpb] - \begin{center} - \includegraphics[scale=0.8]{./images/cpu_arch.pdf} - \end{center} - \caption{Cpu Architecture} - \label{fig:cpuarch} -\end{figure} -アクセスするにはメモリ空間間でコピーしなければならない。 - -OpenCLは host 側で memory buffer を作成してメモリのコピーを行う。 +Host では主に Data を input/output するメモリ資源の確保を行う。 +OpenCL は host 側で memory buffer を作成してメモリのコピーを行う。 これらの処理や Task は Command Queue に enqueue することで実行される。 多次元のデータ構造を扱う計算において高い並列度を保つには、 -それを分割して並列に実行する機能が必要である。 -データ並列実行という。OpenCLはデータ並列実行もサポートしている。 -OpenCL は次元数に対応する index があり、 +多次元データを分割して並列に実行する機能が必要である。 +これをデータ並列実行という。 OpenCL はデータ並列実行もサポートしている。 + OpenCL は次元数に対応する index があり、 OpenCL は一つの記述から異なる index を持つ複数の kernel を自動生成する。 その添字を global\_id と呼ぶ。この時入力されたデータはワークアイテムという処理単位に分割される。 @@ -115,9 +142,22 @@ 例えば 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} @@ -126,19 +166,11 @@ \item CUDA Driver API \end{itemize} CUDA C は GPU 上で動作する、C 言語を拡張したプログラミング言語である。 -CUDA Runtime API も CUDA Driver APIも CUDA C で記述したプログラムを GPU 上で実行させるために -制御用プロセッサが利用するAPIである。 +CUDA Runtime API も CUDA Driver API も CUDA C で記述した Kernel を Queueing するために + Host が利用するAPIである。 Driver API は Runtime APIに比べ、プログラマが管理しなければならないリソースが多くなる代わり、 より柔軟な処理を行う事ができる。 -CUDA も OpenCL と同様、演算用プロセッサ( GPU )を Device 、制御用デバイス側を Host として定義する。 -また、 Device 上で動作するプログラムの事も kernel と呼ぶ。 - -OpenCL における Command 、 CommandQueue に対応するものとして、 CUDA には Operation と Stream がある。 -Stream は Host 側で発行された Operation を一連の動作として Device で実行する。 -Stream に発行された Operation は発行された順序で実行されることが保証されている。 -更に、異なる Stream に発行された Operation も依存関係が存在しない場合、Operationは並列に実行される。 - Stream は cuStreamCreate という Driver API で生成される。 引数に Stream を指定しない API は全て host 側をブロックする同期的な処理となる。 複数の Stream を同時に走らせ、 Operation を並列に実行するためには非同期的な処理を行う API を利用する必要がある。
--- a/paper/chapter2.tex Wed Feb 04 16:38:20 2015 +0900 +++ b/paper/chapter2.tex Mon Feb 09 11:32:28 2015 +0900 @@ -42,6 +42,12 @@ \end{figure} \section{Task の Scheduling} + +GPU や Cell のような Shared Memory でない環境でのプログラミングを行う場合、 +Task の入出力となるデータを転送し、転送が終わってから Task を起動しなければならない。 +転送処理がボトルネックとなり、並列度が低下してしまう。 +そのため、Cerium はパイプライン実行をサポートしている。 + Scheduler に転送された Task はパイプラインで処理される(図:\ref{fig:scheduler})。 Task が全て終了すると Scheduler から TaskManager に mail を通して通知される。 通知に従い依存関係を解決した Task が再び TaskManager から Scheduler に転送される。
--- a/paper/chapter3.tex Wed Feb 04 16:38:20 2015 +0900 +++ b/paper/chapter3.tex Mon Feb 09 11:32:28 2015 +0900 @@ -40,7 +40,7 @@ \begin{figure}[htpb] \begin{center} - \includegraphics[scale=0.7]{./images/sort.pdf} + \includegraphics[scale=0.7]{./images/sort_benchmark.pdf} \end{center} \caption{Bitonic Sort の例} \label{fig:sort}
--- a/paper/chapter4.tex Wed Feb 04 16:38:20 2015 +0900 +++ b/paper/chapter4.tex Mon Feb 09 11:32:28 2015 +0900 @@ -3,7 +3,9 @@ \section{マルチコア上での実行の機構} Cell には MailBox という機能がある。 - MailBox は FIFO キュー構造になっており、Device と Host で双方向のデータの受け渡しが可能である。 +Cell は Shared Memory でないので、 Memory に直接アクセスできない。 +そこで MailBox を用いて双方向のデータの受け渡しを可能にしている。 + MailBox は FIFO キュー構造になっており、Device と Host で この MailBox に対応させる形で Synchronized Queue を用いて MacOSX 、 Linux 用の TaskManager へ MailBox を移植した。 @@ -11,49 +13,11 @@ 各スレッドは Input 用と Output 用として Synchronized Queue を2つ持っており、管理スレッドから Task を受けて 並列に実行する。 -また、 Cell ではデータの受け渡しとして MailBox 以外に DMA 転送を使用する方法がある。 -CPU の代わりに DMA Controller を介してデータ転送を行う方式である。 -Cell は Device と Host でメモリ空間が異なるために DMA 転送を用いている。 -しかしマルチコア上で実行する場合は各 CPU で同じメモリ空間を利用できる。 -よって DMA 転送を用いていた部分をポインタ渡しを行うように修正し、速度の向上を図った。 -\section{ベンチマーク} -今回使用する実験環境を表:\ref{tab:firefly_spec}、表:\ref{tab:dragonfly_spec}に示す. +\section{DMA} +Cell ではデータの受け渡しとして MailBox 以外に DMA 転送を使用する方法がある。 +CPU を介さずに周辺装置とメモリ間でデータ転送を行う方式である。 -\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 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 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 を使用している。 - +Cerium も DMA 転送を用いている箇所がある。 +しかしマルチコア CPU 上で実行する場合は各 CPU で同じメモリ空間を利用できる。 +よって DMA 転送を用いていた部分をポインタ渡しを行うように修正し、 +メモリに直接アクセスさせることで速度の向上が見込める。
--- a/paper/chapter5.tex Wed Feb 04 16:38:20 2015 +0900 +++ b/paper/chapter5.tex Mon Feb 09 11:32:28 2015 +0900 @@ -3,21 +3,154 @@ 現在、GPU のような異なる種類のアーキテクチャを搭載した CPU 、つまりヘテロジニアスな CPU が増えている。 特定の計算に特化した Task の生成やスケジューリングを行い、 GPGPU により高い並列度を出す研究は様々な分野で行われている。 本研究では Cerium を特定の計算に限らず、 GPU を用いて汎用計算を行えるフレームワークに改良する。 -\section{OpenCL による実装} -\subsection{ベンチマーク} -OpenCL を用いた GPU 対応を行った。 -Scheduler と CpuThreads に対応させる形で GpuScheduler と GpuThreads を実装した。 -GpuScheduler 内で OpenCL の API を用いて GPU の制御を行っている。 - -受け取った TaskList からメモリバッファを作成し、 -clEnqueueuWriteBuffer、clEnqueueTask、clEnqueueReadBuffer の順に CommandQueue に Enqueueしていく。 -Task の投入は CommandQueue を2つ用意しパイプライン的に実行を行う。 - -Task の終了は、clWaitForEvent によって検出し、TaskManager 間の通信を担当する同期キューである mail を使って通知する。 - -\section{CUDA による} -\subsection{ベンチマーク} +\section{OpenCL および CUDA による実装} +OpenCL 、CUDA による GPGPU 対応を行った。 +Scheduler と CpuThreads に対応させる形で +OpenCL を用いた GpuScheduler と GpuThreads、 +CUDA を用いた CudaScheduler と CudaThreads を実装した。 +それぞれの Scheduler 内で各フレームワークの API を用いて GPU の制御を行っている。 + + TaskManager から受け取った TaskList をもとに Device 上のメモリバッファを作成する。 +その後 CommandQueue、 Stream といったそれぞれの Queue に Device 制御用の Command を Queueing していく。 + +Command は Queueing した順に実行されるので、以下のように Command を Queueing する。 +\begin{enumerate} +\item Host から Device へのデータ転送 +\item kernel の実行 +\item Device から Host へのデータ転送 +\end{enumerate} + +データの転送や kernel の実行は非同期 API を用いることで並列に行うことができる。 + +通常、フレームワークが依存関係を解決して実行するが、 +非同期 API を用いる場合はユーザが依存関係を考慮する必要がある。 +しかし Task の依存関係は TaskManager が既に解決した状態で送ってくるので、 + Scheduler は依存関係を考慮せずに実行して問題ない。 + +GPGPU 用の Scheduler は CommandQueue を2つ持っており、Task をパイプライン的に実行する。 + +転送されてきた Task が全て終了すると、 +TaskManager 間の通信を担当する同期キューである mail を通して TaskManager に Task の終了を通知する。 +終了が通知されると TaskManager でその TaskList に関する依存関係が解消され、 + +GPGPU の Scheduler 内で Platform や Device ID の取得、 Context の生成、 Kernel の Build と Load等も行っており、OD +並列処理したい計算のみに集中できる。 + \section{データ並列} -\subsection{ベンチマーク} +並列プログラミングにおいて、明示的な並列化部分はループ部分である。 + GPU は数百個のコアを有しており、ループ部分に対してデータ並列で処理を行うことで CPU より高速に演算を行う事ができる。 +プログラムの大部分がループであれば、データ並列による実行だけでプログラムの性能は向上する。 +OpenCL 、 CUDA ともにデータ並列をサポートしている。 +Task を実行する際にデータをどう分割するか指定し、kernel をデータ並列実行用に書き換えることで実現する。 +データ並列実行用の kernel は以下のように記述する。 +2つの input データの積を output データに格納する例題、 multiply を用いる。 + +\begin{lstlisting}[frame=lrbt,label=src:multiply_opencl,caption=Multiply(OpenCL),numbers=left] +__kernel void +multiply(__global const long *params, + __global const float *input1, + __global const float *input2, + __global const float *output) { + + long id = get_global_id(0); + + output[id] = input1[id] * input2[id]; +} +\end{lstlisting} + +\begin{lstlisting}[frame=lrbt,label=src:multiply_cuda,caption=Multiply(CUDA),numbers=left] +__global__ void +multiply(__global const long *params, + __global const float *input1, + __global const float *input2, + __global const float *output) { + + int id = blockIdx.x * blockDim.x + threadIdx.x; + + output[id] = input1[id] * input2[id]; +} +\end{lstlisting} + +このような kernel を分割数分生成する。 +分割数は kernel の生成時にそれぞれのフレームワークが 用意している API を用いて指定する。 +いずれの kernel も +\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 を記述することができる。 +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}節)なため、 +コピーによるオーバーヘッドは少ない。 + +しかし GPU は SharedMemory ではなく、データの転送がオーバーヘッドとなる。 +コピーを減らす事で並列度の向上が見込める。
--- a/paper/chapter6.tex Wed Feb 04 16:38:20 2015 +0900 +++ b/paper/chapter6.tex Mon Feb 09 11:32:28 2015 +0900 @@ -1,4 +1,2 @@ - \chapter{並列処理向けI/O} \section{新たに実装したI/Oの機構} -\section{ベンチマーク}
--- a/paper/chapter7.tex Wed Feb 04 16:38:20 2015 +0900 +++ b/paper/chapter7.tex Mon Feb 09 11:32:28 2015 +0900 @@ -1,5 +1,3 @@ - \chapter{Memory Allocator} \section{現状のMemory Allocator} \section{新しいMemory Allocator} -\section{ベンチマーク}
--- a/paper/master_paper.aux Wed Feb 04 16:38:20 2015 +0900 +++ b/paper/master_paper.aux Mon Feb 09 11:32:28 2015 +0900 @@ -6,19 +6,22 @@ \@writefile{toc}{\contentsline {chapter}{\numberline {第2章}既存のマルチプラットフォームフレームワーク}{3}} \@writefile{lof}{\addvspace {10\p@ }} \@writefile{lot}{\addvspace {10\p@ }} -\@writefile{toc}{\contentsline {section}{\numberline {2.1}OpenCL}{3}} -\@writefile{lof}{\contentsline {figure}{\numberline {2.1}{\ignorespaces Gpu Architecture}}{4}} +\@writefile{toc}{\contentsline {section}{\numberline {2.1}Architecutre}{3}} +\@writefile{toc}{\contentsline {section}{\numberline {2.2}Shared Memory}{3}} +\newlabel{sec:shared_memory}{{2.2}{3}} +\@writefile{lof}{\contentsline {figure}{\numberline {2.1}{\ignorespaces GPU Architecture}}{4}} \newlabel{fig:gpuarch}{{2.1}{4}} -\@writefile{lof}{\contentsline {figure}{\numberline {2.2}{\ignorespaces Cpu Architecture}}{4}} +\@writefile{lof}{\contentsline {figure}{\numberline {2.2}{\ignorespaces CPU Architecture}}{4}} \newlabel{fig:cpuarch}{{2.2}{4}} -\@writefile{lof}{\contentsline {figure}{\numberline {2.3}{\ignorespaces WorkItem ID}}{5}} -\newlabel{fig:workitem_id}{{2.3}{5}} +\@writefile{toc}{\contentsline {section}{\numberline {2.3}OpenCL}{5}} +\@writefile{lof}{\contentsline {figure}{\numberline {2.3}{\ignorespaces WorkItem ID}}{6}} +\newlabel{fig:workitem_id}{{2.3}{6}} \@writefile{lot}{\contentsline {table}{\numberline {2.1}{\ignorespaces kernel で使用する ID 取得の API}}{6}} \newlabel{table:kernel_id_api}{{2.1}{6}} -\@writefile{toc}{\contentsline {section}{\numberline {2.2}CUDA}{6}} -\@writefile{lof}{\contentsline {figure}{\numberline {2.4}{\ignorespaces Calculate Index example}}{7}} -\newlabel{fig:calculateIndex}{{2.4}{7}} -\@writefile{toc}{\contentsline {section}{\numberline {2.3}StarPU}{7}} +\@writefile{toc}{\contentsline {section}{\numberline {2.4}CUDA}{7}} +\@writefile{lof}{\contentsline {figure}{\numberline {2.4}{\ignorespaces Calculate Index example}}{8}} +\newlabel{fig:calculateIndex}{{2.4}{8}} +\@writefile{toc}{\contentsline {section}{\numberline {2.5}StarPU}{8}} \newlabel{src:codelet}{{2.1}{8}} \@writefile{lol}{\contentsline {lstlisting}{\numberline {2.1}codeletの例}{8}} \@writefile{lof}{\contentsline {figure}{\numberline {2.5}{\ignorespaces StarPUにおけるデータ分割}}{9}} @@ -57,41 +60,46 @@ \@writefile{lof}{\addvspace {10\p@ }} \@writefile{lot}{\addvspace {10\p@ }} \@writefile{toc}{\contentsline {section}{\numberline {5.1}マルチコア上での実行の機構}{19}} -\@writefile{toc}{\contentsline {section}{\numberline {5.2}ベンチマーク}{19}} -\@writefile{lot}{\contentsline {table}{\numberline {5.1}{\ignorespaces Ceriumを実行する実験環境1}}{19}} -\newlabel{tab:firefly_spec}{{5.1}{19}} -\@writefile{lot}{\contentsline {table}{\numberline {5.2}{\ignorespaces Ceriumを実行する実験環境2}}{20}} -\newlabel{tab:dragonfly_spec}{{5.2}{20}} -\@writefile{toc}{\contentsline {chapter}{\numberline {第6章}GPGPU への対応}{21}} +\@writefile{toc}{\contentsline {section}{\numberline {5.2}DMA}{19}} +\@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 による実装}{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}} +\newlabel{src:multiply_cpu}{{6.3}{22}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.3}Multiply(CPU)}{22}} +\@writefile{lot}{\contentsline {table}{\numberline {6.1}{\ignorespaces データ並列実行時の index の割り当て}}{22}} +\newlabel{table:dataparallel_index}{{6.1}{22}} +\@writefile{toc}{\contentsline {chapter}{\numberline {第7章}並列処理向けI/O}{24}} \@writefile{lof}{\addvspace {10\p@ }} \@writefile{lot}{\addvspace {10\p@ }} -\@writefile{toc}{\contentsline {section}{\numberline {6.1}OpenCL による実装}{21}} -\@writefile{toc}{\contentsline {subsection}{\numberline {6.1.1}ベンチマーク}{21}} -\@writefile{toc}{\contentsline {section}{\numberline {6.2}CUDA による}{21}} -\@writefile{toc}{\contentsline {subsection}{\numberline {6.2.1}ベンチマーク}{21}} -\@writefile{toc}{\contentsline {section}{\numberline {6.3}データ並列}{21}} -\@writefile{toc}{\contentsline {subsection}{\numberline {6.3.1}ベンチマーク}{21}} -\@writefile{toc}{\contentsline {chapter}{\numberline {第7章}並列処理向けI/O}{22}} +\@writefile{toc}{\contentsline {section}{\numberline {7.1}新たに実装したI/Oの機構}{24}} +\@writefile{toc}{\contentsline {chapter}{\numberline {第8章}Memory Allocator}{25}} +\@writefile{lof}{\addvspace {10\p@ }} +\@writefile{lot}{\addvspace {10\p@ }} +\@writefile{toc}{\contentsline {section}{\numberline {8.1}現状のMemory Allocator}{25}} +\@writefile{toc}{\contentsline {section}{\numberline {8.2}新しいMemory Allocator}{25}} +\@writefile{toc}{\contentsline {chapter}{\numberline {第9章}ベンチマーク}{26}} \@writefile{lof}{\addvspace {10\p@ }} \@writefile{lot}{\addvspace {10\p@ }} -\@writefile{toc}{\contentsline {section}{\numberline {7.1}新たに実装したI/Oの機構}{22}} -\@writefile{toc}{\contentsline {section}{\numberline {7.2}ベンチマーク}{22}} -\@writefile{toc}{\contentsline {chapter}{\numberline {第8章}Memory Allocator}{23}} +\@writefile{lot}{\contentsline {table}{\numberline {9.1}{\ignorespaces Ceriumを実行する実験環境1}}{26}} +\newlabel{tab:firefly_spec}{{9.1}{26}} +\@writefile{lot}{\contentsline {table}{\numberline {9.2}{\ignorespaces Ceriumを実行する実験環境2}}{26}} +\newlabel{tab:dragonfly_spec}{{9.2}{26}} +\@writefile{toc}{\contentsline {chapter}{\numberline {第10章}結論}{27}} \@writefile{lof}{\addvspace {10\p@ }} \@writefile{lot}{\addvspace {10\p@ }} -\@writefile{toc}{\contentsline {section}{\numberline {8.1}現状のMemory Allocator}{23}} -\@writefile{toc}{\contentsline {section}{\numberline {8.2}新しいMemory Allocator}{23}} -\@writefile{toc}{\contentsline {section}{\numberline {8.3}ベンチマーク}{23}} -\@writefile{toc}{\contentsline {chapter}{\numberline {第9章}結論}{24}} -\@writefile{lof}{\addvspace {10\p@ }} -\@writefile{lot}{\addvspace {10\p@ }} -\newlabel{chapter:conclusion}{{9}{24}} -\@writefile{toc}{\contentsline {section}{\numberline {9.1}まとめ}{24}} -\@writefile{toc}{\contentsline {section}{\numberline {9.2}今後の課題}{24}} +\newlabel{chapter:conclusion}{{10}{27}} +\@writefile{toc}{\contentsline {section}{\numberline {10.1}まとめ}{27}} +\@writefile{toc}{\contentsline {section}{\numberline {10.2}今後の課題}{27}} \citation{*} \bibstyle{junsrt} \bibdata{master_paper} -\@writefile{toc}{\contentsline {chapter}{謝辞}{25}} +\@writefile{toc}{\contentsline {chapter}{謝辞}{28}} \bibcite{msgpack:2013}{1} \bibcite{nobuyasu:2013a}{2} \bibcite{shoshi:2011a}{3} @@ -102,5 +110,5 @@ \bibcite{dynamo}{8} \bibcite{deos2013}{9} \bibcite{d_add2013}{10} -\@writefile{toc}{\contentsline {chapter}{参考文献}{26}} -\@writefile{toc}{\contentsline {chapter}{発表文献}{27}} +\@writefile{toc}{\contentsline {chapter}{参考文献}{29}} +\@writefile{toc}{\contentsline {chapter}{発表文献}{30}}
--- a/paper/master_paper.lof Wed Feb 04 16:38:20 2015 +0900 +++ b/paper/master_paper.lof Mon Feb 09 11:32:28 2015 +0900 @@ -1,9 +1,9 @@ \addvspace {10\p@ } \addvspace {10\p@ } -\contentsline {figure}{\numberline {2.1}{\ignorespaces Gpu Architecture}}{4} -\contentsline {figure}{\numberline {2.2}{\ignorespaces Cpu Architecture}}{4} -\contentsline {figure}{\numberline {2.3}{\ignorespaces WorkItem ID}}{5} -\contentsline {figure}{\numberline {2.4}{\ignorespaces Calculate Index example}}{7} +\contentsline {figure}{\numberline {2.1}{\ignorespaces GPU Architecture}}{4} +\contentsline {figure}{\numberline {2.2}{\ignorespaces CPU Architecture}}{4} +\contentsline {figure}{\numberline {2.3}{\ignorespaces WorkItem ID}}{6} +\contentsline {figure}{\numberline {2.4}{\ignorespaces Calculate Index example}}{8} \contentsline {figure}{\numberline {2.5}{\ignorespaces StarPUにおけるデータ分割}}{9} \addvspace {10\p@ } \contentsline {figure}{\numberline {3.1}{\ignorespaces Task Manager}}{11} @@ -16,3 +16,4 @@ \addvspace {10\p@ } \addvspace {10\p@ } \addvspace {10\p@ } +\addvspace {10\p@ }
--- a/paper/master_paper.log Wed Feb 04 16:38:20 2015 +0900 +++ b/paper/master_paper.log Mon Feb 09 11:32:28 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) 4 FEB 2015 16:33 +This is e-pTeX, Version 3.1415926-p3.4-110825-2.6 (utf8.euc) (TeX Live 2013) (format=platex 2013.5.30) 9 FEB 2015 11:32 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 2. 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 30. +(Font) Font shape `JT1/gt/m/n' tried instead on input line 29. 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 30. +(Font) Font shape `JY1/gt/m/n' tried instead on input line 29. 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 30. +(Font) Font shape `JT1/gt/m/n' tried instead on input line 29. 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 30. +(Font) Font shape `JY1/gt/m/n' tried instead on input line 29. LaTeX Font Info: External font `cmex10' loaded for size -(Font) <10.95> on input line 30. +(Font) <10.95> on input line 29. [1 ]) @@ -264,14 +264,7 @@ File: images/emblem-bitmap.eps Graphic file (type eps) <images/emblem-bitmap.eps> [2] 第 2 章(3ページ) -LaTeX Font Info: Try loading font information for OMS+cmr on input line 9. -(/usr/local/share/texlive/2013/texmf-dist/tex/latex/base/omscmr.fd -File: omscmr.fd 1999/05/25 v2.5h Standard LaTeX font definitions -) -LaTeX Font Info: Font shape `OMS/cmr/m/n' in size <12> not available -(Font) Font shape `OMS/cmsy/m/n' tried instead on input line 9. File: ./images/gpu_arch.pdf Graphic file (type pdf) - <./images/gpu_arch.pdf> File: ./images/cpu_arch.pdf Graphic file (type pdf) <./images/cpu_arch.pdf> [3 @@ -280,28 +273,35 @@ File: images/emblem-bitmap.eps Graphic file (type eps) <images/emblem-bitmap.eps> [4] +LaTeX Font Info: Try loading font information for OMS+cmr on input line 66. + (/usr/local/share/texlive/2013/texmf-dist/tex/latex/base/omscmr.fd +File: omscmr.fd 1999/05/25 v2.5h Standard LaTeX font definitions +) +LaTeX Font Info: Font shape `OMS/cmr/m/n' in size <12> not available +(Font) Font shape `OMS/cmsy/m/n' tried instead on input line 66. +File: images/emblem-bitmap.eps Graphic file (type eps) + +<images/emblem-bitmap.eps> [5] File: ./images/workitem.pdf Graphic file (type pdf) <./images/workitem.pdf> File: images/emblem-bitmap.eps Graphic file (type eps) - <images/emblem-bitmap.eps> [5] -File: images/emblem-bitmap.eps Graphic file (type eps) <images/emblem-bitmap.eps> [6] +File: images/emblem-bitmap.eps Graphic file (type eps) + <images/emblem-bitmap.eps> [7] File: ./images/calculateIndex.pdf Graphic file (type pdf) - <./images/calculateIndex.pdf> + +<./images/calculateIndex.pdf> +LaTeX Font Info: Font shape `OMS/cmr/m/n' in size <10> not available +(Font) Font shape `OMS/cmsy/m/n' tried instead on input line 227. +File: images/emblem-bitmap.eps Graphic file (type eps) + <images/emblem-bitmap.eps> [8] +File: ./images/starpu_data_parallel.pdf Graphic file (type pdf) + +<./images/starpu_data_parallel.pdf>) (./chapter2.tex File: images/emblem-bitmap.eps Graphic file (type eps) -<images/emblem-bitmap.eps> [7] -LaTeX Font Info: Font shape `OMS/cmr/m/n' in size <10> not available -(Font) Font shape `OMS/cmsy/m/n' tried instead on input line 195. -File: ./images/starpu_data_parallel.pdf Graphic file (type pdf) - <./images/starpu_data_parallel.pdf>) -(./chapter2.tex -File: images/emblem-bitmap.eps Graphic file (type eps) - <images/emblem-bitmap.eps> [8] -File: images/emblem-bitmap.eps Graphic file (type eps) - <images/emblem-bitmap.eps> -[9] +<images/emblem-bitmap.eps> [9] 第 3 章(10ページ) [10 @@ -315,22 +315,22 @@ <images/emblem-bitmap.eps> [11] LaTeX Font Warning: Font shape `JT1/mc/m/it' undefined -(Font) using `JT1/mc/m/n' instead on input line 66. +(Font) using `JT1/mc/m/n' instead on input line 72. LaTeX Font Warning: Font shape `JY1/mc/m/it' undefined -(Font) using `JY1/mc/m/n' instead on input line 66. +(Font) using `JY1/mc/m/n' instead on input line 72. LaTeX Font Info: External font `cmex10' loaded for size -(Font) <7> on input line 67. +(Font) <7> on input line 73. LaTeX Font Info: External font `cmex10' loaded for size -(Font) <5> on input line 67. -LaTeX Font Info: Try loading font information for OML+cmr on input line 67. +(Font) <5> on input line 73. +LaTeX Font Info: Try loading font information for OML+cmr on input line 73. (/usr/local/share/texlive/2013/texmf-dist/tex/latex/base/omlcmr.fd File: omlcmr.fd 1999/05/25 v2.5h Standard LaTeX font definitions ) LaTeX Font Info: Font shape `OML/cmr/m/n' in size <10> not available -(Font) Font shape `OML/cmm/m/it' tried instead on input line 67. +(Font) Font shape `OML/cmm/m/it' tried instead on input line 73. File: images/emblem-bitmap.eps Graphic file (type eps) <images/emblem-bitmap.eps> [12]) (./chapter3.tex @@ -341,8 +341,8 @@ [14 ] -File: ./images/sort.pdf Graphic file (type pdf) - <./images/sort.pdf> +File: ./images/sort_benchmark.pdf Graphic file (type pdf) + <./images/sort_benchmark.pdf> File: images/emblem-bitmap.eps Graphic file (type eps) <images/emblem-bitmap.eps> [15] File: ./images/wordcount.pdf Graphic file (type pdf) @@ -362,31 +362,36 @@ ] 第 5 章(19ページ) -[19 +) (./chapter5.tex [19 -]) (./chapter5.tex -File: images/emblem-bitmap.eps Graphic file (type eps) - <images/emblem-bitmap.eps> [20] -第 6 章(21ページ) -LaTeX Font Info: Font shape `JT1/mc/bx/n' in size <14.4> not available -(Font) Font shape `JT1/gt/m/n' tried instead on input line 7. -LaTeX Font Info: Font shape `JY1/mc/bx/n' in size <14.4> not available -(Font) Font shape `JY1/gt/m/n' tried instead on input line 7. -) (./chapter6.tex [21 +] +第 6 章(20ページ) +[20 ] -第 7 章(22ページ) -) (./chapter7.tex [22 +File: images/emblem-bitmap.eps Graphic file (type eps) + <images/emblem-bitmap.eps> [21]) (./chapter6.tex +File: images/emblem-bitmap.eps Graphic file (type eps) + +<images/emblem-bitmap.eps> [22] +File: images/emblem-bitmap.eps Graphic file (type eps) + <images/emblem-bitmap.eps> [23] +第 7 章(24ページ) +) (./chapter7.tex [24 ] -第 8 章(23ページ) -) (./conclusion.tex [23 +第 8 章(25ページ) +) (./chapter8.tex [25 ] -第 9 章(24ページ) -) (./thanx.tex [24 +第 9 章(26ページ) +) (./conclusion.tex [26 -]) (./master_paper.bbl [25 +] +第 10 章(27ページ) +) (./thanx.tex [27 + +]) (./master_paper.bbl [28 ] @@ -397,11 +402,11 @@ LaTeX Font Warning: Font shape `JY1/gt/m/it' undefined (Font) using `JY1/gt/m/n' instead on input line 25. -) (./appendix.tex [26 +) (./appendix.tex [29 ]) No file master_paper.ind. -[27 +[30 ] (./master_paper.aux) @@ -409,12 +414,12 @@ ) Here is how much of TeX's memory you used: - 2474 strings out of 494008 - 34809 string characters out of 6154472 - 177731 words of memory out of 5000000 - 5856 multiletter control sequences out of 15000+600000 - 20103 words of font info for 78 fonts, out of 8000000 for 9000 + 2492 strings out of 494008 + 35051 string characters out of 6154472 + 179731 words of memory out of 5000000 + 5871 multiletter control sequences out of 15000+600000 + 18570 words of font info for 72 fonts, out of 8000000 for 9000 745 hyphenation exceptions out of 8191 - 32i,11n,40p,301b,1252s stack positions out of 5000i,500n,10000p,200000b,80000s + 33i,12n,40p,301b,1252s stack positions out of 5000i,500n,10000p,200000b,80000s -Output written on master_paper.dvi (34 pages, 75508 bytes). +Output written on master_paper.dvi (37 pages, 91888 bytes).
--- a/paper/master_paper.lot Wed Feb 04 16:38:20 2015 +0900 +++ b/paper/master_paper.lot Mon Feb 09 11:32:28 2015 +0900 @@ -6,9 +6,11 @@ \contentsline {table}{\numberline {3.2}{\ignorespaces Task 側で使用する API}}{13} \addvspace {10\p@ } \addvspace {10\p@ } -\contentsline {table}{\numberline {5.1}{\ignorespaces Ceriumを実行する実験環境1}}{19} -\contentsline {table}{\numberline {5.2}{\ignorespaces Ceriumを実行する実験環境2}}{20} +\addvspace {10\p@ } +\contentsline {table}{\numberline {6.1}{\ignorespaces データ並列実行時の index の割り当て}}{22} \addvspace {10\p@ } \addvspace {10\p@ } \addvspace {10\p@ } +\contentsline {table}{\numberline {9.1}{\ignorespaces Ceriumを実行する実験環境1}}{26} +\contentsline {table}{\numberline {9.2}{\ignorespaces Ceriumを実行する実験環境2}}{26} \addvspace {10\p@ }
--- a/paper/master_paper.tex Wed Feb 04 16:38:20 2015 +0900 +++ b/paper/master_paper.tex Mon Feb 09 11:32:28 2015 +0900 @@ -84,6 +84,7 @@ \input{chapter5.tex} \input{chapter6.tex} \input{chapter7.tex} +\input{chapter8.tex} \input{conclusion.tex} %謝辞
--- a/paper/master_paper.toc Wed Feb 04 16:38:20 2015 +0900 +++ b/paper/master_paper.toc Mon Feb 09 11:32:28 2015 +0900 @@ -1,9 +1,11 @@ \contentsline {chapter}{\numberline {第1章}研究目的と背景}{1} \contentsline {section}{\numberline {1.1}本論文の構成}{2} \contentsline {chapter}{\numberline {第2章}既存のマルチプラットフォームフレームワーク}{3} -\contentsline {section}{\numberline {2.1}OpenCL}{3} -\contentsline {section}{\numberline {2.2}CUDA}{6} -\contentsline {section}{\numberline {2.3}StarPU}{7} +\contentsline {section}{\numberline {2.1}Architecutre}{3} +\contentsline {section}{\numberline {2.2}Shared Memory}{3} +\contentsline {section}{\numberline {2.3}OpenCL}{5} +\contentsline {section}{\numberline {2.4}CUDA}{7} +\contentsline {section}{\numberline {2.5}StarPU}{8} \contentsline {chapter}{\numberline {第3章}Cerium}{10} \contentsline {section}{\numberline {3.1}Cerium の概要}{10} \contentsline {section}{\numberline {3.2}Cerium TaskManager}{10} @@ -16,24 +18,19 @@ \contentsline {section}{\numberline {4.3}FFT}{18} \contentsline {chapter}{\numberline {第5章}マルチコアへの対応}{19} \contentsline {section}{\numberline {5.1}マルチコア上での実行の機構}{19} -\contentsline {section}{\numberline {5.2}ベンチマーク}{19} -\contentsline {chapter}{\numberline {第6章}GPGPU への対応}{21} -\contentsline {section}{\numberline {6.1}OpenCL による実装}{21} -\contentsline {subsection}{\numberline {6.1.1}ベンチマーク}{21} -\contentsline {section}{\numberline {6.2}CUDA による}{21} -\contentsline {subsection}{\numberline {6.2.1}ベンチマーク}{21} -\contentsline {section}{\numberline {6.3}データ並列}{21} -\contentsline {subsection}{\numberline {6.3.1}ベンチマーク}{21} -\contentsline {chapter}{\numberline {第7章}並列処理向けI/O}{22} -\contentsline {section}{\numberline {7.1}新たに実装したI/Oの機構}{22} -\contentsline {section}{\numberline {7.2}ベンチマーク}{22} -\contentsline {chapter}{\numberline {第8章}Memory Allocator}{23} -\contentsline {section}{\numberline {8.1}現状のMemory Allocator}{23} -\contentsline {section}{\numberline {8.2}新しいMemory Allocator}{23} -\contentsline {section}{\numberline {8.3}ベンチマーク}{23} -\contentsline {chapter}{\numberline {第9章}結論}{24} -\contentsline {section}{\numberline {9.1}まとめ}{24} -\contentsline {section}{\numberline {9.2}今後の課題}{24} -\contentsline {chapter}{謝辞}{25} -\contentsline {chapter}{参考文献}{26} -\contentsline {chapter}{発表文献}{27} +\contentsline {section}{\numberline {5.2}DMA}{19} +\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}{24} +\contentsline {section}{\numberline {7.1}新たに実装したI/Oの機構}{24} +\contentsline {chapter}{\numberline {第8章}Memory Allocator}{25} +\contentsline {section}{\numberline {8.1}現状のMemory Allocator}{25} +\contentsline {section}{\numberline {8.2}新しいMemory Allocator}{25} +\contentsline {chapter}{\numberline {第9章}ベンチマーク}{26} +\contentsline {chapter}{\numberline {第10章}結論}{27} +\contentsline {section}{\numberline {10.1}まとめ}{27} +\contentsline {section}{\numberline {10.2}今後の課題}{27} +\contentsline {chapter}{謝辞}{28} +\contentsline {chapter}{参考文献}{29} +\contentsline {chapter}{発表文献}{30}