こんなサイト始めた
理系のための就活アンテナ
http://kumazaku.com/rikei/
自分で就活で調べた情報とかをまとめていこうかと
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();
よって2.6秒くらいかな??
no error
SET1 : time = 3.144686
no error
SET2 : time = 0.051475
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ベンチマークで遊ぶ
NASAのOpenMPベンチマーク「NPB」
ダウンロードは「http://www.hpcs.cs.tsukuba.ac.jp/omni-openmp/download/download-benchmarks.html」よりできる。
権威のあるベンチマークらしいので、手法の評価のために使えるかも
FFTとか色々入ってる
- 使い方。
- config/make.defを書いて
- make ベンチマーク名
CLASSは問題のサイズを示す。CLASS=S,A,B...とある。
実行結果 FFT 1スレッドの場合
実行結果 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 = 1
Time in seconds = 25.25
Mop/s total = 282.63
Operation type = floating point
Verification = SUCCESSFUL
Version = 2.3
Compile date = 25 Oct 2010Compile options:
CC = gcc-4
CLINK = gcc-4
C_LIB = (none)
C_INC = -I../common
CFLAGS = -fopenmp
CLINKFLAGS = -fopenmp
RAND = randdp
おおむね、2.4倍程度の高速化に成功している
NAS Parallel Benchmarks 2.3 OpenMP C version - FT BenchmarkSize : 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 2010Compile options:
CC = gcc-4
CLINK = gcc-4
C_LIB = (none)
C_INC = -I../common
CFLAGS = -fopenmp
CLINKFLAGS = -fopenmp
RAND = randdp
GPGPU 行列乗算
CPU実行
#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);}
-
- 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 100int 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.636sCUDAヤコビ法
#include
#define N 3000
#define HN 100/*ブロックの数・スレッドの数*/
#ifndef BLOCK_NUM
#define BLOCK_NUM 32
#endif#ifndef THREAD_NUM
#define THREAD_NUM 512
#endifint 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;ireal 0m3.308s
user 0m1.568s
sys 0m1.584s
これだとあまり、効率がよくない。
もっと最適化できる。