これからの並列計算のための GPGPU 連載講座 (II)GPGPU プログラミング環境 CUDA 入門編
大島聡史
東京大学情報基盤センター

1

本編の構成
連載第二回である今回は、GPGPU プログラミング環境 CUDA(CUDA Unified Device Archi-

tecture) について紹介する。
CUDA は NVIDIA 社製 GPU 向けのプログラミング環境であり*1 、C/C++ 言語を元に独自の拡張を行った専用の言語および対応するコンパイラ (nvcc) と実行時ランタイムライブラリ、そし て い く つ か の 数 値 計 算 ラ イ ブ ラ リ か ら 構 成 さ れ て い る 。CUDA は 現 在 NVIDIA 社 製 GPU に対して最も低いレイヤーでアクセスすることが可能なプログラミング環境であり、適切に利用することで GPU の持つ性能を引き出すことができる。逆に使い方を誤ると低い性能しか得ることができない。
今回述べる内容は以下の通りである：• CUDA の導入方法• CUDA プログラミング入門編 1• CUDA のアーキテクチャ概要• CUDA プログラミング入門編 2
CUDA は NVIDIA 社 製 GPU 向 け の プ ロ グ ラ ミ ン グ 環 境 で は あ る も の の 、一 般 的 な PC 環
境—Windows・Linux・MaxOSX—さえあれば CUDA に対応した GPU が搭載されていなくてもプログラムを作成し動作確認をすることが可能である。是非とも実際にプログラムを作成して動かしてみて欲しい。（もちろん、実機でないと評価できないことや実機でないと生じないト
ラブルなどがあるため、なるべく実機で試して欲しい。）

