|
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次元まで展開することが出来る。
|
|
60
|
24 \figref{cudaArchitecture} に thread、block を2次元で展開した例を示す。
|
|
28
|
25
|
|
|
26 \begin{figure}[htbp]
|
|
|
27 \begin{center}
|
|
58
|
28 \includegraphics[scale=0.7]{./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
|
|
|
40 \section{CUDAWorker}
|
|
|
41 CUDAWorker は TaskManager から送信される CUDA用の Task を取得し、実行を行う。
|
|
|
42
|
|
|
43 CUDAWorker は CPUWorker と同じく初期化の際にスレッドが生成される。
|
|
|
44 生成されたスレッドはCUDAライブラリ初期化や後述する CUDAExectuor の生成を行う。
|
|
|
45
|
|
|
46 データ並列用の Task は CUDAWorker に送信する際は Task のコピーを行わず送信する。
|
|
29
|
47 受け取ったデータ並列用の Task は Code Gear のメタレベルで kernel の実行を行う。
|
|
28
|
48
|
|
|
49 \section{CUDAExectuor}
|
|
48
|
50 CUDAExectuor は \coderef{executorInterface} に示す ExecutorInterfaceを実装しており、 Host から Device へのデータの送信(read)、 kernel の実行(exec)、 Device から Host への データの書き出しを行う(write)。
|
|
30
|
51
|
|
|
52 \lstinputlisting[caption=executor Inteface, label=code:executorInterface]{./src/executorInterface.h}
|
|
29
|
53
|
|
|
54 Gears OS では データは Data Gear で表現される。
|
|
|
55 つまり、Host、Device 間でデータのやり取りを行うということは Data Gear を GPU のデータ領域に沿った形に適用する必要がある。
|
|
48
|
56 Host から Device へデータを送信する際、 CUDA では cuMemAlloc 関数を使用してサイズを指定し、Device 側のデータ領域を確保する。
|
|
30
|
57 全ての Data Gear には Meta Data Gear として Data Gear のサイズを持っており、基本的にはこのサイズでデータ領域を取ればよい。
|
|
|
58 しかし、Data Gear によっては内部に更にポインタで Data Gear を持っている場合がある。
|
|
29
|
59 このような Data Gear は Data Gear の実際のサイズではなく、ポインタのサイズで計算されてしまうため、そのままでは Device 用のデータ領域を確保することができない。
|
|
|
60
|
|
|
61 この問題を解決するために、CUDABuffer という CUDA データ送信用の Data Gear を用意した.
|
|
|
62 CUDABuffer には Data Gear の内部にポインタを持たない Data Gear まで展開した Input/Output Data Gear を格納される。
|
|
|
63 Data Gear を CUDABuffer に格納する処理は CUDAExectuor では行わず、実行される Task の stub Code Gear で行われる。
|
|
31
|
64 CUDABuffer に格納されている Data Gear のサイズを参照し、cuMemAlloc 関数で Device のデータ領域を確保する。
|
|
29
|
65
|
|
|
66 Host、Device、CUDABuffer 間の関係を\figref{cudaDataArchitecture} に示す。
|
|
|
67
|
|
|
68 \begin{figure}[htbp]
|
|
|
69 \begin{center}
|
|
|
70 \includegraphics[scale=0.6]{./fig/cudaDataArchitecture.pdf}
|
|
|
71 \end{center}
|
|
|
72 \caption{Host､ Device 間のデータの関係}
|
|
|
73 \label{fig:cudaDataArchitecture}
|
|
|
74 \end{figure}
|
|
|
75
|
|
30
|
76 Host から Device にデータをコピーするには cuMemcpyHtoD 関数を使用して行う。
|
|
|
77 この際に Host で指定するデータは CUDABuffer に格納されている Data Gear となる。
|
|
31
|
78
|
|
30
|
79 kernel の実行後、結果を Device から Host にコピーする際は cuMemcpyDtoH 関数で行われる。
|
|
35
|
80 Host のコピーされたデータは Output Data Gear も含んでいるため、 コピー後は Output Data Gear への書き出す処理に継続する。
|
|
30
|
81
|
|
29
|
82 kernel の実行はcuLaunchKernel 関数で行われる。
|
|
30
|
83 cuLaunchKernel 関数には引数として各次元のblockサイズ、thread サイズ、kernel への引数等を渡す。
|
|
48
|
84 Gears OS ではデータ並列 Task の際は Iterator Interface を持っており、 そこで指定した長さ、次元数に応じて cuLaunchKernel の引数を決定する。
|
|
30
|
85
|
|
31
|
86 % 少ないけどコードはなるべく載せたくない(メタ部分 + 複雑)
|
|
29
|
87 \section{stub Code Gear による kernel の実行}
|
|
86
|
88 Gears OS では stub Code Gear で CUDA による実行を切り替える。
|
|
30
|
89
|
|
31
|
90 stub Code Gear での切り替えの際は CUDABuffer への Data の格納、実行される kernel の読み込みを行う。
|
|
|
91 実際に GPU で実行されるプログラムは \coderef{cudaTwice} のように記述する。
|
|
|
92
|
|
|
93 \lstinputlisting[caption=配列の要素を二倍にする例題, label=code:cudaTwice]{./src/cudaTwice.cu}
|
|
|
94
|
|
86
|
95 通常、stub Code Gear は対応した Code Gear に継続するが、CUDA で実行する際は CUDAExectuor の Code Gear に継続する。
|
|
|
96
|
