Thrustを使おう その1
今回から何回かに分けてThrustのお話をします。
非常に有用なものである分他に資料も多いですが、触れずに置くのはコンセプトからしてナシだと思いましたので。
第1回の今回はまずvectorについてごく簡単に、配列の代わりに使う程度までの話をしましょう。
Thrustとは？
Thrustは端的にいうならばC++におけるSTLに相当するようなライブラリです(違いはもちろん多くありますが)。
CUDA 4.0以降自動的にインストールされるようになっているので使うのにインストールなどは必要ありません。
C++においては特にパフォーマンスを求める場合などを除き「配列ではなくvectorを使っておけ」なんて言われるものですが、CUDAにおいても特別な理由がない限り配列ではなくThrustのvectorを活用するのが良いでしょう。それぐらい有用なものです。
Thrustを使わない場合
これまでCUDAのコードを書いたことがある人ならばcudaMallocを使ったことはあるでしょう。
グローバルメモリを確保しようとした場合こんなコードをまず書くよう学んでいるはずです。
int* dArray=NULL; //グローバルメモリに領域を確保 cudaMalloc((void**)&dArray, 1024*sizeof(int)); //sourceArrayから値をコピー cudaMemcpy(dArray, sourceArray, 1024, cudaMemcpyHostToDevice); //いろいろ処理 //メモリを開放 cudaFree(dArray);
なんとも面倒くさいです。
(void**)て何やねんって感じですし、型を途中で変えてsizeofの中身を書き忘れたりしかねません。
C++にちょっと慣れている方ならテンプレートで楽にしたいと思うところでしょう。
そこに入れるデータも予め作っておいたホスト上の配列からコピーする必要があります。これもまた面倒です。
うっかりcudaMemcpyHostToDeviceとcudaMemcpyDeviceToHostを間違えた事があるのは私だけではないと思います。
何より、cudaFreeをいちいち書き加えるのが面倒で、忘れてしまいがちです。
と、いった感じに非常に面倒であり、それ以上に記述ミスで不具合を生み出してしまいかねない状態なわけです。
Thrustを使った場合
さて、まず百聞は一見にしかずということで使った場合のコードをお見せしましょう。
//コード冒頭 #include "thrust\host_vector.h" #include "thrust\device_vector.h"
//ホスト上のベクトル thrust::host_vector<int> sourceVector(1024); //値を入力 //グローバルメモリ上に領域を確保 thrust::device_vector<int> dVector(1024); //sourceVecの中身をdVectorへコピー dVector=sourceVector; //いろいろ処理 //解放の必要は無い
と、言った感じにシンプルになりました。
サイズ指定も要素数だけで出来ますし、ホストからデバイスへのコピーは代入のみで済みます。
さらに自動的にメモリの解放までしてくれます。至れり尽くせりです。
とりあえず確保とコピーだけお見せしましたが、もう少し詳しく使い方を見てみましょう。
thrust::host_vector<int> sourceVector(1024);
for(int i=0;i<1024;i++)
{
//host_vectorはただのvector,配列のように扱える
sourceVector[i]=i;
}
thrust::device_vector<int> dVector(1024);
dVector=sourceVector;
//カーネルにポインタを渡す
kernel<<<16,64>>>(thrust::raw_pointer_cast(dVector.data()));
まず、host_vectorはただの配列、C++のvectorのようにインデックスでアクセスできます。
ちなみに今は普通にforで代入しましたが今回のような場合はもうちょっといいやり方もあります(それは次々回ぐらいで)。
そして、thrust::raw_pointer_castが何をしているかというと、その名の通りポインタへのキャストを行います。
要はThrustを使わない例でのdArrayと同じ形にしたと思ってもらえばいいです。
なぜこのようなことをするのかというと、カーネル上でvectorを使用できないためです。
そのため、カーネル上ではこれまでと同じようにポインタを受け取って配列として扱う必要があります。
ホスト上ではdevice_vectorとして使い、カーネルから呼ぶ時だけraw_pointer_castでキャストしてやれば問題無いです。
