CUDAの初回呼び出しオーバーヘッド

CUDAは初めてkernelコールを行ったときだけ特別に時間がかかる。
(初期化とかやってるの?)それがどのくらいか計測した。


t3 = gettimeofday_sec();
setMatrixOnGPU(*a);
t4 = gettimeofday_sec();
printf("SET1 : time = %f\n", t4 - t3);
sleep(3);
t3 = gettimeofday_sec();
setMatrixOnGPU(*b);
t4 = gettimeofday_sec();
結果

no error
SET1 : time = 3.144686
no error
SET2 : time = 0.051475
よって2.6秒くらいかな??

Beamerでプレゼン

Latexでプレゼンをやりたい!
そこでbeamerといういいツールがある。
ダウンロードはここ「http://sourceforge.net/projects/latex-beamer/
ちなみにxcolorとpgf2.0「http://sourceforge.net/projects/pgf/
size14.clo「http://www.ctan.org/tex-archive/macros/latex/contrib/extsizes/」が必要なので落としておく。
ここで落としたファイルはwindowsの場合は
C:\tex\share\texmf\tex\latex\pgf
C:\tex\share\texmf\tex\latex\beamer
C:\tex\share\texmf\tex\latex\xcolor
C:\tex\share\texmf\tex\latex\base\size14.clo
にそれぞれコピーする。
あとは
http://windom.phys.hirosaki-u.ac.jp/fswiki/wiki.cgi?page=LaTeX+Beamer%A4%C7%A5%D7%A5%EC%A5%BC%A5%F3%A5%C6%A1%BC%A5%B7%A5%E7%A5%F3
このサイトでも参考にしてコードを書いてコンパイルする。

NASAのOpenMPベンチマークで遊ぶ

NASAOpenMPベンチマークNPB
ダウンロードは「http://www.hpcs.cs.tsukuba.ac.jp/omni-openmp/download/download-benchmarks.html」よりできる。
権威のあるベンチマークらしいので、手法の評価のために使えるかも
FFTとか色々入ってる

CLASSは問題のサイズを示す。CLASS=S,A,B...とある。
実行結果 FFT 1スレッドの場合


NAS Parallel Benchmarks 2.3 OpenMP C version - FT Benchmark

Size : 256x256x128
Iterations : 6
T = 1 Checksum = 5.046735008193e+02 5.114047905510e+02
T = 2 Checksum = 5.059412319734e+02 5.098809666433e+02
T = 3 Checksum = 5.069376896287e+02 5.098144042213e+02
T = 4 Checksum = 5.077892868474e+02 5.101336130759e+02
T = 5 Checksum = 5.085233095391e+02 5.104914655194e+02
T = 6 Checksum = 5.091487099959e+02 5.107917842803e+02
Result verification successful
class = A


FT Benchmark Completed
Class = A
Size = 256x256x128
Iterations = 6
Threads = 1
Time in seconds = 25.25
Mop/s total = 282.63
Operation type = floating point
Verification = SUCCESSFUL
Version = 2.3
Compile date = 25 Oct 2010

Compile options:
CC = gcc-4
CLINK = gcc-4
C_LIB = (none)
C_INC = -I../common
CFLAGS = -fopenmp
CLINKFLAGS = -fopenmp
RAND = randdp

実行結果 FFT 4スレッドの場合


NAS Parallel Benchmarks 2.3 OpenMP C version - FT Benchmark

Size : 256x256x128
Iterations : 6
T = 1 Checksum = 5.046735008193e+02 5.114047905510e+02
T = 2 Checksum = 5.059412319734e+02 5.098809666433e+02
T = 3 Checksum = 5.069376896287e+02 5.098144042213e+02
T = 4 Checksum = 5.077892868474e+02 5.101336130759e+02
T = 5 Checksum = 5.085233095391e+02 5.104914655194e+02
T = 6 Checksum = 5.091487099959e+02 5.107917842803e+02
Result verification successful
class = A


FT Benchmark Completed
Class = A
Size = 256x256x128
Iterations = 6
Threads = 4
Time in seconds = 10.57
Mop/s total = 675.16
Operation type = floating point
Verification = SUCCESSFUL
Version = 2.3
Compile date = 25 Oct 2010

Compile options:
CC = gcc-4
CLINK = gcc-4
C_LIB = (none)
C_INC = -I../common
CFLAGS = -fopenmp
CLINKFLAGS = -fopenmp
RAND = randdp

おおむね、2.4倍程度の高速化に成功している

GPGPU 行列乗算


#include
#include
#include
// Matrices are stored in row-major order:
// M(row, col) = *(M.elements + row * M.width + col)
typedef struct {
int width;
int height;
float* elements;
} Matrix;

// Thread block size
#define BLOCK_SIZE 8
#define MAT_DIM 64*64

// Forward declaration of the matrix multiplication kernel
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);

