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;
     
     
    #define BLOCK_SIZE 16
     
    #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"
     
    //-----------------------------------------------------------------------------
    //-----------------------------------------------------------------------------
    
    __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"
     
    //-----------------------------------------------------------------------------
    //-----------------------------------------------------------------------------
    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"
     
     
    //-----------------------------------------------------------------------------
    //-----------------------------------------------------------------------------
    extern "C"
    {
    
    void CuSetDevice(int argc, char **argv)
    {   
        if(cutCheckCmdLineFlag(argc, (const char**)argv, "device")){
            cutilDeviceInit(argc, argv);
        }
        else{
            cudaSetDevice( cutGetMaxGflopsDeviceId() );
        }
    }
     
    
    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());
    }
     
    
    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)));
     
            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_dlg1.jpg 2965件 [詳細] filerx_cu_simple.zip 2130件 [詳細] filevs2010_cuda4_dlg2.jpg 3059件 [詳細] filevs2010_cuda4_dlg3.jpg 3000件 [詳細] filevs2010_cuda4_dlg7.jpg 2270件 [詳細] filevs2010_cuda4_dlg6.jpg 2447件 [詳細] filevs2010_cuda4_dlg4.jpg 949件 [詳細] filevs2010_cuda4_dlg5.jpg 2398件 [詳細]

トップ   編集 凍結 差分 履歴 添付 複製 名前変更 リロード   新規 一覧 検索 最終更新   ヘルプ   最終更新のRSS
Last-modified: 2022-11-30 (水) 13:48:12