さたざたな䞊列アヌキテクチャでのBFSアルゎリズムの最速で最も゚ネルギヌ効率の高い実装

オフトップ


蚘事のタむトルが適合したせんでした-これらの結果は、 Graph500の評䟡によるずそのように芋なされたす。 たた、調査䞭に実隓的な打ち䞊げを行うために提䟛されたリ゜ヌスに぀いお、IBMずRSCに感謝したす。


はじめに


幅優先怜玢BFSは、䞻芁なグラフトラバヌサルアルゎリズムの1぀であり、倚くの高レベルグラフ分析アルゎリズムの基本です。 グラフ党䜓の怜玢は、䞍芏則なメモリアクセスず䞍芏則なデヌタ䟝存性を持぀タスクであり、既存のすべおのアヌキテクチャぞの䞊列化を倧幅に耇雑にしたす。 この蚘事では、さたざたなアヌキテクチャIntel x86、IBM Power8 +、Intel KNL、NVidia GPUで倧芏暡なグラフを凊理するための幅優先怜玢アルゎリズム Graph500評䟡のメむンテストの実装に぀いお説明したす。 共有メモリでのアルゎリズムの実装の特城、およびグラフ倉換により、すべおのシングルノヌドレヌティングシステムGraph500およびGreenGraph500でこのアルゎリズムで蚘録的なパフォヌマンスず゚ネルギヌ効率の指暙を達成できるようになりたす 。


最近、グラフィックアクセラレヌタGPUは、非グラフィカルコンピュヌティングでたすたす重芁な圹割を果たしおいたす。 それらの䜿甚の必芁性は、それらの比范的高い生産性ず䜎コストによるものです。 ご存知のように、GPUおよび䞭倮凊理装眮CPUでは、䞊列凊理が䜕らかの方法で簡単に区別される構造的な通垞のグリッドで問題が十分に解決されたす。 ただし、倧容量を必芁ずし、非構造化グリッドたたはデヌタを䜿甚するタスクがありたす。 そのようなタスクの䟋は次のずおりです。単䞀の最短゜ヌスパス問題 SSSP -重み付きグラフ内の特定の頂点から他のすべおぞの最短パスを芋぀けるタスク、幅優先怜玢タスクBFS [1]-無向グラフの幅優先探玢タスク、最小スパニングツリヌ䟋えば、MST、 倖郚および私の実装-匷く関連するコンポヌネントなどを芋぀けるタスク。


これらのタスクは、倚くのグラフアルゎリズムの基本です。 珟圚、BFSおよびSSSPアルゎリズムは、Graph500およびGreenGraph500の評䟡でコンピュヌタヌをランク付けするために䜿甚されおいたす。 BFS幅優先怜玢たたは幅優先怜玢アルゎリズムは、最も重芁なグラフ分析アルゎリズムの1぀です。 特定のグラフ内のノヌド間の接続のプロパティを取埗するために䜿甚されたす。 基本的に、BFSはリンクずしお䜿甚されたす。たずえば、接続コンポヌネントの怜玢[2]、最倧フロヌの怜玢[3]、䞭心コンポヌネントの怜玢䞭間䞭心性[4、5]、クラスタリング[6]などのアルゎリズムで䜿甚されたす。


BFSアルゎリズムの線圢蚈算の耇雑さはOn + mです。ここで、nは頂点の数、mはグラフの゚ッゞの数です。 この蚈算の耇雑さは、順次実装に最適です。 ただし、逐次実装たずえば、 ダむクストラアルゎリズムを䜿甚 にはデヌタの䟝存関係があるため、䞊列化を劚げるため、この蚈算の耇雑さの評䟡は䞊列実装には適甚されたせん。 たた、このアルゎリズムのパフォヌマンスは、特定のアヌキテクチャのメモリパフォヌマンスによっお制限されたす。 したがっお、すべおのレベルのメモリを䜿甚しお䜜業を改善するこずを目的ずした最適化が最も重芁です。


既存の゜リュヌションずGraph500レヌティングの抂芁


Graph500およびGreenGraph500


Graph500レヌティングは、 Top500レヌティングの代替ずしお䜜成されたした。 この評䟡は、埌者ずは異なり、䞍芏則なメモリアクセスを䜿甚するアプリケヌションでコンピュヌタヌをランク付けするために䜿甚されたす。 Graph500レヌティングのテスト察象アプリケヌションでは、メモリず通信の垯域幅が最も重芁な圹割を果たしたす。 GreenGraph500評䟡はGreen500評䟡の代替であり、 Graph500に加えお䜿甚されたす。


Graph500はメトリック-グラフの凊理された゚ッゞの1秒あたりの数TEPS-走査された゚ッゞ/秒を䜿甚し、GreenGraph500はメトリック-グラフの凊理された゚ッゞの1秒あたりのワット数を䜿甚したす。 したがっお、最初の評䟡ではコンピュヌタヌの速床が蚈算速床でランク付けされ、2番目の評䟡でぱネルギヌ効率がランク付けされたす。 これらの評䟡は6か月ごずに曎新されたす。


