• トップ
  • G-DEPについて
  • ご購入ガイド
  • サポート
  • お問い合わせ

G-DEPトップ  >  第4回 実際にCUDAを使ってみる

第4回 実際にCUDAを使ってみる

第4回 「実際にCUDAを使ってみる」 

<< 第3回   |   目次に戻る   |   第5回 >>

前回のコラムはいかがでしたか?無事インストールすることできましたでしょうか。

これで、ようやくCUDAを使ってGPUコンピューティングを行う準備ができました。
今回はCUDAをより知ってもらうために簡単なプログラムを書いて、実際に動かしてみたいと思います。
GPUを使いこなすためには、GPUアーキテクチャやCUDA言語についてより深く学ぶ必要がありますが、インストールしたからには少し動かしてみたいですよね!
 
ここではC言語の基礎知識があることが前提で書いていきますが、経験が無い方にもなるべくポイントとなる部分を分かりやすく説明していきます。
ソースをコピペして、コンパイル&実行するだけで構いません。とりあえずやってみましょう!
 
と、問題に入る前に、意外と詰まりやすいのが実際にコーディングするときとコンパイルする方法です。WindowsでVisual Studioを使う場合と、Linux(CentOS5.6)で直接コマンドプロンプトでコンパイルする方法について説明します。
 
 
 準備①Widows編

Visual Studio 2008 Expressの設定

①「C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\src」を開きます。前回GPU Computing SDKブラウザで動かしたサンプルプログラムが格納されています。

②その中の「template」というフォルダがあるので、同じディレクトリにコピーし、名前を「practice」に変更します。この名前は特に何でも構いません。利便性のためCUDAプログラムのテンプレートを利用しています。

③名前を変更したフォルダを開き、「template_vs2008.sln」というファイルを開くと以下のようなウィンドウが出ます。ちなみに左のメニューは「ソリューション>プロジェクト>ソースフォルダ>コードファイル」という階層になっています。

(クリックで拡大)

※ここでエラーが出て開けない場合は、以下の事を行ってください。
C:\Program Files (x86)\NVIDIA GPU Computing Toolkit\CUDA\v4.0\extras\visual_studio_integration\rules」にある4つのファイル(NvCuda~)を全てコピーし、「C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\VCProjectDefaults」に貼り付けてください。
 
④「メニューのツール>オプション>プロジェクトおよびソリューション>VC++ディレクトリ」と順に選択して表示します。ここで3つのパスを設定しておきます。
・「実行可能ファイル」に以下の文を追加(フォルダマークを押すと文を挿入できます。)
$(CUDA_BIN_PATH)C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\bin
・「インクルードファイル」に以下の文を追加
$(CUDA_INC_PATH)C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\inc
・「ライブラリファイル」に以下の文を追加
$(CUDA_LIB_PATH)C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\lib
 
(クリックで拡大)
 
⑤次に、srcフォルダの「template.cu」以外の2つのファイルを削除し、template.cuをダブルクリックして展開します。(適宜、ファイル・プロジェクト名を変更してください。)
 
⑥表示されたソースコードを全て削除し、ここに自分の行いたいプログラムを書き込みます。
 
 
コンパイル&実行
⑦プログラムを書き込んだら、メニューの「ビルド>ソリューションのビルド(F7でも可)」を選択し、ビルドを行います。
 
⑧ビルドが成功したら、メニューの「デバッグ>デバッグなしで開始(Ctrl+F5でも可)」でプログラムが開始されます。
 
備考:キーワード色付けの設定
①「C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\doc\syntax_highlighting\visual_studio_8」の中にある「usertype.dat」を、「C:\Program Files (x86)\Microsoft Visual Studio 9.0\Common7\IDE」にコピー&ペーストします。
 
②Visual Studioを起動し、「ツール>オプション>テキスト エディタ>ファイル拡張子」を選択し、エディタを「Microsoft Visual C++」として拡張子に「cucuh」を追加します。
 
③Visual Studioを再起動し、cuファイルを開いてキーワードが色分けされていれば完了です。

 

 準備②Linux編

①コマンエディタを開きます。

②コンパイルしたいファイルのあるディレクトリに移動します。例:「$ cd Cuda_practice」

③コマンド「nvcc (コンパイルするファイル名).cu」でコンパイルできます。例:「$ nvcc test.cu」
また、実行ファイル名を変えたいときは「$ nvcc -o (実行ファイル名).exe (コンパイルするファイル名).cu」で行います。例:「$ nvcc -o test.exe test.cu」

④できた実行ファイルを起動すると、サンプルが実行されます。例:「$ ./a.exe」 or 「$ ./test.exe」(実行ファイル名を指定して作った場合)

※この方法はWindowsおよびMac X OSのコマンドプロンプトでも同様にコンパイル&実行することができます。

 

 

 

それでは、準備が整ったところで実際にプログラミングしてみましょう!

 Q1. Hello, World !

プログラミング言語の第一問と言えば、コレ!というくらい有名な問題ですね。
内容はというと、単に「Hello, World !」と表示させるだけです。
 
