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