Twitterに投稿 はてなブックマークに追加 Yahoo!ブックマークに追加 Google Bookmarksに追加

目次 >> CUDA >> 初めてのプログラム

CUDA - 初めてのプログラム

何もしないプログラム

まずは、CUDAの初期化、および終了のみをするプログラムを作ってみる。

cuda001.cuという空のテキストファイルを用意し、下記のような内容にする。

#include <stdio.h>
#include <cutil.h>
int main( int argc, char** argv) 
{
    CUT_DEVICE_INIT(argc, argv);

    unsigned int timer = 0;
    CUT_SAFE_CALL( cutCreateTimer( &timer));
    CUT_SAFE_CALL( cutStartTimer( timer));
    CUT_SAFE_CALL( cutStopTimer( timer));
    printf( "Processing time: %f (ms)\n", cutGetTimerValue( timer));
    CUT_SAFE_CALL( cutDeleteTimer( timer));

    CUT_EXIT(argc, argv);
    return 0;
}

CUT_DEVICE_INIT()とCUT_EXIT()だけでも良いのだが、その間のコードは単にタイマーを作成して、かかった時間を計っているだけである。下記のプログラムが完全な最小のプログラムである。(実際のところ、CUT_DEVICE_INIT()は対応したデバイスが存在するか確認しているだけで、CUT_EXIT()は「Press ENTER to exit..」という表示を出すだけであるので、これらがなくても、CUDAは動く。)

#include <stdio.h>
#include <cutil.h>
int main( int argc, char** argv) 
{
    CUT_DEVICE_INIT(argc, argv);
    CUT_EXIT(argc, argv);
    return 0;
}

これをコマンドプロンプトから、

nvcc cuda001.cu -lcutil32D -lkernel32

もしくは、リリース版を使う場合は

nvcc cuda001.cu -lcutil32 -lkernel32

実行ファイルを作成する。実行ファイルを実行し、

cuda001.exe
Processing time: 0.003911 (ms)

Press ENTER to exit...

のように出れば成功である。

GPUで計算をしてみる

cuda001では単に初期化しただけで何も計算をしなかったので、ここではGPUを使って簡単な計算をさせてみる。
まずcuda002.cuとcuda002_kernel.cuという2つの空のテキストファイルを用意する。それからサンプルプログラムsimpleTemplatesのsharedmem.cuhをコピーして同じフォルダに入れる。GPUの計算部分は慣例的に_kernelというファイル名にすることになっている。ここでは、0から99という数字の入った100個のデータを用意し、それをGPUを使って2倍にしてみる。
以下はcuda002.cuの内容である。

#include <stdio.h>
#include <cutil.h>

#include "cuda002_kernel.cu"
int main( int argc, char** argv) 
{
    //デバイスの初期化
    CUT_DEVICE_INIT(argc, argv);

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

    //メインメモリ上にfloat型のデータを100個生成する
    float* h_idata = (float*) malloc(sizeof( float) * 100);
    for( int i = 0; i < 100; i++) 
    {
        h_idata[i] = i;
    }

    //デバイス上(ビデオカードのこと)にも同じくfloat型100個分のメモリを確保する
    float* d_idata;
    CUDA_SAFE_CALL( cudaMalloc( (void**) &d_idata, sizeof( float) * 100 ));
    //メインメモリからデバイスのメモリにデータを転送する
    CUDA_SAFE_CALL( cudaMemcpy( d_idata, h_idata, sizeof( float) * 100 , cudaMemcpyHostToDevice) );

    //デバイス上に結果を格納するfloat型100個分のメモリを確保する
    float* d_odata;
    CUDA_SAFE_CALL( cudaMalloc( (void**) &d_odata, sizeof( float) * 100));

    dim3  grid( 1, 1, 1);
    //100は100個並列であるため
    dim3  threads(100, 1, 1);

    //ここでGPUを使った計算が行われる
    cuda002Kernel<<< grid, threads, sizeof( float) * 100 >>>( d_idata, d_odata);
        
    // 結果をコピーするための領域を確保する
    float* h_odata = (float*) malloc(sizeof( float) * 100);
    //デバイスからメインメモリ上に実行結果をコピー
    CUDA_SAFE_CALL( cudaMemcpy( h_odata, d_odata, sizeof( float) * 100, cudaMemcpyDeviceToHost) );

    //実行結果を表示
    printf("入力データ , 出力データ\n");
    for (int i=0;i<100;i++)
    {
        printf("%f , %f\n",h_idata[i],h_odata[i]);
    }
    //タイマーを停止しかかった時間を表示
    CUT_SAFE_CALL( cutStopTimer( timer));
    printf( "Processing time: %f (ms)\n", cutGetTimerValue( timer));
    CUT_SAFE_CALL( cutDeleteTimer( timer));

    //各種メモリを解放
    free( h_idata);
    free( h_odata);
    CUDA_SAFE_CALL(cudaFree(d_idata));
    CUDA_SAFE_CALL(cudaFree(d_odata));

    //終了処理
    CUT_EXIT(argc, argv);
}

