• 検索結果がありません。

スレッドが一つの天体を計算

i=blockIdx.x*blockDim.x+threadIdx.x;

for(j=0;j<N;j++){

加速度を積算 }

・・・

x[j]

・・・

加速度を求めたい

x[i]

天体

相互作用の計算の ために参照される

GPU による並列処理の方針

 ベクトル和と同様に 1 スレッドが一つの天体を計算

i=blockIdx.x*blockDim.x+threadIdx.x;

for(j=0;j<N;j++){

加速度を積算 }

・・・

x[j]

・・・

加速度を求めたい

x[i]

天体

相互作用の計算の ために参照される

GPU による並列処理の方針

 ベクトル和と同様に 1 スレッドが一つの天体を計算

i=blockIdx.x*blockDim.x+threadIdx.x;

for(j=0;j<N;j++){

加速度を積算 }

・・・

x[j]

・・・

加速度を求めたい

x[i]

天体

相互作用の計算の ために参照される

GPU による並列処理の方針

 ベクトル和と同様に 1 スレッドが一つの天体を計算

i=blockIdx.x*blockDim.x+threadIdx.x;

for(j=0;j<N;j++){

加速度を積算 }

・・・

x[j]

・・・

加速度を求めたい

x[i]

天体

相互作用の計算の ために参照される

GPU による並列処理の方針

 ベクトル和と同様に 1 スレッドが一つの天体を計算

i=blockIdx.x*blockDim.x+threadIdx.x;

for(j=0;j<N;j++){

加速度を積算 }

・・・

x[j]

・・・

加速度を求めたい

x[i]

天体

相互作用の計算の ために参照される

GPU による並列処理の方針

 ベクトル和と同様に 1 スレッドが一つの天体を計算

i=blockIdx.x*blockDim.x+threadIdx.x;

for(j=0;j<N;j++){

加速度を積算 }

・・・

x[j]

・・・

加速度を求めたい

x[i]

天体

相互作用の計算の ために参照される

#define NT 256

#define NB (N/NT) //加速度の計算

__global__ void kernel(float *x,float *y, float *z,float *vx, float *vy, float *vz,float *ax, float *ay,

float *az,float *m){

int i,j;

float rx,ry,rz;

float dist2, dist6, invDist3,s;

i = blockDim.x*blockIdx.x+threadIdx.x;

for(j=0;j<N;j++){

rx=x[j]‐x[i]; 

ry=y[j]‐y[i];

rz=z[j]‐z[i];

//2天体間の距離を計算

dist2 = rx*rx + ry*ry + rz*rz

+ Soften;//軟化パラメータ // m/r^3 の計算

dist6 = dist2*dist2*dist2;

invDist3 = 1.0/sqrt(dist6);

s = m[j]*invDist3;

//天体jによる加速度を加算 ax[i] = ax[i] + rx*s;

ay[i] = ay[i] + ry*s;

az[i] = az[i] + rz*s;

} }

GPU プログラム（ 1 スレッドが天体一つを計算）

kernel2.cu

//時間積分

__global__ void integrate(float *x, float *y, float *z,

float *vx, float *vy, float *vz, float *ax, float *ay, float *az){

int i=blockIdx.x*blockDim.x + threadIdx.x;

x[i] =  x[i] + dt*vx[i];

y[i] =  y[i] + dt*vy[i];

z[i] =  z[i] + dt*vz[i];

vx[i] = vx[i] + dt*ax[i];

vy[i] = vy[i] + dt*ay[i];

vz[i] = vz[i] + dt*az[i];

}

GPU プログラム（ 1 スレッドが天体一つを計算）

kernel2.cu

実行時間

 天体の個数 N = 4096

 スレッド数 NT = 256

カーネル 実行時間

[ms]

1

スレッド実行

33

×

10

3

1

スレッド実行（

if

の消去）

29

×

10

3

1

スレッド

1

天体

7.63

#define NT 256

#define NB (N/NT) //加速度の計算

__global__ void kernel(float *x,float *y, float *z,float *vx, float *vy, float *vz,float *ax, float *ay,

float *az,float *m){

int i,j;

float rx,ry,rz;

float dist2, dist6, invDist3,s;

//加速度のデータをレジスタに置く float r_ax, r_ay, r_az;

i = blockDim.x*blockIdx.x+threadIdx.x;

r_ax = r_ay = r_az = 0.0f;

for(j=0;j<N;j++){

rx=x[j]‐x[i]; 

ry=y[j]‐y[i];

rz=z[j]‐z[i];

//2天体間の距離を計算

dist2 = rx*rx + ry*ry + rz*rz

+ Soften;//軟化パラメータ // m/r^3 の計算

dist6 = dist2*dist2*dist2;

invDist3 = 1.0/sqrt(dist6);

s = m[j]*invDist3;

//天体jによる加速度を加算 r_ax = r_ax + rx*s;

r_ay = r_ay + ry*s;

r_az = r_az + rz*s;

}

ax[i] = r_ax;

ay[i] = r_ay;

az[i] = r_az;

}

