C++におけるCUDAのメモリ管理を楽にする

CUDA関数は、引数が多くて煩雑で、使うのが大変だ（例えばcudaMemcpy2D）

そこで、以下のコードを作ったら、メモリ管理が楽になった
https://github.com/rueiwoqpqpwoeiru/image-process-CUDA/blob/main/Mem.h

（工夫した点）CUDA関数をブラックボックス化して、使わなくて済むようにした
111, 112行目：デバイスメモリと、ホストメモリをペアで持つ (※メモリにゆとりが必要）
118, 120行目：ペアのメモリは同時に確保され、ホスト側はページロックされて転送が高速化(19行目)
　　　　　　（※メモリ確保は時間が掛かるので、1回確保した後は同サイズのデータのみ扱う）

110行目以降

template<typename T> struct D_mem {  // 110行目
	D_data<T> d_data;                // 111行目
	H_data<T> h_data;                // 112行目
	// open memory
	void open_2d(unsigned int w, unsigned int h, unsigned int ch, const cudaTextureFilterMode& interpolation = cudaFilterModePoint) {
		if (nullptr == d_data.ptr) {
		    d_data.w = w; d_data.h = h;	d_data.ch = ch;
		    Stopwatch sw;
		    CUDA_CHECK(cudaMallocPitch(&(d_data.ptr), &(d_data.pitch), sizeof(T) * d_data.w, d_data.h * d_data.ch));  // 118行目（デバイス側のメモリ確保）
		    sw.print_time(__FUNCTION__);
		    h_data.open(d_data.w, d_data.h, d_data.ch);  // 120行目（ホスト側のメモリ確保）
		    this->use_tex(interpolation);
		}
	}
    ...

189, 199行目：データ転送は、h2d()やd2h()を呼ぶだけなので、引数のミスが起こらない

189行目以降

	void h2d() const {  // 189行目（D_memクラスのメンバ関数）
		Stopwatch sw;
		if (1 < d_data.h) {
			CUDA_CHECK(cudaMemcpy2D(d_data.ptr, d_data.pitch, h_data.ptr, sizeof(T) * d_data.w, sizeof(T) * d_data.w, d_data.h * d_data.ch, cudaMemcpyHostToDevice));
		}
		else {
			CUDA_CHECK(cudaMemcpy(d_data.ptr, h_data.ptr, sizeof(T) * d_data.w * d_data.ch, cudaMemcpyHostToDevice));
		}
		sw.print_time(__FUNCTION__);
	}

210, 236行目：画像やcsvを指定するだけで、必要に応じてメモリを確保し、デバイスに転送

210行目以降

	void imread(const std::string& file_name) {  // 210行目（D_memクラスのメンバ関数）
		// for get size
		Img_io img(file_name);
		// alloc memory if empty
		if (nullptr == d_data.ptr) {
			if (img.tag.bit == sizeof(T) * 8) {
				this->open_2d(img.tag.w, img.tag.h, img.tag.ch);  // メモリ確保（デバイスとホスト）
			}
			else {
				assert(false);
			}
		}
		// read img if size is OK
		if (img.tag.w == d_data.w && img.tag.h == d_data.h && img.tag.ch == d_data.ch && img.tag.bit == sizeof(T) * 8) {
			h_data.imread(file_name);
			this->h2d();   // ホストからデバイスに転送
		}
		else {
			assert(false);
		}
	}

231, 252行目：画像やcsvを指定するだけで、ホストにデータを転送してファイルに出力

231行目以降

	void imwrite(const std::string& file_name, const int metric = PHOTOMETRIC_MINISBLACK) const {  // 231行目（D_memクラスのメンバ関数）
		this->d2h();  // デバイスからホストに転送
		h_data.imwrite(file_name, metric);  // ホスト側のimwriteを呼ぶ
	}

　　　　　　　　　画像はtifで、αチャンネルを使うことで、数十チャンネルの画像も入出力可能
　　　　　　　　　（参考）libtiffをソースからビルドする方法は以下に書いた
　　　　　　　　　　　　　https://qiita.com/rueiwoqpqpwoeiru/items/e0eda6752ac259484288
　　　　　　　　　libtiff以外を使用するには、ホスト側のimread（41行目）とimwrite（55行目）を修正

（補足）
　11, 97行目：多チャンネルの画像は、平面内で縦方向に連結したつもりで保持（※メモリは1次元）
　50, 65行目：画素の配置順の変更は、デバイス側で行うべき
104, 106行目：読み取りだけの場合、乗算が1つ少ない分、106行目の方が高速（ここは意外に重要）

上記コードを用いた画像処理について以下に記載
https://qiita.com/rueiwoqpqpwoeiru/items/69a931e4042a4f9d4434