ではまず、下のソースを「hello_world.cu」というファイル名で保存してください。
 
#include <stdio.h>

int main( void ) {
  printf("Hello, World!\n");
  return 0;
}
 
え、これただのC言語じゃないか。とC言語経験者なら思うでしょう。
確かにその通りですが、とりあえずこれを上記の方法でコンパイルしてみてください。
どうでしょう?下のように「Hello, World!」と表示されましたか?
 
ここで大事なのは、拡張子が「cu」であり、「nvcc」というCUDA用のコマンドを使用してコンパイルできた、ということです。nvccコンパイラというのは、ソースの中のCPUで動く部分とGPUで動く部分を分け、CPU用コードをCコンパイラに渡し、GPU用コードをコンパイル(正確にはPTXコードに変換)するものであることを覚えておいてください。
そのため、ソースがC言語のみで書かれていてもコンパイルできますし、その場合実際にコンパイルしているのはCコンパイラ、というわけです。
 
 
では、GPU用のコードを書いて同じことができるのか?実際にやってみましょう。
以下のソースを今度は「hello_gpu.cu」などと名付けて、同じ手順でコンパイル&実行してみてください。
 
#include <stdio.h>

__global__ void kernel( void ) {
}

int main( void ) {
    kernel<<<1,1>>>();
    printf( "Hello, GPU World!\n" );
    return 0;
}
 
 
おや?上と少し違う点がありますね。
・__global __から始まるkernel( void)という名前の空の関数
・<<<1,1>>>で修飾された空の関数の呼び出し
という、C言語では見慣れないコードが出てきました。
 
__global__(←アンダーバー2つ)は、簡単に言えばコンパイラに「以下の関数をGPU側で実行します」と指定するコマンドです。kernel()関数はGPU用コードを処理するコンパイラに渡され、main()関数はCコンパイラに渡されます。
ですが、やっていることは上とほとんど同じです。
 
CUDA言語はC言語と似ており、同じようにプログラムを書くことができることが理解してもらえましたでしょうか。
 
 
 Q2. 99 bottles of beer

こちらはHello, World!より若干マイナーですが、同様によく使われる問題。
99本のビールを一本ずつ減らして、数を表示させていくという内容です。
細かい説明は抜きにして、CPU用とGPU用のソースを載せておきます。
GPU用ソースに関しては「CUDA Information」さんのページを参考にさせて頂きました。(ページ下に参照URL)
 
「99bottles_cpu」
#include <stdio.h>

int main(void)
{       
        int b;
      
        for (b = 99; b >= 0; b--) {
                switch (b) {
                case 0:
                        printf("No more bottles of beer on the wall, no more bottles of beer.\n");
                        printf("Go to the store and buy some more, 99 bottles of beer on the wall.\n");
                        break;
                case 1:
                        printf("1 bottle of beer on the wall, 1 bottle of beer.\n");
                        printf("Take one down and pass it around, no more bottles of beer on the wall\n");
                        break;
                default:
                        printf("%d bottles of beer on the wall, %d bottles of beer.\n", b, b);
                        printf("Take one down and pass it around, %d %s of beer on the wall.\n"
                                ,b - 1
                                ,((b - 1) > 1)? "bottles" : "bottle");
                        break;
                }
        }               
        return 0;
}      
 
 
「99bottles_gpu」
#include <stdio.h>

#define SIZE_TEXT (sizeof(text)-1)
#define SIZE_END (sizeof(end)-1)

__device__ char text[] =
"__ bottles of beer on the wall, __ bottles of beer!\n"
"Take one down, and pass it around, ## bottles of beer on the wall!\n\n";
       
__device__ char end[] =
"01 bottle of beer on the wall, 01 bottle of beer.\n"
"Take one down and pass it around, no more bottles of beer on the wall.\n"
"\n"
"No more bottles of beer on the wall, no more bottles of beer.\n"
"Go to the store and buy some more, 99 bottles of beer on the wall.";


__global__
void bottle99(char *addr){
    int x = threadIdx.x;
       addr += x * SIZE_TEXT;
    int bottle = 99 - x;
       if (bottle == 1) {
        for (int i=0; i<SIZE_END; i++) {
               addr[i] = end[i];
           }
         addr[SIZE_END] = '\0';
           } else {
               char c1 = (bottle/10) + '0';
               char c2 = (bottle%10) + '0';

               char d1 = ((bottle-1)/10) + '0';
               char d2 = ((bottle-1)%10) + '0';
      
           for (int i=0; i<SIZE_TEXT; i++) {
               int c = text[i];
               if (c == '_') {
                  addr[i] = c1;
                  addr[i+1] = c2;
                  i++;
               } else if (c == '#') {

            addr[i] = d1;
               addr[i+1] = d2;
                    i++;
               } else {

                 addr[i] = text[i];
            }
        }
    }
}
      
int main()
{
    char *buffer;
    char *d_buffer;
      
    int size = SIZE_TEXT * 98 + SIZE_END + 1;

    buffer = new char[size];
    cudaMalloc((void**)&d_buffer, size);

    bottle99<<<1, 99>>>(d_buffer);

    cudaMemcpy(buffer, d_buffer, size, cudaMemcpyDeviceToHost);
    cudaFree(d_buffer);

    puts(buffer);
    free(buffer);               
       }
  
 
