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

----
#contents
----

*VS2010+CUDA4.0 [#gb103323]
**CUDAインストール [#z4ef5b2b]
-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プロジェクトの変更 [#s2028291]
-ビルドターゲットの設定~
プロジェクトを右クリックして,「ビルドのカスタマイズ」を選択.
CUDA 4.0(.targets, .props)にチェックを入れる.
#ref(vs2010_cuda4_dlg1.jpg,,50%)

-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ファイルの追加 [#lb436490]
+CUファイルの追加
.cuや.cuhファイルをプロジェクトに追加する.
+CUDA C/C++の設定
CUファイルのプロパティから全般->項目の種類をCUDA C/C++にして,
CUDA C/C++の各項目を設定する.
#ref(vs2010_cuda4_dlg2.jpg,,50%)

私の設定項目は,
-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) -> ホストのランタイムライブラリ設定に合わせる
#ref(vs2010_cuda4_dlg3.jpg,,50%)
#ref(vs2010_cuda4_dlg5.jpg,,50%)
#ref(vs2010_cuda4_dlg7.jpg,,50%)
cuhファイルやデバイス関数(__device__,__global__)のみが書かれたcuファイルに関しては
ビルドから除外を「はい」にする.
#ref(vs2010_cuda4_dlg6.jpg,,50%)

*CUDAファイルの構成例1(カーネル関数を分ける) [#u176bf36]
***ファイル構成 [#lfffc870]
-rx_*_kernel.cuh : カーネルヘッダ,定義などを記述.rx_*_kernel.cu,rx_*.cuh,rx_*.cu,CUDA呼び出し側からインクルードする.
-rx_*_kernel.cu : カーネル関数の実装.rx_*.cuからインクルードする.
-rx_*.cuh : ホスト関数のヘッダ.CUDA呼び出し側からインクルードする.
-rx_*.cu : ホスト関数の実装.CUDA呼び出し側(VCコンパイラを使用するコード)からインクルードしないこと.
***コード例 [#v1fff973]
-rx_*_kernel.cuh
#code(C){{
#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
#code(C){{
#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
#code(C){{
#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
#code(C){{
//-----------------------------------------------------------------------------
// インクルードファイル
//-----------------------------------------------------------------------------
#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"
}}

***ソースコード [#h6329770]
#ref(rx_cu_simple.zip)
実行結果(Core i7 2.93GHz, GeForce GTX 580)
 cpu : 1156 [msec]
 gpu : 256 [msec]

トップ   編集 差分 バックアップ 添付 複製 名前変更 リロード   新規 一覧 単語検索 最終更新   ヘルプ   最終更新のRSS