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]
![[PukiWiki]](image/pukiwiki.png "[PukiWiki]")
