はじめに

動機

数百枚の画像の切り抜きに、cv::Rect と cv::Mat を使用してるが、もっと早く処理したい。GPUを使った画像切り抜きは速いのか？

ググったらそのまま使えそうなコードがあるので Let's Try!

注意

Visual Studio、C++、CUDA、何も分からん状態でChatGPTの協力を受けつつ進めました。アドバイスを頂けるとありがたいです。

作業リポジトリはこちらで公開しています。

環境

CPU: AMD Ryzen 9 3900X 12-Core Processor (24 CPUs), ~3.8GHz
GPU: NVIDIA GeForce RTX 3060

Windows の仕様

エディション	Windows 10 Pro
バージョン	22H2
インストール日	2021/02/02
OS ビルド	19045.4291
エクスペリエンス	Windows Feature Experience Pack 1000.19056.1000.0

CUDA Toolkit

PS C:\Users\owner> nvcc -V
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2024 NVIDIA Corporation
Built on Thu_Mar_28_02:30:10_Pacific_Daylight_Time_2024
Cuda compilation tools, release 12.4, V12.4.131
Build cuda_12.4.r12.4/compiler.34097967_0

開発環境: Visual Studio 2022 version 17.9.6

CUDA Toolkit の導入

CUDA Windows でググるとNVIDIAによる公式ドキュメントが見つかります。ドキュメントに従って CUDA Toolkit をインストールします。

インストール

サンプルプロジェクトの実行

4. Compiling CUDA Programsに従って、リポジトリをCloneし対象のソリューション(.sln)ファイルを開きます。

P:\tmp>git clone https://github.com/NVIDIA/cuda-samples.git
Cloning into 'cuda-samples'...
remote: Enumerating objects: 18291, done.
remote: Counting objects: 100% (3706/3706), done.
remote: Compressing objects: 100% (587/587), done.
remote: Total 18291 (delta 3362), reused 3157 (delta 3119), pack-reused 14585
Receiving objects: 100% (18291/18291), 133.19 MiB | 19.29 MiB/s, done.
Resolving deltas: 100% (16022/16022), done.
Updating files: 100% (4012/4012), done.

P:\tmp>cuda-samples\Samples\1_Utilities\bandwidthTest\bandwidthTest_vs2022.sln

ビルドと実行に成功したので、CUDAが使用できているはずです。

プロジェクトの作成と設定

新しいプロジェクトの作成

Visual Studio 2022 を開き、新しいプロジェクトを作成します。

新しいプロジェクトのテンプレートとして CUDA 12.4 Runtime を選択します。

プロジェクト名は CudaCropTest とします。

サンプルプログラムをデバッグなしで実行すると、正常に動作します。

インクルードディレクトリの設定

Repositoryを公開するときのために、外部ライブラリを専用のフォルダに隔離して作業を行います。

プロジェクトフォルダ（C:\Users\owner\source\repos\CudaCropTest\CudaCropTest）に external フォルダを作成し、これをインクルードディレクトリに追加します。
プロジェクト(P) > CudaCropTest のプロパティの(P) をクリックします。（プロパティの(P) は誤字？）

構成プロパティ > VC++ ディレクトリ > 全般 > インクルードディレクトリの値部分 > <編集...> を選択します。

$(ProjectDir)external (C:\Users\owner\source\repos\CudaCropTest\CudaCropTest\external) を追加します。

dusty-nv/jetson-utils のビルド

IssueやCMakeLists.txtを見る限りWindowsはサポートされていません。頑張りましょう。

`cuda\cudaCrop.cu`

kernel.cu に #include "cudaCrop.h" を追加して生じるエラーを解決していきます。

E1696 ソースファイルを開けません "cudaCropa.h" CudaCropTest

重大度レベル	コード	説明	プロジェクト	ファイル	行	抑制状態	詳細
エラー (アクティブ)	E1696	ソース ファイルを開けません "cudaCropa.h"	CudaCropTest	C:\Users\owner\source\repos\CudaCropTest\CudaCropTest\kernel.cu	7

external フォルダに cudaCrop.h を含む必要なファイル群を追加します。

external フォルダに追加したファイル

commandLine.cpp
commandLine.h
csvReader.h
csvReader.hpp
cudaCrop.cu
cudaCrop.h
cudaUtility.h
imageFormat.h
imageFormat.inl
logging.cpp
logging.h

