WindowsでcuDNNを使用して畳み込みを行う

前回の日記で書いた方法でChainerのコードを調べつつ、WindowsでcuDNNを使用して畳み込みを行うことができたので、方法を示しておく。

使用バージョン

Windows 10
CUDA7.5
cuDNN v5
Visual Studio 2015

cuDNN v5を使用するには、CUDA7.5が必要になる。
CUDA7.5をインストールするにはVisual Studio 2013が必要になるが、cuDNNを使用する場合は間接的に利用することになるので、Visual Studio 2015で開発できる。

プロジェクトの設定

ディレクトリの設定

CUDA7.5をインストールすると、Visual Studio 2013用の新規プロジェクトのテンプレートがインストールされるので、Visual Stuidio 2013で開発する場合はインクルードディレクトリなどの設定は不要だが、Visual Studio 2015を使用する場合は、インクルードディレクトリとライブラリディレクトリは手動で設定する必要がある。

インクルードディレクトリ	$(CUDA_PATH)\include
ライブラリディレクトリ	$(CUDA_PATH)\lib\x64

リンクライブラリ追加

以下のライブラリをリンクの入力に追加する。

cudnn.lib
cudart.lib

cuDNNで畳み込み

準備処理

畳み込み関数の呼び出しの準備として以下の処理が必要になる。

cudnnCreateでハンドル作成
cudnnCreateTensorDescriptor、cudnnSetTensor4dDescriptorで入力、出力の設定を行う
cudnnCreateFilterDescriptor、cudnnSetFilter4dDescriptorでフィルターの設定を行う
cudnnCreateConvolutionDescriptor、cudnnSetConvolution2dDescriptorで畳み込みの設定を行う
cudnnGetConvolutionForwardAlgorithmで畳み込みのアルゴリズムを選択する
cudnnGetConvolutionForwardWorkspaceSizeで作業用メモリサイズを取得する

入力、フィルター、出力

入力は、[ミニバッチ][チャンネル][入力の高さ][入力の幅]の四次元配列となる。*1
フィルターは、[フィルター数][チャンネル][フィルター高さ][フィルター幅]の四次元配列となる。
出力は、[ミニバッチ][フィルター数][入力の高さ][入力の幅]の四次元配列となる。

入力とフィルターの値は、ホスト(CPU)側で配列を作成して、cudaMallocでデバイス(GPU)側のメモリを確保し、cudaMemcpyで転送する。

出力は、cudaMallocでデバイス(GPU)側のメモリを確保しておき、cudaMemcpyでホスト(CPU)側に転送する。

これらは、CUDAの関数を使用した処理になる。

畳み込み

cuDNNでの2次元の畳み込みはcudnnConvolutionForward関数で行える。
上記で準備した変数を引数に与える。

詳細は、CUDNN LIBRARYのUser Guideを参照。

後始末

cudnnDestroyConvolutionDescriptor、cudnnDestroyFilterDescriptor、cudnnDestroyTensorDescriptor、cudnnDestroyでハンドルを破棄する。

コード例

cuDNNTest.cpp

#include <iostream>

#include <cuda.h>
#include <cudnn.h>

#include "error_util.h"

using namespace std;

