さて、CUDAの開発環境の構築は昔から少々面倒で、多少のLinuxの知識と折れない心と英語力が必要です。
自分もこのたびわりと久しぶりに自宅PCにCUDA環境を構築したら、やり方も変わっていて躓いたので備忘録としてメモ。
さかな前線 » CMakeでCUDAの-ccbinの正しい指定方法の姉妹記事です．
さて，Linux上でOpenCVでCUDAサポートをONにしてビルドするとき，GCCのバージョンを指定したいことがあります．
基本的にはデフォルトGCCなんですが，4.8系などは現在のCUDAのnvccではサポートされていませんので，CUDAビルド時に限りGCC4.4や4.6を使うという設定です．
具体的には，CMakeのときに以下のように書きます
$ pwd /home/hogehogehogehogeho/opencv-2.4.6.1/build $ cmake -DWITH_CUDA=ON -DCUDA_HOST_COMPILER=/usr/bin/g++-4.4 .. $ make -j40 -s; sudo make install
WITH_CUDAとともにCUDA_HOST_COMPILERを正しく設定すればいいのですが，CUDA_HOST_COMPILERに指定するgccはPATHが通っていても絶対パスで指定する必要があるようです．
CUDA_HOST_COMPILER=g++-4.4と書くと，ないよ＞＜ って言われます．
以上です．
参考：OpenCV – Bug #2844: build with CUDA 5.0 on macosx – OpenCV DevZone
さて，Linux上のCMakeでCUDAを使っているんですが，CUDAではたいてい最新のGCCはサポートされません．
ということで，nvccの-ccbinオプションによってCUDA関連のコンパイルのみを古いGCCにすることができます．
普段のC++コードは最新のGCCで，CUDA関連のみはサポートされたGCCでご利用いただけます．
CMakeにおいてこれを行うときは少しクセがあり，だいぶ格闘したのでメモ．
例えば次のソースを用意します．
// prog.cu
int main()
{
}
GCCの4.8.1を使ってnvccでコンパイルしてみると．
$ g++ --version
g++ (GCC) 4.8.1
Copyright (C) 2013 Free Software Foundation, Inc.
This is free software; see the source for copying conditions. There is NO
warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
$ nvcc prog.cu
In file included from /usr/local/cuda/bin/../include/cuda_runtime.h:59:0,
from <command-line>:0:
/usr/local/cuda/bin/../include/host_config.h:82:2: error: #error -- unsupported GNU version! gcc 4.7 and up are not supported!
#error -- unsupported GNU version! gcc 4.7 and up are not supported!
^
$ nvcc prog.cu -ccbin g++-4.4
$ ./a.out
(うごいた)
こうなります．
ちなみに今回はCUDA 4.2の環境．
本題．
CMakeだとnvccは自動的に実行されるので，nvccに与えるオプションはちゃんと与える必要があります．
ということで，CMakeLists.txtの最小の例．
上のprog.cuと同じディレクトリに置きます．
#A tiny example of CMakeLists.txt
cmake_minimum_required(VERSION 2.8) # おまじない
project(osakanacuda) # プロジェクト名
set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} -ccbin g++-4.4) # GCCへのパス．PATHが通っていればこのように名前のみの指定でもOK
#list(APPEND CUDA_NVCC_FLAGS -ccbin g++-4.4) # これでもOK
#以下はどちらもダメ．違いがわかりますか？
#set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} -ccbin=g++-4.4)
#list(APPEND CUDA_NVCC_FLAGS=-ccbin g++-4.4)
find_package(CUDA) # CUDA環境を見つけてくれる．やらないとダメ
cuda_add_executable(prog prog.cu) # progっていう実行ファイルをビルドしてくれる設定になる
こんな配置だとすると，
. ├── CMakeLists.txt ├── prog.cu └── build/
次のようにしてビルドします．
$ ls CMakeLists.txt prog.cu build/ $ cd build $ cmake .. $ make
ダメって書いてある方の指定をすると，
nvcc fatal : redefinition of argument 'compiler-bindir' CMake Error at prog_generated_prog.cu.o.cmake:206 (message): Error generating /hogehogehogehogheohgoehgoehoge/build/CMakeFiles/prog.dir//./prog_generated_prog.cu.o
っていう内容を含んだエラーが出ます．
正しい指定と誤った指定の違いは見ての通り，”-ccbin=hogehoge”か”-ccbin hogehoge”の違いなんですが，CMakeのFindCUDAにあるrun_nvcc.cmakeを読んでみると，
“-ccbin”という文節があった場合にその指定を優先するようになっており，
ない場合はCMakeの中でのデフォルトの指定をすることになっています．
なので，”-ccbin=hogehoge”をCUDA_NVCC_FLAGSに指定すると，nvccが呼び出されるときは
nvcc ... -ccbin=hogehoge -ccbin fugafuga ...
となり，「多重指定はダメだっちゃ」と怒られます．
なんちゅー仕様．誰か修正してくださ・・・
なお今回の環境は，
です．
行列といえばGPGPUが活躍する最も重要な分野の一つと言えますね．
んで，そんな行列でよく使われるタスクのひとつに「転置」があります．
画像屋さんなら，線形なフィルタをかけるときには縦方向にフィルタして横方向にフィルタすればいいのですが，実装的には横方向にフィルタして画像全体を転置してまた横方向にフィルタして転置して戻す方法のほうが効率的なことが多いです．特にGPGPUなどでは．
ということでよく使う転置というタスクですが，CUDAでやるにはちょっと工夫しないとメモリアクセス効率が下がります．今日のお題はそれで．
% cuda-gdb cuda-gdb: error while loading shared libraries: libpython2.6.so.1.0: cannot open shared object file: No such file or directory
自分でlibpython2.6を入手して手元でビルドしてもいいんですがそんなねぇ。
ごまかしちゃえばいいみたいです。
% sudo ln -s /usr/lib/libpython2.7.so.1.0 /usr/lib/libpython2.6.so.1.0 % sudo ln -s /usr/lib64/libpython2.7.so.1.0 /usr/lib64/libpython2.6.so.1.0
ちょ。
YOUこんなことしていいのかYO！
% cuda-gdb NVIDIA (R) CUDA Debugger 4.1 release Portions Copyright (C) 2007-2012 NVIDIA Corporation GNU gdb (GDB) 7.2 Copyright (C) 2010 Free Software Foundation, Inc. License GPLv3+: GNU GPL version 3 or later <http://gnu.org/licenses/gpl.html> This is free software: you are free to change and redistribute it. There is NO WARRANTY, to the extent permitted by law. Type "show copying" and "show warranty" for details. This GDB was configured as "x86_64-unknown-linux-gnu". For bug reporting instructions, please see: <http://www.gnu.org/software/gdb/bugs/>. (cuda-gdb)
∧∧
ヽ(･ω･)/ ｽﾞｺｰ
＼(.＼ ノ
､ﾊ,,､ ￣
￣
ちょっと使った範囲だと問題なく動いてはいます。
nvccは限定的ながらC++に対応してるということで、例えばクラスメンバがvirtualだったりするとダメとかみたいな制限はあるんですが、カーネル自体がクラスメンバになれないとは。
class ACUDAClass
{
__global__ void kernel();
public:
void func()
{
kernel<<<1, 1>>>();
}
};
void ACUDAClass::kernel()
{
}
int main()
{
ACUDAClass ob;
ob.func();
}
こんなエラーが出ます。
% nvcc prog.cu prog.cu(3): error: illegal combination of memory qualifiers 1 error detected in the compilation of "/tmp/tmpxft_00003d82_00000000-4_prog.cpp1.ii".
もちろん、こういう風にクラスメンバからカーネルを呼ぶのは何も問題ありません。
__global__ void kernel()
{
}
class ACUDAClass
{
public:
void func()
{
kernel<<<1, 1>>>();
}
};
int main()
{
ACUDAClass ob;
ob.func();
}
理由はどうも単純なもので、thisをカーネルに渡せないからクラスメンバとしては置けないよ！！というアレみたい。
ならばと思ってカーネルをstaticにしても同じエラーがでます。はて。
例えばグローバル変数(ホスト側にある)を参照するとundefined symbolってちゃんとエラー吐けてるから同じ程度の問題のような気もする。
ところが、__device__だとおっけーなんですね。
ふむ。わからん。
誰かホントのところ教えてくだしあ。
あと最初に言ったようにこれCUDA 4.0での話なので4.1ではどうなってるか未確認です。ごめんね∩(・ω・)∩
さて、CUDAコンパイラnvccの吐くオブジェクトファイルは普通にgccの吐くのと同じELF形式ですので、デバイスコードを記述したソース(*.cu)をnvccで-cオプション付きでコンパイルして出力したオブジェクトファイルとgccで出したオブジェクトファイルをリンクすることができます。
mainから直接nvccを通さないために、カーネルを呼び出すフロントエンド関数を用意
void a_cuda_kernerl_frontend();
デバイスコードがこちら
#include "a_cuda_wrapper.h"
__global__ void cuda_kernel()
{
}
void a_cuda_kernerl_frontend()
{
cuda_kernel<<<1, 1>>>();
}
そしてnvccを通さないmainがこちら
#include <iostream>
using namespace std;
#include "a_cuda_wrapper.h"
int main()
{
a_cuda_kernerl_frontend();
}
これで
% g++ -c main.cpp % nvcc -c kernel.cu % nvcc main.o kernel.o % ./a.out
リンクするときにg++じゃなくてnvccを使いましょう。当然ですね。
これで、C++ベースのでかいシステムの一部としてCUDAを用いることもできますね！！
これができることくらい誰でも知ってると思うんだけど、実際にやってるサンプルて見ないなぁと思って(mainまで*.cuに書いてある)
（そうですサンプルにするほどのこともないからです
いずれもPDF注意です．
東京大学情報基盤センター・大島聡史先生による『これからの並列計算のためのGPGPU連載講座』より，
特に2と3は非常に良いです．必読です．
あと，プログラムはちゃんと写経しましょう．ぼけーっと写すだけでもぜんぜん理解が違います．
CUDA4.0のような最新の技法・アーキテクチャ等はともかく，また量的にまだまだいろいろ書かれてないことはあるでしょうけど，基本を押さえるのによいです．
CUDA4.0はUbuntuに関して10.10を正式サポート対象としてるのもあって，ハッキリ言って，導入に躓く要素はないです．
ちなみに僕のメインマシンはモバイルノート(ThinkPad X201s)なんですが，研究室のマシンがi7-920に9500 GTなので，基本的にCUDAは(エミュレーションじゃなく)こっちでやってます．