概要

第二のドワンゴ Advent Calendar 2019 21日目の記事です。
NVIDIA VIDEO CODEC SDK の NvDecoder を使って、デコードされた VRAM 上の映像フレームにアクセスし、ホストに書き戻さずデバイス上で計算してみました。

はじめに

みなさん、GPU を活用していますか。していますよね。

私は専らライブストリーミングサービスの開発を生業としており、個人的な興味において NVENC/NVDEC で生成したメディアストリームをアプリケーションから操作したいという欲求があります。

GPU を使うとき、「VRAM 上のバッファをホストに転送せず、デバイス上で操作を完結させたい」というケースは多く見かけることでしょう。特に HPC 用途ともなれば GPU は高い諸元を誇ります。その計算力を常時発揮させるためにも、不必要な PCIe 通信を防がねばなりません。

NVIDIA VIDEO CODEC SDK は、NVENC/NVDEC のハードウェアエンコード・デコードを容易に行える機能が提供されています。
このうち、NVDECODE APIs はデコード後のフレームバッファを DirectX や OpenGL コンテキストで操作することを可能にします。過去に CUDA Video Decoder API として提供されていた API の後継であり、特に SDK にある NvDecoder クラスは入出力をシンプルに扱えるよう抽象化されています。用途に応じて API を直接使用することができますが、今回のようにカジュアルに検証したいケースでは NvDecoder クラスは大変ありがたい実装です。
これを使用して、デコードされた VRAM 上の映像フレームにアクセスしてみましょう。

なお、NvDecoder クラスのドキュメントはないので実装を読んで理解しましょう。目的の達成のためには些細なことですよね。

実践

実践していきましょう。今回は GPU を搭載した CentOS 7 環境で検証してみます。

環境

使用可能な機能はハードウェア・アーキテクチャに依存するため、ご利用の際はお使いのデバイスと SDK の CAPABILITIES を確認しましょう。

条件	内容
プラットフォーム	CentOS Linux release 7.4.1708 (Core)
GPU Driver	NVIDIA Driver Version: 384.183
CUDA Toolkit	9.0
NVIDIA VIDEO CODEC SDK	8.2.16

本文書で使用する映像は Mazwai より以下のデータを使用させていただきました。ありがとうございます。

raccoon come and go (by benjamin wu)

VRAM 上の Decode した映像フレームを得る

SDK の NvDecoder サンプルに AppDecMem があります。こちらは入力した映像を NVDEC を通して DEMUX/DECODE し、映像フレームごとにホストメモリに書き出し、ついでにファイルに保存するものです。こちらを動作させてみます。

$ ./AppDecMem -i ~/media/benjamin_wu--raccoon_come_and_go.mp4
GPU in use: ***
[INFO ][18:07:15] Media format: QuickTime / MOV (mov,mp4,m4a,3gp,3g2,mj2)
Session Initialization Time: 28 ms
[INFO ][18:07:15] Video Input Information
    Codec        : AVC/H.264
    Frame rate   : 30000/1001 = 29.97 fps
    Sequence     : Progressive
    Coded size   : [1280, 720]
    Display area : [0, 0, 1280, 720]
    Chroma       : YUV 420
    Bit depth    : 8
Video Decoding Params:
    Num Surfaces : 20
    Crop         : [0, 0, 0, 0]
    Resize       : 1280x720
    Deinterlace  : Weave

Invalid return value 0 for stream protocol
Invalid return value 0 for stream protocol
Total frame decoded: 2217
Saved in file out.nv12Session Deinitialization Time: 7 ms

$ ls -l out.nv12
-rw-rw-r-- 1 *** *** 3064780800 Dec 20 18:07 out.nv12 # 出力したデコード済みフレーム

無事デコードは動作しましたね。

このサンプルではデコードしたフレームを常にホストバッファとファイルへ書き出しているため、VRAM 上のバッファを参照することはできません。さて、どのように変更しましょう。

ポイントは NvDecoder コンストラクタの bUseDeviceFrame です。これを true にすることで NvDecode::Decode() の pppFrame 引数は GPU デバイスメモリを返すことができます。
実装としては、フレームサイズのデバイスメモリを新たに確保し、デコード後に device to device でフレームの内容をコピーしているようです。

NvDecoder コンストラクタのインタフェース:

NvDecoder::NvDecoder(
  CUcontext cuContext, int nWidth, int nHeight,
  bool bUseDeviceFrame, // true: デバイスメモリ上にフレームサイズ分のメモリを確保し、コピーする
  cudaVideoCodec eCodec, std::mutex *pMutex = NULL,
  bool bLowLatency = false, bool bDeviceFramePitched = false,
  const Rect *pCropRect = NULL, const Dim *pResizeDim = NULL,
  int maxWidth = 0, int maxHeight = 0)

