【GPGPU】くだすれCUDAスレ part8【NVIDIA】 [無断転載禁止]©2ch.net
ようやくわかりました Relocatable Deice Code をenableでコンパイルが通りました あちがとうございました 今のGeForceは倍精度遅すぎてなんとかならんか?とおもてたら Kelper世代のTeslaが8万円ぐらいで売ってるのな 1.8TFlps c.f. R3090 FP64 0.556TFPS ww なんそれッ k80とか外付けGPUBOXで正常動作すんのかな? 動作報告探してみたけど見つからないんでよくわからん 消費電力大きいし、使いたいときだけつないぎたい 8万円分AWSのGPUインスタンス借りたほうがよくね? 課金がアクセス時間なのか正味のGPU駆動時間かわからんけど 8万だとほぼほぼ1ケ月やね。k80本体は6.7万ってのを見つけたわ 機材レンタルじゃないのでもっと上衣スペックのGPUサービスとか期待してチェックしたけど 目ん玉飛び出るぐらい高いね。もっと安くしてくれ A100ならAWSよりGPUSOROBANが半額以下やね 22万/月なんて個人じゃ出せないけど 仮面ライダー GPGPU https://pc.watch.impress.co.jp/docs/news/1396728.html 700Wて 10年たってジャンクで入手しても個人で稼動して遊べるシロモノじゃなくなったな CUDAてなんで3次元でスレッド管理するの? DRAMが実装構造としては2次元でもアドレスでコーダがrow. columnは吸収して1次元で管理するように スレッドもそうしちゃ駄目な理由を知りたい。 そもそも、3次元ってのとgrid, blockがなんの関係もないんじゃないの? 2次元を block として管理して 3次元の奥行が必要になった時点でgridが登場するならまだ理解できるんだけど。 しかも threads/block = 1024 と少ない。こんなんじゃ32x32行列越えたら一挙にメンドクサイ 3DCG処理やる場合、最初から3次元のID貰った方が楽じゃん? block,grid へのスレッドの割当て方とかでパフォーマンス変わりますか? 1000個のスレッドを起動したいとき、 block内だけで起動する方が速いのかいのか? それとも 10個/grid, 100個/blockと分ける方が速いのか? OpenCLでNVidiaがAMDほどパフォーマンスが上がらないのはスレッドの立ち上げかたで 癖のようなものがあるからってのはないですか? block,gridってパーティションは単にソフトの問題なのかハード上の制約でこうしてるのかが知りたい 変わるけど、今のGPUそんなん意識しなくても速いからいいだろ Keplerまでは色々考えないと速くならんかったけど、Maxwell以降はそんな考えなくてもよくなったからおけ なるほど じゃblockとかgridとかのパーティションはその時々作成するアプリ側の都合で 作りやすいように自由に設定していいってことですね? OpwnCL使うとCUDAより大分速度低下する理由は何ですかね? AMDはOpenCLしかないけどNVidiaでCUDAで書いたのと遜色ないぐらいのベンチマーク出してるサイトがあったんですが。 単にNvidiaがOpenCLやる気ないだけ? 昔ウチで試したときは同じGPUでCUDAとOpenCLはそんなに差は無かったがなぁ。 ものによってはOpenCLの方が速い場合もあったり。 そんなに差が出るのはバリバリにチューニングしたCUDAカーネルをOpenCLにベタ移植したとかかねぇ? 厳密には、ハードウェアを意識するならちゃんと考えたほうがいい あと適当に言ったけどさすがに今の世代でも1スレッドブロックしか立ち上げないとかだとめっちゃ遅いと思う GPUの並列性を埋めるなら、1SMに対して複数のスレッドブロックを当てられるぐらいであり、全SMに対してそれができるぐらい、つまり数万スレッドを立ち上げた方がいい とはいえレジスタ数との兼ね合いもあるからそこはバランスを見て OpenCLとCUDAに本質的な速度の差はないよ。最適化がどれぐらいかかるかの違いじゃない? 一般的な使い方をするならOpenCLはオンラインコンパイルだからあんま最適化かからんけど、CUDAなら最適化がかかるとかそういう感じ nVIDIAがCLやる気ないのはわかりきってることだろうに 一般にはオンラインコンパイラの方がその環境に合わせた最適化ができると思うがな。 NVIDIAがどこまでやってくれているかは知らないがIntel OpenCLは頑張っていた印象。 nVIDIAはclCreateProgramWithIL無いんだったなw 今度のゲフォ fp64で ラデなみの速度でるとか情報ないの? 11.7 にしたらoptix のdirect callがこわれた 今月のIF誌GPUプログラム特集ですなー AMDも含むけど 最近のIF誌Pythonとラスパイばっかでうんざりしてた IF誌読んだけど、 中級以上のCUSAプログラマ諸氏にはなんの価値もないです 意識してコーディングするとたちまち複雑化するシェアードメモリになんらかの言及があると思ったのですが、 サンプルプログラムはシェアードメモリで速度かせいでます。 だけでした そんな解説なら俺でも出来るわ >>246 初級の私には役立ちそうですね‥‥ぽち‥‥ と思うなら読んでみればいいよ Nvidiaのプログラミングガイド以上の情報あったら教えて >>246 読んでねえしCUDA知らんけど、書くにしても チャネルコンフリクトを避けるように書きましょう! たとえばこうやってずれるように・・・ って程度じゃねえかなあ cudaはグラボ2枚差し使えますか? コンシューマ向きのグラボだとvramが24Gですが 48Gになれば購入したいです。 そもそもvram24Gもってる人でどんな処理をしてるのか 知りたいです。例えば深層学習で〇時間を▽というグラボで こんだけ時間がかかるとか。 そういうデータがほしいです! よろしく! >>252 A6000 (48GB) 使ってるけど pix2pix とかやったら画像1000枚でも数日かかるよ。 vgg16 とか mobilenet で数時間の学習が多いけど、画像集めるのに一週間かかる。 積分する関数をCUDAで書こうと思っています リストのIntegrateのように関数テンプレートにすることで, 通常のC++であれば 呼び出し時に 被積分関数(ここではSinc)として 関数へのポインタでも, ラムダ式でも 直接関数名を書いてもうまく動作させることができるのですが, CUDAではうまくいかないようです. もちろんIntegrateの中で被積分関数を直接記述すれば問題なく動作するのですが, これだと毎回エディタでIntegrateの修正が必要になってきます. 呼び出し時の引数として被積分関数を与えるにはどーすればいいんでしょうか? プログラム リストです __global__ void //__device__ //float Sinc() { printf("Ahyahya\n"); //return start; } template< typename Functor> __global__ void Integrate(Functor func){ func <<<1, 1024>>> (); //Sinc << <1, 1024 >> > (); __syncthreads(); } int main() { Integrate <<<1, 1>>> (0, 10, Sinc); //Sinc << <1, 1024 >> > (); } main 中 誤 Integrate <<<1, 1>>> (0, 10, Sinc); 正 Integrate <<<1, 1>>> ( Sinc ); プログラムリストのコピペ時 0,10, を消すの忘れてました すまんが、MacやdGPUのないWindowsノートでコード書けたらいいなと思うんだけどさ CUDAを扱えるオンラインIDEとかエミュレーターとかなんかないもんなのかな? >>256 最近CUDA使ってないから間違ってるかもしれんが、__global__がついた関数からは__device__しか呼べなかったきがする。__global__関数から__global__が呼べるかちゃんとマニュアルで確認しよう。 >>258 google colaboratoryのGPUインスタンスで昔ちょっとしたCudaのコードを動かしたことはある。 基本はpythonで動かすことが前提なのでちょっと面倒なことしないとCのコードは動かせなかった。 後有料アカウントしないとGPUインスタンスに制限時間がつくかも。 >>259 いや__global__から__global__を呼ぶことは何も問題ない >>256 のリストにある関数Integrateの中でコメントアウトしてる部分がそれ, 直接呼ぶのは何も問題ないのに, ファンクタとして間接的に呼べないんで困ってる せんきゅー! Colaboratoryは怪しいことをしてると警告が出てビビっちゃうし、AWSは高いし、コードを書く時点で契約するには勇気がいるな・・・・ >>264 Linux使うのでもよかったら一万円ぐらいのjetson nano買えばCuda使えるよ んなもん使うぐらいならふつーにx64でコード書いたほうがよっぽどいい エンベ用途でもないのに何を好んでわざわざワンボードマイコンで計算せにゃならんのだ >>266 小さいものを愛でる気持ちがないのかな? お金が掛からないしGPIO使えるし。 GPUの性能が今売ってるGPUよりかだいぶ低いけど。 普通のPCからクロスコンパイルもできるらしいよ。 文脈から想像するに>>264 はCPUよりGPU計算が速いことを期待して ノートからGPUを使ってみたいんだろうよ それを単にCUDAが乗ってるだけでx64のマルチコアに遠く及ばないワンボードマイコンごときのGPUを何のためにわざわざ使う必要があるんだ? 組込用途とかでこれまでFPGA起こしたりASIC起こしたりしてた部分で ワンボードマイコンのGPUが代わることもあるだろう. 実際、産業機械にRasPiをまんま組込む用途も増えてる しかし、どーかんがえてもそっち系じゃない>>264 にワンボードマイコンすすめてどーすんだって話 コードを手元で書いて、実際に計算する際は実機を用意するかAWSを借りるつもりでいます ただコードを書くだけの段階でAWSを借りると負担なので、何かいい方法があればいいなと思った次第です >>270 https://www.%61mazon.co.jp/dp/B085TH77RT カードに安いやつかっても10万円コース AWSなくてもいけるが電気代要るわな それとThunderboltがボトルネックになることはほぼ確定 sharedメモリは https://zukaaax.com/archives/530 ここの最後に書いてる >シェアードメモリを使用しなくても高速にできるようにする方針 これホント? Dynamic Parallelism で親スレッドと子スレッドの属するブロックを 同一ブロックに属するような指定はできますか? 同一ブロックにして親スレッドが確保したShared Memory を子スレッドからアクセスしたいんだですが. 訂正です ○同一ブロックにして親スレッドが確保したShared Memory を親スレッドからアクセスしたいんですが. myKernel<<<gridSize, blockSize, nI*sizeof(int))>>>(...); 親スレッドで確保して子スレッドに引き渡すnI*sizeof(int)サイズのshared memoryですが この領域を親スレッドからアクセスできないんでしょーか? 事故レスです そもそも親子スレッド間でグローバルメモリのローカル変数の参照渡しが無理らしい ローカル変数なんかやめにして、大域変数にして__device__つけろとorz 当然のごとくshared memoryもだめっぽい C++でUser I/Fを使いやすく仕上げたいと思ってたけど そーゆーおしゃれなことは いろんな制限が頭に入っるまでは素直にゴリゴリCで書いてたほうがよさそう __global__のなかでnew,delete使えるので もしやと思ってmake_unique使おうとしたらだめだった make_unique以前にstd::unique_ptrの時点でだめだったわ cudaMalloc使ったunique_ptrバージョンはネットで公開してる人いるのめっけた https://proc-cpuinfo.fixstars.com/2019/02/cuda_smart_pointer/ host側からgpuメモリ確保したいならコレでいいけど __global__の中で使うmalloc,freeはGPUメモリを使うバージョンが正しくリンクされるので malloc, free使って各関数は__device__指定すれば, cuda::unique_ptrで正しく使えるのかな? cudaMallocやcudaFreeってのはホスト側からGPUとのインターフェース用として プログラムの中で1度呼ぶだけなんだよね GPUへの初期設定メモリ inmem GPUの計算結果 outmem として ホスト側で kernel<<<1,1>>>(inmem, outmem); を立ち上げる際 この前後でcudaMalloc, cudaMemcpy, cudaFreeを使うだけ kernel関数で一旦制御をGP側に移したあと おもむろにDynamic Parallelism を使ってkernel関数の中からマルチスレッドを起動する方がよっぽど書きやすい new/deleteは__global__関数のなかで自由に使えるので cudaMalloc/cudaFreeにuniqu_ptr使えてもあんま有り難みは少ないな 質問です。 cpuは64コア、gpuは年末のvram48ギガ2つ、メモリはddr5 256g これ以上は経済的事情により無理ですw vram48ギガ二つで深層学習したいんですけど どの規模でどのくらいのデータ処理ができますか? マルチcudaで処理速度はどれくらい高速化できますか? CUDA Python (Numba) を使用して気象分析の分野で 200,000 倍高速化する こういう記事がありました。 このようなデータがありましたが20万倍以上の高速化ってできますか? きになります。もしグラボがvram48一つと二つで差がなければ vram一つにしたいと思います。 >>279 VRAM48Gもあれば、相当の大きさのモデルでも、バッチサイズある程度取れると思う。 2枚にするなら、NVLINKやったかのボード間通信できるやつやらないと、 同時に使ってバッチサイズ倍にするぜってやつが遅い。 一枚でもいいんでない?って気はする。 何と比較して速くしたいのか分からんけど、3090ti(24GBメモリ)1枚でも普通のモデルなら ちゃんと学習できる。長いと数日学習に時間がかかるけど。 爆速学習したいなら、クラウドで大富豪するしかないぞ。 具体的に動かしたいモデルとかあるの? >>282 返信サンキュー!情報ありがとさん!! 気象予測がしたくて(とある記事に20万倍の高速化と書いてあった) クラウドいいですね〜! どうしてもpythonのnumbaがやりたくて! ちなみにpycharmかgoogle coalb pro plusはどちらが良いのでしょうか? >>254 cudaつかっても1週間もかかるんですか? 気象予測て流体力学だろ 学習とかの前に基礎方程式がわかってるものは 粒子フィルタで検討したほうがいいんじゃないの? もっぱら線形微分方程式に使われるカルマンフィルタを拡張した アンサンブルカルマンフィルタで見たことあるけど 今だと線形非線形にかかわらず利用可能な粒子フィルタに頼るほうが簡単だと思うけど RTX 4090 Ti 48GB GDDR6X これ買う人いますか? >>289 その辺のが一般化してきたら albertを使った実装例とかがもっと増えるか ま、もっと高性能なのがクラウドに多数デプロイされそうだけど RTX 4090 Tiを複数差せばおもしろそうやん? 600WのTBPで動作すると予想されます。 だそう。つまり3枚差せば1800wですね。 つまり電源ユニットを2枚で4000ワット以上必要ですね。 電源ユニットが3000wとか4000wとかでないのかな? 困るよね。 .cuファイルで何とか普通のメモリと GPUのVRAMを同居させて使えるなら 容量の大き過ぎる計算でもスピードは落ちるが 普通のCPUだけの時より速く計算できる、 とか出来たらいいんだけど、まず無理なんだろうなあ (沖電気はそういうのを上手くやったらしいが詳細が分からない) 何を問題視してるのかよくわからんけど 例えばmkl入れるとかじゃ駄目なのか? >>294 LinuxならUnifiedMemoryでできるな。 ホスト側のメモリをピンして使うことはできるけど多分そういう用途ではないよなあ >>294 が言ってるのはOversubscriptionだろ。 C#でCUDAを使おうとして悪戦苦闘してようやくこのスレにたどり着いたのだが・・・ GPUで計算した結果をCPUに取り出すには結局どうやったらいいんだ? 検索してもサンプルプログラムはほとんど出てこないし、GPU動かして終わりでその結果を使うやり方が示されてない。 教えろくださいお願いしますだ CUDAでdll作成してC#に読み込ませる こんだけ >>268 C言語で普通にCuda使うときだったら、GPU側にメモリ確保してGPUで計算した結果をそこに書き込む。 GPUでの計算が終わったらGPU側のメモリをCPU側のメモリにコピーするみたいな感じだ。 後unified memoryとかいうのでGPUから直接CPU側のメモリに書き込めるらしい。 C言語だったらCUDA sdkにサンプルコードがたくさん入ってるだけどね。 どういう関数使えばいいかはCUDAのドキュメントをちゃんと読めば見つけられるでしょう。 ありがとうございます。 昨日はとりあえず↓を参考にして以下のように書き換えてみたんですが、これじゃダメってことですよね? (textBox1の結果は0のままです) double temp = 0; var gpu = Gpu.Default; // GPUインスタンス取得 gpu.For(0, 10, i => { temp = 10; }); gpu.Synchronize(); // ここで同期する textBox1.Text = temp.ToString(); Alea GPUライブラリを使ってC#で簡単GPU並列プログラミング ? Crow's eye https://kzmmtmt.pgw.jp/?p=1170 Alea GPUで簡単C# GPUプログラミング - Qiita https://qiita.com/y_miyoshi/items/921903e3499abf18abdd Alea GPUなんて知らなかったな で、もう使われてなくね? それでやる必要ある? GPU使って並列計算できるってのを昨日知った初心者なんで、すいません。 「gpu 並列計算 C#」で検索してもこれしか出てこないんです。 で、あればGPUの使い方、CUDAの使い方を、まず勉強した方が良いのでは? まだGPUとホストのメモリ間のデータ転送とか、さっぱりわからないよね cudaは別にコンパイルしてC#から呼べば?このほうが調べやすそう 趣味でやっているだけなので調べてすぐにできればいいと思ったんですよ。 >>304 でGPU側で計算した結果を渡すだけなのにそれを乗せてるウェブサイトが見つからないなんて、そんな手間な作業なの? >>309 cpuとgpuで別々のメモリを見てるんで >>302 のようなことが必要です CUDA本体はC言語かC++言語で使う前提なのでそれらの言語ならサンプルコードや情報がたくさん見つかるんだけど。 C#からCUDA使うとなるとマイナーなライブラリを使うことになって情報を見つけづらいかもしれない。英語の情報までくまなく調べればそうじゃないのかもしれないが。 C#から使う場合は上にも書いた通り、cuda でdllを作成して C#からはDllImportで読み込む そうなんですね。ありがとうございます。 ライブラリが用意されているのでそれでできないのなら何のためのライブラリなのかと素人的には感じてしまいました。 出来ないかどうかはちゃんと調べないとわからないだろう 簡単に出来るかどうかはライブラリは保証しないよ、特にC#とCUDAみたいな組み合わせだったら。 めちゃくちゃ沢山のグルーコードをC++/CLIで書かなきゃいけないのを省略してくれる、ぐらいなもんでCUDAに対する理解なしに使える代物じゃないと思うけどな。 すいません、↓にテストコードが載っていたのでそれを動かしたところ、GPUの計算結果を取り出せました。 お騒がせしました。 Alea GPUで簡単C# GPUプログラミング - Qiita https://qiita.com/y_miyoshi/items/921903e3499abf18abdd CUDA初心者です。RTX2060 superでCUDA環境構築したいのですが、どうしてもcuda.is_availableの結果がFalseとなってしまいます。(Nvidiaのドライバ、Python・CUDA・Pytorchのバージョンを新しくしたり古くして見たり、CUDNN入れてもダメでした。) python -m torch.utils.collect_envで読み込みした現在の環境は下記の通りとなります。Trueとさせる方法がわからず、もう自分では完全に手詰まりとなっておりますので、ご教授頂けますと大変有難いです。 C:\Users\●●●>python -m torch.utils.collect_env Collecting environment information... PyTorch version: 1.13.0+cu116 Is debug build: False CUDA used to build PyTorch: 11.6 ROCM used to build PyTorch: N/A OS: Microsoft Windows 10 Pro GCC version: Could not collect Clang version: Could not collect CMake version: Could not collect Libc version: N/A Python version: 3.9.13 (tags/v3.9.13:6de2ca5, May 17 2022, 16:36:42) [MSC v.1929 64 bit (AMD64)] (64-bit runtime) Python platform: Windows-10-10.0.19045-SP0 Is CUDA available: False CUDA runtime version: 11.6.124 CUDA_MODULE_LOADING set to: N/A GPU models and configuration: GPU 0: NVIDIA GeForce RTX 2060 SUPER Nvidia driver version: 511.65 cuDNN version: Could not collect HIP runtime version: N/A MIOpen runtime version: N/A Is XNNPACK available: True Versions of relevant libraries: [pip3] numpy==1.23.5 [pip3] torch==1.13.0+cu116 [pip3] torchaudio==0.13.0+cu116 [pip3] torchvision==0.14.0+cu116 [conda] Could not collect pytorchはよく知らんけど toolkit入れた? cuda.is_availableてtorchのメッセージだよね? てな感じで全然情報不足だわな そうです。Torchのメッセージです。 CUDA Toolkitなら入れてますね。 現在は11.6.2をインストールしています。 以下コマンドプロンプトの出力結果。 ■Nvcc -V で表示される内容 nvcc: NVIDIA (R) Cuda compiler driver Copyright (c) 2005-2022 NVIDIA Corporation Built on Tue_Mar__8_18:36:24_Pacific_Standard_Time_2022 Cuda compilation tools, release 11.6, V11.6.124 Build cuda_11.6.r11.6/compiler.31057947_0 ■print(torch.__version__)で表示される内容 1.13.0+cu116 ■torch.cuda.is_available()で表示される内容 False 過去にcuda_11.3.1とか10.2をツールキットからインストールした事もありますが、それでもダメでした。 まずcudaは動いてるの? devicequeryとかsampleのプログラムが動くかどうか確認して問題なければ torch周りの問題かと ありがとうございます! サンプル試したことなかったのでやってみました。 https://github.com/NVIDIA/cuda-samples/releases 上記サイトからCUDA Samples v11.6をダウンロードして、VisualStudio2017.slnをビルド ビルド結果:ビルド: 175 正常終了、11 失敗、0 更新不要、0 スキップ(ただしビルド文書の中にいくつか失敗という文字あり) deviceQueryをコマンドプロンプトから実行したところ・・・。 cudaGetDeviceCount returned 100 -> no CUDA-capable device is detected Result = FAIL という結果だったので、CUDA対応デバイスが検出されていないようです・・・。 あっ、ビルド結果、11個失敗しているって意味なんですね・・・。 GPUがちゃんと刺さって無いとか、電源不足とかドライバがちゃんとインストールされてないとか。 ん〜、GPU自体は認識されてるみたいなんですよね。 GPU-Zで調べてみたらCUDAのところチェック外れてたので、やはりCUDAは認識されていないみたいです。 https://gpuz.techpowerup.com/22/11/27/e6p.png 現在の電源は550Wなのですが、もうちょっと大きい容量のに変えてみて、それでダメならやはりドライバが合っていないという事でしょうか・・・。 gpu-zはハードウェア情報なのでcudaのチェックが外れてるって 何かがかなりおかしい 20年前のgeforceなら分からんでもないがRTXなら 偽物をつかまされたか、ハードウェア故障を疑った方が良い気がする まずはsystem32にnvcuda.dll があるかどうか確認するかな 316です。 Sytem32にnvcuda.dll入ってました。 やはりハードウェアの故障でしょうか? (今のGPUは約2年前にドスパラで買いました) とりあえず電源ユニット注文しちゃったんで、電源交換&GPUも付け直してみて、それでダメだったら新しいGPU購入ですかね(T_T) グラボを3060TIに変えてみましたが、それでもGPU-ZでCUDAにチェックが入っていませんでした・・・。 接続するPCIEスロットも変えてみましたが・・・ダメでした。 グラボ上部にある補助電源?用のケーブルはちゃんと繋いでるはずなのですが、他に必要な接続とかあるんでしょうか? 現状のGPU-Zでの結果 https://gpuz.techpowerup.com/22/12/02/pn6.png OpenGL, DirectX, vulkanとかのグラフィッ関係のAPIを使ったプログラムやゲームも動かないの? 本当にハードウェア関係に問題があるんならそういうプログラムも動かないだろうけど。 他のGPU全部殺してドライバアンインストール 3060Tiのドライバ入れなおせ 316です。 思い切ってCドライブ初期化してみたらCUDA認識されました。 皆さんありがとうございました。 経験上、一番難しいと感じたのはpython listとかタプルとかデータ型が沢山あってこの変換がやたら発生 pandaとnumpyでも型が違うとか torchは経験ないけどtensorflowは理屈が分かればそれほど難しくないと思う >>338 普通にC++の方が難しいと思うけど、どんな言語を試した経験上なの? もしかしてCUDAでC++を覚えようとしてるの? CUDAはCの知識でもできるけどCは? C++は変態過ぎて 経験豊富ならべつにC++かまわんが 経験無いならCからやった方が良い >>341 Cは便利じゃないけど、確実だよね。マイコンとかでC++を使うのは、開発時間が十分に取れるなら(慣れてるなら)いいけど、処理系の変な制限にハマりそうで怖くて使えないわ。 >>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ちゃんねる