【GPGPU】NVIDIA CUDA質問スレッド

[表示 : 全て最新50 1-99 101- 201- 301- 401- 501- 601- 701- 801- 901- 2chのread.cgiへ]
Update time : 04/13 15:36 / Filesize : 252 KB / Number-of Response : 990
[このスレッドの書き込みを削除する]
[＋板最近立ったスレ＆熱いスレ一覧 : ＋板最近立ったスレ／記者別一覧] [類似スレッド一覧]


↑キャッシュ検索、類似スレ動作を修正しました、ご迷惑をお掛けしました

1 名前：デフォルトの名無しさん mailto:sage [2007/09/17(月) 14:54:28 ]: 本家
developer.nvidia.com/object/cuda.html
696 名前：デフォルトの名無しさん [2008/10/31(金) 17:56:27 ]: 最近銀行システムの開発で、6000人のSE集めた超プロジェクト失敗したものねえ。まあ当然だが。
SEが6000人だからねえ。プログラマはもっと多いとかもう想像つかない。
697 名前：デフォルトの名無しさん mailto:sage [2008/10/31(金) 21:10:20 ]: TMPGencのCUDA対応版が出たんでインストールしたんだけど
CUDAの項目にチェックできないのは何故・・？

ドライバは178.24でグラボがASUSの8800GTS(640MB)
698 名前：697 mailto:sage [2008/10/31(金) 21:19:21 ]: スマソ自己解決
g80はダメなんだってねOTL
699 名前：デフォルトの名無しさん mailto:sage [2008/10/31(金) 23:34:51 ]: >>698
ｲｷﾛ。
700 名前：デフォルトの名無しさん mailto:sage [2008/11/01(土) 00:39:03 ]: G80はストリーム系のAPIが使えないからねぇ。
701 名前：デフォルトの名無しさん mailto:sage [2008/11/01(土) 01:49:28 ]: >>696
JRとかの鉄道や、電力といったインフラ系はもっと大きい。
しかし、大きいが故にPJ失敗しまくってる。

人数を増やせば増やす程、集めた人材の質は低下する。
そして頭脳労働の場合、一番質の低い人のレベルに
足並みを揃えなきゃいけなくなるからなぁ。

しかし戦中・戦後に一気に作ったシステムが老朽化して、
銀行どころでなく大規模な改修がどれもこれも必要なのだが。

#mixiで見掛けたよ＞団子の中の人
702 名前：デフォルトの名無しさん mailto:sage [2008/11/01(土) 01:54:45 ]: >>662
HP ML115サーバ機に、GF9400GTあたり刺して、
Linux入れたら？ 1CD-Linuxの knoppix for CUDA
なら、最初からCUDA環境が構築済みで、サンプル
も憑いて来るし。

慶應義塾大学泰岡（やすおか）顕治研究室 Yasuoka Laboratory
www.yasuoka.mech.keio.ac.jp/cuda/

個人的にはGF8200なM/BのオンボでCUDA走れば、
裸M/BのCUDAクラスタ組もうかと思ってるが、
CPUやメモリの値段を考えると、ML115の方が
安上がりなんだよな。
703 名前：662 mailto:sage [2008/11/01(土) 06:57:01 ]: >>702
これはビックリ！こんな激安サーバがあるなんて知らなかった…
激安なのにPCI-Expressとかついてて（x16必須な）nvidiaのグラボもちゃんと動く、
ということでゲームの人達にも人気があると…ふむふむ。

ところで素のML115はメモリ512Mなのだけれど
上記研究室のページによるとknoppix for CUDAの推奨動作環境はメモリ2G以上、とある
ML115を使う場合、安いやつを別に買ってきて刺し換えればよろし、ということですね？
（ML115もhp直販だとメモリ増設オプションはECCつきの高いやつしかない…）

ML115が16k、9400GTが9k、2Gメモリも安いのは3k未満、で30kを切りますな。
個人的にサーバ機もAMD64もknoppixも使ったことがないので、
それらの組み合わせとなると微妙に不安だ（笑）が、いずれにせよこの値段は魅力的

大変参考になりました。ありがとうございます。
704 名前：デフォルトの名無しさん mailto:sage [2008/11/01(土) 07:57:06 ]: デモ機で借りたTeslaC1060使っているんだけど、ホストCPUがAMDのPhenom。
Xeonに較べて遅い遅い。普段使っているXeonに8800GTの組み合わせの方が早いって何さw
705 名前：,,・´∀｀・,,）っ-●◎○ mailto:sage [2008/11/01(土) 08:16:53 ]: つまんない質問だけどGTX2xxの人は電源いくらよ。
＋150Wくらいはマージンとったほうがいいと思うよな？よな？
706 名前：デフォルトの名無しさん mailto:sage [2008/11/01(土) 10:03:17 ]: なんに対して+150?
GTX280ボード単体での消費電力は236W、GTX260でも180Wクラス消費するからね。
ついでに言えば、補助電源用コネクタもGTX280は6ピン+8ピンの特殊コネクタが必要だし。
707 名前：,,・´∀｀・,,）っ-●◎○ mailto:sage [2008/11/01(土) 10:25:14 ]: >>706
システム全体で。500W電源以上推奨って言ってるけどじゃあ500Wで安定するかっていうと
信用できねー
708 名前：デフォルトの名無しさん mailto:sage [2008/11/01(土) 10:25:55 ]: 無理。
709 名前：,,・´∀｀・,,）っ-●◎○ mailto:sage [2008/11/01(土) 10:28:03 ]: とすると、マシン一式組んで貸し出してもらうのがベストだよな？
よし参考になった。
710 名前：デフォルトの名無しさん mailto:sage [2008/11/01(土) 11:59:03 ]: 8800GTなら100Wだし、補助電源コネクタも6ピンだけで済むよ。
711 名前：,,・´∀｀・,,）っ-○◎● mailto:sage [2008/11/01(土) 12:04:23 ]: うん、俺も8800GTまでなら550Wで余裕といえるラインかなと思っている。
712 名前：デフォルトの名無しさん mailto:sage [2008/11/01(土) 14:45:29 ]: GTX280を使うのなら、700Wクラスの電源が欲しいところだね。
713 名前：デフォルトの名無しさん mailto:sage [2008/11/02(日) 15:04:39 ]: CUDAは8800以上のクラスで無ければ意味ない。
8500とかはとりあえず走るだけでパフォーマンスは全然駄目。
714 名前：,,・´∀｀・,,）っ-●◎○ mailto:sage [2008/11/02(日) 17:39:50 ]: 8400GSとかになると額面性能でもCore 2シングルコア以下だな。
715 名前：デフォルトの名無しさん mailto:sage [2008/11/02(日) 19:59:18 ]: 　　そだ　 |------､｀⌒ｰ--､
　　れが　 |ﾊ{{　}} ）））ヽ､l l ハ
　　が　　　|､{ ﾊリﾉﾉﾉﾉﾉﾉ）､ l l
　　い　　　|ヽヽｰ､彡彡ﾉﾉﾉ}　　に
　　い　　　|ヾヾヾヾヾヽ彡彡} 　や
　　!! 　　 /:.:.:.ヾヾヾヾヽ彡彡} l　っ
＼＿＿／{ l ii | l|} ﾊ､ヾ} ﾐ彡ﾄ
彡ｼ ,ｪ､､､ヾ{{ヽ} l|l ｨｪ=ﾘ､シ} |l
lミ{　ﾞｲｼﾓ'ﾃ､ミヽ}ｼｨ=ﾗ'ｧ､ }ﾐ}} l
ヾﾐ　　￣~'ｨ''': |ﾞ:ｰ.￣　 lﾉ/l | |
ヾヾ　　 " 　: :　!､　｀　　lｲﾉ l| |
　>lﾞ､　　　　ｰ､,'ｿ　　 /.|}､ l| |
:.lヽヽ　　ｰ_ ‐-‐ｧ'　 /::ﾉl ﾄ､
:.:.:.:＼ヽ　　二"　 /::// /:.:.l:.:.
:.:.:.:.:.::ヽ:＼　　　　/::://:.:,':.:..:l:.:.
;.;.;.;.;;.:.:.:.＼｀ｰ-- '" //:.:.:;l:.:.:.:l:.:
716 名前：デフォルトの名無しさん mailto:sage [2008/11/03(月) 01:57:15 ]: サブノートPCでCUDA動くようにならんかな。
通勤、出張の途中でいぢってみたい。こんな
時でもないと、仕事に直結しないプログラム
組んでる暇無いからなぁ。

>>703
AMDの場合、メモリコントローラがCPUに内蔵なので、
ECCでもnon-ECCでも使える。安い通常のnon-ECCメモリ
1GBx2枚買ってくればOK。ML115はNTT-Xで買えば、
13800円(送料込)。
717 名前：デフォルトの名無しさん mailto:sage [2008/11/03(月) 02:19:41 ]: つ N10J
718 名前：662 mailto:sage [2008/11/03(月) 06:43:24 ]: >>716
回答ありがとうございます。
719 名前：デフォルトの名無しさん mailto:sage [2008/11/03(月) 11:13:39 ]: >>716
俺はサブノートでソースは書いてるよ。動作確認は自宅に戻ってからだけどね。
一発で動けば気持ちいいもんだ。
720 名前：デフォルトの名無しさん mailto:sage [2008/11/03(月) 11:52:24 ]: 質問です

OpenCLが出たらCUDAはお払い箱ですか？
721 名前：デフォルトの名無しさん mailto:sage [2008/11/03(月) 12:00:27 ]: >>719
エミュは動いている?

>>720
いいえ、画像処理だけがCUDAの使い道ではありません。
722 名前：,,・´∀｀・,,）っ-●◎○ mailto:sage [2008/11/03(月) 12:55:07 ]: Apple主導の言語処理系って流行らんだろ。
GPU版Objective-Cだと思え。

ちなみにNVIDIAから補助もらってる俺は仕事につながるって言えるのかな？
723 名前：デフォルトの名無しさん mailto:sage [2008/11/03(月) 13:08:35 ]: >>722
その仕事、こっちにくれw
情報少なくて、参ってるんだ。
724 名前：デフォルトの名無しさん mailto:sage [2008/11/03(月) 13:10:10 ]: >>722
ObjectiveCは言語仕様からしてクソだったから流行らなかった。
それだけです。
725 名前：,,・´∀｀・,,）っ-●◎○ mailto:sage [2008/11/03(月) 13:25:26 ]: >>723
メールサポートだけもらってるけどマニュアル落として自分でやったほうが早いしなぁ

俺のほうこそ各ptx命令のレイテンシ・スループットの資料欲しいんだけど。
Intelはそういうのまめに出してくれるから助かるんだが
726 名前：,,・´∀｀・,,）っ-●◎○ mailto:sage [2008/11/03(月) 13:29:01 ]: YellowBoxだっけ？
WindowsでもMacでも動くアプリケーションが動くフレームワークとか
大風呂敷広げてあれ結局どうなったっけ？

MicrosoftはDX11があるからOpenCLの標準化なんて破談する可能性大
Appleのフレームワークは地雷ばかりで困る。
727 名前：デフォルトの名無しさん mailto:sage [2008/11/03(月) 15:35:29 ]: >>725
なんだ、ないのか。NVIDIAの日本法人は、ろくに情報持ってないっぽいんだよね。
728 名前：デフォルトの名無しさん mailto:sage [2008/11/03(月) 16:03:03 ]: >>725
ptxは中間言語だろ？
729 名前：デフォルトの名無しさん mailto:sage [2008/11/03(月) 17:02:05 ]: >>726
先入観が身を滅ぼすだろう。
730 名前：デフォルトの名無しさん mailto:sage [2008/11/03(月) 17:55:12 ]: 質問スレッドなので、唐突に質問するわけですが、ごきげんよう

