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 
    #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 
    #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 
    #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 
    //-----------------------------------------------------------------------------
// インクルードファイル
//-----------------------------------------------------------------------------
#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"
↑

ソースコード†

filerx_cu_simple.zip

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

cpu : 1156 [msec]
gpu : 256 [msec]
添付ファイル: filerx_cu_simple.zip 2276件 [詳細] filevs2010_cuda4_dlg1.jpg 3151件 [詳細] filevs2010_cuda4_dlg2.jpg 3222件 [詳細] filevs2010_cuda4_dlg3.jpg 3149件 [詳細] filevs2010_cuda4_dlg4.jpg 1020件 [詳細] filevs2010_cuda4_dlg5.jpg 2553件 [詳細] filevs2010_cuda4_dlg6.jpg 2590件 [詳細] filevs2010_cuda4_dlg7.jpg 2422件 [詳細]
トップ   編集 凍結 差分 履歴 添付 複製 名前変更 リロード   新規 一覧 検索 最終更新   ヘルプ   最終更新のRSS
Last-modified: 2024-03-08 (金) 18:06:10