# HG changeset patch # User Shohei KOKUBO # Date 1424065638 -32400 # Node ID 346b04bec5925c1401bfd046fb223f4a7e020722 # Parent aa8dc80cd3ed9da78458a7c337aac95485e2a90e commit diff -r aa8dc80cd3ed -r 346b04bec592 paper/Benchmark.tex --- a/paper/Benchmark.tex Thu Feb 27 12:51:36 2014 +0900 +++ b/paper/Benchmark.tex Mon Feb 16 14:47:18 2015 +0900 @@ -52,11 +52,59 @@ \caption{WordCount} \end{table} -パイプラインありの方が少しだけ実行結果が向上している。しかし、期待するほどの効果が出ていない。 +パイプラインありの方が OpenCL では1.03倍、CUDA では1.04倍の性能向上が見られた。しかし、期待するほどの効果が出ていない。 これは綺麗なパイプラインで実行されていないことが考えられる。 パイプラインで効果を出すためにはデータ転送と Task の実行が同じ時間で終了することが必要になる。 -一方、データ並列による実行では OpenCL では 1 CPU より速く、CUDA ではほぼ同等という結果になった。 +一方、データ並列による実行ではタスク並列の実行に比べ、OpenCL は93倍、CUDA では73倍の性能向上が見られた。 +\newpage \section{FFT} 次に、フーリエ変換と周波数フィルタによる画像処理を行う例題を利用し測定する。 -使用する画像のサイズは 512*512 で、それに対して High Pass Filter をかけて変換する。 +FFT の Task は構成は以下のようになる。 +\begin{itemize} +\item spinFact +\item bitReverse +\item norm +\item butterfly +\item transpose +\item highPassFilter +\end{itemize} + +spinFact はバタフライ演算で乗算要素として用いられる回転因子を計算する Task である。 +この値はバタフライ演算するたびに毎回計算してもよいが、サンプル数が同じであれば同じ回転因子を使うのであらかじめ計算しておけばよい。 + +bitReverse はバタフライ演算の前に各要素をビット逆順に並び替える Task である。 +データ並列によって実行したいので、index のビットを実際に反転させることで入れ替え先の index を生成します。 + +norm はデータを使いやすくするため正規化する Task である。 + +butterfly はペアリングされた一つのバタフライ演算を計算する Task である。 + +transpose は行と列の index を入れ替える Task である。 + +highPassFilter は高周波数領域のみを通すフィルタをかける Task である。 + +今回、使用する画像のサイズは 512*512 で、それに対して High Pass Filter をかけて変換する。 +表:\ref{table:fft}はその実行時間である。 +\begin{table}[!h] + \begin{center} + \small + \begin{tabular}[t]{c||r} \hline + & Run Time \\ \hline + 1 CPU & 0.45s \\ \hline + 2 CPU & 0.26s \\ \hline + 4 CPU & 0.17s \\ \hline + 8 CPU & 0.11s \\ \hline + 12 CPU & 0.07s \\ \hline + OpenCL & 0.09s \\ \hline + CUDA & 0.21s \\ \hline + \end{tabular} + \end{center} + \label{table:fft} + \caption{FFT} +\end{table} + +OpenCL では 1 CPU の5倍、CUDA では 2.1倍の性能向上が見られた。 +しかし、OpenCL と CUDA で差がある。これは、GpuScheduler に対して CudaScheduler の同期がまだ不十分であることが考えられる。 +また、OpenCL の kernel では cl\_float2 というベクター型を用いている。しかし、CUDA では cl\_float2 を float に変換して演算している。 +OpenCL ではベクターの演算なので、その部分に最適化がかかっている可能性もある。 diff -r aa8dc80cd3ed -r 346b04bec592 paper/Conclusion.tex --- a/paper/Conclusion.tex Thu Feb 27 12:51:36 2014 +0900 +++ b/paper/Conclusion.tex Mon Feb 16 14:47:18 2015 +0900 @@ -1,5 +1,29 @@ \chapter{結論} - \section{まとめ} +本研究では並列プログラミングフレームワーク Cerium の改良を行なった。 +Ceirum Task Manager は PS3 および Linux、Mac OS X 上で動く並列プログラミングフレームワークである。 +CPU/Cell だけでなく、GPU も計算資源として利用できる。 +今回、新たに iterate API を実装し、Multi Core CPU でのデータ並列での実行を可能にした。 +OpenCL による GPGPU に加え、CUDA による GPGPU を可能にした。 +また、host と device 間でのデータ転送をオーバーラップするため GpuSschduler および CudaScheduler にパイプライン実行機構を実装した。 +WordCount および FFT を例題として用いて、実装した GpuScheduler および CudaScheduler の評価を行なった。 +これらの改良によって、パイプライン実行を行うと OpenCL で1.03倍、CUDE で 1.04倍の性能向上に成功した。 \section{今後の課題} +\subsection{CPU と GPU による同時実行のスケジューリング} +現在、Cerium では CPU と GPU の同時実行に対応している。 +しかし、現在 Task を CPU, GPU に対し交互に割り振っているため CPU 単体、GPU 単体で実行するより遅くなる。 +Task の割り当てを最適にスケジューリングすることでより性能向上させることができる。 +スケジューリングの方法として、一度 Task を CPU のみ、GPU のみ実行し、プロファイルを取りそれを元に Task の割り当てを決定するなどが考えられる。 + +\subsection{パイプライン実行の改良} +パイプラインを入れたことに性能は向上したが、期待した通りの効果が出ていない。 +まず、パイプラインで実行されているか調査する必要がある。 +綺麗なパイプラインで実行するためには Task の実行時間とデータ転送が同じ時間で終了することが望ましい。 +Task の粒度や転送するデータのサイズを考える必要がある。 + +\subsection{kernel の自動変換} +現在、kernel の記述は CPU, OpenCL, CUDA でほぼ同じ記述である。 +しかし、修飾子など多少の違いがある。 +どれか1種類 kernel を記述すれば他の kernel も生成できるようにするのが望ましい。 + diff -r aa8dc80cd3ed -r 346b04bec592 paper/Thanks.tex --- a/paper/Thanks.tex Thu Feb 27 12:51:36 2014 +0900 +++ b/paper/Thanks.tex Mon Feb 16 14:47:18 2015 +0900 @@ -1,1 +1,5 @@ -\chapter{謝辞} +\chapter*{謝辞} +本研究を行うにあたって、ご多忙にも関わらず日頃より多くのご助言を、ご指導を頂きました河野真治准教授に心より感謝いたします。 +また、本研究に大変貴重なご意見を下さった情報工学科の先生方に感謝いたします。 +また本研究に大きく役立つ技術的指導を賜りました、大城信康さん、當眞大千さん、谷成雄さん、並びに研究室の配属生の皆様に深く感謝いたします。 +最後に、長年に渡り理解を示し、支援してくださった家族に感謝いたします。 diff -r aa8dc80cd3ed -r 346b04bec592 paper/introduction.tex --- a/paper/introduction.tex Thu Feb 27 12:51:36 2014 +0900 +++ b/paper/introduction.tex Mon Feb 16 14:47:18 2015 +0900 @@ -2,4 +2,20 @@ \pagenumbering{arabic} \section{研究背景と目的} -コンピュータを高速化するため、CPU のビット幅の拡張、動作クロックの高速化、キャッシュの大容量化などが図られてきた。しかし、これらの方法による性能の向上はすでに限界に達している。そこで並列化による高速化が注目された。同一アーキテクチャのプロセッサを複数台使った並列コンピュータ(ホモジニアス)が一般的でしたが、GPU の普及と高速化、GPU の演算資源を画像処理以外の目的にも使用する GPGPU (GPU による汎目的計算)の登場によって異なるアーキテクチャのプロセッサを組み合わせた並列コンピュータ(ヘテロジニアス)が出現した。これにより Many Core CPU とは比較にならないほどの並列化数を実現できるようになった。ただし、GPU の各コアは CPU のコアほど高性能ではなく、演算は高速だが、制御部分は弱い。また、ヘテロジニアスなシステムはホモジニアスなシステムとは異なり、メモリ空間がまったく同一でない場合であることが多い。このような場合、メモリコピーによるオーバーヘッドが大きくなる。よって、GPU を用いて高い並列度を出すためには、特定の計算に特化した Task の生成やスケジューリング、Task をパイプライン実行することでデータ転送をオーバーラップする必要がある。本研究では、OpenCL, CUDA を用いて GPU 上での Task 並列実行と Data 並列実行の両方をサポートし、CPU と GPU の Task をほぼ同じに記述することができ、自動でパイプライン実行を行いデータ転送をオーバーラップするように当研究室で開発した並列プログラミングフレーム Cerium を改良した。WordCount, FFT を用いて実行時間を測定して、その結果から Cerium 上での GPU 実行機構を評価する。また、Task を CPU, GPU 上での同時実行も可能にした。しかし、Task をどの程度の割合で CPU/GPU に割り当てるかというスケジューリング等の問題がある。測定結果からスケジューリング等の問題の解決方法についても考察し、信頼性のある並列プログラミングフレームを目指す。 +PS3 および Linux、Mac OS X 上で動く並列プログラミングフレームワーク、Cerium Task Manager\cite{gongo:2008a}の開発・改良を行なっている。 + +Cell だけでなく、GPU のような異なるアーキテクチャを搭載した CPU、つまりヘテロジニアスな CPU が増えてきた。 +GPU の普及と高性能化にともない、GPU の演算資源を画像処理以外の目的にも利用する GPGPU(GPU による汎目的計算)が注目されている\cite{fft:gpu}。 + +特定の計算に特化した Task の生成やスケジューリングを行い、高い並列度を出すという研究は様々な分野で行われている。 +しかし、Cerium Task Manager は特定の計算に限らず、GPU を用いた汎用計算できるフレームワークを目指している。 + +GPU のような数百個の Core を持つようなプロセッサではデータ並列と呼ばれる実行方法が推奨されている。 +データ並列とは多次元のデータ構造に対して、それを分割して各要素に対して処理を行うことを指す。 +また、GPU は CPU とメモリ空間が異なるため、データに直接アクセスすることができない。 +データにアクセスするためにはメモリ空間ごとコピーする必要がある。 +このデータ転送部分が大きなオーバーヘッドになる。 +これらに対応するために今回では OpenCL と CUDA を用いて GPU 上で実行するための機構を実装した。 +OpenCL と CUDA にはデータ並列を行うための API が存在する。 +また、自動的にデータ転送をオーバーラップするためパイプライン実行するように Scheduler を設計し、実装した。 +WordCaunt および FFT を例題として用いて、評価を行なった。 diff -r aa8dc80cd3ed -r 346b04bec592 paper/master_paper.sty --- a/paper/master_paper.sty Thu Feb 27 12:51:36 2014 +0900 +++ b/paper/master_paper.sty Mon Feb 16 14:47:18 2015 +0900 @@ -72,7 +72,7 @@ \topmargin 0mm \headheight 10mm \headsep 15mm -\textheight 39\baselineskip \addtolength{\textheight}{\topskip} +\textheight 36\baselineskip \addtolength{\textheight}{\topskip} \textwidth 160mm \marginparsep 3mm \marginparwidth 15mm diff -r aa8dc80cd3ed -r 346b04bec592 paper/thesis-paper.bib --- a/paper/thesis-paper.bib Thu Feb 27 12:51:36 2014 +0900 +++ b/paper/thesis-paper.bib Mon Feb 16 14:47:18 2015 +0900 @@ -6,20 +6,11 @@ year = 2008 } -@article{akira:2008a, - author = "神里 晃 and 宮國 渡 and 杉山 千秋 and 河野 真治", - title = "CからCellアーキテクチャを利用したCbCへの変換", - journal = "電子情報通信学会VLSI設計技術研究会", - month = "March", - year = 2008 -} - -@article{akira:2006a, - author = "神里 晃 and 河野 真治", - title = "Continuation based CによるPS3 Cell のシミュレーション", - journal = "情報処理学会システムソフトウェアとオペレーティング・システム研究会", - month = "May", - year = 2006 +@article{fft:GPU, + author = "Yasuhiko OGATA and Toshio ENDO and Naoya MARUYAMA and Satoshi MATSUOKA", + title = "性能モデルに基づくCPU及びGPUを併用する効率的なFFTライブラリ", + school = "東京工業大学", + year = 2008 } @mastersthesis{akira:master, @@ -30,14 +21,6 @@ year = 2008 } -@mastersthesis{gongo:master, - author = "宮國 渡", - title = "Cell用の Fine-Grain Task Manager の実装", - school = "琉球大学理工学研究科情報工学専攻", - month = "Feb", - year = 2009 -} - @article{gongo:2008a, author = "宮國 渡 and 河野 真治 and 神里 晃 and 杉山 千秋", title = "Cell 用の Fine-grain Task Manager の実装", @@ -53,14 +36,6 @@ year = 2008 } -@article{akamine:2010a, - author = "赤嶺 一樹 and 河野 真治", - title = "Meta Engine を用いた Federated Linda の実験", - journal = "日本ソフトウェア科学会第 27 会大会(2010 年度)", - month = "Sep", - year = 2010 -} - @mastersthesis{kaito:master, author = "多賀野 海人", title = "Cell Task Manager Cerium における Task を用いたパイプラインの改良", @@ -137,21 +112,12 @@ year = 2008 } - -@misc{blender, -title = "{blender.org}", -howpublished = "{http://blender.org/}" +@misc{opencl, +title = "{OpenCL}", +howpublished = "{https://www.khronos.org/opencl/}" } - -@misc{opengl, -title = "{opengl}", -howpublished = "{http://www.opengl.org/}" + +@misc{cuda, +title = "{CUDA}", +howpublished = "{https://developer.nvidia.com/category/zone/cuda-zone/}" } - -@article{akamine, -author = "{赤嶺一樹, 河野真治}", -title = "{Meta Engine を用いた Federated Linda の実験}", -journal = "{日本ソフトウェア科学会第 27 会大会(2010 年度)}", -month = "{Sep}", -year = 2010 -} diff -r aa8dc80cd3ed -r 346b04bec592 paper/thesis-paper.pdf Binary file paper/thesis-paper.pdf has changed diff -r aa8dc80cd3ed -r 346b04bec592 presen/index.html --- a/presen/index.html Thu Feb 27 12:51:36 2014 +0900 +++ b/presen/index.html Mon Feb 16 14:47:18 2015 +0900 @@ -28,7 +28,7 @@
担当教官 : 河野 真治
-

