もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
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/
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/
ttp://www.softek.co.jp/SPG/Pgi/TIPS/para_guide.html
まちがえんな
個人blogのは情報古すぎなような…
「cudaさわってみた」はリンク切れだし。
やっぱりCUDAやってる人は荒んでるんだな
なんでだろ
それがAmazonのポリシーなのか出版社のポリシーなのかは知らないけれど。
多分出版社のサイト直接見た方が正確。
現実的じゃないか。どうせemuで正確なエミュレーションができるわけじゃないし。
4.0でdeviceemu使う方法なんてねーよ、ってことだよ。
言わせんな恥ずかしい
Cuda x86
オプションは -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メモリに値を書き込みます。
原因がわかる方、教えていただけないでしょうか。
assumed oldの値をcpu側に持ってきて表示してみたら、
oldの値(atomicCASの返す値)とassumedの値が末尾の数ビットが
異なってました。50回ほどループしてやっとビットレベルで値が一致して
ループを抜けているようでした。その間に過剰に加算が行われていたようです。
四則演算と√と三角関数とかを全ビットを演算器全部チェックしろと云う無茶ぶりがきた。
計算機が計算を間違うなんてありえないのに…
過去スレを読むとわかってもらえると思うが、みんなそれを求めているんだ。
そのチェックにおける最先端はおそらく長崎大学。
誰か調べてみろよ
もしかしてブロック辺りのスレッド数が32以下なら動いたりする？
何がしたいのか分からんがそれは簡単だ。
コアと計算結果の表が出来上がればコア間の相違が分かるじゃん
やり方知ってるなら教えてください
29じゃないけどｗ
入れてCUDAのためのNvidiaのドライバを入れて使ってみたが、OSのカーネル
のアップデートがあるたびに、コンソールが再起動時にブラックアウトした。
そのたびごとに、別のマシンからリモートログインをしてCUDAのための
Nvidiaのドライバを再インストールする必要があった。
CUDA3.xまではそれでなんとかなったが、現在のCUDA4.0になると、
幾らやっても駄目になってしまった。どうやらFedora15ではまともに
サポートされていないようなのだ。Nvidiaの提供しているCUDAのための
ドライバがインストール時におかしなメッセージが出てインストールが
出来ない。
ブロックあたりのスレッド数を32以下にして試してみましたが、
結果は変わらなかったです。
とりあえず、プログラミングガイドのコードそのままだし、
実際アルゴリズム自体には問題ないから、
意図どおり動かないとしたら、ドライバかSDKのバグか
ハードの故障じゃね。
まあ、ここで詳細なバージョンとか晒してみたり、
PTXやアセンブリを貼ってみるのもいいけど、
本家のフォーラムで聞いてみたほうが早いんじゃね。
間違えますって情報だけ垂れ流して具体例出さんって、どういうことなんだべ。
○演算間違い
relocation truncated to fit: R_X86_64_32S against `.bss'
とエラーが出て完了しません。gcc ではデータ量制限を解除するのに-mcmodel=medium というオプションがあるようですが、nvcc でこれにあたるものはないでしょうか。
CUDAは基本グローバルメモリ使用の制約があるから”巨大”は決して得意ではないのだけど
今確認したところ、要素が数十万個の配列を使用していて、それらの合計は200MBほどです。
そのうち、実際にGPUに計算させるためGPUに渡すデータは50MB程度です。
上記のエラーメッセージでググったところ、データ量制限が問題だろうということだったのですが…それほど大きなデータでもないですよね。
なんなんでしょうこのエラー。
プログラミングに関してはまだ知識が浅いためvンカンプンなこと言ってたらすみません。
日本語の間違えより、長崎大の言ってる間違いの詳しい情報くれよぅ…。
配列は何次元？
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];
}
とかにしたほうがいいよ。
gccと同様にオプションでエラー回避できないかな…と思ったのですが、配列の取り方を変えるほうが良さそうですね。
特に>>44は知りませんでした。
ありがとうございました！
-Xlinkerも要るかもしれないが。
普通はありえない。そんなことしたら、ポインタアクセスが一回余計に要るから遅くなるし、
キャッシュアウトしやすくなるから更に遅くなる。
まぁ、GPUに転送するのは1スライス分だけだろうからGPUに任せてCPUは暇こいていられるならいいかも知れんが。
それは無い。
・メモリも無駄に消費
・アクセス性能がた落ち
・バグの温床になる
つーかC++ならvectorで初期化時に
サイズ指定したほうがいろいろ便利
配列の添え字がINT_MAXを越える大きさならそうするしかないと思うが・・
要素のデカイ配列が確保できんって話なわけだし。
ヒープのフラグメンテーションの訳無いだろう。
nvccからgccに自由にオプション渡すやり方は
知らんが、最悪c(++)部分のコードを分割して
gccでコンパイルすりゃ良い。
仮にヒープのフラグメンテーションが酷い場合でも
>>44みたいなやり方は下策だよ。
素人におかしなやり方を覚えさせるだけ。
大きな配列確保する前に小さな物についてShrink-to-fitやればデフラグされるから
その後巨大配列を確保すれば良い。
特にGPGPUで巨大な配列扱うような用途なら
配列の配列なんて転送も困るだろうに。
>その後巨大配列を確保すれば良い。
コレ具体的にどんなコード書けばいいのん？
保持用のポインタなりvectorの変数だけ始めに
ヒープに確保した後、メインの計算前の前処理をする。
この処理で、一時的なデータと最終的に使うデータの
ヒープ確保が入り乱れる所為でフラグメンテーションが発生する。
最終的に使うデータについて、大体処理した順番を守って
ひたすら同容量確保してはコピーを繰り返せば、
こういうデータは全て前方に圧縮されて配置されるようになる。
まあスレ違いなんで他でやってよ。
今やってみたらコンパイルでエラーになった
構造体引数使えるべな
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.
Maximum dimension of a grid( 65535 , 65535 , 65535 )
ってなってるんですがz軸のブロックの取れる数は1じゃなくなったんですか？
例えば１００万回のモンテカルロを実行して、途中の１０万回、２０万回が済んだらその情報を取りたいんだけど…
streamingでできないかな?
SMのSじゃないよね？
俺は>>61を見てそういうのがあったなと思い出したレベルなので使い方は知らないけど。
http://hibari.2ch.net/test/read.cgi/software/1311428038/
すげえな
と逆に不安になるんだが…
なってきているから、CUDAプログラミングの旨味はいくらでもあるだろ。
いまだに開発環境のセッティングの仕方すら調べていない・・・
>>67くらいのコード書ける？
おれがそれくらいのコード書けるようになるころには、自ずとましな使い道が浮かんでくることだろう、
っていうこと？
試してないけど、長い処理の途中でグローバルメモリに書き込んで、
別streamでその値を読むカーネルを起動すれば取れると思う。
キャッシュからグローバルメモリに書き込むのは特別な関数を
使うと思ったけど名前忘れた
思い出した。atomicを使うか、書き込む側で __threadfence()
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です。
どなたかアドバイスください。
compute capabilityはnvccなりcl.exeのオプションに渡すものであって、
-archコマンドなんてものはないよ。
>>77
ありがとうございます。「cuda by example」にもatomic関数を使ったソースには-archコマンドラインオプションを指定する必要があります。
とあるのですが、みなさんはatomicを使ったソースをコンパイルするときなにかコマンドつけませんか＞＜？
あ、もちろん全部64bitです！
というか、VCなら、カスタムビルド規則でCUDA Runtime API Build Ruleを追加したほうが早い
になってるなぁ俺の環境だと。VC2010ならプロジェクトのプロパティでCUDA関係の設定するところがあるんじゃない？
ありがとうございます。
えーと、それはおそらくVC2008の話ではないかと思います。2010にカスタムビルド規則という名前はなくなり、その代りCUDA4.0というビルドカスタマイズに統一されています。
ご返答ありがとうございます。
そのコードはどこに書いていますか？
はい、cuda c/c++という項目にいろいろいじるところはあるのですがコンパイラオプションを設定する項目はないように思えます。
私がオプションを追加しているのはビルドイベントという項目ですね。
以前にもatomic関連の話が出ていたようです。>>22
にすればいいんじゃねえの？
ありがとうございます。
おっしゃるように>>75のように追加するとエラーを吐きます。
追加の仕方が違うのでしょうか。
知らんかったすまん。
とりあえず「CUDA VS2010」でググって出てくる一番上のサイト参考にして、CUDA4.0のビルドカスタマイズファイルを導入した後、プロパティからCode Generationをcompute_20,sm_20にしたらいけた
同じようなことをsandy bridgeとグラボ1枚でも出来るのでしょうか？
Z68マザーなら、マザー側DVIにディスプレイつないでおけばいい話じゃ？
CUDAで計算したデータをintelで表示したいなら、
1.DirectXかOpenGLのリソースをCUDAにバインドしてデータ書き込み
2.メインメモリに転送
3.intel側に転送
って形になるかなぁ。多分。
ゴメン、俺もVC2010じゃないや。VC2008。
ちなみに>>81は構成プロパティの中に「CUDA Runtime API」があってそこの「コマンドライン」にある。
CUDA Runtime APIの中にGUIでGPU Architecture（sm_21）を指定するところがあるから、
コマンドを打ってるわけじゃないんだけど、結果としてnvccには>>81のようなオプションが渡ってる。
>>81
ありがとうございます！code generationを>>80のように書きなおしたらオプション指定することなくコンパイル通りました。
どうやらVS2010からはcommand line指定しなくてもcode generationでオプションを渡せるようです。
本当に助かりました。ありがとうございます！
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にもかなり話題が上がってるっぽいんですが...私のところにも発病してしまいました。
nvccのエラーメッセージ載せないと。
敵性語なんでよくわかんなかったけど、
VisualStudio を管理者として実行したら動く？
プログラミングガイド嘘つきやないかー
ちくしょー
そりゃ命令実行時間は同じだけど、
必要な帯域は違う。
コマンドラインのオプションが何か1つ間違ってた気がする。
コマンドプロンプトで同じコマンド実行するとエラーが出たので、
そのエラーのひとつ前の引数が問題だった、はず。
行った方法は、CUDAのサンプルをRTコンポーネントに
パッケージングする方法。
動作周期にあわせて、行列計算が実行されたのを確認。
TOP カテ一覧 スレ一覧 2ch元 削除依頼 ▲
【Google】Androidアプリ作成part13 (294)
ポインタを難しいと言う奴が理解できない (797)
【アンチ】関数型言語は使えない【玩具】 ２ (226)
関数型言語Part5 (851)
UNIXプログラミング質問すれ Part10 (438)
【至急】助けてください。 (231)
--log9.info------------------
男でレオタードとピンクタイツでバレエ (483)
★東京シティバレエ団 ★ Part 5 (240)
[チャコット４・５階のバレエスクールについて」 (750)
札幌発★ダンス好き集合!! (614)
ロックダンスで使える曲教えて下さい！ (360)
誰かジャズダンスについてかたろウ！！ (311)
☆★ 神奈川のバレエ教室おしえて その4 ★☆ (713)
熊本のバレエ事情を教えて (101)
ＢｏＡ と 温水洋一、どっちがダンスうまいのか？ (168)
【サイン】 THE 出待ち ４ 【握手】 (959)
● 広島のバレエ事情 Part5 ● (540)
舞台化粧品は？ (308)
★★★来日決定ピナ・バウシュ★★★ (862)
＿オイ(ﾟДﾟ)！！NHK！もっとバレエ見せろ！＿ (636)
静岡のバレエ事情 (225)
☆バレエを題材にした漫画を語ろう５☆ (593)
--log55.com------------------
【Youtube】まーさんガレージを語る Part.12【レストア】
【非豪雪地専用】スタッドレスタイヤを語る part15
----車の運転が下手な奴の特徴----その27
☆【一種】教習所で普通免許を取る190号車【二種】
【EV】 電気自動車総合スレ IP有り その 52 【モーター】
＿＿＿＿＿AT免許とMT免許を分けろ！MT免許者にATを運転させる事は危険！＿＿＿＿＿
【ポン付け】純正ソケット用ＬＥＤ専用【市販品】43
うちさぁ、車あんだけど...乗ってかない？192