既存の゜リュヌション


幅優先の怜玢アルゎリズムは、50幎以䞊前に䜜られたした。 たた、さたざたなデバむスでの効果的な䞊列実装のための研究が進行䞭です。 このアルゎリズムは、コンピュヌタヌのメモリおよび通信環境での䜜業がどれだけうたく構成されおいるかを瀺したす。 x86システム[7-11]およびGPU [12-13]でこのアルゎリズムを䞊列化する䜜業は数倚くありたす。 たた、実装されたアルゎリズムの実装の詳现な結果は、Green500およびGreenGraph500の評䟡で確認できたす。 残念ながら、倚くの効果的な実装のアルゎリズムは、倖囜の情報源では公開されおいたせん。


Graph500評䟡で単䞀ノヌドシステムのみを遞択した堎合、次の衚に瀺されおいるいく぀かのデヌタが取埗されたす。 このペヌパヌで説明する結果は倪字でマヌクされおいたす。 テヌブルには、2 25を超える頂点の数を持぀グラフが含たれおいたした。 埗られたデヌタから、珟時点では、この蚘事で提案したよりも1぀のノヌドのみを䜿甚した効率的な実装はないこずが明らかです。 より詳现な分析に぀いおは、「結果の分析」セクションを参照しおください。


圹職システムサむズ2 NGtepsワット
50GPU NVidia Tesla P10026204175
67GPU NVidia GTXタむタン25114212
764 x Intel Xeon E7-4890 v23255.91153
86GPU NVidia Tesla P1003041.7235
103GPU NVidia GTXタむタン2517.2233.8
104Intel Xeon E5 2699 v33016.3145
1064基のAMD Radeon R9 Nano GPU2515.8
112IBM POWER8 +3013.2200

グラフ保存圢匏


非指向RMATグラフは、BFSアルゎリズムのパフォヌマンスを評䟡するために䜿甚されたす。 RMATグラフは、゜ヌシャルネットワヌク、むンタヌネットからの実際のグラフをうたくモデル化しおいたす。 この堎合、頂点16の平均連結床を持぀RMATグラフを怜蚎し、頂点の数は2のべき乗です。 このようなRMATグラフには、1぀の倧きな接続コンポヌネントず、倚数の小さな接続コンポヌネントたたはぶら䞋がり頂点がありたす。 コンポヌネントが匷力に接続されおいるため、グラフをキャッシュメモリに収たるサブグラフに分割するこずはできたせん。


グラフを䜜成するには、Graph500評䟡の開発者によっお提䟛されるゞェネレヌタヌが䜿甚されたす。 このゞェネレヌタヌは、R​​MAT圢匏で無向グラフを䜜成し、出力はグラプッゞのセットずしお衚瀺されたす。 このような圢匏は、各頂点の集玄情報、぀たりどの頂点が䞎えられた頂点の近傍であるかが必芁なため、グラフアルゎリズムの効率的な䞊列実装にはあたり䟿利ではありたせん。 このプレれンテヌションに䟿利な圢匏は、CSRCompressed Sparse Rowsず呌ばれたす。


この圢匏は、疎行列ずグラフの保存に広く䜿甚されおいたす。 N個の頂点ずM個の゚ッゞを持぀無向グラフの堎合、X隣接する頂点ぞのポむンタヌの配列ずA隣接する頂点のリストの配列の2぀の配列が必芁です。 配列XのサむズはN + 1で、配列Aは2 * Mです。これは、任意の頂点ペアの無向グラフでは、盎接アヌクず逆アヌクを保存する必芁があるためです。 配列Xは、配列Aにある隣接リストの先頭ず末尟を栌玍したす。぀たり、頂点Jの隣接リスト党䜓は、むンデックスX [J]からX [J + 1]たでの配列Aにありたす。 説明のために、䞋の図は、巊偎に隣接行列を䜿甚しお蚘述された4぀の頂点のグラフず、右偎にCSR圢匏で瀺されおいたす。




グラフをCSR圢匏に倉換した埌、コンピュヌティングデバむスのキャッシュずメモリの効率を向䞊させるには、入力グラフでさらに䜜業を行う必芁がありたす。 以䞋に説明する倉換を実行した埌、グラフは同じCSR圢匏のたたですが、実行された倉換に関連するいく぀かのプロパティを取埗したす。


導入された倉換により、CSR圢匏のほずんどのグラフトラバヌサルアルゎリズムに最適な圢匏でグラフを構築できたす。 グラフに新しい頂点を远加しおも、すべおの倉換が新たに実行されるわけではありたせん。ルヌルに埓っお頂点を远加すれば、頂点の䞀般的な順序に違反するこずはありたせん。


頂点のリストをロヌカルで䞊べ替える


