• http://gpu-computing.gsic.titech.ac.jp/node/63
  • 主催: GPUコンピューティング研究会
  • 日時: 2012年3月5日(月)13:30〜 (雨)
  • 場所: 東京工業大学(大岡山) 情報ネットワーク演習室 第2演習室

99445e6d.jpg4e244804.jpg9900a9c5.jpg

CUDAはまださわったことがないけど、ディレクティブベースならやってもいいかなとおもって勉強しにいってきた。

会場の演習室はWUXGA(1920x1200)のディスプレイがずらっと並んでいてWindows 7 enterprise。講師の画面のミラーが2人に1台ついてた。豪華だなぁ。飲食禁止。

c8cef21e.jpgb75acfb0.jpg

ひさしぶりに松田さんに合った。HPCをやるのことになりそうなので情報収集に来ているとのこと。

ゴードンベル受賞(2011)研究の紹介 / 青木尊之(東京工業大学・学術国際情報センター)

  • ゴードンベル賞というのはスパコンをどれだけ応用したかに対する賞。
  • 気象計算では賞をとれないとおもっていた。賞をねらうためにはフェーズフィールド法をつかった。効率がいい。
  • 凝固して成長する仮定をシミュレーションするのは2次元がやっとで3次元は1本がやっと(それでも1週間かかってしまったり)
  • 数mmのオーダでシミュレーションしないと特性がわからない。
  • TSUBANE2.0 4node 1ヶ月 10万円しないらしい Amazonより安いよ。
  • GPUが90%の性能をもっている。京は1000億円 vs TSUBAME2.0 32億円。
  • フェーズフィールドφ: 0が液体1が固体とか
  • 隣接格子なのでメモリ律速のようにみえて、実はキャッシュに乗るのでCPU律速になる。非線形の計算なので計算負荷が重い→性能がでやすい。
  • 無理にshared memoryをつかう必要はない。勝手にキャッシュにのる。
  • あえて2次元分割
  • 今はGPU間の転送ができないので一旦CPUに転送してMPI通信。
  • 境界の計算を先にやって通信しつつ、オーバラップして中の計算。。
    • 境界はCPUでやって(GPUで計算すうとメモリコピー必要になる)、中はGPU。
  • 64bitから32bitにしたら早くなったのはアドレス計算の関係。
  • weak scaling 2PFlopsでた。44.5%の実行性能。

OpenACC の紹介 / 丸山直也(東京工業大学・学術国際情報センター)

  • OpenACCのコンパイラを入手できていないので細かな点でまちがいがあるかも
  • PGIとかCrayとかのほかに、CAPS:フランスの会社もつくっている。
  • PGI Accelerator Compierがベースでよく似ている。対象言語はC/C++,Fortran。
  • OpenMPにとりこまれることを目的にしていたが、なかなか大変なので先にOpenACCとしてリリース。
  • Crayのコンパイラをつくっている人が... (ききもらした;だれかおしえて)
  • ポインタがあると並列化がむつかしい。
  • 書き方は ホストプログラム + (GPU)カーネル関数
  • OpenACCではCPUとGPUのメモリは分離されているように見せている。(共有メモリモデルではない)
  • メモリコピーは自動でやってくれる。無駄なコピーを省くのはチューニング項目。
  • CUDAだとすべて明示的に書かないといけない。
  • OpenACC用語: SIMD(thread?;CPUのSIMD演算に対応) < worker(thread?;CPUのcoreに対応) < gang(thread block;CPUのsocketに対応)
    • (ちがってるかも)
  • #pragma acc cache(list)でshared memoryに置く指令(ヒント)
  • NVIDIAもOpenACCに注力しているよう。
  • windowsではどう?

HMPPの概要説明 / 小野寺高之(株式会社JCC ギミック)

  • CAPS: フランスの会社。7人の研究者。10年の歴史がある。
  • HMPP: 最低2行のディレクティブでGPUで実行できるようになる。(性能は別にして)
  • IntelのMICにも対応予定。
  • source-to-sourceコンパイラ: CPUコードのコンパイラそのままつかえる。
    • hmpp gcc myProgram.c
  • pure functionの呼出をGPUにオフロードできる。基本は関数単位で。
  • #pragma hmpp LABEL codelet で関数定義
  • #pragma hmpp LABEL callsiteで関数呼出
  • codelet,target=CでCのコードを生成できる。
  • GPUが利用できない場合にはCPUにフォールバックするが環境変数HMPPRT_NO_FALLBACK=1を定義すると止まるようになる。
  • HMPPRT_LOG_LEVEL=INFOでメモリアロケートとかも表示する。(デバッグ用)
  • aquire/release: ハードウェアの確保/解放
  • allocate/free: コードレットのメモリ割当/解放
  • advancedload/delegatedstore: データ転送
  • コードレットグループを指定すると同じハードウェアを使うようにできる。
  • MAPを指定すると同じ変数であるとヒントを与えられる。
  • リファレンスマニュアル: 英語のみ
  • チュートリアル: スライドなど(2.xは日本語あり、3.xは日本語化中)。スライド以上のものは、いまのところない。
  • プログラミングマニュアル:
  • OepnACCのサポート: 3月中に出るとアナウンスしているが。
  • OpenACCは共通サブセットになっていてPGIのモデルが基本になっていて、他のベンダは拡張機能として入れる。

HMPP によるプログラムの高速化実習

  • TSUBAME2にログインして用意されたソースコードにpragma directiveを入れているだけの簡単なおしごと..
  • 受付でログインアカウントが書いてある紙をもらっているので
  • windowsにログインして
  • cygwinを起動して
  • sshでTSUBAME2にログインして
  • 環境を設定するスクリプトを実行して
  • エディタでpragmaを入れていくのだが、ちゃんとソースコード上に「ここにかいてね」マークがあるので間違うことはないが
  • 入れるディレクティブが長〜いので打ち間違えリスクが高い。
  • 細かくチューニングするために変数を列挙しないといけないのが面倒臭い。
  • でもって実行しようとするとエラーになって動かない。どうやら皆が一斉に実行するのでGPUのキャパを越えてしまうらしい。
  • hmpp --codelet-requiredは便利なオプション。GPUコードが生成できなかったらコンパイルエラーにする。
  • あまりに単調な作業だったので途中で抜けた。(だって頭が働いてないんだもん)

44b96e2c.jpg