Mercurial > hg > Papers > 2015 > yuhi-master
annotate paper/chapter5.tex @ 79:7990a2abbf05 default tip
add file
| author | Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp> |
|---|---|
| date | Mon, 02 Mar 2015 11:21:40 +0900 |
| parents | f9b73e12a52f |
| children |
| rev | line source |
|---|---|
| 13 | 1 \chapter{GPGPU への対応} |
| 77 | 2 Cerium の新たな演算資源として GPU の使用を可能にした。\cite{yuhi:2013a} |
|
10
1519674c30ab
dragon and firefly spec
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
8
diff
changeset
|
3 現在、GPU のような異なる種類のアーキテクチャを搭載した CPU 、つまりヘテロジニアスな CPU が増えている。 |
|
1519674c30ab
dragon and firefly spec
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
8
diff
changeset
|
4 特定の計算に特化した Task の生成やスケジューリングを行い、 GPGPU により高い並列度を出す研究は様々な分野で行われている。 |
|
1519674c30ab
dragon and firefly spec
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
8
diff
changeset
|
5 本研究では Cerium を特定の計算に限らず、 GPU を用いて汎用計算を行えるフレームワークに改良する。 |
| 7 | 6 |
| 7 | |
| 15 | 8 \section{OpenCL および CUDA による実装} |
| 9 OpenCL 、CUDA による GPGPU 対応を行った。 | |
| 10 Scheduler と CpuThreads に対応させる形で | |
| 11 OpenCL を用いた GpuScheduler と GpuThreads、 | |
| 12 CUDA を用いた CudaScheduler と CudaThreads を実装した。 | |
| 13 それぞれの Scheduler 内で各フレームワークの API を用いて GPU の制御を行っている。 | |
| 14 | |
| 15 TaskManager から受け取った TaskList をもとに Device 上のメモリバッファを作成する。 | |
| 16 その後 CommandQueue、 Stream といったそれぞれの Queue に Device 制御用の Command を Queueing していく。 | |
| 17 | |
| 18 Command は Queueing した順に実行されるので、以下のように Command を Queueing する。 | |
| 19 \begin{enumerate} | |
| 20 \item Host から Device へのデータ転送 | |
| 21 \item kernel の実行 | |
| 22 \item Device から Host へのデータ転送 | |
| 23 \end{enumerate} | |
| 24 | |
| 25 データの転送や kernel の実行は非同期 API を用いることで並列に行うことができる。 | |
| 26 | |
| 27 通常、フレームワークが依存関係を解決して実行するが、 | |
| 28 非同期 API を用いる場合はユーザが依存関係を考慮する必要がある。 | |
| 29 しかし Task の依存関係は TaskManager が既に解決した状態で送ってくるので、 | |
| 30 Scheduler は依存関係を考慮せずに実行して問題ない。 | |
| 31 | |
| 32 GPGPU 用の Scheduler は CommandQueue を2つ持っており、Task をパイプライン的に実行する。 | |
| 50 | 33 GpuScheduler のパイプライン処理部分をソースコード:\ref{src:pipeline_gpu}に示す。 |
| 34 \newpage | |
| 35 \begin{lstlisting}[frame=lrbt,label=src:pipeline_gpu,caption=GpuSchedulerにおけるパイプライン処理,numbers=left] | |
| 36 void | |
| 37 GpuScheduler::run() { | |
| 38 for (;;) { | |
| 39 memaddr params_addr = connector->task_list_mail_read(); | |
| 40 // read task list mail from DmaManager | |
| 41 | |
| 42 while (params_addr) { | |
| 43 // since we are on the same memory space, we don't has to use dma_load here | |
| 44 tasklist = (TaskListPtr)connector->dma_load(this, params_addr,sizeof(TaskList), | |
| 45 DMA_READ_TASKLIST); | |
| 46 | |
| 47 for (TaskPtr nextTask = tasklist->tasks; nextTask < tasklist->last(); | |
| 48 nextTask = nextTask->next()) { | |
| 49 | |
| 50 kernel[cur] = clCreateKernel(program, function, &ret); | |
| 15 | 51 |
| 50 | 52 int param = 1; |
| 53 for(int i=0;i<nextTask->inData_count;i++) { | |
| 54 ListElement *input_buf = nextTask->inData(i); | |
| 55 if (input_buf->size==0) break; | |
| 56 createBuffer(&memin[cur], param, context, mem_flag, input_buf->size, &ret); | |
| 57 ret = clEnqueueWriteBuffer(command_queue[cur], memin[cur].buf[param], | |
| 58 CL_FALSE, 0, input_buf->size, | |
| 59 input_buf->addr, 0, NULL, NULL); | |
| 60 ret = clSetKernelArg(kernel[cur], param, sizeof(memaddr), | |
| 61 (void *)&memin[cur].buf[param]); | |
| 62 param++; | |
| 63 } | |
| 64 memin[cur].size = param; // +1 means param | |
| 65 | |
| 66 for(int i = 0; i<nextTask->outData_count;i++) { // set output data | |
| 67 ListElement *output_buf = nextTask->outData(i); | |
| 68 if (output_buf->size==0) break; | |
| 69 createBuffer(&memout[cur], i, context, CL_MEM_WRITE_ONLY, output_buf->size, &ret); | |
| 70 ret = clSetKernelArg(kernel[cur], param, | |
| 71 sizeof(memaddr), (void *)&memout[cur].buf[i]); | |
| 72 param++; | |
| 73 } | |
| 74 memout[cur].size = param - memin[cur].size; | |
| 75 | |
| 76 ret = clEnqueueTask(command_queue[cur], kernel[cur], 0, NULL, NULL); | |
| 77 | |
| 78 for(int i=0;i<nextTask->outData_count;i++) { // read output data | |
| 79 ListElement *output_buf = nextTask->outData(i); | |
| 80 if (output_buf->size==0) break; | |
| 81 GpuBufferPtr mem = memout ; | |
| 82 ret = clEnqueueReadBuffer(command_queue[cur], mem[cur].buf[i0], CL_FALSE, 0, | |
| 83 output_buf->size, output_buf->addr, 0, | |
| 84 NULL,&memout[cur].event[i]); | |
| 85 } | |
| 86 cur++; | |
| 87 if (STAGE <= cur) cur = 0; | |
| 88 wait_for_event(kernel_event, memout, tasklist, cur); | |
| 89 } | |
| 90 reply = (memaddr)tasklist->waiter; | |
| 91 params_addr = (memaddr)tasklist->next; | |
| 92 } | |
| 93 | |
| 94 wait_for_event(kernel_event, memout, tasklist, cur); | |
| 95 | |
| 96 unsigned long long wait = 0; | |
| 97 (*connector->end_dmawait_profile)(&wait, &(connector->start_time), &(connector->stop_time)); | |
| 98 connector->mail_write((memaddr)MY_SPE_STATUS_READY); | |
| 99 } | |
| 100 /* NOT REACHED */ | |
| 101 } | |
| 102 | |
| 103 \end{lstlisting} | |
| 104 | |
| 105 \begin{itemize} | |
| 106 \item 4行目 : DMAManager から tasklist mail を取得 | |
| 107 \item 9行目 : 取得した mail から TaskList を read をする。 | |
| 108 TaskList に格納されている Task 全てを実行し終わるまでループする | |
| 109 \item 15行目 : 二段のパイプラインを形成するため、kernel を2つ持作る | |
| 110 \item 18行目 : パイプラインの read 部分。input データを全て kernel の引数として MemoryBuffer に書き込み、 | |
| 111 kernel に Buffer をset する | |
| 112 \item 31行目 : Output データを書き込む MemoryBuffer を用意し、kernel にset している。 | |
| 113 \item 40行目 : パイプラインの exec 部分。kernel を実行する | |
| 114 \item 42行目 : パイプラインの write 部分。kernel は実行した結果を MemoryBuffer に書き込む。 | |
| 115 MemoryBuffer に書き込まれた値をここで読み出している。 | |
| 116 \item 52行目 : パイプラインのステージの切り替えを行っている。 | |
| 117 wait\_for\_event 内で依存関係を解決したら実行の終わったステージの MemoryBuffer を delete し、次のステージへ移行する | |
| 118 \item 55行目 : 次の task を読み出し、このループを終了する | |
| 119 \end{itemize} | |
| 120 | |
| 121 DMAManager から転送されてきた Task を読み込み、Input/Output データを取り出す。 | |
| 122 データは OpenCL の API を介して GPU の MemoryBuffer に送信され、kernel が実行される。 | |
| 123 実行終了後は MemoryBuffer から戻り値を読み取り、Cerium に Output データとして返している。 | |
| 124 一連の処理は CommandQueue を介して GPU で実行される。 | |
| 125 GpuScheduler は CommandQueue を2つ持っており、二段のパイプラインが形成される。 | |
| 126 | |
| 127 全ての Task が終了すると、 | |
| 15 | 128 TaskManager 間の通信を担当する同期キューである mail を通して TaskManager に Task の終了を通知する。 |
| 50 | 129 終了が通知されると TaskManager でその TaskList に関する依存関係が解消される。 |
| 15 | 130 |
| 50 | 131 Scheduler 内で Platform や Device ID の取得、 Context の生成、 Kernel の Build と Load等も行っており、 |
| 15 | 132 並列処理したい計算のみに集中できる。 |
| 133 | |
| 7 | 134 \section{データ並列} |
| 15 | 135 並列プログラミングにおいて、明示的な並列化部分はループ部分である。 |
|
35
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
136 GPU は数百個のコアを有しており、ループ部分に対してデータ並列で処理を行うことで CPU より高速に演算を行う事ができる。 |
| 15 | 137 プログラムの大部分がループであれば、データ並列による実行だけでプログラムの性能は向上する。 |
|
35
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
138 |
| 77 | 139 OpenCL 、 CUDA ともにデータ並列をサポートしている。\cite{yuhi:2014a}\cite{kkb:2014a} |
|
35
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
140 OpenCL と CUDA はTask を実行する際にデータをどう分割するか指定し、 |
|
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
141 kernel にデータ並列用の処理を加えることで可能となる。 |
|
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
142 \ref{sec:multicore_dataparallel}節で Cerium でマルチコア CPU におけるデータ並列を可能にした。 |
|
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
143 GPGPU においてもデータ並列実行をサポートする。 |
|
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
144 GPU 上でのデータ並列実行もマルチコア CPU と変わらず、iterate API によりデータ並列用の Task を生成することができる。 |
|
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
145 iterate で Task を生成することで Scheduler が OpenCL 及び CUDA の API に適切なパラメタを渡している。 |
|
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
146 Task の生成部分は マルチコア CPU と GPU で完全に同じ形式で記述できる。 |
|
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
147 |
|
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
148 データ並列実行の際、Task は以下のように記述する。 |
|
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
149 なお、例題は multiply を用いている。 |
| 15 | 150 |
| 151 \begin{lstlisting}[frame=lrbt,label=src:multiply_opencl,caption=Multiply(OpenCL),numbers=left] | |
| 152 __kernel void | |
| 153 multiply(__global const long *params, | |
| 154 __global const float *input1, | |
| 155 __global const float *input2, | |
| 156 __global const float *output) { | |
| 157 | |
| 158 long id = get_global_id(0); | |
| 159 | |
| 160 output[id] = input1[id] * input2[id]; | |
| 161 } | |
| 162 \end{lstlisting} | |
| 163 | |
| 164 \begin{lstlisting}[frame=lrbt,label=src:multiply_cuda,caption=Multiply(CUDA),numbers=left] | |
| 165 __global__ void | |
| 166 multiply(__global const long *params, | |
| 167 __global const float *input1, | |
| 168 __global const float *input2, | |
| 169 __global const float *output) { | |
| 170 | |
| 171 int id = blockIdx.x * blockDim.x + threadIdx.x; | |
| 172 | |
| 173 output[id] = input1[id] * input2[id]; | |
| 174 } | |
| 175 \end{lstlisting} | |
| 176 | |
|
35
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
177 このような Task を分割数分生成する。 |
|
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
178 分割数は Task それぞれのフレームワークが用意している API を用いて指定する。 |
|
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
179 |
| 15 | 180 \begin{itemize} |
| 181 \item 自分の計算する範囲を取得(ソースコード\ref{src:multiply_opencl}、ソースコード\ref{src:multiply_cuda}の7行目) | |
| 182 \item 取得した範囲を計算(ソースコード\ref{src:multiply_opencl}、ソースコード\ref{src:multiply_cuda}の9行目) | |
| 183 \end{itemize} | |
|
35
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
184 いずれの Task も上記の手順で処理を行っている。 |
|
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
185 計算する範囲について、 OpenCL では取得用の API を用い、 CUDA では kernel の持つ組み込み変数から算出する。 |
|
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
186 マルチコア CPU では引数としてデータを直接渡していたが、OpenCL 、CUDA では上記の方法でメモリバッファから Load し、計算を行う。 |
|
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
187 値渡しや修飾子等若干の違いはあるが、OpenCL 、CUDA ともにマルチコア CPU(ソースコード:\ref{src:multicore_cpu}) とほぼ同じ形式で kernel を記述することができる。 |
| 15 | 188 CPU、 OpenCL、 CUDA いずれか1つの記述から残りのコードも生成できるようにする事が望ましい。 |
| 189 | |
| 190 データ並列で実行する場合、 Input と Output を各 Task 間で共有するため、少ないコピーに抑えられる。 | |
|
35
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
191 CPU ではメモリ領域を節約する事はできるが、 Task と Manager でメモリ領域が同じ(\ref{sec:shared_memory}節)なため、 |
| 15 | 192 コピーによるオーバーヘッドは少ない。 |
| 193 | |
|
35
7956856211c5
move data multicore-parallel from GPGPU-chapter to Multicore-chapter
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
194 しかし GPU は SharedMemory ではなく、データの転送がオーバーヘッドとなるため、コピーを減らす事で並列度の向上が見込める。 |