【GPGPU】くだすれCUDAスレ【NVIDIA】
at TECH
[前50を表示]
600:デフォルトの名無しさん
09/04/22 16:25:47
ﾓﾅｰ
601:デフォルトの名無しさん
09/04/22 17:03:27
.NETはMONOで動くが、GCCで.NETアプリってコンパイルできたっけ？
602:,,・´∀｀・,,）っ-○◎○
09/04/22 22:19:21
CUDAかC#かって、ベクトルが全然別ですがな
>>601
Mono入れたらmcsってコンパイラが使えるようになるはずだが。。。
貴殿はGCCに入ってないという理由でPerlやPHPをも嫌うのですか？
603:デフォルトの名無しさん
09/04/22 22:45:40
問題はなぜこのスレは勢いがないのかってことだ
604:デフォルトの名無しさん
09/04/22 23:07:34
CUDAとOpenCLの認識の仕方として、
抽象レイヤ的にこんな感じかな？？
APP
--------
C/C++
--------
OpenCL
--------
CUDA
--------
driver and runtime
605:,,・´∀｀・,,）っ-○◎○
09/04/23 06:04:45
>>603
見た目簡単そうに見えて実は使いづらくて、本質はCellよりも更に応用分野は厳しいからね。
「CPUの数十倍とか言ってたけど全然遅いじゃん！」で、使い方を理解しないままみんな匙j投げる
いや、使い方がわかったところで、その正しい使い方が、本質的に目的のアプリケーション向きじゃなかったり。
606:デフォルトの名無しさん
09/04/23 07:22:35
そうそう、その演算だけに絞れば確かに速いんだけど、アプリケーション全体で見るとXeonに勝てなかったりね。
ボードメーカ側も自覚しているらしく、私の客先でのCUDA開発は2チップGPUボード4枚挿しするところまでいってしまっているし。
607:デフォルトの名無しさん
09/04/23 09:11:40
どうせララビーも期待外れに終るさ
608:デフォルトの名無しさん
09/04/23 16:09:19
nv社員乙ｗ
609:デフォルトの名無しさん
09/04/23 20:57:12
Larrabeeは、たかがx86、されどx86だな
Atomに毛が生えたような小規模なx86コアが数十コアあったら何が出来る？
汎用プロセッサとしては程度が知れてる分、逆に落胆しようがない。
良くも悪くも身の丈以上の期待はされてないからな。
610:デフォルトの名無しさん
09/04/23 21:20:24
流れをぶった切るが
GeforceはCELLより変態的な構造って認識でおk？
611:,,・´∀｀・,,）っ-○◎○
09/04/23 21:43:53
餅は餅屋
612:デフォルトの名無しさん
09/04/23 23:47:54
色々調べて見たけど結局CUDAのsuper piはまだ出てないんだな
CPUとGPUの比較が出来ると思ったのに
613:デフォルトの名無しさん
09/04/24 05:23:51
むしろスーチーパイがもっとリアルにぬるぬる動けば…
614:,,・´∀｀・,,）っ-○○○
09/04/24 06:37:53
＞スッチーのπ
まで読んだ
とりあえずPTXの自己コンパイルは最低限だろ
どっかの営業さんが言うにさ
「たとえCPUより速くなくとも、CPUでやってる仕事を肩代わりしてやることが
出来るだけでも使う価値があるんじゃないでしょうか」
いや、それのお膳立てのためにCPU時間食うから本末転倒なのよ。
615:デフォルトの名無しさん
09/04/24 06:44:16
＞「たとえCPUより速くなくとも、CPUでやってる仕事を肩代わりしてやることが
＞出来るだけでも使う価値があるんじゃないでしょうか」
そういうことを臆面もなく語る営業マンを一人知っているんだけどw
同一人物と考えてよさそうだな。
616:デフォルトの名無しさん
09/04/25 00:58:56
肩代わりしてやるなんて大それたことを無理に言い張るから、おかしくなるんだよね。
重要なのは、ＣＰＵとＧＰＵとが各々の得意分野を担当し、住み分けをすることだろう。
ＧＰＵは汎用計算に向いていないのだから無理にＧＰＵを使わずＣＰＵを使えばいいし、
３Ｄゲームや科学技術計算などＧＰＵの方が効率的な計算でＧＰＵを使えばいいんだよ。
いわゆるアインシュタインとタイピストの喩えだ。
アインシュタインが優れた物理学論文を清書してもらうためにタイピストを雇ったら、
なんとまあそのタイピストよりアインシュタインの方がタイプが速かったとしよう。
じゃあ、そのタイピストを解雇すべきか？答えは否だ。タイプはタイピストに任せ、
アインシュタインは少しでも長い時間、優れた物理理論を考え出すことに費やすべきだ。
それが最も効率がいい。
617:デフォルトの名無しさん
09/04/25 01:48:13
>>615
営業ならだいたい同じこと言うんじゃねぇか？
618:デフォルトの名無しさん
09/04/25 03:57:19
>>616
>答えは否だ。
いや、答えは科研費の額によるだろうｗ
619:デフォルトの名無しさん
09/04/25 05:08:45
>>616
タイピストに指示だすのに、タイプするのと同じような時間がかかるから問題なんじゃね？
620:,,・´∀｀・,,）っ-○○○
09/04/25 06:42:29
清書する段階で更に考えても無駄だろう
621:デフォルトの名無しさん
09/04/25 15:20:06
CPUを管理職、GPUを部下に例えてみよう。
CPUからGPUへの指示の中身が足りなかったりすると
CPU-GPU間のやりとりが増えてしまい遅くなる。
CPUから指示する内容がGPUの能力を超えると
なかなか結果が返ってこない。
逆にCPUの能力が低いとGPUへの指示や対応が遅くなる。
GPUの仕事に信用がおけないとCPU側でのチェックが
必要となり負荷となる。
GPUはCPUほど守備範囲は広くないし経験も少ない。
よいCPUやGPUを入手するには予算が必要である。
またCPU、GPUを動かし続けるには経費がかかる。（電気代、冷却設備）
あまり負荷をかけるとうるさくなったり、たまに壊れたりする。
overclockによる故障は保証の範囲外であることに注意。
622:デフォルトの名無しさん
09/04/25 15:22:15
GPUのIPコアが強化されれば良いんだが。
それをしようとして、intelに待ったを掛けられたんかな。
623:デフォルトの名無しさん
09/04/25 16:40:02
チップセットのバスライセンスと何の関係が？
624:デフォルトの名無しさん
09/04/25 18:21:57
たとえ話にすると細部の理解が必要ないから生半可な知識でも初心者が騙せて優越感に浸れてうめぇw
っていつも思う。
625:デフォルトの名無しさん
09/04/25 21:11:34
CPUとGPUは、お互いに交わる方向で
じきに差は無くなるんだろ
GPUいらねでおわりじゃねーの
626:デフォルトの名無しさん
09/04/25 21:13:38
昔GPGPUスレに書いたが、別のパラメータで同じコンテキストの処理をするようなときに
高速に処理できるのがGPUの利点。
別々のコンテキストが必要なら、丁度良いマルチプロセシングの環境を探しなさい。
どちらも歩み寄っているようだけど、ゲーム屋的には、現行世代機的なトランジスタ
バランスのマシンが次世代にも欲しいところ。
627:626
09/04/25 21:15:48
>>625
似たようなコストで作れるならな。
見当付いてるなら特許とって今すぐ始めるのがいいぞ。
628:,,・´∀｀・,,）っ-○○○
09/04/25 22:56:05
GTX280って28SMじゃなかったか？
629:,,・´∀｀・,,）っ-○○○
09/04/25 22:57:21
260のほうか
280は30か
630:デフォルトの名無しさん
09/04/26 17:16:50
>>619
常識的に考えてそんなことは起こらんだろ。
仮にタイピストに指示だすのにタイプするのと同じくらい時間がかかるなら、
それはこの喩えが適用できないケースだってだけの話だろう。
ＧＰＵに命令だすのにＧＰＵで計算するのと同じくらいＣＰＵ時間がかかるなら、
そりゃＧＰＵを使うのが不適切なケースだってだけのこと。
>>620
いや、清書してもらうのは既に考え出した理論であって、アインシュタインは
タイピストが清書してる間に次の理論を考えるんだよ。
631:デフォルトの名無しさん
09/04/26 17:24:31
たとえ話にすると細部の理解が必要ないから生半可な知識でも初心者が騙せて優越感に浸れてうめぇw
っていつも思う。
632:,,・´∀｀・,,）っ-○◎○
09/04/26 17:34:27
>>631に全面的に同意
633:デフォルトの名無しさん
09/04/26 18:23:11
>>624
>>631
634:デフォルトの名無しさん
09/04/26 23:38:55
>>625
ジョンカーマックが昔言ってたわな。それ。
ま、今は宇宙大好きっ子になっちまったが。
635:デフォルトの名無しさん
09/04/27 01:38:03
>>630
GPUを使うのが不適切なケースばっかりなのが問題。
636:デフォルトの名無しさん
09/04/27 09:13:30
アインシュタインは一人しかいないけど、CPUとGPUがアインシュタインとタイピストのような関係なら、
CPU増やせばいいんじゃね？無理にGPUにしなくても。
637:,,・´∀｀・,,）っ-○○○
09/04/27 21:24:06
なんにも出来ないのになんでも出来ますと宣伝してるから問題なわけで。
638:デフォルトの名無しさん
09/04/28 09:16:56
そらあんた、ドラッグレーサーをそれなりに走らせるためには適当なコースとそこまで運ぶためのトランスポーターと
燃料などの消耗品が必要になりますがな。
>>638
>631
639:デフォルトの名無しさん
09/04/28 12:52:26
>>637 団子
自己紹介、乙。
640:,,・´∀｀・,,）っ-○○○
09/04/28 21:36:19
だんごやさんだよ
だんごせんもんてんだよ
641:デフォルトの名無しさん
09/04/29 15:05:31
AviUtlがCUDAに対応するのを待つか…
642:デフォルトの名無しさん
09/04/30 21:33:03
某フィルタでシェーダで書いたより遥かに遅くて駄目だしされたような
643:デフォルトの名無しさん
09/05/03 00:16:21
今日、CUDASDK入れてみた。CudaSetup-2.1とNVIDIA_SDK10_CUDA_2.10の入れる順番なのか
環境変数でコンパイルエラー、何度かやっているうちになんとか、サンプルが起動できるようになった。
SDKのサンプルはどこにインストールしているんだ アホか C:\に持ってきた。
サンプルへのパスを追加してやっとコンパイル、起動できた。
644:デフォルトの名無しさん
09/05/03 00:24:59
>>643
2.1のサンプルって意味不明なとこおかれるよね。
VistaのUAC対策かと勝手に思ってるけど。
645:デフォルトの名無しさん
09/05/03 06:12:04
NVIDIAは昔っから何でもそう。
ドライバも一旦C:\に展開してからインスコしてくださりやがる。
まぁ、GCCやそれ用のライブラリが、スペースの入ったパスを
嫌うからかも試練。
646:デフォルトの名無しさん
09/05/04 00:48:26
2.0はちゃんとProgram Files配下に置かれてたよ。
647:デフォルトの名無しさん
09/05/04 23:00:36
>>635
ＧＰＵはゲームや科学技術計算では実績をあげているので、
不適切なケースばかりではないだろう。
>>636
アインシュタインという不世出の天才物理学者と
タイピストという（当時は）いくらでもいた職業を
比較しているのが、この喩えの肝なんだよ。
ＣＰＵのコアを１個増やすより、ＧＰＵのＳＰを１個
（１０個でもいい）増やす方が、ずっと簡単だろう。
648:,,・´∀｀・,,）っ-○○○
09/05/04 23:23:48
別にCPUコアなんて年間何億個でも量産できるだろ
649:デフォルトの名無しさん
09/05/05 05:38:26
タイピストが何人もいても意味無いだろｗ
650:デフォルトの名無しさん
09/05/05 22:12:00
こんなコア橋の下に捨てますよ!
651:デフォルトの名無しさん
09/05/06 08:41:31
Compute Capability 1.3 の GeForce って、GTX だけ?
GTS とかはだめなの?
ファンがうるさいのはやだな～
652:デフォルトの名無しさん
09/05/06 12:29:08
GTSは9800シリーズのリネーム
653:デフォルトの名無しさん
09/05/09 23:52:17
初歩的なこと聞くけど、これってグラボ一台でもできるよね？
表示用と別にCUDA専用のグラボって必要？
654:デフォルトの名無しさん
09/05/10 06:05:44
>>653
その程度が分からないと厳しいかと思いますよ。
分からなくても、やってみて駄目だったら買い足すってことで問題ないと思うけど。
655:デフォルトの名無しさん
09/05/10 12:20:36
>>653
Wikipedia項目ﾘﾝｸ
656:デフォルトの名無しさん
09/05/10 18:27:35
>>655
赤くなっている。。。
GTS250持っているからできると思ったんだけど、二台必要なのかな？
GPUGRIDに参加しようかと思ったらドライバ入れてるのにCUDA対応のデバイスが見つかりませんっていわれるし・・・・。
657:デフォルトの名無しさん
09/05/10 20:56:39
1台でも使える。
ただ処理中に画面が固まってOSが制御不能になることがある。
658:デフォルトの名無しさん
09/05/10 21:24:34
>>657
うーん、CUDAがちゃんと入ったかどうか確認する方法ってありますか？
659:デフォルトの名無しさん
09/05/11 15:35:15
PyCUDAなんてあるのか、おもしろそう
660:,,・´∀｀・,,）っ-○◎●
09/05/11 20:57:31
RuCUDAが必要だな
661:デフォルトの名無しさん
09/05/11 21:41:12
>>656
参加したいGPUGRIDがどんなのか分からんが、
もし、倍精度浮動小数点の演算が必要なものなら、
GTX200シリーズじゃないと無理。
ちなみに、モニタがつながっているかPhysX指定がされてないと、
CUDAでデバイス列挙されないぽい。
662:デフォルトの名無しさん
09/05/12 01:29:13
Py損とかルビィとか手続き型スクリプト言語は向かんだろ。
ocamlとか、Earlangとかの関数型言語をGPGPU対応に
した方が御利益は大きいんじゃね?
並列計算の場合、副作用とか、計算の依存関係が有ると
性能出ないんで…。
663:デフォルトの名無しさん
09/05/12 22:31:21
Earlang（笑）
664:デフォルトの名無しさん
09/05/16 00:29:01
統計解析ソフト「R」用のパッケージ「gputools」：
URLﾘﾝｸ(cran.r-project.org)
これのWindows版バイナリを作ってくださるネ申はいらっしゃいませんでしょうか。 m(ﾟ-ﾟ；)ｶﾐｻﾏ…
665:デフォルトの名無しさん
09/05/16 10:27:22
>>660
こんなのはあるみたいですが
URLﾘﾝｸ(ruby-opencl.rubyforge.org)
666:デフォルトの名無しさん
09/05/18 17:38:13
GPUの計算部分で
レジスタを多く使っちゃうようにコンパイラで最適化されちゃうんだけど
部分的に無効にする方法はありませんか？
667:,,・´∀｀・,,）っ-○○○
09/05/18 20:36:57
volatile
668:デフォルトの名無しさん
09/05/18 23:05:56
>>667
？volatileは最適化から外すだけで、レジスタには適応されるっしょ
つか>>666 のレジスタ使ったら嬉しくない事ってのが想像できない。。。
669:デフォルトの名無しさん
09/05/18 23:42:16
>>648
ＳＰだって年間何億個でも量産できるけど。
>>649
誰も、ＧＰＵを何個も用意しろとは言ってない。
670:,,・´∀｀・,,）っ-○○○
09/05/18 23:56:12
>>668 volatile __shared__
671:デフォルトの名無しさん
09/05/19 09:21:51
>666の動機によっては__shared__では何の解決にもならないような。
確認していないけれど、恐らくレジスタを使い回さずに消費しまくる方が速いんだろうねぇ。
672:デフォルトの名無しさん
09/05/26 14:36:24
SP2＋未公開パッチで7RC以上に軽くなってるよ
Windows Vista SP3 Part1
ｽﾚﾘﾝｸ(win板:225番)
673:デフォルトの名無しさん
09/05/26 14:38:56
ｺﾞﾊﾞｰｸ
674:デフォルトの名無しさん
09/05/29 18:11:27
SSE ４コアフルに使ったら
最上位品でも大差ないｗ
675:デフォルトの名無しさん
09/05/30 15:06:57
URLﾘﾝｸ(code.google.com)
>Thrust is a CUDA library of parallel algorithms with an interface resembling the C++ Standard Template Library (STL).
676:,,・´∀｀・,,）っ-○○○
09/05/30 18:16:26
きた！STLきた！これで勝つる！
ないない
677:デフォルトの名無しさん
09/05/30 21:28:10
brookみたいだな。
678:デフォルトの名無しさん
09/06/02 17:20:15
仮想マシン上でCUDAのインストールに成功した方はいらっしゃいますか？
当方、ホストOS：Vista、ゲストOS：Ubuntu8.04です。
仮想マシンであるUbuntu上で、NVIDIAドライバ: NVIDIA-Linux-x86-180.22-pkg1.runを起動してみました。
すると、「 You do not appear to have an NVIDIA GPU supported by the 180.22 NVIDIA Linux graphics driver installed in this system」とエラーがでました。
GPUは、GeForce 9800 GTです。どなたか、アドバイスお願いします。
679:デフォルトの名無しさん
09/06/02 17:50:02
仮想マシンは無理
680:デフォルトの名無しさん
09/06/03 02:45:20
>>675
合計なんかCUDAでやって早くなるのか？
681:,,・´∀｀・,,）っ-○○○
09/06/03 02:49:18
分割統治法は並列化の基本だな
682:デフォルトの名無しさん
09/06/03 10:53:10
合計求めるのは苦労したなぁ。
結局、分割数(128とか256とか)置きに足していって、その結果はPCで足した記憶がある。
683:デフォルトの名無しさん
09/06/03 11:18:35
とりあえず公式の３つをインストールしたのですがTMPGEncで確認できませんみたいなことを言われました。
インストールするだけではcudaの恩恵を与れないのでしょうか?
684:デフォルトの名無しさん
09/06/03 11:31:01
すいません、直ぐ解決しましたorz
ｸﾀﾞがちゃんと動いているか確認したいのですが方法はありますでしょうか？
685:デフォルトの名無しさん
09/06/03 11:57:52
>>684
SDKをインストールしたのなら、サンプルもインストールしてビルドしてみよう。サンプルが動けば、大丈夫。
# 特に、deviceQueryは便利。
686:デフォルトの名無しさん
09/06/03 14:12:00
こんにちは。CUDA初心者です。質問があります。
Visual C++ 2008、CUDA tool kit ver2.1、CUDA SDK ver2.1
で、サンプルのsimpleCUBLASをビルドすると、
1>LINK : fatal error LNK1181: 入力ファイル 'cutil32D.lib' を開けません。
と出ます。
そこで、CUDA SDKのlibを調べたところ、
cutil64D.libがあり、32のほうはありませんでした。
この場合、どうすればビルドできるのでしょうか？
687:デフォルトの名無しさん
09/06/03 15:00:01
リリースバージョンをリンクする。
688:デフォルトの名無しさん
09/06/03 16:01:01
リリース構成でビルドしたところ、今度は、
1>LINK : fatal error LNK1181: 入力ファイル 'cutil32.lib' を開けません。
と出ました。
CUDA SDKには、64があり、32はありません。
32と64の違いって一体何なのでしょうか・・・
689:デフォルトの名無しさん
09/06/03 16:17:51
パスが通ってないんだろ
690:デフォルトの名無しさん
09/06/03 17:07:17
OSが64bitだと、32bitのCUDAライブラリはインストールされなかったような。
691:デフォルトの名無しさん
09/06/03 17:07:40
ご回答ありがとうございます。
パスが通っていないということですが、
「パスを通す」について、詳しく説明していただけませんか？
知識不足で申し訳ありません；；
692:デフォルトの名無しさん
09/06/03 17:43:34
環境を名に使ってるかによるが、
Visual Studioだとプロジェクトのプロパティからインクルードするファイルがあるディレクトリのパスと、
libがあるディレクトリのパスをついかする
linuxだとコンパイラのオプションに追加する
詳しくはぐぐれ
693:デフォルトの名無しさん
09/06/03 18:04:37
何度も回答していただきありがとうございます。
リンカの追加ライブラリを調べたところ、
ちゃんと、SDKのcommon/libが指定されてました。
しかし、この中には、cutil32.libではなく、64があります。
ということは、690さんのおっしゃるとおり、
32bitのCUDAライブラリはインストールされなかったということなのでしょうか？
もしそうでしたら、サンプルプログラムは64bitに対応してないが、
自分でプログラムを作る分には、上記のようなエラーはでないということでしょうか？
694:デフォルトの名無しさん
09/06/03 23:53:30
サンプルのリンカ設定を編集して64bitのライブラリをリンクすればいいんじゃね？
695:デフォルトの名無しさん
09/06/04 13:12:53
アクティブソリューションプラットフォームとプラットフォームをWin32からx64へ変更したところ、
エラーがなくなりました。
そのかわり、
1>------ ビルドのスキップ: プロジェクト: simpleCUBLAS ------
1>
========== ビルド: 0 正常終了または最新の状態、0 失敗、1 スキップ ==========
とスキップしてしまいました。
何が原因なのでしょうか・・・
696:デフォルトの名無しさん
09/06/04 18:26:30
threadIdx.xがうまく値を返さなく困っています。
最小のプログラミングだと ちゃんとした値を確認できるんですけど、
規模のあるプログラム書いた物では、threadIdx.xをみると最大でも1000以内の数値が40000を超えていたりします。
かなりラフな書き方していて、グローバル変数使いまくってるのが意見ないのでしょうか？
__device__ kouzoutai[2000];//グローバル変数
とか宣言しまくって搭載メモリーを超えちゃってるかもしれませんが、その時は明確なエラーとか出ますか？
697:デフォルトの名無しさん
09/06/04 19:59:22
エラーは基本的に出ない
698:デフォルトの名無しさん
09/06/05 09:08:24
>>696
threadIdxはプログラムで変更できないと思うが。
どうやってthreadIdx.xを「みる|確認できる」としたのか知らんが、その手段を確認すべきでは?
ptx出力を眺めればthreadIdx.xをコピーした先で壊してたりするのが分かるんじゃない?
グローバル変数にたった2000要素くらいなんてことないと思うが、巨大な構造体でも置いてるの?
あーそれから、メモリをオーバフローしてもコンパイルエラーも実行時エラーも出ないどころか、
突然システムごとフリーズするかもしれないから覚悟しておいてね。
699:デフォルトの名無しさん
09/06/06 17:41:22
>>501
？
∧∧
(´･ω)
_｜⊃／(＿＿＿
／ ヽ_(＿＿＿_／
￣￣￣￣￣￣￣
700:デフォルトの名無しさん
09/06/07 01:25:14
寝た子を起こしたなw
701:デフォルトの名無しさん
09/06/08 08:06:10
>>699
！
∧∧
(･ω･ )
_| ⊃／(＿＿
／ ヽ-(＿＿_／
￣￣￣￣￣￣
702:デフォルトの名無しさん
09/06/08 20:11:26
あ"
∧∧
(◎ω◎)
_| う／(＿＿
／ ヽ-(＿＿_／
￣￣￣￣￣￣
703:デフォルトの名無しさん
09/06/09 23:30:40
>>701
ｵﾊﾖｰ!!
∧∧ ∩
(`･ω･)/
⊂ ノ
(つノ
(ﾉ
＿＿_／(＿＿_
／ (＿＿＿／
￣￣￣￣￣￣
704:デフォルトの名無しさん
09/06/10 20:07:54
質問。
グリッド・ブロック・スレッドのカーネル側での計算への応用方法が
よくわからないのですが・・・・特に動きです。
const unsigned int tid =blockIdx.x * blockDim.x + threadIdx.x;
がスレッドＩＤを示しているというのはわかるのですが・・・
for(int x=0;x<height;x++)
for(int y=0;y<width;y++){
sum=sum+tex2D(tex, x, y)*((float)(cos((M_PI*((x*width+y)+0.5)*tid)/number)));
}
outputdata[tid]=sum;
の中で、tidはどういう働きをしてるんでしょうか？
特に回収する配列であるoutputdataの動きが知りたいです。
わかる方ご教授お願いします・・・
705:デフォルトの名無しさん
09/06/12 04:58:52
>>703
∧,,∧
( `･ω･) ｳｰﾑ…ここは？
/ ∽ |
しー-Ｊ
706:デフォルトの名無しさん
09/06/12 09:52:33
>>704
何が判らんのか判らん。
tidが個別のデータスレッドの番号を指すように働くことが判っているなら、
その番号でoutputdataにアクセスしていることくらい判りそうなもんだが。
仮に、blockDimが4でgridDimが2だとしたら、blockIdxは0か1、threadIdxは0から3になるわけで。
つまりはtidは0から7になるわけだな。
つーか、VIPのWikiにも情報があるよ。
URLﾘﾝｸ(vipprog.net)
707:デフォルトの名無しさん
09/06/13 00:18:28
>>730
カウンター3連打って言うけどスローで見るとカウンター2発に追撃1発って感じ
まあそれでも凄いって言うか、何この異次元映像ｗ
もはや芸術的とか圧倒的とかそういうレベルを超越していて笑いしか出てこないわけだが？ｗｗｗ
しかもメディナは立ち上がるしｗｗ不死身かｗｗｗｗ
この試合はボクシング以外の何か別の競技だわｗｗｗｗｗｗｗｗ
708:デフォルトの名無しさん
09/06/13 00:24:10
誤爆すまそ
俺が代わりに言っといてやったぜ
709:デフォルトの名無しさん
09/06/13 01:35:36
CUBLASに加えてCULAPACKとか用意して欲しい
LU分解もできない環境なんて
710:デフォルトの名無しさん
09/06/13 04:07:09
GPGPUはDirectX11で花開くんだろ。
エンコード高速化対応お待ちしております。
DirectX 11講座（中編）
非ゲーマーのアナタのパソコンも
DirectX 11演算シェーダーでスーパーコンピューターに変身する！
URLﾘﾝｸ(game.watch.impress.co.jp)
GPGPUの場合は、組んだGPGPUプログラムがATIかNVIDIAのどちらかでしか動かない、
踏み絵的な「ATI対NVIDIA」なのだ。
これではGPGPUの進化、発展、普及が望めない。
そこでGPGPUの標準化プラットフォームが強く望まれるようになった。
このGPGPU標準プラットフォームを、Windows環境下で提供しようと出てきたのがDirectX 11ということになる。
DirectX 11は、このGPGPUの標準プラットフォームとして「DirectX Copmute Shader」（演算シェーダー）を提供する。
演算シェーダーでなにができるのか
URLﾘﾝｸ(game.watch.impress.co.jp)
711:デフォルトの名無しさん
09/06/13 08:42:36
>>710
エンコードって意外に並列処理できる箇所少ないんだけど。。。CUDAでどのベンダも対応してこないのはそのため。
Compute Shader出てきても状況は変わらないと気が
712:デフォルトの名無しさん
09/06/14 15:19:59
それマルチ
713:デフォルトの名無しさん
09/06/14 18:56:21
今日本屋見てきたけど
CUDAの入門書とかって無いね。
世の中にはあるのかな。
714:デフォルトの名無しさん
09/06/16 16:56:23
__device__ int particle_position[30][30][50][30];
__device__ int particle_position_num[30][30][50];
：
とか
__device__ __constant__ float K_dWpress;
：
な感じで、結構たくさんグローバル変数をデバイス側に作ってるんだが、
これって、CUDA的にやっていい手法？てか数次元配列ってカーネル関数内で普通のcっぽくアクセスできる？
とりあえずメモリは足りるみたいだから、デバイスホスト間でのデータ転送は基本的にしなくて、
結果のみホスト側に持ってきたいんだけど。
コンパイルエラーはないけど、カーネル起動させたあとエラーが出て、
一瞬画面が消える。
715:714
09/06/16 17:22:05
追記
ちなみに、Emuモードだと問題なく終了できます。
716:デフォルトの名無しさん
09/06/16 17:29:37
>>714
CUDAはメモリ管理なんてしないから、デバイス側の変数がどれだけあるのか自分で把握すること。
一瞬画面が消えるのは表示用とCUDA用を同じGPUでやっている所為だと思われ。
もっと酷いときはフリーズするから要注意で。
717:デフォルトの名無しさん
09/06/16 23:29:13
CUDA初心者です。
サンプルをいじくっているのですが、いくつかわからないことがあり質問させてください。
①テクスチャメモリ
キャッシュが効くからグローバルメモリよりも高速との事ですが、
時間を計ってみるとグローバルメモリと変わらない・・・。
どのような用途で使用するのでしょうか？
②異なるカーネル関数の同時実行
ブロック、スレッド数を決めてカーネル関数を実行しますが、
cudaThreadSyncrinize関数やメモリコピー関数を呼ばなければ
同期待ちはしないですよね？
であれば、カーネル関数を2つ書いたときに同時実行するのでしょうか？
マニュアルを見る限りではできないようですが、会社の先輩が出来るというのでホントかなと。
お分かりの方宜しくお願いいたします。
718:デフォルトの名無しさん
09/06/16 23:32:51
俺も初心者だが。
1. 線形補完できるよ！ 速度面で違わないなら、それ以上はどう違うのか知らん。
2. なんか俺の動かしてる感触だと同時実行してる気がする。 少なくともCPUとGPUで同期はしていないよ。
719:717
09/06/16 23:59:21
>>718
サンクス。
①は線形補間機能があるのは存じていますが、
補間せず生データを使用したいものでして・・・。
②は自分でも本家CUDAZONEで調べてみましたが、どうもできないっぽいです。
URLﾘﾝｸ(forums.nvidia.com)
このスレの前にもありましたね。
よく見てませんでした、すみません。
cudaStream～って何のためにあるんだか。
720:デフォルトの名無しさん
09/06/19 00:40:24
CUDAを始めようかと考えていますが、下記の様な処理はCUDAで速くなりますか？
721:デフォルトの名無しさん
09/06/19 00:55:31
function hoge(const b:string):boolean;
begin
result:=true;
end;
722:デフォルトの名無しさん
09/06/19 00:58:07
>>720
すみません、間違えて書き込んでしまいました。
改めて、
typedef struct _table_t {
float x,float y, float z;
} table_t;
typedef sturct _vec {
float vector[4];
} vec;
typedef union _vtemp_t {
vec V[3], table_t data[4];
} vtemp_t;
table_t DATA[100];(DATAにいろいろと情報をが入っている)
vtemp_t temp;
temp.data[0] = DATA[15];
temp.data[1] = DATA[43];
temp.data[2] = DATA[11];
temp.data[3] = DATA[80];
(tempのメンバdataにＤＡＴＡの値をランダムに代入)
メンバdataがもつ、x､y､z情報をメンバVにx､y､z毎にまとめたいです。例えばこんな感じに。
V[0] = {data[0].x, data[1].x, data[2].x, data[3].x};
V[1] = {data[0].y, data[1].y, data[2].y, data[3].y};
V[2] = {data[0].z, data[1].z, data[2].z, data[3].z};
CellとかSSEならSIMDを使って出来ると思いますが、CUDAでも速くなるなら、
CUDAをやってみようかと思っています。
723:デフォルトの名無しさん
09/06/19 01:25:58
それ計算じゃなくてただのメモリ転送じゃん。 CUDAでやることじゃないよ。
724:デフォルトの名無しさん
09/06/20 14:41:38
Mac でCUDAをやろうとしています。
/Developer/CUDA/lib/にあるライブラリをライブラリのサーチパスに含めるには
どの環境変数を設定すればよいのでしょうか？
LD_LIBRARY_PATHに追加してみたのですが、サーチしてくれません
725:デフォルトの名無しさん
09/06/20 14:43:15
>>724
DYLD_LIBRARY_PATH
726:デフォルトの名無しさん
09/06/20 15:02:09
>>725
早速の回答ありがとうございます。
DYLD_LIBRARY_PATH に追加してみたのですがやはりうまくいきません。
ちなみに-Lで指定すればうまくいきます。
727:デフォルトの名無しさん
09/06/20 18:11:27
Coalesced、Non-Coalescedって簡単にいうとどういうこと？
728:デフォルトの名無しさん
09/06/20 18:14:29
>>108
729:デフォルトの名無しさん
09/06/20 18:18:03
>>728
よくわからん
730:デフォルトの名無しさん
09/06/20 20:19:04
URLﾘﾝｸ(noridon.seesaa.net)
731:デフォルトの名無しさん
09/06/20 21:42:57
大前提として、ブロック内でデバイスメモリにアクセスするアドレスが連続していたら、コアレス。
んで対応アーキテクチャによってスレッド単位でアドレスが連続している必要があったりする。
まぁ簡単に言うと、バス幅有効活用できまっせ適な使い方をイメージすればいいかと。
なんでそのバス幅でアクセスできるアラインを考える事になる。
732:デフォルトの名無しさん
09/06/20 23:08:40
>>726
実行時じゃなくて、ビルド時のリンクのこと？
DYLD_LIBRARY_PATHは関係ないから-Lでやってね。
733:デフォルトの名無しさん
09/06/21 01:03:25
>>726
LIBRARY_PATHでできなかったっけ？
734:デフォルトの名無しさん
09/06/21 06:52:13
配列で最初と最後の変数だけ別の処理をしたいのですが、
どのように書けばよいのでしょうか？
従来のプログラムの場合、配列がa[N]だったばあい、
for(i=1;i<N-1;i++)と書けばよいのですが、
CUDAのカーネルで処理する場合どのようにすればよいのでしょうか？
735:デフォルトの名無しさん
09/06/21 07:45:47
>>734
とりあえず、現状のカーネル書いてくれれば教えやすいのに
736:デフォルトの名無しさん
09/06/21 14:51:04
NVIDIAが配布しているということで、PhysXの質問をさせて下さい。
キネマティックなアクターが物体と接触した時に受ける力の取得方法が分かりません。
ご存知の方がいらっしゃいましたらご教示下さい。
737:734
09/06/21 17:50:38
>>735
例えば、今は
__global__ void test(int thread_size, int one_size, float *res_d)
{
int bx = blockIdx.x;
int tx = threadIdx.x;
int pos = bx * thread_size + tx;
if (pos==0)
res_d[pos] = cos(res_d[pos]);
else
res_d[pos] = sin(res_d[pos]);
}
このように、if文で0とN-1の場合だけ別の処理をするようにしているのですが、if文は遅くなるので、避けたいのですが。
具体的にやりたいのは、3次元の偏微分方程式を解く際の境界条件部分の計算です。
この場合各面、各辺、各頂点合わせてif文が20個以上になってしまいます。
何か良い方法はないでしょうか？
738:デフォルトの名無しさん
09/06/21 18:12:51
>>737
最初と最後だけ別の処理なら、そこはCPUに任せて、
残りをGPUにやらせればいいと思うけど、それじゃだめなん？
posを+1して、並列数をN-2にすればいけそうだと思うけど。
あと、そのソースだと N-1 の処理が別になってない気がするのと、
posがN以上かどうかのチェックが無いのが気になる。
（ソースは単に省略しただけで、チェックはそもそも必要ないように組んでるのかも知れないが）
739:デフォルトの名無しさん
09/06/21 18:53:40
>>737
一般的にアセンブリレベルで分岐をなくしたいときは、こんなテクニックを使ったりする。
// i == Nのときだけsinをcosにする
float s = sin(なんとか);
floac c = cos(なんとか);
int b = i == N; // true が (int)1に変換されることを期待する
return (c * i) + (s * (1 - i));
もちろん3行目は実装依存なので、プロセッサのマニュアル見ながらいちいち確認はした方がいいんだが、
だいたいのプロセッサで使えるテクニック。
740:デフォルトの名無しさん
09/06/21 19:09:54
>>737
それだと全てのパターンでsinとcosの両方が計算されるから
float s = sin(なんとか + b * PI / 2);
にすれば？
741:740
09/06/21 19:11:22
ごめん安価ミスった
>>740は>>739へのレスです
742:デフォルトの名無しさん
09/06/21 19:37:25
CUDAにはsincosねーの？
また三項演算子でいいんじゃねーの？
743:デフォルトの名無しさん
09/06/21 23:01:37
>>737
あらかじめ、特殊条件について位相を補正しときゃいいんじゃないの？と思うのは俺だけ？
744:デフォルトの名無しさん
09/06/22 00:12:40
>>736
PhysXスレがゲーム製作技術にあるからそっち行け
745:736
09/06/22 01:52:13
>>744
誘導ありがとうございます
746:デフォルトの名無しさん
09/06/22 10:53:08
>>742
あるよ。但し、ストリームプロセッサに一個しか超越関数演算機がないから使い方間違うと遅くなるけど。
747:デフォルトの名無しさん
09/06/23 08:26:25
>>746
超越関数演算機なんてハード的に存在したっけ？倍精度演算機じゃなくて？
748:デフォルトの名無しさん
09/06/23 14:47:11
>>747
SFU(Super Function Unit)がある。
ただし、4SPで共有で、>>747のいうように8SPで共有するのは倍精度演算器。
sin()、cos()、tan()は確かに遅いけど、
sinとcosについては__sinf()と__cosf()を使えば、
精度は悪くなるけど1op/clockでできる。
除算以外の算術演算が8op/clockでできることを考えたらそれでも遅いことには変わりないけど。
749:デフォルトの名無しさん
09/06/23 15:22:26
超越関数がレイテンシ1なんて凄いね
SFUだけ100倍のクロック数とか
してるのかなぁ
750:デフォルトの名無しさん
09/06/23 18:30:15
URLﾘﾝｸ(www.khronos.org)
OpenCL 1.0 リリース
751:デフォルトの名無しさん
09/06/23 18:36:49
First OpenCL demo on a GPU
URLﾘﾝｸ(www.youtube.com)
752:デフォルトの名無しさん
09/06/23 18:43:26
URLﾘﾝｸ(www.nvidia.com)
ドライバ一式は登録しないといけないようだ
753:デフォルトの名無しさん
09/06/23 21:47:02
>>748
SFUなんてあったのか、しらなかった。DPみたいに実行はSPと並行してできるのかな？
754:デフォルトの名無しさん
09/06/24 07:13:15
やっとFortran対応の話出てきたけど、有料なんかね
PGI and NVIDIA Team To Deliver CUDA Fortran Compiler
URLﾘﾝｸ(eu.st.com)
755:デフォルトの名無しさん
09/06/24 07:18:14
ていうかいつのまにかに2.3のベータ出てたのか
756:デフォルトの名無しさん
09/06/26 12:01:06
GTX 285でようやくおれのようなへっぽこでも普通にcpuを凌駕するコードが書けるように
なったようだね。
後は（値段はともかく）消費電力が下がってくれれば、大ブレイクしそう。
757:デフォルトの名無しさん
09/06/27 12:19:52
３００Wなんてクアッドコアマシンを４台くらいクラスタリングしたような消費電力だからな
758:デフォルトの名無しさん
09/07/01 03:13:56
>おれのようなへっぽこでも普通にcpuを凌駕する
だったらCPU要らんがな。
むしろCPU<->GPU間の通信がボトルネックなので、CPU無くして
GPUに直接I/O繋げられるようにしろや！
759:デフォルトの名無しさん
09/07/01 21:14:50
>>758
それなんていうTesla?
760:デフォルトの名無しさん
09/07/01 23:30:38
>>759
それTesｌaでもない。
761:デフォルトの名無しさん
09/07/02 07:45:08
CUDAで乱数を効率的に生成する方法を教えてください
762:デフォルトの名無しさん
09/07/02 09:23:34
>>761
コアごとに違う種から乱数生成すればいいんじゃね？
763:デフォルトの名無しさん
09/07/03 00:25:57
質問なんですが、SDKに付いてたサンプルをmakeしようとしたところ
以下のようなエラーが出てうまくいきません、何か足りないものがあるのでしょうか
/usr/bin/ld: cannot find -lXi
collect2: ld はステータス 1 で終了しました
make[1]: *** [../../bin/linux/release/nbody] エラー 1
OSはfedora10です。色々初心者なのでどうかお願いします。
764:デフォルトの名無しさん
09/07/03 00:52:49
libXiというパッケージをインストールできないかね
俺もfedoraはよく記憶にないんだが
765:デフォルトの名無しさん
09/07/03 07:55:20
>>764
libXiというのがインストールされていないということだったのですか
早速インストールしたいと思います。どうもありがとうございました。
766:デフォルトの名無しさん
09/07/05 04:21:12
teslaとgeforceって何が違うの？
767:デフォルトの名無しさん
09/07/05 05:03:33
OpenGL向けかDirectX向けか
768:デフォルトの名無しさん
09/07/05 10:40:00
>>767
それquadroとgeforceでしょ。
>>766
teslaはhpc専用ボードで、ディスプレイの出力ポートがない。
NVIDIAがボードの製造まで一貫して行ってて、テストもしっかり通している…らしい。
あとメモリが4GB乗ってる。
769:デフォルトの名無しさん
09/07/06 09:01:42
CUDAというよりハードウェアの話なんですが質問させてください
TESLA C1060をMM3500に刺したいんですが、この状態だとBIOSが拝めません
ビープ音「－・・（長短短）」って音がするんでAwardで調べたらビデオアダプタ不良らしい
マザーボードだけを他のに変えたら正しく起動するので、電源やTESLAが悪いわけではなさそうです
TESLAなしのMM3500だけならUbuntuのBoot CDからMemtestが走るとこまで行きました
770:デフォルトの名無しさん
09/07/06 09:34:48
スレ違いどころか板違いだ
判ってるなら該当板いけよ
771:デフォルトの名無しさん
09/07/06 22:26:16
>>769
買ったとこに聞いた方が良いかと...
772:デフォルトの名無しさん
09/07/07 00:46:08
>>772
ありがとう まずはVIAにメールしてみます
773:NVIDIAの営業の受け売りw
09/07/07 09:05:04
・GeForce
NVIDIAはパーツ選択にはノータッチ。各社さんが勝手に安い部品を積んでいるから動作しなくても当然。
プロならそんなアキバ的発想はやめて、QuadroFXかTeslaを使ってください。
・QuadroFX
NVIDIAの厳しいテスト基準に合格した高品質の出力を保証しているから決して安くありません。
プロにも安心してお使いいただけます。
・Tesla
QuadroFX以上に品質に重点を置いてテストを行なっています。HPCにはこれ以外の選択肢はありえません。
774:デフォルトの名無しさん
09/07/07 09:07:57
>>768
OpenGL向けもDirectX向けもないよ。>773にもある通り、NVIDIAがテストしているかどうかの違い。
なんせ、GeForceの製造販売を行なっているELSAの営業の前で堂々と>773みたいなことを仰ってる。
10倍の価格差も、当然ということなのでしょう。
775:デフォルトの名無しさん
09/07/07 22:07:57
そりゃカノプーが手を引くわけだよな
どうしたってGeForceは価格競争になっちまう
776:デフォルトの名無しさん
09/07/08 00:28:37
不特定多数の客向けのプログラムでGPGPUを使うコードを書くやつはアホウ、ってことだな。
777:デフォルトの名無しさん
09/07/08 06:57:16
全てアホウと決め付けるのはねぇ
中には超やばいものを開発するかもよ（世界は広い）
特にロシア東欧は要注意！
778:デフォルトの名無しさん
09/07/09 00:40:40
>>777
アホウと言っているのが、nVIDIAの営業ってことだよ。
779:デフォルトの名無しさん
09/07/11 00:02:30
「阿呆」とは言わなかったけれど、「動かなくても当然ですから」くらいのことは言ってたなぁ。
780:デフォルトの名無しさん
09/07/11 00:17:01
今月末に CUDA2.3 が出るとさ。
PGI compiler 9.0 は2週間体験できるから、Fortran で単純な行列積の
コードを書いて試してみたけど、Core2 Quad + GeForce9800GTX で
CUDA 2.0 という古い組み合わせのせいか，
!$acc region の設定の効果が全く現れなかった。
pgf95 -ta=nvidia:cc11 matrix.f
というコマンドでいいと思うんだけど。
781:デフォルトの名無しさん
09/07/11 19:31:18
CUDA用の姫野ベンチをGTX285で実行しようとしたのですけど動きませんでした。
何が原因かわかりません、どなたかわかる方いらっしゃいましたらお願いします。
782:デフォルトの名無しさん
09/07/11 20:08:49
原因はわかりますが、何をお願いされているのかわかりません。
783:デフォルトの名無しさん
09/07/11 21:19:47
その原因について詳しく教えて頂きたいのですが
784:デフォルトの名無しさん
09/07/11 22:37:16
>>783
原因なら簡単ですよ。
あんたが人間じゃなくて人だからです。
785:デフォルトの名無しさん
09/07/11 23:13:59
動くかどうかも保障してないものが普及するとか思ってるのか？
786:デフォルトの名無しさん
09/07/11 23:20:06
>>785
>779のことを言っているのなら、QuadroFXとTeslaがあるから大丈夫だと思っている
頭に蛆が涌いているとしか思えないNVIDIAの営業に言ってあげてください。
無駄だけどw
787:デフォルトの名無しさん
09/07/15 19:53:47
visual profilerについて質問です
プロファイラの項目GPUTimeを見ると全体で3000くらいなのですが、
プログラム中で cutCreateTimer( &timer) を使って計測すると 18ms～60msのバラけた値が出ます
これってどちらが正しい値なのでしょうか？
788:デフォルトの名無しさん
09/07/15 21:31:07
>>781
バイナリのCUDA SDK1.1だからだろ
789:デフォルトの名無しさん
09/07/15 23:02:02
>>787
cutCreateTimer()の実装を見てみた?
790:デフォルトの名無しさん
09/07/20 19:00:37
Fixtarsの資料面白いね、祝日に更新ご苦労さんだわ
この調子で本家資料全部訳してくれないかなｗ
URLﾘﾝｸ(gpu.fixstars.com)
791:デフォルトの名無しさん
09/07/22 08:05:55
2.3来たね、やっと64bit版で32bitコンパイルできるようになった
792:デフォルトの名無しさん
09/07/23 13:44:49
Mac版の2.3.0ドライバーインストールしてみたら、一回目libcuda.dylibだけ入らなかったことに気付かなくて、全然見当違いの所で原因探してて苦労したよ。
一発で入るようにしてください、nVidia様
793:デフォルトの名無しさん
09/07/25 22:10:54
URLﾘﾝｸ(www.nv-info.com)
20年前の高校生は参加しちゃダメかな?
794:デフォルトの名無しさん
09/07/25 22:13:36
Teslaプレゼントしろよ
相変わらず一貫してないな
795:デフォルトの名無しさん
09/07/26 23:27:14
CUDAで動くトリップ検索ツール頂ける方いらっしゃいませんか
796:デフォルトの名無しさん
09/07/27 00:50:11
ﾉｼ
797:デフォルトの名無しさん
09/07/27 03:25:45
【レポート】中高生がCUDAにチャレンジ! NVIDIAがGPUコンピューティングの体験イベント開催 | パソコン | マイコミジャーナル
URLﾘﾝｸ(journal.mycom.co.jp)
798:デフォルトの名無しさん
09/07/27 10:10:10
>>795
あれ？まだ誰も作ってなかったの？
799:デフォルトの名無しさん
09/07/27 10:43:20
あと6年早ければ・・・
800:デフォルトの名無しさん
09/07/27 11:08:28
>>798
らしいです
801:デフォルトの名無しさん
09/07/27 22:53:08
じゃぁ俺が作るか。
まずはCUDA無いバージョンのソースがあるURL教えてｗ
802:デフォルトの名無しさん
09/07/28 00:02:52
まあ、待て屋。
803:デフォルトの名無しさん
09/07/28 00:07:41
>>802
それはラデ用のしかなかったような…
804:デフォルトの名無しさん
09/07/28 00:11:46
CUDA無しを所望のようだからこれで良くね？
CPU用のコードも当然あるよ。
805:デフォルトの名無しさん
09/07/28 00:17:31
>>804
nVIDIAのVGAしか持ってないんだorz
806:デフォルトの名無しさん
09/07/28 00:22:33
>>805は>>795か？
俺は>>801に向けて書いたんだが・・・
807:デフォルトの名無しさん
09/07/28 00:26:42
>>806
あ、はい。
俺は>>805であり、>>795です。
ID出てないと面倒だね
808: ◆hEpdoZ.tHU
09/07/28 06:51:10
>>807
トリップ使ってけれw
809: ◆DokuomoeW.
09/07/28 16:50:48
>>808
おｋ
>>807です
810:,,・´∀｀・,,）っ-○○○
09/07/29 21:19:39
Larrabee向けのソースコードなら多分あるんだがｗｗｗｗｗ
811: ◆DokuomoeW.
09/07/30 12:59:13
まだ手元にある定額給付金で
ラデ買ってくるかな…
それも一つの選択肢
812: ◆DokuomoeW.
09/07/31 21:10:00
かつてTripcode ExplorerのCUDA版があったらしいね
でも、今はもう消えてる…
813:,,・´∀｀・,,）っ-○○○
09/07/31 21:52:43
ギコハハハ
814:デフォルトの名無しさん
09/07/31 22:14:27
>>812
がはは！それはネタだからーw
815: ◆DokuomoeW.
09/07/31 23:11:08
>>814
そうなんだｗｗｗｗｗｗ
816:,,・´∀｀・,,）っ-○○○
09/07/31 23:31:10
ふふふ
817:デフォルトの名無しさん
09/08/01 05:48:13
あらあらふふふ
次ページ最新レス表示スレッドの検索類似スレ一覧話題のニュースおまかせリスト▼オプションを表示暇つぶし2ch
5328日前に更新/215 KB
担当:undef