トップページ⇒tech
1001コメント314KB

【GPGPU】くだすれCUDAスレ pert4【NVIDIA】

■ このスレッドは過去ログ倉庫に格納されています
0001デフォルトの名無しさん2010/12/04(土) 21:57:13
このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
CUDA使いが優しくコメントを返しますが、
お礼はCUDAの布教と初心者の救済をお願いします。

CUDA・HomePage
http://www.nvidia.com/cuda

関連スレ
GPUで汎用コンピューティングを行うスレ
http://hibari.2ch.net/test/read.cgi/tech/1167989627/
GPGPU#5
http://hibari.2ch.net/test/read.cgi/tech/1281876470/

前スレ
【GPGPU】くだすれCUDAスレ【NVIDIA】
http://pc12.2ch.net/test/read.cgi/tech/1206152032/
【GPGPU】くだすれCUDAスレ pert2【NVIDIA】
http://pc12.2ch.net/test/read.cgi/tech/1254997777/
【GPGPU】くだすれCUDAスレ pert3【NVIDIA】
http://hibari.2ch.net/test/read.cgi/tech/1271587710/
0159デフォルトの名無しさん2010/12/18(土) 17:43:36
>>158
行列の計算ってのはね、分割できるんですよ。
0160デフォルトの名無しさん2010/12/18(土) 18:13:31
>>159
確かに。
分割できなければコアが複数あっても意味がない。
0161デフォルトの名無しさん2010/12/18(土) 18:55:39
でもね、いったんメインメモリに移して計算するのはもったいない。

結論C2070を買え。
01621302010/12/19(日) 18:30:10
結局、部分行列作成と計算結果の結合で、処理にかかる時間が逆転してしまいました。

一般論として言えるのかどうか知らないけど、
行列が（計算結果を含めて）ビデオカードのメモリにスッポリ入るなら劇的な高速化が見込めるけど、
そうでなければ、マルチコアCPU＋本体メモリのほうが速度でもメモリ消費量でも有利なようです。
01631282010/12/19(日) 20:32:38
コンスタントメモリはいつデバイスに確保されるのでしょうか
__const__宣言時とは考えにくいので最初にcudaMemcpyToSymbolしたときでしょうか

デバイスにあるということはコンスタントメモリも終了時に解放したほうがいいと思うのですが、
cudaFree()では解放できないようです

コンスタントメモリはどのように解放すればいいのでしょうか。
0164デフォルトの名無しさん2010/12/19(日) 21:03:47
行列計算なら転送を隠せるはず
リンパックベンチでGPU搭載マシーンが上位に来てるんだし
0165デフォルトの名無しさん2010/12/19(日) 23:49:30
>>163
SDKのサンプルみればいいんじゃないか？
0166デフォルトの名無しさん2010/12/20(月) 00:43:41
__constant__を気にするなら
kernelコードだって似たようなものなんじゃ。
0167デフォルトの名無しさん2010/12/20(月) 11:19:24
>>163
定数メモリはcudaMalloc()のように「確保する」ということなしに使うようになっています。
その為、cudaMemcpyToSymbol()のような特殊な関数でしか転送できません。
また、「確保する」ことなしに使える以上、「解放する」手続きは用意されていません。
0168デフォルトの名無しさん2010/12/21(火) 15:33:13
安定ソートが欲しいのですが
CUDA向きの安定ソートってありますか？
16000要素くらいです
0169デフォルトの名無しさん2010/12/21(火) 15:52:44
NVIDIAにあるバイトニック使ってみ
0170デフォルトの名無しさん2010/12/21(火) 15:54:21
SDKにmergesortってあるがそれが同じ奴だ
01711682010/12/21(火) 16:18:22
バイトニックは安定じゃないようなのですが、
マージソートのサンプルがSDKにあったのですね。
チェックしてみます
0172デフォルトの名無しさん2010/12/21(火) 20:58:05
gtx580を4基積んだ．電源が音を上げた．
やれやれ．
0173デフォルトの名無しさん2010/12/21(火) 21:12:11
>>172
何Wで駄目だった？
0174デフォルトの名無しさん2010/12/21(火) 21:46:21
>>173
1400W。ディスク系を外においてもなぁ。
0175デフォルトの名無しさん2010/12/21(火) 21:49:20
SC748あたりなら・・・
0176デフォルトの名無しさん2010/12/21(火) 21:59:47
7046GT-TRFね。
0177デフォルトの名無しさん2010/12/21(火) 22:17:39
1400Wで駄目だとなると電源2つ積めるケースにして連動させるケーブルで繋いでと
いったことをしないと無理そうだな。
0178デフォルトの名無しさん2010/12/21(火) 23:19:45
無茶しやがって…
0179デフォルトの名無しさん2010/12/22(水) 01:04:00
1400Wとか言って売っている時点で粗悪電源確定。
>>172が日本在住でないなら話は別だが。
0180デフォルトの名無しさん2010/12/22(水) 07:14:57
>>179
あれは110V/1400W。80GOLDのリタンダント。
日本じゃ実効1100W。
0181デフォルトの名無しさん2010/12/22(水) 22:34:50
大容量PC電源はエアコン引っこ抜いて200Vから取ろう。
0182デフォルトの名無しさん2010/12/22(水) 23:25:26
>>180
なら、「1100W電源」として売るべきだろ。

