|
0
|
1 \chapter{既存のマルチプラットフォームフレームワーク}
|
|
|
2 \section{OpenCL}
|
|
1
|
3 OpenCL とは、マルチコア CPU と GPU のようなヘテロジニアスな環境を利用した並列計算を支援するフレームワークである。
|
|
|
4
|
|
|
5 OpenCL には主に2つの仕様がある。
|
|
|
6
|
|
|
7 \begin{itemize}
|
|
|
8 \item OpenCL C言語
|
|
|
9 \item OpenCL Runtime API
|
|
|
10 \end{itemize}
|
|
|
11 OpenCL C は演算用プロセッサ上で動作する、 C 言語を拡張したプログラミング言語である。
|
|
|
12 一方で OpenCL Runtime API は OpenCL C で記述したプログラムを演算用プロセッサ上で実行させるため、
|
|
|
13 制御用のプロセッサが利用する API である。
|
|
|
14
|
|
|
15 OpenCL では演算用プロセッサ側を device 、制御用デバイス側を host として定義する。
|
|
|
16 また、 Device 上で動作するプログラムの事を kernel と呼ぶ。
|
|
|
17
|
|
|
18 \subsection{Command Queue}
|
|
|
19 OpenCL では、デバイスの操作に Command Queue を使用する。 Command Queue は Device に命令を送るための仕組みである。
|
|
|
20 Command Queue は clCreateCommandQueue という OpenCL API で作成され、
|
|
|
21 Command Queueが所属するコンテキストや実行対象となるデバイスを指定する。
|
|
|
22
|
|
|
23 kernel の実行、input data への書き込み、 output data の読み込みといった
|
|
|
24 メモリ操作はこの Command Queue を通して行われる。
|
|
|
25
|
|
|
26 \subsection{メモリアクセス}
|
|
|
27 host では主に data を input/output するメモリ資源の確保を行う。
|
|
|
28 GPU のメモリ空間(図:\ref{fig:gpuarch})はマルチコア CPU (図:\ref{fig:cpuarch})と違い、
|
|
|
29 共有メモリでないため host と kernel(Task)間で data の共有ができない。
|
|
|
30
|
|
|
31 \begin{figure}[htpb]
|
|
|
32 \begin{center}
|
|
|
33 \includegraphics[scale=0.4]{./images/gpu_arch.pdf}
|
|
|
34 \end{center}
|
|
|
35 \caption{Gpu Architecture}
|
|
|
36 \label{fig:gpuarch}
|
|
|
37 \end{figure}
|
|
|
38
|
|
|
39 \begin{figure}[htpb]
|
|
|
40 \begin{center}
|
|
|
41 \includegraphics[scale=0.8]{./images/cpu_arch.pdf}
|
|
|
42 \end{center}
|
|
|
43 \caption{Cpu Architecture}
|
|
|
44 \label{fig:cpuarch}
|
|
|
45 \end{figure}
|
|
|
46 アクセスするにはメモリ空間間でコピーしなければならない。
|
|
|
47
|
|
|
48 OpenCLは host 側で memory buffer を作成してメモリのコピーを行う。
|
|
|
49 これらの処理や Task は Command Queue に enqueue することで実行される。
|
|
|
50 \subsection{データ並列}
|
|
|
51 多次元のデータ構造がある場合に高い並列度を保つには、それを分割して並列に実行する機能が必要である。
|
|
|
52 データ並列実行という。OpenCLはデータ並列実行もサポートしている。
|
|
|
53 OpenCL は次元数に対応する index があり、 OpenCL は一つの記述から異なる index を持つ複数の kernel を自動生成する。
|
|
|
54 その添字を global\_id と呼ぶ。この時入力されたデータはワークアイテムという処理単位に分割される。
|
|
|
55
|
|
|
56 OpenCL はワークアイテムに対してそれぞれを識別する ID ( global\_id )を割り当てる。
|
|
|
57 kernel は get\_global\_id API によって ID を取得し、取得した ID に対応するデータに対して処理を行い、
|
|
|
58 データ並列を実現する。
|
|
|
59 この ID によって取得してきたワークアイテムをグローバルワークアイテムという。
|
|
|
60 また、ワークアイテムは3次元までのデータを渡すことができる。
|
|
|
61
|
|
|
62 データ並列による kernel 実行の場合は clEnqueueNDRangeKernel API を使用するが、
|
|
|
63 この関数の引数としてワークアイテムのサイズと次元数を指定することでデータ並列で実行できる。
|
|
|
64
|
|
|
65 \subsection{ワークグループ}
|
|
|
66 前節でワークアイテムという処理単位について述べたが、
|
|
|
67 さらに複数個のグローバルワークアイテムを work\_group という単位にまとめることができる。
|
|
|
68 work\_group 内では同期やローカルメモリの共有が可能となる。
|
|
|
69
|
|
|
70 グローバルワークアイテム(ワークアイテム全体)の個数と、
|
|
|
71 ローカルワークアイテム(グループ一つ辺りのアイテム)の個数を指定することでワークアイテムを分割する。
|
|
|
72 なお、このときグローバルワークアイテム数はローカルアイテム数の整数倍でなければ
|
|
|
73 clEnqueueNDRangeKernel API 呼び出しは失敗する。
|
|
|
74
|
|
|
75 ローカルアイテム数は0を指定することで、コンパイル時に最適化させることができる。
|
|
|
76 したがってローカルアイテムのサイズは0を指定するのが一般的である。
|
|
|
77
|
|
|
78 なお、 work\_group を設定した場合は global\_id の他に work\_group\_id 、local\_id が
|
|
|
79 それぞれの kernel に割り当てられる(図:\ref{fig:workitem_id})。
|
|
|
80
|
|
|
81 \begin{figure}[htpb]
|
|
|
82 \begin{center}
|
|
|
83 \includegraphics[scale=0.65]{./images/workitem.pdf}
|
|
|
84 \end{center}
|
|
|
85 \caption{WorkItem ID}
|
|
|
86 \label{fig:workitem_id}
|
|
|
87 \end{figure}
|
|
|
88
|
|
|
89 なお、work\_groupを設定した場合はglobal\_idの他にwork\_group\_id、local\_idが
|
|
|
90 それぞれのkernelに割り当てられる(図:\ref{fig:workitem_id})。
|
|
|
91
|
|
|
92 kernel 側からそれぞれ ID に対応した API を使用して、各 ID を取得する。
|
|
|
93 取得した ID から自分が担当する index を計算して導く。
|
|
|
94 表:\ref{table:kernel_id_api}は kernel 側で使用できる、 ID を取得するための API となる。
|
|
|
95 \begin{tiny}
|
|
|
96 \begin{table}[htpb]
|
|
|
97 \begin{center}
|
|
|
98 \label{table:kernel_id_api}
|
|
|
99 \small
|
|
|
100 \begin{tabular}[htpb]{c|l}
|
|
|
101 \hline
|
|
|
102 get\_group\_id & work\_group\_id を取得 \\
|
|
|
103 \hline
|
|
|
104 get\_local\_id & local\_id を取得 \\
|
|
|
105 \hline
|
|
|
106 get\_global\_id & global\_id を取得 \\
|
|
|
107 \hline
|
|
|
108 \end{tabular}
|
|
|
109 \caption{kernel で使用する ID 取得の API}
|
|
|
110 \end{center}
|
|
|
111 \end{table}
|
|
|
112 \end{tiny}
|
|
|
113 なお、 local\_id 、global\_id を取得する API は引数に0、1、2の値を set することができる。
|
|
|
114 id は x, y, z 座標があり、それぞれが 0, 1, 2 に対応している。
|
|
|
115 例えば get\_global\_id(1) と呼び出した場合は y 座標の、
|
|
|
116 get\_global\_id(1) と呼び出した場合は z 座標の global\_id を取得する。
|
|
0
|
117
|
|
|
118
|
|
|
119
|
|
|
120
|
|
|
121
|
|
1
|
122
|
|
|
123
|