基本的な構造は、まずメインメモリ上に計算させたいデータを用意する。それをGPUに転送して、計算し、それをメインメモリ上に書き戻すという構造になる。
次に、GPUの計算部分であるcuda002_kernel.cuの内容は次の通り。

#include <stdio.h>
#include "sharedmem.cuh"

__global__ void cuda002Kernel( float* g_idata, float* g_odata) 
{
    SharedMemory<float> smem;
    float* sdata = smem.getPointer();

    // スレッドIDを取得
    const unsigned int tid = threadIdx.x;

    //グローバルメモリから入力データの読み込み
    sdata[tid] = g_idata[tid];
    __syncthreads();

    //ここで計算を行う
    sdata[tid] = (float) 2 * sdata[tid];
    __syncthreads();

    //グローバルメモリに結果を書き込む
    g_odata[tid] = sdata[tid];
}

コンパイルと実行結果は、

>nvcc -ccbin "C:\Program Files\Microsoft Visual Studio 8\VC\bin" -c -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/O2,/Zi,/MT cuda002.obj cuda002.cu
cuda002.cu
tmpxft_00000f64_00000000-3.gpu
tmpxft_00000f64_00000000-7.gpu
tmpxft_00000f64_00000000-3.c
tmpxft_00000f64_00000000-12.i

>link /OUT:cuda002.exe cuda002.obj cudart.lib cutil32D.lib kernel32.lib
Microsoft (R) Incremental Linker Version 8.00.50727.42
Copyright (C) Microsoft Corporation. All rights reserved.

>cuda002.exe
入力データ , 出力データ
0.000000 , 0.000000
1.000000 , 2.000000
2.000000 , 4.000000
3.000000 , 6.000000
4.000000 , 8.000000
.
.
99.000000 , 198.000000
Processing time: 25.315508 (ms)

Press ENTER to exit...


最終更新日


本文中のFC4はFedora ProjectのFedora Core 4を、FC5はFedora Core 5を、FC6はFedora Core 6をopenSUSEはNovellのSUSE Linux OSSを表します。Fedora7以降は、単にFedora7、Fedora8、Fedora9、Fedora10、Fedora11、Fedora12、Fedora13、Fedora14、Fedora15と表示しています。Ubuntuは、必要に応じて17.04、18.04のようにバージョン番号をつけて区別しています。

ここに登場するドメイン名やIPアドレスなどはフィクションです。実在の人物・団体等とは一切関係がありません。
実際に使用する際は、各自の環境に合わせて書き換えてください。
もし何か間違いなどありましたらこちらからご連絡ください
リンクに許可は不要です。
Copyright (C) 2014 Chikuma Engineering Co., Ltd. All Rights Reserved.