排気量2000ccの車を3000ccだと言って売る自動車屋などありえない。
そのあり得ないことを長年やってきたのが、PC電源代理業者だ。
0183デフォルトの名無しさん2010/12/23(木) 00:59:59
>>182
supermicroは日本向けのサイトを持っていない。仕方あるまい。
ちなみに法人で入れたときはこの説明はあった。
www.supermicro.com/products/system/4U/7046/SYS-7046GT-TRF.cfm
0184デフォルトの名無しさん2010/12/24(金) 20:02:02
GPUの分類について教えてください
例えばGTX480は
Fermi-ComputeCapability2.0-GF100-GTX480
でよいのでしょうか
それともアーキテクチャとCCは1対1対応しているのでしょうか
また、Fermiの分類位置はここでよいのでしょうか
0185デフォルトの名無しさん2010/12/24(金) 22:08:13
ComputeCapability2.0-GF100-GTX480
は正しいけど、FermiはC.C.2.X全部を含む
GFの"F"はFermi、GTXの"T"はなぜかTeslaだったはず
0186デフォルトの名無しさん2010/12/24(金) 22:20:21
てめぇGeForce 7800GTXさんに謝れよ
0187デフォルトの名無しさん2010/12/24(金) 23:41:29
>>184
CUDA目的ならComputeCapabilityだけ意識すればよいかと。
あとはパフォーマンスが違うだけだから。
0188デフォルトの名無しさん2010/12/24(金) 23:43:40
>>186
GeForce 7800GTXさんはCUDA対象外じゃ・・・・・・
0189デフォルトの名無しさん2010/12/25(土) 00:54:10
コンパイルされたCUDAプログラムを実行するだけの場合、ドライバのみでいけますか？
toolktのインストールは必須ですか？
LD_LIBRARY_PATH /usr/local/cuda/lib64
を通さなくてはならないということはtoolkit必須なのかな
0190デフォルトの名無しさん2010/12/25(土) 01:00:31
ライブラリ使っててもstaticならいらないはずだし、dllつかってるなら
そのdllをexeと同じフォルダにおけばいいんじゃないか？
きっとlib64内のなんかのライブラリつかってんでしょ。cutilとかかな
0191デフォルトの名無しさん2010/12/25(土) 02:39:52
DLL再配布して良かったっけ
0192デフォルトの名無しさん2010/12/25(土) 09:47:13
Linuxならドライバ（カーネルモジュールを含む）も必要なので再配布だけでは解決しない。
もちろんツールキットのライブラリも必要。
0193デフォルトの名無しさん2010/12/27(月) 19:25:09
すみません、教えてください。
CentOS 5.5(64bit)でCUDAを導入したくて
・Developer Drivers for Linux (260.19.26)
・CUDA Toolkit for RedHat Enterprise Linux 5.5 (64bit版)
をインストールし、とりあえずSDK code samplesをコンパイルして動くか確かめたのですが

