編集：私の最初のコードサンプルは間違っていました。よりシンプルに修正されました。

大きなベクトルと行列の間の代数演算用にC++ライブラリを実装します。x86-x64 CPUで、OpenMP並列ベクトルの追加、ドット積などがシングルスレッドよりも高速ではないことがわかりました。並列操作は、シングルスレッドよりも-1％〜6％高速です。これは、メモリ帯域幅の制限が原因で発生します（私は思います）。

したがって、問題は、次のようなコードに実際のパフォーマンス上の利点があるかどうかです。

このサンプルでは、​​OpenMPを使用することによる重大な欠点はありません。しかし、スパースベクトルとスパース行列を使用してOpenMPで作業している場合、たとえば* .push_back（）を使用することはできません。その場合、質問は深刻になります。（スパースベクトルの要素はデンスベクトルのように連続的ではないため、並列プログラミングには欠点があります。これは、結果要素がいつでも到着する可能性があるためです。インデックスが低いものから高いものへの変換ではありません）

最新のマルチコア プラットフォームでは、メモリ帯域幅が制限されたアプリケーションの並列パフォーマンスは、多くの場合、コアの数に応じて適切にスケーリングされません。通常、ある程度のコア数まで高速化が見られますが、それ以降は性能が飽和します。総合的な例として、有名なSTREAM ベンチマークがあります。これは、達成可能なメモリ帯域幅、つまり飽和点でのメモリ帯域幅を報告するためによく使用されます。

ピーク メモリ帯域幅が 42.7GB/s (DDR3-1333) の単一の Xeon E5-2680 での STREAM ベンチマーク (Triad) の次の結果を考慮してください。

STREAM は 1 コアから 2 コアまで適切にスケーリングしますが、3 コアを超えると、パフォーマンスはほぼ一定になります。

私の質問は、単一の CPU コアで達成できるメモリ帯域幅を決定するものは何ですか? この質問は明らかに範囲が広すぎるため、上記のアーキテクチャに絞り込みます: E5-2680 の仕様から、またはハードウェア カウンターなどを見て、1 スレッドの STREAM で 16 GB/s が得られることをどのように予測できますか? ?

私の CPU の仕様によると、メモリへの帯域幅は 5.336GB/s になるはずです。これをテストするために、大きな配列で memset (または memcpy) を実行し、タイミングを報告する簡単なプログラムを作成しました。memset で 3.8GB/s、memcpy で 1.9GB/s を示しています。 http://en.wikipedia.org/wiki/Intel_Core_(microarchitecture)によると、私の Q9400 は 5.336MB/s になるはずです。どうしたの？

memset または memcpy を割り当てループに置き換えてみました。私は、メモリの配置について学ぶためにグーグルで検索しました。さまざまなコンパイラ フラグを試しました。私はこれに恥ずかしい時間を費やしました。ご協力いただきありがとうございます。

libc-dev バージョン 2.15-0ubuntu10.5 およびカーネル 3.8.0-37-generic で Ubuntu 12.04 を使用しています

コード：

コンパイル コマンド:

サンプル出力:

lshwによると私のハードウェア：

これら 2 つのブロック サイズ (1024x1 と 32x32) は、スレッド スケジューリングとメモリ帯域幅の観点からどのように実行されると予想されますか? これら 2 つのブロック サイズのパフォーマンスに予想される違いはありますか? どちらもブロックごとに 1024 スレッドを使用することに注意してください。

CUDA カーネルがメモリ帯域幅の使用率に関してどの程度優れているかを知りたいです。ECCをオンにしたTesla K40cでそれらを実行します。ユーティリティによって得られた結果はbandwidthTest、達成可能なピークの適切な近似値ですか? そうでなければ、ピーク帯域幅を見つけるために同様のテストを作成するにはどうすればよいでしょうか?

デバイスのメモリ帯域幅を意味します。

タイトルで述べたように、 の関数はcudaMallocPitch()多くの時間をcudaMemcpy2D()消費し、かなりの時間を消費することもわかりました。

私が使用しているコードは次のとおりです。

実装では、Size.widthとSize.heightは両方とも 4800 です。所要時間cudaMallocPitch()は約 150 ～ 160ms (事故の場合は複数のテスト) でcudaMemcpy2D()、約 50ms を消費します。

CPU と GPU 間のメモリ帯域幅がそれほど制限されているとは思えませんが、コードにエラーは見られません。その理由は何ですか?

ちなみに、私が使用しているハードウェアは Intel I7-4770K CPU と Nvidia Geforce GTX 780 (エラーのないかなり良いハードウェア) です。

クラスターの 1 つのノードでSTREAM ベンチマークのMPI バージョンを実行して、さまざまな数の MPI プロセスの持続可能な帯域幅を測定したいと考えています。各ノードは、2 つのインテル® Xeon® プロセッサー E5-2680 v3 (12 コア) で構成されています。

以下では、openmpi v. 1.8.2 とオプションmap-by core. ソースはラッパーを使用してコンパイルされました

コンパイラ オプションを使用して

各コアは長さ 2*10^6 の倍精度配列を使用しており、テストは 50 回繰り返されます。

私が困惑しているのは、2 ～ 6 プロセスの測定された持続可能な帯域幅の停滞です。使用されているプロセッサのターボ ブーストが結果に偏りを与えている可能性があると思います。どうやらTurbo Boostはいくつかのコアが使用されている場合にのみアクティブになりますが、結果を正しく解釈する方法がわかりません.

