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

1 ：11/08/23 〜 最終レス ：12/05/25

このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
CUDA使いが優しくコメントを返しますが、
お礼はCUDAの布教と初心者の救済をお願いします。
CUDA・HomePage
ttp://developer.nvidia.com/category/zone/cuda-zone
関連スレ
GPGPU#5
ttp://hibari.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/

2 ：

関連サイト
CUDA
http://www.nvidia.co.jp/object/cuda_home_new_jp.html
CUDAに触れてみる
http://chihara.naist.jp/people/STAFF/imura/computer/OpenGL/cuda1/disp_content
cudaさわってみた
http://gpgpu.jp/article/61432191.html
CUDA のインストール
http://blog.goo.ne.jp/sdpaninf/e/9533f75438b670a174af345f4a33bd51
NVIDIAの「GeForce 8800 GT(G92)」と次に控える64-bit GPUアーキテクチャ
http://pc.watch.impress.co.jp/docs/2007/1031/kaigai398.htm
CUDAを使う
http://tech.ckme.co.jp/cuda.shtml
NVIDIA CUDAを弄ってみた その2
http://dvd-r.sblo.jp/article/10422960.html
CUDAベンチ
http://wataco.air-nifty.com/syacho/2008/02/cuda_2044.html
KNOPPIX for CUDA
http://www.yasuoka.mech.keio.ac.jp/cuda/

3 ：

宣伝になっちゃうと難なので、一応分けて。いろいろ載っているサイト。
ttp://www.softek.co.jp/SPG/Pgi/TIPS/para_guide.html

4 ：

おい、pertだろ
まちがえんな

5 ：

>>2
個人blogのは情報古すぎなような…
「cudaさわってみた」はリンク切れだし。

6 ：

>>4
やっぱりCUDAやってる人は荒んでるんだな

7 ：

ほんとCUDAらないことで騒ぐよな

8 ：

ケンカするのはやめてCUDAさい＞＜

9 ：

そうだよ。喧嘩はやめ…GeForんｹﾞﾌｫﾝ

10 ：

まだスレが始まっTeslaいないのに

11 ：

ときに口論になるのもみんな自分の技術にそれだKeplerイドを持ってるってことだ。

12 ：

amazon.comで検索して出てくるCUDA本ってことごとく出版がずいぶん先か延期になったりするな
なんでだろ

13 ：

無効とは情報の登録ポリシーが違う…ような気がする。
それがAmazonのポリシーなのか出版社のポリシーなのかは知らないけれど。
多分出版社のサイト直接見た方が正確。

14 ：

cudaの本が必要、って訳でも無いが出版されるなら買ってしまうと思う

15 ：

4.0になってdeviceemuが使えんのだが、CUDA載せてないPCで動かす方法ない？

16 ：

3.2を使う。

17 ：

>>16 アホしかおらんのか

18 ：

>>17
現実的じゃないか。どうせemuで正確なエミュレーションができるわけじゃないし。

19 ：

>>17
4.0でdeviceemu使う方法なんてねーよ、ってことだよ。
言わせんな恥ずかしい

20 ：

>>15
Cuda x86

21 ：

理屈の上ではKnights Ferryでも動かせるのかな？

22 ：

ubuntu 11.04 tesla c2050 cuda sdk 4.0の環境でdoubleのアトミック演算が正しく動作しません
オプションは -arch=compute_20を指定してます。
本家のプログラミングガイドに書いてある、以下のコードを実行してます。
__device__ double atomicAdd(double* address, double val)
{
unsigned long long int* address_as_ull
　　　　　　　　　　　　　　　　　　= (unsigned long long int*)address;
unsigned long long int old = *address_as_ull, assumed;
do {
assumed = old;
old = atomicCAS(address_as_ull, assumed,
　　　　　　　　　　　　　__double_as_longlong(val + __longlong_as_double(assumed)));
} while (assumed != old);
return __longlong_as_double(old);
}
引数のadrees*には共有メモリのアドレスを渡してます。
一つのスレッドのみが共有メモリにアクセスする場合は正しく加算は行われますが、二つ以上のスレッドが衝突する場合は
正しく動作しません。printfで見てみたら、二つ目のスレッドがループを抜けれずに何度も共有メモリの値に
加算しているようでした。上の文をdouble→float,long long →　int　に書き換えた場合は正しく動作しました。
また、atomicCASのところを二つのprintf、
printf("%d %f %f\n",threadIdx.x,__longlong_as_double(assumed) ,__longlong_as_double(old));
ではさむとなぜか正しくループを抜け、正しくsharedメモリに値を書き込みます。
原因がわかる方、教えていただけないでしょうか。

