【GPGPU】くだすれCUDAスレ part8【NVIDIA】 [無断転載禁止]©2ch.net (407ﾚｽ)
【GPGPU】くだすれCUDAスレ part8【NVIDIA】 [無断転載禁止]©2ch.net http://mevius.5ch.net/test/read.cgi/tech/1465969275/
上
下
前次
1-
新
通常表示
512ﾊﾞｲﾄ分割
ﾚｽ栞
抽出解除
必死ﾁｪｯｶｰ(本家)
(べ)
自ID
ﾚｽ栞
あぼーん
ﾘﾛｰﾄﾞ規制
です｡10分ほどで解除するので､
他のﾌﾞﾗｳｻﾞ
へ避難してください｡
358: デフォルトの名無しさん [sage] 2023/05/14(日) 05:25:29.03 ID:EUIAlyu+ test http://mevius.5ch.net/test/read.cgi/tech/1465969275/358
359: デフォルトの名無しさん [sage] 2023/05/14(日) 06:02:02.90 ID:EUIAlyu+ おまえら教えろください 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関数を一切使わず高速処理が可能となります http://mevius.5ch.net/test/read.cgi/tech/1465969275/359
360: デフォルトの名無しさん [sage] 2023/05/14(日) 06:21:14.34 ID:EUIAlyu+ 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__でコールできないので駄目なんですよ http://mevius.5ch.net/test/read.cgi/tech/1465969275/360
ﾒﾓ帳
(0/65535文字)
上
下
前次
1-
新
書
関
写
板
覧
索
設
栞
歴
ｽﾚ情報
赤ﾚｽ抽出
画像ﾚｽ抽出
歴の未読ｽﾚ
AAｻﾑﾈｲﾙ
Google検索
Wikipedia
ぬこの手
ぬこTOP
0.013s