コンパイル後のファイルを実行しようとすると
"cannot execute binary file"のメッセージが出てきてプログラムが実行できません。
この問題を解決したいのですがどこを確認してどう改善したらよいのでしょうか？

ちなみにカーネルのバージョンは2.6.9-55.0.2.ELsmpです。
0194デフォルトの名無しさん2010/12/27(月) 19:40:29
環境設定は？
0195デフォルトの名無しさん2010/12/27(月) 20:29:33
それはCentOS 4.5のようだが。
0196デフォルトの名無しさん2010/12/27(月) 22:11:09
>>193
例えば、ディストリビューションのnvidiaドライバなんかが残っていない？
その上に、最新ドライバをインストールしていて、
コンパイル時と実行時に別の共有ライブラリを参照していたりとか。

ていうか、自分も、ubuntuでx-updatesのドライバと最新ドライバが両方入っていて挙動不審。
ubuntuでもOS再インストールかｗ
01971932010/12/28(火) 00:14:59
>>196

返信ありがとうございます。
ちょっとその辺を調べてみます。
後日、問題が解決したらフィードバックします。
0198デフォルトの名無しさん2010/12/28(火) 07:48:23
スルーされた。勝手にしてくれ。
0199デフォルトの名無しさん2010/12/28(火) 16:15:18
いじけるなよｗ
0200デフォルトの名無しさん2010/12/28(火) 22:47:32
とある処理のスレッド数と時間を関係を調べているのですが、
スレッド数が192を超えたあたりで処理が行われていません。

スレッド数が512より小さくても処理が行われないことがあるのでしょうか？

グラボは9800GTで11レジスタ、共有メモリを36byte使用してます。
0201デフォルトの名無しさん2010/12/28(火) 22:48:39
>>200
「処理が行なわれない」ってどういう意味さ。
0202デフォルトの名無しさん2010/12/28(火) 23:01:14
>>201
説明不足でした。

処理を1000回行い、平均処理時間を計測しているのですが、
スレッド数が192を超えると0(ms)になってしまいます。
0203デフォルトの名無しさん2010/12/28(火) 23:19:14
何かが間違っているんだろうね。
0204デフォルトの名無しさん2010/12/29(水) 00:34:35
まずはエラー確認してからだろ。
0205デフォルトの名無しさん2010/12/29(水) 00:38:37
スレッド数を増やしたことによって確保されていない領域にアクセスするような
エラーが発生するような状況になっていないかどうかを確認・・・・・かな。
これが起きているとそれ以降の（タイマーを含めた）全ての処理がエラーになる。
0206デフォルトの名無しさん2010/12/29(水) 01:31:39
まずは
cudaThreadSynchronize();
cudaGetLastError();
0207デフォルトの名無しさん2010/12/29(水) 14:09:27
>>203-206
助言ありがとうございます。

まずはエラーを確認してみたいと思います。
0208デフォルトの名無しさん2010/12/29(水) 19:33:12
たぶんリソース不足でしょ。自分もよくそうなる。
カーネルの引数を構造体のポインタにまとめると解決する。
0209デフォルトの名無しさん2010/12/30(木) 20:27:57
GTX460ってCUDAは使えるんですかね？
lspci | grep -i nVidia
でPCIバスデバイス情報確認したら
VGA compatible controller: nVidia Corporation: Unknown device …
ってコンピューターさんがおっしゃられたんだが…
0210デフォルトの名無しさん2010/12/31(金) 00:21:26
>>209
/sbin/update-pciids
してからコンピュータさんにもう一度お伺いたててご覧なさい
0211デフォルトの名無しさん2010/12/31(金) 05:27:05
ドライバをインストールしてないんじゃないの？
0212デフォルトの名無しさん2011/01/01(土) 11:59:19
正月セールで狙う1万5000円くらいのグラボでオススメのものをおせーて
0213デフォルトの名無しさん2011/01/04(火) 09:25:33
GeforceとRADEONを2枚挿して
GeforceをCUDA専用、RADEONをグラフィック専用とすることは出来ますか？
0214デフォルトの名無しさん2011/01/04(火) 09:34:04
日本語的に複数の意味に取れるので訂正