int main() {
	const int minibatch_size = 1;
	const int feature_num = 2;
	const int filter_num = 3;
	const int in_size = 8;
	const int filter_size = 5;

	// 入力
	float srcData[minibatch_size][feature_num][in_size][in_size] = {
		{
			{ { 0, 1, 0, 0, 0, 0, 0, 0 }, { 0, 1, 0, 0, 0, 0, 0, 0 }, { 0, 1, 0, 0, 0, 0, 1, 0 }, { 1, 1, 0, 0, 0, 0, 0, 1 }, { 0, 1, 0, 1, 0, 0, 0, 0 }, { 0, 1, 0, 0, 0, 0, 0, 0 }, { 0, 1, 0, 0, 0, 0, 0, 0 }, { 0, 1, 0, 0, 0, 0, 0, 0 } },
			{ { 1, 1, 1, 1, 1, 1, 1, 1 }, { 1, 1, 1, 1, 1, 1, 1, 1 }, { 1, 1, 1, 1, 1, 1, 1, 1 }, { 1, 1, 1, 1, 1, 1, 1, 1 }, { 1, 1, 1, 1, 1, 1, 1, 1 }, { 1, 1, 1, 1, 1, 1, 1, 1 }, { 1, 1, 1, 1, 1, 1, 1, 1 }, { 1, 1, 1, 1, 1, 1, 1, 1 } }
		}
	};

	float *srcData_dev;
	checkCudaErrors(cudaMalloc((void**)&srcData_dev, sizeof(srcData)));

	// Copy input vectors from host memory to GPU buffers.
	checkCudaErrors(cudaMemcpy(srcData_dev, srcData, sizeof(srcData), cudaMemcpyHostToDevice));

	// フィルター係数
	float filterData[filter_num][feature_num][filter_size][filter_size] = {
		{
			{ { 0.1f, 1.0f, 0.1f, 0.5f, 0.2f }, { 0.1f, 1.0f, 0.1f, 0.5f, 0.2f }, { 0.1f, 1.0f, 0.1f, 0.5f, 0.2f }, { 0.1f, 1.0f, 0.1f, 0.5f, 0.2f }, { 0.1f, 1.0f, 0.1f, 0.5f, 0.2f } },
			{ { 1.0f, 0.1f, 0.1f, 0.5f, 1.0f }, { 1.0f, 0.1f, 0.1f, 0.5f, 1.0f }, { 1.0f, 0.1f, 0.1f, 0.5f, 1.0f }, { 1.0f, 0.1f, 0.1f, 0.5f, 1.0f }, { 1.0f, 0.1f, 0.1f, 0.5f, 1.0f } }
		},
		{
			{ { 0.1f, 1.0f, 0.1f, 0.5f, 0.2f }, { 0.1f, 1.0f, 0.1f, 0.5f, 0.2f }, { 0.1f, 1.0f, 0.1f, 0.5f, 0.2f }, { 0.1f, 1.0f, 0.1f, 0.5f, 0.2f }, { 0.1f, 1.0f, 0.1f, 0.5f, 0.2f } },
			{ { 1.0f, 0.1f, 0.1f, 0.5f, 1.0f }, { 1.0f, 0.1f, 0.1f, 0.5f, 1.0f }, { 1.0f, 0.1f, 0.1f, 0.5f, 1.0f }, { 1.0f, 0.1f, 0.1f, 0.5f, 1.0f }, { 1.0f, 0.1f, 0.1f, 0.5f, 1.0f } }
		},
		{
			{ { 0.1f, 1.0f, 0.1f, 0.5f, 0.2f }, { 0.1f, 1.0f, 0.1f, 0.5f, 0.2f }, { 0.1f, 1.0f, 0.1f, 0.5f, 0.2f }, { 0.1f, 1.0f, 0.1f, 0.5f, 0.2f }, { 0.1f, 1.0f, 0.1f, 0.5f, 0.2f } },
			{ { 1.0f, 0.1f, 0.1f, 0.5f, 1.0f }, { 1.0f, 0.1f, 0.1f, 0.5f, 1.0f }, { 1.0f, 0.1f, 0.1f, 0.5f, 1.0f }, { 1.0f, 0.1f, 0.1f, 0.5f, 1.0f }, { 1.0f, 0.1f, 0.1f, 0.5f, 1.0f } }
		}
	};

	float *filterData_dev;
	checkCudaErrors(cudaMalloc((void**)&filterData_dev, sizeof(filterData)));

	// Copy input vectors from host memory to GPU buffers.
	checkCudaErrors(cudaMemcpy(filterData_dev, filterData, sizeof(filterData), cudaMemcpyHostToDevice));

	// 出力
	float dstData[minibatch_size][filter_num][in_size][in_size];
	float *dstData_dev;
	checkCudaErrors(cudaMalloc((void**)&dstData_dev, sizeof(dstData)));


	// 畳み込み
	cudnnHandle_t cudnnHandle;
	cudnnTensorDescriptor_t srcTensorDesc, dstTensorDesc;
	cudnnFilterDescriptor_t filterDesc;
	cudnnConvolutionDescriptor_t convDesc;

	checkCUDNN(cudnnCreate(&cudnnHandle));
	checkCUDNN(cudnnCreateTensorDescriptor(&srcTensorDesc));
	checkCUDNN(cudnnCreateTensorDescriptor(&dstTensorDesc));
	checkCUDNN(cudnnCreateFilterDescriptor(&filterDesc));
	checkCUDNN(cudnnCreateConvolutionDescriptor(&convDesc));

	checkCUDNN(cudnnSetTensor4dDescriptor(srcTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, minibatch_size, feature_num, in_size, in_size));
	checkCUDNN(cudnnSetTensor4dDescriptor(dstTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, minibatch_size, filter_num, in_size, in_size));
	checkCUDNN(cudnnSetFilter4dDescriptor(filterDesc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, filter_num, feature_num, filter_size, filter_size));
	checkCUDNN(cudnnSetConvolution2dDescriptor(convDesc, 2/*pad_h*/, 2/*pad_w*/, 1/*stride_h*/, 1/*stride_w*/, 1, 1, CUDNN_CROSS_CORRELATION));

	cudnnConvolutionFwdAlgo_t algo;
	checkCUDNN(cudnnGetConvolutionForwardAlgorithm(cudnnHandle,
		srcTensorDesc,
		filterDesc,
		convDesc,
		dstTensorDesc,
		CUDNN_CONVOLUTION_FWD_PREFER_FASTEST,
		0,
		&algo
		));

	cout << "Fastest algorithm is Algo " << algo << endl;

	size_t sizeInBytes = 0;
	checkCUDNN(cudnnGetConvolutionForwardWorkspaceSize(cudnnHandle,
		srcTensorDesc,
		filterDesc,
		convDesc,
		dstTensorDesc,
		algo,
		&sizeInBytes));

	void* workSpace = NULL;
	if (sizeInBytes != 0)
	{
		checkCudaErrors(cudaMalloc(&workSpace, sizeInBytes));
	}

	float alpha = 1.0f;
	float beta = 0.0f;
	checkCUDNN(cudnnConvolutionForward(cudnnHandle,
		&alpha,
		srcTensorDesc,
		srcData_dev,
		filterDesc,
		filterData_dev,
		convDesc,
		algo,
		workSpace,
		sizeInBytes,
		&beta,
		dstTensorDesc,
		dstData_dev));

	// Copy output vector from GPU buffer to host memory.
	checkCudaErrors(cudaMemcpy(dstData, dstData_dev, sizeof(dstData), cudaMemcpyDeviceToHost));

	for (int i = 0; i < filter_num; i++) {
		for (int y = 0; y < in_size; y++) {
			cout << "{";
			for (int x = 0; x < in_size; x++) {
				cout << dstData[0][i][y][x] << ", ";
			}
			cout << "}, ";
		}
		cout << endl;
	}

	checkCUDNN(cudnnDestroyConvolutionDescriptor(convDesc));
	checkCUDNN(cudnnDestroyFilterDescriptor(filterDesc));
	checkCUDNN(cudnnDestroyTensorDescriptor(srcTensorDesc));
	checkCUDNN(cudnnDestroyTensorDescriptor(dstTensorDesc));
	checkCUDNN(cudnnDestroy(cudnnHandle));

	return 0;
}