パイプライン実行および CudaScheduler の実装

+

CudaScheduler の実装および OpenCL, CUDA でのパイプライン実行

PS3 および Linux、Mac OS X 上で動く並列プログラミングフレームワーク、 Cerium の開発・改良を行っている @@ -59,8 +59,25 @@ この質問があったので CUDA バージョンを3週間で実装した

+
+

GPU のアーキテクチャ

+ + + + + + + +
+

+ CPU と GPU はメモリ空間が異なるのでメモリの共有ができない。データにアクセスするためには空間ごとコピーする必要がある。この転送部分がネックになる。 +

+

+ GPGPU ではデータ転送が頻繁に起こるような Task は並列度が低くなる。 +

+
-

CUDA と OpenCL の対応(1/2)

+

Porting OpenCL to CUDA(1/2)

用語

@@ -92,40 +109,20 @@
-

CUDA と OpenCL の対応(2/2)

- - -

オブジェクト

- - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
OpenCLCUDA
cl_device_idCUdevice
cl_contextCUcontext
cl_programCUmodule
cl_kernelCUfunction
cl_memCUdeviceptr
cl_command_queueCUstream
+

Porting OpenCL to CUDA(2/2)

+ OpenCL + 
+   cl_mem memA = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(float), NULL, &ret);
+   clEnqueueWriteBuffer(command_queue, memA, CL_FALSE, 0,
+                        sizeof(float), A, 0, NULL, NULL);
+
+ CUDA + 
+   CUdeviceptr devA;
+   cuMemAlloc(&devA, sizeof(float));
+   cuMemcpyHtoDAsync(devA, A, sizeof(float), stream);
+