GeforceとRADEONの2枚を挿して
GeforceをCUDA専用、RADEONをグラフィック専用とすることは出来ますか？
0215デフォルトの名無しさん2011/01/04(火) 09:41:44
no problem
0216デフォルトの名無しさん2011/01/04(火) 10:36:53
ありがとうございます
0217デフォルトの名無しさん2011/01/04(火) 10:52:31

cudaSetDeviceFlags(cudaDeviceBlockingSync);
0218デフォルトの名無しさん2011/01/04(火) 10:53:37
ほとんど同じソースで
cudaSetDeviceFlags(cudaDeviceBlockingSync);
が有効になるプログラムとならないプログラムがあるのですがどこを
調べれば良いでしょうか。
0219デフォルトの名無しさん2011/01/04(火) 12:31:38
ほとんど同じソースならどこが違うのかを書いてください。
そして「有効」かどうかをCPU負荷で判定しているのかどうかも。

cudaSetDevice()を除くすべてのCUDA関数より先にcudaSetDeviceFlags()を
呼び出しているかどうか以外には特に調べるべきところはありません。

0220デフォルトの名無しさん2011/01/05(水) 18:57:39
マニュアル的なpdfが沢山あってどれを読めばいいのかわからないのですがどれから読むのがいいのでしょうか？
CUDA_C_Best_Practices_Guide.pdf
CUDA_C_Programming_Guide.pdf
Introductory CUDA Technical Training Courses VolumeI.pdf
0221デフォルトの名無しさん2011/01/05(水) 20:31:24
Introductory CUDA Technical Training Courses VolumeI.pdfで軽く概要をつかむ。
ただこれ内容が古かったと思うからどんな仕組みで動いてるかを把握するくらいでいいと思う
必要になったらまた読み返せばいい。

次にCUDA_C_Programming_Guide.pdfで基本を学ぶ。
これは読み返すことが多くなると思う。

その次にCUDA_C_Best_Practices_Guide.pdfで良い実装の仕方を学ぶ。
これは実際にプログラム組みながらでいいと思う。

あと余談だがfermi　tuning guideも読むといい。
英語が苦手なら「はじめてのCUDAプログラミング」や「CUDAプログラミング実践講座 - 超並列プロセッサにおけるプログラミング手法」
を読んでから必要なところをPDFで追加で読むと理解が早いと思う。
0222デフォルトの名無しさん2011/01/05(水) 21:02:09
フランス語のはないですかね？
0223デフォルトの名無しさん2011/01/05(水) 21:15:20
ワロタ
0224デフォルトの名無しさん2011/01/05(水) 21:16:37
>>222
残念ながらありません
0225デフォルトの名無しさん2011/01/06(木) 00:39:25
cc2.0でコンパイルすると、cc1.3でコンパイルした時よりレジスタ使用量が
10くらい増えるのですが、なぜですか？
0226デフォルトの名無しさん2011/01/06(木) 22:50:40
なんでだろうねー
0227デフォルトの名無しさん2011/01/08(土) 18:48:55
Fermiで各スレッドで数百バイトのワーキングエリアが必要な場合、
ローカルメモリを使うのが最良なのでしょうか？
シェアードメモリを使うと、同時に立ち上げられるスレッド数が減って
逆に遅くなりました。
0228デフォルトの名無しさん2011/01/09(日) 21:38:49
>>225
そういうのはプログラムによって変わるだろ
とりあえずptxファイル嫁