// Matrix multiplication - Host code
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE
void MatMul(const Matrix A, const Matrix B, Matrix* C)
{
int i,j;
// Load A and B to device memory
Matrix d_A;
d_A.width = A.width; d_A.height = A.height;
size_t size = A.width * A.height * sizeof(float);
cudaMalloc(&d_A.elements, size);
cudaMemcpy(d_A.elements, A.elements, size,
cudaMemcpyHostToDevice);
Matrix d_B;
d_B.width = B.width; d_B.height = B.height;
size = B.width * B.height * sizeof(float);
cudaMalloc(&d_B.elements, size);
cudaMemcpy(d_B.elements, B.elements, size,
cudaMemcpyHostToDevice);

// Allocate C in device memory
Matrix d_C;
d_C.width = C->width; d_C.height = C->height;
size = C->width * C->height * sizeof(float);
cudaMalloc(&d_C.elements, size);

// Invoke kernel
dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
printf(" %d %d ",B.width / dimBlock.x, A.height / dimBlock.y);
dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
MatMulKernel<<>>(d_A, d_B, d_C);
printf("%s\n",cudaGetErrorString(cudaGetLastError()));

// Read C from device memory
cudaMemcpy(C->elements, d_C.elements, size,
cudaMemcpyDeviceToHost);
cudaThreadSynchronize();

// Free device memory
cudaFree(d_A.elements);
cudaFree(d_B.elements);
cudaFree(d_C.elements);
}

// Matrix multiplication kernel called by MatMul()
__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)
{
// Each thread computes one element of C
// by accumulating results into Cvalue
float Cvalue = 0;
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
for (int e = 0; e < A.width; ++e)
Cvalue += A.elements[row * A.width + e]
* B.elements[e * B.width + col];
C.elements[row * C.width + col] = Cvalue;
}
int getRandom(int min,int max)
{
return min + (int)(rand()*(max-min+1.0)/(1.0+RAND_MAX));
}


void setMatrix(Matrix *a)
{
int i=0;
for(i=0;iwidth*a->width;i++)
a->elements[i]=(float)getRandom(0,10);
}

Matrix* createMatrix(int size){
int i=0;
Matrix* a=(Matrix*)malloc(size*sizeof(Matrix));
a->height=size;
a->width=size;
a->elements=(float*)malloc(size*size*sizeof(float));
for(i=0;iheight*a->width;i++)
a->elements[i]=0;

return a;
}

double gettimeofday_sec()
{
struct timeval tv;
gettimeofday(&tv, NULL);
return tv.tv_sec + (double)tv.tv_usec*1e-6;
}

void printMatrix(Matrix *c)
{
int i=0,j=0;
for(i=0;iheight;i++)
{
for(j=0;jwidth;j++)
printf("%f ",c->elements[i*c->width+j]);
printf("\n");
}

printf("\n");
}


void MatMulCPU(Matrix* A,Matrix* B, Matrix* C){

int i,j,k;
for(i=0;iheight;i++)
for(j=0;jwidth;j++)
{
for(k=0;kwidth;k++)
C->elements[i*C->width+j]=C->elements[i*C->width+j]+(A->elements[A->width*i+k]*B->elements[B->width*k+j]);
}

}

int main(){
double t1,t2;
Matrix* a=createMatrix(MAT_DIM);
Matrix* b=createMatrix(MAT_DIM);
Matrix* c=createMatrix(MAT_DIM);

setMatrix(a);
setMatrix(b);

t1 = gettimeofday_sec();
#ifndef HOST
MatMul(*a,*b,c);
#else
MatMulCPU(a,b,c);
#endif
t2 = gettimeofday_sec();
printf("GPU : time = %f\n", t2 - t1);

// printMatrix(a);
// printMatrix(b);
// printMatrix(c);

}

CPU実行

    • 5分経っても終わらず。強制終了

GPU実行
512 512 no error
GPU : time = 9.788575
real 0m11.628s
user 0m9.641s
sys 0m1.832s

GPGPU ヤコビ法

CUDAにより、Ax=bの解をヤコビ法で求めてみる。
CPU版


#include
#include
#define N 3000
#define HN 100

int getRandom(int min,int max)
{
return min + (int)(rand()*(max-min+1.0)/(1.0+RAND_MAX));
}


int main(int argc, char *argv[])
{
int i,j,k;
float x[N],x2[N],a[N*N],b[N];
float temp;

// float a[N*N]={3.0,-6.0,9.0,2.0,5.0,-8.0,1.0,-4.0,7.0};
// float b[N]={6.0,8.0,2.0},x[N]={1.0,1.0,1.0};

/*Set Matrix default value*/
for(i=0;i
実行時間
real 0m7.103s
user 0m5.320s
sys 0m1.636s

CUDAヤコビ法


#include
#define N 3000
#define HN 100

/*ブロックの数・スレッドの数*/
#ifndef BLOCK_NUM
#define BLOCK_NUM 32
#endif

#ifndef THREAD_NUM
#define THREAD_NUM 512
#endif

int getRandom(int min,int max)
{
return min + (int)(rand()*(max-min+1.0)/(1.0+RAND_MAX));
}


__global__ void jacobi(float *a,float *x,float *b)
{
float temp=0.0;
int k;
int j=blockIdx.x*THREAD_NUM+threadIdx.x;
int i=0;

if(j>=N) return;
temp=0.0;
for(k=0;k>>(devA,devX,devB);
cudaThreadSynchronize();
cudaMemcpy(x,devX, N*sizeof(float), cudaMemcpyDeviceToHost);
cudaMemcpy(devX,x, N*sizeof(float), cudaMemcpyHostToDevice);
}

cudaFree((void**)&devA);
cudaFree((void**)&devX);
cudaFree((void**)&devB);

/*print result*/
// for(i=0;i

real 0m3.308s
user 0m1.568s
sys 0m1.584s
これだとあまり、効率がよくない。
もっと最適化できる。