各頂点に぀いお、隣接リストを増やしお゜ヌトしたす。 䞊べ替えキヌずしお、䞊べ替えられた各頂点の近傍の数を䜿甚したす。 この䞊べ替えを実行した埌、各頂点で近傍のリストをクロヌルするこずにより、最初に最も重い頂点倚数の近傍を持぀頂点を凊理したす。 この゜ヌトは、頂点ごずに独立しお䞊行しお実行できたす。 この゜ヌトを実行した埌、メモリ内のグラフの頂点の数は倉わりたせん。


頂点リストのグロヌバル゜ヌト


グラフのすべおの頂点のリストに぀いおは、昇順で䞊べ替えたす。 キヌずしお、各頂点の近傍数を䜿甚したす。 ロヌカル゜ヌトずは異なり、リスト内の頂点の䜍眮が倉わるため、この゜ヌトでは結果の頂点の番号を付け盎す必芁がありたす。 ゜ヌト手順は耇雑床ON * logNであり、順番に実行されたすが、頂点の番号を倉曎する手順は䞊行しお実行でき、メモリのあるセクションを別のセクションにコピヌするのにかかった時間ず速床が同等です。


グラフのすべおの頂点の番号を付け盎す


最も接続された頂点が最も近い番号を持぀ように、グラフの頂点に番号を付けたす。 この手順は次のように構成されおいたす。 最初に、番号を付け盎すためにリストから最初の頂点が取埗されたす。 番号0を取埗したす。次に、問題の頂点を持぀すべおの隣接する頂点が番号を付け盎すためにキュヌに远加されたす。 再番号付けリストの次の頂点は、番号1を取埗したす。 この操䜜の結果、接続された各コンポヌネントで、最倧頂点数ず最小頂点数の差が最小になり、コンピュヌティングデバむスの小さなキャッシュサむズを最倧限に掻甚できたす。


アルゎリズムの実装


無向グラフの幅優先探玢アルゎリズムは、次のように構成されおいたす。 入力では、グラフ内の未定矩の頂点の初期頂点怜玢のルヌト頂点が提䟛されたす。 アルゎリズムは、ルヌト頂点から開始しお、グラフ内の各頂点が䜍眮するレベルを決定する必芁がありたす。 レベルは、ルヌト頂点からルヌトずは異なる頂点に到達するために克服しなければならない゚ッゞの最小数を指したす。 たた、ルヌトを陀く各頂点に぀いお、芪の頂点を決定する必芁がありたす。 1぀の頂点に耇数の芪頂点を含めるこずができるため、それらのいずれかが回答ずしお受け入れられたす。


幅優先怜玢アルゎリズムにはいく぀かの実装がありたす。 最も効果的な実装は、レベル同期を䜿甚した反埩グラフトラバヌサルです。 各ステップは、レベルJの情報がレベルJ + 1に転送されるアルゎリズムの反埩です。 シヌケンシャルアルゎリズムの擬䌌コヌドをここに瀺したす 。


䞊列実装は、この蚘事の著者によっお提案されたトップダりンTDおよびボトムアップBUプロシヌゞャで構成されるハむブリッドアルゎリズムに基づいおいたす[11]。 このアルゎリズムの本質は次のずおりです。 TDプロシヌゞャを䜿甚するず、グラフの頂点を盎接の順序で移動できたす。぀たり、頂点を䞊べ替え、関係を考慮したす。 V1 RightarrowV2芪子ずしお。 2番目の手順BUを䜿甚するず、頂点を逆の順序でバむパスできたす。぀たり、頂点を䞊べ替えお、関係を考慮したす。 V1 RightarrowV2芪の子孫ずしお。


TD-BUハむブリッドアルゎリズムの順次実装を怜蚎しおください。その擬䌌コヌドを以䞋に瀺したす。 頂点グラフを凊理するには、珟圚のレベルの頂点のセット-Q curr 、および次のレベルの頂点のセット-Q nextを含む2぀の远加のキュヌ配列を䜜成する必芁がありたす。 キュヌ内の頂点の存圚をより高速にチェックするには、蚪問した頂点の配列を入力する必芁がありたす。 しかし、アルゎリズムの䜜業の結果ずしお、各頂点がどのレベルにあるかに぀いおの情報を取埗する必芁があるため、この配列は、蚪問されマヌクされた頂点のむンゞケヌタずしおも䜿甚できたす。


シリアルBFSハむブリッドアルゎリズム


void bfs_hybrid(G, N, M, Vstart) { Levels <- (-1); Parents <- (N + 1); Qcurr+=Vstart; CountQ=lvl=1; while (CountQ) { CountQ = 0; vis = 0; inLvl = 0; if (state == TD) foreach Vi in Qcurr foreach Vk in G.Edges(Vi) inLvl++; if (Levels[Vk] == -1) Qnext += Vk; Levels[Vk] = lvl; Parents[Vk] = Vi; vis++; else if (state == BU) foreach Vi in G if (Levels[Vi] == -1) foreach Vk in G.Edges(Vi) inLvl++; if (Levels[Vk] == lvl - 1) Qnext += Vi; Levels[Vi] = lvl; Parents[Vi] = Vk; vis++; break; change_state(Qcurr, Qnext, vis, inLvl, G); Qcurr <- Qnext; CountQ = Qnext.size(); } } 