CUDA プログラミングの助けになる資料やサイトについては、NVIDIA 社の web サイト CUDA
Zone(http://www.nvidia.com/object/cuda home new.html) に 多 く の 資 料 が 公 開 さ れ て い る う
えに、「“CUDA” “プログラミング”」などのキーワードで web 検索すれば日本語で多数の資料を入手可能である。また、後述する CUDA Toolkit をインストールすればいくつかのドキュメントが利用可能になる。特にプログラミングガイドは、チュートリアルをはじめとしてアーキ
テクチャの解説や主な API とその使い方が掲載されており重宝する。(日本語翻訳版も公開されているものの、英語版と比べて公開が遅いことが多いため、最新情報を参照したい場合には注意が必要である。) プログラミングガイド以外に有用な資料や web サイトとしては以下のものが挙げられる：

NVIDIA Forums (NVIDIA フォーラム) NVIDIA 社の公開しているフォーラム。現状では日本語フォーラムは初歩的なプログラミングの質問が多く、英語フォーラムはより深い話題が多*1

名前の示すとおり、プログラミング環境というよりはアーキテクチャそのものであるとも言えるが、ここではこだわらないことにする。

1

くポストされている傾向がある。(http://forums.nvidia.com/, http://forum.nvidia.co.jp/)
はじめての CUDA プログラミング

日本語で書かれた CUDA プログラミングの解説本。書籍の形式をとっているため CUDA の更新に追従できないという弱点はあるものの、日本語で丁寧に解説された本であり、タイトル通り CUDA をはじめて使う人でもそうでない人でも一読の価値がある。(http://www.kohgakusha.co.jp/books/detail/978-4-7775-1477-9)

NVIDIA CUDA Information Site フ ィ ッ ク ス タ ー ズ 社 の 有 志 が 運 営 す る wiki サ イ ト 。
(http://gpu.fixstars.com/)
今 回 の 記 事 に お い て は 特 に 断 り が な い 限 り CUDA に 対 応 し た NVIDIA 社 製 GPU の こ と を 単に GPU と呼ぶことにする。ただし、GPU そのものにもいくつかのバージョンが存在している。
本記事の一部の記述は最新バージョン (Compute Capability 1.3) の GPU にのみ該当する場合があることを断っておく。また、今回述べる内容は CUDA Toolkit version 2.3 を対象としている。

CUDA は比較的頻繁にバージョンアップが行われており、バージョンアップのたびに API の変
更などが起きるため、注意していただきたい。
な お 、本 連 載 中 で 扱 う プ ロ グ ラ ム 類 は 筆 者 の web サ イ ト (http://www.cspp.cc.u-

tokyo.ac.jp/ohshima/) にて公開予定である。

2

CUDA の導入方法
まずは CUDA の導入方法を簡単に説明する。CUDA を利用するには、

1. CUDA 対応ドライバ (CUDA 対応 GPU を利用する場合のみ)
2. CUDA Toolkit
3. CUDA SDKが必要である。いずれも NVIDIA 社の CUDA ダウンロードサイト(http://developer.nvidia.com/object/cuda download.html) にて無償で公開されている。画像処
理ライブラリやプロファイラ等も同様に入手することができる。現在、Windows 向け、Linux
向け、MaxOSX 向けのツールが公開されており、具体的なダウンロード・インストール手順はOS ごとに異なっているものの、いずれも実行環境に対応した実行形式ファイルやスクリプトを実行させるのみの容易な手順で導入することができる。なお、実際に開発を行うためにはそれぞれ開発環境 (VisualStudio や gcc など) が必要となる。さらに、Linux に CUDA ドライバを導入するためにはカーネルのソースコードなども必要となる。各ディストリビューションのパッ
ケージシステムなどを用いることで導入可能なものばかりなので、必要に応じて各自導入していただきたい。
ド ラ イ バ や 開 発 環 境 の イ ン ス ト ー ル が 完 了 し た 後 は 、サ ン プ ル プ ロ グ ラ ム を 実 行
し て み る と 良 い 。サ ン プ ル プ ロ グ ラ ム は SDK に 付 属 し て お り 、た と え ば Linux で あれ ば デ フ ォ ル ト で$HOME/NVIDIA GPU Computing SDK/C/bin/linux/release (た だ し$HOME/NVIDIA GPU Computing SDK/C にて make コマンドを実行する必要がある) 以下に配置されるはずである。サンプルプログラムはソースコードが公開されており、解説ドキュメ
ントが付属しているものもあるため、大いに参考になるだろう。なおサンプルの中には、CUDAに対応した実機の GPU が搭載されていないと実行できないもしくは実行できても著しく性能が2低くなるものや、GUI が利用できないと実行できないものもあるので注意していただきたい。
次章以降では CUDA を用いたプログラムの記述方法や実行結果についていくつか紹介していくが、実行環境としてはスパコン利用者にも馴染みが深いと思われる Linux を用いることにする。基本的に、実行環境 (OS) によって大きく異なる点は無いが、異なる環境で実行する場合には適宜読み替えていただきたい。


メーカー製 PC に対する GPU ドライバの導入について

一部メーカー製の PC や多くのノート PC では、NVIDIA 社の提供する最新の GPU ドライバを導入できない (そのため、最新版の CUDA を利用できない、もしくは CUDA そのものを利用できない) ことがあるので注意が必要である。これはメーカーが独自にカスタマイズ
したドライバを必要とすることがあるためであり、メーカーによってはサポート web サイ
トにてドライバ更新用のプログラムなどを公開していることもあるので、チェックしてみると良いだろう。
また、メーカーのサポートが外れるなどのリスクはあるものの、一部の機種については何らかの手段でドライバを更新できる可能性がある。例えば、PC の機種名と CUDA、もしくは PC の機種名とビデオドライバなどのキーワードを用いて web 検索を行うと、ドライ
バを更新する方法が公開されていることがある。参考にすると良いだろう。(実際に導入する場合には自己責任となるので注意されたい。)



3



CUDA プログラミング入門編 1
続いて、CUDA を用いた初歩的なプログラムの例を示し、CUDA プログラムの基本的な動作の流れおよびプログラム作成方法について解説する。CUDA のアーキテクチャ等は次章で解説することにして、まずはどのようなプログラムを書くとどのように動くのか、雰囲気をつかん
でいただきたい。
なお、CUDA には簡潔な記述で GPU を利用することができる RuntimeAPI と、より低レイヤー
で GPU を細かく制御可能な DriverAPI がある。今回はプログラムの作成が容易な RuntimeAPIを対象として解説を行う。

3.1 CUDA プログラムの例と実行の流れ
まずは、図 1 および図 2 を見ていただきたい。図 1 は、C 言語で記述した単純な配列加算プロ
グラムのソースコードおよび同じ内容を CUDA を用いて記述したソースコードである。また図2 にはソースコードをコンパイル・リンクし実行する手順および実行結果を示している。なお、
CUDA プログラムの拡張子には cu を用いることになっている。
図 1 からは CUDA プログラムは確かに CPU プログラムと比べて様々な記述が必要ではあるものの、特別に難解で手間のかかるプログラムではないことがおわかりいただけるだろうか。(ちなみに余談ではあるが、CUDA が登場する以前のプログラマブルシェーダを用いたプログラム
では本プログラムの 3 倍程度の記述と画像処理プログラミングに関する知識が要求されていた。
そのうえ、直感的な記述でもなかった。) CUDA プログラムには global という関数の接頭辞
や <<< と >>> を用いた記述、cuda で始まる関数群のように、既存の CPU 向けの C 言語プログ3

図 1 C 言語 (上) と CUDA(下) の配列加算プログラム例

4

C 言語でのコンパイルと実行
> ls
cpu.c
> gcc -O3 cpu.c> ls
a.out cpu.c> ./a.out
CPU:
InA: 0.30 0.60 0.70 0.50 0.30 0.50 0.60 0.20 0.90 0.10 0.20 0.70 (以下省略)
InB: 0.00 0.60 0.40 0.60 0.20 0.50 0.80 0.60 0.20 0.80 0.40 0.70 (以下省略)
Out: 0.30 1.20 1.10 1.10 0.50 1.00 1.40 0.80 1.10 0.90 0.60 1.40 (以下省略)>

CUDA でのコンパイルと実行> ls
gpu.cu> nvcc -O3 gpu.cu -I${HOME}/NVIDIA_GPU_Computing_SDK/C/common/inc> lsa.out gpu.cu> ./a.out
GPU:
InA: 0.30 0.60 0.70 0.50 0.30 0.50 0.60 0.20 0.90 0.10 0.20 0.70 (以下省略)
InB: 0.00 0.60 0.40 0.60 0.20 0.50 0.80 0.60 0.20 0.80 0.40 0.70 (以下省略)
Out: 0.30 1.20 1.10 1.10 0.50 1.00 1.40 0.80 1.10 0.90 0.60 1.40 (以下省略)>

図2

配列加算プログラムの実行手順および実行結果

ラムにはない特徴的な記述が含まれていることがわかる。また図 2 からは、CUDA プログラムもCPU プログラムと同様にソースコードを専用のコンパイラ (nvcc) でコンパイル・リンクすることで実行ファイルが生成されること、生成されたファイルを実行すれば GPU 上でプログラムが動作すること (a.out 以外に特殊なファイルを生成したり必要としたりはしないこと) がわかる。

CUDA プログラムにおいては、図 1 中に記したように、プログラムに記述された処理の全てが GPU 上で実行されるわけではない。むしろ明示的に指示された部分以外は CPU 上で実行される。この「明示的な指示」を行うのに用いるのが global などの「関数に対する指示子」と、<<< と >>> を 用 い た「GPU 呼 び 出 し 記 述 」で あ る 。CUDA コ ン パ イ ラ nvcc は こ れ ら の 指 示
子・呼び出し記述を目印としてソースコードを解析・分離し、CPU によって実行される部分はCPU 向け、GPU によって実行される部分は GPU 向けのコンパイルを行い、それぞれを結合した最終的な実行可能プログラムを出力する。
「関数に対する指示子」と記したが、CUDA において GPU に実行させる処理の単位は関数である。これを「GPU カーネル」や「GPU カーネル関数」、もしくは単に「カーネル」などと呼ぶ。計算量が大きいため計算時間が長くなおかつ並列度が高い部分を関数として抽出し、抽出した関数を GPU に実行させることでプログラム全体の実行時間を削減する、というのが CUDAの 基 本 的 な 戦 略 と な る 。(こ れ は CUDA の み な ら ず 、ClearSpeed Advanced Accelerator や Cell

BE における SPE など、「アクセラレータ」と呼ばれるハードウェアを用いる場合の基本的な戦略であると言える。)

CUDA における「関数に対する指示子」には以下の 3 つがある。複数の指示子を指定するこ
5

とも可能である。使い分ける上で特に難しい点はなく、間違えて記述した場合には nvcc に明確に指摘されるだろう。

globalCPU に呼び出されて GPU 上で実行される関数

deviceGPU に呼び出されて GPU 上で実行される関数

host

CPU に呼び出されて CPU 上で実行される関数 (global/device 関数と同名の関数を CPU
上でも実行したい場合に使用する)

また、CPU と GPU は独立したメモリを持つ。CUDA においては CPU から GPU 上のメモリに対する読み書きには制限があり、GPU から CPU 上のメモリを操作することはできない。CPUから GPU に対するメモリの読み書き、すなわち CPU-GPU 間のデータ転送については、カーネル関数の前後で API（cudaMemcpy などの API 関 数）を 用いて行う必要が ある。これは、MPIのようにデータの送受信両方を明示的に記述するプログラミングとも、OpenMP のようにデー
タの送受信が不要なものとも異なっている。他の環境で並列化プログラミングを行ってきた経験がある場合には特に注意されたい。
次に cuda から始まる API 関数について簡単に説明する。今回利用している API は以下の通りである。

cudaSetDevice 使用する GPU の ID を指定する関数
cudaMalloc GPU 上のメモリを確保する関数 (以下の関数では本関数によって得られたポイン
タを用いる)

cudaMemcpy CPU-GPU 間のデータ転送を行う関数 (CPU から GPU と GPU から CPU の両方で同じ関数を使用し転送方向は第 4 引数で制御する)

cudaFree 確保した GPU 上のメモリを解放する関数
これらの関数以外にも、CUDA により様々な関数が提供されている。詳細についてはサンプルプログラムやプログラミングガイドを参照していただきたい。

6

nvcc のコンパイルオプションやその他の機能について

本節では特に断りなく nvcc を gcc と同様のオプションを用いて実行した。しかし、nvcc は

gcc と同じオプションに対応しているわけではない。例えば gcc を使用する際に良く用いられる-Wall オプション (全ての警告を表示するオプション) には対応していない。

nvcc には様々なオプションと機能がある。たとえば-cubin や-ptx といったオプションを利用すると、GPU によって実行される部分のみをコンパイルしバイナリ形式や中間表現 PTX
形式として得ることができる。(これらの機能は RuntimeAPI を利用する場合には使わない
ため、今回は解説しない。)
頻繁に利用するオプションとしては、対象とする GPU の世代を指定する-arch オプションがある。CUDA は世代 (Compute Capability) によって対応する機能などに違いがあり、新しい GPU 向けの機能を使用するプログラムを古い世代の GPU で実行することができない。
各 GPU にどの arch オプションを指定すれば良いかについては、プログラミングガイドを参照していただきたい。

nvcc が対応しているオプションの一覧は-h オプションによって得ることができる。興味があれば確認してみて欲しい。



3.2 CUDA における並列処理
前節では CPU から GPU を操作する方 法や CPU-GPU 間のデータ転送について確 認した。つ
づいて本節では GPU カーネルに注目してみることにする。
前節で用いた図 1 の主要な計算部分を比較すると、CPU プログラムである前者にはループ構造が含まれているのに対して、GPU プログラムである後者にはループ構造が含まれていないにも関わらず配列の各要素に対する計算が行えていた。CUDA プログラムにおいては、GPU
上の多数のプロセッサそれぞれにおいて同一の GPU カーネルが動作する。実行される各イン
スタンスは個別の ID を持つため、ID を元に計算対象のデータを特定し変更することができる(図 3)。動作のイメージとしては、pthread のようにスレッド毎に個別のプログラムを実行する
並列化よりは、rank を用いて処理を分ける MPI にやや近いと言えるだろう。

図 3 CUDA における ID を用いた並列処理

CUDA では各インスタンスに threadId と BlockId という二階層の ID が割り当てられる。また、ID の範囲、すなわち CPU 向けの並列化プログラミングで言うところのスレッド数やプロセス数については、GPU カーネルを呼び出す際に指定する。図 1 における <<< と >>> の間に記述

7

さ れ て い る 値 が こ の 指 定 に 該 当 す る 。図 1 の 例 で は 、16*16=256 の イ ン ス タ ン ス が 生 成 さ れ 、GPU 上のプロセッサに割り当てられて実行されることになる。そのため、長さ 256 の配列に対する計算が明示的なループ記述無しに実行されているわけである。

CUDA におけるインスタンスの割り当ては、既存の CPU 向け並列化プ ログラミングと比べるとやや複雑である。これについては GPU のハードウェアアーキテクチャについて知らなくては理解が困難であるため、次章で説明する。

4

CUDA のアーキテクチャ
本章では、CUDA のハードウェアモデル、実行モデル、メモリモデルについて解説する。前章で紹介した CUDA プログラムの基本的な構 造や CUDA プログ ラミン グの 手順と照ら し合わ
せて理解して欲しい。

4.1 物理的なハードウェア構成
はじめに、CUDA が対象としている NVIDIA 社製 GPU の中でも 2010 年 2 月現在の最新アーキ
テクチャである GT200 アーキテクチャ (以下 GT200) におけるハードウェア構成を示す。
図4はGPUにおける演算器とメモリの構成である。基本的な構成としては、ScalarProcessor(SP)と呼ばれるシンプルなプロセッサが 8 つ集まって MultiProcessor(MP) を構成しており、MP がGPU のグレードによって複数個搭載されている。さらに、MP 毎に独立したメモリ (同一 MP 内の SP でのみ共有されるメモリ) と MP 間で共有されるメモリが搭載されている。各メモリの特徴や使い方については 4.3 節で紹介する。このように、演算器とメモリが階層性を持っているのが GPU の特徴である。(ちなみに、NVIDIA 社製の旧世代 GPU や AMD 社製 GPU も演算器とメモリに階層性を持つアーキテクチャを採用している。)

図 4 NVIDIA 社製 GPU における演算器とメモリの構成

GPU に は 1GPU あ た り 最 大 30 の MP が 搭 載 さ れ て い る た め 、SP 数 は 最 大 で 240 に 達 す る 。
(TeslaS1070 は 1 筐体に 4GPU が搭載されているハードウェアであり、1GPU あたりの SP 数はやはり 240 である。) さらに後述するが、GPU は CPU と異なり物理コア数よりも高い並列度での並 列 処 理 に 適 し た 設 計 と な っ て お り 、1GPU あ た り で は 1000 を 超 え る 並 列 性 に ま で 性 能 が ス
ケールする。ただし、SP はマルチコア CPU における CPU コアのように独立して演算を行える
わけではない。既存の CPU と比較するのであれば、SP は SIMD コアであり、CPU コアに対応す

8

るのは MP ということになる。コア数が 30 ある CPU はもちろん並列度が高い CPU であると言えるものの、単純にコア数が 240 ある CPU と考えるのは現実にそぐわないため注意が必要である。
さて、SP は SIMD コアであると述べたが、実際に同一 MP 上の SP は同時に異なる演算を行うことができない。複数の SP がそれぞれ異なるデータに対して同じ演算を行うデータ並列処理
が CUDA における基本的な並列処理である。また、SP はインオーダーであり分岐予測器も持たないシンプルな計算コアである。そのため、SP は同一クロックの CPU コアと比べると演算性能が低く、GPU は多数の SP を活用可能な並列度の高い問題でなくては高い性能を得ることができない。
一方で MP 単位での並列処理を考えた場合は、異なる MP 間で同時に異なる演算を行うことができるため、データ並列処理にこだわる必要はない。ただし、MP 単位での並列処理では同期をとれる範囲に気をつけなくてはならない。同一 MP 内の SP 同士では容易に同期をとることができる (カーネル関数内で syncthreads 関数を呼び出すだけで同期をとることができる) 一方
で、異なる MP 間の SP 同士で同期をとるためには一度 GPU カーネルを終了して CPU に制御を戻さなくてはならない。GPU 上の共有メモリに対して排他的な演算を行うことができる atomic関数を用いて同期を実現することも不可能ではないが、性能が低下する可能性が大きい。

4.2 物理的な構成と実行モデル
GPU に搭載されている演算器は SP と MP の階層性を持つことを述べた。SP は 8 個ごとに MPを構成しており、MP が GPU のグレードによって複数個搭載されているのが GT200 の物理的な構成である。CPU が CPU コア数以上のインスタンス（スレッドやプロセス）を割り当てられると時分割実行を行うのと同様に、GPU も SP 数や MP 数以上のインスタンスを割り当てられると時分割実行を行うことになる。CPU が CPU コア数よりも多くのインスタンスを実行しようとした場合、一般的には CPU コア数以下のインスタンス数の場合と比べて性能が低下する。
一方で GPU は、コンテキストスイッチのコストが非常に低く、GPU にハードウェア実装されているスケジューラはメモリを読み書きする際に SP が待つ必要がある場合に積極的にコンテ
キストスイッチを行う。そのため、物理的な SP 数や MP 数よりも多くのインスタンスを生成した方がメモリアクセスのレイテンシが隠蔽されて高い性能を得ることができる。

CUDA においては SP 単位のジョブをスレッド (Thread)、MP 単位のジョブをブロック (Block)
もしくはスレッドブロック (Thread Block)、さらにカーネル実行単位をグリッド (Grid) と呼ぶ(図 5)。図中ではいずれも一次元で表現しているが、実際にはブロックは二次元、スレッドは三次 元 の 空 間 を 割 り 当 て る こ と が で き る 。た だ し 、同 一 の グ リ ッ ド 内 に お い て ブ ロ ッ ク ご と に
スレッド数を変更することはできない。同一ブロック内のスレッドは全て同じ MP に割り当てられ、また各ブロックはコンテキストスイッチが行われても異なる MP へと割り当てられることはない。どのブロックがどの MP に割り当てられるかをプログラマが制御することはできない。一般的にスレッド数は 128 程度以上割り当てると良いとされている。ただし、各スレッドがどれだけの資源 (レジスタや共有メモリ) を使うかなどにより最適な値は異なるため、最大の性能を得るためにはアプリケーション毎に数を調整して性能を測定してみる必要がある。

9

図5

SP・MP とスレッド・ブロック

4.3 メモリの分類と使い分け
最後に、GPU 上に搭載されたメモリについて解説する。
図 6 に GPU におけるメモリの分類を示す。GPU には複数種類のメモリが搭載されており、それぞれ異なる特性を備えている。以下に各メモリの概要を示す。

Registers 各 MP ごとに独立して搭載されているレジスタ。スレッド毎に独立したレジスタとして利用可能。GT200 では MP 毎に 16384 本。高速で低レイテンシ。デバイス関数内で宣言された通常の変数は基本的にレジスタへ割り当てられる。CPU から直接値を読み書きすることはできない。

SharedMemory 各 MP ごとに独立して搭載されているメモリ。同一 MP 内の各スレッドからは共有メモリとして利用可能。GT200 では各 MP 毎に 16KB。高速で低レイテンシ。 shared
指示子を用いて宣言された変数が SharedMemory として扱われる。CPU から直接値を読み書きすることはできない。Grid をまたいでデータを保持することはできない。
メモリがバンクに分かれており、同時に複数のスレッドが同一のバンクにアクセスすると性能が劣化する（バンクコンフリクト）ため、最適化の際には注意が必要。(次回の連載にて説明する。)

GlobalMemory GPU 全 体 で 共 有 さ れ る メ モ リ 。全 ブ ロ ッ ク・全 ス レ ッ ド か ら 共 有 メ モ リ と して利用可能。いわゆるビデオメモリ (VRAM) 容量分利用可能だが、TextureMemory と共用。連続アクセスに対しては高速だが高レイテンシ、ランダムアクセスに対しては低速で高レイテンシ。 device 指示子を用いて宣言された変数の他、 global 関数の引数がGlobalMemory として扱われる。CPU から API を介して値を読み書きすることができる。
Grid をまたいでデータを保持することができる。
ConstantMemory GPU 全 体 で 共 有 さ れ る メ モ リ 。GPU 全 体 で 64KB。全 ブ ロ ッ ク・全 ス レ ッ
ド か ら 共 有 メ モ リ と し て 利 用 可 能 だ が 、GPU か ら は 読 む こ と し か で き な い 。MP 毎 にConstantMemory に対するキャッシュが 8KB 搭載されている。 constant 指示子を用いて宣言された変数が ConstantMemory として扱われる。CPU から API を介して値を設定することができる。Grid をまたいでデータを保持することができる。

TextureMemory GPU 全体で共有されるメモリ。いわゆるビデオメモリ (VRAM) 容量分利用可

10能だが、GlobaleMemory と共用。全ブロック・全スレッドから共有メモリとして利用可能
だが、GPU からは読むことしかできない。MP 毎に TextureMemory に対するキャッシュが 6KB か ら 8KB 搭 載 さ れ て い る 。専 用 の 構 造 体 を 用 い て 記 述 す る 必 要 が あ る 。CPU から API を介して値を設定することができる。Grid をまたいでデータを保持することができる。

LocalMemory 実 態 と し て は GlobalMemory で あ り 、ス レ ッ ド 毎 の レ ジ ス タ 数 が 多 す ぎ る 場 合に自動的に割り当てられる。

図6

メモリモデル

各メモリを適切に利用することは GPU を用いて高い性能を得る上で非常に重要である。次章では各メモリを用いたプログラミングの簡単な例を、次回以降の連載では高性能を得るための最適化プログラミングの基本的な考え方と例、および実例を用いた最適化の例を紹介する。

5

CUDA プログラミング入門編 2
これまでに説明してきたように、CUDA は既存の CPU(C 言語プログラミング) よりも複雑な独自の実行モデルやメモリモデルを備えている。そこで本章では、特に使用する機会が多いと思われる GlobalMemory と SharedMemory に ついて 、いくつか 小 さなテスト プログラム を 作成
して基本的な使い方を確認することにする。なお、今回は基本的な使い方の確認を中心とし、性能に関する内容は次回の内容とする。

5.1 GlobalMemory を用いたプログラミング
まずは GlobalMemory を用いたプログラミングについて見てみることにする。

3.1 節 で 示 し た 図 1 の サ ン プ ル プ ロ グ ラ ム で は 、 device の 記 述 こ そ 無 か っ た も の の 既 に
GlobalMemory を利用していた。それは、 global 関数の引数である。4.3 節で述べたように、
global 関数の引数は GlobalMemory として扱われる。これは配列でも配列以外の変数でも同様である。

GlobalMemory の特徴としては、GPU 上の全 SP で共有するメモリであることと、グリッドをまたいで値が保持されることが挙げられる。これらの特徴を確認できるプログラムとして図 7 を作成した。このプログラムは、1 スレッドを持つブロックを 2 つ作成し、1 つ目のグリッド ( global関数) では各スレッドが GlobalMemory への書き込みを行い、2 つ目のグリッド ( global 関数) で11

はもう 1 つのスレッド (異なるブロックのスレッド) が書き込んだ内容を読み出すというプログラムである。こ のプ ログラムから GlobalMemory が ブロック (MP *2 ) を超 えて共有されている ことが確認できるとともに、2 つのグリッド ( global 関数) の間で明示的なメモリコピーを行っていないことから GlobalMemory 値がグリッドをまたいで保持されていることも確認できる。
それでは、GlobalMemory の同一のメモリに対して複数のスレッドが同時に書き込みを行うと ど う な る だ ろ う か 。こ れ ま で に 並 列 処 理 プ ロ グ ラ ム を 書 い た 経 験 が あ れ ば 想 像 で き る よ うに、GPU においても同一のメモリに対する読み書きには気をつけなくてはならない。図 8 は、多数のスレッドから同一の GlobalMemory 上の変数に対して加算を行うという、並列処理においてありがちな「正しく動作しないプログラム」である。本プログラムの実行結果は実行タイ
ミングに依存し、一意に定まらない。これを「正しく動作させる」、すなわちスレッド数分だけ確実に加算させる方法としては、「atomic 関数」を利用するのが容易である。図 8 における

data[0] += 1;を

atomicAdd(&data[0], 1);に書き換えれば意図した結果 (この場合では 16384 という数値) が得られることになる。ただし、多数のスレッドで atomic 関数を使用すると性能低下を招くため多用するべきではない。これに関 連 す る サ ン プ ル と し て は 、各 ス レ ッ ド の 持 つ デ ー タ を 集 め る リ ダ ク シ ョ ン 処 理 の サ ン プ ル(reduction) が参考になるだろう。
以上が GlobalMemory の基本的な使用方法と注意点である。この他、性能について特に重要なこととしてメモリアクセスをできる限り連続アクセス (コアレスなメモリアクセス) になるよ
うにすることなどが挙げられるが、これらについては次回解説する。

5.2 SharedMemory を用いたプログラミング
SharedMemory は低レイテンシで高速な共有メモリではあるものの、局所的な共有メモリであ り 容 量 が 小 さ い た め 、使 用 が や や 難 し い メ モ リ で あ る と 言 え る 。今 回 は 入 門 編 と い う こ とで、SharedMemory を用 いる ことで GlobalMemory へのア ク セス回 数を 減らし て性能を 向 上させる、という基本的な使い方を紹介する。
テストプログラムとして、「各スレッドはブロックごとに GlobalMemory の決められた範囲のデータを収集し、スレッド ID を掛け合わせる」という処理を考えることにする。この処理自体に意味はないが、「スレッド毎に決められた範囲のデータを収集する」「隣接スレッドは互いに近い範囲のデータを処理する」という点については既存の CPU 向けプログラム (並列処理プロ
グラムを含む) において様々な場面で用いられている処理である。(格子点データの時間発展などに応用できるはずである。)
テストプログラム (SharedMemory を用いないカーネル関数と SharedMemory を用いたカーネ
ル関数) を図 9 に示す。fData 配列に注目すると、kernel1 では GlobalMemory 上にある必要なデー*2 ブロックと MP の割り当てを制御することはできないが、今回のようにブロック数が MP 数より少ない場合には各ブ

ロックは別々の MP に割り当てられると考えて良い

12

プログラム (test1.cu)#include <stdlib.h>#include <stdio.h>__device__ float globalarray[2];__global__ void kernel1(){if(blockIdx.x==0){globalarray[0] = 111.11f;}else{
globalarray[1] = 222.22f;}}__global__ void kernel2(float *array){if(blockIdx.x==0){
array[0] = globalarray[1];}else{
array[1] = globalarray[0];}}int main(int argc, char** argv){
int i;
printf("GPU:\n");srand(0);cudaSetDevice(0);float h_Out[2];float *d_Out;
cudaMalloc((void**)&d_Out, sizeof(float)*2);
kernel1<<< 2, 1 >>>();
kernel2<<< 2, 1 >>>(d_Out);cudaMemcpy(h_Out, d_Out, sizeof(float)*2, cudaMemcpyDeviceToHost);
printf("Out: "); for(i=0; i<2; i++)printf(" %.2f", h_Out[i]); printf("\n");cudaFree(d_Out);return 0;}

実行結果>nvcc -O3 -o test1 test1.cu -I/home/ohshima/NVIDIA_GPU_Computing_SDK/C/common/inc>./test1
GPU:
Out: 222.22 111.11>

図 7 GlobalMemory を用いるテストプログラム 1(GlobalMemory の基本動作)

13

プログラム (test2.cu)#include <stdlib.h>#include <stdio.h>__global__ void kernel1(int *data){
data[0] += 1;}
int main(int argc, char** argv){
printf("GPU:\n");srand(0);cudaSetDevice(0);int h_Out = 0;
int *d_Out;
cudaMalloc((void**)&d_Out, sizeof(int));
cudaMemcpy(d_Out, &h_Out, sizeof(int), cudaMemcpyHostToDevice);
kernel1<<< 128, 128 >>>(d_Out); // 128*128=16384parallel
cudaMemcpy(&h_Out, d_Out, sizeof(int), cudaMemcpyDeviceToHost);
printf("Out: "); printf("%d\n", h_Out);cudaFree(d_Out);return 0;}

実行結果>nvcc -O3 -o test2 test2.cu -I/home/ohshima/NVIDIA_GPU_Computing_SDK/C/common/inc>./test2
GPU:
Out: 2>

図 8 GlobalMemory を用いるテストプログラム 2(正しく加算されないプログラム)

タを各スレッドがそれぞれ取得している。そのため、データの取得と足し合わせに合計でスレッド数＊データサイズ分の GlobalMemory アクセスが、各スレッド単位でもデータサイズ分の逐
次的な GlobalMemory アクセスが必要である。一方で kernel2 では途中の計算に SharedMemoryを用いることで、合計の GlobalMemory アクセスはデータサイズ分のみ、各スレッド単位では 1
アクセスのみに削減することができている。kernel2 では SharedMemory へのアクセスを必要とする計算 が追加 され て い るが、SharedMemory は GlobalMemory アク セ スよりもレ イ テンシが小さいため追加された分の時間よりも削減される時間の方が大きく、全体として性能向上を期待することができる。
残念ながら今回のプログラムは、全体のデータ量や並列度が小さすぎる・SharedMemory 上
でのデータ足し合わせも並列化するべきである・GlobalMemory を高速に利用できる問題である、といった事情があり、実際に SharedMemory を利用するメリットが十分に得られる問題ではない。SharedMemory を活用するイメージとして理解していただきたい。

14

// 問題を単純にするため、ブロックあたりのスレッド数は 256 とする// SharedMemory を用いない場合__global__ void kernel1(float *fOut, float *fData){
int i;
float tmp = 0.0f;
int id = blockIdx.x*blockDim.x + threadIdx.x;// 各スレッドが GlobalMemory のデータを取得し足し合わせる
for(i=0; i<256; i++){tmp += fData[i];}// ID を掛け合わせて返す
tmp *= (float)threadIdx.x;
fOut[id] = tmp;}// SharedMemory を用いる場合__global__ void kernel2(float *fOut, float *fData){__shared__ float sData[256];
int i;
float tmp = 0.0f;
int id = blockIdx.x*blockDim.x + threadIdx.x;// 各スレッドが GlobalMemory の部分データを取得し SharedMemory に格納する
sData[threadIdx.x] = fData[threadIdx.x];// スレッド間の同期__syncthreads();// 特定のスレッド上で SharedMemory 上のデータを足し合わせる
if(threadIdx.x==0){
for(i=1; i<256; i++){sData[0] += sData[i];}}// スレッド間の同期__syncthreads();// 各スレッドが ID を掛け合わせて返す
tmp = sData[0] * (float)threadIdx.x;
fOut[id] = tmp;}

図 9 SharedMemory を用いるテストプログラム

以上、第二回の今回は CUDA 対応 GPU アーキテクチャの概要と、CUDA プログラムの概要および実行方法について紹介した。次号では、GPU 上に搭載された階層的な演算器とメモリをより効果的に活用するための最適化プログラミングについて紹介する予定である。

(次号に続く)15