この記事について

CMakeを使って、CUDAアプリケーション開発用プロジェクトを作ります。WindowsとLinuxの両方でビルドできるようにします。
が、今回はLinux側は環境構築していないので、動作未確認です。

CMakeを使うことのメリット
- マルチプラットフォーム対応
- インクルードパスやライブラリ設定を自動化できる
- (特に、GitHub等で他の人に渡すときに便利)

環境

Windows10 64-bit
- Visual Studio 2017 Community
  - Version 15.6.7
  - _MSC_VER = 1913
- CUDA 10.0

CUDAとVisualStudioのバージョン問題

CUDA 9.0だと、_MSC_VER = 1911以降のVisual Studioをサポートしていません。公式に言われているわけではなさそうですが、_MSC_VERのバージョンチェックでエラーになります。
以下のスレッドでず～っと、語られてきた問題ですが、現時点においても明確な解決策がありません。唯一の回避策はVSのバージョンを下げる、ということでした。そうするくらいなら、最新のCUDA10.0を使った方がいいと思います。
https://devtalk.nvidia.com/default/topic/1022648/cuda-9-unsupported-visual-studio-version-error/

ただ、CUDA9とCUDA10の共存は難しいようなので、最近はやりの機械学習ライブラリを使っている方は、そのライブラリがCUDA10.0に対応しているか確認してから乗り換えた方がいいと思います。

高速化するロジック

以下のように、n個の配列の各要素を2倍する関数を、CUDAで実装して高速化することを考えます。

void original_vecDouble(int *in, int *out, const int n)
{
	for (int i = 0; i < n; i++) {
		out[i] = in[i] * 2;
	}
}

CUDAの準備

Windowsの場合

下記からダウンロードします。
https://developer.nvidia.com/cuda-downloads

こちらもご参照ください。https://qiita.com/take-iwiw/items/0bf12bec9c3bade2e331

プロジェクトの作成

└─TestCUDA
    │  CMakeLists.txt
    │  main.cpp

CMakeLists.txt

cmake_minimum_required(VERSION 2.8)
project(MyCudaProject)

# For CUDA
find_package(CUDA REQUIRED)
if(CUDA_FOUND)
	message("Cuda version: " ${CUDA_VERSION})
	CUDA_ADD_EXECUTABLE(Main
		main.cu
	)
endif()

CMakeLists.txtでは、CUDAをfind_packageして、CUDA用の実行ファイル開発プロジェクトをmain.cuから作るようにしています。

main.cu

# include <stdio.h>
# include <stdlib.h>
# include "cuda_runtime.h"

static void vecDouble(int *in, int*out, const int n);

int main()
{
	printf("Hello\n");

	const int n = 10;
	int *in = new int[n];
	int *out = new int[n];
	int *answer = new int[n];

	for (int i = 0; i < n; i++) in[i] = rand() % 100;
	for (int i = 0; i < n; i++) answer[i] = in[i] * 2;

	vecDouble(in, out, n);

	int i;
	for (i = 0; i < n; i++) {
		if (answer[i] != out[i]) {
			printf("error at index = %d\n", i);
			break;
		}
	}
	printf("OK\n");

	delete[] in;
	delete[] out;
	delete[] answer;

	return 0;
}

__global__ void kernel_vecDouble(int *in, int *out, const int n)
{
	int i = threadIdx.x;
	if (i < n) {
		out[i] = in[i] * 2;
	}
}

static void vecDouble(int *hIn, int *hOut, const int n)
{
	int *dIn;
	int *dOut;
	cudaMallocHost((void**)&dIn, n * sizeof(int));
	cudaMallocHost((void**)&dOut, n * sizeof(int));
	cudaMemcpy(dIn, hIn, n * sizeof(int), cudaMemcpyHostToDevice);

	kernel_vecDouble<<<1, n>>>(dIn, dOut, n);
	cudaDeviceSynchronize();

	cudaMemcpy(hOut, dOut, n * sizeof(int), cudaMemcpyDeviceToHost);
	cudaFree(dIn);
	cudaFree(dOut);
}

main.cuは上記の通りで、kernel_vecDouble()がCUDAで実行されるカーネル関数です。それを制御したりメモリ管理する関数がvecDouble()です。
main関数がvecDouble()を呼んでいます。