GPU プログラム（レジスタ利用）

kernel3.cu

実行時間

 天体の個数 N = 4096

 スレッド数 NT = 256

カーネル 実行時間

[ms]

1

スレッド実行

33

×

10

3

1

スレッド実行（

if

の消去）

29

×

10

3

1

スレッド

1

天体

7.63

レジスタ利用

4.49

GPU プログラム（共有メモリによる再利用）

x[i]

x[j]

・・・

共有メモリ

blockIdx.x=0 blockIdx.x=1

i=   0    1   2    3   4    5   6    7 threadIdx.x=   0    1   2    3   0    1   2    3

x[0] x[1] x[2] x[3]   x[0] x[1] x[2] x[3]

j=   0    1   2    3   4    5   6    7

0+threadIdx.x

0+threadIdx.x

GPU プログラム（共有メモリによる再利用）

x[i]

x[j]

・・・

共有メモリ

blockIdx.x=0 blockIdx.x=1

i=0~3

の天体と

j=0~3

の天体で 加速度を計算

i=4~7

の天体と

j=0~3

の天体で 加速度を計算

threadIdx.x=   0    1   2    3   0    1   2    3

x[0] x[1] x[2] x[3]   x[0] x[1] x[2] x[3]

j=   0    1   2    3   4    5   6    7

i=   0    1   2    3   4    5   6    7

GPU プログラム（共有メモリによる再利用）

x[i]

x[j]

・・・

共有メモリ

blockIdx.x=0 blockIdx.x=1

i=0~3

の天体と

j=0~3

の天体で 加速度を計算

i=4~7

の天体と

j=0~3

の天体で 加速度を計算

threadIdx.x=   0    1   2    3   0    1   2    3

x[0] x[1] x[2] x[3]   x[0] x[1] x[2] x[3]

j=   0    1   2    3   4    5   6    7

i=   0    1   2    3   4    5   6    7

GPU プログラム（共有メモリによる再利用）

x[i]

x[j]

・・・

共有メモリ

blockIdx.x=0 blockIdx.x=1

i=0~3

の天体と

j=0~3

の天体で 加速度を計算

i=4~7

の天体と

j=0~3

の天体で 加速度を計算

threadIdx.x=   0    1   2    3   0    1   2    3

x[0] x[1] x[2] x[3]   x[0] x[1] x[2] x[3]

j=   0    1   2    3   4    5   6    7

i=   0    1   2    3   4    5   6    7

GPU プログラム（共有メモリによる再利用）

x[i]

x[j]

・・・

共有メモリ

blockIdx.x=0 blockIdx.x=1

i=0~3

の天体と

j=0~3

の天体で 加速度を計算

i=4~7

の天体と

j=0~3

の天体で 加速度を計算

threadIdx.x=   0    1   2    3   0    1   2    3

x[0] x[1] x[2] x[3]   x[0] x[1] x[2] x[3]

j=   0    1   2    3   4    5   6    7

i=   0    1   2    3   4    5   6    7

GPU プログラム（共有メモリによる再利用）

x[i]

x[j]

・・・

共有メモリ

blockIdx.x=0 blockIdx.x=1

i=   0    1   2    3   4    5   6    7 threadIdx.x=   0    1   2    3   0    1   2    3

x[4] x[5] x[6] x[7]   x[4] x[5] x[6] x[7]

j=   0    1   2    3   4    5   6    7

4(=NT)+threadIdx.x

4(=NT)+threadIdx.x

GPU プログラム（共有メモリによる再利用）

x[i]

x[j]

・・・

共有メモリ

blockIdx.x=0 blockIdx.x=1

i=0~3

の天体と

j=4~7

の天体で 加速度を計算

i=4~7

の天体と

j=4~7

の天体で 加速度を計算

threadIdx.x=   0    1   2    3   0    1   2    3

j=   0    1   2    3   4    5   6    7 i=   0    1   2    3   4    5   6    7

x[4] x[5] x[6] x[7]   x[4] x[5] x[6] x[7]

GPU プログラム（共有メモリによる再利用）

x[i]

x[j]

・・・

共有メモリ

blockIdx.x=0 blockIdx.x=1

i=0~3

の天体と

j=4~7

の天体で 加速度を計算

i=4~7

の天体と

j=4~7

の天体で 加速度を計算

threadIdx.x=   0    1   2    3   0    1   2    3

j=   0    1   2    3   4    5   6    7 i=   0    1   2    3   4    5   6    7

x[4] x[5] x[6] x[7]   x[4] x[5] x[6] x[7]

GPU プログラム（共有メモリによる再利用）

x[i]

x[j]

・・・

共有メモリ

blockIdx.x=0 blockIdx.x=1

i=0~3

の天体と

j=4~7

の天体で 加速度を計算

i=4~7

の天体と

j=4~7

