HPC パフォーマンスの測定

2013.06.22
HPC
MICXeon Phiオフロード

この記事は、インテル® デベロッパー・ゾーンに掲載されている「Measuring performance in HPC」(http://software.intel.com/en-us/articles/measuring-performance-in-hpc) の日本語参考訳です。

この記事は、HPC (ハイパフォーマンス・コンピューティング) におけるインテル® Xeon Phi™ コプロセッサーについて紹介するシリーズの 1 つ目です。インテル® Xeon Phi™ コプロセッサーは、インテル® メニー・インテグレーテッド・コア (インテル® MIC) アーキテクチャーを採用した最初の商用製品です。この記事では、インテル® Xeon Phi™ アーキテクチャーの基本、プログラミング・モデル、およびマイクロベンチマーク用にパフォーマンス (サイクル数) を測定する方法を紹介します。

インテル® Xeon Phi™ コプロセッサー

インテル® Xeon Phi™ コプロセッサーは、インテル® MIC アーキテクチャーの最初の商用製品です。開発コード名は Knights Corner (KNC) で、プロトタイプである Knights Ferry (KNF) の後継にあたります。60 コアを搭載しており、固定クロック速度 1.053GHz で動作します。8GB GDDR5 メモリーを搭載しており、帯域幅は 320 GB/s です。命令キャッシュは 32KB、L1 キャッシュは 32KB です (それぞれ 8 ウェイ、64B キャッシュライン)。L2 キャッシュはコアあたり 512KB ですが、完全にコヒーレントなキャッシュと見なすことができ、総容量はコアあたりのサイズの合計と等しくなります。各コアに使用する情報をコピーすることで高速なローカルアクセスを実現したり、あるいは 1 つのコピーをすべてのコアで利用することでキャッシュの使用量を抑えることができます。L2 キャッシュには、命令とデータの両方が含まれます (L2 キャッシュも 8 ウェイ、64B キャッシュラインです)。

インテル® MIC の命令セットは特殊であることを理解しておく必要があります。x86 命令セットを基にしていますが、広い幅のベクトルユニットを利用する特殊なベクトル命令セットが含まれているため、SIMD プログラミングを効率良く行えます。また、FMA 命令にも対応しています。各スレッドは 1 サイクルおきに命令を実行するので、FMA 命令への最適化を試みると良いでしょう。

インテル® Xeon Phi™ コプロセッサーは、インオーダー・マイクロアーキテクチャーにおいて必然的に生じるレイテンシーを隠蔽するため、各コアでマルチスレッドを活用します。これは、動的実行エンジンを最大限に利用することを目的とするインテル® Xeon® プロセッサーのハイパースレッディングとは異なります。HPC のワークロードでは、多くの場合、ハイパースレッディングは無視されたり、パフォーマンスを低下させないにもかかわらず無効にされます。これは、インテル® Xeon Phi™ コプロセッサーのハードウェア・スレッドには当てはまりません。インテル® Xeon Phi™ コプロセッサーでは、プログラムのマルチスレッド化は無視されるべきではありません。また、ハードウェア・スレッドを無効にすることはできません。

インテル® Xeon Phi™ コプロセッサーは、コアあたり 4 つのハードウェア・スレッドを提供し、コアあたり 1 スレッドでは一般に上限に達することがない十分なメモリー帯域と浮動小数点処理能力を備えています。高度にチューニングされたコードのカーネルでは、2 つのスレッドで飽和状態になる場合がありますが、一般的なアプリケーションでコプロセッサーのすべてのリソースを使い果たすには、最低でもコアあたり 3 ～ 4 つのアクティブなスレッドが必要です。そのため、コアあたりの使用するスレッド数は、アプリケーションでチューニング可能なパラメーターとし、アプリケーションの実行結果に基づいて設定すると良いでしょう。

インテル® Xeon Phi™ コプロセッサーのプログラミング

私たちは、ホストシステムでインテル® Xeon® プロセッサーのプログラミング方法を理解しており、どのようにアプリケーションでインテル® Xeon Phi™ コプロセッサーを活用するかが課題となります。次の 2 つの主要なアプローチがあります。

  • ホスト・プロセッサー中心の「オフロード」モデルは、プログラムをホスト・プロセッサーで実行し、選択した作業をコプロセッサーにオフロードします。
  • 「ネイティブ」モデルは、プログラムをホスト・プロセッサーとコプロセッサーでネイティブに実行し、ホスト・プロセッサーとコプロセッサーはさまざまな方法で通信します。

MPI プログラムは、どちらのモデルも適用できます。例えば、ホスト・プロセッサー側でのみランク付けされたプログラムは、オフロードすることでコプロセッサーのパフォーマンスを利用でき、ホスト・プロセッサーとコプロセッサーの両方でランク付けされているプログラムは、どちらでもネイティブに実行できます。いずれのモードも、実際にはマシンの「モード」ではなく、1 つのプログラムに混在させることができるプログラミング・スタイルです。

オフロードは、一般に細粒度の並列化に使用されるため、プログラムの局所的な変更を伴います。MPI は、粗粒度で行われることが多く、MPI 呼び出しを追加するため、プログラムのさまざまな場所で変更が必要になります。インテル® MPI は、ホスト・プロセッサーおよびコプロセッサー向けにチューニングされているため、リモート・ダイレクト・メモリー・アクセス (RDMA) のようなハードウェア機能を利用できます。

最初に、「オフロード」モデルについて見てみましょう。

オフロードモデルを使用するプログラミング

インテル® Xeon Phi™ コプロセッサー向けのオフロードモデルは充実しています。オフロード言語拡張の構文とセマンティクスには、OpenACC などその他のオフロードモデルにはない機能が含まれています (OpenACC は GPU との互換性により制限されます)。これによって、より優れた OpenMP* との互換性が得られます。また、複数のインテル® Xeon Phi™ コプロセッサー・カードを制御したり、GPU では処理できない複雑なプログラム・コンポーネントをオフロードすることができます。

void doMult(int size, float (* restrict A)[size], float (* restrict B)[size], float (* restrict C)[size]) 
{
#pragma offload target(mic:MIC_DEV) 
                in(A:length(size*size)) in( B:length(size*size))    
                out(C:length(size*size))
  {
    // C 行列をゼロにする
#pragma omp parallel for default(none) shared(C,size)
    for (int i = 0; i < size; ++i)
      for (int j = 0; j < size; ++j)
        C[i][j] =0.f;
     
    // 行列の乗算を計算する
#pragma omp parallel for default(none) shared(A,B,C,size)
    for (int i = 0; i < size; ++i)
      for (int k = 0; k < size; ++k)
        for (int j = 0; j < size; ++j)
          C[i][j] += A[i][k] * B[k][j];
  }
}

上記のように、offload プラグマは、上記のように、offload プラグマは、インテル® Xeon Phi™ コプロセッサーとの間で正しくデータを転送できるように、コンパイラーに注釈を与えます。offload 指示句のスコープには、複数の OpenMP* ループを含めることができます。次に、ほかの節について説明します。

offload プラグマには、ターゲットデバイスへのオフロードに関連する情報を含む節を指定します。ここでは、target(mic:MIC_DEV) が target 節で、ホスト・プロセッサーと指定したオフロードデバイス向けの両方のコードを生成するようにコンパイラーに指示しています。この例では、定数 MIC_DEV で指定された番号に関連付けられたインテル® Xeon Phi™ コプロセッサー・カードがターゲットになります。

in(var-list modifiersopt) 節は、ホストからコプロセッサーへ明示的にデータをコピーすることを指示します。デフォルトでは、デバイスでメモリーが割り当てられ、指示句の範囲を出るときに解放されます。alloc_if(condition) 修飾子と free_if(condition) 修飾子は、この動作を変更できます。

out(var-list modifiersopt) 節は、コプロセッサーからホストへ明示的にデータをコピーすることを指示します。ここでも、デフォルトでは、指示句の範囲を出るときに指定されたメモリーが解放されます。free_if(condition) 修飾子を使用して、このデフォルトの動作を変更できます。

最後に、「ネイティブ」モデルについても見てみましょう。

ネイティブモデルを使用するプログラミング

利用可能なプログラミング・モデルは 3 つあります。すでに簡単に説明したとおり、オフロードモデルは、ホスト・プロセッサーでアプリケーションを実行し、ソースコードでプラグマを用いて、計算負荷の高いコードと関連データをデバイスにオフロードします。2 番目の方法として、ホストで従来の OpenMP* アプリケーションとして実行することができます。これは、インテル® Xeon Phi™ コプロセッサー・カードを必要としないため、ここでは取り上げません。これとは逆のモデルは、ネイティブ・プログラミング・モデルと呼ばれます。このモデルでは、アプリケーション全体がインテル® Xeon Phi™ コプロセッサー・カードで実行されます。

ネイティブ・プログラミングは、メニーコア・プロセッサー向けのプログラムの記述と考えることができます。インテル® Xeon Phi™ コプロセッサー・カードであるかどうかは重要ではありません。ただ 1 つの違いは、最後にインテル® MIC 向けにコードをコンパイルして、デバイスで実行することです。これは手動で行うことも (SCP を使用してすべてを転送し、SSH で接続して実行します)、あるいは micnativeloadex のようなユーティリティーで行うこともできます。インテル® Xeon Phi™ コプロセッサーでもプログラムの依存関係は解決する必要があります。つまり、プログラムが正しく実行されるように、環境変数を適切に設定します。

#include <stdio.h>

void say_hello()
{
  #ifdef __MIC__
    printf("Hello, I am MIC!n");
  #else
    printf("We are still on the host!n");
  #endif
}

int main(int argc, char **argv)
{
  say_hello();
  return 0;
}

このコードをコンパイルするには、”icc -mmic -o hello-mic hello-mic.c” を実行します。micnativeloadex ユーティリティーで実行する場合は、”micnativeloadex hello-mic” を利用します。

オフロード・アプリケーションのビルドと異なり、ここではコンパイルに -mmic オプションを指定しています。オフロード・アプリケーションでは、代わりに offload-build オプションを指定してコンパイルする必要があります。また、オフロード・プログラミング・モデルを使用するプログラムには、前述したプラグマが含まれます。

インテル® Xeon Phi™ コプロセッサー向けプログラミングについては、このシリーズの次の記事で詳しく述べます。

パフォーマンスの測定

時間を秒、ミリ秒、あるいはマイクロ秒で測定する方法はいくつかありますが、サイクル数 (またはナノ秒) で測定するのはそれほど簡単ではありません。アセンブリー・コードで命令数をカウントすることは可能ですが、各命令のサイクル数を把握しておかなければならず、オーバーヘッドは考慮されません。また、理論上の計算値だけではなく、実際のデータも必要になるため、サイクル単位の精度で測定しなければなりません。

幸いにも、すべての x86 プロセッサーは、タイム・スタンプ・カウンター (TSC) と呼ばれる 64 ビット・レジスターを備えています。TSC は、リセットからのサイクル数をカウントします。対応するアセンブリー命令は RDTSC (RD は読み込みを意味する) と呼ばれ、TSC の値を EDX:EAX に格納して返します。

これまで、RDTSC は、CPU のタイミング情報を得るための、高分解能に優れた低オーバーヘッドな方法でした。しかし、マルチコア/ハイパースレッド CPU、マルチ CPU システム、およびオペレーティング・システムの休止状態の登場により、一般に TSC から正確な結果を得ることはできなくなりました。1 つのマザーボードにマルチ CPU が搭載されている場合、細心の注意を払っても RDTSC が同期される保証はありません。

インテル® MIC の場合、すべてのカウンターはほぼ同じになると想定できます。クロックレートも同じです。大きな問題は、最近の CPU はアウトオブオーダー実行をサポートしていることです。そのため、命令が必ずしも実行ファイルの順序で実行されるとは限りません。この問題に対処するには RDTSCP (P は「およびプロセッサー ID」を意味する) と呼ばれる命令を利用します。これは、このバージョンがシリアル化されていること、つまり、順序が保証されていることを示します。

残念ながら、インテル® MIC は RDTSCP 命令をサポートしていません。ただし、このコードは簡単に記述することができます。各コアのカウンターのオフセットは異なるかもしれませんが、これは一定のまま変わらないので (ほとんどの測定では) 無視するか、各カウンターのオフセットがなくなるように測定を調整することができます。

C/C++ で次の関数を作成します。

static inline unsigned long rdtsc()
{
	unsigned int hi, lo;

	__asm volatile (
		"xorl %%eax, %%eax nt"
		"cpuid             nt"
		"rdtsc             nt"
		:"=a"(lo), "=d"(hi)
		:
		:"%ebx", "%ecx"
	);
	return ((unsigned long)hi << 32) | lo;
}

これは、プロセッサー ID も読み込んでいるため、RDTSCP 命令のシリアルバージョンと言えます。最初に、このコードを使って測定の精度/オーバーヘッドを測定します。

static int tsc_overhead()
{
    unsigned long t0, t1;
    t0 = rdtsc();
    t1 = rdtsc();
    return (int)(t1 - t0);
}

通常、オーバーヘッドは O(100) サイクルまたは約 0.1 µs です (クロック周波数により異なります)。そのため、1 サイクルの命令であっても、少なくとも O(100) サイクルは測定しなければなりません。

結論

この記事では、インテル® Xeon Phi™ コプロセッサーの基本アーキテクチャーと利用可能なプログラミング･モデルについて紹介しました。また、実行時間の測定についても簡単に説明しました。ハイパフォーマンス・カウンターを利用する場合であっても、いくつかのルールを考慮する必要があります。有効な数値を得るためには、合理的な小さな命令数の繰り返しを複数回測定すべきです。また、各コアごとに一定のオフセットがある可能性も考慮する必要があります。

参考文献 (英語)

  • インテル® Xeon Phi™ コプロセッサーの概要 (https://www-ssl.intel.com/content/www/us/en/processors/xeon/xeon-phi-detail.html?)
  • ウィキペディアのタイム・スタンプ・カウンター
  • インテル® Xeon Phi™ コプロセッサーのブロック・ダイアグラム (https://www-ssl.intel.com/content/www/us/en/processors/xeon/xeon-phi-coprocessor-block-diagram.html)
  • インテル® Xeon Phi™ コプロセッサーの仕様 (http://ark.intel.com/products/71992/Intel-Xeon-Phi-Coprocessor-5110P-8GB-1_053-GHz-60-core)
  • 最適化とパフォーマンス・チューニング (http://software.intel.com/en-us/articles/optimization-and-performance-tuning-for-intel-xeon-phi-coprocessors-part-2-understanding)
  • オフロードモデル (http://software.intel.com/sites/products/documentation/doclib/stdxe/2013/composerxe/compiler/cpp-lin/GUID-EAB414FD-40C6-4054-B094-0BA70824E2A2.htm)

コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

タイトルとURLをコピーしました