🏞️

GpuMatとPtrStep

に公開
C++
OpenCV
CUDA
tech

PtrStepはcv::cuda::GpuMatの持っているメモリ領域を参照するための軽量の型で、特にカーネルなどデバイス側の処理において便利です。GpuMatそのものはデバイスコードで使用出来ないようなので、必然的にこのクラスを使用するかポインタで頑張るかのどちらかを選択する事になるでしょう・・・

PtrStepは便利なのですがリファレンスが無く何ができるのかよくわかりません。ソースコードを見れば自明なのですが、一々見に行くのもめんどいので記録しておきます。

※ 以下の記述はOpenCV 4.5時点のソースを参考にしています。

cv::cuda::PtrStep

使うときは<opencv2/core/cuda.hpp>をインクルードすれば使えます。つまり、cv::cuda::GpuMatを使ってるときは自動で使えるようになってるはずです。実体は<opencv2/core/cuda_types.hpp>にあります。

PtrStepはクラステンプレートであり、参照するメモリ領域の要素の型Tを受け取ります。

#include <opencv2/core/cuda.hpp>

using cv::cuda::PtrStep;

void f(int* p, std::size_t step) {
  PtrStep<int> ps{p, step};
}

このように構築した場合、pの指す領域を横幅step [byte]の2次元でint型の領域として参照します。

PtrStep<T>は以下の操作が可能です。

操作 説明 constオーバーロード
T& operator()(int y, int x) 参照する領域の(x, y)位置の要素を取得する 戻り値型がconst T&になる
T* ptr(int y = 0) y行目の先頭ポインタを取得する 戻り値型がconst T*になる
size_t elemSize() sizeof(T)を取得する あり、変化なし
operator T*() 参照する領域ポインタへの暗黙変換 戻り値型がconst T*になる
PtrStep() デフォルトコンストラクタ -
PtrStep(T* data_, size_t step_) data_の領域を横幅step_の二次元領域として参照する -

まあ専ら使用するのはoperator()と2つ目のコンストラクタでしょう。

#include <opencv2/core/cuda.hpp>

using cv::cuda::PtrStep;

__global__ void kernel(PtrStep<float> ps) {
  const int x = ...;
  const int y = ...;

  // 読み出し
  int v = ps(y, x);

  // 書き込み
  ps(y, x) = 10;
  
  // 3行目の先頭ポインタを取得
  int* pl = ps.ptr(3);
  
  // デフォルト構築可能
  PtrStep<int> psi{};
  
  // これは出来ない（Tが異なる）
  psi = ps;
  
  // これはできる
  PtrStep<float> ps2 = ps;
}

これらのメンバ関数は全て、ホスト側でもデバイス側でも実行可能です。

また、メンバ変数が公開されており直接アクセス可能です。とはいえこれは知らなかった事にしてアクセスしないことをお勧めします。書き換えると危険なので・・・

変数名 説明
T* data 参照する領域のポインタ
size_t step 参照する二次元領域の横幅 [byte]

テンプレートパラメータTについての制約は何もないので、どんな型を指定することもできます。例えばfloat2とかuint2とかを指定することもできます。

PtrStepの構造

PtrStepは単純には次のような構造体です（これは例示のための単純化であり正確ではありません）。

template<typename T>
struct PtrStep {
  T* data;
  size_t step;

  // 略
};

このことから分かるようにsizeof(PtrStep<T>)はTによらずに決まり、ほとんどの環境では16バイトになるはずです。

このように、PtrStepはとても単純なクラスであり、コピーコンストラクタは宣言されていないので、常にTrivially Copyableです。一方、デフォルトコンストラクタがdefaultではないのでTrivialな型ではなく、実際はメンバがあるクラスを継承しているためStandard Layoutでもありません。

このことは型Tとは無関係に決まります。

template<typename T>
void f() {
  // ok
  static_assert(std::is_trivially_copyable<PtrStep<T>>::value, "");

  // ng
  static_assert(std::is_trivial<PtrStep<T>>::value, "");

  // ng
  static_assert(std::is_standard_layout<PtrStep<T>>::value, "");
}

Trivially Copyableでありサイズがポインタ2つ分程度にしかならないので、PtrStepは参照渡しではなく値渡しでやり取りする事が推奨されます。

GpuMat -> PtrStep

PtrStepは主にGpuMatの領域を参照するのに用いられます。そのため、GpuMatの側ではPtrStepへの暗黙変換を用意しています。

#include <opencv2/core/cuda.hpp>

using cv::cuda::PtrStep;

__global__ void kernel(PtrStep<unsigned char> ps);

void kernel_caller(const cv::cuda::GpuMat &img) {
  // 特に何もせずそのまま渡せる
  kernel<<<1, 1>>>(img);
}

（値渡しをする場合は）これがあるので、PtrStepをどう構築するか（できるか）と言うのは忘れても良いでしょう。GpuMatによらないメモリ領域を参照する時に思い出せば良いです。

注意点としては、GpuMatが実際に確保している型と異なる型によるPtrStepに変換するときでも特に警告やアサートはされない点です。

#include <opencv2/core/cuda.hpp>

using cv::cuda::PtrStep;

// 要素型をfloatとして領域を参照する
__global__ void kernel(PtrStep<float> ps);

void kernel_caller(int width, int height) {
  // unsigned charの領域を確保
  cv::cuda::GpuMat img{cv::Size(width, height), CV_8UC1};

  // 警告も実行時アサートもない
  kernel<<<1, 1>>>(img);
}