23 ：

>>22の続きです
assumed oldの値をcpu側に持ってきて表示してみたら、
oldの値(atomicCASの返す値)とassumedの値が末尾の数ビットが
異なってました。50回ほどループしてやっとビットレベルで値が一致して
ループを抜けているようでした。その間に過剰に加算が行われていたようです。

24 ：

カードの計算が正しいかチェックするベンチマークってありませんか？
四則演算と√と三角関数とかを全ビットを演算器全部チェックしろと云う無茶ぶりがきた。
計算機が計算を間違うなんてありえないのに…

25 ：

Teslaの存在意義全否定デスカ

26 ：

>>24
過去スレを読むとわかってもらえると思うが、みんなそれを求めているんだ。
そのチェックにおける最先端はおそらく長崎大学。

27 ：

1スレッドの実行がどのコアかがわかる仕組みがどのような方法か分かればいいわけだよな
誰か調べてみろよ

28 ：

>>22,23
もしかしてブロック辺りのスレッド数が32以下なら動いたりする？

29 ：

>>27
何がしたいのか分からんがそれは簡単だ。

30 ：

>>29
コアと計算結果の表が出来上がればコア間の相違が分かるじゃん
やり方知ってるなら教えてください

31 ：

煽っても教えてやらん
29じゃないけどｗ

32 ：

NvidiaのCUDAをやろうと思い、Nvidiaのカードの刺さったマシンにFedoraを
入れてCUDAのためのNvidiaのドライバを入れて使ってみたが、OSのカーネル
のアップデートがあるたびに、コンソールが再起動時にブラックアウトした。
そのたびごとに、別のマシンからリモートログインをしてCUDAのための
Nvidiaのドライバを再インストールする必要があった。
CUDA3.xまではそれでなんとかなったが、現在のCUDA4.0になると、
幾らやっても駄目になってしまった。どうやらFedora15ではまともに
サポートされていないようなのだ。Nvidiaの提供しているCUDAのための
ドライバがインストール時におかしなメッセージが出てインストールが
出来ない。

33 ：

ここはおまえの日記帳だから好きに使いなさい

34 ：

>>28
ブロックあたりのスレッド数を32以下にして試してみましたが、
結果は変わらなかったです。

35 ：

>>23
とりあえず、プログラミングガイドのコードそのままだし、
実際アルゴリズム自体には問題ないから、
意図どおり動かないとしたら、ドライバかSDKのバグか
ハードの故障じゃね。
まあ、ここで詳細なバージョンとか晒してみたり、
PTXやアセンブリを貼ってみるのもいいけど、
本家のフォーラムで聞いてみたほうが早いんじゃね。

36 ：

長崎大の言う演算間違えって、具体的にどういう演算を間違えたのかの情報ってないの？
間違えますって情報だけ垂れ流して具体例出さんって、どういうことなんだべ。

37 ：

×演算間違え
○演算間違い

38 ：

