1read 100read
2013年02月プログラム251: GPGPU#5 (276)
TOP カテ一覧 スレ一覧 2ch元 削除依頼 ▼
【論理】Prolog【初心者】 (626)
Python の宿題ここで答えます Part 2 (344)
awkについて語るスレ $2 (733)
文字コードの種類は何故複数あるのでしょうか? (346)
インテルC++コンパイラ9.0発表! (586)
Gtkプログラミング on Windows!!! (353)
GPGPU#5
1 :2010/08/15 〜 最終レス :2012/11/18 GPGPUについて語りましょう 前スレ GPGPU#4 http://hibari.2ch.net/test/read.cgi/tech/1255256230/l50 関連スレ OpenCLプログラミング#1 http://hibari.2ch.net/test/read.cgi/tech/1228891105/l50 【GPGPU】くだすれCUDAスレ pert3【NVIDIA】 http://hibari.2ch.net/test/read.cgi/tech/1271587710/l50 参考リンク 総本山? gpgpu.org http://www.gpgpu.org/ OpenCL http://www.khronos.org/opencl/ NVIDIA CUDA http://developer.nvidia.com/object/cuda.html ATI Stream http://developer.amd.com/gpu/ATIStreamSDK/Pages/default.aspx GPUをCPU的に活用するGPGPUの可能性 http://pcweb.mycom.co.jp/articles/2005/09/06/siggraph2/
2 : >998 :デフォルトの名無しさん [↓] :2010/08/15(日) 21:51:01 >と言いつつAgeiaの中の人も今じゃAMDにいるからなぁ >とんだ詐欺師なのかねあの人 金です。 nvにとっちゃすでに用済みで、要らない子
3 : 専用設計とはいえPPUは58gflopsしかないんだが
4 : 基本的には、GPGPUが得意な処理を "適切なサイズ" に並列分割して その分割された小包の集団をどかっとCUDAに押し込んでやると、分割が上手ければ それなりに速く結果が出る。ただ、GPGPUで効率が出る並列化は簡単ではない。Larrabeeがこけたのもここ。 しかもC++のCUDA方言は不思議挙動だったりで、技術者がCUDAに習熟して十分な速度が 出せるようになるまでの時間を考えると、結構経費がかかる。だから、相当大きな話、というか CUDAのX86@Intel CPUに対するワットパフォーマンス優位性が技術者の勉強代をカバーできる規模で無いと わざわざわけわからん方言を勉強したくない。しかも、この方言は、いつまで有効かも怪しい。 だから、ほとんどの用途では、Nehalem-Ex とか、速いCPU乗せたマシンを増やした方が良い。 他のプログラムが、"全部" 速くなりますからねw
5 : 今後のCPUコアの高速化が鈍化するから その対策として出てきたのがCPUのマルチコア化と グラボのGPGPUとしての活用なわけで・・・ 大部分の人には上位CPUなんて必要ないのと同様に 大部分のアプリにもGPGPUなんて必要ない。 6コアもGPGPUも本当に必要な人・アプリが使えばいいだけ
6 : 大部分って、静的WEBページを見るだけのユーザーのことか?w そんなもん無視でいいだろw
7 : WEBブラウズだろうがオフィスアプリだろうが 音楽・動画再生だろうがゲームだろうが大部分のアプリには 高価な上位CPUも高速なGPGPUも必要じゃないだろ。 そこそこヘビーな自分でも4コア(疑似8コア)や 1TFLOPS以上のGPUをフル活用できるのは全PC作業の1割程度だし
8 : LAMEとかiTunesとかで、GPGPUが効けばもっと広がると思うんだけど… やる気無いですよねぇ
9 : やる気程度で速くなってくれるなら今ごろみんな取り掛かってるだろうよ
10 : LAME(音声の非可逆圧縮)程度じゃ処理が軽すぎるし 条件分岐も少なくないからCPUで計算したほうがいい。 映像編集ソフトですらエフェクト処理がメインでエンコードにはGPGPUが使えなかったりする。 iTunes(映像再生ソフト)にGPGPUとして使うなんて問題外。 大人しくOpenGLやDirect2DなんかでGPUとして活用すべき。 リアルタイムで映像にエフェクト処理を加えながら再生したいなら別だがiTunesの仕事じゃないw
11 : ATI Stream使ってエンコードして負荷軽減してるソフトなかったけか?
12 : >>11 PowerDirector?
13 : エンコードに使うなら売りは速度ではなく品質にすべき。 データ転送がボトルネックなのだから 単位データあたりの演算量を増やさなきゃメリットが無い。
14 : 演算量が増えてもプログラムのフローが複雑になるようでは
15 : >>14 どんだけ複雑になったって、大量に並列実行できればGPGPUにとってアドバンテージがある。 データに対して演算量が少なすぎると転送や処理待ちばかりになってパフォーマンスが上がらない。 だから問題は複雑性よりもデータの相互依存性とデータに対する演算量の少なさ。
16 : 複雑性ってなに?
17 : 文脈から鑑みるに、プログラムの複雑さじゃないの? もっと端的に言ってしまえば分岐命令の数
18 : この場合、相互依存性と複雑性は同義だと思うけどね。
19 : >>18 この場合は違う。 >14 はそのつもりで発言しているのかも知れないが、>13 は違う。
20 : そう言い切るのなら、どう違うかまでを説明せんといかんよ。
21 : >>13 って8x8DCTを4x4DCTにするみたいな話でしょ? 演算回数は増えるがGPUなら並列数を増やせる感じで
22 : S|A What is AMD's Northern Islands? A look at what is coming in October http://www.semiaccurate.com/2010/09/06/what-amds-northern-islands/ ごめんSIって言ってたけど実はNIだったよ。えへ。 だから今度出るのはHD6000ファミリーはNIね。 32nmでNIテープアウトしてたけど40nmで出すよ。 コアは○○な感じで、アンコアは××な感じで強化してるよ。 なんでチップがEvergreenより10-15%大きくなるよ。 リリーススケジュールは10月12日にイベントで25日前後に店頭並ぶよ。 まずはAMDの穴の開いてる$175-250帯のHD6700から始めるよ。 次にHD6800、HD6900、年初にローエンド、28nmまでこのラインナップだよ。 HD6000出たら緑チームはHD5000よりコスト高いのに値下げしなくちゃだし、それでなくても冷め切ったセールスにもろ影響しちゃうよ。 だって、トップエンドは価格維持でHD5000は下がり始めるしね。 Nvidiaの夢と希望を打ち砕いちゃうね。 打つ手もないしね。 AMDはDX11のトーナメント1回戦をHD5000で勝利して、第2回戦もHD6000で勝利しちゃて、Nvidiaには財務的にもパフォーマンスでっかいマージンを取っちゃうよ。 28nmまではNvidiaにチャンスはないね。
23 : 余所に作らせたGPUを使ったプログラムが、CUDA部分でメモリリークくさいエラーを吐いてまともに動かないんですが、 窓から投げ捨てるべきでしょうか?
24 : 窓から投げるべき
25 : 証拠資料を作ろうとしても、「いつ止まるか」の再現性が微妙 やっぱり実績の無いハウスに委託したのが間違いだったか・・・
26 : メモリの確保と解放を繰り返しているんじゃないかな。 弊社ではソースがあればデバッグも承りますw
27 : ソースないっす・・・ その辺だけはしっかりしているという・・・ ていうか、ウチ(受け入れ側)のマネージャーが完全に「ドモホルンリンクル」で どんなゴミを渡されても「努力あるのみ」とかの類の精神論を吐いて話にならないし どっか、受託開発や納入後の展開方法についての客観的な評価をしてくれる コンサルタントはないですかね・・・
28 : CUDAでソースなし納品はありえんやろ いつバージョンアップでバイナリが動かなくなってもおかしくないのに
29 : >>28 コストの問題だろ? ソースを要求すると価格が上がる
30 : いや、将来動かない可能性が低くないのにコストカットされてもw
31 : 将来動かなくなる可能性が高いから値切るんだろうが
32 : gpgpuを使用した場合、 CPUの性能はどの程度影響しますか? teslaを用いた計算機を導入しようとしているのですが、i7-980xにするかi7-930にするか 迷っています。
33 : CUDAやOpen CL以外のCPUコードの実行速度にモロに影響する。 他にもGPGPU用中間コードのコンパイルにも影響するが誤差範囲。
34 : 聞きたいのはCPUの性能によってGPUの性能が変わるかどうかじゃないの
35 : 初心者なんですけどフリーソフトでATI技術に対応してて MP4に変換できるソフトってありますか? あとRADEONのカードってエンコードなら値段と性能みてどれがコスパいいですか?
36 : 板違いです ソフト板か自作板、DTV板へgo
37 : 板違いです。 ここは「ATI技術に対応しててMP4に変換できるソフト」を作る側の板です。
38 : caymanは期待できそうだな。
39 : GPGPU使って何かしたいけどこれっていう何かが見つからないのー Actor とか Map Reduce とか上位層で駆逐されてしまうねん
40 : 俺はいっぱいアイデアあるけどな。
41 : あら、気になるじゃない。聞きたいわ
42 : 突然申し訳ありません cudaやってるんですけど・・・ カーネル関数起動させるところでエラーが出てしまいます サンプルコードでアウトなんです 考えられる可能性を挙げていただきたいです エロい人助けてください ちなみに、 win7professional32 グラボ1:8600gs(出力用) グラボ2:460gtx(→cuda) 開発環境:visual studio 2008 質問あればできるものはすべて答えますんでよろしくお願いします
43 : >>42 エラーメッセージぐらいのせろやカスが
44 : 密かにevergreenのISA仕様書が更新されているな。 ttp://developer.amd.com/gpu/ATIStreamSDK/assets/AMD_Evergreen-Family_Instruction_Set_Architecture.pdf メモリアクセス周りの挙動について言及されているのがなかなか面白い。 コアレス化が余り重要じゃないという話がどういう意味か分かる。 要は、アーキテクチャ的に1スレッドが複数のメモリアクセス命令を同時発行可能で 1wavefront単位で発行された複数のメモリアクセス命令の間だけ キャッシュ無しアクセスでもキャッシュが有効になっているから 複数のメモリアクセス命令間でコアレス化と同様の効果が得られるらしい。
45 : CAL+ILの情報が少ないので、書き込みがあるだけで嬉しい。
46 : >>43 スイマセン CUDA error: Kernel execution failed コンパイルはできていて、 他のマシンでは同じソースコードのプログラムは動かせます 原因は何なんでしょうか
47 : >>46 カーネルの起動に失敗したというのだから、起動要件を満たしていないのだろう。 まさかとは思うが、CUDA用のドライバをインストールしていないというオチではあるまいな。
48 : >>47 ドライバは入っています それと、今日起動することに成功しました。 main()の変数宣言のすぐ後に、 cudaSetDevice(1); を記述したら、それで通りました。 なぜ起動できなかったかは分かりません。 今可能性を探っているのですが、 タイムアウトが起こったのかもしれないと考えています。
49 : なんだ、動いたなら後はcudaスレへ。
50 : >>48 CUDAの命令はよく分からんけど、名前から察するに単純に処理するGPU指定してなかっただけじゃ。 8400GSだってCUDA対応なんだし。
51 : >>48 >>50 デバイスを指定しなければデフォルトのデバイス0、8600GSで実行されるはずだけど。 Capability 2.0以降限定の関数を呼び出しているサンプルコードだったとか? まあCUDAスレあるし、そっちでやるべきかな。
52 : A Fast GEMM Implementation On a Cypress GPU http://galaxy.u-aizu.ac.jp/trac/note/wiki/Fast_GEMM_Implementation_On_Cypress
53 : Cayman GPUではスーパーファンクションユニットが削除されて5VLIWプロセッサーから 4VLIWプロセッサーになるとのことですが、現在のCALでサポートされているsin/cos等の 超越関数命令は、自分で多項式近似計算をしろと言うことなのでしょうか?
54 : >>53 今までtレーンが担当してきた命令は xyzwの複数レーンで1命令を実行するように変更されている。 で、超越関数は3SPを使った1命令で実行される。
55 : >>54 53 です。 CORDICやチェビシェフ多項式をWeb上で漁っていたのですが 安心しました。
56 : チェビシェフは自分で作るもんじゃろ
57 : CORDICって条件分岐ばっかなのでGPGPUには不向きだという先入観があるんだけどどうなの?
58 : それは全部ソフトウェアでやったら、の話だろ。
59 : おやお久しぶり。 ソフトウェア無線専用ハード(微妙に矛盾?)でCORDICを使ってるとか話には聞いたことがありますな。 あと自分で実装してみて気づいたんだけど、分岐と言ってもある数を足すか引くかなので、 分岐しないようにしてビット操作に落とせるんですよね。
60 : GPUは分岐が苦手とはいっても、単純なプレディケートに落とし込めるものならむしろ効率がいいくらいです。 GPUは同一ワープ内で命令ストリームを共有してますから同じ方向にしか進めない。 Cでいうif-elseは一見分岐だけど、GPUではプレディケート情報によって実行・不実行(あるいは結果に反映させない)を 選択する単一の流れに展開されています。 プレディケート自体はそんなに重たくないです。 むしろ分岐先が増えると増えた分だけ処理時間が増えるだけで。
61 : A's Video Converterって、10/31付けで配布サイトが閉鎖されてるな もう手に入らんの?(´・ω・`)今日HD5670GETしたのに・・・
62 : 移転してるよ http://bluesky23.yu-nagi.com/
63 : >>62 うぉ!マジでありがとう ググっても見つけられなかったんだ(T-T)ウレシイ
64 : プレディケートは現行Intel系GPUでは使い物になりません GPGPU向けに機能追加されたSandy Bridge GPUコアの登場を待ちましょう。
65 : そもそもGPGPUできないし
66 : OpenCL対応するんでしょ>>Sandy Bridge GPUコア
67 : L3使えるから規模の割には速いかもな
68 : 主にGPUコアで回してるかAVXつこてるかはインテル任せ だんごに好かれてしまったからGPGPUも端ッパの技術バリエーションの一つに転落決定だな
69 : AVXでGPUくらい速くできるなら寧ろ大歓迎だが。 ただしアセンブリ言語で書くのは嫌。
70 : そんなもん規模しだいだろ
71 : Sandy Bridgeの1EU=4Way-FMACと仮定しても、まだCPU(AVX)のほうが速いですから
72 : ION乗ってるノートでXP入れました!!!! これでCUDAできますよね? 俺の夢かなえられますよね? ひゃっはあああああああああああああああああああ
73 : IONのCUDAベンチ&レビューよろしく
74 : GPU グレートプログラマー初春
75 : P=>パイパン
76 : DSPやFPGA叩いて高速度・複雑なシステム作るよりは CPU+GPU叩いて作ったほうがはるかにましだがなあ。生産効率が桁違いだわ。
77 : e?
78 : >>76 今までDSPやFPGAでやっていた事をCPU+GPUでできる用途ってどんなものがあるの? 考えたけど一つも思い浮かばなかった。
79 : 具体的なカテゴリは勘弁だが、サンプリングしたデータをフィルタで処理して画面に表示 みたいな処理では、GPUでの代替はかなり強烈だよ。 FPGAだとたかだか数百タップの複素FIRフィルタを40〜50MHzの動作速度でさばくのにも 現状だと5〜15万ぐらいのデバイスがいるし。 大量生産するものなら処理をチップ化して安くあげちゃうんだろうけど、 俺のとこみたいな数のでない無線通信製品だとGPGPUはかなり魅力的。 たぶん、画像検査装置みたいな分野でもGPGPUは強力だと思う。
80 : DX11世代だと、本当に何でも出来そうだよな。
81 : Linux版のAMD APP 2.4にCALのサンプルが付属していないのですが、 Windows版は付属していますか?
82 : してません。 それどころかCALは(IL含め)2.5で死滅。 かわりにLLVM IR使え。 そんな感じです。
83 : >>82 ありがとうございます。 そうですか。。 せっかくIL習得したところなんですが、困りましたね。
84 : http://developer.amd.com/gpu/AMDAPPSDK/assets/AMD_APP_SDK_Release_Notes_Samples.pdf ここにはCALのsampleはcalといディレクトリにあるとかいてますが、 旧バージョンから修正されてないだけでしょうか?
85 : これからGPGPUを勉強する場合、どれを勉強しておくのが良いのでしょうか? 無難という意味では、OpenCLですか?
86 : ソリューションは結局、問題や環境が決定するもの。
87 : CUDAでいいんでない
88 : >>85 ソフトをやるのかハードをやるのか? ソフトの場合は上位をやるのか下位をやるのか?
89 : 500万個の3×3行列の固有値を (1)CPU Intel Q9450 (4 posix threads) GCC 4.4.3 (最適化無し、125万個/スレッド) (2)ATI HD4870 + OpenCL (AMD APP SDK 2.4) (最適化無し) で計算させてみた(行列は正定値実対称の素直な行列)。 ハードはCPU、GPUともに定格で使用。OSはUbuntu x64 10.4 LTS , AMDドライバはCatalyst 11.5 GCC4.4.3とOpenCLで使用したソースコードは略同じものを使用(相違点は OpenCL側コードに__global 指定が付いた程度) 時間測定はC言語側の計算ルーチン呼び出し元でgettimeofday()を使用してマイクロ秒単位で測定。
90 : (1) Q9450 4スレッド ------------------------------------------------------------------ Posix thread creation : 0.000124(sec) Thread #1 Exection time: 5.992(sec) Thread #2 Exection time: 6.08448(sec) Thread #3 Exection time: 5.9132(sec) Thread #4 Exection time: 5.91843(sec) Total Exection time : 6.08452(sec) <ーースレッド中の最大値 + α
91 : (2) HD4870 (800スレッド?) ------------------------------------------------------------------ GPU Kernel Compile : 1.6e-05(sec) GPU Kernel Build : 5.02891(sec) GPU Kernel Creation : 6e-06(sec) GPU Kernel Set Args. : 2e-06(sec) GPU Kernel Execution : 4e-05(sec) --- clEnqueueNDRangeKernel ()を挟むgettimeofday()の時間差 Memory mapping(READ MODE) : 5.38053(sec) <この間でデータ読み出し> Memory UnMapping(from READ MODE) : 0.020504(sec) OpenCLソースのビルド&結果データの読み出しまで含めるとGPUが1.7倍遅いが計算実行時間の単純比較だと 6.08452 / 4.0E-5 = 1.5E5 = 15万倍速い! と言う結果になりました。 いくら何でも15万倍は速すぎのような気が・・・・(^_^;;) 以上。
92 : >90 (自己レス) ごめんなさい。 ミススペルしてました。 (1) Q9450 4スレッド ------------------------------------------------------------------ Posix thread creation : 0.000124(sec) Thread #1 Execution time: 5.992(sec) Thread #2 Execution time: 6.08448(sec) Thread #3 Execution time: 5.9132(sec) Thread #4 Execution time: 5.91843(sec) Total Execution time : 6.08452(sec) <ーースレッド中の最大値 + α
93 : GPGPUについては詳しくないんだけど、 (sizeof float)*3*3*5000000≒180[MB] これがシステムメモリとVRAM間で往復するから360[MB] 所要時間が2[ms]だから、1[s]に180[GB]も動いてることになる 何か変だ
94 : ところで、結果は一致してるのか?w
95 : 89 です。 使った行列は 2.000000E+00, 1.000000E+00, -1.000000E+00 1.000000E+00, 3.000000E+00, 2.000000E+00 -1.000000E+00, 2.000000E+00, 4.000000E+00 ただしデータは対称性の為、(2.000000E+00, 3.000000E+00, 4.000000E+00, 1.000000E+00, 2.000000E+00, -1.000000E+00) の6成分のみで、 システムメモリ〜VRAM間の転送量は 3×3行列 sizeof(cl_float)*6*5000000 = 114MB 固有値 sizeof(cl_float)*3*5000000 = 57MB 固有ベクトル sizeof(cl_float)*9*5000000 = 171MB 反復解法の収束回数 sizeof(cl_int )*5000000 = 19MB CPUでの解は Eigen value e1 = 2.855207E-01, Eigen vector1 = ( 6.345872E-01, -5.961551E-01, 4.918314E-01) Eigen value e2 = 3.143277E+00, Eigen vector2 = ( 7.708385E-01, 5.341269E-01, -3.471549E-01) Eigen value e3 = 5.571202E+00, Eigen vector3 = ( 5.574225E-02, -5.994226E-01, -7.984895E-01) <e1,e2> = 5.960464E-08 ← 固有ベクトル間の内積での直交性チェック <e2,e3> = 0.000000E+00 <e3,e1> = 0.000000E+00 GPUでの解は Eigen value e1 = 2.855215E-01, Eigen vector1 = ( 6.345873E-01, -5.961550E-01, 4.918314E-01) Eigen value e2 = 3.143277E+00, Eigen vector2 = ( 7.708384E-01, 5.341271E-01, -3.471551E-01) Eigen value e3 = 5.571201E+00, Eigen vector3 = ( 5.574221E-02, -5.994227E-01, -7.984894E-01) <e1,e2> = -4.470348E-08 <e2,e3> = -5.960464E-08 <e3,e1> = 0.000000E+00 で略一致してます。
96 : 89 です。 同じ問題を maxima の eigens_by_jacobi で解くと (%i1) A:matrix([2,1,-1],[1,3,2],[-1,2,4]); (%i2) eigens_by_jacobi(A); (%o2) [[0.28552125561224, 3.143277321839643, 5.571201422548121], [ 0.63458730239817 0.77083835004074 - 0.055742207899264 ] [ ] [ - 0.59615502320039 0.53412697029887 0.59942269552653 ]] [ ] [ 0.49183141821965 - 0.347155034107 0.79848934767235 ] (こちらは、固有ベクトルの成分が縦方向に並んでいます)
97 : GPUのほうは最適化の有無でガラっと変わるんでそこんとこどうなのよ
98 : 89です。 rtn = clBuildProgram ( pgm, the_number_of_devices, devices, "-cl-opt-disable", NULL, NULL ); でBUILDしています。 KHRONOSのPDFマニュアル p115に -cl-opt-disable This option disables all optimizations. The default is optimizations are enabled. と記述があります。 またkernel 実行は rtn = clEnqueueNDRangeKernel ( CommandQueue, kernel, 1, NULL, &pe_size, // 5000000 &group_size, // 64 0, NULL, NULL // No triger event will be used. ); 今気がついたのですが、p132に clEnqueueNDRangeKernel returns CL_SUCCESS if the kernel execution was successfully 〜〜 queued 〜〜. Otherwise, it returns one of the following errors: とありました。 "Execution Time" と上で書いた時間は実行キューへの登録時間でした。 お騒がせしてすみませんでした。
99 : 89 です。以下の方法で、Kernel実行時間とメモリマッピング時間の計測が可能であることが分かりましたので再計測してみました。 cl_event event; rtn = clEnqueueNDRangeKernel ( CommandQueue, kernel, 1, NULL, &pe_size /* 5000000 */, &group_size /* 64 */, 0, NULL, &event <- イベント追加 ); if( event ){ (void)clWaitForEvents( 1, &event ); (void)clGetEventProfilingInfo ( event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &device_time_counter[0], NULL ); (void)clGetEventProfilingInfo ( event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &device_time_counter[1], NULL ); (void)clReleaseEvent( event ); } 実行時間 device_time_counter[1] - device_time_counter[0] (nsec); GPU Kernel Compile : 1.5e-05(sec) GPU Kernel Build : 5.02459(sec) GPU Kernel Creation : 6e-06(sec) GPU Kernel Set Args. : 1e-06(sec) *GPU Kernel Execution : 0.114895(sec) *C[114MB] memory mapping(READ MODE): 0.0358828(SEC) 3177.01(MB/sec) *E[ 57MB] memory mapping(READ MODE): 0.0179288(SEC) 3179.24(MB/sec) *V[171MB] memory mapping(READ MODE): 0.0537894(SEC) 3179.07(MB/sec) *iter[19MB] memory mapping(READ MODE): 0.00600078(SEC) 3166.26(MB/sec) *はOpenCLのプロファイリング機能で測定した時間。 それ以外はgettimeofday()を使用して呼び出し元から測定した時間。 結局 6.08452 / 0.114895 = 52.96倍 次期 HD7000 が楽しみになってきました (^_^)。
100read 1read 1read 100read TOP カテ一覧 スレ一覧 2ch元 削除依頼 ▲
Silverlight登場で.NET使い大勝利!!! Part2 (492)
Borland Developer Studio 2006 No.13 (279)
MFC相談室 mfc22d.dll (307)
インテルC++コンパイラ9.0発表! (586)
BASICの宿題はお前にまかせた (553)
C言語なら俺たちに聞け パート0001 (220)
--log9.info------------------
トラックで公道を初めて運転した時の感想 (392)
けん引免許に路上試験導入すべき (502)
ショートキャブの上の寝台 (453)
調子こいてる大型海苔 (278)
トラックにナビ取付け (372)
日産ディーゼル・ビッグサム (267)
【大型一種免許】言い訳番長奮闘記【隔離部屋】 (354)
【タイヤが】低床幅広2セット目【いっぱい】 (464)
【ロータリー】除雪車にハァハァする奴のスレ【プラウ】 (339)
80〜90年代『哥麿会』を語る (403)
【アイチ】高所作業車【タダノ】 (645)
【陸軍】旧日本軍・軍服総合スレ2【海軍】 (714)
◆◆最初に読もう!軍事速報&雑談スレ2257◆◆ (998)
☆前の人が嫌いな戦車(AFV)を褒めるスレ I号戦車 (200)
先島諸島防衛 第二十一 (906)
別宮暖朗スレ その12 (870)
--log55.com------------------
【世界を】。゚+.上松秀実 Part.1。゚+.【変える】
【和製R&B】和田アキ子【女王】
☆森高千里★統一スレ【116】 [無断転載禁止]©2ch.net
■ 邦楽女性ソロ板 案内&自治スレ part1■
ラブリーサマーちゃん [マジレス02]
【バラードの】岡本真夜【女王】
【今日はゆっくり】大野愛果 Part.4【話そう】
YUKI PART 176"Blink Blink"