CUDAのSDKに付いてくる
Programming Guide Version2.0の60ページ目の真ん中あたり

For devices of compute capability 1.x,
the warp size is 32 and the number of banks
is 16 (see Section 5.1);
a shared memory request for a warp is split into one request
for the first half of the warp and one request for the second half of the warp.
As a consequence, there can be no bank conflict between a thread belonging to the first
half of a warp and a thread belonging to the second half of the same warp.

が分からない。
何が分からないのかというと、これはShared Memoryの最適なアクセスに関する記述なんだけど、
ワープの中に並列実行できるスレッドが32個あるというのにshared memoryのバンク数は16個しかない。
普通に考えたら２つのスレッドが同時に１つのbankにアクセスするわけで、
思いっきりバンク競合するはずよね？
でも、この記述はバンク競合が起こらないって自信を持って記述されているわけよ
nVidiaの人教えてちょ
731 名前：デフォルトの名無しさん mailto:sage [2008/11/03(月) 18:59:44 ]: Half Warp(つまり16スレッド)ずつスケジューリングされるんじゃなかったかな
だからバンク競合は起きない
nVidiaの人では無いが

なら何でWarp=16スレッドとしないんだろう…というのが俺の疑問
732 名前：,,・´∀｀・,,）っ管 mailto:sage [2008/11/03(月) 20:56:32 ]: 中の人いわく
命令レイテンシ隠蔽のためにクロック毎にインタリーブしてるだけだから細かいことは気にすんな
733 名前：730 mailto:sage [2008/11/03(月) 22:01:55 ]: なるほど～
ワープの正体は16並列と見つけたり
ってことですな！
734 名前：デフォルトの名無しさん [2008/11/04(火) 14:28:34 ]: GeForce 9400MってCUDA使える？
735 名前：デフォルトの名無しさん mailto:sage [2008/11/04(火) 16:04:29 ]: 2.1でサポートできるように頑張ってるけど間に合わないかもしんないって言ってた
736 名前：,,・´∀｀・,,）っ-●◎○ mailto:sage [2008/11/08(土) 18:21:17 ]: 今月中に何かしら動きが・・うわなにをする
くぁwせｄｒｆｔｇｙふじこｌｐ；「’」
737 名前：デフォルトの名無しさん mailto:sage [2008/11/08(土) 18:28:53 ]: 個人的には1.3世代の1スロット厚のGPUボードが欲しいのだけれど……
# 出ますと言ってた奴はその後連絡寄越さないしなぁ。
738 名前：,,・´∀｀・,,）っ-●◎○ mailto:sage [2008/11/09(日) 21:55:53 ]: コードの実行時動的生成（分岐除去とかパラメータの定数化とかってレベルで）って
CUDAではいまんとこ無理なんだよな？
Larrabeeが出たらそういう最適化できる部分はXbyak使おうかなと思ってるんだが

っていうか、SPMDじゃないプログラミングモデルまだー？
739 名前：,,・´∀｀・,,）っ-●◎○ mailto:sage [2008/11/10(月) 01:14:51 ]: >>730-732
に補足。

各SPは最大2issue同時実行なんだけどデコーダは半速。
1SPあたり4スレッドでインターリーブして同じオペレーションを実行するとちょうど命令供給が間に合う構造だな。

1warp＝
16にすると、デコーダは等速か、半速×2にしないといけない。
デコーダの負荷を抑えたかったんじゃないの？
740 名前：デフォルトの名無しさん mailto:sage [2008/11/11(火) 21:20:50 ]: CUDAはじめようと思って調べ始めたんだが、
7xxxシリーズはなんで切り捨てられたのか・・・

今週末に9600GT買いに行かなきゃ
741 名前：デフォルトの名無しさん mailto:sage [2008/11/11(火) 21:26:36 ]: どうせなら260だか280あたりにしといたら
742 名前：,,・´∀｀・,,）っ-○◎● mailto:sage [2008/11/11(火) 22:15:53 ]: 電源とかケースとかの敷居高くない？
743 名前：デフォルトの名無しさん mailto:sage [2008/11/11(火) 22:33:54 ]: >>740
切り捨てられたんじゃなくて、始めから想定されてない。

ttp://journal.mycom.co.jp/column/graphics/index.html
この連載のはじめの方のGPUの進化を追うと、少しは判るかも知れない。
で、どうせならQuadroFX3700をお勧めします。8800GTとほぼ同一仕様でお値段10倍w
744 名前：デフォルトの名無しさん [2008/11/12(水) 22:41:03 ]: CUDA-Zなんて便利なものがありました。
forum.nvidia.co.jp
745 名前：デフォルトの名無しさん mailto:sage [2008/11/13(木) 21:17:18 ]: それよりレイトレベンチマークのほう、Ｂｉｏ１００％が作ったのか！
SuperDepthとかカニミソとかが蘇ってきたぜ

スレ違い済まん
746 名前：デフォルトの名無しさん mailto:sage [2008/11/13(木) 21:20:01 ]: >>745
noridon.seesaa.net/
747 名前：デフォルトの名無しさん mailto:sage [2008/11/13(木) 21:58:30 ]: >>745-746
まだ生きてたのか！
PC-98では大変お世話になりました。

そしてブログを読んでみたら、超わかりやすい！
coalescedの意味とか、8/29のエントリみたいなメモリアクセスが遅い理由とかよく分からなかったんだよ。助かった。
748 名前：デフォルトの名無しさん mailto:sage [2008/11/15(土) 01:15:52 ]: vista sp1にCUDAをインストールしたいんだけど
ドライバ：○
↓
tool kit：×インストールが終了しない。。

で上手くインストール出来ないんですが、誰かしりません？
強制終了したらアンインストールの項目にtool kitの項目があるのにアンインストールするとerror:5001で失敗しやがるし。。。

最悪
749 名前：名無し募集中。。。 mailto:sage [2008/11/15(土) 02:28:14 ]: TMpegEncのCUDA対応は4フィルタだけで今のところあまり効果がないみたい
AviUtilのCUDA対応フィルタもパフォーマンス出ないという理由で公開停止
今からでもチャンスありますかね?
750 名前：デフォルトの名無しさん mailto:sage [2008/11/15(土) 03:36:03 ]: 作りたいなら是非作ってくれ
751 名前：デフォルトの名無しさん mailto:sage [2008/11/15(土) 06:15:48 ]: >>748
管理者権限でやっている?
後柱ね。
752 名前：デフォルトの名無しさん mailto:sage [2008/11/15(土) 21:32:05 ]: なんかLinux向けのドライバにCUDA2.1入ってるらしーよ
ttp://www.nvidia.com/object/linux_display_ia32_180.06.html
753 名前：デフォルトの名無しさん mailto:sage [2008/11/17(月) 19:18:56 ]: >>738
www.nvidia.com/docs/IO/47904/VolumeII.pdf
スライド87
754 名前：,,・´∀｀・,,）っ-○◎● mailto:sage [2008/11/18(火) 17:41:42 ]: >>752
ついに来てしまったか>>736
755 名前：デフォルトの名無しさん mailto:sage [2008/11/18(火) 23:45:16 ]: ラジオシティできるソフトってありますか?できればソース付きで...
756 名前：デフォルトの名無しさん mailto:sage [2008/11/20(木) 23:06:52 ]: Cg勉強しようと思って調べてたら、CUDAってのもあるんだな。
それぞれできる事って、具体的に何が違うの？

とりあえず7600GTしか持ってないんで、CUDAは使えないんだが、
Cg勉強するぐらいなら、CUDA勉強したほうが圧倒的に良いなら
対応グラボ買おうと思うんだが
757 名前：デフォルトの名無しさん mailto:sage [2008/11/20(木) 23:34:27 ]: そういう何処にでも載っていることすら調べられないのならCgにすれば。
758 名前：デフォルトの名無しさん mailto:sage [2008/11/21(金) 13:20:39 ]: CUDA 2.1 beta
ttp://forums.nvidia.com/index.php?showtopic=82290

VC++9とDX10インターオペラビリティがやっと
759 名前：デフォルトの名無しさん [2008/11/21(金) 21:39:18 ]: cuda sdkのサンプルを実行するとtest failedと出て実行できないんですけど。
環境はos xp 64, quadro FX 4600です。
先ほどnvidiaからドライバとツールとSDKをダウンロードして
インストールしました。ドライバは更新されています。
visual studio 2005も入れました。
760 名前：デフォルトの名無しさん [2008/11/22(土) 17:06:06 ]: Teslaを使っているのですが、電源コードを抜く以外の方法で、装置を再起動
する方法はないでしょうか。
バグのあるコードを何度も実行した結果、cudaMalloc()が返ってこない
状態になっています。
761 名前：デフォルトの名無しさん [2008/11/22(土) 18:21:05 ]: たわけた質問だと思いますが、お許しください。
NVIDIA製のカードが入っていないPC上で、
nvemulate.exeを利用してCUDAを使用する事は可能なのでしょうか？
実際の処理に使うのではなく、CUDAプログラミングの練習に使うのが主です。
762 名前：デフォルトの名無しさん mailto:sage [2008/11/22(土) 18:21:19 ]: >>760

Sシリーズならホストを再起動するだけで復活しませんか？
763 名前：デフォルトの名無しさん mailto:sage [2008/11/22(土) 18:48:56 ]: やっぱ大学くらいしかまだ使ってないのかね
764 名前：デフォルトの名無しさん [2008/11/23(日) 00:55:06 ]: 最近発売された、GeForce9300、9400を積んだＭＢ、
少し前のGeForce8200、8300を積んだＭＢでも実用ではないですが、
CUDAのプログラミングをして走らせる事は可能なのでしょうか？
誰もmGPUでCUDAを使っていないので…
765 名前：デフォルトの名無しさん [2008/11/23(日) 16:49:51 ]: みんな何の計算させてるの？
766 名前：デフォルトの名無しさん mailto:sage [2008/11/23(日) 16:53:52 ]: 株価予測をリアルタイムに
767 名前：デフォルトの名無しさん [2008/11/23(日) 19:03:26 ]: 株価の予測はできんだろ。アホか。
768 名前：デフォルトの名無しさん mailto:sage [2008/11/23(日) 20:24:25 ]: >>767
阿呆丸出し乙
769 名前：デフォルトの名無しさん [2008/11/23(日) 20:25:59 ]: 株価の予測ができたって言ってるのは、数年前の慶応が出してた論文ぐらいじゃねーの？
770 名前：デフォルトの名無しさん mailto:sage [2008/11/23(日) 20:35:29 ]: >>768
予想と予測は明確に違うんだぜ？
771 名前：デフォルトの名無しさん mailto:sage [2008/11/23(日) 23:52:44 ]: 論点がづれてるー髪もづれてるー
772 名前：sage [2008/11/25(火) 12:13:34 ]: >> 762
shutdown -> 電源切断 -> 電源投入の手順を踏むと、復活しました。
ただのrebootで良いかどうかは試していません。
773 名前：アク禁中なので纏めてレス mailto:sage [2008/11/25(火) 12:18:30 ]: >>772
色色と掲示板の使い方を間違っているw
で、reboot試してないなら報告しなくていいから。

