10/10/13 15:43:50
・GPUでやるほどの処理なのか疑問。
→実験目的なら納得。
・可読性の悪さをコメントで補うという厄介なコードになっている。読み易いように以下のように書けば充分。
・同様に、elseの前後で複文と単文を混ぜるのは宜しくない。後でコードを追加する場合の備えも兼ねて複文の方が宜しかろう。
・ついでに、c/c++では「一つ増やす」ことを明示的に書けるから積極的に使った方が読み易い。
# この場合、「数え上げ」のために1増やすのであって、1という数値自体に意味があるわけではない。
--
if (stu[id].point >= 60) {
++ends[0];
} else {
++ends[1];
}
--
・周辺コードがないからなんとも言えないが、ちゃんと転送されているのだろうか。
# thrust::copy()かな。
・同様に、endsは初期化されているのだろうか。
659:デフォルトの名無しさん
10/10/13 15:57:16
CUDAは並列動作時に同じアドレスへの書き込みを保証できないから
atomic命令使ってカウントするか、同数のtrue/falseフラグ配列を用意して
カーネル終了後に別途足し併せないと無理
660:デフォルトの名無しさん
10/10/13 16:06:49
そうだ、肝腎なことを忘れていた。
# 馬鹿でぇ＞自分
## つーか、流石に自分自身ではこんなミスはしないのに……
そもそも、複数のスレッドが同じメモリに書き出すのは無理。
ほぼ同時に走るから、1や2になるのも道理。
逆に言えば、巧く並列に実行しているからこそちゃんとカウントできない。
661:デフォルトの名無しさん
10/10/13 16:47:31
>>658
>>・可読性の悪さを～
申し訳ないです。実際endsの初期化とかstudentへの値の代入などその他の処理が結構あったため切り貼りしたのが原因です。
>>・周辺コードが～，同様に～
この部分の動作は一応確認できていたため省いてしまいました。すいません。
>>660
自分も並列につまり同時に計算しているのに同じ場所に加算を行えるのか疑問でしたがやはりムリだったんですね。
確かに逆に考えれば動いているともいえそうです。
実は今は多少違う方法で解決したため問題はなさそうですがEmuで成功した理由がわからなかった+後学のために質問しました。
あともう1つ質問なのですが
thrust::device_vector<hoge> hoge(N);
という宣言を行った場合ホスト側にメモリ領域を取らずデバイス側にメモリ領域を確保できるのでしょうか?
662:デフォルトの名無しさん
10/10/13 16:52:00
すいません肝心のお礼を書いていませんでした。
即レスありがとうございます。参考になりました。
663:デフォルトの名無しさん
10/10/13 17:00:45
>659が真の功労者。お礼はそっちに。
emuで成功するのは、マルチスレッドでエミュレーションしないでシングルスレッドでエミュレーションしているからだと思われ。
マルチスレッドで並列化しても同じように起きる典型例だからね。
thrustは詳しくないけど、ちょっと見てきたところで判断する限り、ホスト側とデバイス側は別々にメモリを確保する必要がある。
従って、device_vectorはデバイス側だけメモリを確保すると思われ。
664:デフォルトの名無しさん
10/10/13 17:09:35
重ね重ねありがとうございます。この板IDが出ないので誰だかわからないですね。
実際に取り扱うデータは10万件くらいのデータ(蓄積しているのでもっと多い)の分析で
分析をC++で書いていたのと並列化に向いたアルゴリズムも存在していたため
CUDAなら乗せれると思って勉強しはじめました。基本はC,C++ライクなのでがんばってみます。
665:デフォルトの名無しさん
10/10/14 00:37:01
runtime versionてどう調べるの？
cudaの
666:デフォルトの名無しさん
10/10/14 00:42:13
devicequery
実行できないなら自分が入れたもんくらい覚えとけよ
667:デフォルトの名無しさん
10/10/14 18:01:10
CUDAでカーネルを10回実行しようと思った場合
以下のような書き方でいいのでしょうか?
for(int i=0; i<10; i++){
kernel<<<n,m>>>(,,,);
}
668:デフォルトの名無しさん
10/10/14 18:22:25
一応。
タイマーとか挟むつもりなら注意が必要だけど
669:デフォルトの名無しさん
10/10/14 21:32:45
>>667
カーネル呼び出しはコストが掛かるから、できれば一発呼びで済ませた方が。
その例なら、kernel<<<dim3(n, 10, 1), m>>>(...)できないだろうか。
670:デフォルトの名無しさん
10/10/15 02:11:53
667ではない，プログラミング初心者ですがコストっていうのはどういう意味ですか．
671:2の方だね
10/10/15 11:36:33
>>670
cost
[名][U][C]
1 代価；原価；支出, 費用, 経費, コスト(running costs)；((?s))（家庭・企業などの）諸経費；((?s))《法律》訴訟費用
the high ［the low］ cost of electricity
高い［低い］電力コスト
the total cost of the scheme
計画の総経費
the cost of production ［＝production costs］
生産費
sell at ［below］ cost
原価［原価以下］で売る
without cost
無料で
cover the cost
元をとる
cut ［reduce］ costs
経費を切り詰める
My car was repaired at a cost of ￥50,000.
車の修理代は5万円だった. ⇒PRICE[類語]
2 ((しばしば the ?))（時間・労働などの）犠牲, 損失
at a person's cost
人の犠牲において, 人に損害をかけて.
672:デフォルトの名無しさん
10/10/15 13:12:33
配列aには100以下の数値がランダムで100個格納されています。
配列bの中身をaと同じにしたい場合以下のように書けばできました。結果も正しく出力できました。
__global__ void changeA(int *a, int *b){
int id = threadIdx.x;
a[id]=b[id]
}
上記は100スレッドで並列計算を行いましたがスレッド数を10個(10,1,1)にした場合正しく計算されませんでした。
ブロック数を10個(10,1,1)スレッド数を10個(10,1,1)にしてidを
int id = blockIdx.x * blockDim.x + threadIdx.x;
とした場合正しく計算できました。こうなる原因は大体わかっているのですが
スレッド数が少ない場合、またはスレッド数が膨大な場合たとえば(1000,1,1)のようにした場合で正しい結果を得ることができるのでしょうか?
スレッドが膨大な場合size(この場合100)を渡しidがsizeを超えれば処理をしないというものです。
この場合無駄ができそうなのでよい方法はないでしょうか?
__global__ void changeA(int *a, int *b, int size){
int id = blockIdx.x * blockDim.x + threadIdx.x;
if(id =< size){
a[id]=b[id]
}
}
673:デフォルトの名無しさん
10/10/15 13:15:30
すいませんスレッドの場合、総計512が最大でした。ブロック数が(1000,1,1)です。
674:デフォルトの名無しさん
10/10/15 14:23:27
>スレッド数を10個(10,1,1)にした場合正しく計算されませんでした。
意味が分からん。んなわけないだろ
分岐はウォープ内で同じ方向だったら対して無駄にならないんじゃない？
675:デフォルトの名無しさん
10/10/15 14:23:37
a[id] = b[id]じゃ説明が逆だろうけどそれはそれとして。
単にコピーしているだけとして、これでいいと思う。
--
__global__ void changeA(int * a, int * b, int size)
{
for (int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < size; idx += gridDim.x * blockDim.x) {
a[id] = b[id];
}
}
--
勿論gridDim.x * blockDim.xがsizeよりも大きければ>672と同じになるし、sizeが大きくても破綻しない。
ブロック数が過剰な場合に無駄は出ないかという点だけど、それについてはホスト側で判定すればいいだけじゃないかと。
--
const int NumThreads = 512; // 実際には192以上の32の倍数で調整すべきかも
const int NumBlocks = std::min((size + NumThreads - 1) / NumThreads * NumThreads, 1000); // block数の最大を1000としてみた
changeA<<<dim3(NumBlocks, 1, 1), dim3(NumThreads, 1, 1)>>>(a, b, size)
676:デフォルトの名無しさん
10/10/15 14:25:43
>>674
スレッド数が10だとsizeが100なのに10件しか処理しないという自明なことを書いているだけかと思った。
677:デフォルトの名無しさん
10/10/15 17:56:27
>>676
あー。。。なるほどサンクス
そりゃそーだろう...
678:デフォルトの名無しさん
10/10/15 20:22:23
>>670
この文脈だと、カーネルをコールする時のオーバーヘッドの事