Mercurial > hg > Papers > 2018 > parusu-master
annotate paper/gpu.tex @ 48:87c3ca104689
Fix
| author | Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> |
|---|---|
| date | Tue, 06 Feb 2018 06:29:31 +0900 |
| parents | 9fbe922723e1 |
| children | 7d72512ac2e8 |
| rev | line source |
|---|---|
| 28 | 1 % TODO |
| 2 % CUDABuffer はいらないかも | |
| 3 % CUDA のアーキテクチャの図がいる? | |
| 4 % block とか grid とかの関係 | |
| 5 | |
| 29 | 6 \chapter{CUDA への対応} |
| 30 | 7 Gears OS では GPU での実行もサポートする\cite{ikkun-sigos}。 |
| 8 また、CPU、GPU の実行環境の切り替えは Meta Code Gear、つまり stub Code Gear で切り替えを行う。 | |
| 28 | 9 |
| 10 本章では、 Gears OS での CUDA 実行のサポートについて説明する。 | |
| 11 | |
| 12 \section{CUDA} | |
| 13 CUDA\cite{cuda} とは NVIDA 社が提供する GPU コンピューティング向けの総合開発環境である。 | |
| 14 | |
| 15 CUDA は演算用プロセッサ(GPU) を Device、制御用デバイス側(CPU) を Host として定義する。 | |
| 16 また、Device 上で実行するプログラムのことを kernel と呼ぶ。 | |
| 17 | |
| 18 また、CUDA には CUDA Runtime API と CUDA Driver API の2種類存在する。 | |
| 19 Driver API は Runtime API に比べて管理すべきリソースが多いが、 Runtime API より柔軟な処理を行うことが出来る。 | |
| 20 Gears OS では Driver API を用いて GPU 実行の実装を行う。 | |
| 21 | |
| 22 CUDA では処理の最小の単位を thread と定義しており、それをまとめたものを block と呼ぶ。 | |
| 23 block と thread は それぞれ3次元まで展開することが出来る。 | |
| 24 \figref{cudaArchitecture} に thread, block を2次元で展開した例を示す。 | |
| 25 | |
| 26 \begin{figure}[htbp] | |
| 27 \begin{center} | |
| 29 | 28 \includegraphics[scale=0.5]{./fig/cudaArchitecture.pdf} |
| 28 | 29 \end{center} |
| 30 \caption{blockサイズ(3,3)、 threadサイズ(3,3) に展開} | |
| 31 \label{fig:cudaArchitecture} | |
| 32 \end{figure} | |
| 33 | |
| 34 kernel を起動すると、各 thread に対して block ID と thread ID が付与される。 | |
| 35 この ID は blockIdx、threadIdx といった組み込み変数で取得できる。 | |
| 36 これらの変数は3次元のベクター型になっており、blockIdx.x とすると x座標の block ID が取得でき、 threadIdx.x とするとx座標の thread Id を取得できる。 | |
| 37 また、block 内の thread 数は blockDim という組み込み変数で取得でき、これも3次元のベクター型になっている。 | |
| 38 CUDA では これらの組み込み変数から thread が対応するデータを割り出し、データ並列の処理を行う。 | |
| 39 | |
|
32
ebcf093795f3
Add twice examples
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
31
diff
changeset
|
40 \newpage |
|
ebcf093795f3
Add twice examples
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
31
diff
changeset
|
41 |
| 28 | 42 \section{CUDAWorker} |
| 43 CUDAWorker は TaskManager から送信される CUDA用の Task を取得し、実行を行う。 | |
| 44 | |
| 45 CUDAWorker は CPUWorker と同じく初期化の際にスレッドが生成される。 | |
| 46 生成されたスレッドはCUDAライブラリ初期化や後述する CUDAExectuor の生成を行う。 | |
| 47 | |
| 48 データ並列用の Task は CUDAWorker に送信する際は Task のコピーを行わず送信する。 | |
| 29 | 49 受け取ったデータ並列用の Task は Code Gear のメタレベルで kernel の実行を行う。 |
| 28 | 50 |
| 51 \section{CUDAExectuor} | |
| 48 | 52 CUDAExectuor は \coderef{executorInterface} に示す ExecutorInterfaceを実装しており、 Host から Device へのデータの送信(read)、 kernel の実行(exec)、 Device から Host への データの書き出しを行う(write)。 |
| 30 | 53 |
| 54 \lstinputlisting[caption=executor Inteface, label=code:executorInterface]{./src/executorInterface.h} | |
| 29 | 55 |
| 56 Gears OS では データは Data Gear で表現される。 | |
| 57 つまり、Host、Device 間でデータのやり取りを行うということは Data Gear を GPU のデータ領域に沿った形に適用する必要がある。 | |
| 48 | 58 Host から Device へデータを送信する際、 CUDA では cuMemAlloc 関数を使用してサイズを指定し、Device 側のデータ領域を確保する。 |
| 30 | 59 全ての Data Gear には Meta Data Gear として Data Gear のサイズを持っており、基本的にはこのサイズでデータ領域を取ればよい。 |
| 60 しかし、Data Gear によっては内部に更にポインタで Data Gear を持っている場合がある。 | |
| 29 | 61 このような Data Gear は Data Gear の実際のサイズではなく、ポインタのサイズで計算されてしまうため、そのままでは Device 用のデータ領域を確保することができない。 |
| 62 | |
| 63 この問題を解決するために、CUDABuffer という CUDA データ送信用の Data Gear を用意した. | |
| 64 CUDABuffer には Data Gear の内部にポインタを持たない Data Gear まで展開した Input/Output Data Gear を格納される。 | |
| 65 Data Gear を CUDABuffer に格納する処理は CUDAExectuor では行わず、実行される Task の stub Code Gear で行われる。 | |
| 31 | 66 CUDABuffer に格納されている Data Gear のサイズを参照し、cuMemAlloc 関数で Device のデータ領域を確保する。 |
| 29 | 67 |
| 68 Host、Device、CUDABuffer 間の関係を\figref{cudaDataArchitecture} に示す。 | |
| 69 | |
| 70 \begin{figure}[htbp] | |
| 71 \begin{center} | |
| 72 \includegraphics[scale=0.6]{./fig/cudaDataArchitecture.pdf} | |
| 73 \end{center} | |
| 74 \caption{Host､ Device 間のデータの関係} | |
| 75 \label{fig:cudaDataArchitecture} | |
| 76 \end{figure} | |
| 77 | |
| 30 | 78 Host から Device にデータをコピーするには cuMemcpyHtoD 関数を使用して行う。 |
| 79 この際に Host で指定するデータは CUDABuffer に格納されている Data Gear となる。 | |
| 31 | 80 |
| 30 | 81 kernel の実行後、結果を Device から Host にコピーする際は cuMemcpyDtoH 関数で行われる。 |
| 35 | 82 Host のコピーされたデータは Output Data Gear も含んでいるため、 コピー後は Output Data Gear への書き出す処理に継続する。 |
| 30 | 83 |
| 29 | 84 kernel の実行はcuLaunchKernel 関数で行われる。 |
| 30 | 85 cuLaunchKernel 関数には引数として各次元のblockサイズ、thread サイズ、kernel への引数等を渡す。 |
| 48 | 86 Gears OS ではデータ並列 Task の際は Iterator Interface を持っており、 そこで指定した長さ、次元数に応じて cuLaunchKernel の引数を決定する。 |
| 30 | 87 |
| 31 | 88 % 少ないけどコードはなるべく載せたくない(メタ部分 + 複雑) |
| 29 | 89 \section{stub Code Gear による kernel の実行} |
| 30 | 90 Gears OS では stub Code Gear で CUDA による実行の切り替える。 |
| 91 | |
| 31 | 92 stub Code Gear での切り替えの際は CUDABuffer への Data の格納、実行される kernel の読み込みを行う。 |
| 93 実際に GPU で実行されるプログラムは \coderef{cudaTwice} のように記述する。 | |
| 94 | |
| 95 \lstinputlisting[caption=配列の要素を二倍にする例題, label=code:cudaTwice]{./src/cudaTwice.cu} | |
| 96 | |
| 97 stub Code Gear は通常はその stub に対応した Code Gear に継続するが、 CUDA で実行する際は CUDAExectuor の Code Gear に継続する。 |