【GPGPU】NVIDIA CUDA質問スレッド

1 名前：デフォルトの名無しさん mailto:sage [2007/09/17(月) 14:54:28 ] 本家 developer.nvidia.com/object/cuda.html 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