CUDAを使ってみた

SL policy networkをChainerを使って学習した結果を、囲碁プログラムに組み込む際、C++のプログラムからDCNNの実行のためにpythonプログラムを呼ぶとオーバーヘッドが大きい。
できれば、DCNNをC++から実行したい。
順伝播のみであれば、ディープラーニングフレームワークを使用しなくて実装は難しくない。
そこで、CUDAを直接利用して実装することを考えている。

CUDAの勉強のために、サンプル的なプログラムを作って実行速度など確認してみた。

まずは、2つのfloatのベクトルの要素同士の積を計算するプログラムで試してみた。
比較のために、CPUで実行した場合と、AVX命令を使った場合も試してみた。

ベクトルの要素数は、1024 * 1024 * 128で、実行速度を測った結果は以下の通り。

CPU	164 ms	1スレッド
AVX	157 ms	1スレッド
GPU	372 ms	メモリ転送含む
GPU	15 ms	メモリ転送除く

※CPUは、Core i7-6700K 4GHz
※GPUは、GeForce GTX 760 2GBを使用

メモリ転送量は、1ベクトル512MBあるので、入力2ベクトル、出力1ベクトルで合計1.5GBになる。

このプログラムではメモリ転送に時間がかかっており、GPUの方が遅いという結果になった。
しかし、メモリ転送を除く演算部分のみであれば、GPUはCPUより、10倍以上高速である。

DCNNを実行する場合、メモリ転送量の多いネットワークのパラメータは固定であるため、起動時に1回転送すればよく、GPUの方が速くなる可能性があるが、囲碁プログラムはルート並列化して動かしているので、CPU側を並列にした場合に、GPUが1つだと待ち合わせが発生するので、微妙かもしれない。
GPUの処理をキューにして管理するなど複雑なことを行う必要がありそうだ。

測定に使ったプログラムコードは以下の通り。

CPU

#include <iostream>
#include <random>
#include <chrono>

using namespace std;

const size_t SIZE = 1024 * 1024 * 128;
_declspec(align(32)) float a[SIZE];
_declspec(align(32)) float b[SIZE];
_declspec(align(32)) float c[SIZE];

int main() {
	// ランダムな値を生成
	random_device seed;
	mt19937 rnd(seed());
	for (size_t i = 0; i < SIZE; i++) {
		a[i] = (float)rnd() / rnd.max();
		b[i] = (float)rnd() / rnd.max();
	}

	auto start = chrono::system_clock::now();
	// 要素同士の積
	for (size_t i = 0; i < SIZE; i ++) {
		c[i] = a[i] * b[i];
	}
	auto end = chrono::system_clock::now();
	auto msec = chrono::duration_cast<std::chrono::milliseconds>(end - start).count();

	cout << "duration = " << msec << " msec" << endl;

	for (size_t i = 0; i < 8; i++) {
		cout << c[i] << endl;
	}

	return 0;
}

AVX

#include <immintrin.h>
#include <iostream>
#include <random>
#include <chrono>

using namespace std;

const size_t SIZE = 1024 * 1024 * 128;
_declspec(align(32)) float a[SIZE];
_declspec(align(32)) float b[SIZE];
_declspec(align(32)) float c[SIZE];

int main() {
	// ランダムな値を生成
	random_device seed;
	mt19937 rnd(seed());
	for (size_t i = 0; i < SIZE; i++) {
		a[i] = (float)rnd() / rnd.max();
		b[i] = (float)rnd() / rnd.max();
	}

	auto start = chrono::system_clock::now();
	// 要素同士の積
	for (size_t i = 0; i < SIZE; i += 8) {
		*(__m256*)(c + i) = _mm256_mul_ps(*(__m256*)(a + i), *(__m256*)(b + i));
	}
	auto end = chrono::system_clock::now();
	auto msec = chrono::duration_cast<std::chrono::milliseconds>(end - start).count();

	cout << "duration = " << msec << " msec" << endl;

	for (size_t i = 0; i < 8; i++) {
		cout << c[i] << endl;
	}

	return 0;
}

GPU(CUDA)

#include <iostream>
#include <random>
#include <chrono>

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

using namespace std;

const size_t SIZE = 1024 * 1024 * 128;
_declspec(align(32)) float a[SIZE];
_declspec(align(32)) float b[SIZE];
_declspec(align(32)) float c[SIZE];

float *dev_a = 0;
float *dev_b = 0;
float *dev_c = 0;

