WindowsでcuDNNを使用して畳み込みを行う(活性化関数追加)

前回の日記で実装した畳み込みのコードに活性化関数を追加しました。

活性化関数は、cudnnCreateActivationDescriptorでハンドルを作成し、cudnnSetActivationDescriptorでどの関数を使用するか選択し、cudnnActivationForwardで実行できます。

活性化関数として使用可能なのは、

CUDNN_ACTIVATION_SIGMOID	シグモイド関数
CUDNN_ACTIVATION_RELU	ReLU(ランプ)関数
CUDNN_ACTIVATION_TANH	tanh(双曲線正接)関数
CUDNN_ACTIVATION_CLIPPED_RELU	Clipped-ReLU(ランプ)関数

です。

cudnnActivationForwardの入力と出力は、Users Guideに「In-place operation is allowed for this routine」と書かれているので同じメモリでも大丈夫です。

入力と出力のディメンションは同じであれば、どのような形でもよいようです。
Chainerのソースでは、[ミニバッチ数][フィルター数×高さ×幅][1][1]の4次元配列となっていましたが、[ミニバッチ数][フィルター数][高さ][幅]としても問題ありませんでした。

前回のコードに活性化関数を追加したコードを以下に示します。

コード例

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 dstData[minibatch_size][filter_num][in_size][in_size];
	float *dstData_dev;
	checkCudaErrors(cudaMalloc((void**)&dstData_dev, sizeof(dstData)));

	// フィルター係数
	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 biasData[filter_num] = {
		0, 0.5f, -5.0f
	};
	float *biasData_dev;
	checkCudaErrors(cudaMalloc((void**)&biasData_dev, sizeof(biasData)));
	// Copy input vectors from host memory to GPU buffers.
	checkCudaErrors(cudaMemcpy(biasData_dev, biasData, sizeof(biasData), cudaMemcpyHostToDevice));


	// 畳み込み準備
	cudnnHandle_t cudnnHandle;
	cudnnTensorDescriptor_t srcTensorDesc, dstTensorDesc, biasTensorDesc;
	cudnnFilterDescriptor_t filterDesc;
	cudnnConvolutionDescriptor_t convDesc;
	cudnnActivationDescriptor_t  activDesc;

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

	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(cudnnSetTensor4dDescriptor(biasTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, 1, filter_num, 1, 1));
	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));
	checkCUDNN(cudnnSetActivationDescriptor(activDesc, CUDNN_ACTIVATION_RELU, CUDNN_PROPAGATE_NAN, 0.0/*reluCeiling*/));

	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));

	// バイアス
	alpha = 1.0f;
	beta = 1.0f;
	checkCUDNN(cudnnAddTensor(cudnnHandle, &alpha, biasTensorDesc, biasData_dev, &beta, dstTensorDesc, dstData_dev));

	// 活性化関数
	alpha = 1.0f;
	beta = 0.0f;
	checkCUDNN(cudnnActivationForward(cudnnHandle,
		activDesc,
		&alpha,
		dstTensorDesc,
		dstData_dev,
		&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(cudnnDestroyActivationDescriptor(activDesc));
	checkCUDNN(cudnnDestroyConvolutionDescriptor(convDesc));
	checkCUDNN(cudnnDestroyFilterDescriptor(filterDesc));
	checkCUDNN(cudnnDestroyTensorDescriptor(srcTensorDesc));
	checkCUDNN(cudnnDestroyTensorDescriptor(dstTensorDesc));
	checkCUDNN(cudnnDestroyTensorDescriptor(biasTensorDesc));
	checkCUDNN(cudnnDestroy(cudnnHandle));

	return 0;
}

実行結果

{6.3, 5.4, 11.1, 8.4, 8.3, 8.6, 5.2, 4.6, }, {8.5, 8.2, 14.9, 11.2, 11, 11.5, 7.4, 5.9, }, {10.6, 10.2, 19.1, 14.1, 14.7, 14.3, 9.1, 7.1, }, {10.6, 10.2, 19.1, 14.1, 14.7, 14.3, 9.1, 7.1, }, {10.6, 10.2, 19.1, 14.1, 14.7, 14.3, 9.1, 7.1, }, {10.6, 10.2, 19.1, 14.1, 14.5, 13.8, 9, 6.1, }, {8.4, 7.4, 15.3, 11.3, 11.8, 10.9, 6.8, 4.8, }, {6.3, 5.4, 11.1, 8.4, 8.1, 8.1, 5.1, 3.6, },
{6.8, 5.9, 11.6, 8.9, 8.8, 9.1, 5.7, 5.1, }, {9, 8.7, 15.4, 11.7, 11.5, 12, 7.9, 6.4, }, {11.1, 10.7, 19.6, 14.6, 15.2, 14.8, 9.6, 7.6, }, {11.1, 10.7, 19.6, 14.6, 15.2, 14.8, 9.6, 7.6, }, {11.1, 10.7, 19.6, 14.6, 15.2, 14.8, 9.6, 7.6, }, {11.1, 10.7, 19.6, 14.6, 15, 14.3, 9.5, 6.6, }, {8.9, 7.9, 15.8, 11.8, 12.3, 11.4, 7.3, 5.3, }, {6.8, 5.9, 11.6, 8.9, 8.6, 8.6, 5.6, 4.1, },
{1.3, 0.4, 6.1, 3.4, 3.3, 3.6, 0.2, 0, }, {3.5, 3.2, 9.9, 6.2, 6, 6.5, 2.4, 0.899999, }, {5.6, 5.2, 14.1, 9.1, 9.7, 9.3, 4.1, 2.1, }, {5.6, 5.2, 14.1, 9.1, 9.7, 9.3, 4.1, 2.1, }, {5.6, 5.2, 14.1, 9.1, 9.7, 9.3, 4.1, 2.1, }, {5.6, 5.2, 14.1, 9.1, 9.5, 8.8, 4, 1.1, }, {3.4, 2.4, 10.3, 6.3, 6.8, 5.9, 1.8, 0, }, {1.3, 0.4, 6.1, 3.4, 3.1, 3.1, 0.0999994, 0, },

ReLU(ランプ)関数が適用されてマイナスの値がちゃんと0になっています。