これはある意味で便利ではありますが、そこそこ危険です。特にこのように型のサイズが異なる場合に、バッファオーバーランの恐れがあります。

なお、変換先と元の要素型がStrict Aliasing Ruleに沿わない型同士の場合未定義動作になります。nvccがそれをどうするのかは知りませんが・・・

入力と出力を明示する

前項の危険性を有効に利用すると、PtrStepによる参照を入力用と出力用に型レベルで区別可能とする事ができます。

#include <opencv2/core/cuda.hpp>

using cv::cuda::PtrStep;

// 要素型をfloatとして領域を参照する
__global__ void kernel(PtrStep<const uchar> in, PtrStep<uchar> out) {
  const int x = ...;
  const int y = ...;
  
  // コンパイルエラー、書き込み不可
  in(y, x) = 0;

  // ok
  const auto pix = in(y, x);
  
  // ok
  out(y, x) = pix;
}

void kernel_caller(const cv::cuda::GpuMat& in_img, cv::cuda::GpuMat& out_img) {
  kernel<<<1, 1>>>(in_img, out_img);
}

このように、入力用のPtrStep<T>をPtrStep<const T>のように宣言しておく事で、その参照する領域に書き込みをしないことを明示し、書き込まれた場合にコンパイルエラーとなるようにする事ができます。

この場合は何ら未定義動作を踏まないはずです。

とはいえ、メンバ関数がきちんとconst修飾されているため、PtrStep自身をconstにしても同様になります。

// とはいえこっちでも同じになるはず・・・
__global__ void kernel(const PtrStep<uchar> in, PtrStep<uchar> out) {
  const int x = ...;
  const int y = ...;
  
  // コンパイルエラー、書き込み不可
  in(y, x) = 0;

  // ok
  const auto pix = in(y, x);
  
  // ok
  out(y, x) = pix;
}

PtrStepは巷でviewとかspanとか呼ばれるもののため、それ自身のconst性と参照先のconst性が同期する事に違和感を覚える人もいるかもしれません（そんな人いるんですか？）。そんな場合はテンプレートパラメータのTをconstにすると良いでしょうか・・・？

PtrStepファミリ

PtrStepの定義されているcuda_types.hppには、親族とも言うべきクラスがいくつか定義されています。それらは継承関係を持ち、実はPtrStepもそこに組み込まれています。

// ポインタのラッパ
template<typename T>
struct DevPtr {
  T* data;
};

// ポインタとステップ（横幅）のラッパ
template<typename T>
struct PtrStep : public DevPtr<T> {
  size_t step;
};

// ポインタとステップと領域サイズ（width × height）のラッパ
template<typename T>
struct PtrStepSz : public PtrStep<T> {
  int cols;
  int rows;
};

// ポインタと領域サイズのラッパ
template<typename T>
struct PtrSz : public DevPtr<T> {
  size_t size;

  template <typename U>
  explicit PtrStepSz(const PtrStepSz<U>& d);
};

DevPtrは他の3つの定義を簡略化するための共通部分なので直接使用することはないでしょう。

PtrStepSzとPtrSzは両方とも、参照領域のポインタとその領域のサイズを持ちますが、PtrStepを経由しない分PtrSzの方がサイズが小さくなります。PtrStepSzはサイズが大きくなりますが領域の縦横サイズを持っており、PtrStepでできることは全てできます。

ほとんどの環境では、sizeof(PtrStepSz<T>) == 24、sizeof(PtrSz<T>) == 16となるはずです。PtrStepSzは値渡しと参照渡しの損益分岐点の微妙な所にありますが、多くの場合は値渡しで問題ないでしょう。

PtrStepで可能な操作のうち次の2つはDevPtrで定義されているため、この2つはPtrSzでも使用可能です。

操作 説明 constオーバーロード
size_t elemSize() sizeof(T)を取得する あり、変化なし
operator T*() 参照する領域ポインタへの暗黙変換 戻り値型がconst T*になる

PtrStepSz<T>はcols, rowsというint型の二つのメンバをpublicに持つ事に加えて、T* -> U*に変換可能なPtrStepSz<U>からの変換コピーコンストラクタを備えています。ただ、このコンストラクタは特に制約されていないのに加えて、初期化時にCスタイルキャストを使用するために多分なんでも変換できます・・・

GpuMatはPtrStepSzに対する暗黙変換も備えているため、カーネルに渡すときなどはそのまま渡す事ができます。

PtrStepSzはカーネル中で領域アクセスに加えて画像の縦横サイズを参照したい時にPtrStepの代わりに使用すると良いでしょうか。PtrSzはポインタとその領域サイズのペアを単に持ち運ぶ時に活用できそうです。

エイリアス

PtrStepとPtrStepSzには、よく使うであろう型についてエイリアスが用意されています。

typedef PtrStepSz<unsigned char> PtrStepSzb;
typedef PtrStepSz<unsigned short> PtrStepSzus;
typedef PtrStepSz<float> PtrStepSzf;
typedef PtrStepSz<int> PtrStepSzi;

typedef PtrStep<unsigned char> PtrStepb;
typedef PtrStep<unsigned short> PtrStepus;
typedef PtrStep<float> PtrStepf;
typedef PtrStep<int> PtrStepi;

接尾辞に慣れないと便利かは微妙ですが、引数宣言が長くなる時などに使ってみると良いかもしれません。

参考文献

Discussion