メニーコアと格闘ブログ:CUDAで実装熱伝導解析(その1)とりあえずGPUで計算する改造

出典: トータル・ディスクロージャ・サイト(事実をありのままに)

[デバイスとホスト間のデータコピー]

さて、CUDA実装でもMPI実装と同じく、並列計算の肝になるのは“solve関数”です。 solve関数内で必要になる変数はT,T_newです。T_newはTの値を基に上書きされてしまうので、ここで重要になる変数はTということになります。

まずは初期値Tをデバイスに“送る”、“戻す”のやり方から実装してみます。デバイスで起動する全てのスレッドからアクセスできる“グローバルメモリ”とデータを交換します。

※ホストからはデバイスの“グローバルメモリ”しかアクセスできません。



ここではまず、

  • Tに初期値を代入
  • 熱源設定

を終えた後、デバイス上のglobalメモリに初期配列を転送します。

   // デバイス側メモリを確保(ここではホスト側T,T_newに対してT_d、T_new_dを動的に作成)
   cudaMalloc((void**)&T_d,sizeof(double)*NH*NW);
   cudaMalloc((void**)&T_new_d,sizeof(double)*NH*NW);

   // デバイスへデータを転送
   // 配列Tの内容1列分をtmpへ確保して、
   // デバイスへコピーするループを、配列の列数分だけ繰り返す
   for (int jdex = 0; jdex < NW; jdex++) {
       double* tmp = T[jdex];
       cudaMemcpy(((char*)T_d + sizeof(double) * jdex * NH), tmp, sizeof(double) * NH, cudaMemcpyHostToDevice);

       tmp = T_new[jdex];
       cudaMemcpy(((char*)T_new_d + sizeof(double) * jdex * NH), tmp, sizeof(double) * NH, cudaMemcpyHostToDevice);
   }

ホスト、デバイス間でデータをコピーするAPIはcudaMemcpyです。フォーマットは、以下のようになります。

      cudaMemcpy(宛先アドレス、送元アドレス、サイズ、どこからどこへ)

“どこからどこへ”の部分には、以下のいずれかが入ります。

  • cudaMemcpyHostToDevice : ホストからデバイス
  • cudaMemcpyDeviceToHost:デバイスからホスト
  • cudaMemcpyDeviceToDevice:デバイス同士のメモリコピー


デバイス側の1次元配列に対して、ホスト側の2次元配列内のデータをコピーしています。

C言語では、配列をarr[a][b]と宣言すると、メモリ上の情報はbが先に動くように並び、デバイス側に確保した1次元配列には、図のようにデータが並ぶことになります。

※CUDAのAPIには2次元配列を確保し、データ送受信もあるようですが、使用方法を現在調査中です。


ファイル:CUDAで実装熱伝導解析(その1)とりあえずGPUで計算する改造01.jpg


次に、デバイスからホストにデータをコピーする方法ですが、単純に上の逆をたどることになります。

// まず、ホスト側にデータ受け入れ領域を確保しておく
double T_res[NW*NH];
// cudaMemcpy()を使ってT_dに入っているデータをT_resにコピー。
// ここでは、1次元配列のものを1元配列にコピーすることを考えたので、
// NHxNWのサイズを指定。コピー終了後、ホスト側で、2次元配列に戻す。
cudaMemcpy(&T_res, T_d, sizeof(double)*NH*NW, cudaMemcpyDeviceToHost); 

for(i=0;i<NW;i++){
    for(j=0;j<NH;j++){
        T[i][j]=T_res[i*NH+j]; 
    }
}

効率悪いですが、これでデバイスとデータの交換ができるようになりました。


ナン 2009年8月28日 (金) 11:44 (JST)


この記事へのコメントをお寄せください

  • サイトへの書き込みに差し支えございましたら トータルディスクロージャーサイトサポート係へメールをお送りください
  • トータル・ディスクロージャ・サイトに投稿された文章と画像は、すべてその著作権がHPCシステムズ株式会社に帰属し、HPCシステムズ株式会社が著作権を所有することに同意してください。
  • あなたの文章が他人によって自由に編集、配布されることを望まない場合は、投稿を控えてください。
  • コメントを書き込む場合は名前にひらがなを織り交ぜてください。
  • あなたの投稿する文章と画像はあなた自身によって書かれたものであるか、パブリック・ドメインかそれに類する自由なリソースからの複製であることを約束してください。あなたが著作権を保持していない作品を許諾なしに投稿してはいけません!

<comments hideform="false" />


Comments

ノート:メニーコアと格闘ブログ:CUDAで実装熱伝導解析(その1)とりあえずGPUで計算する改造

個人用ツール