CUDA による GPGPU への対応

@@ -148,46 +145,26 @@

-

Cerium Task Manager

- - - - - - - -
-

- TaskManager と各 Threads は mail によって同期を取っている。 -

-

- 実行するデバイスに対応した Threads が Task を受け取り、並列実行する。 -

-
+

Cerium Task Manager(1/2)

+

- 生成した Task は TaskManager で依存関係をチェックし、依存関係が解消されたあと、各 Scheduler に転送される。 + 生成した Task は TaskManager で依存関係をチェックされる。依存関係がないと ActiveTaskList に移され、実行するデバイスの Scheduler に転送される。

-
-

GPU のアーキテクチャ

- - - - - - - -
-

- CPU と GPU はメモリ空間が異なるのでメモリの共有ができない。データにアクセスするためには空間ごとコピーする必要がある。この転送部分がネックになる。 -

-

- GPGPU ではデータ転送が頻繁に起こるような Task は並列度が低くなる。 -

+
+

Cerium Task Manager(2/2)

+ +

+ Scheduler に転送されるとそこで並列実行される。Task が終了すると同期キューである mail を使用し、TaskManager に終了した Task を通知し、再び依存関係をチェックする。 +

Multi Stream による並列実行

+ OpenCL の Command Queue は queue を生成するとき、プロパティを設定することで単一の queue でも Operation を並列実行するようにできる。 + しかし、ほとんどのデバイスでこのプロパティは使用することができない。 +

