【GPGPU】くだすれCUDAスレ part8【NVIDIA】 [無断転載禁止]©2ch.net (407ﾚｽ)
上下前次1-新
抽出解除 必死ﾁｪｯｶｰ(本家) (べ) 自ID ﾚｽ栞 あぼーん

---

ﾘﾛｰﾄﾞ規制です｡10分ほどで解除するので､他のﾌﾞﾗｳｻﾞへ避難してください｡

---

358: デフォルトの名無しさん [sage] 2023/05/14(日) 05:25:29.03 ID:EUIAlyu+(1/3) AAS
test

---

359(1): デフォルトの名無しさん [sage] 2023/05/14(日) 06:02:02.90 ID:EUIAlyu+(2/3) AAS
おまえら教えろください
cudaDeviceSynchronizeが将来サポート打ち切りになることが決定し,
CUDA 12では既にオプションを指定しないとコンパイルできなくなっています.
これの代替APIは何でしょうか?
__global__関数内での待ち合わせ処理はどうすればいいのでしょう?

例
グローバルメモリに1M(1kx1k)の数値データがあります.そこから最大値を検索したい.

プログラムとしては
1Mデータを1kごとにグループ分割します.合計1k個のグループができます.
1つのグループを1つのスレッドに割り当て最大値を検索します. 当然スレッド数も1k個立ち上げます.
各スレッドは割り当てられたグループの1k個の中から最大値を検索し,見つけた値をシェアードメモリに書き込みます.
シェアードメモリもスレッド数と同じ1k個の配列から成リ, 各スレッドIDに紐づけされたアドレスに書き込みます.
例えばシェアードメモリを配列submaxとしたとき,
submax[threadIDx.x] = each_group_max;
みたいな感じです.

このとき各スレッドの書き込み完了を待ち合わせるのにcudaDeviceSynchronizeを使います.
書き込み完了後, 特定スレッドIDをもつスレッド(例えばthreadIDx.x==0)だけが1k個のsubmax中の最大値を検索することで
1M個データの最大値が決まります
これですとatomic関数を一切使わず高速処理が可能となります

---

360: デフォルトの名無しさん [sage] 2023/05/14(日) 06:21:14.34 ID:EUIAlyu+(3/3) AAS
Dynamic parallelismでも使える実行時間測定関数でも
キャプチャ直前にcudaDeviceSynchronize();をコールしています

__device__
long long int
CaptureTimer() noexcept {
auto timer = (long long int)0; //コンパイラが文句いうので初期化しました
cudaDeviceSynchronize();
asm volatile("mov.u64 %0, %globaltimer;" : "=l"(timer));

return timer;
}

cudaStreamSynchronize()じゃ__global__でコールできないので駄目なんですよ

---

上下前次1-新書関写板覧索設栞歴

---

ｽﾚ情報 赤ﾚｽ抽出 画像ﾚｽ抽出 歴の未読ｽﾚ AAｻﾑﾈｲﾙ

---

ぬこの手 ぬこTOP 0.013s