こんにちは、キャスレーコンサルティング 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段階となっています。
基本的なコードとしては、以下のような順序でコードを記載します。
1.cudaMalloc
デバイス側のメモリ領域確保です。
66行目から76行目にかけてホストからデバイスへデータ転送する前のメモリ領域を確保を行っています。
2.cudaMemcpy
ホスト、デバイス間のデータ転送を行います。
79行目、89行目でホストからデバイスへデータのコピーを行っています。
110行目では、デバイスから計算結果をホスト側で受け取っています。
3.__global__
デバイスでの計算式を記載します。
本稿では、8行目からの関数がデバイス側で実行される関数となります。
実行は、92行目で行っています。
4.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/*行列1辺の数*/ #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側は、今回データ転送も処理としてタイムに計測しています。
ただし、実処理で使用する場合、ホストとデバイスのメモリ転送が多いと、それがボトルネックと場合もあるので注意しましょう。
最後に
今回は入門ということで、ほぼサンプルの紹介レベルとなっています。
また次の機会があれば、少し発展した応用編を書きたいと思います。
それでは、また技術ブログでお会い出ればと思います。