>>771
髪はずれないと思うぞ、髪は。

>>765
私の所では、最近はFFTWの代わりにCUFFTでFFTを計算させている。

>>764
NVIDIAのサイトのCUDA ZONEでリストアップされていれば、使える。

>>763
んなこたーない。
774 名前：デフォルトの名無しさん [2008/11/25(火) 15:29:23 ]: --device-emulationでは正しく動くけれども、実機では動かないときには
ソースコードをにらむしかないのでしょうか。

nvcc --device-debug (-G) というオプションがあったので、これをつけて
コンパイルすると、ptxas が Parsing error を出して失敗します。

forums.nvidia.com/index.php?showtopic=66291
の会話を見ると、--device-debug は今年5月の段階ではまだ使えなかった
らしく、私の場合と現象が似ているので、以前としてまだ使えないままか
と思ったのです。
775 名前：デフォルトの名無しさん mailto:sage [2008/11/25(火) 18:33:51 ]: >>769
そりゃ、予想は出来る罠
ただ、外乱はいつも不明だし、確定解は得られない。
つまり、最尤推定しかできないし、当然推定結果が大ハズレってことも、
初めから推定理論に謳われてる
776 名前：デフォルトの名無しさん mailto:sage [2008/11/25(火) 19:24:20 ]: 株価予想が正確になればみんなそれを信じて買うようになるでしょ
予測自体が株価に影響を与えだして本来の予測とは違う値動きを始める
そして的中率は下がる
つまり一定以上の正確な予測を行うことは不可能なのだ
777 名前：デフォルトの名無しさん mailto:sage [2008/11/25(火) 20:36:28 ]: ここには、当たり前の簡単なことを、必死に難しく言おうとしてる
能無しがたくさん居るようだねｗ
778 名前：デフォルトの名無しさん mailto:sage [2008/11/26(水) 02:25:54 ]: １つの.cuの中で実装しているglobal関数の個数によって、
Kernel呼び出しのターンアラウンドタイムが変わるという奇妙な現象に遭遇してます。
特にKernelで処理するデータが少ない時に顕著になります。
関数の数を５～１０個で変えてみると、ターンアラウンドタイムは
最悪値で80μsec、最良で30μsecでした。
この値は
timer.start();
for (int i=0;i<100;i++) test_kernel<<grid,thread>>(test);
cudaThreadSynchronize();
timer.end();
みたいな書き方で調べてます。

9個目、10個目あたりで底があるようなのですが
こういう現象について、合理的な説明はありますか？
僕にはさっぱり見当がつかないのであります。
779 名前：デフォルトの名無しさん mailto:sage [2008/11/26(水) 02:28:14 ]: 何かを勘違いしている
780 名前：アク禁中（以下略 mailto:sage [2008/11/26(水) 20:28:00 ]: >>778
再現できるソースを貼ってくれたら解析するじょ。
781 名前：778 mailto:sage [2008/11/26(水) 23:00:43 ]: #include <windows.h>
#include <stdio.h>
#include <cuda_runtime.h>
//Round a / b to nearest higher integer value
int cuda_iDivUp(int a, int b) {return (a + (b - 1)) / b;}
#define BLOCK_DIM ( 32)
template <unsigned int loops>
__global__ void testcuuuuKernel(float* d_h0, unsigned int size)
{
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < size) {
float d = d_h0[i];
for (int j = 0; j < loops; j++) {d -= j * 0.1; d += 0.9;}
d_h0[i] =d ;
}
}

void dummy() {
dim3 block(BLOCK_DIM, 1, 1); dim3 grid(1, 1, 1);
testcuuuuKernel<4><<<grid, block>>>(NULL, 0);
//testcuuuuKernel<5><<<grid, block>>>(NULL, 0);
//testcuuuuKernel<6><<<grid, block>>>(NULL, 0);
}
int main(int argc, char* argv[]) {
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, 0);
cudaSetDevice(0);
#define DATA_SIZE ( 100)
unsigned int byte_size = DATA_SIZE * sizeof(float);
782 名前：778 mailto:sage [2008/11/26(水) 23:02:57 ]: float* data = new float[DATA_SIZE];
for (int i = 0; i < DATA_SIZE;i++) {data[i] = i;}
float* d_data; cudaMalloc((void **)&d_data, byte_size );
cudaMemcpy(d_data, data, byte_size, cudaMemcpyHostToDevice);
LARGE_INTEGER nFreq, nBefore, nAfter; //TIMER初期化
DWORD dwTime;
memset(&nFreq, 0x00, sizeof nFreq);
memset(&nBefore, 0x00, sizeof nBefore);
memset(&nAfter, 0x00, sizeof nAfter);
dwTime = 0;
QueryPerformanceFrequency(&nFreq);
#define LOOPNUM 100
dim3 block(BLOCK_DIM, 1, 1);
dim3 grid(cuda_iDivUp(DATA_SIZE, block.x), 1, 1);
for (int k = 0; k < 10; k++) { //試行の試行
//start!
QueryPerformanceCounter(&nBefore);
for (int i=0; i<LOOPNUM;i++) { testcuuuuKernel<3><<<grid, block>>>(d_data, DATA_SIZE); }
cudaError err=cudaThreadSynchronize();
//stop!!
QueryPerformanceCounter(&nAfter);
cudaMemcpy(data,d_data,byte_size,cudaMemcpyDeviceToHost) ;
dwTime = (DWORD)((nAfter.QuadPart-nBefore.QuadPart) * 1000000 / nFreq.QuadPart / LOOPNUM);
printf("%d usec for %d times kernel launch\n", dwTime, LOOPNUM);
Sleep(400); //ちょっと待つ
}
cudaFree(d_data); delete [] data; getchar(); return 0; }
783 名前：778 mailto:sage [2008/11/26(水) 23:07:45 ]: 再現できるコードを書いてみました。
Dummyという関数でテンプレート展開されている__global__関数の数を調整してみてください。
ちなみに使っているチップはGTX260です
784 名前：アク禁明けw mailto:sage [2008/11/26(水) 23:21:32 ]: >>783
面倒だから動かしてもじっくり読んでもいないのだけれど、
カーネル関数はGPUに都度転送することになるから
一回の呼び出し粒度が小さいと転送コストが目立つことになるよ。
その位だと、恐らくは起動コストも無視できないからもっと処理させるべき。
つーか、カーネル呼び出し(<<<>>>)をループで包んだらそりゃ遅いって。
785 名前：778 mailto:sage [2008/11/26(水) 23:28:05 ]: >>784
もちろんそれは分かるのですが、カーネル呼び出しの処理の内容は、<<<>>>の中の次元数に束縛される
傾向にあると思います。
データ構造などが異なる色々な処理を連携して実行する場合は、カーネル内の分岐では限界があるように思われます。
なので、１回当たりのカーネルのレイテンシを正確に把握しておきたいわけです。
786 名前：デフォルトの名無しさん mailto:sage [2008/11/26(水) 23:35:04 ]: ＞データ構造などが異なる色々な処理を連携して実行する場合は、カーネル内の分岐では限界があるように思われます。
何にも判ってないと思われ。
787 名前：,,・´∀｀・,,）っ-○◎● mailto:sage [2008/11/26(水) 23:39:07 ]: っていうか、分岐ってさ、プレディケートで全部実行するんだよな？

if (cond) { //ここの条件は要素ごとに変わる
　　funcA();
} else {
　　funcB();
}

だったら、funcAとfuncBをインライン展開して全部プレディケートつき実行する感じだと思ってるんだが。
788 名前：デフォルトの名無しさん mailto:sage [2008/11/26(水) 23:42:08 ]: 困ったことに、団子に同意。
789 名前：デフォルトの名無しさん mailto:sage [2008/11/27(木) 00:03:15 ]: 条件分岐したら負けかなと思ってる by GPU
790 名前：デフォルトの名無しさん mailto:sage [2008/11/27(木) 00:32:04 ]: そいえばCUDAって１つのカーネルのサイズが制限されてない？
でかいやつがまったく動かなくて苦労したんだけど
791 名前：デフォルトの名無しさん mailto:sage [2008/11/27(木) 01:39:04 ]: どの位かは知らんが、そりゃぁ制限はあるだろうねぇ。
792 名前：,,・´∀｀・,,）っ-○◎● mailto:sage [2008/11/27(木) 01:54:09 ]: 64Kのコンスタントメモリがあるじゃん。
コンスタントメモリは自分自身では中身の入れ替えは不可能。
あとはわかるよな？
793 名前：デフォルトの名無しさん mailto:sage [2008/11/27(木) 02:23:46 ]: cudaで自己書き換えプログラムってできますか？
794 名前：,,・´∀｀・,,）っ-○◎● mailto:sage [2008/11/27(木) 02:37:41 ]: GPUのカーネルコード自身で書き換えるって意味なら無理。
PTXのバイナリコードを動的生成とかなら何かやれば可能かもしれない。

その辺の資料を中の人に要求したら

「機密事項ですのでお答えできません」
795 名前：デフォルトの名無しさん [2008/11/27(木) 17:49:14 ]: 大学の研究室にCUDAプログラミング用のコンピュータが導入された！と喜んでいたら、
HP ML115 + げふぉ8400GSカードだった…orz 学習用仕様で萎えた…
796 名前：デフォルトの名無しさん mailto:sage [2008/11/27(木) 17:56:05 ]: まぁ、学部生なら十分だろ。
どうせ大した論文も書かないくせに。
797 名前：デフォルトの名無しさん mailto:sage [2008/11/27(木) 18:59:33 ]: >>796
お前みたいなの、必ず湧くよなｗ
人を馬鹿にしたら、自分が偉くなるとでも思ってんの？
798 名前：デフォルトの名無しさん mailto:sage [2008/11/27(木) 19:42:03 ]: でも俺もそこまでは言わないにしても
学習用仕様で萎えたってセリフは贅沢だと思うよ。
799 名前：名無し募集中。。。 mailto:sage [2008/11/27(木) 19:46:05 ]: 値段を見て萎えたって意味だろうが性能的には十分な気がするんだけど
800 名前：デフォルトの名無しさん mailto:sage [2008/11/27(木) 19:54:22 ]: そもそもGPGPGUはなんなんだ？8800はVGA用途だろ
801 名前：デフォルトの名無しさん mailto:sage [2008/11/27(木) 20:23:28 ]: >>797
この程度で反応するなよｗ
社会でやっていけないぞ
802 名前：デフォルトの名無しさん mailto:sage [2008/11/27(木) 22:49:14 ]: >>795
それなんて俺
CUDA使ってみたくて買い換えを考えていたけど
PCI Express x16バス搭載で1万だったのでML115と
玄人志向の8400GSカード買ってきてCUDA環境を手に入れたぜ
メインで使ってるマシンより性能がよくてこのままメインになりそうな予感
803 名前：デフォルトの名無しさん mailto:sage [2008/11/27(木) 22:51:16 ]: 事前に断っておくと俺は批判されるとチンコがたつよ
804 名前：デフォルトの名無しさん mailto:sage [2008/11/27(木) 23:04:29 ]: 自宅でCUDA使ってる人のコンピュータスペックってどんなもん？
805 名前：デフォルトの名無しさん mailto:sage [2008/11/27(木) 23:05:58 ]: >>804
c2d e6320
806 名前：デフォルトの名無しさん mailto:sage [2008/11/27(木) 23:11:07 ]: >>804
PenD920 GeForce9600GT
807 名前：デフォルトの名無しさん mailto:sage [2008/11/27(木) 23:31:34 ]: >>804
Core i7 920 + GeForce GTX 280
808 名前：デフォルトの名無しさん mailto:sage [2008/11/28(金) 00:15:30 ]: >>804
Xeon + GeForce GTX280
809 名前：798 mailto:sage [2008/11/28(金) 03:29:12 ]: >>799
いや、新技術半可通の俺が察するにGeF8000系は
CUDA対応ハードの底辺だから文句言ってるんだよ。
研究室で導入する(自腹じゃない)んだからもっと良いのよこせってことだろ？