ソースファイルをProjectへ追加し、コンパイルされるようにします。
ソリューションエクスプローラー >すべてのファイルを表示を選択して表示を切り替え、externalフォルダを右クリック > プロジェクトに含める(J) を選択するとやりやすいかもしれません。cuファイルは個別に追加します。

E1696 ソースファイルを開けません "strings.h" CudaCropTest

重大度レベル	コード	説明	プロジェクト	ファイル	行	抑制状態	詳細
エラー (アクティブ)	E1696	ソース ファイルを開けません "strings.h"	CudaCropTest	C:\Users\owner\source\repos\CudaCropTest\CudaCropTest\external\imageFormat.inl	26

<strings.h> はUNIX系のシステムで使用されるヘッダーファイルであり、Windowsにはありません。strcasecmpとstrncasecmpを定義したexternal\strings.hを作成し、インクルードされるようにコードを変更します。

external\strings.h

#pragma once
#include <string.h>
#define strcasecmp _stricmp
#define strncasecmp _strnicmp

-#include <strings.h>
+#include "strings.h"

C2010 '.': マクロのパラメーターリスト内で予期されていません

重大度レベル	コード	説明	プロジェクト	ファイル	行	抑制状態	詳細
エラー	C2010	'.': マクロのパラメーター リスト内で予期されていません	CudaCropTest	C:\Users\owner\source\repos\CudaCropTest\CudaCropTest\external\logging.h	144

external\logging.h

#define GenericLogMessage(level, format, args...) if( level <= Log::GetLevel() ) fprintf(Log::GetFile(), format, ## args)

可変個引数マクロ ## args でエラーが生じています。コンパイルできるように書き換えます。

CudaCropTest\external\logging.h

-#define GenericLogMessage(level, format, args...) if( level <= Log::GetLevel() ) fprintf(Log::GetFile(), format, ## args)
+#define GenericLogMessage(level, format, ...) if( level <= Log::GetLevel() ) fprintf(Log::GetFile(), format,  __VA_ARGS__)

-#define LogError(format, args...)              GenericLogMessage(Log::ERROR, LOG_COLOR_RED LOG_LEVEL_PREFIX_ERROR format LOG_COLOR_RESET, ## args)
+#define LogError(format, ...)          GenericLogMessage(Log::ERROR, LOG_COLOR_RED LOG_LEVEL_PREFIX_ERROR format LOG_COLOR_RESET,  __VA_ARGS__)

-#define LogWarning(format, args...)    GenericLogMessage(Log::WARNING, LOG_COLOR_YELLOW LOG_LEVEL_PREFIX_WARNING format LOG_COLOR_RESET, ## args)
+#define LogWarning(format, ...)        GenericLogMessage(Log::WARNING, LOG_COLOR_YELLOW LOG_LEVEL_PREFIX_WARNING format LOG_COLOR_RESET,  __VA_ARGS__)

-#define LogSuccess(format, args...)    GenericLogMessage(Log::SUCCESS, LOG_COLOR_GREEN LOG_LEVEL_PREFIX_SUCCESS format LOG_COLOR_RESET, ## args)
+#define LogSuccess(format, ...)        GenericLogMessage(Log::SUCCESS, LOG_COLOR_GREEN LOG_LEVEL_PREFIX_SUCCESS format LOG_COLOR_RESET,  __VA_ARGS__)

-#define LogInfo(format, args...)               GenericLogMessage(Log::INFO, LOG_LEVEL_PREFIX_INFO format, ## args)
+#define LogInfo(format, ...)           GenericLogMessage(Log::INFO, LOG_LEVEL_PREFIX_INFO format,  __VA_ARGS__)

-#define LogVerbose(format, args...)    GenericLogMessage(Log::VERBOSE, LOG_LEVEL_PREFIX_VERBOSE format, ## args)
+#define LogVerbose(format, ...)        GenericLogMessage(Log::VERBOSE, LOG_LEVEL_PREFIX_VERBOSE format,  __VA_ARGS__)

-#define LogDebug(format, args...)              GenericLogMessage(Log::DEBUG, LOG_LEVEL_PREFIX_DEBUG format, ## args)
+#define LogDebug(format, ...)          GenericLogMessage(Log::DEBUG, LOG_LEVEL_PREFIX_DEBUG format,  __VA_ARGS__)

E0029 式が必要です CudaCropTest

重大度レベル	コード	説明	プロジェクト	ファイル	行	抑制状態	詳細
エラー (アクティブ)	E0029	式が必要です	CudaCropTest	C:\Users\owner\source\repos\CudaCropTest\CudaCropTest\kernel.cu	93

