読者です 読者をやめる 読者になる 読者になる

TadaoYamaokaの日記

山岡忠夫 Home で公開しているプログラムの開発ネタを中心に書いていきます。

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

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

使用バージョン

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で畳み込み

準備処理

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

  1. cudnnCreateでハンドル作成
  2. cudnnCreateTensorDescriptor、cudnnSetTensor4dDescriptorで入力、出力の設定を行う
  3. cudnnCreateFilterDescriptor、cudnnSetFilter4dDescriptorでフィルターの設定を行う
  4. cudnnCreateConvolutionDescriptor、cudnnSetConvolution2dDescriptorで畳み込みの設定を行う
  5. cudnnGetConvolutionForwardAlgorithmで畳み込みのアルゴリズムを選択する
  6. 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の略と思われる。