Mercurial > hg > Papers > 2015 > nozomi-sigos
comparison paper/cuda.tex @ 0:0127effb8fcd
first commit
author | Nozomi Teruya <e125769@ie.u-ryukyu.ac.jp> |
---|---|
date | Tue, 05 May 2015 15:36:41 +0900 |
parents | |
children |
comparison
equal
deleted
inserted
replaced
-1:000000000000 | 0:0127effb8fcd |
---|---|
1 \section{CUDA} | |
2 CUDA とは、半導体メーカー NVIDIA 社が提供する GPU コンピューティング向けの総合開発環境でコンパイラ、ライブラリ、デバッガなどから構成される。プログラミング言語である CUDA C は C 言語ベースに拡張を加えたものである。 | |
3 | |
4 CUDA には CUDA Runtime API と CUDA Driver API の2種類がある。 | |
5 Driver API は Runtime API に比べてプログラマが管理すべきリソースが多い。 | |
6 しかし、Runtime API より柔軟な処理を行うことができる。 | |
7 今回は Driver API を使用して実装した。 | |
8 | |
9 CUDA も OpenCL と同様に、制御を行う CPU 側を host、GPU 側を device と定義している。 | |
10 また、device 上で動作するプログラムも OpenCL と同様に kernel と呼ぶ。 | |
11 | |
12 \subsection{Stream} | |
13 CUDA には OpenCL の CommandQueue と似たような仕組みとして Stream がある。 | |
14 Stream は host 側で発行された Operation を一連の動作として device で実行する。 | |
15 Stream に発行された Operation は発行された順序で実行されることが保証されている。 | |
16 異なる Stream に発行された Operation に依存関係が存在しない場合、Operation を並列に実行することができる。 | |
17 | |
18 Stream は cuStreamCreate という Driver API で生成される。 | |
19 引数に Stream を指定しない API はすべて host 側をブロックする同期的な処理となる。 | |
20 複数の Stream を同時に走らせ Operation を並列に実行するためには非同期な処理を行う API を利用する必要がある。 | |
21 | |
22 \subsection{Data Parallel Execution} | |
23 CUDA では OpenCL の WorkItem に相当する単位を thread と定義している。 | |
24 この thread をまとめたものを block と呼ぶ。 | |
25 CUDA でデータ並列による kernel 実行をする場合、cuLaunchKernel API を使用する。 | |
26 この関数は引数として各座標の block 数と各座標の block 1つ当たりの thread 数を指定することでデータ並列で実行できる。 | |
27 | |
28 cuLaunckKernel で kernel を実行すると各 thread に対して block ID と thread ID が割り当てられる。 | |
29 CUDA には OpenCL とは異なり、ID を取得する API は存在しない。 | |
30 代わりに、kernel に組み込み変数が準備されており、それを参照し、対応するデータに対し処理を行うことでデータ並列を実現する。 | |
31 組み込み変数は以下の通りである。 | |
32 | |
33 \begin{itemize} | |
34 \item uint3 blockDim | |
35 \item uint3 blockIdx | |
36 \item uint3 threadIdx | |
37 \end{itemize} | |
38 | |
39 各組み込み変数はベクター型で、blockDim.x とすると x 座標の thread 数を参照することができる。 | |
40 blockIdx.x とすると x 座標の block ID が参照でき、threadIdx.x とすると x 座標の thread ID を参照することができる。 | |
41 blockDim.x * blockIdx.x + threadIdx.x で OpenCL の get\_global\_id(0) で取得できる ID に相当する ID を算出することができる。 | |
42 例として、ある kernel で get\_global\_id(0) の返り値が13の場合、CUDA では図:\ref{fig:calculate_index}のようになる。 | |
43 | |
44 \begin{figure}[!h] | |
45 \begin{center} | |
46 \includegraphics[scale=0.4]{./images/culculate_index.pdf} | |
47 \end{center} | |
48 \caption{Calculate Index} | |
49 \label{fig:calculate_index} | |
50 \end{figure} |