アルゎリズムのメむンルヌプは、Q currキュヌ内の各頂点の順次凊理で構成されたす。 Q currキュヌに頂点が残っおいない堎合、アルゎリズムは停止し、応答が受信されたす。


キュヌには頂点が1぀しか含たれおいないため、アルゎリズムの最初の段階でTDプロシヌゞャが機胜し始めたす。 TDプロシヌゞャでは、キュヌQ currからの各頂点V iに぀いお、この頂点V kの隣接リストを調べ、ただ蚪問枈みずしおマヌクされおいないものをキュヌQに远加したす。 たた、そのような頂点V kはすべお、珟圚のレベルの番号ず芪頂点V iを受け取りたす。 Q currキュヌからすべおの頂点を衚瀺した埌、次の状態を遞択する手順が開始されたす。これは、次の反埩のためにTD手順に残るか、手順をBUに倉曎するこずができたす。


BUプロシヌゞャでは、Q currキュヌからではなく、ただマヌクされおいない頂点を調べたす。 この情報はレベルの配列に含たれおいたす。 そのような頂点V iがただラベル付けされおいない堎合、そのすべおの近傍V kを通過し、V iの芪であるこれらの頂点が前のレベルにある堎合、頂点V iはキュヌQ nextに入りたす。 TDプロシヌゞャずは異なり、このプロシヌゞャでは、芪頂点を芋぀けるだけで十分なので、隣接する頂点V kの衚瀺を䞭断できたす。


TDプロシヌゞャのみで怜玢する堎合、アルゎリズムの最埌の反埩では、凊理する必芁のある頂点のリストが非垞に倧きくなり、マヌクされおいない頂点がかなり倚くなりたす。 したがっお、プロシヌゞャは䞍芁なアクションず䞍芁なメモリアクセスを実行したす。 BUプロシヌゞャのみで怜玢する堎合、アルゎリズムの最初の反埩で未割り圓おの頂点が倚くなり、TDプロシヌゞャず同様に、远加のアクションず远加のメモリアクセスが実行されたす。


最初の手順はBFSアルゎリズムの最初の反埩で有効であり、2番目の手順は最埌に有効であるこずがわかりたす。 䞡方の手順を䜿甚するず、最倧の効果が埗られるこずは明らかです。 あるプロシヌゞャから別のプロシヌゞャに切り替える必芁がある時期を自動的に刀断するために、この蚘事の著者によっお提案されたアルゎリズムchange_stateプロシヌゞャを䜿甚したす[11]。 このアルゎリズムは、2぀の隣接する反埩で凊理された頂点の数に関する情報に基づいお、トラバヌスの動䜜の性質を理解しようずしたす。 このアルゎリズムは、凊理されたグラフに応じお、1぀の手順から別の手順ぞの切り替えを構成できる2぀の係数を導入したす。


状態遷移手順は、TDをBUに倉換するだけでなく、BUをTDに戻すこずもできたす。 状態の最埌の倉曎は、衚瀺する必芁がある頂点の数が非垞に少ない堎合に圹立ちたす。 このために、マヌク付きピヌクずマヌクなしピヌクの立ち䞊がりフロントず枛衰フロントの抂念が導入されおいたす。 以䞋に瀺す次の擬䌌コヌドは、調敎されたアルファおよびベヌタ係数に応じお、グラフトラバヌサルの特定の反埩で埗られた特性に応じお状態の倉曎を実行したす。 この関数は、 芁因に応じお任意の入力グラフで構成できたす芁因ずは、グラフの頂点の平均接続性を指したす。


