HOME サイトマップ  
大学研究機関等パワーユーザの方々を対象に
主に科学計算用 Custom-Made PCを製作販売しております

トップページ > GPUクラスタ構築例 > MPI+CUDA 姫野ベンチマーク

当社製GPUクラスタの構成例です


MPI + CUDA 姫野ベンチマーク

MPIの姫野ベンチマークのCUDAポーティングを行ったものを使用して、MPI+CUDAの開発環境におけるコンパイルおよび
アプリケーション実行を動作確認しています.
姫野ベンチマークをGPU上で実行して高い性能が得られることは多くの研究などで報告されていますが、GPUの性能を
十分に発揮させるにはチューニングが必要となります.
ここではCPUコードを極力変更せず、簡単なCUDAポーティングのみでどの程度の性能が出るかを紹介します.

コードは姫野ベンチマークのC言語・MPI版を使用しています.あまりCPUコードに変更を加えないようにし、ホスト側の
配列サイズの調整や、カーネル実行とデータ転送のオーバーラップなどは行っていません.また、境界要素に関する
プロセス間のデータ転送の部分も簡略化するために、計算領域の分割はI方向のみに限定しています.これによって
ホスト−デバイス間もしくはホスト−ホスト間でのデータ転送時に、不連続な配列要素が現れなくなります.
CUDA化に関する主な追加・変更点は以下の3箇所になります.

(1)デバイスメモリの割り当て・解放、データ転送
カーネルで使用するデバイスメモリの確保・解放や初期値のホストからデバイスへの転送を行うためのAPIを追加しました.
また、プロセス間の通信を行う際には一度デバイスからホストへデータ転送を行うようにしているため、そこで必要な
ホストメモリの確保をcudaHostMallocにより行っています. 以下はコードの抜粋となります.

 if(id==0){
  printf(" Start rehearsal measurement process.\n");
  printf(" Measure the performance in %d times.\n\n",nn);
 }
 ・・・
 // デバイスメモリの確保
 cudaMalloc((void **)&p_d, arrsize * sizeof(float));
 cudaMalloc((void **)&a_d, 4 * arrsize * sizeof(float));
 ・・・
 // MPI通信に使うホストメモリ確保
 float *send_buf_u, *recv_buf_u, *send_buf_l, *recv_buf_l;
 cudaMallocHost( (void**) &send_buf_u, sizeof(float)*MKMAX*MJMAX);
 cudaMallocHost( (void**) &recv_buf_u, sizeof(float)*MKMAX*MJMAX);

 // 初期値のホストからデバイスへの転送
 cudaMemcpy(p_d, p, arrsize * sizeof(float), cudaMemcpyHostToDevice);
 cudaMemcpy(a_d, a, 4*arrsize*sizeof(float), cudaMemcpyHostToDevice);
 ・・・
 MPI_Barrier(MPI_COMM_WORLD);

 // GPU版のjacobi関数の呼び出し
 cpu0= MPI_Wtime();
 gosa = jacobi_dev(nn, p_d, a_d, b_d, c_d, wrk1_d, wrk2_d, bnd_d, omega, ndx, ndy, ndz, send_buf_u, recv_buf_u, send_buf_l, recv_buf_l);
 cpu1= MPI_Wtime() - cpu0;



(2)反復計算の中のCUDAカーネル化
関数jacobiに関して、CUDAカーネルの呼び出しや、プロセス間のMPIデータ転送を行う部分の追加を行っています.
CUDAカーネルは、配列wrk2と変数ssの計算部分と配列pの更新部分を分けて実装しています.
どちらもCUDAカーネルについてはK,J方向の2次元のCUDA Gridを作成し、各スレッドでI方向の反復計算を行うように
しています.
この計算部分でも特にチューニングは行わず、単純にCPUコードをCUDAカーネルに記述しているだけです.
以下はwrk2_dを計算する部分の抜粋です.


