追記06/09

http://int.main.jp/files/waifu2x-converter_x64_staticgcc.exe

本家のブロック分割をマージした。大きいファイルでも動くはず。

その影響でちょっと遅くなったのでそれをリカバーするためにOpenCVごとGCCでビルドしなおした。

960x540 でブロック分割しないと10.6[sec]、512x512のブロック分割ありで11.3[sec]ぐらい。


VCは、

v00 = MADD<have_fma>(_mm256_loadu_ps(&w[0*VEC_WIDTH]), i00, v00);
v01 = MADD<have_fma>(_mm256_loadu_ps(&w[0*VEC_WIDTH]), i01, v01);

v00 = MADD<have_fma>(_mm256_loadu_ps(&w[1*VEC_WIDTH]), i01, v00);
v01 = MADD<have_fma>(_mm256_loadu_ps(&w[1*VEC_WIDTH]), i02, v01);

v00 = MADD<have_fma>(_mm256_loadu_ps(&w[2*VEC_WIDTH]), i02, v00);
v01 = MADD<have_fma>(_mm256_loadu_ps(&w[2*VEC_WIDTH]), i03, v01);

こういうコードで _mm256_loadu_ps(&w[X*VEC_WIDTH]) のロードをスケジューリングしてしまって、レジスタスピルする悲しいコードを出していた。(GCCは想定どおりレジスタ14個使うコードを出していた)


32bit版はOpenCVビルド面倒なのでちょっと待って。欲しい人いるなら連絡もらえればやる気出てやると思う。

追記

32bit版 http://int.main.jp/files/waifu2x-converter_x86.exe

あとAVXあればFMAなしでも動くようになってるはず(Sandy以降)。

あんまチューンしてないので期待しないで。

まーチューンしてないというなら、VCがGCCより悲しいコードを出していてそもそもWindows用にチューンできてないという感じはある。GCC+Linuxならさらに20%くらいはやい。

更新06/13

http://int.main.jp/files/waifu2x-converter_x64_20150613.exe

R9 290Xとか使ってる廃人環境向けに変更した。

具体的には最初と最後のイテレーションOpenCLにしている。最初と最後は全体の10%以下かつメモリネックなので、あんまりやる気が出なかったのだけど、R9 290Xとかだと、この部分が30〜40%ぐらいになっていたので、試しに入れてみた。

CPU強くてGPU強くない環境だと遅くなってるかも。実際Kaveri環境だとわずかに悪くなっている。メモリネックなのでキャッシュがでかいほうが有利になっていて、KaveriのGPUの帯域ではそれを挽回できてない。
まあ計測結果見る限りまだ5GB/sしか出てなくて、あと3倍ぐらいはよくできるから真面目にブロッキングすれば不可能ではない問題なのだけど、2時間くらい頑張って良くならなかったからもういいや…


とりあえずcaffe版はGTX780で700x700が1.6秒とからしいので、R9 290Xでそれと同じになるぐらいまではやる(理論性能は40%増しぐらいなので可能なら1秒切るべき)。それができたら、

  • CUDA版
  • DLL化
  • 私もUIつくる
  • 説明

で終わりにするか…

i7-2600K + R9-290X @ Linux で 700x700 ノイズ+拡大を動かした参考値

use GPU: Hawaii
Iteration #1...(6.20262[ms], 47.3417[GFLOPS], 10.8491[GB/s])
Iteration #2...(11.3175[ms], 830.27[GFLOPS], 11.5315[GB/s])
Iteration #3...(14.3321[ms], 1311.26[GFLOPS], 13.6589[GB/s])
Iteration #4...(22.1963[ms], 1693.35[GFLOPS], 11.7594[GB/s])
Iteration #5...(37.765[ms], 1990.54[GFLOPS], 10.3674[GB/s])
Iteration #6...(68.2035[ms], 2204.36[GFLOPS], 7.65402[GB/s])
Iteration #7...(4.055[ms], 289.659[GFLOPS], 64.8717[GB/s])
total : 0.172079[sec], 1701.32[GFLOPS]
start scaling
#1 2x scaling...
Iteration #1...(11.3071[ms], 101.852[GFLOPS], 23.3412[GB/s])
Iteration #2...(36.5203[ms], 1009.11[GFLOPS], 14.0153[GB/s])
Iteration #3...(51.6715[ms], 1426.43[GFLOPS], 14.8586[GB/s])
Iteration #4...(84.0687[ms], 1753.47[GFLOPS], 12.1768[GB/s])
Iteration #5...(133.917[ms], 2201.53[GFLOPS], 11.4663[GB/s])
Iteration #6...(267.429[ms], 2204.87[GFLOPS], 7.6558[GB/s])
Iteration #7...(27.189[ms], 169.429[GFLOPS], 37.9451[GB/s])
total : 0.644018[sec], 1782.86[GFLOPS]
process successfully done! (all:1.21433[sec], 1186.62[GFLOPS], filter:0.776175[sec], 1856.49[GFLOPS])