+

Stream に投入された Operation は投入された順序で実行されることが保証されている

@@ -199,17 +176,17 @@

CUDA におけるデータ並列

- 同一の kernel を生成し、各 kernel に対して ID を割り振る。 + 同一の kernel を複数生成し、各 kernel に対して index を割り振る。

- kernel に割り振られた ID から担当する範囲を求め、処理を行う。 + 割り振られた index から自分が担当する範囲(GPU 上でのメモリ領域)を求め、各 kernel が並列に処理を行う。

- ID は組み込み変数から算出することができる。 + CUDA では index は組み込み変数から算出することができる。OpenCL では get_global_id という API を用いることで index を取得できる。

-

Cerium におけるデータ並列のための新しい API

+

Many Core CPU によるデータ並列実行のための新しい API

@@ -225,40 +202,49 @@
  • 1つの記述から複数のTaskを生成する
  • 生成した複数のTaskにIDとInput/Output Dataを割り当てる
    • +
  • Multi Core CPU での index の取得は引数として渡される SchedTask に格納されている
-

Cerium における CUDA でのデータ並列

+

Cerium における CUDA でのデータ並列(1/2)

+ 
+  void
+  HTask::iterate(long x, long y, long z) {
+    this->flag.dim_count = 1;
+    TaskList *tl = (TaskList*)rbuf;
+    tl->self = this;
+    tl->set_last(last);
+    tl->dim=3;
+    tl->x=x;
+    tl->y=y;
+    tl->z=z;
+    mimpl->spawn_task(this);
+  }
+

