環境
Windows Vista Home Premium Service Pack 2
プロセッサ : Intel(R) Core(TM)2 Duo CPU T8100 @2.10GHz 2.10GHz
メモリ(RAM) : 4.00GB
システムの種類 : 32ビット オペレーティング システム
姫野ベンチマーク(実行形式、M) : 1182.665 MFLOPS (1回測定)
まずは、ダウンロードページを開き、
CUDA Toolkit (nvccなどが入っている)とCUDA SDK code samples (ヘッダなどが入っている)をダウンロードします。
CUDA Toolkit 2.3 Downloads | NVIDIA Developer Zone
ダウンロードしたら、適当にインストールします。
とりあえずnvccを実行させると、cl.exeに依存しているようです。
とりあえずgcc.exeをcl.exeにコピーしてみてもうまくいかないようだったので(当たり前)、
諦めておとなしくこのサイトからVisual C++のisoをダウンロードし、インストールしました。
NonSoft - Visual Studio 2008 Expressのダウンロードとインストール
gizmoでマウントしてもうまくいかなかったですが、7-zipで展開したらインストールできました。
インストールしてもnvccが「cl.exeが見つからない」と文句を言ってきますが、
普通のコマンドプロンプトではなく、スタートメニューの中にある「Visual Studio 2008 コマンド プロンプト」を使用するとうまくいきます。
CUDA入門・サンプル集
ゼロから始めるGPU Computingの一覧 | G-DEP
ここらへんのサイトを参考に、いくつかのサンプルを動かしたあと、自分で考えてプログラムを組んでみます。
今回は離散フーリエ変換のプログラムを用意しました。
まずはCUDAを使わない、従来のプログラム(rfh_normal.c)
#include
#include
#include
#define THE_PI 3.1415926535897932384626433832795028841f
void risanfuuriehenkan(float out[],const float in[],int N) {
int x;
#ifdef _OPENMP
#pragma omp parallel for
#endif
for(x=0;x
#include
#include
#define THE_PI 3.1415926535897932384626433832795028841f
__global__ void risanfuuriehenkan(float out[],const float in[],int N) {
float a=0,b=0;
int i=blockIdx.x*512+threadIdx.x;
int j;
for(j=0;j65535*512+511 || (N>512 && N%512!=0)) {
puts("invalid N");
return 1;
}
if(N%32!=0) {
puts("warning: N should be a multiple of 32");
}
if(N>512) {
bnum=N/512;
snum=512;
} else {
bnum=1;
snum=N;
}
in=(float*)malloc(sizeof(float)*N);
out=(float*)malloc(sizeof(float)*N);
if(in==NULL || out==NULL) {
puts("malloc error");
if(in)free(in);
if(out)free(out);
return 1;
}
for(i=0;i>>(deviceOut,deviceIn,N);
CUDA_SAFE_CALL(cudaMemcpy(out,deviceOut,sizeof(float)*N,cudaMemcpyDeviceToHost));
for(i=0;i
#include
#include
int main(int argc,char* argv[]) {
int N=512,i;
srand((unsigned int)time(NULL));
if(argc>=2)sscanf(argv[1],"%d",&N);
printf("%d\n",N);
for(i=0;i<N;i++)printf("%g\n",(double)rand()/RAND_MAX);
return 0;
}
.PHONY: all
all: rfh_cuda.exe rfh_normal.exe rfh_omp.exe ;
rfh_cuda.exe: rfh_cuda.cu
nvcc -O2 -o rfh_cuda.exe rfh_cuda.cu -I "C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK\C\common\inc" -deviceemu
rfh_normal.exe: rfh_normal.c
gcc -O2 -o rfh_normal.exe rfh_normal.c
rfh_omp.exe: rfh_normal.c
gcc -O2 -fopenmp -o rfh_omp.exe rfh_normal.c
今回は、N=1024,2048,4096,6144,8192の入力を用意しました。
それぞれ1回ずつ測定します。
測定結果(時間の単位はms)
[table=border:1px solid white;border-collapse: collapse;][tr=][td=border:1px solid white;padding: 0.5em;]N[/td][td=border:1px solid white;padding: 0.5em;]工夫なし[/td]
[td=border:1px solid white;padding: 0.5em;]OpenMP[/td][td=border:1px solid white;padding: 0.5em;]CUDA[/td]
[td=border:1px solid white;padding: 0.5em;]工夫なし/OpenMP[/td]
[td=border:1px solid white;padding: 0.5em;]工夫なし/CUDA[/td][/tr]
[tr=][td=border:1px solid white;padding: 0.5em;]1024[/td][td=border:1px solid white;padding: 0.5em;]339.053776[/td]
[td=border:1px solid white;padding: 0.5em;]205.399778[/td][td=border:1px solid white;padding: 0.5em;]985.876976[/td]
[td=border:1px solid white;padding: 0.5em;]1.650701765[/td][td=border:1px solid white;padding: 0.5em;]0.343910837[/td][/tr]
[tr=][td=border:1px solid white;padding: 0.5em;]2048[/td][td=border:1px solid white;padding: 0.5em;]1099.708832[/td]
[td=border:1px solid white;padding: 0.5em;]613.226402[/td][td=border:1px solid white;padding: 0.5em;]3496.804412[/td]
[td=border:1px solid white;padding: 0.5em;]1.793316185[/td][td=border:1px solid white;padding: 0.5em;]0.31448966[/td][/tr]
[tr=][td=border:1px solid white;padding: 0.5em;]4096[/td][td=border:1px solid white;padding: 0.5em;]4111.300528[/td]
[td=border:1px solid white;padding: 0.5em;]2167.149805[/td][td=border:1px solid white;padding: 0.5em;]13477.011743[/td]
[td=border:1px solid white;padding: 0.5em;]1.897100292[/td][td=border:1px solid white;padding: 0.5em;]0.305060247[/td][/tr]
[tr=][td=border:1px solid white;padding: 0.5em;]6144[/td][td=border:1px solid white;padding: 0.5em;]9296.543695[/td]
[td=border:1px solid white;padding: 0.5em;]4994.841244[/td][td=border:1px solid white;padding: 0.5em;]30281.241261[/td]
[td=border:1px solid white;padding: 0.5em;]1.861229064[/td][td=border:1px solid white;padding: 0.5em;]0.307006692[/td][/tr]
[tr=][td=border:1px solid white;padding: 0.5em;]8192[/td][td=border:1px solid white;padding: 0.5em;]16038.909065[/td]
[td=border:1px solid white;padding: 0.5em;]8806.296471[/td][td=border:1px solid white;padding: 0.5em;]53111.161684[/td]
[td=border:1px solid white;padding: 0.5em;]1.821300148[/td][td=border:1px solid white;padding: 0.5em;]0.30198754[/td][/tr][/table] OpenMPを使用すると平均で1.804729491倍の高速化ができたのに対し、
CUDAを使用すると平均でなんと0.314490995倍の高速化に成功しました!
しかし、CUDA版の出力は、工夫なし版と比べて誤差が出ていました。
OpenMPでも誤差が出るので、今回は少数点以下を出力しないことで誤差を隠そうとしたにもかかわらず、です。
といっても、出力のほとんどの数値は一致しているので、やはり誤差のようです。
私のプログラムが遅いのはどう考えてもPCが悪い!