とりあえずこれで、CUDAアプリケーション開発が出来ます。

プロジェクトを分ける

続いて、将来的にプロジェクトが大きくなることを見越して、CUDA用コードと呼び出し用コードを別プロジェクトにしてみます。
全体として見通しが良くなるのと、CUDAコードでアルゴリズムをがりがり実装する人とUI等のアプリケーションを実装する人でプロジェクトを分けられる点がメリットだと思います。
上手く工夫すれば、CUDAコードの単体テストもやりやすくなると思います。

├─TestCuda2
│  │  CMakeLists.txt
│  │
│  ├─CudaVecDouble
│  │      CMakeLists.txt
│  │      CudaVecDouble.cu
│  │      CudaVecDouble.h
│  │
│  └─Main
│          main.cpp

Mainプロジェクト

最上位のCMakeLists.txt

cmake_minimum_required(VERSION 2.8)
project(MyCudaProject)

# Create Main project
set(SOURCES
	Main/main.cpp
)

set(HEADERS
	# Main/main.h
)

add_executable(Main
	${SOURCES}
	${HEADERS}
)

# Add sub project
add_subdirectory(CudaVecDouble)
target_include_directories(Main PUBLIC CudaVecDouble)
target_link_libraries(Main CudaVecDouble)

最上位のCMakeLists.txtでは、main.cppを持つMainプロジェクトを作成します。
サブプロジェクトとしてCudaVecDoubleを追加して、インクルードパス設定とリンク追加します。

main.cpp

# include <stdio.h>
# include <stdlib.h>
# include "CudaVecDouble.h"

int main()
{
	printf("Hello\n");

	const int n = 10;
	int *in = new int[n];
	int *out = new int[n];
	int *answer = new int[n];

	for (int i = 0; i < n; i++) in[i] = rand() % 100;
	for (int i = 0; i < n; i++) answer[i] = in[i] * 2;

	vecDouble(in, out, n);

	int i;
	for (i = 0; i < n; i++) {
		if (answer[i] != out[i]) {
			printf("error at index = %d\n", i);
			break;
		}
	}
	printf("OK\n");

	delete[] in;
	delete[] out;
	delete[] answer;

	return 0;
}

main.cppはvecDouble()関数(CUDAコード呼び出しのwrapper)を呼ぶだけのシンプルな形になりました。
このプログラムはCUDAとは関係ない純粋なC/C++プログラムであり、コンパイラも通常のC/C++が使われます。

CudaVecDoubleプロジェクト

CMakeLists.txt

find_package(CUDA REQUIRED)
if(CUDA_FOUND)
	message("Cuda version: " ${CUDA_VERSION})
	CUDA_ADD_LIBRARY(CudaVecDouble STATIC
		CudaVecDouble.cu
		CudaVecDouble.h
	)
endif()

CUDA側プロジェクトのCMakeLists.txtはこれだけです。今回は、CUDA_ADD_LIBRARYを使って静的ライブラリを作るようにしています。

CudaVecDouble.cu

# include "cuda_runtime.h"
# include "CudaVecDouble.h"

__global__ void kernel_vecDouble(int *in, int *out, const int n)
{
	int i = threadIdx.x;
	if (i < n) {
		out[i] = in[i] * 2;
	}
}

void vecDouble(int *hIn, int *hOut, const int n)
{
	int *dIn;
	int *dOut;
	cudaMallocHost((void**)&dIn, n * sizeof(int));
	cudaMallocHost((void**)&dOut, n * sizeof(int));
	cudaMemcpy(dIn, hIn, n * sizeof(int), cudaMemcpyHostToDevice);

	kernel_vecDouble<<<1, n>>>(dIn, dOut, n);
	cudaDeviceSynchronize();

	cudaMemcpy(hOut, dOut, n * sizeof(int), cudaMemcpyDeviceToHost);
	cudaFree(dIn);
	cudaFree(dOut);
}

CudaVecDouble.h

# ifndef CUDA_VEC_DOUBLE_H
# define CUDA_VEC_DOUBLE_H

void vecDouble(int *hIn, int *hOut, const int n);

# endif	/* CUDA_VEC_DOUBLE_H */

CudaVecDouble.cu/hはどちらも、CUDAで実装するロジックだけが記述されるようになりました。