【GPGPU】くだすれCUDAスレ part8【NVIDIA】 [無断転載禁止]©2ch.net

**デフォルトの名無しさん** · 2023/06/02(金) 23:38:04.08

>>366
同期する範囲がぜんぜん違ぇだろうが

**デフォルトの名無しさん** · 2023/06/03(土) 01:32:33.61

submit済みカーネルが2並列同じキューで動いていて、両方が完了待ちしたらデッドロック起こすなｗｗｗ

**デフォルトの名無しさん** · 2023/06/03(土) 10:34:44.84

プログラムわかってないやつがいきなりの攻撃性を発揮するのガイジっぽくていいね

**デフォルトの名無しさん** · 2023/06/03(土) 16:01:42.94

>>368
ランタイムAPIｗｗｗｗ
こっち見とけ
https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

**デフォルトの名無しさん** · 2023/06/10(土) 08:16:47.35

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>>>();
}

**デフォルトの名無しさん** · 2023/06/10(土) 08:17:55.93

つづき
ホスト側からcuda用のエントリポイント cuda_main()を一つ起動して
そこで計算に必要なgpuメモリをnew/mallocで確保
delete/freeもしくはデストラクタでメモリ解放してほとんどふつーのC++プログラム作成
エントリポイントとなるcuda_mainスレッドを一つ起動するだけでほぼgpu側だけで閉じたふつーのc++コードになる

gpuが計算した結果をどー出力するか?ってのはあるんだが、
__global__関数内のprintfが標準出力にちゃーんと出力してくれる
つまり、リダイレクトでファイルに落とせる
コマンドラインから
> cuda_program.exe >> file.dat
て感じ

**デフォルトの名無しさん** · 2023/06/10(土) 08:20:34.56

つづき
こういうコードスタイルにすることで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があっても上のようなコードは書けない

**デフォルトの名無しさん** · 2023/06/10(土) 10:12:15.01

>>374
同じstream使ってんのにいちいちそんなところで同期してたら効率悪くね？

**デフォルトの名無しさん** · 2023/06/10(土) 11:53:19.12

まだやってたのかよw>>359

**デフォルトの名無しさん** · 2023/06/16(金) 10:37:03.14

>>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側に制御戻す必要なく処理速度が落ちるわけじゃない.
それどころかはるかにプログラムの見通しもよくなるんだから

**デフォルトの名無しさん** · 2023/06/16(金) 20:00:03.41

丸々一ヶ月経ってもまだ躓いてるのかw
せっかくアドバイス貰っても理解する気がないならしょうがないな。

**デフォルトの名無しさん** · 2023/06/16(金) 23:41:50.06

>>379
dynamic parallelismがどうたらイキってた輩が実はstreamすらろくに理解していなかったという

**デフォルトの名無しさん** · 2023/06/17(土) 09:32:02.85

そもそもShared Memoryが云々の話とこれ全く関係ないし、unique_ptrがどうとか誰も聞いてないし、ただイキってクソコード書いてるnoobにしか見えん

**デフォルトの名無しさん** · 2023/06/30(金) 06:45:59.93

>>381
streamてCPUとGPUの協調のこと言ってるの?
協調はこの際,まったく関係ないんだが,
CPUからGPUに制御を移して完全にCPUとは分離する話をしてんの
そのほうが遥かにプログラムが楽に簡潔に書けるからそれを言ってる.
GPUのほうが8倍から1桁速度が速いので実験的な計算ならこれで十分な速度が得られる
CPUとの協調とか難しいプログラムを書くんじゃなく
GPUだけで計算を完結させるためにDynamic Parallelismを使う
Dynamicに並列化したいわけじゃなくGPUのプロセスを一つ立ち上げたいためにそれを使う

**デフォルトの名無しさん** · 2023/06/30(金) 06:53:13.13

>>380
何のアドバイス?ww
GPUプロセスの中で全プロセスどうやってwaitするのさw
ま, cudaのユーザサイトには代替案とか書いてる人いるよ
cudaDeviceSynchronize()と比較して速度は期待できないが一応使える.
せめてそれぐらいのアドバイスできるようになってから言いな.
俺のようなコード書いてGPUプロセスの中からcudaDeviceSynchronize()コールしてるユーザは一定数いるってこった
どーせ__global__でnew/deleteふつーに使えること知らんかったんだろお前ww

**デフォルトの名無しさん** · 2023/06/30(金) 07:16:01.83

>>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

**デフォルトの名無しさん** · 2023/06/30(金) 08:24:20.23

俺はCUDA使ったこと無いんだけど

＞__global__でnew/deleteを使ってGPUだけで完結処理

これって思い込みじゃねえかなあ

**デフォルトの名無しさん** · 2023/06/30(金) 11:03:40.92

>>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しか使えないけどな.

**デフォルトの名無しさん** · 2023/06/30(金) 11:07:50.35

>>386
ウソだと思うなら,>>374のコードを実際cudaでコンパイルして確認してみろや
実際__global__内でnew/malloc使えて__global__から__global__を起動できることだけがポイントなんだから
ちなみに, dynamic parallelismサポートしてないような古いグラボは使用禁止な.

