Visual Studio 2010のプロジェクトにCUDAコードを追加する手順のまとめ.
CUDAのバージョンは4.0とする.



VS2010+CUDA4.0

CUDAインストール

  • CUDAツールキットをインストールする.
    32ビットの場合は,
    cudatoolkit_4.0.*_win_32.msi
    cudatools_4.0.*_win_32.msi
    64ビットの場合は,
    cudatoolkit_4.0.*_win_64.msi
    cudatools_4.0.*_win_64.msi
    注意として,64ビットOS上でも32ビットアプリケーション開発している場合は32ビットCUDAをインストールする. 両方開発している場合は両方インストールすることもできる. また,64ビットCUDAで32ビットアプリケーションを開発する方法は 「x64版CUDA上で32bitアプリケーションのビルド」を参照.
  • GPU Computing SDKのインストールする. 32ビットの場合は,
    gpucomputingsdk_4.0.*_win_32.exe
    64ビットの場合は,
    gpucomputingsdk_4.0.*_win_64.exe
    インストール後,環境変数を適用するために一度ロブオフしておく. ちなみに設定される環境変数(32ビットの場合)は,
    CUDA_PATH=C:\Program Files (x86)\NVIDIA GPU Computing Toolkit\CUDA\v4.0\
    CUDA_PATH_V4_0=C:\Program Files (x86)\NVIDIA GPU Computing Toolkit\CUDA\v4.0\
    
    CUDA_BIN_PATH=%CUDA_PATH%\bin
    CUDA_INC_PATH=%CUDA_PATH%\include
    CUDA_LIB_PATH=%CUDA_PATH%\lib\Win32
  • ルールファイルをコピー
    C:\Program Files (x86)\NVIDIA GPU Computing Toolkit\CUDA\v4.0\extras\visual_studio_integration\rules\
    にある
    NvCudaDriverApi.v4.0
    NvCudaRuntimeApi.v4.0
    C:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\VCProjectDefaults
    にコピーする.
  • cutilのビルド
    C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\cutil_vs2010.sln
    をVS2010で開いてビルドする(必要に応じてランタイムライブラリを変更する). 32ビットの場合,
    C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\lib\Win32\
    にcutil32.lib, cutil32.dllができるので,VS2010から見られるライブラリディレクトリに移動する or 上記ディレクトリを設定する.

VS2010プロジェクトの変更

  • ビルドターゲットの設定
    プロジェクトを右クリックして,「ビルドのカスタマイズ」を選択. CUDA 4.0(.targets, .props)にチェックを入れる.
    vs2010_cuda4_dlg1.jpg
  • vcxprojファイルの修正
    ビルドのカスタマイズでCUDAを追加した状態でビルドすると
    エラー: 要素 <UsingTask> 内の属性 "AssemblyFile" の値 "$(CudaBuildTasksPath)" を評価した結果 "" は無効です。 
      C:\Program Files (x86)\MSBuild\Microsoft.Cpp\v4.0\BuildCustomizations\CUDA 4.0.targets
    といったエラーが出る.
    C:\Program Files (x86)\MSBuild\Microsoft.Cpp\v4.0\BuildCustomizations\
    にある"CUDA 4.0.targets"ファイルを見ると 12行目でCudaBuildRulesPathが参照されているが事前に定義されていないようである. この定義は"CUDA 4.0.targets"と同じフォルダにある"CUDA 4.0.props"に書かれているので, これを参照するようにvcxprojファイルを書き換える.
    vcxprojファイルをテキストエディタで開いて,
     <Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
     <ImportGroup Label="ExtensionSettings">
     </ImportGroup>
    の部分を
     <Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
     <ImportGroup Label="ExtensionSettings">
       <Import Project="$(VCTargetsPath)\BuildCustomizations\CUDA 4.0.props" />
     </ImportGroup>
    のように書き換えることで "CUDA 4.0.porps" を追加する
  • プロジェクトの設定 インクルードフォルダ
    $(CUDA_INC_PATH)
    ライブラリフォルダ
    $(CUDA_PATH)/lib/$(PlatformName)
    ライブラリファイル
    cudart.lib cutil32.lib
    を追加.

CUファイルの追加

  1. CUファイルの追加 .cuや.cuhファイルをプロジェクトに追加する.
  2. CUDA C/C++の設定 CUファイルのプロパティから全般->項目の種類をCUDA C/C++にして, CUDA C/C++の各項目を設定する.
    vs2010_cuda4_dlg2.jpg