の天体で 加速度を計算

threadIdx.x=   0    1   2    3   0    1   2    3

j=   0    1   2    3   4    5   6    7 i=   0    1   2    3   4    5   6    7

x[4] x[5] x[6] x[7]   x[4] x[5] x[6] x[7]

GPU プログラム（共有メモリによる再利用）

x[i]

x[j]

・・・

共有メモリ

blockIdx.x=0 blockIdx.x=1

i=0~3

の天体と

j=4~7

の天体で 加速度を計算

i=4~7

の天体と

j=4~7

の天体で 加速度を計算

threadIdx.x=   0    1   2    3   0    1   2    3

j=   0    1   2    3   4    5   6    7 i=   0    1   2    3   4    5   6    7

x[4] x[5] x[6] x[7]   x[4] x[5] x[6] x[7]

#define NT 256

#define NB (N/NT) //加速度の計算

__global__ void kernel(…){

int i,j;

float rx,ry,rz;

float dist2, dist6, invDist3,s;

//加速度と座標のデータをレジスタに置く float r_ax,r_ay,r_az,r_x,r_y,r_z;

__shared__ float 

s_x[NT],s_y[NT],s_z[NT],s_m[NT];

i = blockDim.x*blockIdx.x+threadIdx.x;

r_ax = r_ay = r_az = 0.0f;

r_x = x[i];

r_y = y[i];

r_z = z[i];

for(j=0;j<N;j+=NT){

s_x[threadIdx.x] = x[j+threadIdx.x];

s_y[threadIdx.x] = y[j+threadIdx.x];

s_z[threadIdx.x] = z[j+threadIdx.x];

s_m[threadIdx.x] = m[j+threadIdx.x];

__syncthreads();

for(js = 0;js<NT;js++){

rx=s_x[js]‐r_x;

ry=s_y[js]‐r_y;

rz=s_z[js]‐r_z;

//2天体間の距離を計算

dist2 = rx*rx + ry*ry + rz*rz

+ Soften;//軟化パラメータ // m/r^3 の計算

dist6 = dist2*dist2*dist2;

invDist3 = 1.0/sqrt(dist6);

s = s_m[js]*invDist3;

//天体jによる加速度を加算 r_ax = r_ax + rx*s;

r_ay = r_ay + ry*s;

r_az = r_az + rz*s;

} }

ax[i] = r_ax;

ay[i] = r_ay;

az[i] = r_az;

}

GPU プログラム（共有メモリによる再利用）

kernel4.cu

実行時間

 天体の個数 N = 4096

 スレッド数 NT = 256

カーネル 実行時間

[ms]

1

スレッド実行

33

×

10

3

1

スレッド実行（

if

の消去）

29

×

10

3

1

スレッド

1

天体

7.63

レジスタ利用

4.49

共有メモリ利用

4.27

#define NT 256

#define NB (N/NT) //加速度の計算

__global__ void kernel(…){

int i,j;

float rx,ry,rz;

float dist2, dist6, invDist3,s;

//加速度と座標のデータをレジスタに置く float r_ax,r_ay,r_az,r_x,r_y,r_z;

__shared__ float 

s_x[NT],s_y[NT],s_z[NT],s_m[NT];

i = blockDim.x*blockIdx.x+threadIdx.x;

r_ax = r_ay = r_az = 0.0f;

r_x = x[i];

r_y = y[i];

r_z = z[i];

for(j=0;j<N;j+=NT){

s_x[threadIdx.x] = x[j+threadIdx.x];

s_y[threadIdx.x] = y[j+threadIdx.x];

s_z[threadIdx.x] = z[j+threadIdx.x];

s_m[threadIdx.x] = m[j+threadIdx.x];

#pragma unroll

for(js = 0;js<NT;js++){

rx=s_x[js]‐r_x;

ry=s_y[js]‐r_y;

rz=s_z[js]‐r_z;

//2天体間の距離を計算

dist2 = rx*rx + ry*ry + rz*rz

+ Soften;//軟化パラメータ // m/r^3 の計算

dist6 = dist2*dist2*dist2;

invDist3 = 1.0/sqrt(dist6);

s = s_m[js]*invDist3;

//天体jによる加速度を加算 r_ax = r_ax + rx*s;

r_ay = r_ay + ry*s;

r_az = r_az + rz*s;

} }

ax[i] = r_ax;

ay[i] = r_ay;

az[i] = r_az;

}

GPU プログラム（ループアンロール）

kernel5.cu

実行時間

 天体の個数 N = 4096

 スレッド数 NT = 256

カーネル 実行時間

[ms]

1

スレッド実行

33

×

10

3

1

スレッド実行（

if

の消去）

29

×

10

3

1

スレッド

1

天体

7.63

レジスタ利用

4.49

共有メモリ利用

4.27

共有メモリ利用

+

ループアンロール

3.50

今ダウンロードする "N 体問題 長岡技術科学..."

Outline

関連したドキュメント