ぶっちゃけ研究といっても今の時期じゃ全体的に未成熟だし
そのスペックでさえ限界動作なんぞせんだろと思って贅沢だなと。
810 名前：デフォルトの名無しさん mailto:sage [2008/11/28(金) 03:39:20 ]: >>800
そんなことはない。GPGPUは対応ハードならグラボ1枚でも2枚でも一応使えるんだよ。
実例は押さえてないんでなんともいえないが
競合さえおきなきゃRadeon(VGA)・GeF8800以上(GPGPU専用)ということも出来るはず。

用途と価格の兼ね合いもあるだろうけど、
とりあえず試したいとかコストパフォーマンス的に有利な部分もあるかもしれん。

4gamerの記事が参考になるよ。
www.4gamer.net/games/050/G005004/20080615001/
811 名前：デフォルトの名無しさん mailto:sage [2008/11/28(金) 04:14:57 ]: >>809
底辺どころかメインストリーム
812 名前：798 mailto:sage [2008/11/28(金) 06:47:41 ]: >>811
ラインナップ上は確かGeF8世代からの対応だから一応底辺だろう。
>>795の萎える理由に対しての推測だからこの説明でいいんじゃね？

もっとも現状ではSP数の差による性能差がそこまで決定的じゃないから
2x0とかQuadroを期待してるんだったら贅沢なんじゃねって話。
813 名前：デフォルトの名無しさん mailto:sage [2008/11/28(金) 07:50:33 ]: GT2x0:性能は兎も角、(電源などの)要求仕様が一般的でない。
QuadroFX:同程度の性能のGeforceの数倍の価格と言う段階で、論外。
# なんて書くと、「アキバ的発想云々」って言われるんだけどねw
もしこれらが理由なら、見識不足と言わざるを得ない。
単に、>795はML115が気に入らないんだろw

