【GPGPU】くだすれCUDAスレ part8【NVIDIA】
1 :デフォルトの名無しさん :2016/06/15(水) 14:41:15.11 ID:d2Xou3GL.net このスレッドは、他のスレッドでは書き込めない超低レベル、 もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。 CUDA使いが優しくコメントを返しますが、 お礼はCUDAの布教と初心者の救済をお願いします。 CUDA・HomePage ttp://developer.nvidia.com/category/zone/cuda-zone 関連スレ GPGPU#5 http://peace.2ch.net/test/read.cgi/tech/1281876470/ 前スレ 【GPGPU】くだすれCUDAスレ【NVIDIA】 ttp://pc12.2ch.net/test/read.cgi/tech/1206152032/ 【GPGPU】くだすれCUDAスレ pert2【NVIDIA】 ttp://pc12.2ch.net/test/read.cgi/tech/1254997777/ 【GPGPU】くだすれCUDAスレ pert3【NVIDIA】 ttp://hibari.2ch.net/test/read.cgi/tech/1271587710/ 【GPGPU】くだすれCUDAスレ pert4【NVIDIA】 ttp://hibari.2ch.net/test/read.cgi/tech/1291467433/ 【GPGPU】くだすれCUDAスレ part5【NVIDIA】 http://toro.2ch.net/test/read.cgi/tech/1314104886/ 【GPGPU】くだすれCUDAスレ part6【NVIDIA】 ttp://peace.2ch.net/test/read.cgi/tech/1348409867/ 【GPGPU】くだすれCUDAスレ part7【NVIDIA】 http://echo.2ch.net/test/read.cgi/tech/1416492886/
359 :デフォルトの名無しさん :2023/01/23(月) 20:21:43.20 ID:VfNttDi/.net NVIDIA HPC SDK使ってる人いる? windows版がないんだけどそのうちサポートされるんかね?
360 :デフォルトの名無しさん :2023/01/23(月) 20:47:20.55 ID:1CzktcoW.net HPC SDKになった2020年にはWindows版は翌年公開予定と言っていたけど, その記載もなくなったし永遠に出ない可能性が高いと思われる 当時とは違ってCUDA on WSL2でHPC SDKも使えるようになったから Windows需要もそこで解消されるし
361 :デフォルトの名無しさん :2023/01/25(水) 11:08:54.16 ID:/YL2yMwg.net >>360 詳しい人降臨キター 旧PGIのユーザはほぼLinuxばっかだったようだしWindowsは見捨てられたのかと思ったけど必ずしもそうじゃないのかな WSLで本当に性能出るの?ってのは気にはなるけど ありがとうございました
362 :デフォルトの名無しさん :2023/03/05(日) 20:30:51.25 ID:skhIF3To.net てst
363 :デフォルトの名無しさん :2023/03/05(日) 21:06:34.13 ID:skhIF3To.net 亀レス >>354 あるっしょ いくらでも, このまえいびられて死んだ三菱電機の社員は電気系でも物性系の修士出てるのに 会社に入ってC++まともに書けなくて死んだ 物性選んだ時点でプログラム苦手なのにいきなりC売り物用のC++コード書け, しかも,意地の悪い上司が適当なサンプルコードも見せてくれなきゃ死にたくなるかもな
364 :デフォルトの名無しさん :2023/05/14(日) 05:25:29.03 ID:EUIAlyu+.net test
365 :デフォルトの名無しさん :2023/05/14(日) 06:02:02.90 ID:EUIAlyu+.net おまえら教えろください 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関数を一切使わず高速処理が可能となります
366 :デフォルトの名無しさん :2023/05/14(日) 06:21:14.34 ID:EUIAlyu+.net 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__でコールできないので駄目なんですよ
367 :デフォルトの名無しさん :2023/05/28(日) 07:37:34.81 ID:rAN5cbTU.net 何もレスないけど cudaDeviceSynchronize なくても困らないコード書いてるの? shared memory使うときどーしても使うと思うんだがこのAPI
368 :デフォルトの名無しさん :2023/05/30(火) 17:45:24.83 ID:GWh7Fu6m.net ふつうは __syncthreads() 使うなあ。 というかカーネル内から全スレッド同期出来たとは知らなかった。 やっぱりいろいろ無理があったから廃止になったんじゃね?
369 :デフォルトの名無しさん :2023/06/01(木) 19:34:35.47 ID:woZHz0xe.net カーネルの実行完了を待つ関数をカーネル内から呼び出せるのはおかしいだろ
370 :デフォルトの名無しさん :2023/06/02(金) 21:14:34.84 ID:GpIK8Zp9.net >>368 >>369 お前らマジでレベル低いわ シェアードメモリとかまったく使いこなせてないだろ シェアードメモリってのはブロック内スレッドでのみ共通のメモリ カーネル関数内で待ち合わせ処理できずにどうやって使うんだこんなもん
371 :デフォルトの名無しさん :2023/06/02(金) 21:16:18.40 ID:GpIK8Zp9.net >>369 にいたってはCUDA Dynamic Parallelizmすら理解してない
372 :デフォルトの名無しさん :2023/06/02(金) 21:29:40.49 ID:GpIK8Zp9.net >>368 ふつうは? __syncthreads()自体とっくに廃止されてそのかわりのcudaDeviceSynchronize()だろ
373 :デフォルトの名無しさん :2023/06/02(金) 21:32:58.33 ID:GpIK8Zp9.net >>368 CUDA Runtime API rev.1 をみても __syncthreadsなんて載ってないんだが, どこに載ってるか教えてくれ
374 :デフォルトの名無しさん :2023/06/02(金) 22:14:21.99 ID:GpIK8Zp9.net >>368 https://docs.nvidia.com/cuda/pdf/CUDA_Runtime_API.pdf ↑から最新のAPIマニュアルをダウソして検索したが __syncthreadsなんてものは存在しない お前の書いたコードが最新のcuda 12.1.1コンパイラで使えるかどうか確認してみろ とっくにサポートが打ち切られたAPIを使うと嘯吹いてただけじゃねーのか? お前はとんでもなく長い間cudaコードを書いてなかったんじゃないのか?
375 :デフォルトの名無しさん :2023/06/02(金) 23:26:41.73 ID:zbQswG7U.net カーネルからRuntime API使うのがそもそも違和感あるんだがイマドキのCUDAはそういうものなんか?
376 :デフォルトの名無しさん :2023/06/02(金) 23:38:04.08 ID:44SgMOSU.net >>372 同期する範囲がぜんぜん違ぇだろうが
377 :デフォルトの名無しさん :2023/06/03(土) 01:32:33.61 ID:f1mb9frl.net submit済みカーネルが2並列同じキューで動いていて、両方が完了待ちしたらデッドロック起こすなwww
378 :デフォルトの名無しさん :2023/06/03(土) 10:34:44.84 ID:hgrjzBME.net プログラムわかってないやつがいきなりの攻撃性を発揮するのガイジっぽくていいね
379 :デフォルトの名無しさん :2023/06/03(土) 16:01:42.94 ID:f1mb9frl.net >>374 ランタイムAPIwwww こっち見とけ https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html
380 :デフォルトの名無しさん :2023/06/10(土) 08:16:47.35 ID:gJM3u8Zc.net cudaDeviceSynchronizeはこんなときに使う. Ki=1024, Mi = Ki*Ki, Gi = Ki*Miとでもして __global__ void cuda_main(){ double *idata = new [Mi]; double *odata1 = new [Mi]; double *odata2 = new [Mi]; body1<<< Ki, Ki >>> ( idata, odata1 ); //マルチスレッド実体 cudaDeviceSynchronize(); body2<<< Ki, Ki >>> ( odata1, odara2 ); //マルチスレッド実体 cudaDeviceSynchronize(); for( int i = 0; i < Mi; i++ ){ cudaDeviceSynchronize(); // *** こいつは毎回要ったと思う printf(" %d %e\n", i, odata2[i]); //計算結果表示 } cudaDeviceSynchronize(); // 上の***だけで こいつはなくても構わなかったと思う delete[] odata2; delete[] odata1; delete[] idata; } main(){ //cuda 内newで確保するメモリが8MBを超える場合は設定要 cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t(Gi + Gi)); //printf fifoを16Miにしてみた cudaDeviceSetLimit(cudaLimitPrintfFifoSize, 16 * Mi); cuda_main<<<1、1>>>(); }
381 :デフォルトの名無しさん :2023/06/10(土) 08:17:55.93 ID:gJM3u8Zc.net つづき ホスト側からcuda用のエントリポイント cuda_main()を一つ起動して そこで計算に必要なgpuメモリをnew/mallocで確保 delete/freeもしくはデストラクタでメモリ解放してほとんどふつーのC++プログラム作成 エントリポイントとなるcuda_mainスレッドを一つ起動するだけでほぼgpu側だけで閉じたふつーのc++コードになる gpuが計算した結果をどー出力するか?ってのはあるんだが、 __global__関数内のprintfが標準出力にちゃーんと出力してくれる つまり、 リダイレクトでファイルに落とせる コマンドラインから > cuda_program.exe >> file.dat て感じ
382 :デフォルトの名無しさん :2023/06/10(土) 08:20:34.56 ID:gJM3u8Zc.net つづき こういうコードスタイルにすることでcudamallocで確保するgpuメモリはホスト<->gpuのインターフェースに使う最低限のメモリに限定することができて コーディングが格段に楽になる ほぼふつーのC++プログラムと変わらん この書き方で問題なのは__global__内でunique_ptrが使えないことなんだが、 数値計算でunique_ptr使えなくて困る場合ってどんな場合か逆に聞きたいわ. それより ホスト側でgpu内部処理にしか使わないメモリまで確保して、そいつをcuda::unique_ptrなんか作って管理するよりよほど楽だと思うが? こういうのはディレクトリの外からディレクトリ操作するのに似た感覚でとても耐えられんし、 cuda Dynamic Parallelismと__global__関数内でnew/deleteを書けるようになってるお陰でふつーのC++コードが書けるようになってるんだわ. んで肝心の cudaDeviceSynchronize() だが マルチスレッド実態は body1、 body2だがbody2の実行と計算結果odata2の表示はbody1とbody2の計算完了が前提としてる それらの待ち合わせ処理としてを使うんだが、 ここで cudaDeviceSynchronize() がないと dynamic Parallelismがあっても上のようなコードは書けない
383 :デフォルトの名無しさん :2023/06/10(土) 10:12:15.01 ID:5lxShvGN.net >>380 同じstream使ってんのにいちいちそんなところで同期してたら効率悪くね?
384 :デフォルトの名無しさん :2023/06/10(土) 11:53:19.12 ID:Y9DL3tjK.net まだやってたのかよw>>365
385 :デフォルトの名無しさん :2023/06/16(金) 10:37:03.14 ID:VMczRTMU.net >>383 いちいちも何も body1<<< Ki, Ki >>> ( idata, odata1 ); //マルチスレッド実体 cudaDeviceSynchronize(); body2<<< Ki, Ki >>> ( odata1, odara2 ); //マルチスレッド実体 cudaDeviceSynchronize(); body2の処理開始はbody1の計算終了が前提としたプログラムのつもりで書いてるのでこの場合は必ず必要なわけで この2つのの同期は マルチスレッド起動箇所がホスト側ならホストで マルチスレッド起動箇所がgpu側ならgpuでやることが必須なんだが gpuの側でやる意味はふつーに_global__関数の中で, ローカルメモリの管理にnew/delete使えること cpuのアクセラレータとしてGPUを使うんじゃなくGPUだけで完結するならこう書くでしょーよ わざわざcpu側に制御戻す必要なく処理速度が落ちるわけじゃない. それどころかはるかにプログラムの見通しもよくなるんだから
386 :デフォルトの名無しさん :2023/06/16(金) 20:00:03.41 ID:duxPCvpi.net 丸々一ヶ月経ってもまだ躓いてるのかw せっかくアドバイス貰っても理解する気がないならしょうがないな。
387 :デフォルトの名無しさん :2023/06/16(金) 23:41:50.06 ID:T6+41XgS.net >>385 dynamic parallelismがどうたらイキってた輩が実はstreamすらろくに理解していなかったという
388 :デフォルトの名無しさん :2023/06/17(土) 09:32:02.85 ID:6mG7lpSl.net そもそもShared Memoryが云々の話とこれ全く関係ないし、unique_ptrがどうとか誰も聞いてないし、ただイキってクソコード書いてるnoobにしか見えん
389 :デフォルトの名無しさん :2023/06/30(金) 06:45:59.93 ID:HGq0NS3h.net >>387 streamてCPUとGPUの協調のこと言ってるの? 協調はこの際,まったく関係ないんだが, CPUからGPUに制御を移して完全にCPUとは分離する話をしてんの そのほうが遥かにプログラムが楽に簡潔に書けるからそれを言ってる. GPUのほうが8倍から1桁速度が速いので実験的な計算ならこれで十分な速度が得られる CPUとの協調とか難しいプログラムを書くんじゃなく GPUだけで計算を完結させるためにDynamic Parallelismを使う Dynamicに並列化したいわけじゃなくGPUのプロセスを一つ立ち上げたいためにそれを使う
390 :デフォルトの名無しさん :2023/06/30(金) 06:53:13.13 ID:HGq0NS3h.net >>386 何のアドバイス?ww GPUプロセスの中で全プロセスどうやってwaitするのさw ま, cudaのユーザサイトには代替案とか書いてる人いるよ cudaDeviceSynchronize()と比較して速度は期待できないが一応使える. せめてそれぐらいのアドバイスできるようになってから言いな. 俺のようなコード書いてGPUプロセスの中からcudaDeviceSynchronize()コールしてるユーザは一定数いるってこった どーせ__global__でnew/deleteふつーに使えること知らんかったんだろお前ww
391 :デフォルトの名無しさん :2023/06/30(金) 07:16:01.83 ID:HGq0NS3h.net >>388 誰も聞いてないunique_ptrて CPU側でGPUのマルチスレッド実体を起動しようとすると, GPU内部でしか利用しないローカルメモリまでCPUから管理する羽目になる そのときcudamalloc/cudafreeなんか使ってたんじゃメモリ管理が大変で, gpu::unique_ptrでも作らないとCコード書いてるのと変わらんことになるから言ってる そんなことせずとも__global__でnew/deleteがデフォで使えるんだからメモリ管理はC++03並には書ける __global__でnew/deleteを使ってGPUだけで完結処理するには 親スレッドの__global__内で子スレッドの待ち合わせ処理がどうしても必要になる場合があるから言ってる できるもんなら>>380 のcudaDeviceSynchronize(); を同等処理に置き換えてみ あ, すでに cudaのユーザーサイトに投稿してる人が居るんで探せばしまいだけどねww
392 :デフォルトの名無しさん :2023/06/30(金) 08:24:20.23 ID:pBAX8tCp.net 俺はCUDA使ったこと無いんだけど >__global__でnew/deleteを使ってGPUだけで完結処理 これって思い込みじゃねえかなあ
393 :デフォルトの名無しさん :2023/06/30(金) 11:03:40.92 ID:HGq0NS3h.net >>392 作ったことないんだろww こっちは書いてるから言ってる. >>380 を見ろ この構成の場合がそれに該当する CPUはgpu内でnew/deleteに使う最大値のみ設定する.デフォルトは小さいんでな 例えば 1.GPUで用意してる数学関数だけでできる数値計算 2.信号もノイズも乱数から作るモンテカルロ・シミュレーション 要はファイルやなんらかのデバイス,ネットから読み込んでその処理をGPUに転送する場合は CPUは無関係じゃいられないが, 上のような場合はGPUだけで計算可能だ. 具体例として 1. f(x) = ∫_0^x sin(t)/t dt 0<= x <= 10 を計算して計算結果をファイルにセーブする 2. ある変調方式のエラーレートをシミュレーションで計算し. ファイルにセーブする. これらならCPUとほぼ無関係に計算できる ファイルや外部デバイスからデータを読み込んでくる必要がないんでね. あと__global__でのprintfはふつーにリダイレクトが有効. GPU内でファイルオープンとかの必要もなくファイルに落とせる ま,これらは極端な例だが, 最低限, 初期値とか処理データ, GPUで計算終了後の GPU-CPU間のインターフェース用メモリのみcudamalloc/dudafreeで管理してやればよくて GPU内部だけで使うようなメモリはCPU側で一切管理する必要がなくなる. gpu内ではC++03レベルのnew/deleteしか使えないけどな.
394 :デフォルトの名無しさん :2023/06/30(金) 11:07:50.35 ID:HGq0NS3h.net >>392 ウソだと思うなら,>>380 のコードを実際cudaでコンパイルして確認してみろや 実際__global__内でnew/malloc使えて__global__から__global__を起動できることだけがポイントなんだから ちなみに, dynamic parallelismサポートしてないような古いグラボは使用禁止な.
395 :デフォルトの名無しさん :2023/06/30(金) 11:38:34.15 ID:HGq0NS3h.net >>392 new/deleteが__global__内でふつーに使えるってことなら cuda c programming guide 最新版の 7.34. Dynamic Global Memory Allocation and Operations どっちが思い込みかね?ww
396 :デフォルトの名無しさん :2023/06/30(金) 18:28:44.41 ID:h/tbEWPQ.net もしかしてNVIDIAのフォーラムでモデレータに I’m confused とか言われてんのコイツ?
397 :デフォルトの名無しさん :2023/06/30(金) 22:03:38.67 ID:nvcNe2IT.net 炎上学習法かとも思ったが全然学習してる素振りもないのよなあ >streamてCPUとGPUの協調のこと言ってるの? ちょっとはググるなりしたらどうかね > body1<<< Ki, Ki >>> ( idata, odata1 ); //マルチスレッド実体 > cudaDeviceSynchronize(); > body2<<< Ki, Ki >>> ( odata1, odara2 ); //マルチスレッド実体 body1とbody2は同じstreamだから基本的にそんなところにsynchronizeはいらんのよ
398 :デフォルトの名無しさん :2023/06/30(金) 22:07:49.01 ID:nvcNe2IT.net >そのときcudamalloc/cudafreeなんか使ってたんじゃメモリ管理が大変で, これも意味不明だなあ malloc/freeと比べてどう大変だと言うんだろう
399 :デフォルトの名無しさん :2023/06/30(金) 23:46:28.63 ID:lTu2cFop.net 病気の人に関わってもいいことないですよ
400 :デフォルトの名無しさん :2023/07/02(日) 11:33:55.45 ID:ZzvN3CPt.net 同期オブジェクトが出てないんだから、そりゃシリアルに動くやろうなあ CUDA知らん
401 :デフォルトの名無しさん :2023/07/02(日) 14:47:53.06 ID:tUgy2gS2.net >>396 ゲラゲラ それどこだよww お前のプロファイル推定正しいかどうか見てから発表してやるよ さらせよソコをよww >>397 12.2 プログラムガイド pp47 For code that is compiled using the --default-stream legacy compilation flag, the default stream is a special stream called the NULL stream and each device has a single NULL stream used for all host threads. The NULL stream is special as it causes implicit synchronization as described in Implicit Synchronization. For code that is compiled without specifying a --default-stream compilation flag, --default-stream legacy is assumed as the default. 11.4以降 --default-streamは非推奨. 当然このオプションはデフォルトでなくなり, 暗黙の同期ストリームであるNull streamはデフォルトではなくなった. つまり, ストリームは何も指定しなければ非同期ストリームとなった 同期ストリームとなることを保証したい場合, 当然cudaStreamSynchronize()で挟むだろうが そのすぐ下の6.2.8.5.3 Explicit Synchronizationも読んどけよ コロコロ変わるデフォのコンパイルオプションに頼るお前 >>398 freeが毎回書いてられないといってるわけ. それとcudafreeが同じなのは当たり前な だからホスト側ではgpu用のgpu::unique_ptrとかユーザ側でこさえてるんだろが. これを書いてるサイトはいくらでもある. __global__内ではこういうmake_uniqueは動かないのでc++11レベルでは書けないが, 普通にnew/deleteやC++03レベルのデストラクタが動作するので, ホスト側でgpuのローカルメモリの解放を手でやらかすよりよほど楽だと言ってるのだよ.
402 :デフォルトの名無しさん :2023/07/02(日) 19:20:21.21 ID:nj7sKJew.net >>401 いや、それストリーム間の同期の話であって同じストリームに投入したカーネルの実行の話とは関係ないんだが。 いろいろ検索したりしたんだろうけどここ勘違いするようなレベルでイキりまくられても。 もともとは cudaDeviceSynchronize がデバイス上で非推奨になるんで代替策をどうするかって話だったと思うけど、 「自分は間違ってない」と強弁するだけならそりゃ2ヶ月近く解決しないわけだわ。
403 :デフォルトの名無しさん :2023/10/20(金) 12:42:50.23 ID:/M3RKJCH.net NVIDIA中國から撤退
404 :デフォルトの名無しさん :2023/12/07(木) 06:26:05.21 ID:R1AO2r8W.net PCくそど初心者で、オーディオをやっている者です。 現在PC(linux)でHQplayerというソフトで音楽ききてます。 そのPCにグラボ刺したら音質上がるなんてガセネタつかまされたんで、騙されてみようと思うのですが、その際CUDAの設定しないとGPUが使われないということのようなのです。 調べたら、ドライバー、ツールキット、cuDNNの3つをインストールするみたいなのですが、最後のcuDNNの意味がわかりません。 有識者の方、どうかご教授下さい!
405 :デフォルトの名無しさん :2023/12/10(日) 12:43:06.09 ID:9z8kD1aM.net >>404 他人の作った何をやっているかわからないプログラムを使うより 音とプログラミングとCUDAについて勉強して 自分で音質を上げるプログラムを作ったほうがよいと思うよ。 本屋とかアマゾンでそういう専門書も探せば見つかるでしょ。
406 :デフォルトの名無しさん :2023/12/13(水) 06:47:58.97 ID:YKWD9gfa.net >>405 398です。 なるほど、そういう方法もあるのですね。 自分で勉強してプログラムのことまで理解するなんて大変そうですが、根本を理解していないと使う意味ないですもんね。 大変そうですが、チャレンジしてみます。 ありがとうございました。
407 :デフォルトの名無しさん :2024/01/13(土) 07:53:38.21 ID:Ce28ohDx.net RTX4090よりA4000をおすすめしてる所がありますが、これは長期稼働の安定性と低消費電力が理由でしょうか。 lstmで出来るだけ早く学習させたいのですが4090の方が早そうですがいかがですか。
408 :デフォルトの名無しさん :2024/03/14(木) 19:33:05.60 ID:LpPl9eeZ.net 自己解決 webページ作成者に聞いた 4090の方が速いが、例えば3時間が4時間になったところでたいした違いはないでしょうとか いくつものパターンを試したいので少しでも早いものをってことで4090を買った
113 KB
新着レスの表示
掲示板に戻る
全部
前100
次100
最新50
read.cgi ver 2014.07.20.01.SC 2014/07/20 D ★
本文 スレッドタイトル 投稿者