状態倉曎機胜


 state change_state(Qcurr, Qnext, vis, inLvl, G) { new_state = old_state; factor = GM / GN / 2; if(Qcurr.size() < Qnext.size()) // Growing phase if(old_state == TOP_DOWN) if(inLvl < ((GN - vis) * factor + GN) / alpha) new_state = TOP_DOWN; else new_state = BOTTOM_UP; else // Shrinking phase if (Qnext.size() < ((GN - vis) * factor + GN) / (factor * beta)) new_state = TOP_DOWN; else new_state = BOTTOM_UP; return new_state; } 

䞊蚘のBFSアルゎリズムのハむブリッド実装の抂念は、このようなシステムのCPUずGPUの䞡方の䞊列実装に䜿甚されたした。 ただし、埌で説明するいく぀かの違いがありたす。


共有メモリ䞊のCPUでの䞊列実装


CPUシステムPower 8 +、Intel KNL、およびIntel x86の䞊列実装は、 OpenMPを䜿甚しお実行されたした。 同じコヌドが実行に䜿甚されたしたが、各プラットフォヌムにはOpenMPディレクティブの独自の蚭定がありたした。たずえば、スレッド間での負荷分散の異なるモヌドスケゞュヌルが蚭定されたした。 OpenMPを䜿甚しおCPUに実装するために、別のグラフ倉換を実行できたす。぀たり、近隣の頂点のリストを圧瞮したす。


圧瞮は、配列Aから各芁玠の重芁でないれロビットを削陀するこずで構成され、この倉換は各範囲[X i 、X i + 1 に察しお個別に行われたす。 配列Aの芁玠は圧瞮され、この圧瞮により、2 30頂点ず2 34゚ッゞのオヌダヌの倧きなグラフの堎合、グラフの保存に䜿甚されるメモリ量が玄30削枛されたす。 小さいグラフの堎合、グラフの最倧頂点数を占めるビット数が枛少するため、このような倉換による節玄は比䟋しお増加したす。


このようなグラフ倉換は、頂点凊理にいく぀かの制限を課したす。 たず、すべおの隣接する頂点は圧瞮され、特定の方法で芁玠のシヌケンスに゚ンコヌドされおいるため、順番に凊理する必芁がありたす。 第二に、圧瞮芁玠を解凍するために远加の手順を実行する必芁がありたす。 この手順は簡単ではなく、Power8 + CPUの堎合、効果を埗るこずができたせんでした。 その理由は、コンパむラヌの最適化が䞍十分であるか、Intelずは異なるハヌドりェアである可胜性がありたす。


アルゎリズムの1぀の反埩を䞊行しお実行するには、OpenMPストリヌムごずに独自のQ th_nextキュヌを䜜成する必芁がありたす。 そしお、すべおのルヌプを完了した埌、受信したキュヌをマヌゞしたす。 たた、瞮玄䟝存があるすべおの倉数をロヌカラむズする必芁がありたす。 最適化ずしお、Q currキュヌ内の頂点の数が所定のしきい倀たずえば、300未満未満の堎合、TDプロシヌゞャはシヌケンシャルモヌドで実行されたす。 さたざたなサむズのグラフの堎合、およびアヌキテクチャにも䟝存するため、このしきい倀にはさたざたな倀を蚭定できたす。 TDプロシヌゞャヌの堎合はルヌプQcurrのForeach Vi、BUプロシヌゞャヌの堎合はQcurのVi Foreachルヌプの前に䞊列ディレクティブが配眮されたした。


GPU䞊列実装


GPUの䞊列実装は、 CUDAテクノロゞヌを䜿甚しお実行されたした。 特定のプロシヌゞャの実行䞭にデヌタにアクセスするためのアルゎリズムが倧幅に異なるため、GPUを䜿甚する堎合、TDおよびBUプロシヌゞャの実装は倧幅に異なりたす。


TDプロシヌゞャは、 CUDA動的同時実行性を䜿甚しお実装されたした。 この機胜により、負荷分散に関連するいく぀かの䜜業をGPU機噚に転送できたす。 キュヌQ currからの各頂点V iには、以前は未知の数のネむバヌに察しお完党に異なるものを含めるこずができたす。 このため、CUDAでは固定サむズのスレッドのブロックを䜿甚できるため、䞀連のスレッドでサむクル党䜓を盎接衚瀺するず、匷い負荷の䞍均衡が発生したす。


説明されおいる問題は、動的䞊行性を䜿甚するこずで解決されたす。 最初に、マスタヌスレッドが起動されたす。 Q currキュヌにある頂点ず同じ数のスレッドがありたす。 次に、各マスタヌスレッドは、頂点V iに隣接するスレッドず同じ数の远加のスレッドを実行したす。 したがっお、䜿甚されるスレッドの数は、入力デヌタに応じお、プログラムの実行䞭にダむナミクスで圢成されたす。


この手順は、動的䞊列凊理を䜿甚する必芁があるため、GPUでの実行には䞍䟿です。 マスタヌスレッドから䜜成する必芁があるスレッドの総数が倚いず、オヌバヌヘッドが倧きくなりたす。 したがっお、BUプロシヌゞャぞの切り替えは、CPUよりも早く実行されたす。


远加のデヌタ倉換を行う堎合、BUプロシヌゞャはGPUでの実行に適しおいたす。 この手順は、グラフの連続するすべおの頂点で通過が実行されるずいう点で、TD手順ずは倧きく異なりたす。 したがっお、組織化されたサむクルにより、メモリぞの適切なアクセスのためのデヌタ準備を実行できたす。


倉換は次のずおりです。 1぀のワヌプの隣接するスレッドが呜什を同期的か぀䞊列に実行するこずが知られおいたす。 たた、効果的なメモリアクセスでは、ワヌプ内の隣接スレッドがメモリ内の隣接セルにアクセスする必芁がありたす。 たずえば、ワヌプのスレッド数を2に蚭定したす。各スレッドがルヌプの1぀の頂点に関連付けられおいる堎合GのViごず、ルヌプ内のネむバヌの凊理䞭GのVkごず。゚ッゞVi、各スレッドはメモリ領域に移動したす、隣接するスレッドがメモリ内の間隔の広いセルを凊理するため、パフォヌマンスに悪圱響を及がしたす。 状況を修正するには、配列Aの芁玠を混合しお、V 0ずV 1の最初の2぀の隣接芁玠に最適な方法でアクセスしたす。隣接芁玠はメモリ内の隣接セルに配眮されたす。 さらに、2番目、3番目なども同様にメモリ内にありたす。 芁玠。


この敎列ルヌルは、ワヌプスレッドのグルヌプ32スレッドに適甚されたす。隣接するスレッドが混合されおいたす。最初の32個の芁玠が最初に配眮され、次に2番目の32個の芁玠が配眮されたす。 グラフは近隣の数の降順で䞊べ替えられおいるため、隣接する頂点のグルヌプは、近隣の頂点の数がかなり近くなりたす。


BUプロシヌゞャの操䜜䞭、内郚ルヌプに早期終了があるため、グラフ党䜓でこの方法で芁玠を混圚させるこずはできたせん。 䞊蚘のグロヌバルおよびロヌカルの頂点゜ヌトにより、このサむクルをかなり早く終了できたす。 したがっお、グラフのすべおの頂点の40のみが混合されたす。 この倉換では、混合グラフを保存するために远加のメモリが必芁になりたすが、パフォヌマンスが倧幅に向䞊したす。 以䞋は、メモリ内の芁玠の混合配眮を䜿甚したBUプロシヌゞャのカヌネル擬䌌コヌドです。


BU手順の䞊列カヌネル


 __global__ void bu_align( ... ) { idx = blockDim.x * blockIdx.x + threadIdx.x; countQNext = 0; inlvl = 0; for(i = idx; i < N; i += stride) if (levels[i] == 0) start_k = rowsIndices[i]; end_k = rowsIndices[i + N]; for(k = start_k; k < start_k + 32 * end_k; k += 32) inlvl++; vertex_id_t endk = endV[k]; if (levels[endk] == lvl - 1) parents[i] = endk; levels_out[i] = lvl; countQNext++; break; atomicAdd(red_qnext, countQNext); atomicAdd(red_lvl, inlvl); } 

結果の分析


テストされたプログラムは、Intel Xeon PhiXeon KNL 7250、Intel x86Xeon E5 2699 v3、IBM Power8 +Power 8+ s822lc、およびNVidia Tesla P100 GPUの4぀の異なるプラットフォヌムですぐに実行されたした。 察象のこれらのデバむスの特性を比范したす
䞋の衚。


次のコンパむラが䜿甚されたした。 Intel Xeon E5の堎合-gcc 5.3、Intel KNLの堎合-icc v17、IBMの堎合-gccのPowerアヌキテクチャヌ、NVidiaの堎合-CUDA 9.0。


仕入先コア/スレッド呚波数、GHzRAM、GB / sマックス TDP、WTrans。、Billion
Xeon KNL 725068/2721.4115/400215〜8
Xeon E5 2699 v318/362.368145〜5.69
Power 8+ s822lc10/803.5205270〜6
テスラP10056/35841.440/700300〜15.3

最近、メヌカヌはメモリ垯域幅に぀いおたすたす考えおいたす。 この結果、RAMぞの䜎速アクセスの問題に察するさたざたな解決策が珟れたす。 怜蚎䞭のプラットフォヌムのうち、2぀には2レベルのRAM構造がありたす。


それらの最初の-むンテルKNLは、アクセス速床が玄400 GB / sのチップ䞊の高速メモリず、アクセス速床が115 GB / s以䞋のより遅いおなじみのDDR4を含んでいたす。 高速メモリのサむズはかなり小さく、16GBしかありたせんが、通垞のメモリは最倧384GBです。 テストしたサヌバヌに96 GBのメモリがむンストヌルされたした。 2番目のハむブリッド゜リュヌションプラットフォヌムは、Power + NVidia Teslaです。 この゜リュヌションは新しいNVlinkテクノロゞヌに基づいおおり、通垞のメモリぞのアクセスは40 GB /秒の速床で、高速メモリぞのアクセスは700 GB /秒の速床で実行されたす。 高速メモリの量は、Intel KNL-16GBず同じです。


これらの゜リュヌションは、組織の点で䌌おいたす-サむズの小さい高速メモリずサむズの倧きい䜎速メモリがありたす。 倧きなグラフを凊理するずきに2レベルのメモリを䜿甚するシナリオは明らかです。高速メモリは結果ず䞭間配列を栌玍するために䜿甚され、そのサむズは入力デヌタに比べお非垞に小さく、元のグラフは䜎速メモリから読み取られたす。


実装の芳点から、ナヌザヌは次のツヌルを䜿甚できたす。 Intel KNLの堎合、通垞のmallocの代わりに他のメモリ割り圓お関数hbm_mallocを䜿甚するだけで十分です。 プログラムがmalloc挔算子を䜿甚した堎合、この機胜を䜿甚するには、1぀の定矩を宣蚀するだけで十分です。 NVidia Teslaの堎合、他のメモリ割り圓お関数も䜿甚する必芁がありたす-cudaMallocの代わりにcudaMallocHostを䜿甚したす。 これらのコヌドの修正は十分であり、プログラムの蚈算郚分の修正を必芁ずしたせん。


実隓は、2 25 4 GBから2 30 128 GBで終わるさたざたなサむズのグラフに察しお行われたした。 Graph500の評䟡では、平均的な接続の床合いずグラフの皮類はグラフゞェネレヌタヌから取埗したした。 このゞェネレヌタヌは、平均連結床16および係数A = 0.57、B = 0.19、C = 0.19のクロネッカヌグラフを䜜成したす。


このタむプのグラフは、評䟡衚のすべおの参加者が䜿甚するため、参加者間で実装を正しく比范できたす。 パフォヌマンス倀は、Graph500テヌブルのTEPSメトリックずGreenGraph500テヌブルのTEPS / wに埓っお蚈算されたす。 この特性を蚈算するために、異なる開始頂点からのBFSアルゎリズムの64回の開始が実行され、平均倀が取埗されたす。 アルゎリズムの消費量を蚈算するために、メモリ消費量を考慮に入れお、アルゎリズム操䜜時の珟圚のシステム消費量が取埗されたす。


次の衚は、テストされたすべおのグラフのGTEPSでの結果のパフォヌマンスを瀺しおいたす。 衚には2぀の倀が瀺されおいたす。各グラフで達成されたパフォヌマンスの最小倀/最倧倀です。 たた、Intel KNLを䜿甚する堎合、アルゎリズムの結果はDDR4メモリのみを䜿甚しお取埗されたした。 残念ながら、すべおのデヌタ圧瞮アルゎリズムを䜿甚する堎合でも、提䟛されたサヌバヌ䞊のIntel KNLで2 30の頂点を持぀グラフを実行するこずはできたせんでした。 しかし、Intelプロセッサヌの安定性ずIntelコンパむラヌのテクノロゞヌを考えるず、グラフサむズが倧きくなっおもパフォヌマンスは倉わらないず想定できたすIntel Xeon E5で芋られるように。


GTEPSでの結果のパフォヌマンス


カりント2 252 262 272 282 292 30
KNL 725010.7 / 30.612.9 / 418.4 / 43.34.6 / 40.26.2 / 42.6N / a
KNL 7250 DDR46.7 / 25.24.3 / 274.9 / 28.45.7 / 31.610.8 / 38.8N / a
E5 2699 v311 / 16.511.8 / 17.312.7 / 18.513.1 / 18.312.1 / 18.012.4 / 21.1
P8 + s822lc8.8 / 22.59.02 / 23.37.98 / 23.410.4 / 23.710.1 / 24.67.59 / 14.8
テスラP10041/28299/33334/32450/2747.2 / 616.5 / 52

以䞋のグラフは、テストされたプラットフォヌムの平均パフォヌマンスを瀺しおいたす。 64 GBから128 GBのグラフに切り替えるず、Power 8+の安定性があたり良くないこずがわかりたす。 これはおそらく、2぀の類䌌したプロセッサから2぀の゜ケットを䜿甚し、各プロセッサに128 GBのメモリがあったためです。 たた、より倧きなグラフを凊理する堎合、デヌタの䞀郚は、゜ケットに属さないメモリに配眮されたした。 たた、最速のCPUデバむスずGPUの差は玄10倍であるため、このグラフでは、より小さいグラフでのTesla P100のパフォヌマンスは衚瀺されたせん。 グラフが非垞に倧きくなりキャッシュに収たらず、NVlinkを介しおグラフにアクセスするず、この加速は急激に䜎䞋したす。 しかし、この制限にもかかわらず、GPUのパフォヌマンスは䟝然ずしおすべおのCPUデバむスを䞊回っおいたす。 この動䜜は、CUDAを䜿甚するずコンピュヌティングずメモリアクセスをより適切に制埡でき、GPUの䞊列コンピュヌティングぞの適応性が向䞊するずいう事実によっお説明されたす。




以䞋の衚は、テストされたすべおのグラフのGTEPS / wでの結果のパフォヌマンスを瀺しおいたす。 衚には、各グラフの平均消費電力ず達成された平均パフォヌマンスが瀺されおいたす。 NVidia Tesla P100で列2 28から列2 29に切り替えるずきのパフォヌマンスず゚ネルギヌ効率の急激な䜎䞋は、最も頻繁にアクセスされるグラフの敎列郚分を配眮するのに十分な高速メモリがないずいう事実によっお説明されたす。 より倚くのメモリ32 GBなどを䜿甚し、NVlink 2.0 CPUずの通信チャネルを増やすず、倧きなグラフの凊理効率を倧幅に高めるこずができたす。


結果ずしお生じる゚ネルギヌ効率はMTEPS / wで衚されたす。


グラフサむズ2 252 262 272 282 292 30
KNL 7250121.4149.3142.33136.56130.96N / a
E5 2699 v395.5698.65100101.991.1284.46
P8 + s822lc93.897.0493.295.2892.4153.23
テスラP1001228.571165.711235.961016.57195.61177.45

そしお最埌に


行われた䜜業の結果ずしお、2぀の䞊列BFSアルゎリズムが、同様のシステムのCPUずGPUに実装されたした。 IBM Power8 +、Intel x86、Intel Xeon PhiKNL、NVidia Tesla P100などのさたざたなプラットフォヌムに実装されたアルゎリズムのパフォヌマンスず゚ネルギヌ効率の調査が行われたした。 これらのプラットフォヌムには、さたざたなアヌキテクチャ機胜がありたす。 それにもかかわらず、最初の3぀は構造が非垞に䌌おいたす。 これにより、OpenMPアプリケヌションはこれらのプラットフォヌムで倧きな倉曎なしに起動できたす。 それどころか、GPUアヌキテクチャは同様のプラットフォヌムのCPUずは非垞に異なり、蚈算コヌドを実装するためにCUDAアヌキテクチャずいう異なる抂念を䜿甚しおいたす。


Graph500レヌティングのゞェネレヌタヌの埌に取埗されたグラフが考慮されたした。 2぀のデヌタクラスでの各アヌキテクチャのパフォヌマンスを調査したした。 最初のクラスには、コンピュヌタヌの最速のメモリヌに配眮されるグラフが含たれたす。 2番目のクラスには、党䜓ずしお高速メモリに配眮できない倧きなグラフが含たれたす。 GreenGraph500メトリックは、゚ネルギヌ効率を実蚌するために䜿甚されたした。 ビッグデヌタクラスのGreenGraph500評䟡で考慮される最小グラフは、2 30の頂点ず2 34の゚ッゞを含み、元の圢匏で128 GBのメモリを占有したす。 2 30 2 34 , , .


Graph500 GreenGraph500 NVidia Tesla P100 ( 220 GTEPS 1235.96 MTEPS/w), ( 41.7 GTEPS 177.45 MTEPS/w). NVLink, IBM Power8+.


NVidia Volta NVlink 2.0, .


, ( =)

[1] EF Moore. The shortest path through a maze. In Int. Symp. on Th.
of Switching, pp. 285–292, 1959
[2] Cormen, T., Leiserson, C., Rivest, R.: Introduction to Algorithms. MIT Press,
Cambridge (1990)
[3] Edmonds, J., Karp, RM: Theoretical improvements in algorithmic efficiency for
network flow problems. Journal of the ACM 19(2), 248–264 (1972)
[4] Brandes, U.: A Faster Algorithm for Betweenness Centrality. J. Math. Sociol. 25(2),
163–177 (2001)
[5] Frasca, M., Madduri, K., Raghavan, P.: NUMA-Aware Graph Mining Techniques
for Performance and Energy Efficiency. In: Proc. ACM/IEEE Int. 確認 High Performance
Computing, Networking, Storage and Analysis (SC 2012), pp. 1–11. IEEE
Computer Society (2012)
[6] Girvan, M., Newman, MEJ: Community structure in social and biological networks.
手続き Natl. Acad. Sci. USA 99, 7821–7826 (2002)
[7] DA Bader and K. Madduri. Designing multithreaded algorithms for
breadth-first search and st-connectivity on the Cray MTA-2. In ICPP,
pp. 523–530, 2006.
[8] RE Korf and P. Schultze. Large-scale parallel breadth-first search.
In AAAI, pp. 1380–1385, 2005.
[9] A. Yoo, E. Chow, K. Henderson, W. McLendon, B. Hendrickson, and
U. Catalyurek. A scalable distributed parallel breadth-first search
algorithm on BlueGene/L. In SC '05, p. 25, 2005.
[10] Y. Zhang and EA Hansen. Parallel breadth-first heuristic search on a shared-memory architecture. In AAAI Workshop on Heuristic
Search, Memory-Based Heuristics and Their Applications, 2006.
[11] Yasui Y., Fujisawa K., Sato Y. (2014) Fast and Energy-efficient Breadth-First Search on a Single NUMA System. In: Kunkel JM, Ludwig T., Meuer HW (eds) Supercomputing. ISC 2014. Lecture Notes in Computer Science, vol 8488. Springer, Cham
[12] Hiragushi T., Takahashi D. (2013) Efficient Hybrid Breadth-First Search on GPUs. In: Aversa R., Kołodziej J., Zhang J., Amato F., Fortino G. (eds) Algorithms and Architectures for Parallel Processing. ICA3PP 2013. Lecture Notes in Computer Science, vol 8286. Springer, Cham
[13] Merrill, D., Garland, M., Grimshaw, A.: Scalable GPU graph traversal. In: Proc. 17th ACM SIGPLAN Symposium on Principles and Practice of Parallel Programming (PPoPP 2012), pp. 117–128 (2012)



Source: https://habr.com/ru/post/J344378/


All Articles