- 1 名前：a36 ◆K0BqlCB3.k [2008/12/10(水) 15:38:25 .net]
- さてついにOpenCLの仕様が公開されました。
www.khronos.org/opencl/
公式ページにはAPIのヘッダファイルが公開されており、
まだ実際に動かす事はできないもののプログラミングすることは可能となっています。
ということで、公開に先んじてプログラミングを始めてしまいましょう。
- 516 名前：デフォルトの名無しさん mailto:sage [2011/10/17(月) 07:38:26.26 .net]
- C++のラッパーを最近使い始めたけどかなり使いやすいね。
とくにメモリ解放が楽になった。
- 517 名前：デフォルトの名無しさん mailto:sage [2011/11/02(水) 21:22:05.07 .net]
- OpenCLは構造体のメモリオブジェクトを作成できますか？
CUDAだったら
typedef struct {
float *num;
} DATA;
DATA data;
cudaMalloc( &data.num, sizeof(float) * 1024 );
みたいにできるんですけど
- 518 名前：デフォルトの名無しさん mailto:sage [2011/11/02(水) 23:38:01.05 .net]
- >>512
例が意味不明
それじゃGPU上に単なるfloat型の配列を確保して、
そのdeviceポインタをCPUの構造体メンバに代入しているだけ。
OpenCLでもGPU上にfloat型の配列をbufferとして確保して、
それをCPUの構造体メンバに代入することは出来る。
メンバ変数の型はfloat*では無く、cl_memだけどね。
でも、その配列自体をGPU側にコピーして、
間接アクセスしようとするとCUDAとOpenCLでは全く違う。
OpenCLではGPU上のポインタは1つのカーネル呼び出し内でしか
一貫性が保証されないから、ポインタを保存しておいて
次のカーネル呼び出しで使うという事が不可能
(cl_memはハンドルに過ぎず、OpenCLのランタイムは
GPU上のオブジェクトを再配置する可能性があるから)
まあ、やるとしたら大きなbufferをメモリプールとして生成して、offsetをポインタの代わりに保存
後でアクセスするときにはbufferを引数に渡して、それにoffset足した位置を触るという
形にするしかない。
まあ、CUDAでも実際にはこの方がCPUとGPUで同じデータを扱えるし、
性能面でも悪くないやり方なんだけど。
- 519 名前：デフォルトの名無しさん mailto:sage [2011/11/03(木) 01:23:08.95 .net]
- ラデ外付けGPUに大量にデータを送りたいんだけど、1/4までって制限どうにかならないの？256MBまでしか送れん
- 520 名前：デフォルトの名無しさん mailto:sage [2011/11/03(木) 10:05:29.53 .net]
- いや、1/4なのではなく 256MBまでという制限。
sizeof(float4) * 4096 * 4096.
- 521 名前：デフォルトの名無しさん mailto:sage [2011/11/03(木) 10:38:23.81 .net]
- >>515
これって変えられないの？
- 522 名前：デフォルトの名無しさん mailto:sage [2011/11/04(金) 11:46:02.78 .net]
- 馬鹿には無理
- 523 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 02:56:58.68 .net]
- 並列化させるのも結構苦労するよね。
簡単な演算ならいいけど、データ依存がちょっとでも複雑になると
動かすカーネルの順番とか数とか
気にしなきゃいけないことｲﾊﾟｰｲ
- 524 名前：デフォルトの名無しさん mailto:sage [2011/11/15(火) 23:57:17.03 .net]
- OpenCL 1.2
www.khronos.org/news/press/releases/khronos-releases-opencl-1.2-specification
- 525 名前：デフォルトの名無しさん mailto:sage [2011/11/16(水) 00:44:39.04 .net]
- NVIDIAのドライバがカオスになるな
ｇｄｇｄの果てに漸く1.0対応が落ち着いたと思ったら
28x世代の1.1対応でまたおかしくなってまだ終息してないのに
- 526 名前：デフォルトの名無しさん mailto:sage [2011/11/16(水) 01:54:02.46 .net]
- 彼らにはCUDAがあるからなぁ。。
頑張る必要が無いのだろう。
- 527 名前：デフォルトの名無しさん mailto:sage [2011/11/19(土) 02:59:11.80 .net]
- nVidiaはどこに向かっているんだ…
OpenACC : 新しい並列コンピューティングのためのプログラミング環境
www.shader.jp/?p=466
- 528 名前：デフォルトの名無しさん mailto:sage [2011/11/19(土) 19:26:18.05 .net]
- それはOpenMPのGPU版みたいなものなので、CUDAやOpenCLとは衝突しない
- 529 名前：デフォルトの名無しさん mailto:sage [2011/11/19(土) 20:50:05.10 .net]
- AlteraがFPGAでOpenCLを、とか言い出してて面白そうな感じ。
- 530 名前：デフォルトの名無しさん mailto:sage [2011/11/19(土) 21:07:31.87 .net]
- C++AMPのNVIDIA版ぽいね
まぁC++AMPはNVIDIAでも動くけど
インライン記述の世代でどれが主導権を握れるかは気になるところ
- 531 名前：デフォルトの名無しさん mailto:sage [2011/11/19(土) 23:38:32.61 .net]
- MSは最終的にC++AMPをオープンにすると言ってはいるが、結局オプソ系コミュニティがどう動くかだな
CPUの並列ライブラリはMSはOpenMPからPPL推し、オプソ系は主にTBBと別れてしまっているので、
GPUではどうにかして歩調を揃えてもらいたいところ
ただ、一応オープン化を標榜するC++ AMPに介入するわけでもなくかといってCUDAの様に自社GPU専用に囲い込むわけでもなく、
立ち位置の被るオープン規格を立ち上げたNVIDIAの意図が分からんといえば分からん
GPUに全てを賭けるメーカーとしては握れる手綱は全て握っておきたい、という事なのかな
- 532 名前：デフォルトの名無しさん mailto:sage [2011/11/20(日) 01:50:41.74 .net]
- >>523
使う側は思いっきり衝突するだろ
- 533 名前：デフォルトの名無しさん mailto:sage [2011/11/20(日) 02:32:04.15 .net]
- >>522
PGI Accelerator が元になっているのかな?
www.softek.co.jp/SPG/Pgi/Accel/index.html
- 534 名前：デフォルトの名無しさん mailto:sage [2011/11/22(火) 14:27:24.23 .net]
- インタビューで簡単にOpenACCについてふれてる
insidehpc.com/2011/11/21/cuda-reaches-5th-birthday-openacc-ramps-up/
- 535 名前：デフォルトの名無しさん mailto:sage [2011/11/22(火) 23:50:42.88 .net]
- カーネルの実行順位はイベントで指定できる
OpenCLのバイナリコンパイルと読み込みうまくできない・・・
バイナリなしだと環境自由になるけど
ソース丸出しになるから計算高速化くらいしか使い道ないね
- 536 名前：デフォルトの名無しさん mailto:sage [2011/11/25(金) 00:37:00.99 .net]
- 1.2の新機能
www.streamcomputing.eu/blog/2011-11-19/difference-between-opencl-1-2-and-1-1/
- 537 名前：デフォルトの名無しさん mailto:sage [2011/11/25(金) 22:22:16.51 .net]
- >>531
ダイナミックに追加が変更があったのって、DirectX関係だけだな。。。
- 538 名前：デフォルトの名無しさん mailto:sage [2011/11/28(月) 10:27:25.45 .net]
- なんかもーGPUメーカーは独自に動いてるし、開発する気無いだろｗ
- 539 名前：デフォルトの名無しさん mailto:sage [2011/11/29(火) 19:35:56.69 .net]
- intel CPU制限多すぎ・・・SSE4.1対応って書いておいてくれよ・・・
core2 quad全部対応してるかのようなのはやめてほしい・・・
XPで使えないのはちょっと困る・・・
- 540 名前：デフォルトの名無しさん mailto:sage [2011/11/29(火) 23:11:26.28 .net]
- SSE見逃してたごめん
- 541 名前：デフォルトの名無しさん mailto:sage [2011/11/30(水) 00:12:08.31 .net]
- SSEぐらいOpenCL使わなくってもいいじゃん
- 542 名前：デフォルトの名無しさん mailto:sage [2011/11/30(水) 00:50:42.75 .net]
- >>536
OpenCLの機能実現するのに都合がいい命令がSSE4.1にあるからインテルの開発ツールはSSE4.1対応の世代以降でないと使えないんだよ
どうせ普及する頃には古い世代のCPUいなくなってるよねって方針なんだろ
- 543 名前：デフォルトの名無しさん mailto:sage [2011/11/30(水) 00:58:02.95 .net]
- GPUが使えない環境での互換性用と割り切って広くサポートしてくれたほうがまだ有用なのにな
- 544 名前：デフォルトの名無しさん mailto:sage [2011/12/05(月) 06:55:31.93 .net]
- 構造体そのままカーネルに放り投げられない・・・
x.yに配列分けなくちゃだめか・・・
GPU正直CUDAのほうが楽だよね・・・
CPUとGPU同時並列に魅力感じてたけど
両方ともスレッドが少なすぎて・・・
- 545 名前：デフォルトの名無しさん mailto:sage [2011/12/05(月) 09:48:13.49 .net]
- なにいってるんだ？
- 546 名前：デフォルトの名無しさん mailto:sage [2011/12/05(月) 11:51:31.75 .net]
- 構造体 a
int x;
int y;
の配列をそのままメモリバッファにコピーしても
うまくいかなかったので・・・
AMDのカーネルアナライザーでエラーがでてたんです
if(a[id].x-a[id].y){}
結局配列分割してコピーしました・・・
- 547 名前：デフォルトの名無しさん mailto:sage [2011/12/05(月) 12:04:37.78 .net]
- 組み込みのint2でいいじゃない。
- 548 名前：デフォルトの名無しさん mailto:sage [2011/12/05(月) 16:36:44.62 .net]
- そうですね
それで組み込んでみます
- 549 名前：542 mailto:sage [2011/12/05(月) 21:49:58.78 .net]
- >>543
まず無いだろうけど、別スレッドで
同じ添え字のxとyをバラバラに更新すると嵌るよ。
- 550 名前：543 mailto:sage [2011/12/06(火) 04:05:54.65 .net]
- >>544
xを固定してyの値すべて計算して次のxへという
九九を生成するようなマニアックな使い方をしてるんです・・・
- 551 名前：542 mailto:sage [2011/12/06(火) 07:51:19.62 .net]
- それなら分けた方がいいと思う。
- 552 名前：543 mailto:sage [2011/12/06(火) 09:23:20.52 .net]
- そうしてみます！
- 553 名前：デフォルトの名無しさん mailto:sage [2011/12/07(水) 22:41:57.13 .net]
- ふう・・・ついに完成しました
ちまちま25%使うよりフルロードはいいですね
ただ、オーバーヘッドがいくらあるからはわかりませんが・・・
- 554 名前：デフォルトの名無しさん mailto:sage [2011/12/08(木) 07:31:00.39 .net]
- 256MBの制限に引っかかって処理が止まるorz
これの上限増やせないのか?
- 555 名前：デフォルトの名無しさん mailto:sage [2011/12/08(木) 10:09:58.25 .net]
- >>549
グローバルメモリが足りない・・・のかな？
ttp://www.ozone3d.net/gpu_caps_viewer/
これで確認してみて
それかワークスレッドの設定が悪いのかも
- 556 名前：デフォルトの名無しさん mailto:sage [2011/12/08(木) 11:03:39.84 .net]
- >>550
なんでダメか分かった、ありがとう
ローカルメモリ32kbしかないのに2mb使おうとしてたwww
- 557 名前：デフォルトの名無しさん mailto:sage [2011/12/08(木) 14:28:54.91 .net]
- >>551
使い捨てる変数の宣言くらいがちょうどいいよ
NVIDIAはローカルメモリ使わないと倍ぐらい遅くなるけど
他はそうでもないから(OpenCL入門に比較があったよ)
グローバルで問題ないと思うよ
- 558 名前：デフォルトの名無しさん mailto:sage [2011/12/15(木) 20:39:38.90 .net]
- www.4gamer.net/games/076/G007660/20111214033/
おまいらこれからもOpenCLにしがみついて行くの？
- 559 名前：デフォルトの名無しさん mailto:sage [2011/12/15(木) 20:44:38.54 .net]
- CUDAはGPUしかできないけど一度により多くの処理ができる
OpenCLはCPUとGPUを同時並列処理ができるのが魅力
どっちにも特化した特性があるからプログラムしだいだよ
- 560 名前：デフォルトの名無しさん mailto:sage [2011/12/15(木) 22:16:31.17 .net]
- GPUには留まらない図が載ってるよ、ヌフォの所に
- 561 名前：デフォルトの名無しさん mailto:sage [2011/12/15(木) 22:40:09.72 .net]
- もうGPU固有じゃなくなってきたのね・・・
カーネル丸出しのCLは論文すら少ない・・・
- 562 名前：デフォルトの名無しさん mailto:sage [2011/12/26(月) 03:28:48.61 .net]
- 引数の渡し方が面倒なんだよなー
思わぬところでバグが出たりする
- 563 名前：デフォルトの名無しさん mailto:sage [2012/01/18(水) 10:01:49.46 .net]
- openclを実用的に使うにはどんな環境がおすすめでしょうか
- 564 名前：デフォルトの名無しさん mailto:sage [2012/01/18(水) 10:11:34.51 .net]
- いつでもCUDAに逃げられる環境
- 565 名前：デフォルトの名無しさん mailto:sage [2012/01/18(水) 22:05:36.88 .net]
- nVidiaってこと？
- 566 名前：デフォルトの名無しさん mailto:sage [2012/01/19(木) 09:51:18.68 .net]
- グラフィックボードは倍精度不動小数点数（cl_amd_fp64）が利用可能な
Radeon HD 7900、6900、5900/5800 シリーズがおすすめ
blog.tommy6.net/archives/74
- 567 名前：デフォルトの名無しさん mailto:sage [2012/01/19(木) 16:52:31.76 .net]
- 宗教戦争が始まりそうだな
- 568 名前：デフォルトの名無しさん mailto:sage [2012/01/19(木) 17:05:38.23 .net]
- >>561
紹介thx
- 569 名前：デフォルトの名無しさん mailto:sage [2012/01/19(木) 18:25:09.26 .net]
- どうせおまえらめんどくさい組み込み関数なんか使わないだろってことか？
- 570 名前：デフォルトの名無しさん mailto:sage [2012/01/23(月) 19:34:34.31 .net]
- GCNでアセンブリが変わったのか、動かなくなってしまった・・・
- 571 名前：デフォルトの名無しさん mailto:sage [2012/01/25(水) 02:58:13.00 .net]
- 変わる以前にGCN対応ドライバ＆ランタイムまだ出てないでしょ
CCC12.1Previewが出回ってるからそれ付属のSDKランタイムなら動くんじゃないの
- 572 名前：デフォルトの名無しさん mailto:sage [2012/01/25(水) 03:08:02.37 .net]
- Kernel AnalyzerもTahiti対応版出てないしな
- 573 名前：デフォルトの名無しさん [2012/02/01(水) 22:06:59.00 .net]
- RADEONとGeforce、ガチンコ対決ではごっちが速いの？
- 574 名前：デフォルトの名無しさん mailto:sage [2012/02/01(水) 22:33:40.43 .net]
- ごっちかな
- 575 名前：デフォルトの名無しさん mailto:sage [2012/02/01(水) 23:52:53.73 .net]
- そもそもガチンコな応用でOpenCLってまともに実績あるの?
- 576 名前：デフォルトの名無しさん mailto:sage [2012/02/02(木) 00:02:24.64 .net]
- 無ければなんだっての？
- 577 名前：デフォルトの名無しさん mailto:sage [2012/02/02(木) 15:35:32.54 .net]
- いや別に問題ない
- 578 名前：デフォルトの名無しさん [2012/02/10(金) 23:47:32.03 .net]
- とりあえず齧る分には公式のプログラミングガイド買っとけばおｋ？
- 579 名前：デフォルトの名無しさん mailto:sage [2012/02/13(月) 02:50:23.08 .net]
- OK
C言語と同じ
- 580 名前：デフォルトの名無しさん mailto:sage [2012/02/22(水) 23:40:42.01 .net]
- サブルーチンみたいにカーネルに直接引数渡して処理できればなぁ
アドレスを渡す時とかすげえ面倒。
- 581 名前：デフォルトの名無しさん mailto:sage [2012/02/23(木) 01:38:17.77 .net]
- そんな気軽にあちこちで使うようなもんじゃないっしょ
- 582 名前：デフォルトの名無しさん mailto:sage [2012/03/08(木) 16:45:33.65 .net]
- カーネルソース丸出しか
特定デバイス用にコンパイルしておかないと
いけないのがなあ
- 583 名前：デフォルトの名無しさん mailto:sage [2012/03/08(木) 16:57:57.20 .net]
- 大したことないやつほど隠したがる銭湯の法則
- 584 名前：デフォルトの名無しさん mailto:sage [2012/03/08(木) 22:38:56.05 .net]
- BOINC的なもので使うと不正対策必要で、既知の答えも一緒に計算させたりして対策するんだけどエコじゃないて悩ましいよ
- 585 名前：デフォルトの名無しさん mailto:sage [2012/03/11(日) 01:23:31.48 .net]
- それなら同じWU他のやつに配って結果一致したやつだけ採用、
合わなかったら更に配布して一致した方採用でいいんでないかね
どのみち普通にやったって計算エラーになるのもいるし
- 586 名前：デフォルトの名無しさん mailto:sage [2012/03/11(日) 02:45:56.59 .net]
- つーかカーネルがソースコードだから
改竄される恐れがなんて懸念するような
プログラムなら、どんな形態で配るにせよ
ディスアセンブルされりゃ同じだって。
- 587 名前：デフォルトの名無しさん mailto:sage [2012/03/17(土) 16:17:03.96 .net]
- 個人PCでスパコンの1/1000の計算速度（10テラFLOPS）だってよ。スゲーな。
dualsocketworld.blog134.fc2.com/blog-date-201202.html
- 588 名前：デフォルトの名無しさん mailto:sage [2012/03/18(日) 17:23:23.96 .net]
- 10TFLOPSっていうと10年前のスパコンの性能と同じくらいだな
- 589 名前：デフォルトの名無しさん mailto:sage [2012/03/18(日) 20:53:01.85 .net]
- 10年後に京の性能がご家庭に来るかというと、さすがにそうでもない感。
- 590 名前：デフォルトの名無しさん mailto:sage [2012/03/18(日) 21:28:00.64 .net]
- 物理的限界が近づいてきたからこれまでとは事情が違うよね。
無知な人は技術でなんとかなるって言っちゃうだろうけどw
- 591 名前：デフォルトの名無しさん mailto:sage [2012/03/18(日) 23:15:06.83 .net]
- 一般家庭にベクトル演算ってそんなに必要ない気が。。。
- 592 名前：デフォルトの名無しさん mailto:sage [2012/03/19(月) 10:47:53.21 .net]
- O(n^2)の直接法のN体とか本当にベンチマーク以上の意味はないんだがな
実用コードはツリー法やFMM法を使う
- 593 名前：デフォルトの名無しさん mailto:sage [2012/03/19(月) 17:14:11.16 .net]
- O(n^2)がツリー法だとO(n logn)に、FMMだとO(n)になるそうだ
www.kurims.kyoto-u.ac.jp/~kyodo/kokyuroku/contents/pdf/1084-12.pdf
- 594 名前：デフォルトの名無しさん mailto:sage [2012/04/13(金) 00:11:16.74 .net]
- AdobeのCS6でOpenCLが使われているという噂を聞いたんだが、カーネルのソースコードとかどうやってるんだろうか
- 595 名前：デフォルトの名無しさん mailto:sage [2012/04/13(金) 03:29:06.78 .net]
- >>589
Adobeのライセンス認証部分のコードだって、一緒にディスクに入ってるわけですよ。
どうにでもなる。
- 596 名前：デフォルトの名無しさん [2012/04/29(日) 15:57:09.88 .net]
- Intel OpenCL SDKを使って開発をしようとしているのですが、CPUとGPUを非同期で
走らせる方法が分かりません。
サンプルなどないでしょうか?
- 597 名前：デフォルトの名無しさん mailto:sage [2012/04/29(日) 20:57:55.38 .net]
- >>587
必ずしもそうでもない。もちろん限られたケースだが、空間ごとに時間刻みを変えたい場合とか、ツリー方とかだと複雑になりすぎてまだまだ難しい場合も多い。
そもそも、ツリーやFMMでも近傍場の計算は直接法だし、サンプルコードとして無意味ではない。
- 598 名前：デフォルトの名無しさん [2012/05/02(水) 10:33:58.88 .net]
- >>591
CPU側を並列化してGPUを扱うスレッド/プロセスとCPU側の計算をするものに分ける
並列化、同期はお好みの方法で
- 599 名前：デフォルトの名無しさん mailto:sage [2012/06/27(水) 01:05:42.85 .net]
- 質問なんだけど今はAMD/nVidia/IntelどのSDKでビルドしても
AMD/nVidiaのドライバが入っている環境ならGPUが利用できるって理解でok？
- 600 名前：デフォルトの名無しさん mailto:sage [2012/06/27(水) 21:29:31.15 .net]
- >>594
違うというか、根本的に違う
- 601 名前：デフォルトの名無しさん mailto:sage [2012/06/30(土) 22:46:38.04 .net]
- >>594
理論上は合っているよ。
まあ、微妙に非互換だったり、ランタイムの
インストール方法など気を使わざるを得ないのが現実だが。
- 602 名前：デフォルトの名無しさん mailto:sage [2012/07/01(日) 01:24:59.44 .net]
- >>595
- 603 名前：デフォルトの名無しさん mailto:sage [2012/07/01(日) 12:56:18.02 .net]
- >>595
- 604 名前：デフォルトの名無しさん mailto:sage [2012/07/05(木) 01:16:06.38 .net]
- GT440でclinfoしたら
Max compute units: 2
と出たんだけど、この２とGT440のCUDAコア96とはどういう関係にあるの？
- 605 名前：デフォルトの名無しさん mailto:sage [2012/07/05(木) 12:03:19.25 .net]
- NVIDIAの場合
Max Compute unit = Streaming Multiprocessor(SM) = Streaming Multiprocessor eXtreme(SMX)
- 606 名前：デフォルトの名無しさん mailto:sage [2012/07/27(金) 10:28:17.75 .net]
- CPUデバイスのランタイムを有効にするにはAMDかインテルのSDKをインストールしてもらうしかないの？
ランタイムだけのパッケージってある？
- 607 名前：デフォルトの名無しさん mailto:sage [2012/07/27(金) 11:10:16.02 .net]
- >>601
Intel の Windows 用ならあるね。
AMD は知らん。
- 608 名前：デフォルトの名無しさん mailto:sage [2012/07/28(土) 14:23:52.35 .net]
- software.intel.com/en-us/articles/vcsource-tools-opencl-sdk/
のダウンロードリストボックスのOpenCL CPU only runtime(Windows* OS)ってやつか。
見落としてた。ありがとう。
- 609 名前：デフォルトの名無しさん mailto:sage [2012/07/30(月) 20:27:19.07 .net]
- CPUでCLやっても遅いだけで意味無いなぁ
- 610 名前：デフォルトの名無しさん mailto:sage [2012/07/30(月) 22:03:29.12 .net]
- >>604
どんなプログラムなのかしらないけど、ボクが試したのは充分速かったけどなぁ。
いやまあ、GPUと絶対値を比べればもちろん遅いけど。ｗ
- 611 名前：デフォルトの名無しさん mailto:sage [2012/07/30(月) 22:41:12.50 .net]
- CPUでOpenCL使うとお手軽にSIMD+マルチコア使えるな。
- 612 名前：デフォルトの名無しさん mailto:sage [2012/07/30(月) 22:48:37.20 .net]
- ベンチも取らずに知ったか
- 613 名前：デフォルトの名無しさん mailto:sage [2012/07/31(火) 00:12:55.35 .net]
- ベンチと言えば Intel と AMD で比べてみたら、Intel のほうがかなり速かった。
いやまあ、その PC の CPU は Intel だったんだけど。ｗ
AMD な CPU だと逆転したりするのかなぁ。
- 614 名前：デフォルトの名無しさん mailto:sage [2012/07/31(火) 07:48:58.72 .net]
- >>606
お手軽じゃねーっつの。
CPUでSIMDとマルチコアを使いたいためだけにOpenCLを使つかうならアホだわ。
- 615 名前：デフォルトの名無しさん mailto:sage [2012/07/31(火) 10:50:46.32 .net]
- phoronixのベンチ見る限りだとAMDのSDKの方が速そうだったけど
- 616 名前：デフォルトの名無しさん mailto:sage [2012/07/31(火) 18:30:47.41 .net]
- 少なくともカーネル部分はお手軽だよ。
素のCで使うと準備が面倒だが
C++のラッパーなら大した事は無い。
|
 |
