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でキャストしてやれば問題無いです。
CUDAのメモリの種類が多くて面倒だという話
はじめに
この記事はKMCアドベントカレンダー2013の5日目の記事です。
tyage氏の書いた昨日の記事はこちらです。
ちなみにKMCにはガチ情報系な人が多く居ますが、意外とCUDA書いてる人って全然居ないんですよ。
私も普段ブログを書いてはいるのですが、あまりにも技術的な話と程遠い日常的な愚痴ばかり書いているブログなのでこのイベントに参加させるのは場違いなのでは、というのがこのブログを作るきっかけだったりします。
ついでにこのブログが大学の研究室の人に知れても大丈夫なようにわざわざアカウント分けたりとかも、ね。
さてさて、今回はCUDAのメモリ管理がめんどいというお話を今回はしていきましょう。
デバイスから見られるメモリの種類
デバイス(GPU)から扱えるメモリにはいくつかの種類があります。
それぞれどういった特徴を持つか見てみましょう。
グローバルメモリ
ホスト(CPU)側から書き込みができるメモリです。
容量も非常に大きい(数GBとか)です。
その分読み書きには手間がかかります。
シェアードメモリ
ブロックごとに積まれているメモリです。
容量が小さい、ブロックごとにしか共有できない、カーネルから出ると使えない(同じIDでも同じブロックになるとは限らない)と厄介な性質を持ちますが読み書きは速いです。
レジスタ
スレッドごとに用いるメモリです。
非常に読み書きは速いですが非常に小さく、スレッド間でレジスタのデータをやりとりする場合には昨日触れたWarp Shuffleやシェアードメモリの使用が必要になります。
キャッシュ
通常のキャッシュの他にコンスタントキャッシュ、テクスチャキャッシュが有ります。
いずれも明示的に用いることは出来ませんが後半の特殊なものについてはある程度意図的に用いることになります。
ここまでが実際に存在するメモリです。
さらにこれに加えてコード上から仮想的に指定するメモリが存在します。
これらは実際はグローバルメモリ上にありますがコンスタントキャッシュ、テクスチャキャッシュを活用して高速で読み込みを行うことが出来ます。
テクスチャメモリ
テクスチャキャッシュを用いるためのメモリです。
1～3次元の格子状にデータを配置してそのインデックスを用いて読み込みを行います。
1度読み込みを行うとその周辺の値をまとめてテクスチャキャッシュに読み込み、そこから読み込みができるようになります。
そのため、近接する場所の値を連続して読み込むのは高速になります。
ただし、キャッシュに読み込んだデータを返せないため書き込みはできません。
コンスタントメモリ
コンスタントキャッシュを用いるためのメモリです。
ホストからは書き込めますがデバイスからは読み込みしか出来ません。
コンスタントメモリの中身は全てコンスタントキャッシュに移すため、使える容量は小さいです。
サーフェスメモリ
テクスチャキャッシュを使いつつもデバイス側から書き込みができるメモリ……らしいのですが使っている例が全く見つからずよくわかりません。
いずれ自分で調べてここに書きたいです。
と、まぁ沢山の種類があるわけです。
CUDAで効率化を図る場合はこれらを適切に使い分ける必要があるわけです。
どう使い分ける？
グローバルメモリの使用を極力避けるのが肝になります。
何度も読み込むデータはシェアードメモリに移しましょう。
シェアードメモリは明示的に扱えるキャッシュのようなものとしてガンガン使いましょう。
定数はコンスタントメモリに入れましょう。これはシンプルにできます。
テクスチャメモリは使いどころが限られますが、1～3次元上に並んだデータに連続アクセスする場合には大きな効果を発揮します。
無理して使うほどではないですが、もともとそういう形のデータを扱うなら使わない手はないでしょう。
最後に
KMCアドベントカレンダー2013、明日はohai氏による「Rubyでバイナリデータを取り扱う方法について」らしいですよ。
他にもいろいろあるので興味があるテーマについて話しそうな日があったら見てみて下さいね。
Warp Shuffle
Warp Shuffleとは
Warp Shuffleは同Warp内の別スレッドが持つレジスタの値を受け渡すための命令です。
これを用いずにレジスタの値をスレッド間で共有するためにはシェアードメモリなどのメモリを用いる必要があります。
同Warp内(32のスレッド)でしかやりとりが出来ないので汎用性は劣りますが速度は向上します。
Warp Shuffleの使い方
WarpShuffleを行うメソッドは4種類あります。
//指定したレーン(Warp内のスレッド番号)のvarの値を受け取ります T __shfl(T var, int srcLane, int width=warpSize); //指定した数離れたスレッドのvarの値を受け取ります //upは指定した数上の番号、downは下の番号のスレッドから受け取ります //Warpの外に出てしまった場合は自身のスレッドでの値が返ってきます T __shfl_up(T var, unsigned int delta, int width=warpSize); T __shfl_down(T var, unsigned int delta, int width=warpSize); //自分のレーン番号と指定したlaneMaskをXORした結果の番号のスレッドから受け取ります //laneMaskごとにレジスタの中身を交換することが出来ます T __shfl_xor(T var, int laneMask, int width=warpSize);
渡す値はfloatもしくはintに限られます。
いずれも別のスレッドでvarに渡した値を返り値として受け取る形になります。
交換する形でなく、1つのスレッドから他のスレッドに値をブロードキャストする用途でも使えます。
高速化される？
メモリを介して値をやりとりした場合と比べてWarp Shuffleが速いのはなぜでしょうか？
メモリを介しての値をやり取りをする場合、メモリに値を書き込む→同期をとる→そのメモリから値を読み出すというプロセスが必要になります。
それに対してWarp Shuffleは1回の書き込みのみで済むのが強みです。
同Warp内であれば同期の必要もないですしね。
用途は？
調べて書いては見たものの、使い所がなかなか難しいです。
最初にも述べたように同Warp内でのやりとりに限られるのでちょっと扱いにくいかもしれません……。
とりあえず今日調べたことをまとめておきたくて書いたけどちょっとサンプルとか少なくて読みにくかったかもしれません。
ある理由から明日までに1つ記事を書いておきたかったというのもありますけどね。
その理由は明日に。
はじめに
このブログ is 何？
このブログは情報学の知識も大してないのにCUDAを書くことになった人のブログです。
厳密性よりはわかりやすさ重視のゆるっとふわっとなブログです。
ガチでプログラミングしてる人よりは、何となくCUDAを書いてみたい人/上の方針で書かされている人なんかが対象です。
ただ、基礎の基礎のお話しは(多分)しないのでコードの書き方そのものとかは他所で勉強してきてくださいね。
記事の順番も基礎からとかではなく、何となく勉強した順に書いていきますので必要に応じて探して下さいね。
CUDA is 何？
念のためCUDAについてちょろっとだけ触れておきましょう。
CUDAは(nVIDIA社製の)GPU上で動作するプログラムを記述するための言語です。
CPUで動作する部分はC/C++でそのまま、GPUで動くカーネル部分も同様にCの記法で書くことが出来ます。
GPU is 何？
GPUはGraphics processing unitの略です。
元々は画面の表示を司る部分でしたが、時が進むに連れ3Dレンダリングを担当するようになり、さらにその機能を活かして別の計算にも使われるようになっています。
並列計算に特化した構造を持っており、最大パフォーマンスはCPUを上回ります。
ただ、並列化を考えなければならない他にもいろいろと考えるべきことが多く、CPUで動作するプログラムを作るよりは手間がかかります。
そんな手間にまつわるいろいろなことをこのブログに書いていこうと思うわけです。
とりあえずちょっとずつ書いては行きますが更新は不定期になると思います。
少しでも同じような境遇の人の助けになればと思います。
