カーネル関数を分離して C++11 コードをビルド.

nvcc でないとビルドできないのはカーネル関数なので, ここと通常の C++ コードを分離してそれぞれビルドしたオブジェクトをリンクすれば, 普通に C++11 を使ったコードが(カーネル関数以外で)書けるようになる.

後, CUDA を使うために必要なライブラリとか, リンクするのが面倒なので nvcc にやってもらう.

以下で検証.

CUDA 4.2
gcc 4.6.3

// kernel.h
#ifndef _KERNEL_H_
#define _KERNEL_H_

void call_kernel(int block, int thread, float * x);

#endif

// kernel.cu
__global__
void kernel_func(float * x) {
  // do something.
}

void call_kernel(int block, int thread, float * x) {
  kernel_func<<<block, thread>>>(x);
}

#include "kernel.h"

// カーネル関数を呼び出すコード.
call_kernel(block, thread, x);

かなりはしょったけど, こんな感じ, cudaMemcpy などのランタイム API は cuda_runtime_api.h に宣言されているので, /include をインクルードパスに追加する.

Makefile で適当に書くとこんな感じになるはず.

CC = g++
NVCC = nvcc
CPPFLAG = -std=gnu++0x -O3 -Wall
NVCCFLAG = -O3 -Xcompiler -Wall
INCLUDE = -I/usr/local/cuda/include
OBJS = kernel.o main.o
NAME = a.out
.SUFFIXES : .cu

$(NAME) : $(OBJS)
  $(NVCC) $(NVCCFLAG) -o $(NAME) $(OBJS)

.cpp.o :
  $(CC) $(INCLUDE) $(CPPFLAG) -c $<

.cu.o :
  $(NVCC) $(INCLUDE) $(NVCCFLAG) -c $<

(ハイライトが微妙におかしい...)

clean とかは省略.

例えば, カーネル関数が float, double など複数の型で欲しいときは, どうしようもないのでカーネル関数自体は template にして call_kernel といった関数を必要な型の分定義するしかない.

そういう面倒なところがあっても, C++ コードを C++11 でかけるのは利点がでかい. 例えばデバイスメモリを RAII で保存する為に unique_ptr とか, move semantics とか, ラムダ式とか. auto が使えるだけでも個人的には嬉しい. 若干メタいコードを良く書くので.

結構備忘録.