【GPGPU】くだすれCUDAスレ part8【NVIDIA】 [無断転載禁止]©2ch.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/ >>343 お気持ちC++実装でフル規格満たしてないとか普通にある。調べるのが手間すぎる。 フル規格を満たしてない事が普通にあるのはCも同じ 恐くて使えないのは経験が少ないからだな MSVCで作っておけば大体オッケー g++依存は死ぬ 当然 高速処理が要求されるようなところはC/C++だよ 質問です。 cudaの9.0バージョンで nvcc -V このようなことをうっても、一切反応しないです。 誰か教えてえらいひとーー! 無修正を観ても、一切反応しないです。 誰か教えてえろいひとーー! NVIDIA HPC SDK使ってる人いる? windows版がないんだけどそのうちサポートされるんかね? HPC SDKになった2020年にはWindows版は翌年公開予定と言っていたけど, その記載もなくなったし永遠に出ない可能性が高いと思われる 当時とは違ってCUDA on WSL2でHPC SDKも使えるようになったから Windows需要もそこで解消されるし >>354 詳しい人降臨キター 旧PGIのユーザはほぼLinuxばっかだったようだしWindowsは見捨てられたのかと思ったけど必ずしもそうじゃないのかな WSLで本当に性能出るの?ってのは気にはなるけど ありがとうございました 亀レス >>348 あるっしょ いくらでも, このまえいびられて死んだ三菱電機の社員は電気系でも物性系の修士出てるのに 会社に入ってC++まともに書けなくて死んだ 物性選んだ時点でプログラム苦手なのにいきなりC売り物用のC++コード書け, しかも,意地の悪い上司が適当なサンプルコードも見せてくれなきゃ死にたくなるかもな おまえら教えろください 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関数を一切使わず高速処理が可能となります 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__でコールできないので駄目なんですよ 何もレスないけど cudaDeviceSynchronize なくても困らないコード書いてるの? shared memory使うときどーしても使うと思うんだがこのAPI ふつうは __syncthreads() 使うなあ。 というかカーネル内から全スレッド同期出来たとは知らなかった。 やっぱりいろいろ無理があったから廃止になったんじゃね? カーネルの実行完了を待つ関数をカーネル内から呼び出せるのはおかしいだろ >>362 >>363 お前らマジでレベル低いわ シェアードメモリとかまったく使いこなせてないだろ シェアードメモリってのはブロック内スレッドでのみ共通のメモリ カーネル関数内で待ち合わせ処理できずにどうやって使うんだこんなもん >>363 にいたってはCUDA Dynamic Parallelizmすら理解してない >>362 ふつうは? __syncthreads()自体とっくに廃止されてそのかわりのcudaDeviceSynchronize()だろ >>362 CUDA Runtime API rev.1 をみても __syncthreadsなんて載ってないんだが, どこに載ってるか教えてくれ >>362 https://docs.nvidia.com/cuda/pdf/CUDA_Runtime_API.pdf ↑から最新のAPIマニュアルをダウソして検索したが __syncthreadsなんてものは存在しない お前の書いたコードが最新のcuda 12.1.1コンパイラで使えるかどうか確認してみろ とっくにサポートが打ち切られたAPIを使うと嘯吹いてただけじゃねーのか? お前はとんでもなく長い間cudaコードを書いてなかったんじゃないのか? カーネルからRuntime API使うのがそもそも違和感あるんだがイマドキのCUDAはそういうものなんか? submit済みカーネルが2並列同じキューで動いていて、両方が完了待ちしたらデッドロック起こすなwww プログラムわかってないやつがいきなりの攻撃性を発揮するのガイジっぽくていいね 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>>>(); } つづき ホスト側からcuda用のエントリポイント cuda_main()を一つ起動して そこで計算に必要なgpuメモリをnew/mallocで確保 delete/freeもしくはデストラクタでメモリ解放してほとんどふつーのC++プログラム作成 エントリポイントとなるcuda_mainスレッドを一つ起動するだけでほぼgpu側だけで閉じたふつーのc++コードになる gpuが計算した結果をどー出力するか?ってのはあるんだが、 __global__関数内のprintfが標準出力にちゃーんと出力してくれる つまり、 リダイレクトでファイルに落とせる コマンドラインから > cuda_program.exe >> file.dat て感じ つづき こういうコードスタイルにすることで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があっても上のようなコードは書けない >>374 同じstream使ってんのにいちいちそんなところで同期してたら効率悪くね? >>377 いちいちも何も 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側に制御戻す必要なく処理速度が落ちるわけじゃない. それどころかはるかにプログラムの見通しもよくなるんだから 丸々一ヶ月経ってもまだ躓いてるのかw せっかくアドバイス貰っても理解する気がないならしょうがないな。 >>379 dynamic parallelismがどうたらイキってた輩が実はstreamすらろくに理解していなかったという そもそもShared Memoryが云々の話とこれ全く関係ないし、unique_ptrがどうとか誰も聞いてないし、ただイキってクソコード書いてるnoobにしか見えん >>381 streamてCPUとGPUの協調のこと言ってるの? 協調はこの際,まったく関係ないんだが, CPUからGPUに制御を移して完全にCPUとは分離する話をしてんの そのほうが遥かにプログラムが楽に簡潔に書けるからそれを言ってる. GPUのほうが8倍から1桁速度が速いので実験的な計算ならこれで十分な速度が得られる CPUとの協調とか難しいプログラムを書くんじゃなく GPUだけで計算を完結させるためにDynamic Parallelismを使う Dynamicに並列化したいわけじゃなくGPUのプロセスを一つ立ち上げたいためにそれを使う >>380 何のアドバイス?ww GPUプロセスの中で全プロセスどうやってwaitするのさw ま, cudaのユーザサイトには代替案とか書いてる人いるよ cudaDeviceSynchronize()と比較して速度は期待できないが一応使える. せめてそれぐらいのアドバイスできるようになってから言いな. 俺のようなコード書いてGPUプロセスの中からcudaDeviceSynchronize()コールしてるユーザは一定数いるってこった どーせ__global__でnew/deleteふつーに使えること知らんかったんだろお前ww >>382 誰も聞いてないunique_ptrて CPU側でGPUのマルチスレッド実体を起動しようとすると, GPU内部でしか利用しないローカルメモリまでCPUから管理する羽目になる そのときcudamalloc/cudafreeなんか使ってたんじゃメモリ管理が大変で, gpu::unique_ptrでも作らないとCコード書いてるのと変わらんことになるから言ってる そんなことせずとも__global__でnew/deleteがデフォで使えるんだからメモリ管理はC++03並には書ける __global__でnew/deleteを使ってGPUだけで完結処理するには 親スレッドの__global__内で子スレッドの待ち合わせ処理がどうしても必要になる場合があるから言ってる できるもんなら>>374 のcudaDeviceSynchronize(); を同等処理に置き換えてみ あ, すでに cudaのユーザーサイトに投稿してる人が居るんで探せばしまいだけどねww 俺はCUDA使ったこと無いんだけど >__global__でnew/deleteを使ってGPUだけで完結処理 これって思い込みじゃねえかなあ >>386 作ったことないんだろww こっちは書いてるから言ってる. >>374 を見ろ この構成の場合がそれに該当する 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しか使えないけどな. >>386 ウソだと思うなら,>>374 のコードを実際cudaでコンパイルして確認してみろや 実際__global__内でnew/malloc使えて__global__から__global__を起動できることだけがポイントなんだから ちなみに, dynamic parallelismサポートしてないような古いグラボは使用禁止な. >>386 new/deleteが__global__内でふつーに使えるってことなら cuda c programming guide 最新版の 7.34. Dynamic Global Memory Allocation and Operations どっちが思い込みかね?ww もしかしてNVIDIAのフォーラムでモデレータに I’m confused とか言われてんのコイツ? 炎上学習法かとも思ったが全然学習してる素振りもないのよなあ >streamてCPUとGPUの協調のこと言ってるの? ちょっとはググるなりしたらどうかね > body1<<< Ki, Ki >>> ( idata, odata1 ); //マルチスレッド実体 > cudaDeviceSynchronize(); > body2<<< Ki, Ki >>> ( odata1, odara2 ); //マルチスレッド実体 body1とbody2は同じstreamだから基本的にそんなところにsynchronizeはいらんのよ >そのときcudamalloc/cudafreeなんか使ってたんじゃメモリ管理が大変で, これも意味不明だなあ malloc/freeと比べてどう大変だと言うんだろう 同期オブジェクトが出てないんだから、そりゃシリアルに動くやろうなあ CUDA知らん >>390 ゲラゲラ それどこだよww お前のプロファイル推定正しいかどうか見てから発表してやるよ さらせよソコをよww >>391 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も読んどけよ コロコロ変わるデフォのコンパイルオプションに頼るお前 >>392 freeが毎回書いてられないといってるわけ. それとcudafreeが同じなのは当たり前な だからホスト側ではgpu用のgpu::unique_ptrとかユーザ側でこさえてるんだろが. これを書いてるサイトはいくらでもある. __global__内ではこういうmake_uniqueは動かないのでc++11レベルでは書けないが, 普通にnew/deleteやC++03レベルのデストラクタが動作するので, ホスト側でgpuのローカルメモリの解放を手でやらかすよりよほど楽だと言ってるのだよ. >>395 いや、それストリーム間の同期の話であって同じストリームに投入したカーネルの実行の話とは関係ないんだが。 いろいろ検索したりしたんだろうけどここ勘違いするようなレベルでイキりまくられても。 もともとは cudaDeviceSynchronize がデバイス上で非推奨になるんで代替策をどうするかって話だったと思うけど、 「自分は間違ってない」と強弁するだけならそりゃ2ヶ月近く解決しないわけだわ。 PCくそど初心者で、オーディオをやっている者です。 現在PC(linux)でHQplayerというソフトで音楽ききてます。 そのPCにグラボ刺したら音質上がるなんてガセネタつかまされたんで、騙されてみようと思うのですが、その際CUDAの設定しないとGPUが使われないということのようなのです。 調べたら、ドライバー、ツールキット、cuDNNの3つをインストールするみたいなのですが、最後のcuDNNの意味がわかりません。 有識者の方、どうかご教授下さい! >>398 他人の作った何をやっているかわからないプログラムを使うより 音とプログラミングとCUDAについて勉強して 自分で音質を上げるプログラムを作ったほうがよいと思うよ。 本屋とかアマゾンでそういう専門書も探せば見つかるでしょ。 >>399 398です。 なるほど、そういう方法もあるのですね。 自分で勉強してプログラムのことまで理解するなんて大変そうですが、根本を理解していないと使う意味ないですもんね。 大変そうですが、チャレンジしてみます。 ありがとうございました。 RTX4090よりA4000をおすすめしてる所がありますが、これは長期稼働の安定性と低消費電力が理由でしょうか。 lstmで出来るだけ早く学習させたいのですが4090の方が早そうですがいかがですか。 自己解決 webページ作成者に聞いた 4090の方が速いが、例えば3時間が4時間になったところでたいした違いはないでしょうとか いくつものパターンを試したいので少しでも早いものをってことで4090を買った read.cgi ver 07.5.5 2024/06/08 Walang Kapalit ★ | Donguri System Team 5ちゃんねる