+ iterate を使用するとデータ並列を行う Task として登録される。この時点で tasklist にパラメータが設定され、TaskManager から Scheduler に送られる。 +

+
+

Cerium における CUDA でのデータ並列(2/2)

CudaScheduler は受け取った tasklist をもとに GPU の制御を行う
tasklist は制御に必要なパラメータを持っている

+ データ転送 + 
+  TaskPtr nextTask = tasklist->tasks;
+  ListElement* input_buf = nextTask->inData;
+  cuMemcpyHtoDAsync(memin, input_buf->addr, input->buf->size, stream);
+
+ kernel の実行 
    cuLaunchKernel(function, tasklist->x, tasklist->y, tasklist->z,
                   1, 1, 1, 0, stream[cur], kernelParam, NULL);
- - -

taskList の持つメンバ変数

- - - - - - - - - - - - - -
xx 座標のブロック数
yy 座標のブロック数
zz 座標のブロック数
-

WordCount の測定

+

WordCount の測定(1/2)

@@ -286,10 +272,6 @@ - - - - @@ -297,6 +279,43 @@ + +
46.74s
OpenCL(Data Parallel)0.50s
CUDA(no pipeline) 55.71s
CUDA(pipeline) 53.30s
+ + +

100MB のテキストに対するタスク並列で wordcount

+ +

実験環境

+ OS : MacOS 10.9.1
+ CPU : 2*2.66GHz 6-CoreIntel Xeon
+ GPU : NVIDIA Quadro K5000 4096MB
+ Compiler : Apple LLVM version 5.0(clang-500.2.79)(based on LLVM 3.3svn)
+ +

結果

+ パイプラインありの方が OpenCL では1.03倍、CUDA では1.04倍の性能向上が見られた
+ しかし、CPU に比べてまだ60倍以上遅い + + + + +
+
+

WordCount の測定(2/2)

+ + + + @@ -418,6 +436,27 @@ 60inch4Kくらいの買ってください。

+
+

kernel の差異

+ OpenCL + 
+  __kernel void
+  multiply(__global float* A, __global float* B, __global float* C) {
+    long i = get_global_id(0);
+    C[i] = A[i] * B[i];
+  }
+
+ CUDA + 
+  extern "C" {
+    __global__ void
+    multiply(float* A, float* B, float* C) {
+      long i = blockIdx.x*blockDim.x*+threadIdx.x;
+      C[i] = A[i] * B[i];
+    }
+  }
+
+
+ + + + + + + + + + @@ -305,7 +324,7 @@
Time
OpenCL(Data Parallel)0.50s
CUDA(Data Parallel) 0.73s
 -

100MB のテキストファイルに対する wordcount

+

100MB のテキストに対するデータ並列で wordcount

実験環境

OS : MacOS 10.9.1
@@ -315,8 +334,7 @@ Compiler : Apple LLVM version 5.0(clang-500.2.79)(based on LLVM 3.3svn)

結果

- パイプラインありの方が少しだけ性能が向上した
- データ並列による実行では、OpenCL 版は 1 CPU より速く、CUDA 版はほぼ同等という結果になった + データ並列による実行ではタスク並列の実行に比べ、OpenCL は93倍、CUDA では73倍の性能向上が見られた