vpp-delogo 自動フェード/NRのCUDA高速化

NVEnc 4.12で実装したvpp-delogoの自動フェードは、残念ながら遅いという問題があったので、高速化を試みた。

いろいろやってだいぶ速くなったので、その方法とか。

どのくらい遅いの?



ざっくり言うと、450fpsぐらい出ていたエンコードが80fpsぐらいしか出なくなっちゃうぐらいには遅い。

環境
Win10 (17134) x64
i9 7980XE (18C/36T @ 4.0GHz)
GeForce GTX 1060 6GB (1280 cores, 1708 MHz)[397.64]

入力ファイル
mpegts, 1440x1080, 29.97fps, 2696frame (1:30)

自動フェード&NRを入れる前
encoded 2696 frames, 450.31 fps, 6195.04 kbps, 66.43 MB
encode time 0:00:05, CPU: 5.5%, GPU: 8.4%, VE: 84.3%

Vpp Filter Performance
cspconv: 41.6 us
delogo: 8.7 us
cspconv: 41.4 us

自動フェード&NRを入れたあと
encoded 2696 frames, 79.47 fps, 5262.17 kbps, 56.43 MB
encode time 0:00:33, CPU: 5.4%, GPU: 24.8%, VE: 20.1%

Vpp Filter Performance
cspconv: 42.1 us
delogo: 12065.7 us
cspconv: 41.9 us

というわけで、vpp-delogoを入れたとたん、速度がガタ落ちしてしまう。悲しい。



なんで遅いの?



これをなんとか速くするには、遅い原因を探らないといけない。vpp-delogoの処理の大半はGPU上で動くCUDAのコードなので、まず気になるのはGPU使用率なのだけど、これが24.8%と低い。

つまり、GPUでの演算が重いから遅いということではないらしい。

実際、どうなっているのか、NVIDIAのVisual Profilerで確認すると…


クリックで拡大

うーん。思った通りではあったのだけど、GPUで小刻みに計算が行われ、その合間に計算を何もしていない「隙間」の時間が大量にあることがわかった。

この隙間の時間によって、GPU使用率自体は低いのだけど、計算終了まで時間がかかってしまって遅い、という状態になっていると予想できる。

逆に、この「隙間」を詰めてやって、GPUが連続的に動くようにしてやれば、大幅な高速化が期待できる。



処理のどこが原因?



これは答えははっきりしていて、最適なロゴ消しの強さ(フェード値)を求める部分が原因となっていた。

そもそも、最適なフェード値を求める方法は、ざっくり言うとこんな感じ。

1. ロゴ除去の処理をロゴ除去の強さ(フェード値)を変えて行ってみる
2. 1.の処理後の画像でどの程度ロゴが消えている度合い(評価値)をチェック
3. 一番ロゴが消えてるっぽいフェード値が答え

で、フェード値を変えてのところがミソで、ありうるフェード値を総当たりでやっているとさすがに遅い。そこで、オリジナル版では、

ありうるフェード値の範囲を分割して、まずは4つのフェード値でチェック
 ↓
4つの評価結果から、分割した範囲の中から一番よさそうな範囲を見つける
 ↓
その範囲をまた細かく分割して、区切った場所のフェード値でチェック
 ↓
繰り返し…

とやっていくことで、評価回数を抑えて高速化する工夫がされていた。

これをほぼそのままGPUに載せたのだが、ループ等による反復構造を図にするとこんな感じになっていた。


クリックで拡大

反復の中でCPU - GPU間のやりとりが必要となり、相当な回数、CPU - GPU間のやり取りを行っていた。

CPU - GPU間のやりとりはとてもとても遅いので、その時間が隙間の時間となってしまっている。ていうか、計算はすぐ終わるので、計算してる時間より隙間の時間のほうが圧倒的に長い。



これを解消するには、ともかくCPU - GPU間のやりとりを減らしてあげる必要がある。

そこで、厳密に最もよくロゴが消えるフェード値を探すのはあきらめ、一度複数(今回は17通り)のフェード値について評価結果を得たら、その結果を使って最適なフェード値を推定するようにした。最初に行う複数のフェード値での計算をGPUで一度に行うようにすれば、CPU - GPU間は1往復で済むようになる。


クリックで拡大

具体的にどう推定するかというと、まず17通りのフェード値でロゴ除去→消え具合の評価を行う。そうするとこんな感じ。フレームによって、評価結果が変わってくるので、典型的な2パターンを載せてみた。

パターンA


パターンB


この図では、横軸にフェード値、縦軸がその時のロゴの消え具合で、値が低いほど「よく消えている」。

フェード値が低すぎるとロゴが残ってしまい、逆に強すぎるとロゴが反転してしまってまたロゴが見えてくるので、あるフェード値で最も値が低くなる、という感じで、結局最小値を探索する問題になる。図を人間が見れば、最小値のありそうな場所ははっきりわかるのだけど、これをどう計算するかという問題。

最小値/最大値の探索は、評価関数f(x)が明確であれば、ニュートン法やら勾配法やらを使えばよいのだけど、今回の場合は反復計算したくないのもあってそういうわけにもいかない。