kernel.cu

    // Launch a kernel on the GPU with one thread for each element.
    addKernel<<<1, size>>>(dev_c, dev_a, dev_b);

<<< の部分がエラーと表示されています。<<<...>>>はCUDAのSyntaxです。
（実はサンプルプログラムの時点で表示されていました）

このエラーはVisual Studioに表示されるだけで、ビルドには影響しないため無視します。

`image\imageIO.cpp`

cuda\cudaCrop.h で定義されている cudaCrop 関数には、uint8_t*,float*,uchar3*,uchar4*,float3*,float4*型の画像を渡す必要があります。リポジトリを散策すると、画像の読み書きを扱う image\imageIO.h を見つけました。

kernel.cu に #include "imageIO.h" を追加し、生じるエラーを解決していきます。

C1083 include ファイルを開けません。'imageIO.h':No such file or directory CudaCropTest

重大度レベル	コード	説明	プロジェクト	ファイル	行	抑制状態	詳細
エラー	C1083	include ファイルを開けません。'imageIO.h':No such file or directory	CudaCropTest	C:\Users\owner\source\repos\CudaCropTest\CudaCropTest\kernel.cu	7

external フォルダに imageIO.h を含む必要なファイル群を追加します。

external フォルダに追加したファイル

alphanum.h
cudaBayer.cpp
cudaBayer.h
cudaColorspace.cpp
cudaColorspace.h
cudaGrayscale.cu
cudaGrayscale.h
cudaMappedMemory.h
cudaRGB.cu
cudaRGB.h
cudaYUV-NV12.cu
cudaYUV-YUYV.cu
cudaYUV-YV12.cu
cudaYUV.h
filesystem.cpp
filesystem.h
imageIO.cpp
imageIO.h
stb\stb_image_resize.h
stb\stb_image_write.h
stb\stb_image.h

C1083 include ファイルを開けません。'glob.h':No such file or directory CudaCropTest

重大度レベル	コード	説明	プロジェクト	ファイル	行	抑制状態	詳細
エラー	C1083	include ファイルを開けません。'glob.h':No such file or directory	CudaCropTest	C:\Users\owner\source\repos\CudaCropTest\CudaCropTest\external\filesystem.cpp	33

<strings.h> と同様に、Windowsに glob.h は存在しません。画像の入出力をしているコードと関連を調べていくと、glob.h をインクルードする必要はないと分かります。

filesystem.h と filesystem.cpp の不要なインクルード文（glob.h を含む）と関数をコメントアウトし、エラーを解消します。
ただし、fileExists 関数は画像の読み書きに使用されるので対応します。

external\filesystem.cpp

// fileExists
bool fileExists( const std::string& path, uint32_t mask )
{
	return fileIsType(path, mask);
}

external\filesystem.cpp

+#include <filesystem>

// fileExists
bool fileExists( const std::string& path, uint32_t mask )
{
-   return fileIsType(path, mask);
+   return std::filesystem::exists(path);
}

filesystemを使用するため、構成プロパティ > C/C++ > 言語 > C言語標準を 既定 (ISO C++14 標準) から ISO C++17 標準 (/std:c++17) に変更します。

LNK2019 未解決の外部シンボル nppiCFAToRGB_8u_C1C3R が関数 "enum cudaError cdecl cudaBayerToRGB(unsigned char ,struct uchar3 ,unsigned int64,unsigned __int64,enum imageFormat)" (?cudaBayerToRGB@@YA?AW4cudaError@@PEAEPEAUuchar3@@_K2W4imageFormat@@@Z) で参照されました CudaCropTest

重大度レベル	コード	説明	プロジェクト	ファイル	行	抑制状態	詳細
エラー	LNK2019	未解決の外部シンボル nppiCFAToRGB_8u_C1C3R が関数 "enum cudaError __cdecl cudaBayerToRGB(unsigned char *,struct uchar3 *,unsigned __int64,unsigned __int64,enum imageFormat)" (?cudaBayerToRGB@@YA?AW4cudaError@@PEAEPEAUuchar3@@_K2W4imageFormat@@@Z) で参照されました	CudaCropTest	C:\Users\owner\source\repos\CudaCropTest\CudaCropTest\cudaBayer.obj	1

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.4\include\nppi_color_conversion.h

NppStatus 
nppiCFAToRGB_8u_C1C3R(const Npp8u * pSrc, int nSrcStep, NppiSize oSrcSize, NppiRect oSrcROI, 
                            Npp8u * pDst, int nDstStep, NppiBayerGridPosition eGrid, NppiInterpolationMode eInterpolation);