Turbo Boost を「オフ」にするには、次のように STREAM ベンチマークを変更することが考えられます。

• 使用可能な 24 個のコアすべてを常に使用して、ベンチマーク全体でプロセッサの負荷をほぼ一定に維持します
• ベンチマークの MPI コミュニケーターMPI_COMM_WORLDを 2 つの別々のものに分割します。1 つのコミュニケーターは、実際にベンチマークを実行するN 個のプロセスのグループに関連付けられます。もう 1 つのグループには、残りの24-N 個のプロセスが含まれます。これらのプロセスは、コアをビジー状態に保ち、ターボ ブーストを防止することのみを目的としています。
• 24-Nプロセスは、L1 キャッシュからのデータのみを使用して、他のプロセスへの影響を最小限に抑えます。
• ベンチマークの実行時間を数分に延長します (ターボ ブーストの効果は数秒以上持続しませんか?)

最初の実装では、L1 キャッシュ (32kByte) に収まる必要があるデータ セットの残りの24-Nプロセスにトライアド カーネルのみを使用します。ベンチマークの繰り返し回数は 5000 に引き上げられ、データ配列の長さのサイズは 1*10^7 に引き上げられました。この関数fooは、ループのアンロールを防止するために (成功したか?) 呼び出されます。先行する if 条件はfalse、プログラムの実行全体を通じて行われます。関連するコードは次のとおりです

a、b、および c はstatic double a[LEN],b[LEN],c[LEN]、main の外部として宣言されています。また、 static 修飾子とそのベンチマークへの影響に関するこの議論も読みました。通話fooが本当に必要かどうかはわかりません。

残念ながら、私のアプローチでは、ベンチマークとは異なる結果にはなりませんでした。

私のアプローチの何が問題になっていますか？代替手段について知っていますか、またはベンチマークをどのように変更する必要があるかについての提案はありますか?

編集：

あるユーザーが私がベンチマークに使用したコードに興味を示したので:

PETSc に付属する STREAM ベンチマークの修正版を使用しました。

http://www.mcs.anl.gov/petsc/

システムに PETSc をインストールすると、次の場所にベンチマークが表示されます。src/benchmarks/streams

ソースコードのみに興味がある場合は、GitHub ページで最新バージョンを見つけることができます。

https://github.com/petsc/petsc/tree/master/src/benchmarks/streams

x86 movdqu 命令 (movdqu 命令は、xmm- との間の 16 バイト データの移動をサポートします) を使用して、あるソース メモリ ロケーションからある宛先メモリ ロケーションに反復ごとに 64 バイトを移動するループで見られる結果の説明を求めています。整列されていないメモリ位置との間のレジスタ)。これは、memcpy()/java.lang.System.arraycopy() に似た関数を実装するコードの一部です。

コピーを実装しようとした 2 つの異なるパターンがあります。

パターン1

このパターンでは、rsi はソース (src) アドレスを保持し、rdi は宛先 (dst) アドレスを保持し、xmm レジスターは一時レジスターとして使用されます。このコードは、copylen_in_bytes/64 回繰り返されます。ご覧のとおり、ここでは ld-ld-st-st-ld-ld-st-st ロード-ストア パターンに従います。

パターン2

pattern2 では、ld-st-ld-st-ld-st-ld-st パターンに従います。

観察

このコードを数百回実行すると、src と dst がさまざまな 8 バイト境界で整列され、次のことがわかります。

Westmere の場合 (Xeon X5690)

パターン 1 は、実行ごとに非常に高い分散を示します。

パターン 2 はほとんど分散を示しません。

パターン 2 の最小時間 (観測された最速時間) は、パターン 1 の最小時間よりも (~8%) 長くなります。

Ivybridge で (Xean E5-2697 v2)

パターン 1 は、実行ごとに非常に高い分散を示します。

パターン 2 はほとんど分散を示しません。

パターン 2 の最小時間は、パターン 1 の最小時間よりも高くなっています (~20%)。

ハスウェル (Core i7-4770)

パターン 1 は、非常に高い実行ごとの分散を示しません。

パターン 2 はほとんど分散を示しません。

パターン 2 の最小時間は、パターン 1 の最小時間よりも高くなっています (~20%)。

奇妙なことに、Westmere と Ivybridge では、src/dest のアラインメントと悪い結果 (高い分散の原因) との間に相関関係がないようです。同じ src/dest アラインメントに対して良い数値と悪い数値が表示されます。

質問

キャッシュラインにまたがる movdqu のパフォーマンスが、キャッシュラインにまたがらない movdqu よりも悪いことは理解していますが、次のことはわかりません。

1) パターン 1 が Westmere と Ivybridge で高い分散を示すのはなぜですか? ロードストアの順序はどのように違いますか?

2) 異なるアーキテクチャ間で、パターン 2 の最小時間がパターン 1 よりも遅いのはなぜですか?

この長い投稿をお読みいただきありがとうございます。

カーシック

SDRAM はロードとストアを同時に実行できますか?

もしそうなら、保存と読み込みを同時に行うと帯域幅にどのような影響がありますか?

• 両方を実行すると、ストア帯域幅とロード帯域幅の合計としての全体的な帯域幅が低下しますか?
• 全体の合計は同じですが、ロードとストアで分割されていますか?
• または、記載されている帯域幅全体を保存に使用し、記載されている帯域幅全体を保存に使用できますか??