CUDAを使用するプログラムで巨大なデータを扱いたいのですが、nvcc でのコンパイル時に
relocation truncated to fit: R_X86_64_32S against `.bss'
とエラーが出て完了しません。gcc ではデータ量制限を解除するのに-mcmodel=medium というオプションがあるようですが、nvcc でこれにあたるものはないでしょうか。

39 ：

巨大データってどれくらい？
CUDAは基本グローバルメモリ使用の制約があるから”巨大”は決して得意ではないのだけど

40 ：

>>39
今確認したところ、要素が数十万個の配列を使用していて、それらの合計は200MBほどです。
そのうち、実際にGPUに計算させるためGPUに渡すデータは50MB程度です。
上記のエラーメッセージでググったところ、データ量制限が問題だろうということだったのですが…それほど大きなデータでもないですよね。
なんなんでしょうこのエラー。
プログラミングに関してはまだ知識が浅いためvンカンプンなこと言ってたらすみません。

41 ：

どうせデータをソース中に書き込んでるんだろw

42 ：

>37
日本語の間違えより、長崎大の言ってる間違いの詳しい情報くれよぅ…。

43 ：

要素数が大きすぎるなら、mallocを使ってみたら？

44 ：

>40
配列は何次元？
CやC++だと配列は連続したメモリ領域と規定されてるので、
確保可能なメモリ量は物理メモリの絶対量ではなく、
アドレスを連続で確保できるメモリ量に規定されるよ。
例えば大きめの２次元配列を確保する場合は
double* pBuf = new double[row * column];
とするより
double** ppBuf = new (double*)[row];
for(int i = 0; i< row; ++i){
　　　ppBuf[i] = new double[column];
}
とかにしたほうがいいよ。

45 ：

ファイル読み込みで、cudaMallocHost(だったかな？？)でメモリ確保すればいいじゃん

46 ：

>>43>>44
gccと同様にオプションでエラー回避できないかな…と思ったのですが、配列の取り方を変えるほうが良さそうですね。
特に>>44は知りませんでした。
ありがとうございました！

47 ：

良く分からんが、-Xcompiler -mcmodel=mediumじゃ駄目なん？
-Xlinkerも要るかもしれないが。

48 ：

>>44
普通はありえない。そんなことしたら、ポインタアクセスが一回余計に要るから遅くなるし、
キャッシュアウトしやすくなるから更に遅くなる。
まぁ、GPUに転送するのは1スライス分だけだろうからGPUに任せてCPUは暇こいていられるならいいかも知れんが。

49 ：

>>44
それは無い。
・メモリも無駄に消費
・アクセス性能がた落ち
・バグの温床になる
つーかC++ならvectorで初期化時に
サイズ指定したほうがいろいろ便利

50 ：

配列でやる場合、
配列の添え字がINT_MAXを越える大きさならそうするしかないと思うが・・

51 ：

キャッシャがアクセスがとドヤ顔されても・・・。
要素のデカイ配列が確保できんって話なわけだし。

52 ：

コンパイル時の話なのに
ヒープのフラグメンテーションの訳無いだろう。
nvccからgccに自由にオプション渡すやり方は
知らんが、最悪c(++)部分のコードを分割して
gccでコンパイルすりゃ良い。
仮にヒープのフラグメンテーションが酷い場合でも
>>44みたいなやり方は下策だよ。
素人におかしなやり方を覚えさせるだけ。
大きな配列確保する前に小さな物についてShrink-to-fitやればデフラグされるから
その後巨大配列を確保すれば良い。
特にGPGPUで巨大な配列扱うような用途なら
配列の配列なんて転送も困るだろうに。

53 ：

>大きな配列確保する前に小さな物についてShrink-to-fitやればデフラグされるから
>その後巨大配列を確保すれば良い。
コレ具体的にどんなコード書けばいいのん？

54 ：

メインの計算に必要な必要な小容量配列
保持用のポインタなりvectorの変数だけ始めに
ヒープに確保した後、メインの計算前の前処理をする。
この処理で、一時的なデータと最終的に使うデータの
ヒープ確保が入り乱れる所為でフラグメンテーションが発生する。
最終的に使うデータについて、大体処理した順番を守って
ひたすら同容量確保してはコピーを繰り返せば、
こういうデータは全て前方に圧縮されて配置されるようになる。

55 ：

？
まあスレ違いなんで他でやってよ。

56 ：

kernelへの引数に構造体は使えんの？
今やってみたらコンパイルでエラーになった

57 ：

あ、間違えてただ
構造体引数使えるべな

58 ：

>>15
http://code.google.com/p/gpuocelot/
Ocelot currently allows CUDA programs to be executed on NVIDIA GPUs, AMD GPUs, and x86-CPUs at full speed without recompilation.

59 ：

cuda4.0、gtx580で
Maximum dimension of a grid( 65535 , 65535 , 65535 )
ってなってるんですがz軸のブロックの取れる数は1じゃなくなったんですか？

60 ：

デバイス上の処理が長時間になる場合、途中経過を取得する方法ってないかな？
例えば１００万回のモンテカルロを実行して、途中の１０万回、２０万回が済んだらその情報を取りたいんだけど…

61 ：

>>60
streamingでできないかな?

62 ：

streamingってなに？
SMのSじゃないよね？

63 ：

カーネル呼び出し時の<<<>>>の第3引数がstreamingの何かじゃなかったっけ？
俺は>>61を見てそういうのがあったなと思い出したレベルなので使い方は知らないけど。

64 ：

第3パラメータはsharedMemoryのサイズだよ。そこで指定しなくても使えるけど。

65 ：

ごめん第4引数だったみたい。それとAsync系関数がうんたらかんたら。

66 ：

あー、その辺だね。

67 ：

【トリップ検索】CUDA SHA-1 Tripper【GeForce】
http://hibari.2ch.net/test/read.cgi/software/1311428038/
すげえな

68 ：

個人でCUDAプログラミングやってうれしいのはその程度なのか？
と逆に不安になるんだが…

69 ：

世の中なんでも膨大なデータをいかに早く処理するかが勝負の世界に
なってきているから、CUDAプログラミングの旨味はいくらでもあるだろ。

70 ：

いつかやってみようと思って半年くらいこのスレROMってるが
いまだに開発環境のセッティングの仕方すら調べていない・・・

71 ：

>>68
>>67くらいのコード書ける？

72 ：

>>71
おれがそれくらいのコード書けるようになるころには、自ずとましな使い道が浮かんでくることだろう、
っていうこと？

73 ：

>>60
試してないけど、長い処理の途中でグローバルメモリに書き込んで、
別streamでその値を読むカーネルを起動すれば取れると思う。
キャッシュからグローバルメモリに書き込むのは特別な関数を
使うと思ったけど名前忘れた

74 ：

>>73
思い出した。atomicを使うか、書き込む側で __threadfence()

75 ：

失礼します、先輩方お助けを。
atomicを使ったソースを"-arch compute_20"オプションでコンパイルすると次のようなエラーが出てしまいます。
「'-arch' は、内部コマンドまたは外部コマンド、操作可能なプログラムまたはバッチ ファイルとして認識されていません。
C:\Program Files (x86)\MSBuild\Microsoft.Cpp\v4.0\Microsoft.CppCommon.targets(103,5): error MSB3073: コマンド "-arch compute_20
error MSB3073: :VCEnd" はコード 9009 で終了しました。」
-archコマンドはプロパティからビルド前イベントのコマンドラインに追加してあります。
-archコマンドがあるディレクトリへパスが通ってないのかなと思うのですが、このようなコマンドがあるディレクトリの追加の方法がよくわかりません。
環境はcuda4.0 + VC2010 + teslac2070(cc2.0) + windows7 proです。
どなたかアドバイスください。

76 ：

-archコマンドってなんだそら？
compute capabilityはnvccなりcl.exeのオプションに渡すものであって、
-archコマンドなんてものはないよ。

77 ：

自分でオプションでって書いてるけど、メッセージはそれで実行した系だよな

78 ：

>>76
>>77
ありがとうございます。「cuda by example」にもatomic関数を使ったソースには-archコマンドラインオプションを指定する必要があります。
とあるのですが、みなさんはatomicを使ったソースをコンパイルするときなにかコマンドつけませんか＞＜？

79 ：

>>78
あ、もちろん全部64bitです！

80 ：

CUDA以前の話だな。コンパイラオプションの意味がわかってない。
というか、VCなら、カスタムビルド規則でCUDA Runtime API Build Ruleを追加したほうが早い

81 ：

-gencode=arch=compute_20,code=\"sm_21,compute_20\"
になってるなぁ俺の環境だと。VC2010ならプロジェクトのプロパティでCUDA関係の設定するところがあるんじゃない？

82 ：

>>80
ありがとうございます。
えーと、それはおそらくVC2008の話ではないかと思います。2010にカスタムビルド規則という名前はなくなり、その代りCUDA4.0というビルドカスタマイズに統一されています。

83 ：

>>81
ご返答ありがとうございます。
そのコードはどこに書いていますか？
はい、cuda c/c++という項目にいろいろいじるところはあるのですがコンパイラオプションを設定する項目はないように思えます。
私がオプションを追加しているのはビルドイベントという項目ですね。
以前にもatomic関連の話が出ていたようです。>>22

84 ：

-arch=compute_20
にすればいいんじゃねえの？

85 ：

>>84
ありがとうございます。
おっしゃるように>>75のように追加するとエラーを吐きます。
追加の仕方が違うのでしょうか。

86 ：

VS2010は色々変わってるんだな。
知らんかったすまん。
とりあえず「CUDA VS2010」でググって出てくる一番上のサイト参考にして、CUDA4.0のビルドカスタマイズファイルを導入した後、プロパティからCode Generationをcompute_20,sm_20にしたらいけた

87 ：

GPUのリソースを全部CUDAに割り当てるため，グラボを2枚さしてディスプレイ出力用とCUDA用に分けるというのはありますが
同じようなことをsandy bridgeとグラボ1枚でも出来るのでしょうか？

88 ：

>>87
Z68マザーなら、マザー側DVIにディスプレイつないでおけばいい話じゃ？
CUDAで計算したデータをintelで表示したいなら、
1.DirectXかOpenGLのリソースをCUDAにバインドしてデータ書き込み
2.メインメモリに転送
3.intel側に転送
って形になるかなぁ。多分。

89 ：

>>83
ゴメン、俺もVC2010じゃないや。VC2008。
ちなみに>>81は構成プロパティの中に「CUDA Runtime API」があってそこの「コマンドライン」にある。
CUDA Runtime APIの中にGUIでGPU Architecture（sm_21）を指定するところがあるから、
コマンドを打ってるわけじゃないんだけど、結果としてnvccには>>81のようなオプションが渡ってる。

90 ：

そもそもなぜCUDA Runtime APIという項目があるのかというとサンプルのプロジェクトをコピったから。

91 ：

>>80
>>81
ありがとうございます！code generationを>>80のように書きなおしたらオプション指定することなくコンパイル通りました。
どうやらVS2010からはcommand line指定しなくてもcode generationでオプションを渡せるようです。
本当に助かりました。ありがとうございます！

92 ：

最近巷で流行ってる
error MSB3721: コマンド ""C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\bin\nvcc.exe" -gencode=arch=compute_20,code=\"sm_20,compute_20\" --use-. ... template.cu"" はコード 2 で終了しました。
のエラーなんですけど解決法どなたか知りませんか??
CUDA FORUMにもかなり話題が上がってるっぽいんですが...私のところにも発病してしまいました。

93 ：

そこだけ書かれても分からないと思う。
nvccのエラーメッセージ載せないと。

94 ：

>>92
敵性語なんでよくわかんなかったけど、
VisualStudio を管理者として実行したら動く？

95 ：

Visual C studioは下手にGUIに頼るよりも、コマンドラインでやった方がはやいこともある。コマンドラインでnvccやってみよう。

96 ：

単精度のmadってmulより遅いやないかー
プログラミングガイド嘘つきやないかー
ちくしょー

97 ：

>>96
そりゃ命令実行時間は同じだけど、
必要な帯域は違う。

98 ：

>>92
コマンドラインのオプションが何か1つ間違ってた気がする。
コマンドプロンプトで同じコマンド実行するとエラーが出たので、
そのエラーのひとつ前の引数が問題だった、はず。

99 ：

CUDAとRTミドルウェアを組み合わせてみた。
行った方法は、CUDAのサンプルをRTコンポーネントに
パッケージングする方法。
動作周期にあわせて、行列計算が実行されたのを確認。

100read 1read

1read 100read

TOP カテ一覧 スレ一覧 2ch元 削除依頼 ▲

【至急】助けてください。 (237)
制御系なら俺に聞いてもいいぜ（１０） (489)
C++相談室 part95 (482)
【消しゴム】MONOを使ってみるスレ4【じゃない】 (150)
【バグ管理】 BTS使ってる？【追跡】 3 (900)
Win32API質問箱 Build107 (934)
--log9.info------------------
街の自転車屋専用スレ　75店目 (503)
【ATB/CTB】MTBルック総合スレ 52【NO! OFF ROAD】 (187)
スリックＭＴＢ愛好家のスレ52 (638)
荒川ｻｲｸﾘﾝｸﾞﾛｰﾄﾞ 河口から267km (173)
荒川サイクリングロード 河口から267km (318)
岡山のショップ・コース情報スレ12 (696)
【実況】おい、今テレビでやってるぞ！ 45【禁止】 (310)
4万以下のクロスバイク46台目 (881)
【2万〜5万円】 格安折りたたみ＆小径車スレ 【02】 (304)
5LINKS (839)
パーツやる/売る/交換/欲しい part41 (799)
　　　　　　　今日の　ﾌ　ﾞﾁ　ｷ　ﾚ (242)
【100】百円ショップの自転車用品を語るスレ31 (208)
自転車の「あるあるｗｗｗｗ」★2 (809)
[ボトルゲージに]高知の自転車乗りPart3[リープル] (130)
【待てど暮らせど】サーベロCervelo19届かない】 (227)
--log55.com------------------
質問デス
蚊に刺されたらageるスレ
樹海で白骨死体とコンニチワしたらageるスレ
クマゼミは害虫です
◆◆材飼育◆◆
クワガタ、カブトの値段議論
【神戸】とにかくクワガタ＆昆虫が好きな人集まれ！
熊本クワガタ・カブト採集ポイント