体感的な表示速度は変わらないかと思いますが、GPU用ソースでは並列計算部分をGPUで処理しているため、CPUで計算するよりも高速で演算しています。
 
んーせっかくCPUと、GPUの速さの違いが分かると思ったのに…
と思った方へ、次のプログラムです。
 
 
 Q3. Matrix
CPUとGPUの演算能力の差が目に見えて分かるのが行列演算です。
実際にそれを体感して頂くため、以下のようなプログラムを用意しました。
 
ここではn×nの正方行列、A、B、Cを用意し、A、Bにランダムの初期値を与えた後
C = A * B
という演算を行っており、行列の1辺などの各条件は同じです。
ではまずCPUのみで演算するプログラムから見てみましょう。
 
「Matrix_CPU」
#include <stdio.h>
#include <stdlib.h>
#include <malloc.h>
#include <time.h>

/*n正方行列のサイズを定義*/
#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);

    return 0;
}

これをコンパイル&実行してみてください。マシンスペックによっては数秒~1分ほどかかると思われます。
 
ちなみに、筆者のPCで行うと9秒でした。この値を覚えていてください。
 
 
そしてGPU演算プログラムがこちらです。
「Matrix_GPU」
#include <stdio.h>
#include <malloc.h>
#include <stdlib.h>
#include <time.h>
#include <cutil_inline.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;
 
/*デバイスメモリ領域の確保*/
  cutilSafeCall(cudaMalloc((void**)&dMatrixA, matrixSize));
  cutilSafeCall(cudaMemcpy(dMatrixA, hMatrixA, matrixSize, cudaMemcpyHostToDevice));
  cutilSafeCall(cudaMalloc((void**)&dMatrixB, matrixSize));
  cutilSafeCall(cudaMemcpy(dMatrixB, hMatrixB, matrixSize, cudaMemcpyHostToDevice));
  cutilSafeCall(cudaMalloc((void**)&dMatrixC, matrixSize));

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

/*タイマーを作成して計測開始*/
  unsigned int timer = 0;
  CUT_SAFE_CALL( cutCreateTimer( &timer));
  CUT_SAFE_CALL( cutStartTimer( timer));

/*カーネルの起動*/
  matrixMul<<<grid, block>>>(dMatrixA, dMatrixB, dMatrixC);
  cudaThreadSynchronize();

/*結果の領域確保とデバイス側からのメモリ転送*/
  hMatrixC = (int*)malloc(matrixSize);
  cutilSafeCall(cudaMemcpy(hMatrixC, dMatrixC, matrixSize, cudaMemcpyDeviceToHost));

/*タイマーを停止しかかった時間を表示*/
  CUT_SAFE_CALL( cutStopTimer( timer));
  printf("Processing time: %f (msec)\n", cutGetTimerValue( timer));
  CUT_SAFE_CALL( cutDeleteTimer( timer));

/*ホスト・デバイスメモリの開放*/
  free(hMatrixA);
  free(hMatrixB);
  free(hMatrixC);
  cutilSafeCall(cudaFree(dMatrixA));
  cutilSafeCall(cudaFree(dMatrixB));
  cutilSafeCall(cudaFree(dMatrixC));
 
/*終了処理*/
  cudaThreadExit();
  cutilExit(argc, argv);
 }
 
__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;

}
 
結果がコチラ↓
 
99.48… (msec) = 0.099秒
先ほど、CPUで演算させた場合は9秒かかっていたものが、今度は0.1秒かかっていません。
 
つまりナント、GPUで演算させると約100倍も高速化できていることになります。
 
GPUの並列処理能力はすごいですね!
もちろんGPUコード化すれば全てのプログラムが100倍も速くなるというわけではないですが、それでも計算内容によってはCPUとは比べ物にならない処理速度を発揮することができます。
 
ちなみにこの行列演算プログラムは「CUDA 高速GPUプログラミング入門(発行:秀和システム)」の例題を参考に、一部改編して作成したものを載せています。良ければこちらの本もぜひ参考にしてください。
 
 
 

いかがでしたでしょうか?コードが並ぶばかりで少し見づらかったかもしれません。

今回は、一先ずコードの勉強は置いて、実際にCUDAのコンパイル&実行の仕方を覚えてもらうための回でした。

また、これで改めてGPUコンピューティングに興味が湧い頂けたら何よりです。

 

次回はまた戻ってGPUの構造のお話になります。

「GPUの構造を知らずしてCUDAはマスターできず」

ということで、一緒に勉強していきましょう!

 

<< 第3回   |   目次に戻る   |   第5回 >>


【参考ページ】(新しいウィンドウで開きます)

簡単な並列プログラムを書いてみよう(CUDA Information)

Visual StudioでのCUDAの設定(筑波大学 藤澤誠 Physics-Based Computer Graphics Lab. -Wiki)

(執筆 G-DEP Associate Research Engineer 東京大学大学院工学系研究科 岡安優)