__global__ void mulKernel(float *c, const float *a, const float *b)
{
	unsigned int i = threadIdx.x + blockIdx.x * blockDim.x;
	while (i < SIZE) {
		c[i] = a[i] * b[i];
		i += blockDim.x * gridDim.x;
	}
}

int main()
{
	// ランダムな値を生成
	random_device seed;
	mt19937 rnd(seed());
	for (size_t i = 0; i < SIZE; i++) {
		a[i] = (float)rnd() / rnd.max();
		b[i] = (float)rnd() / rnd.max();
	}

	cudaError_t cudaStatus;

	// Choose which GPU to run on, change this on a multi-GPU system.
	cudaStatus = cudaSetDevice(0);
	if (cudaStatus != cudaSuccess) {
		cerr << "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?" << endl;
		return 1;
	}
	
	// Allocate GPU buffers for three vectors (two input, one output)    .
	cudaStatus = cudaMalloc((void**)&dev_c, SIZE * sizeof(float));
	if (cudaStatus != cudaSuccess) {
		cerr << "cudaMalloc failed!" << endl;
		goto Error;
	}

	cudaStatus = cudaMalloc((void**)&dev_a, SIZE * sizeof(float));
	if (cudaStatus != cudaSuccess) {
		cerr << "cudaMalloc failed!" << endl;
		goto Error;
	}

	cudaStatus = cudaMalloc((void**)&dev_b, SIZE * sizeof(float));
	if (cudaStatus != cudaSuccess) {
		cerr << "cudaMalloc failed!" << endl;
		goto Error;
	}

	auto start1 = chrono::system_clock::now();

	// Copy input vectors from host memory to GPU buffers.
	cudaStatus = cudaMemcpy(dev_a, a, SIZE * sizeof(float), cudaMemcpyHostToDevice);
	if (cudaStatus != cudaSuccess) {
		cerr << "cudaMemcpy failed!" << endl;
		return cudaStatus;
	}

	cudaStatus = cudaMemcpy(dev_b, b, SIZE * sizeof(float), cudaMemcpyHostToDevice);
	if (cudaStatus != cudaSuccess) {
		cerr << "cudaMemcpy failed!" << endl;
		return cudaStatus;
	}

	auto start2 = chrono::system_clock::now();

	// 要素同士の積
	// Add vectors in parallel.
	// Launch a kernel on the GPU with one thread for each element.
	mulKernel <<<1024, 1024>>>(dev_c, dev_a, dev_b);

	// Check for any errors launching the kernel
	cudaStatus = cudaGetLastError();
	if (cudaStatus != cudaSuccess) {
		cerr << "addKernel launch failed: " << cudaGetErrorString(cudaStatus) << endl;
		goto Error;
	}

	// cudaDeviceSynchronize waits for the kernel to finish, and returns
	// any errors encountered during the launch.
	cudaStatus = cudaDeviceSynchronize();
	if (cudaStatus != cudaSuccess) {
		cerr << "cudaDeviceSynchronize returned error code " << cudaStatus << " after launching addKernel!" << endl;
		goto Error;
	}

	auto end1 = chrono::system_clock::now();

	// Copy output vector from GPU buffer to host memory.
	cudaStatus = cudaMemcpy(c, dev_c, SIZE * sizeof(float), cudaMemcpyDeviceToHost);
	if (cudaStatus != cudaSuccess) {
		cerr << "cudaMemcpy failed!" << endl;
		goto Error;
	}

	auto end2 = chrono::system_clock::now();
	auto msec = chrono::duration_cast<std::chrono::milliseconds>(end2 - start1).count();
	cout << "duration = " << msec << " msec" << endl;

	msec = chrono::duration_cast<std::chrono::milliseconds>(end1 - start2).count();
	cout << "duration(only gpu) = " << msec << " msec" << endl;

	for (size_t i = 0; i < 8; i++) {
		cout << c[i] << endl;
	}

Error:
	cudaFree(dev_c);
	cudaFree(dev_a);
	cudaFree(dev_b);
	
	// cudaDeviceReset must be called before exiting in order for profiling and
    // tracing tools such as Nsight and Visual Profiler to show complete traces.
    cudaStatus = cudaDeviceReset();
    if (cudaStatus != cudaSuccess) {
		cerr << "cudaDeviceReset failed!" << endl;
        return 1;
    }
	
	return 0;
}