【GPGPU】くだすれCUDAスレ【NVIDIA】
at TECH
[前50を表示]
200:デフォルトの名無しさん
08/08/20 20:57:08
URLリンク(jp.youtube.com)
530万ドルのスーパーコンピュータ「CalcUA」で一時間
一台デスクトップPCで30秒
URLリンク(www.atmarkit.co.jp)
ベルギーのアントワープ大学では、それまで使っていたAMDベースの256ノードのクラスタサーバ「CalcUA」の性能を、8GPUを使ったデスクトップPCシステム「FASTRA」が上回った。「CalcUAは530万ドルのスーパーコンピュータ、FASTRAは7000ドルのデスクトップだ」(グプタ氏)。
NVIDIAによればGPUを使った同様のクラスタサーバは、NCSA、イリノイ大学、ノースカロライナ大学、マックスプランク研究所など、すでに十数の組織で使われているという。
201:194
08/08/21 15:43:36
>>196 さん 、>>197 さん
お返事ありがとうございます。
cutilのライブラリは確認済みです。CUDAインストール時に作成される、
〜/NVIDIA_CUDA_SDK/lib の中にlibcutil.aがありました。
OpenGLは前からインストールしており、またそのプログラムは動くので、
GLUT等はインストールされているはずなのですが。。
202:デフォルトの名無しさん
08/08/22 00:33:34
/home/usr1/NFS/NVIDIA_CUDA_SDK/と〜/NVIDIA_CUDA_SDK/libが違うって落ちじゃないよな。
203:デフォルトの名無しさん
08/08/22 01:02:50
自分からテスト専門です、って宣言してる派遣テスターって何なの?
将来プログラマとかSEになりたい、とかならわかるけど。
向上心ないよね、頑固だし。
そういう派遣テスターって、仕様書は読めない、
テスト仕様書も作れない、テストプログラムも作れない
やれることは「テキトーにプログラムを触る」ことだけ。
俺は派遣プだけどさ、こういう派遣テスターがいると
派遣全体がバカにされるんだよ。
テスト専門派遣なんて氏んで欲しいよ、まったく。
今日も正社員の人が派遣テスターに仕様書を読んで
テスト仕様書を作ってください、って説教してたよ。
その派遣は頑固に「何故、仕様書が必要なんですか?」って
反論してたから、きっとテスト専門派遣テスターだな。
仕様書も読まず、テスト仕様書も作らず、ただテキトーに
プログラム触るだけで給料もらおうなんて頭おかしいんじゃねーの?
あ〜あ、あの派遣テスターが3ヵ月後に切られるまで、
仕様書も読まねーでテキトーにテストしたバグ票がまわってくんのかよ。
そんな糞なもん、読んで処理する派遣プの身にもなってくれよ。
うわ〜、しかもそいつが切られる3ヵ月以内に中間納品あるじゃねーか!
テスト仕様書もなしにテキトーにテストして納品か。
中間納品後にソッコウクレームでデスマ必至だな。俺の休みも返上かよ。
派遣専門テスターさんよ、少しは向上心持てよ!
頑固な性格直して仕様書読めよ!テスト仕様書作れよ!
204:デフォルトの名無しさん
08/08/23 14:37:16
うちの客先はテスト専門部署があるけど
環境構築とかテストプログラム作ったり
海外のテスト専門会社へ行って研究したりしてるぞ
205:デフォルトの名無しさん
08/08/23 18:12:10
バカの作ったものをテストするなんて最悪な仕事の1つなんだぜ
206:デフォルトの名無しさん
08/08/23 18:24:28
以前入社してすぐやめた会社でテスターをやらされたが、
そこの会社、3Dの知識ないやつが、OpenGLで懸命に2Dの描画プログラム作ってた
まさにバカの作ったものでテストするのがバカらしくなって会社を辞めた
207:デフォルトの名無しさん
08/08/23 19:33:34
判ったからマ痛に行け。
208:デフォルトの名無しさん
08/08/23 19:54:30
>そこの会社、3Dの知識ないやつが、OpenGLで懸命に2Dの描画プログラム作ってた
こういうバカ、たまにいるよな
2Dだけ描画すればいい仕様なのにOpenGLしか知らないからとりあえずOpenGLで書いちゃうという
すぐに辞めさせないとプロジェクトが悲惨な結果になる
209:デフォルトの名無しさん
08/08/24 01:26:41
>>208
AutoCADは2DもOpenGLで描画していますが。
210:デフォルトの名無しさん
08/08/24 07:16:25
>>209
>207
211:デフォルトの名無しさん
08/08/24 10:49:13
3Dも描画するならいいだろ
2Dしか描画しないのにOpenGLはバカ
212:デフォルトの名無しさん
08/08/24 12:06:13
GDIで描画したらLinuxでWineが必要になっちゃうじゃん
213:デフォルトの名無しさん
08/08/24 12:27:07
いや、2Dだけ描くのにOpenGL使うようなバカはWindowsだけだろ
214:デフォルトの名無しさん
08/08/24 17:34:26
このスレでうだうだ3D話を続ける奴って頭弱いの?
それとも只の構ってちゃん?
しかも切っ掛けはコピペだって辺りが痛過ぎるよね。
215:デフォルトの名無しさん
08/08/24 17:41:37
きっかけはコピペでのバカの作ったものをバカにされるべき
GPU使っているはずなのにBITBLTしかしてないとかー
もーバカ杉
216:デフォルトの名無しさん
08/08/24 21:55:53
BITBLTの後ろ三つはベーコン・レタス・トマトの略だが、前のBITは何の略なんだ?
217:デフォルトの名無しさん
08/08/24 22:58:34
マルチプラットフォームで2Dを描くのに仕事で使えるレベルの表現力と速度を兼ね備えたライブラリって存在するの?
あればどんなものか教えてください。有償無償問わず。
218:デフォルトの名無しさん
08/08/24 23:02:35
>>217
OpenGL
219:デフォルトの名無しさん
08/08/24 23:23:16
マルチプラットフォームの仕事なんて普通ないだろ
プロならプラットフォームに最適化されたライブラリを選択しろよ
220:デフォルトの名無しさん
08/08/24 23:34:53
OpenGL専門学校の学生さんにマジレスするなよ
221:デフォルトの名無しさん
08/08/25 00:54:10
>>219
>マルチプラットフォームの仕事なんて普通ないだろ
?
222:デフォルトの名無しさん
08/08/25 00:56:55
>>217
URLリンク(d.hatena.ne.jp)
223:デフォルトの名無しさん
08/08/25 01:00:30
>>221
これから増えるだろう
お前は近視か?
224:デフォルトの名無しさん
08/08/25 01:01:51
間違った。>>219だ。
225:デフォルトの名無しさん
08/08/25 01:06:43
>>219
openofficeはいったい何なんだ?マルチプラットフォームのソフトウェアじゃないのか?
226:デフォルトの名無しさん
08/08/25 02:40:58
お前は仕事でopenofficeにかかわったのか?
違うだろ
派遣先のプラットフォームだろ
お前の会社がマルチプラットフォームの製品を開発してんのか?
違うだろ
派遣先のプラットフォームだろ
現実見ようぜ
227:デフォルトの名無しさん
08/08/25 02:51:05
>>226
まず、仮定からおかしい。
・マルチプラットフォームのソフトウェアはopenofficeだけではない。
・>>225の仕事の内容を断定する根拠が見当たらない。
・なぜ無数にある可能性から>>226の仮定を導き出したのか不明。
説明をお願いします。
228:デフォルトの名無しさん
08/08/25 02:56:05
>>227
>・マルチプラットフォームのソフトウェアはopenofficeだけではない。
まず仮定からおかしい。
実際に仕事でマルチプラットフォームを開発しているのかを問題にしている。
>>225がopenofficeを開発したというのなら謝る。おみそれした。
>・>>225の仕事の内容を断定する根拠が見当たらない。
まず自分で開発してそうにないopenofficeを出してくるあたり、
ただのバカ=派遣ってこと
>・なぜ無数にある可能性から>>226の仮定を導き出したのか不明。
派遣に可能性なんてない
229:デフォルトの名無しさん
08/08/25 03:04:17
荒らしは放置の方向で。
230:あぼーん
08/08/25 08:29:00
219のほうがバカだと思います。
マルチぷラットフォームの仕事はあると思います。
231:デフォルトの名無しさん
08/08/25 09:16:36
だからお前はマルチプラットフォームの仕事をしたことあるのかと
232:デフォルトの名無しさん
08/08/25 11:23:15
サターンとプレステで同じゲームを同時に作ったぞ
同時に作っただけだけど
233:デフォルトの名無しさん
08/08/25 12:41:48
>>231
Firefoxはマルチプラットフォームじゃないのかね。
234:デフォルトの名無しさん
08/08/25 12:43:22
マルチプラットフォームの要件はオープンソースではかなり多い。
オープンソース系の企業ならそういう仕事していてもおかしくないね。
235:デフォルトの名無しさん
08/08/25 12:48:10
日本語読めないの?
聞かれてるのは「お前は仕事でマルチプラットフォームをやったことがあるのか?」だぞ
236:デフォルトの名無しさん
08/08/25 12:51:51
>>235
話のすりかえ。
俺自身は学生時代にFirefoxの開発に携わったが、
今の話題は「マルチプラットフォームの仕事があるかないか」。
お前が間違っていることに気づけ。
237:デフォルトの名無しさん
08/08/25 12:53:16
そもそも、ここの住人が開発に携わったことがあるかどうかなんて誰も興味を持っていないし、
信憑性もない。
もっと有益は議論をすべきだと思うけどね。
238:デフォルトの名無しさん
08/08/25 12:54:24
スレ違い お前らネットイナゴは佃煮にされて死んでしまえ
239:デフォルトの名無しさん
08/08/25 12:58:38
派遣先企業で派遣先企業のプラットフォームに合わせたものを実装するのがおまえら派遣プログラマの仕事だろ
派遣先正社員からマルチプラットフォーム対応なんて指示されてないだろ
くだらないこと考えて余計な工数使うなよ
240:デフォルトの名無しさん
08/08/25 13:13:32
>派遣先企業で派遣先企業のプラットフォームに合わせたものを実装するのがおまえら派遣プログラマの仕事だろ
派遣先がパッケージベンダってこともあるだろう
なにムキになってるの?
241:デフォルトの名無しさん
08/08/25 13:23:54
>>239
何故派遣にこだわる?
242:デフォルトの名無しさん
08/08/25 13:49:24
ID出すとしょうもない素人以下の奴がいろんなスレに粘着してた
243:デフォルトの名無しさん
08/08/25 14:17:22
>>242
この板はID出ないじゃん
244:デフォルトの名無しさん
08/08/25 14:18:21
なんだみんな派遣だったのか
245:デフォルトの名無しさん
08/08/25 15:59:39
派遣に異常なコンプレックスを持っている人がいるようですが、
池沼の方なので基本放置でおねがいします。
246:デフォルトの名無しさん
08/08/25 16:06:10
夏の間だけ2chもアカウント認証式にすればいいと思う。
247:デフォルトの名無しさん
08/08/25 21:59:49
正社員にコンプレックスでしょ?
派遣のどこにコンプレックスをもてと
248:デフォルトの名無しさん
08/08/25 23:45:38
ひどい脱線ぶりだな。
249:デフォルトの名無しさん
08/08/25 23:49:36
派遣野郎はスレ違い
250:デフォルトの名無しさん
08/08/26 00:13:24
>>232
それありそう。PS3 と Xbox360 とか。どう似せるかに力が注がれる・・・
251:デフォルトの名無しさん
08/08/26 00:59:55
ここはなんのスレなんだ
252:デフォルトの名無しさん
08/08/26 04:09:43
くだすれ
253:デフォルトの名無しさん
08/08/26 06:23:56
CUDAはマルチプラットフォーム。
254:デフォルトの名無しさん
08/08/26 07:35:34
アホかい。
255:デフォルトの名無しさん
08/09/25 14:29:37
━━━━━━━━━━━━━━━━━━━
■ 今週の名言
───────────────────
「HPCは“ハイ・プロダクティビティ・コンピューティング”」
日本SGIのHPC・サービス事業本部で本部長を務める田坂隆明執行役員)
256:デフォルトの名無しさん
08/09/28 22:50:51
CUDAってSIMD並列といいうことは、過去のスパコン用のコード
移植したら激速ってことでOK?
257:デフォルトの名無しさん
08/09/29 11:43:54
>>256
SIMDじゃないよ。仮にそうだとしても、そんなに単純な話じゃない。
258:デフォルトの名無しさん
08/09/30 05:16:33
SIMDだよw
259:デフォルトの名無しさん
08/09/30 16:35:15
SIMDじゃないなら、何なのか問いたい。
260:デフォルトの名無しさん
08/09/30 17:38:27
SIMDだったりMIMDだったり
261:デフォルトの名無しさん
08/09/30 21:12:21
実行はSIMDだけど
コードはスカラだから面倒臭いよ
262:デフォルトの名無しさん
08/10/01 10:37:52
NVIDIAはSIMTと言っているね。SIMD+αくらい。
オンボメモリはSX8i並だからやっぱ速いんじゃねーの。
263:デフォルトの名無しさん
08/10/01 13:58:23
理屈はいいからなんか作れよ
>>過去のスパコン用のコード移植したら激速ってことでOK?
OKの一言で終わることをなにぐだぐだやってんだか
264:デフォルトの名無しさん
08/10/01 18:46:54
要するに、「移植できるもんならしてみろ」ってことだろw
265:デフォルトの名無しさん
08/10/01 23:50:07
だね。単純に移植できるとは思わない方がいい。
266:デフォルトの名無しさん
08/10/02 00:27:50
たしかにブロック切り分けの所もう少しスマートにならないのかと思うのだが
あんなもんデバイスの種類で最適化した組み合わせを自動算出すりゃいいだろ
267:デフォルトの名無しさん
08/10/02 00:29:52
>>266
CUDA2でストリームプロセッサ数なんかも取得できるようになったから、
是非とも自動算出関数を作ってみてくれ。
268:デフォルトの名無しさん
08/10/02 01:08:48
いや作るのは簡単なんだけどCUDAってクラスも使えないしラップ出来ないからだめじゃん
269:デフォルトの名無しさん
08/10/02 01:11:17
あいや出来るのか
使うのなら作るぞ
270:デフォルトの名無しさん
08/10/06 20:37:38
CUDAに興味があっていろいろ調べてるんだけど
URLリンク(slashdot.jp)
こんなにメモリアクセスがきついのか…
へたれの俺はCPUで頑張ろうって気になってきた
このあたりの条件が和らぐのはどれくらい後だろう?
271:デフォルトの名無しさん
08/10/06 21:10:43
それでもビデオカードは一世代速いメモリを積んでいるから、レイテンシを誤魔化せれば、CPUより多少はいいんじゃないかな。
CPUもGPUも、演算能力よりメモリのスピードがネックなんだよねえ。
272:デフォルトの名無しさん
08/10/06 22:18:57
>>270
そうそう、グローバルメモリはコヒーレントな読み出しなら4クロックなのに
ランダムアクセスすると100倍のクロックを喰ってしまう。
スラッシュドットの記事の書き方はちょっと言葉が足りてなくて、
実際にはいかにメモリをコヒーレントな読み書きで済ませるかが鍵ということになる。
例えば、CUFFTの2Dを1000回回すのと1Dバッチと自前の転置を2回ずつ動かすのとでは、
64x64辺りだと余り変わらないのに128x128になると途端に2Dの方が速くなる。
プロファイラで調べると、CUFFTの内部で用意されている転置関数はコヒーレントな読み書きしかしていない。
一方、自前の転置はコヒーレントな読み出しとインコヒーレントな書き出しになっている。
つまり、1000回呼び出すコストをインコヒーレントな書き出しコストが上回ってしまうということだ。
273:デフォルトの名無しさん
08/10/07 01:53:43
CPUより遅くなるということはない
電気代の割りに効果が少ないというのはあるw
274:デフォルトの名無しさん
08/10/11 06:33:45
TMPGEncで軽い処理はCPUより遅い。
VRAMとのメモリ転送あるし。
275:デフォルトの名無しさん
08/10/15 22:02:44
メインメモリと共有するタイプのGPUで処理した方が速い事も状況によってはあるのかな
276:デフォルトの名無しさん
08/10/15 22:26:03
CUDAでは有り得ないからスレ違い。GPGPUスレにでもどうぞ。
277:デフォルトの名無しさん
08/10/26 18:43:49
すいません、超低級な質問です。
dim3 threads(128, 1, 1);
dim3 grid(128, 1, 1);
hoge_kernel <<< grid, threads >>> (d_ptr, 128);
とかで関数を呼び出したんですが、ホスト側のスレッド数が128個生成
されるみたいです。これって、GPU内にスレッドが出来るんだと思って
たんですが、違うんでしょうか?
ちなみに、-deviceemu はつけてません。コンソールに以下のように
出るので、GPUにはいってると思います。
Using device 0: GeForce 8600 GT
278:デフォルトの名無しさん
08/10/26 20:58:06
>>277
>ホスト側のスレッド数が128個生成されるみたいです。
そう判断した根拠は?
それとは別になるが、128x128が妥当かどうかは検証が必要だと思う。
279:デフォルトの名無しさん
08/10/27 17:11:09
質問です。
__device__ void hoge_kernel(void)
{
}
extern "C" __global__ void hoge(void)
{
dim3 threads(16, 1, 1);
dim3 grid(16, 1, 1);
hoge_kernel <<< grid, threads >>> ();
}
上をコンパイルすると、
「error: call can not be configured」
となるんですが、これはどういうことなんでしょうか。
280:デフォルトの名無しさん
08/10/27 17:47:00
__global__な関数から__device__の関数を<<<>>>で呼ぶことはできない。
281:デフォルトの名無しさん
08/11/01 09:40:52
ちょっと質問させてくらさい。。
8ストリーミングプロセッサ(SE)単位で構成されている
ストリーミングマルチプロセッサ(SM)でのことなのですが、
グローバル(ローカル)メモリ上へのデータのロード、
ストア命令は100サイクル以上かかってしまいますよね。
その間、8つのストリーミングプロセッサはおお休みしている
のでしょうか?
それとも、他のWarpのインストラクションが割り当てられたり
するのでしょうか。そうであるといいのですが、
そうするとSharedメモリがあとからのWARPにのっとられてしまって
まずそうですよね???
レジスタ的には32bit×2048本あるようなので(280の場合)
OKそうですが、Sharedメモリは16KBytesしかないし。
実際LD、STに160サイクルぐらいくったら、
160命令ぐらい無駄にしちゃって効率をかなり落としてる気がするので、
なんかやです。
24WARPぐらい、バッファリングしていて、割り当てられるWARPを
out of oderで実行するという記事も、どっかで見たような
気がするのですがSharedメモリがじゃまして無理があるような・・・
なんか良い方法があるのかな???
っていうか、なにか根本的な部分で俺の勘違い?
282:デフォルトの名無しさん
08/11/01 10:12:52
・グローバルメモリアクセスは、最大400(?)クロック掛かるが、最短では4クロックで済む。
# そのためには、coalescedにアクセスできるように工夫する必要がある。
・各ストリーミングプロセッサは、独立して動作する。Sharedメモリも同様。
例えば、行列の転置のような処理の場合、普通に書くとcoalescedに読んでincoherentに書かざるを得ない。
# 或いはその逆か。
そこで、CUFFT内で行なっている転置処理では、(プロファイルで見る限り)一旦共有メモリにおいて同期を取ることで、
読み書き共にcoalescedアクセスを維持しているようだ。
283:デフォルトの名無しさん
08/11/01 18:40:52
>>282
早速のレス、さんくす。
なるほど、最短4クロックですか。coalescedにしてもレイテンシー
だからだめかなって思ったけど、よくよく考えると、DRAMへの直接
アクセスに数百クロックってのはおかしいことに気づいた。
各ストリーミングプロセッサは、独立だけど1インストラクション単位で
同期してるんじゃないんかな(んなことぁない?)
・・・っておもったけどブランチがあるから、TPXレベルから見ると独立か???
最初見たときSharedメモリも他のSPのregも0レイテンシで使える
ようだったんで、演算と独立にグローバルメモリとのロード、ストアができん
16K程度のSharedメモリ、あまり意味ないじゃんとかおもったけど、
そんなことなさそだね。
なかなか面白そうなチップだ。
分岐粒度が32って実際、汎用でどうなんだろ。
ベクトルつかったことないんでわからんが、
適当にセルオートマタでも乗っけて遊んでみまつ。
284:デフォルトの名無しさん
08/11/12 22:42:34
CUDA-Zなんて便利なものがありました。
forum.nvidia.co.jp
285:デフォルトの名無しさん
08/11/12 22:55:32
>>284
kwsk
286:デフォルトの名無しさん
08/11/12 23:19:24
URLリンク(cuda-z.sourceforge.net)
287:デフォルトの名無しさん
08/11/14 19:39:30
あげてすまん。
CUDAのためにGTX280つきのPCかったのだが
ほかにもう一枚グラボ買わないといけんのかな?
あとPCIバスしか残ってないのだが・・・
LINUXで使う予定。
288:デフォルトの名無しさん
08/11/14 22:57:10
>>287
別に一枚でも大丈夫。二枚挿しても(開発には)却って面倒だったり。
まぁ、二枚あれば表示に足引っ張られる心配なくなるから安定はするけどね。
289:デフォルトの名無しさん
08/11/15 00:41:03
>>288
そうなのか、安心した。
ありがd
290:デフォルトの名無しさん
08/11/15 17:39:36
質問です。
float にてゴリゴリ計算して、結果を返すプログラムを書いてみたのですが、
普通のアルゴリズム、CUDA+GPU、CUDA+CPU(deviceemu) の2つで比較して
みたところ、思ったより差が大きいのですが、こんなもんですか?
もしくは、機種依存があったりするんでしょうか?
291:290
08/11/15 17:40:23
>>290
すいません、間違えました。
「2つで比較して」->「3つで比較して」
292:デフォルトの名無しさん
08/11/15 19:43:01
何の差?
演算精度なら、そりゃぁあるさ、floatなんだもの。
293:デフォルトの名無しさん
08/11/15 20:51:43
ホストは、デフォルトではfloatをfloatで計算してくれないからな。
SSEを使うなりで、ホストがちゃんとfloatで計算すれば、結果は一致するんジャマイカ?
294:デフォルトの名無しさん
08/11/15 20:57:47
>>293
一致する保証はないよ。CUDAのドキュメントにもあるけれど、超越関数はGPU内部の組み込み版を使うと若干誤差が残る。
いずれにしてもfloatの想定の範囲内だから、実用上は問題にならないけどね。
295:,,・´∀`・,,)っ-●◎○
08/11/15 22:26:10
x87 SSE CellB.E. CUDA の浮動小数サポートの対比表みたいなのがCUDAのマニュアルにあったな。
確かに完全じゃない。
糞まじめに準拠してるのはx87くらいだ
296:290
08/11/16 02:56:28
>>292-295
なるほど。出てきた結果が、CPU上の double で計算した結果と、
GPUの float で計算した結果が、最大1%程度違ったから、正直
驚きました。普段doubleしか使って無くて、誤差なんかほとんど
気にしなくてよかったので。
誤差が蓄積してくようなタイプのアルゴリズムではないと思っていた
だけに、少し驚きました。
297:デフォルトの名無しさん
08/11/20 18:53:51
Cで使っていた自作ライブラリは、
nvccでコンパイルし直さなきゃダメなの?
298:,,・´∀`・,,)っ-○◎-
08/11/20 18:58:40
食べてみたけどおいしかったよ。
一つどうぞ、あーんして
299:,,・´∀`・,,)っ-●◎○
08/11/20 19:44:47
飾りだ、食うなボケ。
300:,,・´∀`・,,)っ----
08/11/20 21:36:00
僕のもうないです
>>299から貰ってね
301:デフォルトの名無しさん
08/11/20 23:10:32
>>297
ホスト(CPU)上で実行させるものならVC++と一緒にリンクできる。
(cppIntegrationサンプル参照)
GPU上で実行させたいならそれなりにいじくらないと無理。
302:デフォルトの名無しさん
08/12/01 15:16:29
-deviceemuでは動くプログラムが、GPUを使うと誤動作します。
__device__関数に引数が一部しか渡らなくなります(float4のz成分だけしか渡せない)
ループ回数を極端に減らすと改善されるのですが、これはregisterメモリがパンクしてしまったということでしょうか?
303:デフォルトの名無しさん
08/12/01 17:15:56
ソースを見ないとなんとも言えませんが、ループ回数とregisterは依存関係にはありません。
nvcc -ptxでptxファイルを出力して眺めてみては如何でしょう。
304:デフォルトの名無しさん
08/12/01 23:02:30
ptxファイルですか。。。
正直どこをどう見たらいいのかわからないので敬遠していましたが、
遂に避けられないとこまで来てしまいましたかねえ。
見るべきポイントとかありましたら、よろしければ教えてください。
305:,,・´∀`・,,)っ-●◎○
08/12/01 23:09:56
ptxもネイティブコードじゃなしに中間コードの断片だからな。
レジスタが何本割り当てられるかとかそういった情報すら持ってないから困る。
その点Larrabeeは16本+16本+αとかレジスタ本数が決まってて
コード生成時点で割り当てが決まってしまう。
結果的に静的な最適化がしやすい。
良し悪し。
306:デフォルトの名無しさん
08/12/02 13:02:50
リョウテニ●ヽ( ・∀・)ノ● CUDA!
307:デフォルトの名無しさん
08/12/03 17:58:11
.oとかにしてgccでリンクすることってできますか?
308:デフォルトの名無しさん
08/12/04 00:09:51
>>307
何を? どこで? なんで質問もろくにできないの?
ちなみに、エスパー募集ならお門違い。
309:デフォルトの名無しさん
08/12/04 00:52:30
>>307
普通にできますよ。
-l/usr/local/cuda/lib -lcudart
等のオプションは当然自分で追加する必要があります。
310:デフォルトの名無しさん
08/12/04 06:13:34
Windowsで、とかいう話だったりしてね。
311:デフォルトの名無しさん
08/12/04 10:05:39
.objじゃなくて.oって書いてあるんだからそれはないんじゃない?
312:デフォルトの名無しさん
08/12/04 22:31:14
それだったらわざわざ聞くかねぇ。まぁいいか。
313:デフォルトの名無しさん
08/12/12 18:47:36
【GPGPU】CUDA/ATI STREAM 速度・画質検証スレ
スレリンク(jisaku板)
314:デフォルトの名無しさん
08/12/14 19:30:29
これからCUDAを始めようといろいろサイトを巡回しているのですが、
グリッドサイズとブロックサイズはどのようにして決めたらいいのでしょうか?
同期をとるかとらないかで変わると思いますが、とりあえずとらない場合でお願いします。
315:デフォルトの名無しさん
08/12/14 20:41:30
>>314
BlockSize(スレッド数)は32の倍数で余り大きくない辺り。GridSize(ブロック数)は充分大きく。
316:デフォルトの名無しさん
08/12/14 21:31:49
>>315
ありがとうございます。
もう一つ質問させていただきたいのですが、たとえば、512x512のようにブロックサイズ、グリッドサイズ
ともに最大値未満の場合、<16, 32, 1>のように複数の次元に分けるのと、<512, 1, 1>のように1つの次元
にまとめるのとどちらがいいのでしょうか?
多分通常はブロックサイズは各次元の最大値<512, 512, 64>をオーバーすることが多いと思うので、こちらは
分けることになると思いますが、グリッドサイズの最大値は<65536, 65536, 1>もあるので、
次元を分けなくてもいいときは分けない方がいいのでしょうか?
317:デフォルトの名無しさん
08/12/14 22:01:51
threadサイズが実際の分散処理の単位になる
つまりthread数だけ並列演算が行われる
例えばthread(1)にしてblock(10)とかだと順次処理と変わらない
ただしthreadは256までしか出来ないのでblockという単位が容易されてる
blockは単にプログラムがやりやすいように用意されただけのもので
実際に分散処理には影響しない
318:デフォルトの名無しさん
08/12/14 22:15:42
>>316
言いたいことがよく判らんが、演算量が512x512ならブロックを8x8にしてグリッドを64x64でいいんでない?
或いはxyの区別が重要でないならブロックを64、グリッドを4096とか。
私は後者の方針でptxファイルの出力とプロファイルの結果を睨みながら決定することが多いけど。
# 横着するときは前者だなぁw
319:デフォルトの名無しさん
08/12/14 22:17:30
>>316
単に内部ループが
for(x=0;x<?;x++){}
_global_
か
for(x=0;x<?;x++){
for(y=0;y<?;y++){
_global_
}
}
になるかの違い
320:デフォルトの名無しさん
08/12/14 22:24:34
元の位置を特定する為にglobal関数でインデックスを計算することになるから
もしxやyのどちらかで収まるなら1つだけ使うのが良い
global関数内部の計算コストが一番ネックになる部分だからね
インデックスの計算部分が一番邪魔で削れるだけ削った方が良い
321:デフォルトの名無しさん
08/12/14 22:35:25
>>317,319
ありがとうございます。ようやく喉のつっかえが取れました。
>>318
試しに昔書いた画像処理のコードを移植しようとしていたのですが、入力によって
いろいろサイズが変わるので、512x512で固定するのはちょっとやりづらいかなと…
ただ、ここまでの話からするとchannel(RGB)*height*widthの三次元配列に画像をつっこんで、
grid = <height, channel, 1>
block = <32, width / 32, 1>
で処理を回せそうですね。画像の幅をを32の倍数に合わせる必要はありそうですが。
>>320
ありがとうございます。実際に使うときは上のような感じにブロックとグリッドの数を決めようかなと
思いますが、どうでしょうか?
322:デフォルトの名無しさん
08/12/14 22:59:47
heightは超えてもいいの?
範囲は内部でインデックスが範囲かどうか判定するよりは
端数を含めたメモリを確保して必要無い部分も計算させてしまった方が速いよ
323:デフォルトの名無しさん
08/12/14 23:22:38
>>322
そこまで大きな画像を入れるかどうかは分かりませんが、確かに高さが65536を
超える場合は分割してやる必要がありますね。
>端数を含めたメモリを確保して必要無い部分も計算させてしまった方が速いよ
ということは、配列にコピーするときに幅を合わせてコピーするのではなく、後ろに空白部分を
まとめた方がいいということでしょうか?
#これって多次元配列を引数に渡すのってどうするんだろう…
324:デフォルトの名無しさん
08/12/15 00:01:31
いや、最終的には全て一次元の配列でそれが適切な間隔でアクセスされるように並んでいるのが一番いいことになる。
だから、画像なんかの場合だと座標値に意味がないなら詰め込んで構わない。
実際には、コンボリューションフィルターなどで隣の画素を参照しないといけないことが多いだろうけど、
その場合はどうせ共有メモリを使ってアクセスを減らすなどの工夫が必要だからどの途32の倍数に拘る必要はない気がする。
325:デフォルトの名無しさん
08/12/15 00:50:40
論文用のならともかく普通の画像処理で画像サイズが固定されてることはまずないし
サイズも共有メモリじゃとても足りない
共有メモリを使うならCPUサイドであらかじめ画像を分割して渡すという方法しかないけど効率悪いよ
毎回隣接した画素情報も含めないといけないから
326:デフォルトの名無しさん
08/12/15 01:06:12
>>324
カーネルサイズにもよりますが、sharedメモリはちょっとサイズ的にきつそうですね…
constantメモリならブロックサイズを調整すれば何とかなるかも?
>>325
単なる趣味グラムなので、ぶっちゃけ全部globalメモリからとってきてもいいような気がしたのですが、
逆にそれだとCPUで処理するよりも遅くなりそうですね
327:デフォルトの名無しさん
08/12/16 01:54:49
遅くなりそうというかどんだけがんばっても最終的にはglobalメモリを使うのが一番早いという結論になるよ
まあそれがGPGPUの一番のネックなんだけどね
CPUサイドのメモリを共用するようにいずれはなるんだろうけど今の限界はそこ
328:デフォルトの名無しさん
08/12/16 06:11:12
一番(手っ取り)早いのは同意するが、一番速いとは限らない。
それこそ、思考放棄じゃないかな?
329:デフォルトの名無しさん
08/12/19 00:31:55
詰まってます、助けて・・・
私はVS2005で開発しようとカスタムウィザードを入れ、ツールキットとSDKをインスコし
パスも通したのにLNK1181エラーでるんです・・・存在しないオブジェクトファイルを指定するのはなぜでしょう?
ライブラリパスがみつからないってどういうことなんでしょうか?
330:デフォルトの名無しさん
08/12/20 01:37:50
"Program" "Files"と認識してるとか
331:デフォルトの名無しさん
08/12/20 02:38:31
Windowsってなんで重要なディレクトリ名にスペースなんか入れるんだろう。
MSって頭悪いの?
332:デフォルトの名無しさん
08/12/20 04:16:06
低脳を排除するためかもねw
333:デフォルトの名無しさん
08/12/20 05:06:25
プログラマお断り。ずっと消費者で居てくだちい。
334:デフォルトの名無しさん
08/12/20 06:27:59
デスクトップやマイドキュメントのスペースや半角カナフォルダはなくなったが、Program Filesはそのままだな。
335:デフォルトの名無しさん
08/12/20 09:43:22
>>331
その昔16ビットOSが主流だったころM$は8文字ファイル名しか扱えなかった
そしてUnix系OSは既にロングファイルネームに対応していた
そこで登場したのが32bit Window 95だった
長いファイル名が使えるんだぞ、凄いんだぞと新発明でもしたかのように宣伝していた
そして自慢げにフォルダ名を無駄に長くした
336:,,・´∀`・,,)っ-●◎○
08/12/20 12:23:53
VistaはUNIXerには優しくなったな。
上位エディション限定だけどSUAも使えるし
337:デフォルトの名無しさん
08/12/20 20:39:06
CUDAをさ、VS2005で使えてる人っているの?
つか、なんで英語のマニュアルしかないんだよ。わかんねーよ
338:,,・´∀`・,,)っ-○◎●
08/12/21 00:44:02
むしろ2005以外でどうやって使うんだ?
2.1βで2008対応したけど
339:デフォルトの名無しさん
08/12/21 00:58:06
>>338
おお、それはいいこと聞いた
やっと2005アンインストールできる
340:デフォルトの名無しさん
08/12/21 01:37:02
>>338
どこか日本語で環境設定を教えてくれるサイト知ってたら教えてくれ・・・
341:デフォルトの名無しさん
08/12/21 01:55:48
環境設定なんか何も要らんがな。スクリプト走らせるだけ。
bashのバージョンによってはウォーニング出るけど、ちゃんと動く。
スクリプト猿にも成れないとは、どんだけゆとり脳なんだよ?スポンジなの?
342:,,・´∀`・,,)っ-●◎○
08/12/21 10:35:26
どのみちあの程度の英語が読めない人に優れたプログラミング書くのは無理だと思うんだ。
俺の私的な業務の手伝いしてくれるなら日本法人さんに代わりにチュートリアルとかの要求
してあげてもいいんだけど、職権乱用か。
>>340 ここで解らなかったらもう諦めたら?
URLリンク(chihara.naist.jp)
343:デフォルトの名無しさん
08/12/21 11:55:12
どうでもいいけど、↑のURLは「disp_content」を削らないとエラーになった。
344:デフォルトの名無しさん
08/12/21 12:58:13
>>340
サンプルプロジェクトを開いて設定を見て勉強するのがいいよ
345:デフォルトの名無しさん
08/12/24 12:51:08
英語読めないのにプログラミング言語してる男の人って・・・
346:デフォルトの名無しさん
08/12/24 18:35:50
英語に優越感もってるやつって、さもしいな
347:デフォルトの名無しさん
08/12/24 19:03:01
>>345
可愛いもんだよねv
348:デフォルトの名無しさん
08/12/24 23:58:39
英語なんて誰でも1ヶ月もあればある程度は読めるようになるよ
出来ないと思い込んでやろうとしないだけで
349:デフォルトの名無しさん
08/12/26 12:46:32
vs2005の環境で拡張子cuの場合、定義参照はできるのですが、
インテリセンスが機能しません。freeとかvcのランタイムも同様に
機能しないのですが、やり方分かる方いますか?
ツール->オプション->テキストエディタ->ファイル拡張子にcu c++として
登録するとできるのは上記だけです
350:デフォルトの名無しさん
08/12/26 22:27:18
いっそ、nvccの拡張機能以外は全部cppでやったら?
351:デフォルトの名無しさん
08/12/27 05:03:22
それがスマートだね
352:デフォルトの名無しさん
08/12/27 17:53:02
質問です。
VC++2005で開発する際に、コマンドプロンプトにデバイス名も含め一切文字を表示させたくないのですが、どのようにすればよいでしょうか?
353:デフォルトの名無しさん
08/12/27 18:38:27
>>352
長いパス名が邪魔なだけなら set PROMPT=$G とか?
354:デフォルトの名無しさん
08/12/27 18:47:25
>>352
CUTILを使うと標準出力があるって話なら、リダイレクトするかCUTILを使わなければいいと思うのだが。
355:デフォルトの名無しさん
08/12/28 01:33:08
プロンプト出したくないって話なら_tmainじゃね
356:352
08/12/28 02:53:51
>>353-355
回答ありがとうございます。
>>353
質問があいまいですみません。
パス名が邪魔、ってわけではないです。
>>354
おそらくこの指摘にあてはまるのだと思います。
>リダイレクト
リダイレクトしても、コマンドプロンプトに表示が出てしまいます。
>CUTIL
cutil.hを使わずに、cutil.hで提供されている関数を使わない方法、もしくは、代替可能なヘッダーファイルはありますか?
>>355
_tmainを使ってしまうとWindowsでしか使えないソースコードになってしまうので、別の方法があれば教えてください。
357:352
08/12/28 02:55:28
自己レスですw
誤:cutil.hで提供されている関数を使わない方法
正:cutil.hで提供されている関数を使う方法
の間違いです。
358:デフォルトの名無しさん
08/12/28 03:10:48
system("cls");
とか
359:デフォルトの名無しさん
08/12/28 10:49:28
>>356
> _tmainを使ってしまうとWindowsでしか使えないソースコードになってしまうので、別の方法があれば教えてください。
Windows かどうかで #ifdef するくらいは許されるんじゃないの?
360:デフォルトの名無しさん
08/12/28 14:29:28
プロジェクトの設定でWindowsアプリにすればいいだけでしょ
361:デフォルトの名無しさん
08/12/28 14:32:32
>>_tmainを使ってしまうとWindowsでしか使えないソースコードになってしまうので、別の方法があれば教えてください。
mainだろうがWindows依存部分のコードが発生するのは当たり前でしょ
CUDAの部分とC++の部分とWindows依存部分とファイルを分離するのが鉄則だよ
複数の環境で使う場合は#ifdefプラグマを使って分離したWindows依存部分の
#includeを切り替えるようにするだけ
362:デフォルトの名無しさん
08/12/28 17:29:43
コンパイラのオプションで指定できるでしょ
環境ごとにmakefile書くだけじゃね
363:デフォルトの名無しさん
08/12/28 17:48:49
>>361
#ifdefはいつからpragmaになりましたか?
364:デフォルトの名無しさん
08/12/28 17:53:47
>>363
#ifdefディレクティブと言いたかったんじゃないの?
で、使い慣れない言葉なんでつい間違えたと。
句読点もちゃんと使えないようだし、日本語に慣れていない三国人なんでしょ。
365:デフォルトの名無しさん
08/12/28 22:36:27
cudaのMDコードどこかに追ってませんか?
366:デフォルトの名無しさん
08/12/29 12:11:23
たった1行でここまで意味不明なのも凄いな。
367:デフォルトの名無しさん
08/12/29 12:24:23
×追ってませんか?
○オッドアイいませんか?
368:デフォルトの名無しさん
08/12/29 13:39:38
>>365氏はcudaで書かれたMD(分子動力学)コードどこかに落ちてませんか
と聞いてるのかと。当方も知らないので答えられませんが。
Fortranで書いたものを全部はアレなのでGPU上で実行したいサブルーチンだけ
Cに変えてCUDAで動かしたいのですが、そんな例とかは落ちてないですかね。
mainルーチンその他関係ないところまで全部Cに移植するのが嫌ってだけな
んですが。あ、当方はIntelFortran使用。
当方まだCUDA触りたて、試しにSTREAM BMTのtriadだけ手で適当に書いて
GeForce9600GTで40GB/s弱(効率7割弱)のメモリバンド幅。あ、じゃもうでき
るじゃんとか勝手に思い、Fortranのコードに挑もうとしてあえなく止まってますorz
369:デフォルトの名無しさん
08/12/29 14:08:06
つURLリンク(www.easize.jp)
370:368
08/12/29 14:20:59
>>369
ありがとうございます。m(_ _)m
多次元配列はどのみちGPU上では一次元化するつもりでしたがなるほど。
参考にさせていただきます。
371:デフォルトの名無しさん
08/12/29 14:45:46
>>369
そのページ、ポインタの扱いがメタメタだな
わかってるなら別にいいが
372:デフォルトの名無しさん
08/12/29 14:49:43
うん、ぶっちゃけ「"cuda fortran"でぐぐって一番上に来たリンク貼り付けただけ」なんだ、すまない
373:デフォルトの名無しさん
08/12/29 14:50:01
何か普段使うものをCUDAに移植しようと思いつつ適当な物が見当たらない
374:デフォルトの名無しさん
08/12/30 02:15:09
ぶっちゃけ、俺もそうなんだよね。おまけで付いて来たサンプル書き換えて
動かしてFLOPSベンチしてるだけ。
CUDA動かす環境をコスパ良く構築するには?と考えて色々やってみたが、
構築した環境で動かすモノって、結局サンプルの改造ばっか。まんまと
nVidiaの販促戦略に乗せられたぜw
375:デフォルトの名無しさん
08/12/30 08:50:29
zip圧縮とかjpg圧縮とかを移植したらライセンスの関係はどうなるの?
376:デフォルトの名無しさん
08/12/30 12:52:04
圧縮アルゴリズムに関しては問題ないんじゃね?
377:デフォルトの名無しさん
09/01/19 15:08:44
CUDAプログラミングモデルの概要
URLリンク(http.download.nvidia.com)
CUDAプログラミングの基本(パート I - ソフトウェアスタックとメモリ管理)
URLリンク(http.download.nvidia.com)
CUDAプログラミングの基本(パート II - カーネル)
URLリンク(http.download.nvidia.com)
378:デフォルトの名無しさん
09/01/21 03:21:23
デバイス側に確保したメモリにホスト側から
cudaMemcpyのように一括でコピーするのではなく、
一部分だけコピーする、又は書き換える良い方法がありましたら教えてください。
379:デフォルトの名無しさん
09/01/21 03:28:19
・cudaMemcpy()で一部分だけコピーする。
・cudaMemset()で(ry
・そういうカーネル関数を用意して呼び出す。
380:デフォルトの名無しさん
09/01/21 12:27:16
一部分だけってのはたいてい順次処理になるからCPUにやらせたほうが有利だよ
381:デフォルトの名無しさん
09/01/22 01:33:54
CUDAのプログラミングでZIPアーカイブのパスワード解析とか、早くなりませんかねぇ。
エンコード/デコードに使えるんだから、どうなんでしょう?
382:デフォルトの名無しさん
09/01/22 06:38:19
CUDAのfortranサポート予定ってGPUカーネルの部分をfortranライクに
書けるようになるって理解でいいのかな?2.0からって見かけたけど、それっぽい記述が全く無い…
383:デフォルトの名無しさん
09/01/22 08:47:26
>>379-380
ありがとうございました。
参考にさせていただきます。
384:デフォルトの名無しさん
09/01/22 19:18:17
>>381
総当りか辞書型か知らんが演算能力よりI/Oの問題だろ普通。
385:デフォルトの名無しさん
09/01/22 20:12:54
>>381
早くなるよ
386:デフォルトの名無しさん
09/01/23 00:20:22
>>384
ファイル数が多い場合はI/Oも問題になるかもしれませんが、
ある程度のサイズであればWindowsでもメモリに読み込むので、
それほど問題にならないのでは?
毎回物理的に読みに行くわけではないし。
387:デフォルトの名無しさん
09/01/26 09:28:38
ファイルI/OじゃなくメモリI/Oなのでは?
388:デフォルトの名無しさん
09/01/28 11:49:27
>>387
お前は何を言っているんだ
389:デフォルトの名無しさん
09/01/28 12:47:12
>>388
pciバスつかうんだろ?
立派なioとおもうが
つかクーダってほんとマイナーなんだね。
390:デフォルトの名無しさん
09/01/28 13:18:14
そりゃNVIDIA限定だからだろ?グラボを選ぶのなら汎用的なアプリは作れない。
CUDAよりもっといろんなグラボに対応しているOpenCLへのつなぎに過ぎないよ。
391:デフォルトの名無しさん
09/01/29 00:08:14
ZIPのパスワード解析にZIPファイル全体へのアクセスが必要だと勘違いしてる馬鹿w
392:384
09/01/29 08:00:22
>>381-391
ああ、俺がぼけてた。なんか知らんけど自分でもなんであんなレスつけたんだろ orz
わけわからない流れにしちまって、すまん。
正しくはこうだな。↓
パス解析では総当りにしても辞書にしてもforなどのループ速度が求められるだけで、
CUDAによる計算は意味があるのかい?と思った。
MD5とかデータから導き出す数値ならば本当の意味での解析だから意味ありそうだけど。
393:デフォルトの名無しさん
09/01/29 23:02:34
こんな記事もあるしな
GPUを利用した無線LANのパスワードクラッキングが主流に
URLリンク(itpro.nikkeibp.co.jp)
394:デフォルトの名無しさん
09/01/30 03:18:20
>>393の言ってるのはまさに本当の意味での解析ってやつだしな。
WPA/WPA2-PSKからの復元は総当り/辞書比較と違ってGPGPUにする意味がある。
395:デフォルトの名無しさん
09/02/01 15:41:49
Catalystだけど、トリップ検索は随分速くなるみたい
URLリンク(sourceforge.jp)
396:デフォルトの名無しさん
09/02/02 21:51:40
ホスト側でconstantメモリを動的に確保できないのでしょうか?
方法がありましたら教えてください。
397:デフォルトの名無しさん
09/02/02 22:44:30
>>396
目的が良く分からないのでなんとも言えませんが、
どうせ64KBしかないのだから静的に64KB確保しておくのではだめなのでしょうか。
398:デフォルトの名無しさん
09/02/02 22:54:45
>>396
あるけど、なんでprograming guide読まないの??
399:デフォルトの名無しさん
09/02/02 23:59:28
>>397
静的に確保しておいてもよかったんですが,綺麗にかけるならとw
>>398
一度目を通したつもりでしたが,素通りしていたようです。
もう一度読み直してきますm(_ _)m
400:デフォルトの名無しさん
09/02/06 00:08:13
誰かZIPの暗号解析ツール作って><
それか解析のループ部分だけうpしてくれたら、カーネルは俺が書きます。
401:デフォルトの名無しさん
09/02/06 06:54:02
総当りするだけならどっかのオープンソース実装の
パスワードチェックと復号ルーチンとってくればいいだろ
402:デフォルトの名無しさん
09/02/06 10:20:56
ただでさえ消費電力がデカイのにパスワード解析に何週間も動かしてられっかw
403:デフォルトの名無しさん
09/02/07 16:57:19
超解像のような考え方は昔からあるようだな
フリーソフトImageD2
URLリンク(www.tiu.ac.jp)
これは時間軸方向にも参照するものだ。
NVIDIAのMOTIONDSPと同じ考え方だな。
西川善司の大画面☆マニア 第113回 超解像
URLリンク(av.watch.impress.co.jp)
この方法は数フレームを参照することで低解像な映像のブレから情報量をアップさせるようだが、
この方法だとシーンチェンジやめまぐるしく動く映像では逆効果でめちゃくちゃな映像になるのではないか?
(これは上記のフリーソフトの別ページでも注意点として載っていた)
でも東芝などの日本の各社がやってる1フレームだけで行う超解像はそもそも無理がある。だから不自然な画質になったり、情報量が逆に消失したりする。
Lanczosなどでそのままアプコンしたほうがずっと情報量あるし自然な画質だ。比較してみれば一発で分かる。
URLリンク(plusd.itmedia.co.jp)
ここの元画像を720×480にし、それをAVIUTLなどでLanczosでフルHDにしたもののほうがずっと綺麗。
超解像は単純な処理だから柵とか崩れてるし、文字も駄目だし、元からあった情報を処理によって消しちゃう副作用のほうが強い。
超解像、超解像と目新しくいって盛り上げようとしたいのは分かるが、こんなのはまやかしだよ。
URLリンク(www1.axfc.net)
比較用画像もうpしておいた
一方、数フレームでやる方法は計算が大変だが、シーンチェンジや盛大な動きの問題さえクリアすればかなり使えそうではある。
MOTIONDSPや↓はそのあたりちゃんとクリアしているのだろうか?
URLリンク(www.flashbackj.com)
404:デフォルトの名無しさん
09/02/07 17:01:19
10年ぐらい前にカーネギメロン大学で勉強してた頃に超解像のプログラム作れっていう課題を出されたことがある
405:デフォルトの名無しさん
09/02/07 17:34:44
ちょー解像は判ったからマルチすんなや。
406:デフォルトの名無しさん
09/02/07 20:13:09
>>403
スレ違いじゃない?
GPGPU 向きな処理なのは確かだし、そもそも Cell なり SpursEngine は GPU じゃないし。
画像処理スレとかあったと思うよ?
次ページ最新レス表示スレッドの検索類似スレ一覧話題のニュースおまかせリスト▼オプションを表示暇つぶし2ch
5327日前に更新/215 KB
担当:undef