まぁ、今時なら8400GSでも8600GTでも9600GTでも大して値段変わらんと思うが。
814 名前：デフォルトの名無しさん mailto:sage [2008/11/28(金) 11:07:07 ]: ML115の電源じゃパフォーマンスレンジのGPUは無理ね
815 名前：名無し募集中。。。 mailto:sage [2008/11/28(金) 11:14:47 ]: ML115の電源はサーバー向けだけあってそれなりに良い電源だよ
電源投入後の全開爆音は凄いけど
ちなみに容量は370W
816 名前：デフォルトの名無しさん mailto:sage [2008/11/28(金) 11:52:57 ]: それじゃ精々8600GTSか9600GT止まり。8400GSは案外無難な選択じゃないのかな。
817 名前：デフォルトの名無しさん mailto:sage [2008/11/28(金) 13:12:49 ]: ちょっと誰か日本の公式フォーラムのcudaggさんの日本語をなんとかしてあげて
818 名前：デフォルトの名無しさん mailto:sage [2008/11/28(金) 13:51:35 ]: Ubuntuの話なら、一応意味は判るでしょ。
つーか、ここでフォローしても意味ないしw
819 名前：デフォルトの名無しさん mailto:sage [2008/11/28(金) 18:56:55 ]: >>813
GT2x0で一般的でないとか言ってたらCUDA自体一般的じゃないだろjk
ML115が不服と言うのでも十分贅沢……。俺が行ってた某理系大学なんか(ry
820 名前：流石理系大学出は日本語読むことさえ放棄しているようだな。 mailto:sage [2008/11/28(金) 18:59:40 ]: >>819
GT280は200Wの電源とそれに見合う通風環境が必要になるってこった。
この際、CUDAが一般的かどうかが問題じゃない。
821 名前：考えること放棄してるやつよりマシだよ。 mailto:sage [2008/11/28(金) 19:53:01 ]: >>820
CUDAスレでの話だからマシン云々というよりは
単に出たばっかのGT2x0系じゃなかったことに萎えてるのかと思っただけじゃん。

そもそも研究室とかだったら要求仕様が一般的でないなんてことは些細なことだ。
例えばOracle高価だから絶対使いませんなんて真似しないだろ。
そういうところは必要あれば用意するだろうよ。
GT2x0載せるとしたらそりゃML115じゃ無理だろうけど、それに載せろとは誰も言ってないし。
一般的云々言われなきゃ俺だってCUDAが一般的かどうかなぞ持ち出さんわ。
822 名前：デフォルトの名無しさん mailto:sage [2008/11/28(金) 19:55:38 ]: グラボの方に萎えてると思ったもうひとつの理由が、
CUDAプログラミング用ということだから
マシン自体のスペックはそれほど高くなくても良いはずなんだ。
電源さえ足りてればな。深読みしすぎた俺が悪かったよ……。
823 名前：デフォルトの名無しさん [2008/11/29(土) 01:30:32 ]: そこいらでレスをつけてるやつらの中に、8000系では一部使えない命令がある件を指摘するやつが居ないのは何でだぜ？
お前ら実は妄想だらけでなんもしてないんじゃないの？
824 名前：デフォルトの名無しさん mailto:sage [2008/11/29(土) 01:46:30 ]: 8400GSならCompute Capability 1.1だから
GT200で追加された命令でもなければ使えるわけだが。
825 名前：デフォルトの名無しさん mailto:sage [2008/11/29(土) 02:59:51 ]: >>823
1.0世代は8800GTX,GTS,ULTRAに積まれたG80くらいで、8600GTや8400GSに積まれたG84やG86は
1.1世代だと言う知識もなしにここ見てたの? 他人を妄想だらけなんて指摘してられる状況じゃないじゃんw
826 名前：デフォルトの名無しさん mailto:sage [2008/11/29(土) 03:20:49 ]: CUDAを単純に使いたいやつは、

HP ML115 (nttxstore.jp/_II_HP12591344)
と
G98の8400GS買えばおｋだよ。
827 名前：デフォルトの名無しさん mailto:sage [2008/11/29(土) 19:48:41 ]: グラボ早い順に並べるとどんな感じ？
やっぱりSP数×クロックなの？
828 名前：デフォルトの名無しさん mailto:sage [2008/11/29(土) 21:13:40 ]: 仕様を限定せずに一般的に言うなら、そりゃそうだ。他にどんな要素が来ると?
特定の応用でメモリ転送が重視されるのならメモリ帯域やバス仕様にも依存するし、
メモリ量が多い方が高速なアルゴリズムを採用できる応用ならメモリ量にも依存するわけだけど。

あー、1.0世代だとストリームが使えないと言う大きな欠点はあるね。
829 名前：デフォルトの名無しさん mailto:sage [2008/11/30(日) 00:11:00 ]: チップ外は考えてなかったφ（．．）ﾌﾑﾌﾑ…

後からSPが減ってクロックが上がった製品が出るじゃない？
だいたい値段は一緒だけど、劇的に性能が違うのかなと…
830 名前：デフォルトの名無しさん mailto:sage [2008/11/30(日) 01:46:04 ]: 同じ名前のSP減ったのは、恐らく歩留まりの関係で在庫処分しているだけだと思う。
使い方にも拠るけど、クロックが利く応用なら結構変わるかもね。
# SP数が利く応用なら逆に遅くなりかねない。
831 名前：デフォルトの名無しさん mailto:sage [2008/12/01(月) 03:44:29 ]: やっぱみんなPDEとか立てたり解けたりできるの？
832 名前：デフォルトの名無しさん mailto:sage [2008/12/01(月) 06:49:36 ]: そうだね
SDEとかね
833 名前：デフォルトの名無しさん mailto:sage [2008/12/01(月) 23:36:29 ]: >>795
あぅっ、それはウチかも知れない。

まぁどーせ学生の作るプログラムなんて、レポート見てると
自分で作ってるのは15%ぐらいで、残りの人は友人のコピペ
改変か、ググって見付けたページのコピペ改変だからな。
予算の都合ってやつで、1人1台用意するならML115+GF8400
に成ってしまう。許せ。
おっ、これは?!と思うプログラムが出てきたら、GTX280で
走らせてやるから頑がってくれ。
834 名前：デフォルトの名無しさん mailto:sage [2008/12/01(月) 23:40:13 ]: >>816
PCI-Ex補助電源コネクタ無いので、GF8600GT
か、GF9500GTまでが無難なところ。
835 名前：デフォルトの名無しさん mailto:sage [2008/12/02(火) 00:15:15 ]: texture<float, 1, cudaReadModeElementType> tex;
のテンプレートで
> error C2018: 文字 '0x40' は認識できません。
ってエラーが大量に出るんだが、みんな出ないですか。
オレッスカ
textureのテンプレート使わなければでないのですが。

CUDA 2.1 + VC++ Express 2005
836 名前：デフォルトの名無しさん mailto:sage [2008/12/02(火) 00:20:06 ]: >>835
PDFからコピペしてない？
837 名前：デフォルトの名無しさん mailto:sage [2008/12/02(火) 00:36:23 ]: >>836
一回行コピペしてエラー出たので消して書いたのですが、
ファイルに変なコードが残るとかありますかね。
でもその行消すとエラーなしで、手書きで書いたら、
↓のようになります。

1>.cu(54) : error C2018: 文字 '0x40' は認識できません。
...
1>.cu(54) : error C2018: 文字 '0x40' は認識できません。
1>.cu(54) : error C2065: 'COMPILER' : 定義されていない識別子です。
1>.cu(54) : error C2146: 構文エラー : ';' が、識別子 'ERROR' の前に必要です。
1>.cu(54) : error C2065: 'ERROR' : 定義されていない識別子です。
1>.cu(54) : error C2143: 構文エラー : ';' が '<template-id>' の前にありません。
1>.cu(54) : error C2146: 構文エラー : ';' が、識別子 'tex' の前に必要です。
1>.cu(54) : error C2275: 'texture<T,dim,__formal>' : この型は演算子として使用できません
1> with
1> [
1> T=float,
1> dim=1,
1> __formal=cudaReadModeElementType
1> ]
1>nv_som_gpu.cu(54) : error C2065: 'tex' : 定義されていない識別子です。
838 名前：837 mailto:sage [2008/12/02(火) 00:46:22 ]: texture<float1, 1, cudaReadModeElementType> *tex = new texture<float1, 1, cudaReadModeElementType>;
ならOKなので、スタックにtextureを置くとダメな条件があるのでしょうか。
839 名前：837 mailto:sage [2008/12/02(火) 01:31:47 ]: kernelに引数で渡せないので
スタックにおいたらダメみたいでした。
解決。
840 名前：デフォルトの名無しさん mailto:sage [2008/12/02(火) 23:13:52 ]: まだ、初めてみようかと考えている初心者以前のものです。
現状の自作数値解析プログラムはMPIで各ノードの各コアにプロセスを振っております（ノード内もMPIで並列）。
これの自然な拡張としては、MPIの各プロセスがGPGPUを使うという形になるかと思います。
しかし、マルチCPUのノードの場合、ひとつのGPUを共有することになります。
一つのプロセスがGPUのシェーダのうち1/4だけ使って、4プロセスから同時にGPUを使うなんてことは可能なのでしょうか？
841 名前：デフォルトの名無しさん mailto:sage [2008/12/02(火) 23:47:13 ]: >>839
テクスチャはコンパイルするとただの数値みたいになるので、
グローバル変数としてそのまま使う以外のことはほとんど不可能、らしいよ。

関数にパラメータとしてを渡すのも、ポインタを得るのも無理。
842 名前：デフォルトの名無しさん mailto:sage [2008/12/03(水) 00:02:52 ]: >>840
CUDAを使う場合、4プロセスから同時に使うことも可能。
但し、GPUをどう割り振るかはCUDAドライバの御心次第。
厳密に制御するには、GPU1枚ごとに担当スレッドを設けることになる。
# その場合、プロセス間通信でJOB型にするか1プロセスだけでGPUを占有するかは設計次第。

それはいいけど、プロセス全部分けるとマルチスレッドに較べて効率落ちないかい?
843 名前：840 mailto:sage [2008/12/03(水) 00:20:11 ]: >>842
ありがとうございます。
やはりリソースの競合が起きるようですね。

マルチスレッドは単純に覚えることが増えるのでやってませんw
スパコンのマニュアルでMPIを覚えたので、WSクラスタでもそのままの手法を持ってきてます。
一応、ノード内では共有メモリを使って通信するようなオプションでmpichをインスコしてるので多少はマシでしょうという感覚です。
844 名前：名無し募集中。。。 mailto:sage [2008/12/03(水) 02:19:05 ]: NVIDIA、PhysX/CUDAを活用する「パワーパック」の第2弾を提供開始
journal.mycom.co.jp/news/2008/12/03/001/index.html
845 名前：デフォルトの名無しさん mailto:sage [2008/12/03(水) 03:31:23 ]: 0x40でググったら
>文字 '0x40' は認識できません。原因：ソースコード中に全角空白 ' ' が使われています。

まさか、な
846 名前：デフォルトの名無しさん [2008/12/04(木) 01:11:01 ]: こんな感じで網羅的に点を回転させるプログラムを書いているのですが、
Thread数が320を越えるあたりでrotate関数に全ての引数が渡らなくなってしまいます。
その時でも計算量が少ないからか、roll_axisのz成分などは、引数として機能しています。
roll_axisのx,y成分も適当な値に変えると引数として働くようになります。
rotate(&coord, &pitch_axis,ANGLE);をコメントアウトすれば,roll_axisのx,yはThread数が多くても引数として渡ります。
計算量が多くなると(特に三角関数？)起きる気がするのですが、何が回避する方法はあるのでしょうか？

--ptxas-options=-vはUsed 42 registers, 56+28 bytes lmem, 2080+32 bytes smem, 8024 bytes cmem[0], 88 bytes cmem[1]て出ます。
-deviceemuでは正常に動作します。
847 名前：846 [2008/12/04(木) 01:13:18 ]: >>846のソースコードです。
####kernel(Thread１つが回転させる点１つに対応)####
__device__ runDevice～～の一部
float4 yaw_axis = make_float4(0, 0, -1, 0);
for(iyaw = 0; iyaw < limY; iyaw++){
float4 pitch_axis = make_float4(-sin(radian*iyaw), cos(radian*iyaw), 0, 0);
for(ipitch = 0; ipitch < limP; ipitch++){
float4 roll_axis = make_float4(cos(radian * iyaw) * cos(radian * ipitch), sin(radian * iyaw) * cos(radian * ipitch), sin(radian * ipitch), 0);
for(iroll = 0; iroll < limR; iroll++){
if(tid == iroll + ipitch * limR + iyaw * limR *limP)
g_mem[tid] = coord.x; //回転できているか確かめる。

rotate(&coord, &roll_axis, ANGLE); //回転させる点の座標と、回転軸と角度を与える。
}
rotate(&coord, &pitch_axis, ANGLE);
}
rotate(&coord, &yaw_axis, ANGLE);
}

__device__ void rotate(float4 *coord, float4 *axis, int angle){
coord->x = axis->x;
coord->y = axis->y; //とりあえず現段階では引数が渡るか確かめてるだけ
coord->z = axis->z;
}
848 名前：デフォルトの名無しさん mailto:sage [2008/12/05(金) 14:25:22 ]: 誰か、すげぇ簡単単純だけど、真理を付いてる様な神サンプル晒してよ
849 名前：デフォルトの名無しさん mailto:sage [2008/12/05(金) 15:02:45 ]: >>848
どんなの? SDKのサンプルやbioのblog辺りじゃお気に召さない?
具体的なテーマがあって、実装が難しくなさそうなら作ってもいいけどね。
850 名前：デフォルトの名無しさん mailto:sage [2008/12/05(金) 21:35:18 ]: ATI Streamから来ました。
スイマセン、場違いな場所に来てしまった。
AもまだなのにCなんて出来ません
851 名前：デフォルトの名無しさん mailto:sage [2008/12/05(金) 23:37:57 ]: 共産主義者の書いたマンガによると、Aより簡単らしいぞ。
852 名前：デフォルトの名無しさん mailto:sage [2008/12/06(土) 00:23:10 ]: 最近の若い連中はしょっぱながCだからな
853 名前：デフォルトの名無しさん mailto:sage [2008/12/06(土) 19:46:07 ]: cudaってさ
C-daにしておけばギャグっぽくてよかったんじゃね？
854 名前：デフォルトの名無しさん mailto:sage [2008/12/07(日) 12:03:03 ]: それならCarracudaの方がｗ
855 名前：デフォルトの名無しさん mailto:sage [2008/12/08(月) 22:50:38 ]: teslaの話題が少なくて絶望した！
856 名前：デフォルトの名無しさん mailto:sage [2008/12/08(月) 23:03:23 ]: えー、QuadroFX5600のアナログ回路をとっ外しただけの代物の、何を語れと言うのさ。
聞いてくれたら答えるけど。
あ、4桁シリーズは白根。どうせ事情は一緒だと思うけど。
857 名前：デフォルトの名無しさん mailto:sage [2008/12/08(月) 23:44:28 ]: >>855-856
このスレの住人は　GPGPUの用途にQuadroFX使ってるの・・・？
ていうかそもそもみんな何使っているの？
858 名前：デフォルトの名無しさん mailto:sage [2008/12/08(月) 23:58:17 ]: ML115最強伝説
859 名前：デフォルトの名無しさん mailto:sage [2008/12/09(火) 00:12:18 ]: Tesla D870 が至高
860 名前：デフォルトの名無しさん mailto:sage [2008/12/09(火) 00:59:33 ]: >>857
このスレともう一つのCUDAスレを見ると、QuadroFXのメリットが書かれていると思うが。
# 普通、使うわけないだろって。
8800GTX、QuadroFX5600、Tesla C870の価格を見る限り、Teslaは未だましだと思う。

>>859
D870ってNVIDIA謹製ミニタワーにC870を二枚入れた代物だっけ?
あれだったら同じ筐体のQuadroPlexの方が潰しが利くと思うのだけど。
861 名前：デフォルトの名無しさん mailto:sage [2008/12/09(火) 01:51:03 ]: GeForce GTX 280 極上
862 名前：デフォルトの名無しさん mailto:sage [2008/12/09(火) 20:16:56 ]: ９６００GT買ったらCUDAでエラー出まくりで問い合わせたら
CUDAのようなメーカー付属のドライバで対応してない機能は保障外とのこと
CUDAをやるために買ったのにCUDAが動くことを保障してないなんて
はっきり言って詐欺だよ
863 名前：デフォルトの名無しさん mailto:sage [2008/12/09(火) 20:51:59 ]: DELLにVisual Basicで組んだプログラムでエラー出るんだけど？
って聞いてるようなもんだな。
864 名前：デフォルトの名無しさん mailto:sage [2008/12/09(火) 20:54:42 ]: >>862
nvidiaのドライバだとどうなんだ？
マジに教えてくれ
865 名前：デフォルトの名無しさん mailto:sage [2008/12/09(火) 21:07:30 ]: >>863
CUDA SDKのテストプログラムでエラーが出まくる
実行する度に計算結果が違ってくる

>>862
どういう意味？NVIDIAのドライバ以外にCUDAが動くの？
付属のCDのドライバ以外はNVIDIAの公式ドライバでも保障しないと言われた
866 名前：デフォルトの名無しさん mailto:sage [2008/12/09(火) 23:35:15 ]: >>862
一体どこの糞メーカのPCなんだ?
867 名前：デフォルトの名無しさん mailto:sage [2008/12/10(水) 23:15:44 ]: >>865
>付属のCDのドライバ以外はNVIDIAの公式ドライバでも保障しないと言われた
それが普通ですよ。販売メーカとしては、自分のところでテストした
ドライバ以外は保証してないよ。

だから公式ドライバ入れる時とかもよく言われるでしょ。
トラブっても自己責任で、ってね。
868 名前：デフォルトの名無しさん mailto:sage [2008/12/11(木) 00:00:11 ]: まあ当たり前なのは分かるんだけど、そこんとこ盲点だよ。
付属のバージョンだと根本的にCUDA機能が入ってないからね。
必然的にアップデートしないといけないのに保障外でしょ。
CUDAは素晴らしいとか大々的に宣伝してるけど動くかどうかは運次第という対応だからね。
これが世間に知れたらCUDAなんて誰も相手にしなくなると思うけど。
869 名前：デフォルトの名無しさん mailto:sage [2008/12/11(木) 00:16:59 ]: tesla売るためだろ
870 名前：デフォルトの名無しさん mailto:sage [2008/12/11(木) 00:19:56 ]: そうだろうね。TESLAは唯一保障された製品だからCUDAやりたい奴はTESLA買えってことなんだろう。
一般人には何の恩恵もない誇大広告だったというわけですか。
871 名前：デフォルトの名無しさん mailto:sage [2008/12/11(木) 00:23:48 ]: メーカー純正フロッピーディスク以外使わない奴なんていなかっただろ。
みんなそうやって保証外の物を使ってきているんだ。今更何を言っているんだか。
872 名前：デフォルトの名無しさん mailto:sage [2008/12/11(木) 00:27:34 ]: で、そんなふざけた事言うのはどこなんだ？
873 名前：デフォルトの名無しさん mailto:sage [2008/12/11(木) 00:27:39 ]: どこかの誰かが同様にCUDAが使えないとなると自分がCUDAを使う意味がない。
だからCUDAは必要なくなるので別に何も問題ない。
誰もが使えないプログラムなんて作っても時間の無駄でしょ。
保障がないってのはそういう事よ。個人的な恨みとかじゃなく。
874 名前：デフォルトの名無しさん mailto:sage [2008/12/11(木) 01:57:49 ]: ノートPCなんてNVIDIAが公式ドライバを提供しないんだぜ。
メーカーも提供しないから、アマチュアの作ったものを探してくるしかないっていう。
875 名前：デフォルトの名無しさん mailto:sage [2008/12/11(木) 04:03:59 ]: あー、WindowsでCUDA試してないな、うちのMac
876 名前：デフォルトの名無しさん mailto:sage [2008/12/11(木) 21:46:41 ]: 保証が無いからどうだとかいうのは自作板でやってくれ。
877 名前：デフォルトの名無しさん [2008/12/12(金) 02:15:59 ]: >>857
犬板にも書いたけど、JetWay HA05-GTのオンボッボで動きますた。
USBメモリ起動で、薄いCUDA nodeの出来上がりです。
オンボッボももっと沢山SPU載る日はいつかなぁ。
878 名前：デフォルトの名無しさん mailto:sage [2008/12/12(金) 18:46:10 ]: 【GPGPU】CUDA/ATI　STREAM　速度・画質検証スレ
pc11.2ch.net/test/read.cgi/jisaku/1228764782/
　
879 名前：デフォルトの名無しさん [2008/12/12(金) 20:35:23 ]: これからCUDAを使おうと思っているのですが、visualC++でうまくビルドできません。どなたか解決方法を教えてください。
環境はXPx64Edition,QuadroFX570でx64版のバージョン2.0のドライバとToolkitとSDKを入れました。
VisualC++は2005Expressです。
880 名前：デフォルトの名無しさん [2008/12/12(金) 20:39:35 ]: すみません、879ですが、ビルドできないのはSDKの中にあるprojectsのsimpleTemplateです。
webで調べてみると皆さん問題なく実行できるようなのですが。。。
881 名前：デフォルトの名無しさん mailto:sage [2008/12/12(金) 22:05:56 ]: どんなエラーが出るとかの情報はないの？
882 名前：879 [2008/12/12(金) 23:26:03 ]: 失礼いたしました。下のようなエラーがでます。
nvcc fatal : Visual Studio configuration file '(null)' could not be found for installation at 'C:/Program Files (x86)/Microsoft Visual Studio 8/VC/bin/../..'
>>444さんと同じような問題と思われます。VisualC++のコンパイラとの相性が悪いのでしょうか。詳しい方、よろしくお願いいたします。
883 名前：デフォルトの名無しさん mailto:sage [2008/12/13(土) 06:12:02 ]: こんな連中ばっかりだし、クマには生きづらい世の中になったぜ
884 名前：デフォルトの名無しさん mailto:sage [2008/12/13(土) 12:41:18 ]: C++もろくに使えない人がいきなりCUDAに手を出すのもどうか
885 名前：デフォルトの名無しさん mailto:sage [2008/12/13(土) 14:55:07 ]: なんか64bit環境でもVCのclは32bit版使わないとだめみたいなことCUDAフォーラム（英語のほう）に書いてあったよーな
886 名前：デフォルトの名無しさん mailto:sage [2008/12/13(土) 14:55:41 ]: nvccから呼ばれるほうのclのことね＞32bit版
887 名前：デフォルトの名無しさん mailto:sage [2008/12/14(日) 01:21:53 ]: >>882
パスが通ってないんじゃね？
888 名前：879 [2008/12/17(水) 18:34:27 ]: PATHは通っていると思います。消してみると、
　　nvcc fatal : Cannot find compiler 'cl.exe' in PATH
となりました。platformSDKのclにすると↓のようになりました。
　　nvcc fatal : nvcc cannot find a supported cl version. Only MSVC 7.1 and MSVC 8.0 are supported
結局、CUDAtoolkitを32bitにしましたところOKとなりました。885さん、886さん、887さんありがとうございました。
889 名前：デフォルトの名無しさん mailto:sage [2008/12/17(水) 20:02:14 ]: 最適化してregisterの数を減らしたいんですが、
なんか心がけることとかコツってありますか？
890 名前：デフォルトの名無しさん mailto:sage [2008/12/17(水) 23:45:04 ]: ptx出力を読め。以上。
891 名前：デフォルトの名無しさん mailto:sage [2008/12/18(木) 00:41:24 ]: なんでわからないのでご教授を・・・・
VS 2005にカスタムウィザードを組み込んでパスも通して空の.cuファイルに書き込んだのに
1>LINK : fatal error LNK1181: 入力ファイル '.\Debug\sample.obj' を開けません。
と出るのですが・・・・何か環境設定間違っているのでしょうか？
892 名前：デフォルトの名無しさん mailto:sage [2008/12/18(木) 00:49:51 ]: sourceforge.net/projects/cudavswizard
どうぞ
893 名前：デフォルトの名無しさん mailto:sage [2008/12/18(木) 01:47:35 ]: >>892
それ使ってるんですが・・・readmeを嫁ということですね
894 名前：デフォルトの名無しさん mailto:sage [2008/12/18(木) 03:39:45 ]: ptx出力がどうなるように最適化していけばいいんですか？
895 名前：デフォルトの名無しさん mailto:sage [2008/12/18(木) 07:42:38 ]: >>894
>889なら、registerの数を減らしたいんだろ。ptxの出力を読んで少なくなっていればOKだw
896 名前：デフォルトの名無しさん mailto:sage [2008/12/20(土) 00:21:41 ]: >>862
NVIDIA、モバイルGPUドライバを配布開始
～ノートでPhysXなどCUDAアプリケーションが利用可能に

12月18日(現地時間)配布開始

　米NVIDIAは18日(現地時間)、モバイルGPU用の汎用ドライバを同社ホームページ上で配布開始した。
対応GPUはGeForce 8M/9M、およびQuadro NVS 100M/300Mシリーズ。

　従来、NVIDIA製モバイルGPUのドライバは、ノートPC本体のメーカーが提供していたが、
今回よりNVIDIAが汎用ドライバを提供することになった。対応OSはWindows XP/Vista。

　バージョンはデスクトップ版の180代に近い179.28 BETAで、CUDAに対応。これにより、PhysXや、
CUDA対応の動画編集ソフト、Photoshop CS4などでアクセラレーションが効くようになる。
ただし、PhysXについてはビデオメモリが256MB以上必要となる。

　なお、ソニーVAIO、レノボThinkPad、デルVostro/Latitudeシリーズ、およびHybrid SLI構成の
製品については対象外となっている。

　現バージョンはベータ版で、WHQL準拠の正式版は2009年初頭の公開を予定している。

pc.watch.impress.co.jp/docs/2008/1219/nvidia2.htm
897 名前：デフォルトの名無しさん [2008/12/21(日) 00:11:21 ]: なんでvaioは対象外ですか？いじめですか？？
898 名前：デフォルトの名無しさん mailto:sage [2008/12/21(日) 07:03:38 ]: vaioとか一部のノートはリコールされてるから
899 名前：デフォルトの名無しさん mailto:sage [2008/12/21(日) 19:00:56 ]: それ以前にノートだとメーカーが色々やってるから
一般的にドライバ更新さえノートのメーカー対応待ちだろ。
900 名前：デフォルトの名無しさん mailto:sage [2008/12/21(日) 22:30:31 ]: シェーダプロセッサってエラーがあってもゲームでは無視されて
認識出来ないレベルで画面が崩れるだけだから
一見正常に動いてると思っててもＣＵＤＡは動かない人が多いかも
ドライバが対応してるとか以前にハードが用件を満たしてない
901 名前：デフォルトの名無しさん mailto:sage [2008/12/24(水) 10:44:05 ]: どこでエラーになっているのかデバッグしろよ
902 名前：デフォルトの名無しさん [2008/12/25(木) 15:53:55 ]: ちょいちょいCUDAをいじり始めてるんですが、OpenCLの規格が決まったことで、
CUDAの仕様も大幅に変わってしまうんですかね。
903 名前：デフォルトの名無しさん mailto:sage [2008/12/25(木) 18:15:24 ]: MacBookProなんかだと自前でドライバ突っ込めばCUDA使えたんだけど，
VaioやThinkPadだとどうですか？
使っている人がいれば情報求む．
904 名前：デフォルトの名無しさん [2008/12/25(木) 18:21:26 ]: 5秒制限の条件とか回避の仕方とかがよくわからないのですが
教えてもらえないでしょうか？
905 名前：デフォルトの名無しさん mailto:sage [2008/12/26(金) 01:44:56 ]: ThinkPad T61はOpenGL関係以外は使えた。
906 名前：デフォルトの名無しさん mailto:sage [2008/12/26(金) 03:48:04 ]: OpenCLはAPIのみだからCUDAとは別にライブラリが用意されるだけだと思うよ
こっちは専用コンパイラも必要ない
907 名前：デフォルトの名無しさん mailto:sage [2008/12/26(金) 14:22:18 ]: >>904
Gridを何回も起動するしかないのでは。だから、途中結果をメモリにのこしとかないといけない。
面倒だよね。
908 名前：デフォルトの名無しさん mailto:sage [2008/12/26(金) 14:35:42 ]: >>904
Xを使わない、グラフィック機能を使わなければいけるはず。

…何のためのグラフィックカードやねんって感じだが
909 名前：デフォルトの名無しさん [2008/12/26(金) 17:45:27 ]: >>907 >>908
やっぱりそれぐらいしかないですか。
現状でグラフィック機能止めて使うのは無理なので
Gridの分割を考えてみます。

返答ありがとうございます。
910 名前：デフォルトの名無しさん mailto:sage [2008/12/26(金) 18:15:35 ]: Vistaだとレジストリ弄ればもうちょっと長くできるよーみたいなこと言ってたような、言ってなかったような
911 名前：デフォルトの名無しさん mailto:sage [2008/12/26(金) 22:25:09 ]: 使うグリッド数を制限すればなんとかならなくもないんだが、効率は落ちるんだよね。
つーか、効率重視ならもう一枚挿すしかない希ガス。
912 名前：,,・´∀｀・,,）っ-○◎● mailto:sage [2008/12/27(土) 00:47:07 ]: こんなの教えて貰った
compview.titech.ac.jp/Members/endot/adv-app-hpc/2008schedule
913 名前：デフォルトの名無しさん mailto:sage [2008/12/27(土) 05:05:39 ]: どのみちCUDA実行中は画面の描画が止まるから割り込み入れる工夫は必要でしょ
914 名前：デフォルトの名無しさん mailto:sage [2008/12/27(土) 08:47:20 ]: 描画要求の方がGridよりも優先されているようにしか見えないんだけど。
止まってくれたらどんだけ効率が改善することかw
915 名前：デフォルトの名無しさん mailto:sage [2008/12/27(土) 11:52:08 ]: >>912
面白そうだけど、東工大は遠いなぁ。ワープアSEに新幹線は高い。
わがまま言うと、神戸でやって欲しい。
916 名前：デフォルトの名無しさん mailto:sage [2008/12/27(土) 14:13:48 ]: 大岡山の東工大なら歩いて行けるぜ！
917 名前：,,・´∀｀・,,）っ-○◎● mailto:sage [2008/12/27(土) 23:27:30 ]: >>915
六甲台か？
あの急勾配の坂は登りたくないな。
918 名前：デフォルトの名無しさん mailto:sage [2008/12/29(月) 05:06:23 ]: >>912
これは紹介レベルだから、たいした内容ではないよ
919 名前：デフォルトの名無しさん mailto:sage [2008/12/29(月) 11:44:34 ]: CUDAをクアッドコアで動くように出来ればいいのに
920 名前：デフォルトの名無しさん mailto:sage [2008/12/29(月) 12:15:26 ]: >>919
どういう意味?

CUDAをQuadCoreCPUで動かすことなら何の問題もなくできるのだけど。
921 名前：デフォルトの名無しさん mailto:sage [2008/12/29(月) 12:44:32 ]: っOpenCL
922 名前：デフォルトの名無しさん [2009/01/05(月) 20:09:11 ]: VIPPERの諸君　立ち上がれ！
以前韓国を打ち破ったのは誰か知らない、小学生や、他サイトの管理人がVIPを馬鹿にしている！

全員集結せよ。そして全員で打ち倒すのだ！

もし時間がない人は、この文章をそのまま他のスレに貼ってほしい！

対策本部
takeshima.2ch.net/test/read.cgi/news4vip/1231149782/

今こそ　Vipperの意地を見せつけるのだ！
923 名前：デフォルトの名無しさん [2009/01/08(木) 17:49:55 ]: nc = 100*100
bs = 50
dim3 dimBlock(bs,bs)
dim3 dimGrid(sqrt(nc)/dimBlock.x,sqrt(nc)/dimBlock.y)
kernel<<<dimGrid, dimBlock>>>(idata, odata, sqrt(nc)

__global__ void kernel(float* idata, float*, odata, int nc)
{
index=blockIdx.x * blockDim.x + threadIdx.x +
+(blockIdx.y * blockDim.y + threadIdx.y) * nc
}
この時のイメージは、Gird:2x2、Block:50x50でよいのでしょうか?
それとこのままグローバルメモリで計算するのはできるのですが、一旦
シェアードメモリに退避して計算してグローバルメモリに戻す方法が
サンプルを見てもうまくいきません。どういう感じになるのでしょうか?
924 名前：デフォルトの名無しさん [2009/01/08(木) 18:11:37 ]: いつも疑問に思いつつ使っているのですがグローバルメモリ使うときも
共有メモリ使うときも
float* data;
CUDA_SAFE_CALL(cudaMalloc((void**)&data, sizeof(float) * 100)
CUDA_SAFE_CALL(cudaMemcpy(data1, h_data, sizeof(float) * 100, cudaMemcpyHostToDevice))
になるのでしょうか?これはグローバルメモリにとっているのですよね。
共有メモリの話は、kernelのほうで使用するかどうかですよね?

コンスタントメモリ使うときだけどうも
CUDA_SAFE_CALL(cudaMemcpyToSymbol(data, h_data, sizeof(float) * 100))
としてホストからコピーするんですかね?
925 名前：デフォルトの名無しさん mailto:sage [2009/01/08(木) 23:27:14 ]: >>924
cudaMemcpyToSymbolは定数メモリの他、グローバルメモリでも使えるよ。
要は、device空間で定義されたメモリと言うニュアンスでしかない。
逆に言えば、cudaMemcpyToSymbolを使わないとhost側のメモリ空間に割り付けられたアドレスを利用しようとしてしまう。
# この辺が判り難いのよね。

共有メモリはグリッド実行中しか使えないからhost側からは為す術がない。

>>923
どのサンプル見てどうやったらどう巧くいかないのか具体的に。
少なくとも、>923のような処理なら共有メモリの使いようがないんだけど。
926 名前：デフォルトの名無しさん mailto:sage [2009/01/09(金) 01:50:13 ]: warp単位32スレッドで同じ命令を行って、
実際に一度に動くのは8スレッド

こういう考え方でいいのですか？
927 名前：デフォルトの名無しさん [2009/01/09(金) 18:21:54 ]: >>925
結局現状では、cudaMemcpyXXXの違いは理解できそうにないです。
まあそれはよいとして(全然良くないですが)
>>923
のidataは、実は、一次元配列なんです。なのにカーネルには2次元で利用しよう
としているんです。
float host_idata[nc]
float* idata
cudaMalloc((void**) &idata, sizeof(float) * nc)
cudaMemcpy(idata, host_idata, sizeof(float) * nc, cudaMemcpyHostToDevice)
こんな感じのグローバル変数がidata。
としたとき、ncとbsの設定値によってはうまく処理できないことがあるのです。
これが謎です。わかる方いませんか?もちろんncはsqrtで整数になる値、設定する
値は、割り切れる値です。nc=1000*1000,bs=500で処理結果がおかしくなっています。
kernelの読んだあとエラーでinvalid configuration argument.になってます。

それと、共有メモリのほうですがこれは分かりました。
kernel<<<dimGrid, dimBlock, sizeof(float) * nc>>>(idata, odata, sqrt(nc))
第3引数の共有メモリのサイズを定義しないとだめだった。ちゃんと説明書を
読んでいなかったというレベルでした。これは解決です。
928 名前：デフォルトの名無しさん [2009/01/09(金) 19:05:25 ]: >>927
ブロックは512x512x64と出ているけどいちお3次元は分からないので512x512
としてxで512使い切るとどうもyではもう1しかとれないようです。
ncは、そのままだとエラーが出ないレベルは、bs=20つまり20x20=400ブロック
使用。21,22は割り切れないので設定できない。みたいです。
とすると一次元配列を使用するのにわざわざ、x,yの次元のブロックを使いきる
テクニックは必要なさそうですね。と推察できると思います。結果からですが
ただ、これがバグであるなら次のCUDAのバージョンアップで使えるようになる
かもですね。
929 名前：デフォルトの名無しさん mailto:sage [2009/01/09(金) 20:53:02 ]: >>928
次元とは別に1ブロックあたり512スレッドが上限じゃなかったっけ？
どっかにきちんと書いてあったはずだからバグではないよ。
930 名前：デフォルトの名無しさん mailto:sage [2009/01/09(金) 22:23:41 ]: 私は1次元でやってしまうことが多いなぁ。
で、スレッド数についてはdeviceQuery参照で。
931 名前：デフォルトの名無しさん mailto:sage [2009/01/10(土) 21:32:30 ]: スレッドの最大数はAPIで取得できる。全部512だと思う。
ブロックの最大は65535だから65535*512以上は1次元のグリッドでは扱えない。
あと最大は512だけど、192か256くらいのほうがパフォーマンスが出る。
932 名前：デフォルトの名無しさん mailto:sage [2009/01/14(水) 01:57:36 ]: CUDA 2.1 Release
ttp://forums.nvidia.com/index.php?showtopic=85832
933 名前：デフォルトの名無しさん [2009/01/14(水) 21:34:37 ]: 情報サンクス。早速、Linux版ゲットした。
934 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 01:08:35 ]: 9600GTでVideoDecodeのサンプル動かそうとしたら正常に動かなくて、
結局SDK2.0+178.24まで戻さないとダメだった。

2.1にもソースは付いてるけど、2.0のままみたいだね。
935 名前：,,・´∀｀・,,）っ-●◎○ mailto:!sage [2009/01/15(木) 07:55:42 ]: >>931
1 warp = 32 threadで、GeForce 8/9が24warp/block、GT200で32warp/blockが最大だから
768thread/blockと1024thread/blockが最大なんだけどね。本来は。
CUDAのドライバ側が512でリミッタかけてるんだ。罠としか言いようが無い。
逆に言うとCUDAを経由しなきゃ目いっぱい使えるかもね。

ただ、スレッドインターリーブすると１スレッドあたりで使えるレジスタ本数が減っちゃうんだよね。
メモリレイテンシを隠蔽するならスレッドを目いっぱい使ったほうがいいし
逆に一時変数を何度も再利用する場合は、thread/blockを減らして１スレッドあたりの仕事量を増やしたほうがいい。
936 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 15:42:19 ]: CUDAってシーユーディーエーって読んでいいの？
それともクダとか、なんか読み方あったりする？
937 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 15:57:04 ]: >>936
クーダ
938 名前：,,・´∀｀・,,）っ-●◎○ mailto:sage [2009/01/15(木) 15:58:34 ]: ネイティブの人は「クーダ」って発音してたよ。
Cellは「くた」だから注意な。
939 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 16:01:48 ]: ありがとう！明日から会社でクゥゥゥウーダ！！って叫びまくるわ
940 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 16:02:44 ]: きゅーだ　って発音してる俺は異端ですかね？
941 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 16:23:46 ]: >940
そんな人いたんだ。
942 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 16:29:33 ]: 英語的に　クー　よりも　キュー　だと思わね？
943 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 16:56:19 ]: >>941
異端だな
944 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 16:57:46 ]: どっちでもいいんじゃね