error_util.h

#pragma once

#include <sstream>

#define FatalError(s) {                                                \
    std::stringstream _where, _message;                                \
    _where << __FILE__ << ':' << __LINE__;                             \
    _message << std::string(s) + "\n" << __FILE__ << ':' << __LINE__;\
    std::cerr << _message.str() << "\nAborting...\n";                  \
    cudaDeviceReset();                                                 \
    exit(EXIT_FAILURE);                                                \
}

#define checkCUDNN(status) {                                           \
    std::stringstream _error;                                          \
    if (status != CUDNN_STATUS_SUCCESS) {                              \
      _error << "CUDNN failure\nError: " << cudnnGetErrorString(status); \
      FatalError(_error.str());                                        \
	    }                                                                  \
}

#define checkCudaErrors(status) {                                      \
    std::stringstream _error;                                          \
    if (status != 0) {                                                 \
      _error << "Cuda failure\nError: " << cudaGetErrorString(status); \
      FatalError(_error.str());                                        \
	    }                                                                  \
}

#define checkCublasErrors(status) {                                    \
    std::stringstream _error;                                          \
    if (status != 0) {                                                 \
      _error << "Cublas failure\nError code " << status;        \
      FatalError(_error.str());                                        \
	    }                                                                  \
}

比較用にPython(Chainer)のコードを作成して出力が一致することを確認した。

import numpy as np
import chainer
from chainer import cuda, Function, Variable, optimizers, function, link
from chainer import Link, Chain
import chainer.functions as F
import chainer.links as L