私の設定項目は,

  • Debug,Release共通
    • Common -> Source Dependencies : *.cuhファイルを追加
    • Common -> Additional Include Directories : [SDKインストールフォルダ]\inc
    • Device -> Code Generation : 必要に応じて compute_13,sm_13 や compute_20,sm_20 にする
  • Debug
    • Host -> Runtime Libray : Multi-Threaded Debug (/MTd) -> ホストのランタイムライブラリ設定に合わせる
  • Release
    • Host -> Use Fast Math : はい (-use_fast_math)
    • Host -> Optimization : Full Optimization (/Ox)
    • Host -> Runtime Libray : Multi-Threaded (/MT) -> ホストのランタイムライブラリ設定に合わせる
      vs2010_cuda4_dlg3.jpg
      vs2010_cuda4_dlg5.jpg
      vs2010_cuda4_dlg7.jpg
      cuhファイルやデバイス関数(__device__,__global__)のみが書かれたcuファイルに関しては ビルドから除外を「はい」にする.
      vs2010_cuda4_dlg6.jpg

CUDAファイルの構成例1(カーネル関数を分ける)

ファイル構成

  • rx_*_kernel.cuh : カーネルヘッダ,定義などを記述.rx_*_kernel.cu,rx_*.cuh,rx_*.cu,CUDA呼び出し側からインクルードする.
  • rx_*_kernel.cu : カーネル関数の実装.rx_*.cuからインクルードする.
  • rx_*.cuh : ホスト関数のヘッダ.CUDA呼び出し側からインクルードする.
  • rx_*.cu : ホスト関数の実装.CUDA呼び出し側(VCコンパイラを使用するコード)からインクルードしないこと.