>>227
>各スレッドで数百バイトのワーキングエリアが必要な場合
これが今以上減らせないならそもそもCUDAに向いてない
0229デフォルトの名無しさん2011/01/10(月) 17:57:08
グローバルメモリのアクセスは400~500clock、レジスタやシェアードメモリ、L1キャッシュはほぼ0らしいですが
L2キャッシュは大体何clockかかりますか？
0230デフォルトの名無しさん2011/01/10(月) 19:32:12
L1は大体15
L2は大体150
0231デフォルトの名無しさん2011/01/11(火) 09:01:34
4枚発注ー。
0232デフォルトの名無しさん2011/01/13(木) 11:27:07
CPU向けコードとGPU向けコードを同一ファイルで作成したいと思っています。
nvccはcuでないと、cudaコードとしてコンパイルしてくれないのですが、
拡張子がcppのファイルをcudaソースファイルとしてコンパイルするオプションはありませんか?
0233デフォルトの名無しさん2011/01/13(木) 19:39:51
obj同士はnvccでリンクできるから、それぞれコンパイルすればＯＫだよ
0234デフォルトの名無しさん2011/01/13(木) 20:42:06
>>141
亀レスですが，Driver APIを用いて出来ました．
ありがとうございました．
0235デフォルトの名無しさん2011/01/14(金) 20:47:45
>>234
概略をのせてくれまいか？

cudaAAA()で○○を設定
cudaBBB()でXを取得して
pthread_AAA()を使ってXXXする

くらいの事でもいいから
02362342011/01/15(土) 13:34:59
>>235
簡単にで良いなら．

// 前処理
#pragma omp parallel
{
　　cuInit(); // 初期化
　　cuDeviceGet(); // GPUへのハンドルを取得
　　cuCtxCreate(); // コンテキストを作成し，ホストスレッドにバインド

　　mallocとか色々

　　cuCtxPopCurrent(); // コンテキストをアンバインド
}

glutMainLoop();

// カーネル実行
#pragma omp parallel
{
　　cuCtxPushCurrent(); // バインド

　　カーネル実行

　　cuCtxPopCurrent(); // アンバインド
}

関数の引数は省略しました．
0237デフォルトの名無しさん2011/01/15(土) 17:57:20
>>236
ありがとう。参考にさせてもらいます
0238デフォルトの名無しさん2011/01/15(土) 21:30:19
Fermiもグローバルメモリアクセスはコアレッシングにしたほうがよくって
キャッシュはおまけ程度って考えでおk？
ていうか各スレッドが違うグローバルメモリにアクセスするなら同じカーネル内ではキャッシュ効かないよね
0239デフォルトの名無しさん2011/01/15(土) 22:14:42
CPUのマルチコア（巨大キャッシュ、低並列）のマルチスレッドですら
キャッシュ汚染はひどいからな

マルチスレッドとキャッシュは相性が非常に悪い
0240デフォルトの名無しさん2011/01/15(土) 23:23:54
スレッド番号とグロバールメモリーのアドレスが並んでいないと使い物にならないな。
速度が１０倍は違う。本当に使いにくいよ。
メモリアクセスがネックの（単純な）アプリで、shared memoryの使いまわしのできないアプリでは、
マルチスレッドを上回るのは難しい。
0241デフォルトの名無しさん2011/01/15(土) 23:49:58
しかし、いままでGPUと相性の悪いといわれるプログラムでも
FermiならCPUの2倍くらい出るってことは結構ある。

非Fermiならローカルメモリ使ったら即使い物にならない速度だったけど、
Fermiならキャッシュがなんとかしてくれたり。
0242デフォルトの名無しさん2011/01/16(日) 18:58:37
cublasDgemmにおいて、係数のalphaを(double)1に、
betaを(double)0に指定しても、
内部では真面目に係数をかけているのでしょうか？

