準備

提供: メモ帳@fmaj7b5.info
2012年11月11日 (日) 17:53時点における白飯 (トーク | 投稿記録)による版

(差分) ←前の版 | 最新版 (差分) | 次の版→ (差分)
移動: 案内、 検索

2周くらい遅れのCUDAをやってみる。 ここではWindowsを使うけど、CUDAはOSに依存しないので他でも大丈夫・・・だと思う。

目次

プログラム環境を整える

以下の二つをインストール：

  • Visual C++ 2010 Express
  • CUDA Toolkit

ググれば出てくるはずだけど、CUDAのダウンロード先で迷ったことがあるのでヒントを少々。CUDA周りのツール類は「NVIDIA > CUDA Zone > Download」と辿っていけば見つかる。とりあえずCUDA ZoneでDownloadとかToolkitとか書いてあるところをポチポチしよう。

細かいところは先人のwebページを参考にしてね。

コンパイル方法

最低限必要なのは以下の三つ：

  • *.cuをCUDAコンパイラ(nvcc)でコンパイルするように指定する
  • CUDAコンパイラに対象にするGPUのアーキテクチャ(-arch)とcompute capability(-code)を指定する
  • cudart.lib (linuxだとlibcudart)をくっつけるよう指定する

Visual C++

GUIは説明が面倒だ・・。

  1. プロジェクトを追加
    • コンソールアプリケーションとか、適当に
  2. ビルドのカスタマイズ > CUDA x.x にチェック
  3. *.cuソースを追加
  4. プロパティ > CUDA C/C++ > Device > Code Generation で対象GPUを指定
  5. リンカー > 入力 > 追加の依存ファイル にcudart.libを追加

まあ、詳しくはググるべし。

CMake

Linuxな人はこっちが簡単かも。WindowsでVC++な人でも使えるので、便利な人には便利。 

CMakeLists.txt:

cmake_minimum_required (VERSION 2.8)
project (HelloProject)

find_package (CUDA REQUIRED)

cuda_add_executable (hello hello.cu)

cmake_minimum_requiredのバージョンは手持ちの環境に合わせただけなので、もっと低くても大丈夫かも知れない。

GNU Make

いや、漢(おとこ)ならmakeだ。という人の場合。CMakeもLinuxなら最終的にmakeするじゃん、というツッコミは無しで。 

cuda.mk:

# Set a default location of CUDA
CUDA_HOME ?= /usr/local/cuda

# Set an architecture
ifeq ($(ARCH),)
  ARCH = $(shell uname -i | tr '[:upper:]' '[:lower:]')
endif

# Binary path to CUDA
CUDA_BINDIR = $(CUDA_HOME)/bin

# Set library dir and nvcc command
ifeq ($(ARCH),x86_64)
  CUDA_LIBDIR = $(CUDA_HOME)/lib64
  NVCC = $(CUDA_BINDIR)/nvcc
else
  CUDA_LIBDIR = $(CUDA_HOME)/lib
  NVCC = $(CUDA_BINDIR)/nvcc -m32
endif

# Set default parameters for CUDA libraries
CUDA_LDFLAGS ?= -L$(CUDA_LIBDIR)
CUDA_LDLIBS ?= -lcudart


# Build rules
%.o: %.cu
	$(NVCC) $(NVCC_FLAGS) -c -o $@ $^

%.ptx: %.cu
	$(NVCC) $(NVCC_FLAGS) --ptx -o $@ $^
%.ptx: %.gpu
	$(NVCC) $(NVCC_FLAGS) --ptx -o $@ $^

%.gpu: %.cu
	$(NVCC) $(NVCC_FLAGS) --gpu -o $@ $^


Makefile:

include ./cuda.mk

LDFLAGS = $(CUDA_LDFLAGS)
LDLIBS = $(CUDA_LDLIBS)

NVCC_FLAGS = --generate-code arch=compute_30,code=\'compute_30,sm_30\'

TARGETS = hello

all: $(TARGETS)

clean:
	$(RM) $(TARGETS) *.o *~

hello: hello.o

適当なので、あくまでも参考までに。大体、Makefileを手書きする変態さんならもっといいのを既に書いてるでしょう・・。 あ、よく見たらGettingなんちゃらガイドにMakefileのサンプルが。

基本いろいろ

Host／device