ヘッダーファイルは存在するのに、関連したフォルダ内を調べてもc/cpp/cuファイルがありません。Copilotに尋ねると、nppi.libのようなlibファイルへのリンクが正しく設定されていないと教えてくれました。

どのlibファイルを追加したら良いのか分からないのでググると、NVIDIA NPPのキュメントを発見でき、

NPPICC, color conversion and sampling functions in nppi_color_conversion.h,

と書かれています。nppicc.lib をリンカーに追加します。

プロジェクト(P) > CudaCropTest のプロパティの(P) をクリックします。

構成プロパティ > リンカー > 入力 > 追加の依存ファイルに nppicc.lib を追加します。
（構成プロパティ > リンカー > 全般 > 追加のライブラリディレクトリに $(CudaToolkitLibDir)=C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.4\lib\x64 があるので、上記の対応のみでOKです）

以上で kernel.cu を再びビルドできるようになりました。

画像処理

サンプル画像の用意

https://placehold.jp/ を使用してサンプル画像を用意します。

https://placehold.jp/3d4070/ffffff/2000x1500.png

この画像を CudaCropTest\images\sample-2000x1500.png として保存しました。

画像の読み書き

関数シグネチャを参照してコードを書きます。

kernel.cu

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

#include "imageIO.h"
#include "cudaCrop.h"

int main()
{
    const char* input_path = R"(C:\Users\owner\source\repos\CudaCropTest\CudaCropTest\images\sample-2000x1500.png)";
    const char* output_path = R"(C:\Users\owner\source\repos\CudaCropTest\CudaCropTest\images\test.png)";

    uchar3 *input_image = NULL;
    int width = 0, height = 0;

    // Load image
    if (!loadImage(input_path, &input_image, &width, &height)) {
        fprintf(stderr, "loadImage failed!");
		return 1;
    }
    printf("Image loaded: %d x %d\n", width, height);

    // Save image
    if (!saveImage(output_path, input_image, width, height)) {
		fprintf(stderr, "saveImage failed!");
        return 1;
    }

    return 0;
}

実行すると入力画像と同じ見た目の CudaCropTest\images\test.png が保存されました。

画像の切り抜き

cudaCrop 関数を使用して画像を切り抜きます。

kernel.cu

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

#include "imageIO.h"
#include "cudaCrop.h"
#include "cudaMappedMemory.h"

class CudaMemory {
public:
	CudaMemory(uchar3* ptr) : ptr_(ptr) {}
	~CudaMemory() { cudaFree(ptr_); }
private:
	uchar3* ptr_;
};

int main()
{
	const char* input_path = R"(C:\Users\owner\source\repos\CudaCropTest\CudaCropTest\images\sample-2000x1500.png)";
	const char* output_path = R"(C:\Users\owner\source\repos\CudaCropTest\CudaCropTest\images\test.png)";

	uchar3* input_image = NULL;
	uchar3* output_image = NULL;
	CudaMemory input_memory(input_image), output_memory(output_image);

	int width = 0, height = 0;

	// Load image
	if (!loadImage(input_path, &input_image, &width, &height)) {
		fprintf(stderr, "loadImage failed!");
		return 1;
	}
	printf("Image loaded: %d x %d\n", width, height);

	int crop_width = width / 2;
	int crop_height = height / 3;

	// Allocate output image
	if (!cudaAllocMapped(&output_image, sizeof(uchar3) * crop_width * crop_height)) {
		fprintf(stderr, "cudaAllocMapped failed!");
		return 1;
	}

    // Crop image
	int4 roi = {
		width / 4,
		height / 3,
		width / 4 + crop_width,
		height / 3 + crop_height
	};
	cudaError_t cudaStatus = cudaCrop(
		input_image,
		output_image,
		roi,
		width,
		height
	);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaCrop failed! cudaStatus=%d", cudaStatus);
		return 1;
	}

	// Save image
	if (!saveImage(output_path, output_image, crop_width, crop_height)) {
		fprintf(stderr, "saveImage failed!");
		return 1;
	}

	return 0;
}

実行すると、切り抜き画像が生成されました。

あとがき

ChatGPTとライブラリのおかげで無事にCUDAを使用して、画像を切り抜くことができました。

そのうち以下に取り組もうと思います。

CPU と GPU での切り抜き速度の比較
CUDA上での並列処理
パラメーターの調整（現在の cudaCrop 関数はパラメーターを調整できない: リンク）