ちなみにnudeはニュードとヌードの両方の発音が有
945 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 17:13:01 ]: キュもクも似てるし、まあ仮に間違って言ったとしても気づかれないんじゃない？
946 名前：デフォルトの名無しさん [2009/01/15(木) 19:38:28 ]: ホムペにはクーダと嫁
と書かれてたはず
947 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 22:24:41 ]: "The cuda and my wife"
948 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 23:40:36 ]: たしか英語版のfaqには
kuh-da
と書いてあったはず
949 名前：デフォルトの名無しさん mailto:sage [2009/01/15(木) 23:53:57 ]: 2.1はJITサポートがミソ？
950 名前：デフォルトの名無しさん [2009/01/16(金) 14:10:31 ]: 2.1をつついた人に質問
2.1ではカーネルに渡す引数でポインタ配列は使えますか?
951 名前：デフォルトの名無しさん mailto:sage [2009/01/16(金) 18:48:04 ]: 共有メモリからグローバルメモリに連続領域をコピーするにはどうすればいいでしょうか
カーネルから呼び出せるmemcpyみたいなのがあれば教えてください
952 名前：デフォルトの名無しさん [2009/01/16(金) 18:58:33 ]: cudaMemcpyが長いようで、the launch timed out and was terminated.が
でてしまいす。これなんとかするほうほうありましたっけ?
953 名前：デフォルトの名無しさん mailto:sage [2009/01/16(金) 21:53:12 ]: >>952
転送量を減らす。

