GPUを使用した計算処理

こんにちは、キャスレーコンサルティング　SD(システム・デザイン)部の江本です。
今回は、CUDAによる計算処理について紹介したいと思います。

目次

• はじめに
• GPUの優位性
• 環境
• インストール
• チェック
• プロジェクト作成
• 計算してみよう
• 最後に

はじめに

皆さんは、GPU演算処理をご存知でしょうか？
人工知能にも使用されており、Deeplearning(ディープラーニング)でも有名になっていますよね。
他では、車の自動運転技術などにも使われている技術でもあります。
GPU演算処理の開発環境 CUDAを提供しているNVIDIA社のページでは、以下のように説明されています。

CUDA（クーダ）とはCompute Unified Device Architectureの略称で、半導体メーカーNVIDIA社が提供するGPUコンピューティング向けの統合開発環境です。
GPUコンピューティングおよびCUDAについて

本稿では、CUDAの開発環境セットアップを行い、サンプルソースを使用してCPUでの処理実行速度、GPUでの処理速度の比較を行いたいと思います。

GPUによる汎用計算の概要

昨今CPUは、複数コアを有するCPUが主流となってきています。
普段皆さんが使用しているPCのCPUですと、2コア～4コアくらいではないでしょうか？
しかし、GPUは、比較的安価な値段のGPUでもコアを200個以上積んでいる場合が多いです。
単純にコア数が多いからGPU＞CPUが成り立つわけではありませんが、特定の処理に関してはGPUはCPUに勝る処理性能を出すことができます。
GPU処理がCPU処理に勝るその特定の処理とは・・・並列処理です。
CUDAでは、ホスト（CPU/メインメモリ）からデバイス（GPU側）に命令を出し、実行することでGPUの特徴である多数のコアを同時に操ることができます。
簡単なGPUを使用した処理の概略図は、以下の通りとなります。

それでは、概要を説明したところで、実際にCUDA開発環境の構築から行っていきましょう。

環境

本稿は以下の環境で構築を行いました。

• Windows10 Pro 64bit
• Visual Studio 2013
• C++
• CUDA7.5
• Geforce750Ti

インストール

本稿では、前提条件として以下の2点の注意が必要です。

GPUが搭載されているPCを使用していること

VisualStudio2013がインストールされていること

上記前提条件をクリアされていることを確認して、インストーラからインストールを実施します。
インストーラーは、こちらからダウンロードしてください。
各環境に合わせたインストーラをダウンロードしてください。
環境で説明した通り、私の環境はWindows10であるため、以下のような設定でダウンロードを行いました。

2016/6/10現在では、最新バージョンは、cuda_7.5.18となっています。
ダウンロードしたインストーラを実行します。
以下のような順序で、インストールを実施します。

チェック

インストールが完了したら、次はGPUのスペックをチェックしましょう。
通常インストールを行うとWindows10では、以下のパスにソリューションファイルが作成されます。

“C:\ProgramData\NVIDIA Corporation\
CUDA Samples\v7.5\1_Utilities\deviceQuery\deviceQuery_vs2013.sln”

このソリューションファイルはいわゆるプロジェクトファイルとなっており、
このプロジェクトを開いて実行するとGPUのスペックチェックを行うことができます。
一度実行してみましょう。

先ほどインストールを行ったCUDAのバージョンが7.5であることが確認できます。
その他、GPUに関連したスペックも表示されており、
私のPCでは、CUDACoreが640個積んであることがわかります。
次の章では、プロジェクト作成とサンプルプロジェクトの説明を行います。

プロジェクト作成

それでは、CUDA用のプロジェクトで実際にプログラムを動かしてみましょう。
まずは、プロジェクトの作成からです。
VisualStudio2013を開いてプロジェクト作成してみましょう。

プロジェクト名は、ここでつけてください。
私は今回、入門ということでCudaStartedProjectと名づけてみました。

プロジェクト開くと以下のようなSampleソースが既に存在しています。

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

#include <stdio.h>

cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);

__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
}

int main()
{
    const int arraySize = 5;
    const int a[arraySize] = { 1, 2, 3, 4, 5 };
    const int b[arraySize] = { 10, 20, 30, 40, 50 };
    int c[arraySize] = { 0 };

    // Add vectors in parallel.
    cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addWithCuda failed!");
        return 1;
    }

    printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",
        c[0], c[1], c[2], c[3], c[4]);

    // 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) {
        fprintf(stderr, "cudaDeviceReset failed!");
        return 1;
    }

	getchar();

    return 0;
}

// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(int *c, const int *aint, const int *bint, unsigned int size)
{
    int *dev_a = 0;
    int *dev_b = 0;
    int *dev_c = 0;
    cudaError_t cudaStatus;

    // Choose which GPU to run on, change this on a multi-GPU system.
    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }

    // Allocate GPU buffers for three vectors (two input, one output)    .
    cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    // Copy input vectors from host memory to GPU buffers.
    cudaStatus = cudaMemcpy(dev_a, aint, size * sizeof(int), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

    cudaStatus = cudaMemcpy(dev_b, bint, size * sizeof(int), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

    // Launch a kernel on the GPU with one thread for each element.
    addKernel<<<1, size>>>(dev_c, dev_a, dev_b);

    // Check for any errors launching the kernel
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
        goto Error;
    }

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

    // Copy output vector from GPU buffer to host memory.
    cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

Error:
    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);

    return cudaStatus;
}

ここでは、配列a、bの要素をそれぞれ足し算して表示するようになっています。
ソース量は少しありますが、基本としては、以下の4段階となっています。

基本的なコードとしては、以下のような順序でコードを記載します。

１．cudaMalloc

デバイス側のメモリ領域確保です。
66行目から76行目にかけてホストからデバイスへデータ転送する前のメモリ領域を確保を行っています。

２．cudaMemcpy

ホスト、デバイス間のデータ転送を行います。
79行目、89行目でホストからデバイスへデータのコピーを行っています。
110行目では、デバイスから計算結果をホスト側で受け取っています。

３．__global__

デバイスでの計算式を記載します。
本稿では、8行目からの関数がデバイス側で実行される関数となります。
実行は、92行目で行っています。

４．cudaFree

cudaMallocで確保したデバイス側のメモリ領域を開放します。

それでは実行してみましょう。

計算してみよう

それでは実際に計算処理を作成して実行してみましょう。
今回は、1024行、1024列となる二つの行列を作成し、二つの行列の積を計算する処理を作って、
CPU処理、GPU処理それぞれの実行速度の差を確認します。
行列の初期数値は、ランダム関数を使用して適当に設定します。
（参考：行列の積）

それではまず、CPUの処理速度計測を行います。
C++でコーディングした以下の処理を実行します。

#include <stdio.h>
#include <stdlib.h>
#include <malloc.h>
#include <time.h>

/* 正方行列のサイズを定義 */
#define MATRIX_SIZE 1024

int main(int argc, char** argv){
	unsigned int col_idx, row_idx, scan_idx;
	int* matA;
	int* matB;
	int* matC;

	/* タイマー作成 */
	time_t Start, Stop;

	/* int型のn×n領域をメモリに確保 */
	matA = (int*)malloc(sizeof(int)* MATRIX_SIZE * MATRIX_SIZE);
	matB = (int*)malloc(sizeof(int)* MATRIX_SIZE * MATRIX_SIZE);
	matC = (int*)malloc(sizeof(int)* MATRIX_SIZE * MATRIX_SIZE);

	for (col_idx = 0; col_idx < MATRIX_SIZE; col_idx++) {
		for (row_idx = 0; row_idx < MATRIX_SIZE; row_idx++) {
			matA[col_idx * MATRIX_SIZE + row_idx] = rand() % (MATRIX_SIZE * MATRIX_SIZE);
			matB[col_idx * MATRIX_SIZE + row_idx] = rand() % (MATRIX_SIZE * MATRIX_SIZE);
			matC[col_idx * MATRIX_SIZE + row_idx] = 0;
		}
	}

	time(&Start);

	for (col_idx = 0; col_idx < MATRIX_SIZE; col_idx++) {
		for (row_idx = 0; row_idx < MATRIX_SIZE; row_idx++) {
			for (scan_idx = 0; scan_idx < MATRIX_SIZE; scan_idx++) {
				matC[col_idx * MATRIX_SIZE + row_idx] += matA[col_idx * MATRIX_SIZE + scan_idx] *
					matB[scan_idx * MATRIX_SIZE + row_idx];
			}
		}
	}

	time(&Stop);
	printf("Processing time: %d (sec)\n", Stop - Start);

	/* メモリを解放 */
	free(matA);
	free(matB);
	free(matC);

	getchar();

	return 0;
}

上記プログラムでの処理結果は、以下のようになりました。

CPU処理では、処理時間に6秒かかったことがわかります。
次は、CUDAを使って処理を書いてみましょう。

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