NvDecoder::Decode() のインタフェース:

bool NvDecoder::Decode(
  const uint8_t *pData, int nSize,
  uint8_t ***pppFrame, // フレームの出力先バッファ
  int *pnFrameReturned, uint32_t flags = 0, int64_t **ppTimestamp = NULL,
  int64_t timestamp = 0, CUstream stream = 0)

bUseDeviceFrame = true のときの pppFrame は cudaMalloc() で確保したデバイスメモリになるため、ホストからは参照できません。NvDecoder::Decode() 呼び出し後は VRAM 上に確保したバッファにコピーされた状態となります。

VRAM 上の映像フレームを編集する

この状態ならば、 CUDA Kernel 関数は VRAM 上の pppFrame を入力とすることができます。

フレームを適当な CUDA テクスチャオブジェクトと加算合成してみましょう。大変適当ですが、以下のような Kenrel 関数とその呼び出し部を実装しました。

デバイスコード:

// デバイス側の CUDA テクスチャオブジェクトの加算 Kernel 関数
__global__ void add_texture(unsigned int *dest, cudaTextureObject_t texture,
                            int sx, int sy, int sw, int ssize) {
    auto x = blockDim.x * blockIdx.x + threadIdx.x;
    auto y = blockDim.y * blockIdx.y + threadIdx.y;
    // TODO: stride
    auto dest_index = (sy + y) * sw + (sx + x);
    if ((dest_index < 0) || (dest_index > ssize)) return;
    dest[dest_index] += tex2D<unsigned int>(texture, x, y);
}

ホストコード:

// ホスト側の Kenrel 関数の呼び出し
void add(const texture &texture, std::size_t x, std::size_t y) {
    auto sw = static_cast<int>(texture.get_width());
    auto sh = static_cast<int>(texture.get_height());
    auto dw = static_cast<int>(surface_.get_width());
    auto dh = static_cast<int>(surface_.get_height());
    add_texture << <make_grid(dim_block_, sw, sh), dim_block_ >> >(
        pppFrame.data(), texture.data(), x, y, sw, sw * sh);
    auto result = cudaGetLastError();
    if (result != cudaSuccess) {
        THROW_ERROR(cudaGetErrorString(result));
    }
}

千超コアの演算処理性能を発揮できるとは言い難い処理の上にその実装も適当極まりないですが、要はデバイス上のフレームにアクセスし編集できるということが重要なのです。
ここから更に NvEncoderを使えばデバイス上で映像に戻すことが可能です。

ひとまずフレームを cudaMemcpy() でデバイスからホストへ書き出し、確認してみましょう。

raccoon come and go (by benjamin wu) との加算合成：

NVDEC からホストを一切介さず、映像フレームに対して加算合成を行うことができました。最もシンプルな（そして実装が楽な） CUDA API で検証しましたが、DirectX/OpenGL で任意のポストプロセスを行うことも同様に可能となるでしょう。

メディアストリームを扱う際は、YUV - RGBA 変換や Straight Alpha - Premultiplied Alpha 変換などのピクセル操作・フィルタ処理を合わせて行うケースが多いと思います。GPU が存在し、 NVENC/NVDEC を利用する場合は、せっかくだからデバイス上でできることをやってしまうのがよいでしょう。実用度が比較的高く、計算資源を有効に扱えるケースではないでしょうか。

おわりに

VRAM 上のメディアストリームをアプリケーションから操作する目的を達成できました。

ところで、GPU の話につきものですが、VRAM 上の映像をアクセスすることが有効かどうかは時と場合によります。今回の実践では全く触れていませんが、機能・非機能要求を踏まえ、以下のようなポイントを考えねばなりません。

特定のベンダ・デバイス・ドライバ・SDK バージョンに強く依存する問題
非同期ストリーミング処理を行うためのパイプライン制御、GPU デバイス上のバッファリング機構の設計
NVENC/NVDEC と GPU デバイス上の操作を同居させるためのリソースのやりくり
あるプロセスでデバイスの CUDA コンテキストを得ると他のプロセスからデバイスを操作することができない
保守開発コスト、というかデバッグとチューニングに費やす時間が爆発しがち

皆様は是非、十分な知識と覚悟を持って GPU の活用を検討してみましょう。
私は知識と覚悟が足りず疲れました。良いお年を。