feature_num = 2
k = 3

w_data = [
		[
			[ [ 0.1, 1.0, 0.1, 0.5, 0.2 ], [ 0.1, 1.0, 0.1, 0.5, 0.2 ], [ 0.1, 1.0, 0.1, 0.5, 0.2 ], [ 0.1, 1.0, 0.1, 0.5, 0.2 ], [ 0.1, 1.0, 0.1, 0.5, 0.2 ] ],
			[ [ 1.0, 0.1, 0.1, 0.5, 1.0 ], [ 1.0, 0.1, 0.1, 0.5, 1.0 ], [ 1.0, 0.1, 0.1, 0.5, 1.0 ], [ 1.0, 0.1, 0.1, 0.5, 1.0 ], [ 1.0, 0.1, 0.1, 0.5, 1.0 ] ]
		],
		[
			[ [ 0.1, 1.0, 0.1, 0.5, 0.2 ], [ 0.1, 1.0, 0.1, 0.5, 0.2 ], [ 0.1, 1.0, 0.1, 0.5, 0.2 ], [ 0.1, 1.0, 0.1, 0.5, 0.2 ], [ 0.1, 1.0, 0.1, 0.5, 0.2 ] ],
			[ [ 1.0, 0.1, 0.1, 0.5, 1.0 ], [ 1.0, 0.1, 0.1, 0.5, 1.0 ], [ 1.0, 0.1, 0.1, 0.5, 1.0 ], [ 1.0, 0.1, 0.1, 0.5, 1.0 ], [ 1.0, 0.1, 0.1, 0.5, 1.0 ] ]
		],
		[
			[ [ 0.1, 1.0, 0.1, 0.5, 0.2 ], [ 0.1, 1.0, 0.1, 0.5, 0.2 ], [ 0.1, 1.0, 0.1, 0.5, 0.2 ], [ 0.1, 1.0, 0.1, 0.5, 0.2 ], [ 0.1, 1.0, 0.1, 0.5, 0.2 ] ],
			[ [ 1.0, 0.1, 0.1, 0.5, 1.0 ], [ 1.0, 0.1, 0.1, 0.5, 1.0 ], [ 1.0, 0.1, 0.1, 0.5, 1.0 ], [ 1.0, 0.1, 0.1, 0.5, 1.0 ], [ 1.0, 0.1, 0.1, 0.5, 1.0 ] ]
		]
	]

w = np.array(w_data, dtype=np.float32)

model = Chain(
    layer1=L.Convolution2D(in_channels = feature_num, out_channels = k, ksize = 5, pad = 2, nobias = True, initialW = w))

model.to_gpu()

src_data = [
        [
            [ [ 0, 1, 0, 0, 0, 0, 0, 0 ], [ 0, 1, 0, 0, 0, 0, 0, 0 ], [ 0, 1, 0, 0, 0, 0, 1, 0 ], [ 1, 1, 0, 0, 0, 0, 0, 1 ], [ 0, 1, 0, 1, 0, 0, 0, 0 ], [ 0, 1, 0, 0, 0, 0, 0, 0 ], [ 0, 1, 0, 0, 0, 0, 0, 0 ], [ 0, 1, 0, 0, 0, 0, 0, 0 ] ],
            [ [ 1, 1, 1, 1, 1, 1, 1, 1 ], [ 1, 1, 1, 1, 1, 1, 1, 1 ], [ 1, 1, 1, 1, 1, 1, 1, 1 ], [ 1, 1, 1, 1, 1, 1, 1, 1 ], [ 1, 1, 1, 1, 1, 1, 1, 1 ], [ 1, 1, 1, 1, 1, 1, 1, 1 ], [ 1, 1, 1, 1, 1, 1, 1, 1 ], [ 1, 1, 1, 1, 1, 1, 1, 1 ] ]
        ]
    ]

src = Variable(cuda.to_gpu(np.array(src_data, dtype=np.float32)))

dst = model.layer1(src)

dst_data = cuda.to_cpu(dst.data)

print(dst_data)

*1:「Users Guide」ではNCHWという用語が使われているが、2.3.2. 4-D Tensor Descriptorに、N,C,H,Wは、それぞれbatch number, the number of feature maps, the height and the widthという説明がある。CはChannelsの略と思われる。

TadaoYamaokaの開発日記

個人開発しているスマホアプリや将棋AIの開発ネタを中心に書いていきます。