>>951
cudaMemcpy
954 名前：デフォルトの名無しさん mailto:sage [2009/01/17(土) 00:16:47 ]: >>951
スレッド内で連続的にメモリにアクセスするようなコードは望ましくない。
coalscedにアクセスするとなると、自分でコードを書くしかない。
共有メモリから転送を始める前に、同期を取ることも忘れずに。

>>952
参考までに、どんなハードウェアの組み合わせでどれくらい転送しようとしている?
8800GT辺りだと転送できるだけ全部転送してもタイムアウトしないのだけれど。
955 名前：デフォルトの名無しさん [2009/01/17(土) 02:52:44 ]: >>954
カーネル実行後の結果データコピーでエラーになっているのですが
sizeof(float)*1000バイトを100回連続でコピーしています。
100個とも別のアドレスに結果をコピーしたいためそうしてます。やはり
いっぺんにメインメモリにいったんコピーしてからやったほうがいいので
しょうか?
956 名前：デフォルトの名無しさん [2009/01/17(土) 03:03:49 ]: >>955
sizeof(float)*10000バイトを100回連続コピーでした。
957 名前：デフォルトの名無しさん [2009/01/17(土) 07:16:40 ]: >>954
＞ >>952
＞どんなハードウェアの組み合わせ
958 名前：デフォルトの名無しさん mailto:sage [2009/01/17(土) 07:19:52 ]: >>956
同一データ部分を100箇所にコピー？
独立の100箇所をコピー？
959 名前：デフォルトの名無しさん mailto:sage [2009/01/17(土) 07:55:29 ]: >>956
それ、本当にメモリのコピーでタイムアウトしてる？
カーネル実行後に同期取らずにすぐにメモリのコピー始めて
100回コピーのどこかでカーネル実行がタイムアウトしてるとか。