さらに注意点として、
・図の両方のケース(パターンA/パターンB)に対応したい
・最小値はプロット点の間になることもある

というわけで、お手軽な方法として、プロット点から近似曲線(2次)を求めて、近似曲線から最小値を得ることとした。ここで、プロット上の最小値を挟んで、左側と右側に分けてそれぞれ近似曲線も求めることで、パターンA/パターンBの両方に対応することができた。


パターンA


パターンB


最小二乗法は3次までは比較的楽なので、近似曲線は簡単に求めることができる。(4次以上になると行列解くコード書くのが面倒)

この方法を採用することで、オリジナルと比べて精度はやや落ちるものの、大幅に高速化できた。

自動フェード改良後
encoded 2696 frames, 455.41 fps, 6194.73 kbps, 66.43 MB
encode time 0:00:05, CPU: 5.6%, GPU: 22.5%, VE: 85.1%, GPUClock: 1838MHz, VEClock: 1654MHz

Vpp Filter Performance
cspconv: 41.8 us
delogo: 816.3 us
cspconv: 41.6 us

エンコード速度は、自動フェードなしの場合とほぼ変わらない速度が出るようになったし、delogo自体の速度も大幅に改善した。とはいえ、まだなんか隙間があるんですよねー。


クリックで拡大



残りの隙間は何かというと、自動NRを使用した場合に、NR強度を0~4まで変えて、それぞれ自動フェード値を求める必要があるので、5回のループが見えてしまっている、ということになる。


クリックで拡大

ところが、このNR強度を変えて…という部分は実は同時に行うことも可能だ。CUDAのstreamという機能を使うと、この同時処理が実現できる。


クリックで拡大

隙間が少しづつ埋まってくると、ほかのところも気になってくる。まだ残る隙間についても、CPUで行っていた処理の一部をGPUで行うようにすることで、CPU-GPUのやり取りをなくした。

そうすると、こんな感じ。


クリックで拡大

まだ、隙間は残っているけど、とりあえず十分速くなったので、これでよしとした。



まあ、それでもちょっと気になったのが、なんか周りのkernelと比べて異常に時間がかかっているkernelがあること。コードを書いた感じでは、このkernelが他と比べて遅い、というのは考えにくかった。

いろいろ調べていって、最終的にkernelコードの逆アセンブリを見ると、定数配列を一度Local変数に展開してから使うようなコードになっていて、これが非常に遅いようだった。

定数配列をconstantメモリというものに載せてあげると、Local変数を使わなくなり、大幅に高速化した。(220us → 14us)


クリックで拡大

最終的な高速化結果
encoded 2696 frames, 458.89 fps, 6194.73 kbps, 66.43 MB
encode time 0:00:05, CPU: 5.4%, GPU: 13.9%, VE: 83.0%, GPUClock: 1840MHz, VEClock: 1646MHz

Vpp Filter Performance
cspconv: 41.8 us
delogo: 531.8 us
cspconv: 41.6 us

delogoの処理時間はさらに減り、初めから比べると20倍超速くなった。



まとめ



・CPU-GPUの転送を減らすことで大きく高速化した。CPU-GPUの転送は、転送量が少なくてもレイテンシが辛いので、遅くなる原因になりやすい。

・stream機能を使うと、GPU上で並列処理が可能になる。

・kernelの配列は、うまく最適化してくれないときがあるので、気を付ける。(コードの書き方が悪い?)

・NVIDIA Visual Profilerはかなり使いやすい。CPUでもこういうのがあるといいのだけど、難しいのかな…。
※VTuneはすごくいいけどお値段が高すぎてとても個人的に買う値段じゃない (そんなお金あったらCPU買ったりfgoに課金したり…)



速いCUDAを書くのは難しい…



ループ構造をCUDAで実装することそのものは決して難しくないように思う。

ただ、一般的に言って、今回みたいに、素直にCPUのコードをCUDAに焼き直しただけだとまるで速度が出ないということが結構多い気がする。で、GPUで遅いなら別にわざわざGPUでやる必要はないので、GPUで動かす以上、速いコードを書かないと…という強迫観念がある。

速度を出そうとすると、やれstreamを使えだの、やれsharedメモリを使えだの、まあなかなか面倒だったり難しかったりな話になるので、コードが複雑化して、非常に難しい。

まあ、GPUの特性を考えれば、もちろん仕方ないことだとは思うし、それでもCUDAはもっと面倒なOpenCLと比べれば全然マシだったりするけど、それでももうちょっとなんとかならないの、と思う。

え、そもそも数十usのkernelを連投するのが間違いだって?

そんなこと言われてもねえ。



スポンサーサイト



コメントの投稿

非公開コメント

No title

ですのでGPU使っても早くならないことが多く特定のアルゴリズム以外、割に合わないですね。
CPUで並列処理しやすくかつ、ループ中の計算量が非常に多い場面で考えます。

constexprとかthrustとか

ところで __device__ constexprと__constant__どっちが速いのだろう・・・。