コード例

  • rx_*_kernel.cuh
      1
      2
      3
      4
      5
      6
      7
      8
      9
     10
     11
     12
     13
     14
     15
     16
     17
     18
     19
     20
     21
     22
     23
     24
     25
     26
    
    #ifndef _RX_TEST_KERNEL_CUH_
    #define _RX_TEST_KERNEL_CUH_
     
    //-----------------------------------------------------------------------------
    // インクルードファイル
    //-----------------------------------------------------------------------------
    #include "vector_types.h"
     
    //-----------------------------------------------------------------------------
    // 定義
    //-----------------------------------------------------------------------------
    typedef unsigned int uint;
    typedef unsigned char uchar;
     
     
    // 1ブロックあたりのスレッド数(/次元)
    #define BLOCK_SIZE 16
     
    // 1ブロックあたりの最大スレッド数
    #define THREAD_NUM 256
     
    // 円周率
    #define M_PI 3.141592653589793238462643383279502884197
    #define M_2PI 6.283185307179586476925286766559005768394        // 2*PI
     
    #endif // _RX_TEST_KERNEL_CUH_
    
  • rx_*_kernel.cu
      1
      2
      3
      4
      5
      6
      7
      8
      9
     10
     11
     12
     13
     14
     15
     16
     17
     18
     19
     20
     21
     22
     23
     24
     25
     26
     27
     28
     29
     30
     31
     32
     33
     34
     35
     36
     37
     38
     39
     40
     41
     42
     43
     44
     45
    
    #ifndef _RX_TEST_KERNEL_CU_
    #define _RX_TEST_KERNEL_CU_
     
    //-----------------------------------------------------------------------------
    // インクルードファイル
    //-----------------------------------------------------------------------------
    #include <stdio.h>
    #include <math.h>
     
    #include "cutil_math.h"
    #include "math_constants.h"
     
    #include "rx_test_kernel.cuh"
     
    //-----------------------------------------------------------------------------
    // デバイス関数
    //-----------------------------------------------------------------------------
    /*!
     * [a,b]にクランプ
     * @param[in] x クランプしたい数値
     * @param[in] a,b クランプ境界
     * @return クランプされた数値
     */
    __device__
    float CuClamp(float x, float a, float b)
    {
        return max(a, min(b, x));
    }
     
    //-----------------------------------------------------------------------------
    // カーネル関数
    //-----------------------------------------------------------------------------
    /*!
     * 平方根を計算
     */
    __global__
    void calSquareRoot(float* s, unsigned long n)
    {
        unsigned long idx = __mul24(blockIdx.x, blockDim.x)+threadIdx.x;
        if(idx >= n) return;
     
        s[idx] = sqrt((float)idx);
    }
     
    #endif // #ifndef _RX_TEST_KERNEL_CU_
    
  • rx_*.cuh
      1
      2
      3
      4
      5
      6
      7
      8
      9
     10
     11
     12
     13
     14
     15
     16
     17
     18
     19
     20
     21
     22
    
    #ifndef _RX_TEST_CUH_
    #define _RX_TEST_CUH_
     
    //-----------------------------------------------------------------------------
    // インクルードファイル
    //-----------------------------------------------------------------------------
    #include "rx_test_kernel.cuh"
     
    //-----------------------------------------------------------------------------
    // CUDA関数
    //-----------------------------------------------------------------------------
    extern "C"
    {
    void CuSetDevice(int argc, char **argv);
    void CuSetDeviceByID(int id);
     
    void CuDeviceProp(void);
     
    void CuSquareRoot(float *s, unsigned long n);
    }    // extern "C"
     
    #endif // #ifdef _RX_TEST_CUH_
    
  • rx_*.cu
      1
      2
      3
      4
      5
      6
      7
      8
      9
     10
     11
     12
     13
     14
     15
     16
     17
     18
     19
     20
     21
     22
     23
     24
     25
     26
     27
     28
     29
     30
     31
     32
     33
     34
     35
     36
     37
     38
     39
     40
     41
     42
     43
     44
     45
     46
     47
     48
     49
     50
     51
     52
     53
     54
     55
     56
     57
     58
     59
     60
     61
     62
     63
     64
     65
     66
     67
     68
     69
     70
     71
     72
     73
     74
     75
     76
     77
     78
     79
     80
     81
     82
     83
     84
     85
     86
     87
     88
     89
     90
     91
     92
     93
     94
     95
     96
     97
     98
     99
    100
    101
    102
    103
    104
    105
    106
    107
    108
    109
    110
    111
    112
    113
    114
    115
    116
    117
    118
    119
    120
    121
    122
    123
    124
    125
    126
    127
    128
    129
    130
    131
    132
    133
    134
    
    //-----------------------------------------------------------------------------
    // インクルードファイル
    //-----------------------------------------------------------------------------
    #include <cstdio>
     
    #include <cutil_inline.h>
    #include <cutil.h>
     
    #include "rx_test_kernel.cu"
     
     
    //-----------------------------------------------------------------------------
    // CUDA関数
    //-----------------------------------------------------------------------------
    extern "C"
    {
    /*!
     * 実行時引数を用いたデバイスの設定
     * @param[in] argc コマンドライン引数の数
     * @param[in] argv コマンドライン引数
     */
    void CuSetDevice(int argc, char **argv)
    {   
        if(cutCheckCmdLineFlag(argc, (const char**)argv, "device")){
            cutilDeviceInit(argc, argv);
        }
        else{
            cudaSetDevice( cutGetMaxGflopsDeviceId() );
        }
    }
     
    /*!
     * IDを指定してデバイスを設定
     * @param[in] id デバイスID
     */
    void CuSetDeviceByID(int id)
    { 
        int device_count = 0;
        cudaGetDeviceCount(&device_count);
        if(id < 0 || id >= device_count){
            id = cutGetMaxGflopsDeviceId();
            cudaSetDevice(id);
        }
        else{
            cudaSetDevice(id);
        }
    }
     
    /*!
     * すべてのデバイスプロパティを表示
     */
    void CuDeviceProp(void)
    {
        int n;    //デバイス数
        cutilSafeCall(cudaGetDeviceCount(&n));
     
        for(int i = 0; i < n; ++i){
            cudaDeviceProp dev;
     
            // デバイスプロパティ取得
            cutilSafeCall(cudaGetDeviceProperties(&dev, i));
     
            printf("device %d\n", i);
            printf(" device name : %s\n", dev.name);
            printf(" total global memory : %d (MB)\n", dev.totalGlobalMem/1024/1024);
            printf(" shared memory / block : %d (KB)\n", dev.sharedMemPerBlock/1024);
            printf(" register / block : %d\n", dev.regsPerBlock);
            printf(" warp size : %d\n", dev.warpSize);
            printf(" max pitch : %d (B)\n", dev.memPitch);
            printf(" max threads / block : %d\n", dev.maxThreadsPerBlock);
            printf(" max size of each dim. of block : (%d, %d, %d)\n", dev.maxThreadsDim[0], dev.maxThreadsDim[1], dev.maxThreadsDim[2]);
            printf(" max size of each dim. of grid  : (%d, %d, %d)\n", dev.maxGridSize[0], dev.maxGridSize[1], dev.maxGridSize[2]);
            printf(" clock rate : %d (MHz)\n", dev.clockRate/1000);
            printf(" total constant memory : %d (KB)\n", dev.totalConstMem/1024);
            printf(" compute capability : %d.%d\n", dev.major, dev.minor);
            printf(" alignment requirement for texture : %d\n", dev.textureAlignment);
            printf(" device overlap : %s\n", (dev.deviceOverlap ? "ok" : "not"));
            printf(" num. of multiprocessors : %d\n", dev.multiProcessorCount);
            printf(" kernel execution timeout : %s\n", (dev.kernelExecTimeoutEnabled ? "on" : "off"));
            printf(" integrated : %s\n", (dev.integrated ? "on" : "off"));
            printf(" host memory mapping : %s\n", (dev.canMapHostMemory ? "on" : "off"));
     
            printf(" compute mode : ");
            if(dev.computeMode == cudaComputeModeDefault) printf("default mode (multiple threads can use) \n");
            else if(dev.computeMode == cudaComputeModeExclusive) printf("exclusive mode (only one thread will be able to use)\n");
            else if(dev.computeMode == cudaComputeModeProhibited) printf("prohibited mode (no threads can use)\n");
            
        }
     
        printf("Device with Maximum GFLOPS : %d\n", cutGetMaxGflopsDeviceId());
    }
     
    /*!
     * a/bを計算して結果を切り上げる
     * @param[in] a,b 分子,分母
     * @return ceil(a/b)
     */
    uint Ceil(uint a, uint b)
    {
        return (a%b != 0) ? (a/b+1) : (a/b);
    }
     
    /*!
     * すべてのデバイスプロパティを表示
     */
    void CuSquareRoot(float *hS, unsigned long n)
    {
        float *dS = 0;
        cutilSafeCall(cudaMalloc((void**)&dS, n*sizeof(float)));
     
        // 1スレッド/計算
        uint block;    // ブロック内スレッド数
        dim3 grid;    // グリッド内ブロック数
        block = THREAD_NUM;
        grid = dim3(Ceil(n, block), 1, 1);
        if(grid.x > 65535){
            grid.y = grid.x/32768;
            grid.x = 32768;
        }
     
        printf("grid : (%d, %d, %d), block : %d\n", grid.x, grid.y, grid.z, block);
     
        // カーネル実行
        calSquareRoot<<< grid, block >>>(dS, n);
     
        cutilCheckMsg("calSquareRoot kernel execution failed");    // カーネル実行エラーチェック
        cutilSafeCall(cudaThreadSynchronize());        // 全てのスレッドが終わるのを待つ
     
        cutilSafeCall(cudaMemcpy((void*)hS, (void*)dS, n*sizeof(float), cudaMemcpyDeviceToHost));
     
        if(dS) cutilSafeCall(cudaFree(dS));
    }
     
    }   // extern "C"
    

ソースコード

実行結果(Core i7 2.93GHz, GeForce GTX 580)

cpu : 1156 [msec]
gpu : 256 [msec]

添付ファイル: filevs2010_cuda4_dlg6.jpg 1839件 [詳細] filevs2010_cuda4_dlg4.jpg 730件 [詳細] filevs2010_cuda4_dlg5.jpg 1935件 [詳細] filevs2010_cuda4_dlg7.jpg 1766件 [詳細] filevs2010_cuda4_dlg3.jpg 2464件 [詳細] filerx_cu_simple.zip 1684件 [詳細] filevs2010_cuda4_dlg2.jpg 2446件 [詳細] filevs2010_cuda4_dlg1.jpg 2353件 [詳細]

トップ   編集 凍結 差分 バックアップ 添付 複製 名前変更 リロード   新規 一覧 単語検索 最終更新   ヘルプ   最終更新のRSS
Last-modified: 2011-10-27 (木) 15:09:18 (3172d)