cudaMemcpyの戻り値は
Note that this function may also return error codes from previous, asynchronous launches.
ということだし。
960 名前：954 mailto:sage [2009/01/17(土) 08:44:25 ]: >>955=956
えーと、都合数MB程度ってことか。それだったら>959の言うようにカーネル関数のタイムアウトじゃないかな。
DEBUGつきでCUT_SAFE_CALL(間違ってたら失敬)を使うか、自前でcudaThreadsSynchronize()を呼んで戻り値を見てみて。
カーネル関数起動後は、復帰を待たずにホストの処理に戻るからカーネル関数自体のエラーはカーネル関数起動では捕捉できないのよ。
で、まさかとは思うけど同一データを100回コピーならバス越しにやってはダメ。別のデータならホスト側の用件次第かな。
# でも多分、一旦ホスト側に転送してから分配した方が速い気がする。CPUとメモリ周りがよっぽど遅くなければ。

>>957
下手な突っ込みはお郷が知れるよ。この場合、CPUが遅いんじゃないかとかバスがx1なんじゃないかとかGPUがモニタ表示か何かで
よっぽど負荷が掛かっているんじゃないかとかGPU側メモリが巨大なSTELAなのかとか、組み合わせ次第で色々条件が変わってくるのよ。
961 名前：デフォルトの名無しさん [2009/01/17(土) 09:52:40 ]: >>960
＞下手な突っ込みはお郷が知れるよ。

質問者が答えてないのを指摘しただけなんだが…
誤解させてすまんかった
962 名前：デフォルトの名無しさん mailto:sage [2009/01/17(土) 11:09:08 ]: そういう意味か。あの書き方じゃぁ、誤解されるよ。
まぁ、大量転送じゃないと言う意味で、今回は関係なさそうだね。
963 名前：デフォルトの名無しさん [2009/01/17(土) 15:38:06 ]: >>952の話
>>957
ATIの統合チップセットでオンボードのATIビデオカードを無効
PCI-E 1.0 x16バスにGeforce9600GT 512MB
CPUはAMDデュアルプロセッサ(結構遅いやつ)
964 名前：デフォルトの名無しさん [2009/01/17(土) 15:55:29 ]: >>952の話
>>959
カーネル実行
a:cudaのエラーチェック(カーネルのエラーのありなし)
cudaThreadsSynchronize
for(100回
結果データをデバイスからホストにコピー
b:cudaのエラーチェック(直前のcudaMemcpyのエラーのありなし)
)
bのとこでthe launch timed out and was terminatedが出てるんです。

先週は時間がきてあきらめて帰ったのでここまでです。
言われたようにcudaThreadsSynchronizeの戻り値見てみたほうがよさそう
ですね。
965 名前：デフォルトの名無しさん mailto:sage [2009/01/17(土) 18:45:09 ]: >>951
共有メモリはカーネル起動時に動的に確保されるメモリ領域だから
カーネルが終了したら消えるし、１つの実行スレッドから全部コピーなんて論理的におかしいし
共有メモリはあくまで高速動作と同期が取れる作業用のメモリとして考えたほうがいい
共有メモリを使って計算した結果は１個づつ普通のグローバルメモリに入れてやるのが正しいやり方

>>952
10000を一度に転送して実行しても
1を10000回繰り返して転送しても
実行時間は大差ないんですよ
CUDAで実行する部分は出来るだけコンパクトにまとめて
呼び出す本体の方で特殊な演算関数くらいの感覚でループさせて使うのが正解
966 名前：デフォルトの名無しさん mailto:sage [2009/01/18(日) 23:59:40 ]: JCublasについて
おしえてエロイ人
967 名前：デフォルトの名無しさん [2009/01/19(月) 11:04:50 ]: >>925の話
cudaThreadsSynchronizeの戻り値をチェックしたら
the launch timed out and was terminatedが出ていました。結果コピーで
エラーで落ちてたのだけどエラーデータは、その前に実行していた
cudaThreadsSynchronizeの問題だったようです。

(cudaThreadsSynchronizeが正常になるまで待つとしても配列を
100x1000回、100x10000回、回すと1分待っても同期とれないようです。)

つまりカーネルに同期できない処理が書かれていたのが原因だと思います。
NVIDIA_CUDA_Programming_Guide_2.0.pdfのp34-35のようなコードで
かつそこのカーネルのdevPtrを参照するだけでなく書き戻す処理をや
っているのでそれがだめなのだと思います。恐らくこういうピッチ処理
の場合は、参照のみが許されるのでは?と思います。
問題のコードをこのpdfの変数に直して下記に書いておきます。
__global__ void myKernel(float* devPtr, int pitch)
{
for (int r = 0; r < height - 1; r++) {
float* row1 = (float*)((char*)devPtr + r * pitch);
float* row2 = (float*)((char*)devPtr + (r + 1) * pitch);
for (int c = 1; c < width - 1; c++) {
row2[c] = row1[c]とrow1[c + 1]とrow1[c -1]やらの計算
}
}
}

このp34-35を見て思うのはindex使わずにこんな処理書いてGPUでやる意味
あるの?と思いますが。自分で書いてなんなのですが
968 名前：デフォルトの名無しさん mailto:sage [2009/01/19(月) 17:28:20 ]: お前はあほかｗ
969 名前：デフォルトの名無しさん mailto:sage [2009/01/19(月) 17:49:31 ]: タイムアウトになる原因はそのでかいループのせい
せいぜいミリ秒単位でタイムアウトを判断してるから
ミリ秒以下で応答しないようなカーネルは絶対に書いては駄目

cudaThreadsSynchronizeは発行したカーネルがすべて終了するのを待つだけ

グローバルメモリは読み書きは出来るが前後は保障されないので
１スレッドが書き込みする箇所は限定する必要がある

共有メモリを使って他のスレッドが書き込んだ内容を参照して利用する場合に
カーネル内部で___syncthreadを使う
これが本来の同期の意味
970 名前：デフォルトの名無しさん [2009/01/19(月) 19:44:51 ]: >>952の話
>>968のように言われるのは分かって書いてみたんだけど
NVIDIA_CUDA_Programming_Guide_2.0.pdfのp34-35のようなコードが
マニュアルに書いてあるのがおかしいと思う。
__global__ void myKernel(float* devPtr, int pitch){}
そもそもこんな書き方じたいが書けるけど間違えな使い方だと。
この書き方しているとこにやらないようにこの部分に×印つけてほしい。

あとはコンパイラがえらくなったらfor多重ループをうまく処理する
アセンブラぐらい作ってほしい。(OpenMPでパラレルfor指定すると#のタグだけ
でプロセッサ使ってきって高速化なるんだけどね。そこまでぐらい将来
的にはしてほしい。)
971 名前：デフォルトの名無しさん mailto:sage [2009/01/19(月) 20:53:12 ]: 文句言っても何にもならない。
そう悟る時がくるまで気長に待ちましょうよ。
そうすれば成長しまっせ。
972 名前：デフォルトの名無しさん mailto:sage [2009/01/19(月) 23:38:15 ]: それよりも先ず、日本語を何とかしてくれ。
＞間違えな使い方
＞しているとこにやらないように
＞プロセッサ使ってきって
973 名前：デフォルトの名無しさん mailto:sage [2009/01/19(月) 23:41:25 ]: そのサンプルはあくまでもcudaMallocPitch()がどういう風にメモリを確保するかの説明だからなぁ。
きっとライターは、まさかそのまま動かそうとされるは思っちゃいまいよ。
974 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 02:02:49 ]: ところでドライバ180なんだが2Dモードだかなんだかで消費電力抑える為に
ドライバが自動的にクロックを下げて3Dゲーム動作させると上がる仕様なんだけど
CUDA実行しても上がらないぞｗ
975 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 02:08:30 ]: ドライバを替えてくださいw
つーか、それって169.09のときに指摘されて直したんじゃなかったのか?＞NVIDIA
976 名前：デフォルトの名無しさん [2009/01/20(火) 03:40:32 ]: d_a[0][blockIdx.x][blockIdx.y][threadIdx.x]
って使おうとしたら"expression must have arithmetic or enum type"
ってでたんだけど何がいけないんすか？
教えてください。
977 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 07:58:55 ]: >>976
d_aがポインタなら、途中で解決できない状況があるのかも。
978 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 08:41:39 ]: 共有メモリで構造体を使いたいのだけど
例えば
struct data {
int a;
char b;
}

func<<<dim3(), dim3(), SIZE * sizeof(data)>>>();

__global__ kernel(){
__shared__ data d[SIZE];
}

こんな感じでやるとCUDAとC側のパック構造が違うせいだと思うけどおかしくなる
どうやれば出来るの？
979 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 08:59:56 ]: パックが違うのだけが理由なら、
struct data {int a; char b; char dummy[3];}とでもしておけばいいじゃん。
980 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 13:53:14 ]: ただでさえ少ない共有メモリをそんな無駄に使えない
981 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 20:44:37 ]: >>977
世の中夜勤帰りで朝から寝てる人だっているんだよ？
引っ越しの時ちゃんと挨拶行った？
顔合わせたら軽く会話するとかしてちゃんとコンタクト取り続けてる？
日頃からそういうコミニュケーションが取れてればいつ洗濯機を回していいのか
いつ静かにしなければならないのか
迷惑を掛けないように生活出来るはずなんだが
982 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 23:36:58 ]: >>980
マジで言っているのなら、設計が悪い。
どうしてもパディングしたくないくらい逼迫しているなら、int一個を切り分けて3バイトと1バイトで使え。
983 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 00:22:25 ]: 共有メモリが制限されてるのに無駄な領域作って
ほとんどをグローバルメモリに追いやる方がよっぽど設計が悪いでしょｗ
984 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 00:25:51 ]: GPUのメモリレイテンシって１２とかの世界だぞ
CPU用のDDR2で５だからな
intを内部でcharとして扱うプロセスを考慮しても共有メモリ使った方が早いんだよ
985 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 00:35:27 ]: CUDAで共有メモリを使うこと自体、無駄な同期待ちが発生するから避けたいところだと思うが。
coaxschedな読み書きができるなら、共有メモリより遅くないぞw
986 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 02:24:30 ]: >>975
(エンバグの)歴史は繰り返す。
987 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 12:18:47 ]: >>CUDAで共有メモリを使うこと自体、無駄な同期待ちが発生するから避けたいところだと思うが。
え？
共有メモリってCPUで言う所のただのレジスタみたいなもんで同期は指定しないと発生しないと思うけど

ところでローカル変数をすべて共有メモリに置いた場合のが早いんだけど
ローカル変数ってデバイス上のメモリに確保されるだけで共有メモリより遅いレベルのメモリを使ってるっぽいね
988 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 18:11:15 ]: あー、表現が悪かった。共有メモリを共有目的に使うと同期を取る必要が出てくるから同期待ちが発生するということ。
レジスタみたいに使うのなら確かに関係なかったね。

で、レジスタよりも速いかどうかについてはptx見てみたいところ。
989 名前：,,・´∀｀・,,）っ-○◎● mailto:sage [2009/01/21(水) 18:18:05 ]: 見た感じリードレイテンシはこれくらい

レジスタ＞＞Const Memory＞Shared Memory＞＞＞＞DRAM

[ 新着レスの取得/表示 (agate) ] / [ 携帯版 ]

read.cgi ver5.27 [feat.BBS2 +1.6] / e.0.2 (02/09/03) / eucaly.net products.
担当:undef