あとthrustは今回使っていない感じですかね・

No title

>>2018-08-19(09:49)さま
たしかに、行列積など、GPUに向いたアルゴリズムなら比較的容易に高速化できますよね。

ご指摘のように、今回はロゴというのがもともとピクセル数が少ないためループ内の計算量が少なく、本来GPU向きの計算ではありません。なので、そういう悪条件の中でどう対処するか…という感じになってしまっています。

> yumetodo様
GPUで__constant__というと、ご存知のように定数読み出しに適した専用のキャッシュ付きの領域に置かれます。今回の用途についてはかなり適した使い方になっていると考えています。もちろん、同じ定数的なデータでもメモリアクセスパターンによっては不適切となりえるので、ケースバイケースではあります。

一方、__device__ constexprのほうは、自分の理解が正しければ、どのように扱われるかはCUDAコンパイラ次第で、コンパイラの判断とそのデータに対するメモリアクセスパターン次第となると思います。

今回については、__constant__のほうがよい可能性が高いですから、やや使うのが面倒ですが、__constant__を使う価値はあると思います。

今回、thrustは使ってないです。ソートなど、thrustのライブラリにあるアルゴリズムを使いたい場合は有用ですが、今回はそういうわけではないので…。

thrustを使うとメモリ管理が簡単になるという話もあるんですが、別にthrustを使わなくてもメモリ管理が難しいとは感じないのですよね…。
プロフィール

rigaya

Author:rigaya
アニメとか見たり、エンコードしたり。
連絡先: rigaya34589@live.jp
github twitter

最新記事
最新コメント
カテゴリ
月別アーカイブ
カウンター
検索フォーム
いろいろ
公開中のAviutlプラグインとかのダウンロード

○Aviutl 出力プラグイン
x264guiEx 3.xx
- x264を使用したH264出力
- x264guiExの導入紹介動画>
- x264guiExの導入
- x264guiExのエラーと対処方法>
- x264.exeはこちら>

x265guiEx
- x265を使用したH.265/HEVC出力
- x265guiExの導入>
- x265.exeはこちら>

QSVEnc + QSVEncC
- QuickSyncVideoによるHWエンコード
- QSVEnc 導入/使用方法>
- QSVEncCオプション一覧>

NVEnc + NVEncC
- NVIDIAのNVEncによるHWエンコード
- NVEnc 導入/使用方法>
- NVEncCオプション一覧>

VCEEnc + VCEEncC
- AMDのVCE/VCNによるHWエンコード
- VCEEnc 導入/使用方法>
- VCEEncCオプション一覧>

svtAV1guiEx
- SVT-AV1によるAV1出力
- svtAV1guiExの導入>
- SVT-AV1単体はこちら>

VVenCguiEx
- VVenCによるVVC出力
- VVenCguiExの導入>

ffmpegOut
- ffmpegを使用した出力
- ffmpegOutの導入>


○Aviutl フィルタプラグイン
自動フィールドシフト (ミラー)
- SSE2~AVX512による高速化版
- オリジナル: aji様

clfilters 
- OpenCLベースの複数のGPUフィルタ集
- 対応フィルタの一覧等はこちら

エッジレベル調整MT (ミラー)
- エッジレベル調整の並列化/高速化
- SSE2~AVX512対応
- オリジナル: まじぽか太郎様

バンディング低減MT (ミラー)
- SSE2~AVX512による高速化版
- オリジナル: まじぽか太郎様

PMD_MT
- SSE2~AVX512による高速化版
- オリジナル: スレ48≫989氏

透過性ロゴ (ミラー)
- SSE2~FMA3によるSIMD版
- オリジナル: MakKi氏

AviutlColor (ミラー)
- BT.2020nc向け色変換プラグイン
- BT.709/BT.601向けも同梱

○その他
Amatsukaze改造版
- AmatsukazeのAV1対応版

x264afs (ミラー)
- x264のafs対応版

aui_indexer (ミラー使い方>)
- lsmashinput.aui/m2v.auiの
 インデックス事前・一括生成

auc_export (ミラー使い方>)
- Aviutl Controlの
 エクスポートプラグイン版
 エクスポートをコマンドから

aup_reseter (ミラー)
- aupプロジェクトファイルの
 終了フラグを一括リセット

CheckBitrate (ミラー, 使い方, ソース)
- ビットレート分布の分析(HEVC対応)

チャプター変換 (使い方>)
- nero/appleチャプター形式変換

エッジレベル調整 (avisynth)
- Avisynth用エッジレベル調整

メモリ・キャッシュ速度測定
- スレッド数を変えて測定
- これまでの測定結果はこちら

○ビルドしたものとか
L-SMASH (ミラー)
x264 (ミラー)
x265 (ミラー)
SVT-AV1 (ミラー)

○その他
サンプル動画
その他

○読みもの (ミラー)
Aviutl/x264guiExの色変換
動画関連ダウンロードリンク集
簡易インストーラの概要

○更新停止・公開終了
改造版x264gui
x264guiEx 0.xx
RSSリンクの表示
リンク
QRコード
QR