#define MATRIX_SIZE 1024/*行列１辺の数*/
#define BLOCK_SIZE 16

__global__ void matrixMul(int* inMatrixA, int* inMatrixB, int* inMatrixC);

int main(int argc, char** argv){
	unsigned int matrixSize = sizeof(unsigned int)* MATRIX_SIZE * MATRIX_SIZE;

	int* hMatrixA;
	int* hMatrixB;
	int* hMatrixC;
	hMatrixA = (int*)malloc(matrixSize);
	hMatrixB = (int*)malloc(matrixSize);

	/* 初期化 */
	unsigned int col_idx, row_idx;
	for (col_idx = 0; col_idx < MATRIX_SIZE; col_idx++){
		for (row_idx = 0; row_idx < MATRIX_SIZE; row_idx++){
			hMatrixA[col_idx * MATRIX_SIZE + row_idx] = rand() % (1024 * 1024);
			hMatrixB[col_idx * MATRIX_SIZE + row_idx] = rand() % (1024 * 1024);
		}
	}

	/* デバイス変数用ポインタ */
	int* dMatrixA;
	int* dMatrixB;
	int* dMatrixC;

	/* デバイス側のメモリ領域確保 */
	cudaMalloc((void**)&dMatrixA, matrixSize);
	cudaMemcpy(dMatrixA, hMatrixA, matrixSize, cudaMemcpyHostToDevice);
	cudaMalloc((void**)&dMatrixB, matrixSize);
	cudaMemcpy(dMatrixB, hMatrixB, matrixSize, cudaMemcpyHostToDevice);
	cudaMalloc((void**)&dMatrixC, matrixSize);

	/* GPUのブロックサイズとグリッドサイズの設定 */
	dim3 block(BLOCK_SIZE, BLOCK_SIZE);
	dim3 grid(MATRIX_SIZE / BLOCK_SIZE, MATRIX_SIZE / BLOCK_SIZE);

	/* 計測用タイマー */
	cudaEvent_t start;
	cudaEvent_t stop;
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
	/* タイマー開始 */
	cudaEventRecord(start, NULL);

	/* デバイス側の処理実行 */
	matrixMul <<<grid, block >>>(dMatrixA, dMatrixB, dMatrixC);
	cudaThreadSynchronize();

	/* 処理結果取得 */
	hMatrixC = (int*)malloc(matrixSize);
	cudaMemcpy(hMatrixC, dMatrixC, matrixSize, cudaMemcpyDeviceToHost);

	/* タイマーを停止 */
	cudaEventRecord(stop, NULL);
	cudaEventSynchronize(stop);

	/* 時間計測 */
	float msecTotal = 0.0f;
	cudaEventElapsedTime(&msecTotal, start, stop);

	printf("Processing time: %f (msec)\n", msecTotal);

	/*ホスト・デバイスメモリの開放*/
	free(hMatrixA);
	free(hMatrixB);
	free(hMatrixC);

	cudaFree(dMatrixA);
	cudaFree(dMatrixB);
	cudaFree(dMatrixC);

	/* 終了処理 */
	cudaThreadExit();

	getchar();

	exit(1);
}

/* 演算処理 */
__global__ void matrixMul(int* inMatrixA, int* inMatrixB, int* inMatrixC){
	unsigned int col_idx = blockIdx.x * blockDim.x + threadIdx.x;
	unsigned int row_idx = blockIdx.y * blockDim.y + threadIdx.y;
	unsigned int scan_idx;
	unsigned int target = 0;

	/* 行列の掛け算を行い、結果を詰める */
	for (scan_idx = 0; scan_idx < MATRIX_SIZE; scan_idx++) {
		target += inMatrixA[col_idx * MATRIX_SIZE + scan_idx] * inMatrixB[scan_idx * MATRIX_SIZE + row_idx];
		__syncthreads();
	}
	inMatrixC[col_idx * MATRIX_SIZE + row_idx] = target;
}

上記プログラム、即ちGPUでの処理結果は、以下のようになりました。

CPU処理が6秒、GPU処理が0.6秒という結果から、GPUでの処理はCPUでの処理の10倍の速度という事がわかります！！
GPU側は、今回データ転送も処理としてタイムに計測しています。

ただし、実処理で使用する場合、ホストとデバイスのメモリ転送が多いと、それがボトルネックと場合もあるので注意しましょう。

最後に

今回は入門ということで、ほぼサンプルの紹介レベルとなっています。
また次の機会があれば、少し発展した応用編を書きたいと思います。
それでは、また技術ブログでお会い出ればと思います。