全体1.2秒で、フィルタ以外が0.5秒とかなので周辺をもう少し頑張ったほうがよい気がする。
(ブロック分割無しで動かした結果。いまのところブロックサイズはコード変更しないと変えられないので、上のexeではこの性能は出ないです)


caffe版の時間内訳わからないけど、これなら同等の効率になってると言っていいんではないかな…

waifu2x はやくした

なんか社内チャットで

https://github.com/WL-Amigo/waifu2x-converter-cpp

をはやくしろというメッセージを受信したのでやった。

https://github.com/tanakamura/waifu2x-converter-cpp

http://int.main.jp/files/waifu2x-converter_x64.exe

コードの大半は、@WL_Amigo さんのを流用していて、コマンドラインの入出力方法は一緒なので、使う場合はexeを置き換えてもらうだけでよいです。@WL_Amigo さんの最新では修正されてるでかい入力でメモリ使い果たす問題の修正はマージできてないので、でかい入力には使えないです。

FMA 使っているので、Haswell 以降のCPUが必須です。なんか32bit版はcl.hpp がエラー出てめんどいから見てない。

ni7-4700MQの8スレッドで、noise + x2 で

入力サイズ 960x540 256x256
もと 167sec 11.3sec
15.6sec 3.2sec
倍率 10.1x 3.5x

ぐらいまでいった。


もとのは、スレッド化してある + OpenCV(つまりSSE化済み)のなので、そこからFMA にしたから、4倍ぐらいは妥当で、そっから2倍ぐらい頑張ったという感じかな…

OpenCL 版は途中までやったけどまだ遅いからもうちょっと待って…

あとけっこう頑張ったので解説もそのうち書きたい。

(なお、この文章は一昨日書いてたのだけどアップロードする直前で @logicmachine がよりはやい版(https://github.com/logicmachine/waifu2x-opt)をだしてきたので今もうちょっと頑張った)

追記06/11 : waifu2x OpenCL 版

http://int.main.jp/files/waifu2x-converter_x64_20150611_2.exe

多分良い感じになったはず。Linuxでしか動作確認していないので動かない可能性は残っている。(Kaveri 用のWindows入れたHDDどれかわからなくなってるので…)

親切な人いたら、動いたかどうか(とわかるなら、GPU名+最後に出てくるFLOPS値)をコメントで教えてもらえると助かります。


NVIDIAGPUで動かすとなんか止まるので、今はAMD GPU でしか有効にしていない。


効率は、理論値 737[GFLOPS] の A10-7850K で、260[GFLOPS](35%) ぐらい出ている。OpenCL部分だけならピークで360[GFLOPS](48%) ぐらい。caffe 版は効率30%〜40%ぐらいに見えるので、同じくらいの GPU なら Radeon でも caffe 版ぐらいの性能になってると予想される。


あと、http://d.hatena.ne.jp/w_o/20150609#1433854645 が気になったので、結局スレッドプール入れた。カスペル付いてる環境を使ってる人もマシになったはず。



色々考えたけど、結局、気合いブロッキングと気合いチューニングだけでなんとかした。

https://github.com/tanakamura/waifu2x-converter-cpp/blob/3843eac74730ac608bcc71a163be1df4307c260e/src/modelHandler_OpenCL.cl#L242

このへんあたりがお気に入りかな…