**デフォルトの名無しさん** · 2023/06/30(金) 11:38:34.15

>>386

new/deleteが__global__内でふつーに使えるってことなら
cuda c programming guide 最新版の
7.34. Dynamic Global Memory Allocation and Operations

どっちが思い込みかね?ww

**デフォルトの名無しさん** · 2023/06/30(金) 18:28:44.41

もしかしてNVIDIAのフォーラムでモデレータに I’m confused とか言われてんのコイツ？

**デフォルトの名無しさん** · 2023/06/30(金) 22:03:38.67

炎上学習法かとも思ったが全然学習してる素振りもないのよなあ

>streamてCPUとGPUの協調のこと言ってるの?

ちょっとはググるなりしたらどうかね

>　body1<<< Ki, Ki >>> ( idata, odata1 ); //マルチスレッド実体
>　cudaDeviceSynchronize();
>　body2<<< Ki, Ki >>> ( odata1, odara2 ); //マルチスレッド実体

body1とbody2は同じstreamだから基本的にそんなところにsynchronizeはいらんのよ

**デフォルトの名無しさん** · 2023/06/30(金) 22:07:49.01

>そのときcudamalloc/cudafreeなんか使ってたんじゃメモリ管理が大変で,

これも意味不明だなあ
malloc/freeと比べてどう大変だと言うんだろう

**デフォルトの名無しさん** · 2023/06/30(金) 23:46:28.63

病気の人に関わってもいいことないですよ

**デフォルトの名無しさん** · 2023/07/02(日) 11:33:55.45

同期オブジェクトが出てないんだから、そりゃシリアルに動くやろうなあ
CUDA知らん

**デフォルトの名無しさん** · 2023/07/02(日) 14:47:53.06

>>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のローカルメモリの解放を手でやらかすよりよほど楽だと言ってるのだよ.

**デフォルトの名無しさん** · 2023/07/02(日) 19:20:21.21

>>395
いや、それストリーム間の同期の話であって同じストリームに投入したカーネルの実行の話とは関係ないんだが。
いろいろ検索したりしたんだろうけどここ勘違いするようなレベルでイキりまくられても。

もともとは cudaDeviceSynchronize がデバイス上で非推奨になるんで代替策をどうするかって話だったと思うけど、
「自分は間違ってない」と強弁するだけならそりゃ2ヶ月近く解決しないわけだわ。

**デフォルトの名無しさん** · 2023/10/20(金) 12:42:50.23

NVIDIA中國から撤退

**デフォルトの名無しさん** · 2023/12/07(木) 06:26:05.21

PCくそど初心者で、オーディオをやっている者です。

現在PC（linux）でHQplayerというソフトで音楽ききてます。

そのPCにグラボ刺したら音質上がるなんてガセネタつかまされたんで、騙されてみようと思うのですが、その際CUDAの設定しないとGPUが使われないということのようなのです。

調べたら、ドライバー、ツールキット、cuDNNの３つをインストールするみたいなのですが、最後のcuDNNの意味がわかりません。

有識者の方、どうかご教授下さい！

**デフォルトの名無しさん** · 2023/12/10(日) 12:43:06.09

>>398
他人の作った何をやっているかわからないプログラムを使うより
音とプログラミングとCUDAについて勉強して
自分で音質を上げるプログラムを作ったほうがよいと思うよ。
本屋とかアマゾンでそういう専門書も探せば見つかるでしょ。

**デフォルトの名無しさん** · 2023/12/13(水) 06:47:58.97

>>399

398です。
なるほど、そういう方法もあるのですね。
自分で勉強してプログラムのことまで理解するなんて大変そうですが、根本を理解していないと使う意味ないですもんね。

大変そうですが、チャレンジしてみます。

ありがとうございました。

**デフォルトの名無しさん** · 2024/01/13(土) 07:53:38.21

RTX4090よりA4000をおすすめしてる所がありますが、これは長期稼働の安定性と低消費電力が理由でしょうか。

lstmで出来るだけ早く学習させたいのですが4090の方が早そうですがいかがですか。

**デフォルトの名無しさん** · 2024/03/14(木) 19:33:05.60

自己解決
webページ作成者に聞いた

4090の方が速いが、例えば3時間が4時間になったところでたいした違いはないでしょうとか

いくつものパターンを試したいので少しでも早いものをってことで4090を買った

**デフォルトの名無しさん** · 2024/10/12(土) 23:57:12.31

faster_whisperがGTX1660で6.8倍に高速化した
動画のハードエンコードみたいなことしてる訳でもないのに、なんでこんなに効くんだろう

**デフォルトの名無しさん** · 2025/05/03(土) 21:22:39.00

CUDAを使うにはH100が必要なの？

**デフォルトの名無しさん** · 2025/05/03(土) 21:57:53.32

>>404
nvidiaのGPUなら多少古くてもcuda使えるよ
A100でもA6000 A4000 RTX3060でもいい

**デフォルトの名無しさん** · 2025/05/12(月) 20:36:49.62

clvkってなんか普通に記事あるけど魔境なやつ？

**デフォルトの名無しさん** · 2025/06/05(木) 17:19:33.94

wayland対応はよう