Mercurial > hg > Papers > 2015 > yuhi-master
annotate paper/chapter1.tex @ 19:79d16cee0afd
Benchmark gpu.minor fix abstruct
| author | Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp> |
|---|---|
| date | Fri, 13 Feb 2015 09:29:41 +0900 |
| parents | 712576635154 |
| children | f0060254db3f |
| rev | line source |
|---|---|
| 0 | 1 \chapter{既存のマルチプラットフォームフレームワーク} |
| 15 | 2 マルチプラットフォームでプログラムを動作させる場合、そのアーキテクチャを意識する必要がある。 |
| 3 マルチプラットフォームにはマルチコア CPU 、 GPU や Cell といったヘテロジニアスマルチコアのような | |
| 4 様々な構成がある。 | |
| 5 | |
|
19
79d16cee0afd
Benchmark gpu.minor fix abstruct
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
15
diff
changeset
|
6 \section{Architecture} |
| 15 | 7 本研究では、 CPU の他に GPU 上でのプログラミング (GPGPU) にも対応する。 |
| 8 | |
| 9 GPU(Graphics Processing Unit) は PC の画像処理を担当するユニットで、 | |
| 10 レンダリングに特化したプロセッサが多く集まった構造を持つ。 | |
| 11 一つ一つのプロセッサの構造は単純で、その機能は CPU に比べて限定的ではあるが | |
| 12 大量のデータを複数のプロセッサで並列処理することに長けている。 | |
| 13 | |
| 14 GPGPU (General Purpose computing on Graphics Processing Units) とは、 | |
| 15 GPU の高い演算能力を画像処理ではなく汎用計算に使用することである。 | |
| 16 | |
| 17 \section{Shared Memory} | |
| 18 \label{sec:shared_memory} | |
| 19 計算機にはメモリ空間が別の計算機と、共有メモリ(Shared Memory)な計算機がある。 | |
| 20 GPU のメモリ空間(図:\ref{fig:gpuarch})はマルチコア CPU (図:\ref{fig:cpuarch})と違い、 | |
| 21 共有メモリ(shared mermoy)でないので Host と Device 間で Data の共有ができない。 | |
| 22 そのためマルチプラットフォーム環境に対応したフレームワークには、 | |
| 23 Device と Host 間でデータの転送を行う API 備わっている。 | |
| 24 しかし、異なる Device 間でデータの転送を行うとネックになる。 | |
| 25 そのためデータの入出力を行う回数を減らす、入出力の処理をパイプライン処理にするなどの工夫が必要になる。 | |
| 26 | |
| 27 | |
| 28 \begin{figure}[htpb] | |
| 29 \begin{center} | |
| 30 \includegraphics[scale=0.4]{./images/gpu_arch.pdf} | |
| 31 \end{center} | |
| 32 \caption{GPU Architecture} | |
| 33 \label{fig:gpuarch} | |
| 34 \end{figure} | |
| 35 | |
| 36 \begin{figure}[htpb] | |
| 37 \begin{center} | |
| 38 \includegraphics[scale=0.8]{./images/cpu_arch.pdf} | |
| 39 \end{center} | |
| 40 \caption{CPU Architecture} | |
| 41 \label{fig:cpuarch} | |
| 42 \end{figure} | |
| 43 \newpage | |
| 44 %-------- | |
| 45 % OpenCL | |
| 46 %-------- | |
| 47 | |
| 0 | 48 \section{OpenCL} |
| 3 | 49 OpenCL とは、 Khronos Group の提供するマルチコア CPU と GPU といった、 |
| 50 ヘテロジニアス環境を利用した並列計算を支援するフレームワークである。 | |
| 1 | 51 |
| 15 | 52 OpenCL では演算用プロセッサ側を Device 、制御用デバイス側を Host として定義する。 |
| 53 また、 Device 上で動作するプログラムの事を kernel と呼ぶ。 | |
| 54 | |
| 55 OpenCL では、デバイスの操作に Command Queue を使用する。 | |
| 56 Command Queue は Device に命令を送るための仕組みである。 | |
| 57 Command Queue は clCreateCommandQueue という OpenCL API で作成され、 | |
| 58 Command Queue が所属するコンテキストや実行対象となる Device を指定する。 | |
| 59 | |
| 60 kernel の実行、input data への書き込み、 output data の読み込みといった | |
| 61 メモリ操作はこの Command Queue を通して行われる。 | |
| 62 | |
| 1 | 63 OpenCL には主に2つの仕様がある。 |
| 64 | |
| 65 \begin{itemize} | |
| 66 \item OpenCL C言語 | |
| 67 \item OpenCL Runtime API | |
| 68 \end{itemize} | |
| 69 | |
| 15 | 70 OpenCL C は演算用プロセッサ上で動作する、 C 言語を拡張したプログラミング言語である。 |
| 71 一方で OpenCL Runtime API は OpenCL C で記述した kernel を Queuing するために Host が利用する API である。 | |
| 1 | 72 |
| 15 | 73 Host では主に Data を input/output するメモリ資源の確保を行う。 |
| 74 OpenCL は host 側で memory buffer を作成してメモリのコピーを行う。 | |
| 1 | 75 これらの処理や Task は Command Queue に enqueue することで実行される。 |
| 13 | 76 |
| 2 | 77 多次元のデータ構造を扱う計算において高い並列度を保つには、 |
| 15 | 78 多次元データを分割して並列に実行する機能が必要である。 |
| 79 これをデータ並列実行という。 OpenCL はデータ並列実行もサポートしている。 | |
| 80 OpenCL は次元数に対応する index があり、 | |
| 2 | 81 OpenCL は一つの記述から異なる index を持つ複数の kernel を自動生成する。 |
| 1 | 82 その添字を global\_id と呼ぶ。この時入力されたデータはワークアイテムという処理単位に分割される。 |
| 83 | |
| 84 OpenCL はワークアイテムに対してそれぞれを識別する ID ( global\_id )を割り当てる。 | |
| 85 kernel は get\_global\_id API によって ID を取得し、取得した ID に対応するデータに対して処理を行い、 | |
| 86 データ並列を実現する。 | |
| 87 この ID によって取得してきたワークアイテムをグローバルワークアイテムという。 | |
| 88 また、ワークアイテムは3次元までのデータを渡すことができる。 | |
| 89 | |
| 90 データ並列による kernel 実行の場合は clEnqueueNDRangeKernel API を使用するが、 | |
| 91 この関数の引数としてワークアイテムのサイズと次元数を指定することでデータ並列で実行できる。 | |
| 92 | |
| 93 前節でワークアイテムという処理単位について述べたが、 | |
| 94 さらに複数個のグローバルワークアイテムを work\_group という単位にまとめることができる。 | |
| 95 work\_group 内では同期やローカルメモリの共有が可能となる。 | |
| 96 | |
| 97 グローバルワークアイテム(ワークアイテム全体)の個数と、 | |
| 98 ローカルワークアイテム(グループ一つ辺りのアイテム)の個数を指定することでワークアイテムを分割する。 | |
| 99 なお、このときグローバルワークアイテム数はローカルアイテム数の整数倍でなければ | |
| 100 clEnqueueNDRangeKernel API 呼び出しは失敗する。 | |
| 101 | |
| 102 ローカルアイテム数は0を指定することで、コンパイル時に最適化させることができる。 | |
| 103 したがってローカルアイテムのサイズは0を指定するのが一般的である。 | |
| 104 | |
| 105 なお、 work\_group を設定した場合は global\_id の他に work\_group\_id 、local\_id が | |
| 106 それぞれの kernel に割り当てられる(図:\ref{fig:workitem_id})。 | |
| 107 | |
| 108 \begin{figure}[htpb] | |
| 109 \begin{center} | |
| 110 \includegraphics[scale=0.65]{./images/workitem.pdf} | |
| 111 \end{center} | |
| 112 \caption{WorkItem ID} | |
| 113 \label{fig:workitem_id} | |
| 114 \end{figure} | |
| 115 | |
| 116 なお、work\_groupを設定した場合はglobal\_idの他にwork\_group\_id、local\_idが | |
| 117 それぞれのkernelに割り当てられる(図:\ref{fig:workitem_id})。 | |
| 118 | |
| 119 kernel 側からそれぞれ ID に対応した API を使用して、各 ID を取得する。 | |
| 120 取得した ID から自分が担当する index を計算して導く。 | |
| 121 表:\ref{table:kernel_id_api}は kernel 側で使用できる、 ID を取得するための API となる。 | |
| 122 \begin{tiny} | |
| 123 \begin{table}[htpb] | |
| 124 \begin{center} | |
| 125 \small | |
| 126 \begin{tabular}[htpb]{c|l} | |
| 127 \hline | |
| 128 get\_group\_id & work\_group\_id を取得 \\ | |
| 129 \hline | |
| 130 get\_local\_id & local\_id を取得 \\ | |
| 131 \hline | |
| 132 get\_global\_id & global\_id を取得 \\ | |
| 133 \hline | |
| 134 \end{tabular} | |
| 135 \caption{kernel で使用する ID 取得の API} | |
| 2 | 136 \label{table:kernel_id_api} |
| 1 | 137 \end{center} |
| 138 \end{table} | |
| 139 \end{tiny} | |
| 140 なお、 local\_id 、global\_id を取得する API は引数に0、1、2の値を set することができる。 | |
| 141 id は x, y, z 座標があり、それぞれが 0, 1, 2 に対応している。 | |
| 142 例えば get\_global\_id(1) と呼び出した場合は y 座標の、 | |
| 143 get\_global\_id(1) と呼び出した場合は z 座標の global\_id を取得する。 | |
| 0 | 144 |
| 15 | 145 |
| 146 %------ | |
| 147 % CUDA | |
| 148 %------ | |
| 2 | 149 \section{CUDA} |
| 150 CUDA とは、半導体メーカーNVIDIA社が提供するGPUコンピューティング向けの総合開発環境である。 | |
| 0 | 151 |
| 15 | 152 CUDA も OpenCL と同様、演算用プロセッサ (GPU) を Device 、制御用デバイス側を Host として定義する。 |
| 153 また、 Device 上で動作するプログラムの事も kernel と呼ぶ。 | |
| 154 | |
| 155 OpenCL における Command と CommandQueue に対応するものとして、 CUDA には Operation と Stream がある。 | |
| 156 Stream は Host 側で発行された Operation を一連の動作として Device で実行する。 | |
| 157 Operation は発行された順序で実行されることが保証されている。 | |
| 158 更に、異なる Stream に発行された Operation も依存関係が存在しない場合、Operationは並列に実行される。 | |
| 159 更に依存関係が存在しない、異なる Stream に発行された Operation は並列に実行される。 | |
| 160 | |
| 2 | 161 CUDAには主に3つの仕様がある。 |
| 162 | |
| 163 \begin{itemize} | |
| 164 \item CUDA C | |
| 165 \item CUDA Runtime API | |
| 166 \item CUDA Driver API | |
| 167 \end{itemize} | |
| 168 CUDA C は GPU 上で動作する、C 言語を拡張したプログラミング言語である。 | |
| 15 | 169 CUDA Runtime API も CUDA Driver API も CUDA C で記述した Kernel を Queueing するために |
| 170 Host が利用するAPIである。 | |
| 2 | 171 Driver API は Runtime APIに比べ、プログラマが管理しなければならないリソースが多くなる代わり、 |
| 172 より柔軟な処理を行う事ができる。 | |
| 173 | |
| 174 Stream は cuStreamCreate という Driver API で生成される。 | |
| 175 引数に Stream を指定しない API は全て host 側をブロックする同期的な処理となる。 | |
| 176 複数の Stream を同時に走らせ、 Operation を並列に実行するためには非同期的な処理を行う API を利用する必要がある。 | |
| 177 | |
| 178 CUDA では OpenCL の WorkItemに相当する単位を thread として定義している。 | |
| 179 この thread をまとめた単位として block がある。 | |
| 180 | |
| 181 CUDAでデータ並列による kernel 実行を行う場合、cuLaunchKernelAPIを使用する。 | |
| 182 この関数は引数として各座標の block 数、 | |
| 183 各座標の block 1つ辺りの thread 数を指定することによりデータ並列実行を行う。 | |
| 184 | |
| 185 cuLaunchKernel で kernel を実行すると各 thread に対して blockID と threadID が割り当てられる。 | |
| 186 CUDA には OpenCLと異なり、IDを取得するAPIが存在しない。 | |
| 187 それに代わり、 kernel に組み込み変数が準備されている。 | |
| 188 その組み込み変数を参照し、対応するデータに対し処理を行うことでデータ並列実行を実現する。 | |
| 189 組み込み変数は以下の3つである。 | |
| 190 | |
| 191 \begin{itemize} | |
| 192 \item uint3 blockDim | |
| 193 \item uint3 blockIdx | |
| 194 \item uint3 threadIdx | |
| 195 \end{itemize} | |
| 196 | |
| 197 3つの組み込み変数はベクター型で、 blockDim.x とすると x 座標の thread 数を参照することができる。 | |
| 198 同じように blockID 、 threadID の x 座標を参照することができる。 | |
| 199 blockDim.x * blockIdx.x + threadIdx.x とする事で OpenCL における get\_global\_id(0) で | |
| 200 取得できる ID に相当する値を算出する事ができる。 | |
| 201 | |
| 202 例としてある kernel で get\_global\_id(0) の値が 8 の時、 | |
| 203 CUDA では 図\ref{fig:calculateIndex}のように算出する。 | |
| 0 | 204 |
| 205 | |
| 2 | 206 \begin{figure}[htpb] |
| 207 \begin{center} | |
| 208 \includegraphics[scale=0.5]{./images/calculateIndex.pdf} | |
| 209 \end{center} | |
| 210 \caption{Calculate Index example} | |
| 211 \label{fig:calculateIndex} | |
| 212 \end{figure} | |
| 1 | 213 |
| 2 | 214 \section{StarPU} |
| 3 | 215 StarPUはフランス国立情報学自動制御研究所 (INRIA) の StarPU 開発チームの提供する、 |
| 216 ヘテロジニアス環境向けのフレームワークである。 | |
| 217 GPU の制御に OpenCL と CUDA を用いており、どちらかを選択することで GPU 上で実行することができる。 | |
| 218 | |
| 219 OpenCL と CUDA における実行の単位は kernel だったが、 StarPU では実行の単位を Task と定義している。 | |
| 220 | |
| 221 StarPU では Task を制御するためにcodeletと呼ばれる構造体を使う。 | |
| 222 codelet を Task 生成時にポインタ渡しすることで、 | |
| 223 演算を行うリソースや実行する関数等を指定することができる。 | |
| 224 CPU と GPU で並列に実行する例を\ref{src:codelet}に示す。 | |
| 225 | |
| 226 \begin{lstlisting}[frame=lrbt,label=src:codelet,caption=codeletの例,numbers=left] | |
| 227 starpu_codelet codelet = { | |
| 228 .where = STARPU_CPU|STARPU_CUDA, | |
| 229 .cpu_func = cpu_function, | |
| 230 .cuda_func = cuda_function, | |
| 231 }; | |
| 232 \end{lstlisting} | |
| 233 | |
| 234 計算に必要なデータは、 StarPU のデータプールに登録されている必要がある。 | |
| 235 StarPU ではデータを starpu\_data\_handle という型で登録する。 | |
| 236 Task はこの handle を参照することで値を参照することができる。 | |
| 237 | |
| 238 \begin{figure}[htpb] | |
| 239 \begin{center} | |
| 240 \includegraphics[scale=0.5]{./images/starpu_data_parallel.pdf} | |
| 241 \end{center} | |
| 242 \caption{StarPUにおけるデータ分割} | |
| 243 \label{fig:data_partition} | |
| 244 \end{figure} | |
| 245 | |
| 246 図:\ref{fig:data_partition}に StarPU におけるデータ並列実行の流れを示す。 | |
| 247 StarPU では配列の初期化や代入を行った後、 | |
| 248 starpu\_data\_register 関数を使って StarPU のデータプールに登録する。 | |
| 249 | |
| 250 データ並列で実行する場合、更にデータを分割する必要がある。 | |
| 251 starpu\_data\_partition 関数を用いる事で分割を行うことができる。 | |
| 252 分割数を指定することで、データプールに登録したデータを chunk と呼ばれる単位に分割する。 | |
| 253 starpu\_task\_submit 関数により chunk を CPU や GPU に割り当てることができる。 |