参考(PDF注意)
http://developer.download.nvidia.com/compute/cuda/3_2_prod/toolkit/docs/CUBLAS_Library.pdf
0243デフォルトの名無しさん2011/01/16(日) 19:01:42
FORTRANのソースを自分で追っかけて確認しろ
0244デフォルトの名無しさん2011/01/16(日) 19:13:17
速度を測ればすぐに分かりそう
0245デフォルトの名無しさん2011/01/16(日) 21:42:18
>>243
ありがとうございます。
FORTRANのコードを見る限り、ちゃんと係数の値によって
処理変えているようでした。

>>244
ありがとうございます。
実際に測定したのですが、係数の値を変えても
計算時間に明確な差は現れませんでした。

GTX480 1枚, CUDA 3.2, Windows 7 64bit
A=7500x8000の行列,B=8000x7500,C=7500x7500の行列で
C=alpha * A*B + beta * Cを計算

alpha = 1, beta = 0を15回で85.079秒
alpha = 10, beta = 10を15回で85.072秒

FORTRANのコードのような実装になっていない気がするのですが。
0246デフォルトの名無しさん2011/01/16(日) 21:45:25
「はじめてのCUDAプログラミング」に出てるコード（もしくは同じ出版社からインターネット上に出てる奴）
をひたすらマイナーチェンジしていくのが一番早いと思う。
スクラッチはマイナーチェンジの経験を積んでからで十分。
0247デフォルトの名無しさん2011/01/17(月) 05:28:45
PTXって機械語なの？
GPUがPTX読むみたいな事が書いてあるんだけど
0248デフォルトの名無しさん2011/01/17(月) 06:02:22
ptxはgpuのアセンブラ言語じゃないの？？
0249デフォルトの名無しさん2011/01/17(月) 07:27:52
PTXはあくまで中間言語。
0250デフォルトの名無しさん2011/01/17(月) 10:47:31
TMPGEncで使おうと思ったらG80非対応なんだな。。
0251デフォルトの名無しさん2011/01/17(月) 20:15:09
GTX580を4枚積んでサンプルコード走らせたら電力不足来たー。
0252デフォルトの名無しさん2011/01/19(水) 09:33:50
CUDAツールキットを導入する際にまずやることって何？
どう頑張ってもnvccでコンパイルしたサンプルファイルが実行できませーん
って開発以前の問題で精神折れそうだわ…
0253デフォルトの名無しさん2011/01/19(水) 13:17:05
list.cは以下の通りです．
#include <stdio.h>
#include <stdlib.h>
#include <string.h>

struct record { /* １人分のデータ */
char name[32]; /* 氏名 */
int score; /* 得点 */
struct record *next; /* 次のデータへのポインタ */
};

void printlist( struct record *list ) {
if ( list != NULL ) { /* リストが空でなければ */
printf( "%16s:%4d\n", list->name, list->score );
printlist( list->next ); /* 先頭の次からリストの表示 */
}
}

0254デフォルトの名無しさん2011/01/19(水) 13:59:28
>>252
実行できないならできないなりにエラーメッセージが帰ってくるのではないでしょうか


CUDAコアとはスカラプロセッサのことでいいのでしょうか
FermiのみCUDAコアという呼び方をするのですか？
0255デフォルトの名無しさん2011/01/19(水) 14:09:09
>>252
実行できないならできないなりにエラーメッセージが帰ってくるのではないでしょうか


CUDAコアとはスカラプロセッサのことでいいのでしょうか
FermiのみCUDAコアという呼び方をするのですか？
0256デフォルトの名無しさん2011/01/19(水) 21:03:35
Fermiで倍精度浮動小数点の三角関数計算に
SFUを使えないのでしょうか

倍精度浮動小数点の計算は単精度浮動小数点演算機が2クロックで計算するのですよね
同じようにSFUが単精度計算分の2週でやってくれるということはないのでしょうか
0257デフォルトの名無しさん2011/01/19(水) 21:10:31
geforceは駄目だろ
0258デフォルトの名無しさん2011/01/19(水) 21:57:03
ドライバに　266.58 WHQL　が来たね
■ このスレッドは過去ログ倉庫に格納されています