CPUとGPUが別々になっているPCだと、CPU+メインメモリ・GPU+VRAM(?)という構成でそれぞれ管理されている。前者をhost、後者をdeviceという。 イメージとしては、hostでプログラムが動いていて面倒な計算をするときにdeviceへ丸投げする感じ。なので、device側はGPUに限らず面倒な計算を引き受けてくれるものなら何でもいい。

基本的にはお互いのメモリに直接アクセスできないので、

  1. Hostのメモリ内容をdevice側にコピー
  2. Deviceで計算
  3. Deviceメモリにある計算結果をhost側にコピー

という処理になる。このコピーの処理がオーバーヘッドになるので、いくら並列計算向きな処理でも軽い処理ならhost側で普通に計算した方が速いってこともある。反対にdevice側で計算していて途中の処理が並列計算に不向きでも、host側に渡して計算してdevice側に書き戻すよりそのまま計算した方が速いことも。

最近のGPUだと計算と次に使うメモリのコピーを同時にできるらしいからスループットを上げられるし、CPUとGPUがメモリを共有しているヤツとかコピーが不要だったりするけど、まあ、host <-> device間の転送はなるべく避ける方がいい。

*.cuの運命

CUDAはhost用とdevice用のコードがまぜこぜになっているので、中で何が起きているのか分かりづらい。なので、何となくこんな感じかな～っと。

__host__または指定無し
*.cu -> CUDAっぽいキーワード削除・__global__関数の呼び出しコード追加 (*.c/*.cpp) -> 普通のC/C++コンパイラ(gccとかcl.exeとか)に渡す
__global__または__device__
*.cu -> GPU側コードの抽出 (*.gpu) -> 仮想アーキテクチャ用コード生成 (*.ptx) (-> 実アーキテクチャ用コード生成 (*.cubin) )

たぶん、__host__とか__device__のキーワードを見て生成するソースに含めるかどうか決めているんだと。なので、__host__ __device__という風に両方ついているものはhost側、device側両方のソースに入れられる。ちなみにどちら側でコンパイルされているのかは__CUDA_ARCH__のあるなし、仮想アーキテクチャの種類はその値で判断できるようだ。

仮想／実アーキテクチャ

nvccのコンパイルオプションで迷ったのが、--gpu-architecture (-arch)と--gpu-code (-code)。どっちも引数は同じようなのがとれるし、ヘルプも長ったらしいし。で、たぶんこういう事かなというのが分かった気がしたので書いておく。

まず、2種類の引数について。

compute_*
仮想アーキテクチャ。どういう命令セットを使うかを表す。数字が大きい方がより多くの便利な命令が使えるけど、古い環境では対応できなくなる。
sm_*
実アーキテクチャ。バイナリコードを生成する対象を表す。当然、仮想アーキテクチャのバージョン以上でないとダメ。

それからオプションの意味について。

--gpu-architecture (-arch)
*.cuが対象にしている仮想アーキテクチャを指定する。なので、基本的にはcompute_*の中から選択する。
--gpu-code (-code)
仮想アーキテクチャのコード(*.ptx)から生成して出力に加えるアーキテクチャを指定する。PTXを含めたいときはcompute_*、特定のGPU向けバイナリを含めたいときはsm_*を選択する。

出力にPTXを含めておくと、自分向けのバイナリが見つからなかった時にその場でコンパイル (JIT)してくれるらしい。 オプションを2個も指定するのが面倒な人向けに--gpu-architecture (-arch)にsm_*を指定すると、対応する仮想アーキテクチャと出力対象を指定したことにしてくれる便利機能がある。 

# 同じ
--gpu-architecture=compute_13 --gpu-code=sm_13
-arch=compute_13 -code=sm_13

# 未来のGPUでも動作するよう、出力にPTXコードを含める
-arch=compute_13 -code=compute_13,sm_13
-arch=sm_13

# バイナリをたくさん入れる
-arch=compute_13 -code=compute_13,sm_13,sm_21,sm_30

# 複数のアーキテクチャに対応する
--generate-code arch=compute_10,code=compute_10 --generate-code arch=compute_13,code=\'compute_13,sm_13\' ...

最後の例みたいに、複数の仮想アーキテクチャに対応するには--generate-code (-gencode)で指定する。ソース側は前述のとおり__CUDA_ARCH__でどのアーキテクチャ向けか判定する。

「http://www.fmaj7b5.info/wiki/index.php?title=準備&oldid=58」から取得