__global__ void jacobi_kernel ( int imax, int jmax, int kmax, float *p_d, float *a_d, float *b_d, float *c_d,
float *wrk1_d, float *wrk2_d, float *bnd_d, float *reduce_tmp, float omega)
{
float s0, ss;
float tmp_t = 0;

 int k = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
int i;

 if ( j == 0 ) goto __end;
if ( k == 0 ) goto __end;
if ( j >= jmax - 1 ) goto __end;
if ( k >= kmax - 1 ) goto __end;

for ( i = 1 ; i < imax - 1 ; i++){
s0 = a_d[IDX_COEF(0,i,j,k)] * p_d[IDX(i+1,j ,k )]
+ a_d[IDX_COEF(1,i,j,k)] * p_d[IDX(i ,j+1,k)]
+ a_d[IDX_COEF(2,i,j,k)] * p_d[IDX(i ,j ,k+1)]
・・・



(3)MPI通信
前述の通り領域分割をI方向のみに限定することで、プロセス間のデータ転送に不連続なメモリ領域が含まれないように
しています.
また、ここではCUDA Aware MPIを使わずに、CUDAでデバイス−ホスト間、MPIでホスト−ホスト間と別々にデータ転送を
行っています.
以下はデータ転送部分の抜粋です.


if ( ndx > 1) {
 if ( npx[0] != -2) cudaMemcpy(&send_buf_l[0], &p_d[IDX(1,0,0)], data_jkplane, cudaMemcpyDeviceToHost);
 if ( npx[1] != -2) cudaMemcpy(&send_buf_u[0], &p_d[IDX(imax-2,0,0)], data_jkplane, cudaMemcpyDeviceToHost);

 MPI_Status st[4];
 MPI_Request req[4];
 MPI_Barrier(MPI_COMM_WORLD);

 MPI_Irecv(&recv_buf_u[0], MJMAX * MKMAX, MPI_FLOAT, npx[1], 1, mpi_comm_cart, req);
 MPI_Irecv(&recv_buf_l[0], MJMAX * MKMAX, MPI_FLOAT, npx[0], 2, mpi_comm_cart, req+1);

 MPI_Isend(&send_buf_l[0], MJMAX * MKMAX, MPI_FLOAT, npx[0], 1, mpi_comm_cart, req+2);
 MPI_Isend(&send_buf_u[0], MJMAX * MKMAX, MPI_FLOAT, npx[1], 2, mpi_comm_cart, req+3);

 MPI_Waitall(4, req, st);

 if( npx[1] != -2) cudaMemcpy(&p_d[IDX(imax-1,0,0)], &recv_buf_u[0], data_jkplane,  cudaMemcpyHostToDevice);
 if( npx[0] != -2) cudaMemcpy(&p_d[IDX(0,0,0)], &recv_buf_l[0], data_jkplane,  cudaMemcpyHostToDevice);
}


以上のような追加・修正を行ったコードをGPUで実行した結果を図3に示します.
ここではオリジナルバージョン(CPUのみ)とGPUバージョンのスコアを示しており、グラフ中の横軸がノード数、
縦軸がスコアとなります. CPU計算の場合は1ノードあたり16コア、GPU計算の場合は1ノードあたり4GPU(CPU4コア)を
使用するようにしています.

このグラフから、GPUでは1ノードあたりの性能が高いために、ノード数の増加に対するスコアの増加量も大きいことが
わかります.また、今回の測定においてはノード数の増加に伴ってスコアもほぼ直線的に増加しています.4ノードでの
スコアを比較するGPUが543Gflops、オリジナルが143Gflopsとなっており、GPUの方が約3.8倍高速になっています.
この性能差はGPUのハード性能などから考えると大きいとは言えませんが、本項で紹介したような簡単なCUDA
ポーティングのみでもこの程度の性能が出せると言うことができます

さらに、メモリパディングなどのグローバルメモリのサイズの調整、CUDAグリッドの構成、カーネルコードの
チューニング、カーネル実行とデータ転送のオーバーラップなど、GPUコードもしくはCPUコードに手を加えることによって、
GPU上での性能はここに示したものより大きく向上させることが可能です.

OTB製CUDAクラスター機におけるC姫野ベンチマークの結果例

(C)Copyright OTB Transnational Inc. 2017 All rights reserved.