JP2016119084A - 効率的な疎行列表現及び処理のためのコンピュータ実装システム及び方法 - Google Patents

効率的な疎行列表現及び処理のためのコンピュータ実装システム及び方法 Download PDF

Info

Publication number
JP2016119084A
JP2016119084A JP2015240835A JP2015240835A JP2016119084A JP 2016119084 A JP2016119084 A JP 2016119084A JP 2015240835 A JP2015240835 A JP 2015240835A JP 2015240835 A JP2015240835 A JP 2015240835A JP 2016119084 A JP2016119084 A JP 2016119084A
Authority
JP
Japan
Prior art keywords
matrix
index
row
array
zero
Prior art date
Legal status (The legal status is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the status listed.)
Granted
Application number
JP2015240835A
Other languages
English (en)
Other versions
JP2016119084A5 (ja
JP6630558B2 (ja
Inventor
ロン・チョウ
Chou Ron
Current Assignee (The listed assignees may be inaccurate. Google has not performed a legal analysis and makes no representation or warranty as to the accuracy of the list.)
Palo Alto Research Center Inc
Original Assignee
Palo Alto Research Center Inc
Priority date (The priority date is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the date listed.)
Filing date
Publication date
Application filed by Palo Alto Research Center Inc filed Critical Palo Alto Research Center Inc
Publication of JP2016119084A publication Critical patent/JP2016119084A/ja
Publication of JP2016119084A5 publication Critical patent/JP2016119084A5/ja
Application granted granted Critical
Publication of JP6630558B2 publication Critical patent/JP6630558B2/ja
Active legal-status Critical Current
Anticipated expiration legal-status Critical

Links

Images

Classifications

    • GPHYSICS
    • G06COMPUTING; CALCULATING OR COUNTING
    • G06FELECTRIC DIGITAL DATA PROCESSING
    • G06F17/00Digital computing or data processing equipment or methods, specially adapted for specific functions
    • G06F17/10Complex mathematical operations
    • G06F17/16Matrix or vector computation, e.g. matrix-matrix or matrix-vector multiplication, matrix factorization

Landscapes

  • Engineering & Computer Science (AREA)
  • Physics & Mathematics (AREA)
  • General Physics & Mathematics (AREA)
  • Mathematical Physics (AREA)
  • Pure & Applied Mathematics (AREA)
  • Mathematical Analysis (AREA)
  • Mathematical Optimization (AREA)
  • Computational Mathematics (AREA)
  • Data Mining & Analysis (AREA)
  • Theoretical Computer Science (AREA)
  • Computing Systems (AREA)
  • Algebra (AREA)
  • Databases & Information Systems (AREA)
  • Software Systems (AREA)
  • General Engineering & Computer Science (AREA)
  • Complex Calculations (AREA)
  • Memory System Of A Hierarchy Structure (AREA)

Abstract

【課題】元の構造を乱すことなく疎行列処理の速度を向上させるコンピュータ実装システム及び方法を提供する。【解決手段】疎行列が処理される速度は、行列の改善された圧縮表現を使用することによって増加させることができる。構造化圧縮表現は、キャッシュがランダムにアクセスされなければならない回数を削減することによって行列処理中に経験するキャッシュミス数を低減する。さらに、非ゼロエントリのそれらの数に基づいて行列の行及び列を分割して再グループ化する行列の表現は、行列のこれらの部分を処理するための最も適切なカーネル関数を割り当てるのを可能とし、GPUベースのハードウェアの限界を克服する。【選択図】図6

Description

本特許出願は、一般に、行列データを処理することに関し、特に、効率的な疎行列表現及び処理のためのコンピュータ実装システム及び方法に関する。
疎行列は、要素の大部分がゼロである行列である。そのような行列を使用する操作は、様々なアプリケーションを有し、通常、そのようなアプリケーションの最も計算集約的な部分である。例えば、疎行列−ベクトル乗算(SpMV)及び疎行列転置ベクトル乗算(SpMTV)、疎線形代数(SLA)における基本的な操作は、検索結果を提供するときにウェブページをランク付けするためのグーグル(登録商標)によって使用されているページランク(登録商標)アルゴリズムなどのランク付けアルゴリズムを実行するために使用されている。SpMV及びSpMTVは、そのようなアプリケーションの最も計算集約的な部分であり、行列が使用可能な速度は、SpMV及びSpMTVによって制限される。
疎行列処理の速度を向上させるための試みが行われてきたが、そのような努力は、まだ改善のための重要な余地を残している。例えば、速度を高めるために、行列は、非ゼロエントリの行列内の値及び位置に関する情報の複数の配列を含む圧縮フォーマットで符号化されており、ゼロエントリに関する情報を省略している。例えば、圧縮疎行列の行フォーマットは、非ゼロエントリの値を有する配列、非ゼロエントリが配置された列、及び、各行における第1の非ゼロエントリの第1の配列内のインデックスを保持する配列を含む。圧縮疎列フォーマットは、同様の配列を含む。そのような配列は、配列データに対する高速アクセスを可能にするために計算を実行するプロセッサのキャッシュに記憶される。しかしながら、より大きな行列の場合、圧縮フォーマット配列は、キャッシュに収まらないことがあり、プロセッサが計算の単一ステップを実行するためにメインメモリにおける行列を表す異なる配列にアクセスすることを必要とする。そのような構成において、中央処理装置(CPU)及びグラフィックス処理ユニット(GPU)を含む現代のコンピュータプロセッサは、計算中にキャッシュミスやキャッシュから必要なデータを取得するためにプロセッサによる障害を経験する可能性がある。キャッシュミスの後に計算を終了することは、プロセッサがメインメモリから失われたデータを取得することを必要とし、はるかに遅くなることがある。
さらに、GPUがSpMV及びSpMTVなどの行列演算を実行するために使用される場合に追加の懸念が存在する。GPUは、そのほとんどの要素が非ゼロエントリである密行列の処理などの密計算用により良好に設計されて最適化されている。そのようなハードウェアは、一般に、行列データを処理するために単一のカーネル関数を実行する。その結果、ハードウェアは、異なる行又は列などの行列の異なる部分における非ゼロエントリの数の大幅な変動に応答することができない。例えば、行列の単一の行又は列を処理するために単一のスレッドを割り当てるカーネルは、最も密な行又は列を処理するために割り当てられたスレッドに依存する全体の処理時間により、負荷不均衡を被ることがある。一方、単一の行又は列を処理するために複数のスレッドを割り当てるカーネルは、処理に関与しないいくつかの割り当てられたスレッドにより、割り当てられたスレッドの数が行又は列における非ゼロエントリの数未満である場合にハードウェアリソースの浪費を被る。
したがって、キャッシュミスの可能性を低減させ且つ行列の異なる部分における非ゼロエントリの数の変化に応答することができるような方法で疎行列を表現する必要がある。
疎行列が処理される速度は、行列の改善された圧縮表現を使用することによって増加させることができる。構造化圧縮表現は、キャッシュがランダムにアクセスされなければならない回数を削減することにより、行列処理中に経験したキャッシュミスの数を低減する。さらに、それらの非ゼロエントリ数に基づいて行列の行及び列を分割し且つ再グループ化するという行列の表現は、行列のこれらの部分を処理するための最も適切なカーネル関数を割り当てることができ、GPUベースハードウェアの限界を克服する。その結果、処理速度は、行列の元の構造を乱すことなく増加させることができる。
1つの実施形態は、構造化疎行列表現を取得するためのコンピュータ実装システム及び方法を提供する。各部分が行及び列の一方を含む行列の部分において1つ以上の順序で配置された1つ以上の非ゼロエントリを含む行列の構造化圧縮表現が取得され、各要素が非ゼロエントリのいずれかとその非ゼロエントリを含む部分のいずれかのインデックスとを含む1つ以上の要素を含む複合配列を取得することと、その順序のうちの1つ以上における最初である非ゼロエントリを含む各要素の複合配列におけるインデックスを含み、さらに行列における複数の非ゼロエントリを含むインデックス配列を取得することとを備える。
さらなる実施形態は、効率的な疎行列表現及び処理のためのコンピュータ実装システム及び方法を提供する。1つ以上の非ゼロエントリを含む疎行列の圧縮表現が取得され、各部分が行列の行及び列の一方を含む行列の部分が行列におけるそれらの部分に基づいてインデックス付けされる。行列の部分についての複数のパーティションが定義される。各部分における複数の非ゼロエントリは、圧縮表現を使用して取得される。各部分は、その部分における複数の非ゼロエントリに基づいてパーティションの1つに関連付けられる。各パーティションに関連する全ての部分のリストが作成され、部分はそれらのインデックスの順序でリスト化される。リストを含むマッピング配列を含む行列のパーティション化された圧縮表現が作成される。
GPUベースのハードウェアについて、GPUは、通常、CPUよりもはるかに柔軟性が低い単一命令複数スレッドの実行モデルを想定していることから、パーティション化された(構造化)CSR/CSC符号化は、非パーティション化のものよりも好ましい。このモデルは、GPUをCPUよりも負荷不均衡の問題に対してより脆弱にする。PSCSR及びPSCSCなどのパーティション化された疎行列符号化は、非ゼロのそれらの数に基づいて同様の計算特性を有する行又は列をともにグループ化することによってGPUに対する負荷不均衡を効果的に低減することができる。
本発明のさらに他の実施形態は、以下の詳細な説明から当業者にとって容易に明らかになるであろう。本発明を実施するための最良の形態を説明することによって本発明の実施形態が記載される。理解されるように、本発明は、他の及び異なる実施形態が可能であり、そのいくつかの詳細は、全て本発明の精神及び範囲から逸脱することなく様々な明白な点において変更が可能である。したがって、図面及び詳細な説明は、本質的に例示であり、限定されるものではないとみなされるべきである。
図1は、1つの実施形態(従来技術)にかかる、圧縮された疎行フォーマットの符号化において符号化された行列に対してSpMVを実行するための方法を示すフロー図である。 図2は、1つの実施形態(従来技術)にかかる、圧縮された疎列フォーマットの符号化において符号化された行列に対してSpMTVを実行するための方法を示すフロー図である。 図3は、1つの実施形態にかかる、疎行列の効率的な表現及び処理のためのコンピュータ実装システムを示すブロック図である。 図4は、一例として隣接行列を示すグラフである。 図5は、図3のシステムの様々なハードウェアセットアップのための様々な符号化の強みをまとめたグラフである。 図6は、1つの実施形態にかかる構造化疎行列表現を取得するためのコンピュータ実装方法を示すフロー図である。 図7は、1つの実施形態にかかる図6の方法において使用するための構造化CSR符号化を使用して符号化された行列に対してSpMVを実行するためのルーチンである。 図8は、1つの実施形態にかかる図6の方法において使用するための構造化CSC符号化を使用して符号化された行列に対してSpMTVを実行するためのルーチンである。 図9は、1つの実施形態にかかる効率的な疎行列パーティション及び処理のためのコンピュータ実装方法を示すフロー図である。 図10は、1つの実施形態にかかる図9の方法において使用するための圧縮表現で表される行列の行の順序維持パーティションを実行するためのルーチンを示すフロー図である。 図11は、1つの実施形態にかかる図9の方法において使用するための圧縮表現で表される行列の列の順序維持パーティションを実行するためのルーチンを示すフロー図である。 図12は、1つの実施形態にかかる図9の方法において使用するためのマッピング配列にパーティションをマージするためのルーチンを示すフロー図である。 図13は、1つの実施形態にかかる図9の方法において使用するためのパーティション化された圧縮表現で符号化された行列に対してSpMVを実行するためのルーチンを示すフロー図である。 図14は、1つの実施形態にかかる図9の方法において使用するためのパーティション化された圧縮表現で符号化された行列に対してSpMTVを実行するためのルーチンを示すフロー図である。 図15は、1つの実施形態にかかる図9の方法において使用するためのパーティション化された圧縮表現処理のためのカーネル関数を起動するためのルーチンである。 図16は、1つの実施形態にかかる図15のルーチンにおいて使用するための選択されたカーネル関数の起動引数を設定するためのルーチンである。 図17は、1つの実施形態にかかる図15のルーチンにおいて使用するためのf1T1Rカーネル関数によってSpMVを実行するためのルーチンを示すフロー図である。 図18は、1つの実施形態にかかる図15のルーチンにおいて使用するためのf1T1Rカーネル関数によってSpMTVを実行するためのルーチンを示すフロー図である。 図19Aは、1つの実施形態にかかる図15のルーチンにおいて使用するためのf1W1Rカーネル関数によってSpMVを実行するためのルーチンを示すフロー図である。 図19Bは、1つの実施形態にかかる図15のルーチンにおいて使用するためのf1W1Rカーネル関数によってSpMVを実行するためのルーチンを示すフロー図である。 図20Aは、1つの実施形態にかかる図15のルーチンにおいて使用するためのf1W1RカーネルによってSpMTVを実行するためのルーチンを示すフロー図である。 図20Bは、1つの実施形態にかかる図15のルーチンにおいて使用するためのf1W1RカーネルによってSpMTVを実行するためのルーチンを示すフロー図である。 図21Aは、1つの実施形態にかかる図15のルーチンにおいて使用するためのf1B1Rカーネル関数によってSpMVを実行するためのルーチンを示すフロー図である。 図21Bは、1つの実施形態にかかる図15のルーチンにおいて使用するためのf1B1Rカーネル関数によってSpMVを実行するためのルーチンを示すフロー図である。 図22Aは、1つの実施形態にかかる図15のルーチンにおいて使用するためのf1B1Rカーネル関数によってSpMTVを実行するためのルーチンを示すフロー図である。 図22Bは、1つの実施形態にかかる図15のルーチンにおいて使用するためのf1B1Rカーネル関数によってSpMTVを実行するためのルーチンを示すフロー図である。 図23は、1つの実施形態にかかる図6及び図9の方法において使用するためのべき乗法を実行するためのルーチンを示すフロー図である。
上述したように、疎行列は、圧縮された疎行(CSR)フォーマットで符号化されることができる。Aをe個の非ゼロエントリを有するm×nの疎行列とする。A、A及びAを、それぞれ、長さe、e及び(m+1)の3つの1次元配列とする。CSRフォーマットにおいて、Aは、<A、A、A>として符号化され、
・値配列Aは、行優先順序でAにおける全ての非ゼロエントリの値を保持し、
・列配列Aは、行優先順序でAにおける全ての非ゼロエントリの列を保持し、
・行インデックス配列Aは、Aにおける各行の第1の非ゼロエントリのAにおけるインデックスを保持し、A配列におけるエントリの総数であるA配列における最後のエントリを有する。
例えば、以下の4×4行列
Figure 2016119084
は、<A=[1,0.5,0.5,0.5,0.5,1]、A=[3,0,3,0,1,2]、A=[0,1,3,5,6]>としてCSRにおいて符号化されることができる。
以下に記載される本例及び他の例において、行及び列は、それらのidとして機能するインデックスによってインデックス付けされる。インデックスは、ゼロから始まり、行列の左から右へと行くと列が増加し、上から下へと行くと行が増加する。それゆえに、上記行列Aは、0から3まで行及び列のidを有する。
x及びyを、それぞれ、サイズn及びmの2つの密ベクトルとする。SpMVのタスクは、y=Axを計算することであり、Aは疎行列である。図1は、1つの実施形態(従来技術)にかかる、圧縮された疎行フォーマット符号化において符号化された行列に対してSpMVを実行するための方法10を示すフロー図である。i=0からm−1のidを有する行である行列の全ての行を処理する反復処理ループ(ステップ11〜19)が開始される(ステップ11)。A配列におけるi行目のエントリが配置され、行列におけるその行についての第1の非ゼロエントリのA配列におけるインデックスを識別し、識別されたインデックスは、変数jとして設定される(ステップ12)。A配列における次の(i+1)行目のエントリが配置され、変数jmaxとして設定される(ステップ13)。i行目が行列における最後の行でない限り、A配列における次のエントリは、(i+1)行目における第1の非ゼロエントリのA配列におけるインデックスである。i行目が行列における最後の行である場合、A配列における次のエントリは、A配列におけるエントリの総数である。以下のステップ16において記載される非ゼロ配列の値の乗算結果をともに加算する関数である総和計算部(以下の擬似コードにおける総和積算器とも称される)は、ゼロにおいて合計値を設定することによって初期化される(ステップ14)。jがjmax未満である場合(ステップ15)、以下の式による計算
sum←sum+A[j]×x[A[j]]
が実行される(ステップ16)。計算において、インデックスjを有するA配列における値はxの要素によって乗算され、そのインデックスは、j番目のインデックスを有するA配列の数である。乗算結果は、ステップ16の前の反復中に実行された乗算結果の合計に加算される。上記ステップ14において設定されたように、本方法におけるステップ16の第1の反復中において合計はゼロである。計算が終了すると、jの値に1が加算され、加算結果はjに設定され(ステップ17)、その行における次の列のエントリに処理を移動する。本方法は、上述したステップ15に戻り、i行目における非ゼロ値が処理されるまでステップ15〜17を繰り返す。jがjmax以上である場合(ステップ15)、ループ15〜17における反復中に乗算結果を加算した合計は、密ベクトルyに格納される(ステップ18)。反復処理ループは、次の行に移動し(ステップ19)、全ての行が処理されるまでループ(ステップ11〜19)を介して処理を継続し、方法10は終了する。
図1の方法10はまた、以下の擬似コードを使用して表すことができる。
for i=0からm−1 /*疎行列Aのm行に対するループ*/
j←A[i] /*j:i行目における第1の非ゼロエントリのAにおけるインデックス*/
max←A[i+1] /*jmax:(i+1)行目における第1の非ゼロエントリのAにおけるインデックス*/
sum←0 /*総和積算器を初期化*/
while (j<jmax) /*行の終わりに到達したかどうかのテスト/*
sum←sum+A[j]×x[A[j]] /*yi=ΣjAi,j×xjを計算*/
j←j+1 /*i行目における次の非ゼロの列に移動*/
while文終了
y[i]←sum /*yに結果を格納*/
for文終了
SpMVの最も計算集約的な部分は、上記ステップ16において記載された総和行に起因する。
sum←sum+A[j]×x[A[j]]
ここで、それぞれインデックスj、j及びA[j]を有するA、A及びxの3つの配列がアクセスされる。A及びAは、双方とも、通常はSpMVにおいてA及びxのサイズよりもはるかに大きいe個の要素を有することに留意されたい。CPU及びGPUの双方を含む現代のプロセッサにおけるこのサイズ差の意味は、上記総和行がこれらの配列のサイズに応じて2つ又は3つのキャッシュミスを引き起こす可能性が最も高いということである。3つの配列のうち、x(入力密ベクトルを格納する配列)は、通常最小であり、それゆえに、それは、プロセッサのL2キャッシュに適合する最善の機会を有する。もちろん、行列Aが十分に小さい場合、全てが適合し、自明なケースである。しかしながら、非自明のSpMVの問題について、A又はAのいずれもL2に収まることを仮定するべきではない。換言すれば、A及びAの双方にアクセスすることは、2つの別個のキャッシュミスをトリガすることができ、SpMVの性能に悪影響を与えることがある。キャッシュミスは、SpMVの速度を大幅に低減することができ、大抵の場合にメモリ参照の局所性をほとんど示さない。
疎行列はまた、大抵の場合に「転置」のCSRと思われる圧縮された疎列(CSC)フォーマットで符号化されることができる。Aをe個の非ゼロエントリを有するm×nの疎行列とする。A’、A’及びA’を、それぞれ、長さe、e及び(n+1)の3つの1次元配列とする。CSCフォーマットにおいて、Aは、<A’,A’,A’>として符号化され、ここで、
・値配列A’は、列優先順序でAにおける全ての非ゼロエントリの値を保持し、
・列配列A’は、列優先順序でAにおける全ての非ゼロエントリの行を保持し、
・行インデックス配列A’は、Aにおける各列の第1の非ゼロエントリのA’におけるインデックスを保持し、A’配列におけるエントリの総数であるA’配列における最後のエントリを有する。
先に示された同じ4×4行列
Figure 2016119084
は、<A’=[0.5,0.5,0.5,1,1,0.5]、A’=[1,2,2,3,0,1]、A’=[0,2,3,4,6]>としてCSCにおいて符号化されることができる。
x’及びy’を、それぞれ、サイズm及びnの2つの密ベクトルとする。SpMTVのタスクは、y’=ATx’を計算することであり、Aは元の転置されない疎行列である。図2は、1つの実施形態(従来技術)にかかる、圧縮された疎列フォーマット符号化において符号化された行列に対してSpMTVを実行するための方法20を示すフロー図である。j=0からn−1のidを有する列である行列の全ての列を処理する反復処理ループ(ステップ21〜29)が開始される(ステップ21)。A’配列におけるj列目のエントリが配置され、行列におけるその列についての第1の非ゼロエントリのA’配列におけるインデックスを識別し、識別されたインデックスは、変数iとして設定される(ステップ22)。A’配列における次の(j+1)列目のエントリが配置され、変数imaxとして設定される(ステップ23)。j列目が行列における最後の列でない限り、A’配列における次のエントリは、(j+1)列目における第1の非ゼロエントリのA’配列におけるインデックスである。j列目が行列における最後の列である場合、A’配列における次のエントリは、A’配列におけるエントリの総数である。以下のステップ26において記載される非ゼロ配列の値の乗算結果をともに加算する関数である総和計算部は、ゼロにおいて合計値を設定することによって初期化される(ステップ24)。iがimax未満である場合(ステップ25)、以下の式による計算
sum←sum+A’[i]×x’[A’[i]]
が実行される(ステップ26)。計算において、インデックスiを有するA’配列における値はx’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’配列の数である。乗算結果は、ステップ26の前の反復中に実行された乗算結果の合計に加算される。上記ステップ24において設定されたように、本方法におけるステップ26の第1の反復中において合計はゼロである。計算が終了すると、iの値に1が加算され、加算結果はiに設定され(ステップ27)、その列における次の行のエントリに処理を移動する。方法20は、上述したステップ25に戻り、j列目における非ゼロ値が処理されるまでステップ25〜27を繰り返す。iがimax以上である場合(ステップ25)、ループ25〜27における反復中に乗算結果を加算した合計は、密ベクトルy’に格納される(ステップ28)。反復処理ループは、次の列に移動し(ステップ29)、全ての列が処理されるまでループ(ステップ21〜29)を介して処理を継続し、方法20は終了する。図2の方法20はまた、以下の擬似コードを使用して表すことができる。
for j=0からn−1 /*疎行列Aのn列に対するループ*/
i←A’[j] /*i:j列目における第1の非ゼロエントリのA’におけるインデックス*/
max←A’[j+1] /*imax:(j+1)列目における第1の非ゼロエントリのA’におけるインデックス*/
sum←0 /*総和積算器を初期化*/
while (i<imax) /*列の終わりに到達したかどうかのテスト/*
sum←sum+A’[i]×x’[A’[i]] /*y’j=ΣiAi,j×x’iを計算*/
i←i+1 /*j列目における次の非ゼロの行に移動*/
while文終了
y’[j]←sum /*y’に結果を格納*/
for文終了
上記総和行sum←sum+A’[i]×x’[A’[i]]は、それぞれインデックスi、i及びA’[i]を有するA’、A’、x’の3つの配列にアクセスする。CSRの場合と同様に、3つの配列のアクセスは、図2の方法20のステップ25〜27の1回の反復において3つのキャッシュミスを引き起こす可能性がある。
一般に疎行列のキャッシュミスの量及び処理速度は、疎行列の改善された符号化を使用することによって低減することができる。図3は、1つの実施形態にかかる、疎行列の効率的な表現及び処理のためのコンピュータ実装システム30を示すブロック図である。システム30は、1つ以上の疎行列32を記憶するデータベースなどのコンピュータ記憶装置31を含む。記憶装置は、行列32の圧縮表現35を用意する行列エンコーダ34を実行する1つ以上のサーバ33に接続されている。各表現は、行列における非ゼロエントリに関する情報を記憶する1つ以上のアレイを含む。以下の説明において、用語「表現」及び「符号化」は、互換的に使用される。表現35は、記憶装置31に記憶されることができる。各圧縮表現35は、行列32にゼロエントリを明示的に含まない圧縮フォーマットにおける行列32の符号化である。例えば、そのような表現35は、上述したCSR及びCSCフォーマットにおける符号化を含むことができる。行列エンコーダ34はまた、処理中におけるキャッシュミスの数を低減し、それゆえによりキャッシュフレンドリーである行列のより効率的な表現35を符号化することができる。
これらのより効率的な符号化の1つは、構造化CSR(SCSR)符号化と称することができる。上述したCSR符号化において、A及びAは、Aにおける要素がAにおける要素と同じバイトサイズを有しない場合にそれらの空間的要件が異なるものの、それらにおける同数の要素を有する。同数の要素を正確に有する値(A)及び列(A)配列は、それらが疎行列における非ゼロエントリの同じセットの異なる態様を記載することから一致しない。その結果、それらは、いかなるCSR符号化においても同数の要素を有する必要がある。
2つの別個の配列に同じ非ゼロエントリの値及び列を格納する代わりに、構造化CSRの符号化方式は、i番目の要素が疎行列におけるi番目の非ゼロエントリの値のみならず列も記憶するように、対(値、列)を含む単一の統一された配列で値及び列を置く。換言すれば、CSR符号化における配列A及びAは、以下では値−列配列と称する複合配列Avcを形成するように互いに結合されている。行インデックス配列Aは、Aにおける各行の第1の非ゼロエントリのAvcにインデックスを記憶し、A配列における最後のエントリが配列Avcにおける複合エントリの総数である。それゆえに、構造化CSR符号化は、2つの配列A=<Avc,A>を含む。SCSR符号化の名称は、値−列配列Avcにおける各要素は、浮動小数又は整数などの基本的なタイプである代わりに、構造(すなわち、複合データタイプ)であることを反映する。
上記に示された例としての行列
Figure 2016119084
は、以下のように構造化CSRにおいて符号化されることができる。A=<Avc,A>、ここで、Avc=[<1,3>,<0.5,0>,<0.5,3>,<0.5,0>,<0.5,1>,<1,2>]及びA=[0,1,3,5,6]。
同様に、エンコーダ34は、CSCと比較してSpMTVなどの処理中にキャッシュミスの数を低減することができる構造化CSC(SCSC)符号化を形成することができる。構造化CSCにおいて、CSCの値配列A’及び行配列A’は、値−列配列と称することができる単一の複合配列A’vrにマージされる。値−列配列A’vrの各要素は、行列の左端列から始まる列優先順序に基づく非ゼロ行列エントリの値及び列の双方を記憶する<値,行>対である。列インデックス配列A’は、Aにおける各列の第1の非ゼロエントリのA’vrにインデックスを保持し、A’配列における最後のエントリは、A’vr配列における複合エントリの総数である。
例えば、先に示された同じ行列
Figure 2016119084
は、以下のように構造化CSCにおいて符号化されることができる。A=<A’vr,A’>、ここで、A’vr=[<0.5,1>,<0.5,2>,<0.5,2>,<1,3>,<1,0>,<0.5,1>]及びA’=[0,2,3,4,6]。
サーバ33は、さらに、行列32の圧縮表現35に対する計算を行うことによって行列32を処理する計算モジュール36を含む。それゆえに、計算モジュール36は、さらに図7及び図8を参照して後述するように、行列32の構造化CSC及び構造化CSR符号化に対してSpMV及びSpMTVを実行することができる。また、計算モジュール36は、多くの用途に向けてSpMV及びSpMTVの結果を適用することができる。例えば、SpMV及びSpMTVの結果の1つの適用は、さらに図23を参照して記載されるようにページランク(登録商標)アルゴリズムなどのランク付けアルゴリズムを実行することができる。ページランクアルゴリズムについてのグラフと疎行列との間の接続が以下に簡単に検討される。グラフは、通常は疎であり、それゆえに大抵の場合にCSR/CSCフォーマットで符号化される隣接行列として表現することができることがよく知られている。図4は、一例として、隣接行列を示すグラフ40である。例えば、行列
Figure 2016119084
は、図4におけるグラフ40の行列表現としてみることができる。上記グラフは、以下の4つの頂点及び6つの重み付けエッジを有する。<0,3,1.0>,<1,0,0.5>,<1,3,0.5>,<2,0,0.5>,<2,1,0.5>,<3,2,1.0>、ここで、フォーム<u,v,w>のエッジは、uがエッジの元頂点であり且つvがエッジの先頂点であるように、重みwを有する頂点uからvまでのリンクを示す。これらの6つのエッジが<行,列,値>トリプルの形態で疎行列の非ゼロエントリとして表示される場合、それらが正確にAにおける非ゼロエントリであることを確認することができる。それゆえに、重み付きグラフと疎行列との間に1対1の対応がある。グラフ40は、エッジの重みがその元頂点から出るリンクの数の逆数であるという興味深い性質を有する。そのようなグラフは、頂点間の遷移確率をモデル化し、それは、さらに図23を参照しながら後述するようなページランク(登録商標)アルゴリズムを計算するために使用することができる。他のアルゴリズムもまた、さらに後述するように計算モジュール36によって実行することができる。
構造化CSC及び構造化CSRは、任意のタイプのハードウェアを使用して行列32の処理を高速化することを可能とする。図3を参照すると、エンコーダ34はまた、処理のためにGPUを使用するサーバ33において行列処理を高速化するために特に適している圧縮表現35の他のタイプを形成することが可能である。SpMV又はSpMTVなどのGPUベースの疎行列処理について、サーバ33は、SpMV処理カーネル、後述するものなどのGPUの処理アルゴリズムを実装するサーバにおける1つ以上のGPUにおいて実行される処理スレッドによって実現される機能を必要とする。しかしながら、上述したように、実験は、単一のカーネルが大抵の場合には疎行列行を有することができる非ゼロエントリの数における大きな変動に応答しないことから、単一のカーネルが最適から遠いことを示している。
サーバ33は、スレッドの並列グループ内で実行することが可能である。例えば、サーバ33は、ワープと称される単位にグループ化された複数のGPUスレッドを含むことができる。例えば、カリフォルニア州サンタクララのNvidia(登録商標)社製のGPUの並列スレッドは、同じワープにおける全てのスレッドが単一のストリーミングマルチプロセッサ(SM)を共有し且つ単一命令複数スレッド(SIMT)実行モデルを想定するように、ワープと称される32の単位にグループ化される。スレッドがカリフォルニア州サニーベールのアドバンスト・マイクロ・デバイス(登録商標)社製のGPUにおいて実行される場合、ワープと同等の概念は、64個のGPUコアのグループにおいて実行される64スレッドを現在含むウェーブフロントと称される。その結果、同じワープ(又はウェーブフロント)内の異なるスレッドは、それらが常に並行して異なるデータを処理するのが可能とされるものの、同時に異なる命令を実行することができない。本特許出願の目的のために、用語「ワープ」及び「ウェーブフロント」は、単一のストリーミングマルチプロセッサを共有し且つ同じ命令を実行する32又は64の複数の処理スレッドを並列に実行するGPUコアの集合を指すように以下において互換的に使用される。さらなる実施形態において、他のスレッド数は、ワープにおいて可能である。さらに、スレッドの大きなグループ化は、ワープのグループ化によって可能である。そのような各グループは、スレッドブロックと称され、スレッドの各ブロックは、(実施形態に応じて32又は64以上)ワープよりも多くのスレッドを含む。各ブロック内のスレッド数は、ブロックサイズ(又は以下の説明においてはBLOCKSIZE)と称される一方で、ワープ内のスレッド数は、ワープサイズ(又は以下の説明においてはWARPSIZE)と称される。エンコーダ34は、サーバ33が異なる数の非ゼロエントリを有する行列32の部分について異なるカーネルを実行するのを可能とする行列32の符号化を形成することによってサーバ33のハードウェアの限界を克服するのを可能とする。エンコーダ34は、複数のパーティションを含むパーティション化された圧縮表現で行列32のパーティション化された圧縮表現を作成する。各パーティションは、所定範囲内の複数の非ゼロエントリを有する行又は列などの行列32の部分のグループのリストである。例えば、1つのパーティションは、1〜31の非ゼロエントリを有する行列32の部分を有することができ、2番目のパーティションは、32から1024の非ゼロエントリを有する行列32の部分を有することができ、3番目のパーティションは、1024以上の非ゼロエントリを有する行列のエントリ部分を有することができる。非ゼロエントリの他の範囲及びパーティションの他の数も可能である。エンコーダ34は、以下のようにパーティション化された圧縮表現を定義する。以下の説明は、行であるとしてパーティション化されている行列32の部分に言及しているが、行列の列は、同じ方法を準用してパーティション化されることができる。
kをゼロから始まるパーティションのインデックスとし、少数の非ゼロエントリを有する行のパーティションが多い非ゼロエントリを有するパーティションよりも低いインデックスを有する。それゆえに、32から1024の非ゼロエントリを有する行のパーティションは、1024超の非ゼロエントリを有するパーティションよりも低いインデックスkを有するであろう。pをパーティションの総数とし、Asを、以下のように、各パーティションにおいて許容される行の非ゼロエントリの最小及び最大数を指定する行パーティション区切り配列と称される(p+1)要素の整数配列とする。
・As[0]=1
・k=0,1,2,・・・,p−1について、As[k]<As[k+1]
・As[p]=∞
1つの実施形態において、As[0]及びAs[p]の値は、それらが行列32の行が有することができる非ゼロエントリの最大及び最小数であることから、記憶装置に記憶されない。さらなる実施形態において、値は記憶装置31に記憶される。
パーティションは、順序を維持する。ei〜を疎行列の行iにおける非ゼロエントリの数とする。Rkとして示されるk行目の行ベースのパーティションにおける行のセットは、Rk={i|As[k]≦ei〜<As[k+1]}のように記載することができる。元の行列32の内容を変更しないことが大抵の場合には望ましいことから、以下ではマッピング配列とも称される1次元の順序維持置換配列(A又はA’)は、行のランク又は列のランクと称されるパーティションベースの行idからそれぞれ元の行又は列idへのマッピングを格納するために追加される。ランクのそれぞれは、行列の部分のマッピング配列におけるインデックスである。例えば、行が3のid及び0のランクを有する場合、値3は、マッピング配列にリスト化された第1の値である。通常の置換配列とは異なり、順序維持置換配列は、以下のように同じパーティションに割り当てられた行の相対的な順序を維持するために必要とされる。ri及びrjを置換配列における行i及びjのランクとする。pi及びpjをこれら2行i及びjが属するパーティションとする。ランクri及びrjは、以下の制約が満たされた場合に限り、順序を維持している。
・∀pi<pj,ri<rj−パーティションjのインデックスがパーティションiのインデックスよりも大きい場合、パーティションjにおける行のランクは、パーティションiにおける行のランクよりも大きくなければならない。
・∀pi>pj,ri>rj−パーティションiのインデックスがパーティションjのインデックスよりも大きい場合、パーティションiにおける行のランクは、パーティションjにおける行のランクよりも大きくなければならない。
・∀pi=pj∧i≠j,ri<rj⇔i<j(又は同等に、ri>rj⇔i>j)
上記最後の制約は、小さいidを有する行が同じパーティションにおいてより大きなidを有する行の前にアクセスされるように、元の行列32における同じパーティションの行の相対的な順序を尊重するように設計されている。制約は、構造化及び非構造化CSR符号化の双方においてAにアクセスする際のランダム性を低減させることから、制約は、SpMVなどのSLAアルゴリズムのキャッシュ性能を向上させる。有益な副作用として、順序維持の制約を強制することもまた、CSR行列を完全にパーティション化するために行インデックス配列Aの単一のキャッシュフレンドリーな線形スキャンを必要とするのみであることから、パーティション化アルゴリズムの時間的複雑性を低減する。一方、ソートベースのアルゴリズムは、かなりの数のキャッシュミスをもたらすことがあるO(mlogm)比較を平均的に負担する。ここで、mは行数である。実験は、順序維持パーティションがはるかに高速に計算できることのみならず、それらはまた、通常順序維持ではないソートベースのパーティションよりも大幅に高速なSpMVをもたらすことを示す。
行列32の順序維持のパーティション化された圧縮表現は、CSC、CSR、構造化CSC及び構造化CSRなどの行列32の様々な既存の圧縮表現に基づいて図9から始めて以下に記載されるようなエンコーダ34によって作成することができる。異なる圧縮表現に基づくパーティション化された圧縮表現は互いに異なる。例えば、疎行列Aは、構造化CSRについてA=<Avc,A,A,Ao,Ap,As>として表現されることができる。ここで、Avc及びAは、前と同じであり、Aは、行パーティションマッピング配列であり、Aoは、パーティションのそれぞれの第1の部分のランク及びマッピング配列におけるエントリの総数を含む行パーティションオフセット配列であり、Apは、行ベースのパーティションの数であり、Asは、(必要に応じてパーティション化が終了した場合)行パーティション区切り配列である。行ベースのパーティションを有する構造化CSRは、この時点からパーティション化された構造化CSR(PSCSR)と称される。非構造化CSR行列は、同様に図9を参照して以下に記載されるルーチンを使用してパーティション化されることができるため、得られた符号化は、疎行列Aが<A,A,A,A,Ao,Ap,As>として符号化されるパーティション化CSR(PCSR)と称される。
同様に、疎行列Aは、構造化CSCについてA=<A’vr,A’c,A’,A’,A’,A’>として表現することができる。ここで、A’vr及びA’は、前と同じであり、A’は、行パーティションマッピング配列であり、A’は、行パーティションオフセット配列であり、A’は、列ベースのパーティション数であり、A’は、(必要に応じてパーティションが終了した場合)列パーティション区切り配列である。列ベースのパーティションを有する構造化CSCは、この時点からパーティション化された構造化CSC(PSCSC)と称される。非構造化CSC行列は、同様に同じアルゴリズムを使用してパーティション化されることができるため、得られた符号化は、疎行列Aが<A’,A’,A’,A’,A’,A’,A’>として符号化されるパーティション化CSC(PCSC)と称される。
例えば、同一の例として先に示された行列
Figure 2016119084
を考える。
1.第1のパーティションが単一の非ゼロエントリを有する行のみを含み、
2.第2のパーティションが複数の非ゼロエントリを有する行を含む
ように、Aについて2つのパーティションを作成したい場合、A=A’ =2及びA=A’=[1,2,∞]である。他の配列は、以下のとおりである。
・PCSR及びPSCSRについて、A=[0,3,1,2]及びA=[0,2,4]
・PCSC及びPSCSCについて、A’=[1,2,0,3]及びA’=[0,2,4]
システム30のハードウェアリソースが表1に関連して以下に説明される式を使用して決定することができることを考えると、パーティション化された圧縮表現を作成するかどうかは実用的である。パーティション化された圧縮符号化表現は、サーバ33が1つ以上のGPU又は1つ以上のCPUのみを含むかどうかにかかわらず、図9を参照して始めて以下に記載されるものなどの計算モジュール36によって処理することができる。しかしながら、所定数の非ゼロエントリを有する行列32の部分をリスト化するパーティションの作成は、サーバ33がGPUを含む場合、計算モジュール36がそれらの部分を処理するときに最も適切な処理カーネルを適用するのを可能とする。多くの処理カーネルが可能であるが、計算モジュールは、少なくとも以下の3つのカーネルを適用することができる。
1.行又は列を処理するために単一の処理スレッドを割り当てる1スレッド1行(1T1R)カーネル
2.行又は列を処理するためにワープと称されるユニット内の全てのスレッドを割り当てる1ワープ1行(1W1R)カーネル
それ自体では、いずれのカーネルも、行列32の全ての部分について理想的ではない。上述したように、総ランタイムは、最大数の非ゼロエントリを有する行(又は列)に対応する最も遅いスレッドに依存することから、(以下ではf1T1Rカーネルとも称されることができる)1T1Rカーネルは負荷不均衡を被る。行(又は列)における非ゼロエントリ数が32(又は64)未満のワープ(又はウェーブフロント)サイズである場合、(以下ではf1W1Rカーネルとも称されることができる)1W1Rカーネルは、ハードウェアリソースの浪費を被る。単一の疎行列は、少数の行及び多くの非ゼロエントリを有する行の双方を有することができることに留意されたい。それゆえに、いずれのタイプのカーネルにコミットすることも、めったに問題を解決しない。さらにまた、実験は、同じ行に対して動作するように単一のワープにおけるものよりも多くのスレッドを有する利益があることを示す。これは、第3のタイプを追加する必要性を引き起こす。
3.行又は列を処理するためにスレッドのブロック(32又は64より多い)を割り当てる1ブロック1行(1B1R)カーネル
1W1Rカーネルと同様に、(以下ではf1B1Rカーネルとも称されることができる)1B1Rカーネルはまた、スレッドの数が行列のその部分における非ゼロエントリの数よりも大きい場合にハードウェアリソースの浪費を被ることがある。しかしながら、行列の特定のパーティションと各カーネルを相関させることにより、計算モジュールは、さらに図15を参照して以下に記載されるように、特定の行又は列を処理するための最良のカーネルを使用することができる。
カーネル関数のさらに他の種類が可能である。
f1T1R、f1W1R及びf1B1R又は効率的に他のカーネルを混合するために、SpMVアルゴリズムは、各カーネルが得意である行を迅速に区別する必要がある。高速な行分類について、疎行列パーティション化アルゴリズムは、図9を参照して始めて以下に記載され、同じカーネルによって最良に処理される行列の行をグループ化する。より正確には、さらに図9を参照して始めて以下に記載されるように、パーティション化アルゴリズムは、入力として、各パーティションにおける行についての非ゼロエントリの最小及び最大数をとり、各行又は各列が属するパーティションを決定するために使用することができるマッピングを生成する。さらなる実施形態において、カーネル関数は、パーティションにおける非ゼロエントリ数以外の要因に基づいてパーティションに割り当てることができる。
1つ以上のサーバ33は、インターネット又はセルラーネットワークなどのローカルネットワーク又はインターネットワークとすることができるネットワーク37に接続され、ネットワークを介して少なくとも1つのユーザ装置38と通信することができる。ユーザ装置38は、デスクトップコンピュータとして示されているが、ユーザ装置38はまた、ラップトップコンピュータ、スマートフォン、メディアプレーヤ及びタブレットを含むことができる。ユーザ装置38のさらに他の種類も可能である。ユーザ装置38は、ネットワーク15を介してサーバ33と通信し、計算実行からコマンドを受信し、ユーザ装置に計算結果を出力し返すことができる。
SpMV及びSpMTVは、(PS)CSR及び(PS)CSC符号化の上に構築することができる多くのSLAオペレーションの2つだけである。一般に、疎行列を扱う任意のアルゴリズムは、特に、アルゴリズムが行列内の非ゼロエントリの値及び位置の双方に同時にアクセスする必要がある場合、パーティション化された(P)及び/又は構造化された(S)CSR/CSC符号化から利益を得ることができる。疎行列−行列乗算(SpMM)などの他のSLAオペレーションは、同様のアクセスパターンを有し、それゆえに、ここで紹介したのと同じ符号化方式から利益を得ることができる。
しかしながら、従来のCSR/CSC符号化が構造化することなしでさえも良好に作動することができ且つシステム30が従来のCSR/CSC符号化を使用可能であるタスクがある。例えば、以下のように定義される疎行列Aのフロベニウスノルム演算
Figure 2016119084
は、非ゼロエントリの位置にアクセスするためのアルゴリズムを必要とせず、それゆえに、同じ配列におけるそれらの値及び位置の混合は常に役立たないことがある。他の例は、非ゼロエントリの位置ではなくそれらの値にアクセスすることのみを必要とする疎行列が対角であるか否かをテストすることである。一般に、非ゼロ行列エントリの値又は位置(双方ではない)のいずれかにアクセスする必要があるのみである任意のアルゴリズムは、構造化CSR/CSCから利益を受けないことがある。上記の双方の例は、データプリフェッチが有効な場合、いくつかのキャッシュミスを通常もたらすか又はもたらさない優れたメモリ参照局所性を既に有することに留意されたい。換言すれば、最悪の場合でも、それらの構造化されていない相手方を失うことは構造化CSR/CSCについて多くはなく、一般に構造化CSR/CSC符号化は、これらのシナリオにおいては非構造化CSR/CSC符号化を実行する。それにもかかわらず、構造化CSR/CSCは、一般に、CPUにおいて疎行列を符号化するための最も効率的且つ堅牢なデータ構造である可能性が高く且つ大部分の状況において最も有用であることから、従来のCSR/CSCを使用するために(既存のコードを除く)理由はほとんどない。図5は、図3のシステム30の様々なハードウェアセットアップのための様々な符号化の強みをまとめたグラフ50である。グラフ50は、大部分のシナリオにおいて、上記符号化方式が従来の非構造化CSR/CSC符号化よりも優れていることを図示している。グラフ50は、(1)x軸としてハードウェアの柔軟性及び(2)y軸としてメモリアクセスパターンの2つの直交次元を含む。y軸について、グラフ50は、(下方において「値
Figure 2016119084
位置」としてラベル付けされた(シンボル
Figure 2016119084
排他的OR関係を意味する))非ゼロエントリの値又は位置のいずれか(双方ではない)にアクセスする必要があるのみであるパターンに対して、(上方において「値+位置」としてラベル付けされた)非ゼロ行列エントリの値及び位置の双方に同時にアクセスする必要があるメモリアクセスパターンを区別する。
グラフ50は、4つの象限のうち3つが本特許出願において導入された疎行列符号化を好み、非ゼロエントリについてメモリアクセスパターンが値のみ又は位置のみのいずれかである場合及びハードウェアがより多くのCPUのようである場合には従来の非構造化CSR/CSCフォーマットを好む唯一のシナリオがあることを示している。先に説明したように、従来のCSR/CSCフォーマットによって達成可能な性能向上は、全ての場合に非常に限られている。一方、他の3つの象限、特に右上隅におけるその性能低下は大幅であり得て、全ての状況について上述した符号化の使用を正当化する。
図3を参照すると、1つ以上のサーバ33は、他の要素が可能であるものの、1つ以上のCPU及びGPU及びSM、メモリ、入力/出力ポート、ネットワークインターフェース並びに不揮発性記憶装置などのプログラム可能な計算装置において従来みられる要素を含むことができる。サーバ33は、それぞれ、本願明細書に開示される実施形態を実施するための1つ以上のモジュールを含むことができる。モジュールは、従来のプログラミング言語におけるソースコードとして書かれたコンピュータプログラム又はプロシージャとして実装することができ、それは、オブジェクト又はバイトコードとして中央処理装置による実行のために提示される。あるいは、モジュールはまた、集積回路として又は読み出し専用メモリ要素に焼き付けられたものとしてハードウェアに実装されることができ、サーバ33のそれぞれは、専用コンピュータとして機能することができる。例えば、モジュールがハードウェアとして実装される場合、その特定のハードウェアは、その目的のために使用することができないハードウェアなしでパーティション化及び他のコンピュータを実行するために特化されている。ソースコード及びオブジェクト及びバイトコードの様々な実装は、フロッピー(登録商標)ディスク、ハードドライブ、ディジタルビデオディスク(DVD)、ランダムアクセスメモリ(RAM)、読み出し専用メモリ(ROM)及び同様の記憶媒体などのコンピュータ読み取り可能な記憶媒体に保持することができる。他の種類のモジュール及びモジュール機能のみならず、他の物理的なハードウェア要素も可能である。
上述したように、構造化CSR及び構造化CSCは、SpMVなどの行列処理中におけるキャッシュミス数を低減するのを可能とする。図6は、1つの実施形態にかかる構造化行列表現を取得するためのコンピュータ実装方法60を示すフロー図である。他の実装が可能であるものの、方法60は、図3のシステム30に実装することができる。最初に、疎行列がアクセスされ、行列の部分、行及び列における非ゼロエントリについての情報が取得される(ステップ61)。行列の構造化圧縮表現、行列の構造化CSC又は構造化CSRのいずれかは、上述したように、複合配列(Avc又はA’vr)及びインデックス配列(A又はA’)に非ゼロエントリに関する情報を符号化することによって作成される(ステップ62)。他の種類の処理も可能であるが、図7を参照して後述するSpMV、又は、図8を参照して後述するSpMTVを実行することなどによって構造化圧縮表現が処理される(ステップ63)。処理結果は、処理結果の他の種類のアプリケーションも可能であるが、ページランク(登録商標)又は他のランク付けアルゴリズムを実行することなどによって適用される(ステップ64)。方法60は終了する。
行列の構造化圧縮表現に対して実行されるSpMVは、キャッシュミスの確率を低減する。図7は、図6の方法60において使用するために構造化CSR符号化を使用して符号化された行列に対してSpMVを実行するためのルーチン70である。avc∈Avcを値−列配列の要素とする。avcが構造であることから、以下のように配列メンバにアクセスするためにC++及びJava(登録商標)などの一般的なプログラミング言語におけるようなドット演算子が使用可能である。avc・vは、avcの「値」フィールドを返し、avc・cは、avcの「列」フィールドを返し、それらは後述するステップ76における計算に使用される。値及び列フィールドにアクセスする他の方法も可能である。ルーチン70及び後述する後続のSpMVルーチンを説明するために、上記例において使用された同じm×n行列及びサイズn及びmの2つの密ベクトルx及びyがそれぞれ使用される。
行列の全ての行であるi=0からm−1のidを有する行を処理する反復処理ループ(ステップ71〜79)が開始される(ステップ71)。A配列におけるi行目についてのエントリが配置され、行列のその行についての第1の非ゼロエントリのAvc配列におけるインデックスを識別し、識別されたインデックスは、変数jとして設定される(ステップ72)。A配列における次の(i+1)エントリが配置され、変数jmaxとして設定される(ステップ73)。i行目が行列における最後の行でない限り、A配列における次のエントリは、(i+1)行目における第1の非ゼロエントリのAvc配列におけるインデックスであり、i行目が行列における最後の行である場合、A配列における次のエントリは、Avc配列におけるエントリの総数である。ステップ76において後述する非ゼロ配列の値の乗算結果を加算する関数である総和計算部は、ゼロにおける合計の値を設定することによって初期化される(ステップ74)。jがjmax未満である場合(ステップ75)、以下の式
sum←sum+Avc[j].v×x[Avc[j].c]
にかかる計算が行われる(ステップ76)。計算において、インデックスjを有するAvc配列の要素に格納された値は、そのインデックスがj番目のインデックスを有するAvc要素に格納された列のインデックスであるxの要素によって乗算される。乗算結果は、ステップ76の前の反復中に実行される乗算結果の合計に加算され、ルーチンにおけるステップ76の最初の反復中において、合計は、上記ステップ74において設定されたようにゼロである。計算が終了すると、jの値に1が加算され、加算結果は、jとして設定され(ステップ77)、その行における次の列のエントリに処理を移動する。方法は、上述したステップ75に戻り、i行目における非ゼロ値が処理されるまで、ステップ75〜77を繰り返す。jがjmax以上である場合(ステップ75)、ループ75〜77における反復中の乗算結果を加算した合計は、密ベクトルyに格納される(ステップ78)。反復処理ループは、次の行に移動し(ステップ79)、全ての行がループ(ステップ71〜79)を介して処理されるまで、ループを介した行の処理(ステップ71〜79)が継続し、その後にルーチン70は終了する。ルーチン70はまた、図1を参照して上記に示された擬似コードとの差異を示すテキストボックスにより、以下の擬似コードを使用して表すことができる。
for i=0からm―1 /*疎行列Aのm行にわたるループ*/
j←A[i] /*j:i行目における第1の非ゼロエントリのAvcにおけるインデックス*/
max←A[i+1] /*jmax:(i+1)行目における第1の非ゼロエントリのAvcにおけるインデックス*/
sum←0 /*総和積算器を初期化*/
while (j<jmax) /*行の終わりに到達したかどうかをテスト*/
Figure 2016119084
j←j+1 /*i行目における次の非ゼロ列まで移動*/
while文終了
y[i]←sum /*yに結果を格納*/
for文終了
構造化CSRに対して実行されるSpMVと同様に、構造化CSCに対してSpMTVを実行することは、処理が非構造化CSCに対して実行される場合よりもキャッシュミスが発生する可能性を低減する。図8は、図6の方法60において使用するために構造化CSC符号化を使用して符号化された行列に対してSpMTVを実行するためのルーチン80である。a’vr∈A’vrを値−行配列の要素とする。a’vrが構造であることから、以下のように配列メンバにアクセスするためにC++及びJava(登録商標)などの一般的なプログラミング言語におけるようなドット演算子が使用可能である。a’vr・vは、a’vrの「値」フィールドを返し、a’vr・rは、a’vrの「行」フィールドを返し、それらは後述するステップ86における計算に使用される。値及び行フィールドにアクセスする他の方法も可能である。ルーチン80及び以下の他のSpMTVルーチンを説明するために、上記例において使用された同じm×n行列及びサイズm及びnの2つの密ベクトルx’及びy’がそれぞれ使用される。
行列の全ての列であるj=0からn−1のidを有する列を処理する反復処理ループ(ステップ81〜89)が開始される(ステップ81)。A’配列におけるj列目についてのエントリが配置され、行列のその列についての第1の非ゼロエントリのA’vr配列におけるインデックスを識別し、識別されたインデックスは、変数iとして設定される(ステップ82)。A’配列における次の(j+1)エントリが配置され、変数imaxとして設定される(ステップ83)。j列目が行列における最後の列でない限り、A’配列における次のエントリは、(j+1)列目における第1の非ゼロエントリのA’vr配列におけるインデックスであり、j列目が行列における最後の列である場合、A’配列における次のエントリは、A’vr配列におけるエントリの総数である。ステップ86において後述する非ゼロ配列の値の乗算結果を加算する関数である総和計算部は、ゼロにおける合計の値を設定することによって初期化される(ステップ84)。iがimax未満である場合(ステップ85)、以下の式
sum←sum+A’vr[i].v×x’[A’vr[i].r]
にかかる計算が行われる(ステップ86)。計算において、インデックスiを有するA’vr配列の要素に格納された値は、そのインデックスがi番目のインデックスを有するA’vr要素に格納された行のインデックスであるx’の要素によって乗算される。乗算結果は、ステップ86の前の反復中に実行される乗算結果の合計に加算され、本方法におけるステップ86の最初の反復中において、合計は、上記ステップ84において設定されたようにゼロである。計算が終了すると、iの値に1が加算され、加算結果は、iとして設定され(ステップ87)、その列における次の行のエントリに処理を移動する。ルーチン80は、上述したステップ85に戻り、j列目における非ゼロ値が処理されるまで、ステップ85〜87を繰り返す。iがimax以上である場合(ステップ85)、ループ85〜87における反復中の乗算結果を加算した合計は、密ベクトルy’に格納される(ステップ88)。反復処理ループは、次の列に移動し(ステップ89)、全ての列がループ(ステップ81〜89)を介して処理されるまで、ループを介した列の処理(ステップ81〜89)が継続し、その後にルーチン80は終了する。図8のルーチン80はまた、図2を参照して上記に示された擬似コードとの差異を示すテキストボックスにより、以下の擬似コードを使用して表すことができる。
for j=0からn―1 /*疎行列Aのn列にわたるループ*/
i←A’[j] /*i:j列目における第1の非ゼロエントリのA’vrにおけるインデックス*/
max←A’[j+1] /*imax:(j+1)列目における第1の非ゼロエントリのA’vrにおけるインデックス*/
sum←0 /*総和積算器を初期化*/
while (i<imax) /*列の終わりに到達したかどうかをテスト*/
Figure 2016119084
i←i+1 /*j列目における次の非ゼロ行まで移動*/
while文終了
y’[j]←sum /*y’に結果を格納*/
for文終了
構造化された符号化は、GPU及び排他的CPUベースのハードウェアの双方について有益であるが、処理速度のさらなる向上は、GPUベースのハードウェアについての圧縮符号化の順序維持パーティションを行うことによって得ることができる。ハードウェアシステムのハードウェアリソースが表1に関連して以下に説明する式を使用して決定することができることを考えると、パーティション化された圧縮表現を作成するかどうかは実用的である。図9は、1つの実施形態にかかる効率的な疎行列パーティション及び処理のためのコンピュータ実装方法90を示すフロー図である。他の実装が可能であるものの、方法90は、図3を参照して上述したシステムを使用して実装することができる。疎行列の圧縮表現が取得される(ステップ91)。そのような圧縮表現は、CSC符号化、CSR符号化、構造化CSC符号化又は構造化CSR符号化とすることができる。表現は、サーバ33によって記憶装置32からアクセスされるか又は他のソースから取得されることができる。サーバ33は、圧縮表現で表される行列の部分についてのパーティションを定義する(ステップ92)。具体的には、サーバ33は、作成されることになるパーティションの数(Ap及びA’)及び図3を参照して上述したAs配列における値を定義し、パーティションにおける行列の部分が有することができる非ゼロエントリの範囲を指定する。パーティションが定義されると、さらに図10を参照して後述するように圧縮表現の順序維持パーティションが実行される(ステップ93)。作成されたパーティションはマージされ、さらに図10を参照して記載されたように上述したマッピング配列(A又はA’)を形成し、マッピング配列に基づいてオフセット配列(Ao及びA’)を定義し、それゆえに、行列の順序維持パーティション圧縮表現を完成する(ステップ94)。そして、行列のパーティション化された圧縮表現は、図17〜図23を参照して後述するように、ページランク(登録商標)アルゴリズムなどの順次適用されるそのような処理の結果を用いたSpMV及びSpMTVなどの処理のために使用することができ(ステップ95)、方法90を終了する。
パーティション化は、それらの順序を維持しながら同様の数の非ゼロエントリを有する行列の部分をグループ化するのを可能とする。図10は、1つの実施形態にかかる図9の方法90において使用するための圧縮表現で表される行列の行の順序維持パーティションを実行するためのルーチン100を示すフロー図である。圧縮表現は、CSR及び構造化CSRの双方とすることができる。行列の全ての行であるi=0からm−1のidを有する行を処理する反復処理ループ(ステップ101〜107)が開始される(ステップ101)。A配列におけるi行目についてのエントリが配置され、初期の圧縮表現がCSR又は構造化CSRであったかどうかに応じて、行列におけるその行についての第1の非ゼロエントリのAv又はAvc配列におけるインデックスを識別し、識別されたインデックスが変数jとして設定される(ステップ102)。A配列における次の(i+1)エントリが配置され、変数jmaxとして設定される(ステップ103)。i行目が行列における最後の行でない限り、A配列における次のエントリは、(i+1)行目における第1の非ゼロエントリのAvc(又はA)配列におけるインデックスであり、i行目が行列における最後の行である場合、A配列における次のエントリは、Avc(又はAv)配列におけるエントリの総数である。i行目における非ゼロエントリの数は、jmaxからjを差し引くことによって決定され(ステップ104)、非ゼロエントリの数は、ei〜として示される。ei〜がゼロよりも大きい場合(ステップ105)、ei〜の値に基づいて及びパーティションにおいて許容される非ゼロエントリの最大及び最小数に基づいて定義されたパーティションのいずれかに行iが割り当てられ、パーティションkにおける行のリストの最後に行iの行idが追加される(ステップ106)。それゆえに、行iが属するインデックスkを有するパーティションは、As[k]≦ei〜<As[k+1]のようにみられる。1つの実施形態において、パーティションkは、bと設定されたei〜を有し、b以上である昇順配列Asの第1の要素のインデックスを返すlower_bound(As,b)として示される関数を使用して求めることができる。パーティションを求めるための他の方法も可能である。反復処理ループは、次の行に移動し(ステップ107)、全ての行が処理されるまでループ(ステップ101〜107)を介した行の処理が継続する。ei〜がゼロよりも大きくない場合(ステップ105)、ルーチン100は、ステップ107に直接移動する。行列の全ての行がループ(ステップ101〜107)を介して処理された後、ルーチン100は終了する。図9を参照して示されるルーチン100はまた、以下の擬似コードを使用して表現することができる。
for i=0からm―1 /*疎行列Aのm行にわたるループ*/
j←A[i] /*j:i行目における第1の非ゼロエントリのインデックス*/
max←A[i+1] /*jmax:(i+1)行目における第1の非ゼロエントリのインデックス*/
ei〜←jmax−j /*ei〜:i行目における非ゼロエントリ数*/
(ei〜>0)の場合
k←lower_bound(As,ei〜) /*パーティションk s.t.As[k]≦ei〜<As[k+1]をみつける*/
partition[k].push_back(i) /*パーティションkの終わりに行idiを追加(すなわち、Rk)*/
if文終了
for文終了
同様に、CSC又は構造化CSCについての順序維持行列パーティション化アルゴリズムを設計することができる。e〜jを疎行列の列jにおける非ゼロエントリ数とする。Ckと表されるk番目の列ベースのパーティションにおける列の設定は、Ck={j|As[k]≦e〜j<As[k+1]}と記載することができる。相対的順序を尊重することは、列インデックス配列A’にアクセスする際のランダム性を低減することから、パーティション化アルゴリズムは、元のCSC行列における同じパーティション列の相対的順序を尊重し、構造化及び非構造化CSC符号化の双方のキャッシュ性能を向上させる。図11は、1つの実施形態にかかる図9の方法90において使用するための圧縮表現で表される行列の列の順序維持パーティションを実行するためのルーチン110を示すフロー図である。圧縮表現は、CSC及び構造化CSCの双方とすることができる。
行列の全ての列であるj=0からn−1のidを有する列を処理する反復処理ループ(ステップ111〜117)が開始される(ステップ111)。A’配列におけるj列目についてのエントリが配置され、行列におけるその列についての第1の非ゼロエントリについて、元の圧縮符号化が構造化CSC又はCSCであるかどうかに応じて、A’vr又はA’配列におけるインデックスを識別し、識別されたインデックスは、変数iと設定される(ステップ112)。A’配列における次の(j+1)エントリが配置され、変数imaxと設定される(ステップ113)。j列目が行列における最後の列でない限り、A’c配列における次のエントリは、(j+1)列目における第1の非ゼロエントリのA’vr(又はA’)配列におけるインデックスであり、j列目が行列における最後の列である場合、A’配列における次のエントリは、A’vr(又はA’)配列におけるエントリの総数である。j列目における非ゼロエントリの数e〜jは、imaxからiを差し引くことによって決定される(ステップ114)。e〜jがゼロよりも大きい場合(ステップ115)、e〜jの値に基づいて及びパーティションにおいて許容される非ゼロエントリの最大及び最小数に基づいて定義されたパーティションのいずれかに列jが割り当てられ、パーティションkにおける列のリストの最後に列jの列idが追加される(ステップ116)。それゆえに、列jが属するインデックスkを有するパーティションは、A’s[k]≦e〜j<A’[k+1]のようにみられる。1つの実施形態において、パーティションkは、bと設定されたe〜jを有し、b以上である昇順配列A’の第1の要素のインデックスを返すlower_bound(A’)として示される関数を使用して求めることができる。パーティションを求めるための他の方法も可能である。反復処理ループは、次の列に移動し(ステップ117)、全ての列が処理されるまでループ(ステップ111〜117)を介した列の処理が継続する。e〜jがゼロよりも大きくない場合(ステップ115)、ルーチン110は、ステップ117に直接移動する。行列の全ての列が処理された後、ルーチン110は終了する。図11を参照して示されるルーチン110はまた、以下の擬似コードを使用して表現することができる。
for j=0からn―1 /*疎行列Aのn列にわたるループ*/
i←A’[j] /*i:j列目における第1の非ゼロエントリのインデックス*/
max←A’[j+1] /*imax:(j+1)列目における第1の非ゼロエントリのインデックス*/
e〜j←imax−i /*e〜j:j列目における非ゼロエントリ数*/
(e〜j>0)の場合
k←lower_bound(A’e〜j) /*パーティションk s.t.A’[k]≦e〜j<A’[k+1]をみつける*/
partition[k].push_back(j) /*パーティションkの終わりに列idjを追加(すなわち、Ck)*/
if文終了
for文終了
パーティションが作成されると、パーティションは、マッピング配列にマージされることができる。図12は、1つの実施形態にかかる図9の方法90において使用するためのマッピング配列にパーティションをマージするためのルーチン120を示すフロー図である。最初に、マッピング配列(A又はA’)は、元の圧縮表現に応じて、配列のサイズ、配列における行列の部分の数をゼロに設定することによって初期化される(ステップ121)。作成された全てのパーティションであるインデックスk=0からp−1を有するパーティションを処理する反復処理ループ(ステップ122〜126)が開始される(ステップ122)。パーティションkについて、kよりも小さいインデックスを有する全てのパーティションの積算サイズが計算され、kが0である場合、積算サイズもまた0である(ステップ123)。k未満のインデックスを有するパーティションは、配列におけるパーティションkに先行することから、前のパーティションの積算サイズは、マッピング配列の最後がどこにあるかを示し、積算サイズを決定すると、パーティションkは、マッピング配列の終わりに挿入される(ステップ124)。必要に応じて、パーティションkによって占有される記憶装置31内の任意のメモリは(125)まで解放される。ルーチン120は、次のパーティションに移動し(ステップ126)、全てのパーティションが処理されるまでループ(ステップ122〜126)を介したパーティションの処理が継続する。全ての作成されたパーティションを処理すると、全てのパーティションの積算サイズが計算されてAo[p]と設定され、ルーチン120を終了する。図12のルーチン120はまた、以下の擬似コードを使用して示すことができる。
.size←0 /*A:(順序維持置換)マッピング配列*/
for k=0からp―1 /*p個のパーティションにわたるループ*/
Ao[k]←A.size() /*Ao[k]:全ての前のパーティションの積算サイズ<k*/
minsert(partition[k]) /*Aの終わりにパーティションkを挿入*/
partition[k]を削除 /*パーティションkによって使用される空きメモリ(必要に応じて)*/
for文終了
Ao[p]←A.size() /*Ao[p]:全てのパーティションの総サイズ*/
図12を参照して示されたルーチン120に戻ると、Ao[k]は、パーティションkの第1の行(又は列)のランクを格納する。パーティションkの最後の行(又は列)のランクは、k=0,1,・・・,p−1について、Ao[k+1]−1によって与えられる。配列オフセットAo[p]の最後の要素は、CSR(又はCSC)についての行(又は列)の数である置換配列の要素数に常に等しい。
行列のパーティション化された圧縮表現は、CPUのみを含むものと同様にCPU及びGPUの双方を含むサーバ33による処理のために使用することができる。図13は、1つの実施形態にかかる図9の方法90において使用するためのパーティション化された圧縮表現で符号化された行列に対してSpMVを実行するためのルーチン130を示すフロー図である。ルーチンは、PSCSR及びPCSR表現の双方に適用することができる。ルーチン130は、CPUのみを含むものと同様にCPU及びGPUの双方を含むサーバによって使用することができる。作成された全てのパーティションであるインデックスk=0からAp−1を有するパーティションを処理する外側の反復処理ループ(ステップ131〜142)が開始される(ステップ131)。k番目のパーティションにおける全ての行を処理する内側の反復処理ループ(ステップ132〜141)が開始され、マッピング配列における行のランクは、r=Ao[k]からAo[k+1]−1である(ステップ132)。r番目にランク付けされた行のidは、iと識別されて設定される(ステップ133)。A配列におけるi行目についてのエントリが配置され、初期の圧縮表現がCSR又は構造化CSRであるかどうかに応じて、行列におけるその行についての第1の非ゼロエントリのAv又はAvc配列におけるインデックスを識別し、識別されたインデックスが変数jとして設定される(ステップ134)。A配列における次の(i+1)エントリが配置され、変数jmaxと設定される(ステップ135)。i行目が行列における最後の行でない限り、A配列における次のエントリは、(i+1)行目における第1の非ゼロエントリのAvc(又はAv)配列におけるインデックスであり、i行目が行列における最後の行である場合、A配列における次のエントリは、Avc(又はAv)配列におけるエントリの総数である。ステップ138において後述する非ゼロ配列の値の乗算結果を加算する関数である総和計算部は、ゼロにおける合計の値を設定することによって初期化される(ステップ136)。jがjmax未満である場合(ステップ137)、SpMVが実行される符号化がPSCSR又はPCSRであるかどうかに応じた動作で、jの値に対して乗算及び加算計算が実行される(ステップ138)。符号化がPSCSRである場合、計算は、以下の式にしたがって行われる。
sum←sum+Avc[j].v×x[Avc[j].c]
計算において、インデックスjを有するAvc配列の要素に格納された値はxの要素によって乗算され、そのインデックスは、j番目のインデックスを有するAvc要素に格納された列のインデックスであり、乗算結果は、ステップ138の前の反復中に実行された乗算結果の合計に加算される。あるいは、符号化がPCSRである場合、計算は、以下の式にしたがって行われる。
sum←sum+Av[j]×x[A[j]]
ここで、インデックスjを有するAv配列における値は、xの要素によって乗算され、そのインデックスは、j番目のインデックスを有するA配列における数であり、乗算結果は、ステップ138の前の反復中に実行された乗算結果の合計に加算される。いずれかの式に基づいて計算が終了すると、jの値に1が加算され、加算結果がjと設定され(ステップ139)、その行における次の列のエントリに処理を移動する。ルーチン130は、上述したステップ137に戻り、i行目における非ゼロ値が処理されるまでステップ137〜139を繰り返す。jがjmax以上である場合(ステップ137)、ループ137〜139における反復中に乗算結果を加算した合計は、密ベクトルyに格納される(ステップ140)。反復処理ループは、次の行に移動し(ステップ141)、パーティションの全ての行がステップ132〜141を介して処理されるまで内側ループ(ステップ132〜141)を介した行の処理が継続する。k番目のパーティションの全ての行が処理されると、ルーチン130は、次のパーティションに移動し(ステップ142)、全てのパーティションが処理されるまで外側処理ループ(ステップ131〜142)を介したパーティションの処理が継続する。全てのパーティションがステップ131〜142において処理されると、ルーチン130は終了する。ルーチン130はまた、PSCSR符号化に対してSpMVを実行することを含む以下の擬似コードを使用して表現することができる。
for k=0からAp―1 /*Ap個の行ベースのパーティションにわたるループ*/
for r=Ao[k]からAo[k+1]−1 /*k番目のパーティションにおける行にわたるループ*/
i←A[r] /*i:r番目にランク付けされた行のid*/
j←A[i] /*j:i行目における第1の非ゼロエントリのAvcにおけるインデックス*/
max←A[i+1] /*jmax:(i+1)行目における第1の非ゼロエントリのAvcにおけるインデックス*/
sum←0 /*総和積算器の初期化*/
while(j<jmax) /*行の終わりに到達したかどうかのテスト*/
sum←sum+Avc[j].v×x[Avc[j].c] /*yi=ΣjAi,j×xjの計算
j←j+1 /*i行目における次の非ゼロ列に移動*/
while文終了
y[i]←sum /*yに結果を格納*/
for文終了
for文終了
ルーチン130はまた、準用するPCSR符号化に対してSpMVを実行するための擬似コードを使用して表現することができる。
SpMTVは、CPUのみを含むものと同様にCPU及びGPUの双方を含むサーバを使用してパーティション化された圧縮表現に対して行うことができる。図14は、図9の方法90において使用するための行列のパーティション化された圧縮表現に対してSpMTVを実行するためのルーチン150を示すフロー図である。ルーチンは、PSCSC及びPCSC表現の双方に適用することができる。ルーチン150は、CPUのみを含むものと同様にCPU及びGPUの双方を含むサーバを使用することができる。作成された全てのパーティションであるインデックスk=0からA’−1のパーティションを処理する外側の反復処理ループ(ステップ151〜162)が開始される(ステップ151)。k番目のパーティションにおける全ての列を処理する内側の反復処理ループ(ステップ152〜161)が開始され、処理されるマッピング配列における列のランクは、r=A’[k]からA’[k+1]−1である(ステップ152)。r番目にランク付けされた列のidは、jと識別されて設定される(ステップ153)。A’配列におけるj列目についてのエントリが配置され、圧縮表現がCSC又は構造化CSCであるかどうかに応じて、行列におけるその列についての第1の非ゼロエントリのA’又はA’vr配列におけるインデックスを識別し、識別されたインデックスが変数iとして設定される(ステップ154)。A’配列における次の(j+1)エントリが配置され、変数imaxと設定される(ステップ155)。j列目が行列における最後の列でない限り、A’配列における次のエントリは、(j+1)列目における第1の非ゼロエントリのA’vr(又はA’)配列におけるインデックスであり、j列目が行列における最後の列である場合、A’配列における次のエントリは、A’vr(又はA’)配列におけるエントリの総数である。ステップ158において後述する非ゼロ配列の値の乗算結果を加算する関数である総和計算部は、ゼロにおける合計の値を設定することによって初期化される(ステップ156)。iがimax未満である場合(ステップ157)、SpMTVが実行される符号化がPSCSC又はPCSCであるかどうかに応じた動作で、iの値に対して乗算及び加算計算が実行される(ステップ158)。符号化がPSCSCである場合、計算は、以下の式にしたがって行われる。
sum←sum+A’vr[i].v×x’[A’vr[i].r]
ここで、インデックスiを有するA’vr配列の要素に格納された値は、x’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’vr要素に格納された行のインデックスであり、乗算結果は、ステップ158の前の反復中に実行された乗算結果の合計に加算される。あるいは、符号化がPCSRである場合、計算は、以下の式にしたがって行われる。
sum←sum+A’[i]×x’[A’[i]]
ここで、インデックスiを有するA’配列における値は、x’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’配列における数であり、乗算結果は、ステップ158の前の反復中に実行された乗算結果の合計に加算される。いずれかの式に基づいて計算が終了すると、iの値に1が加算され、加算結果がiと設定され(ステップ159)、その列における次の行のエントリに処理を移動する。ルーチン150は、上述したステップ157に戻り、j列目における非ゼロ値が処理されるまでステップ158〜159を繰り返す。iがimax以上である場合(ステップ157)、ループ157〜159における反復中に乗算結果を加算した合計は、密ベクトルy’に格納される(ステップ160)。反復処理ループは、次の列に移動し(ステップ161)、全ての列がステップ152〜161を介して処理されるまで内側処理ループ(ステップ152〜161)を介した列の処理が継続する。k番目のパーティションの全ての列が処理されると、ルーチン150は、次のパーティションに移動し(ステップ162)、全てのパーティションが処理されるまで外側処理ループ(ステップ151〜162)を介したパーティションの処理が継続する。全てのパーティションがステップ151〜162において処理されると、ルーチン150は終了する。ルーチン150はまた、PSCSR符号化に対してSpMTVを実行することを含む以下の擬似コードを使用して表現することができる。
for k=0からA’―1 /*A’個の列ベースのパーティションにわたるループ*/
for r=A’[k]からA’[k+1]−1 /*k番目のパーティションにおける列にわたるループ*/
j←A’[r] /*j:r番目にランク付けされた列のid*/
i←A’[j] /*i:j列目における第1の非ゼロエントリのA’vrにおけるインデックス*/
max←A’[j+1] /*imax:(j+1)列目における第1の非ゼロエントリのA’vrにおけるインデックス*/
sum←0 /*総和積算器の初期化*/
while(i<imax) /*列の終わりに到達したかどうかのテスト*/
sum←sum+A’vr[i].v×x’[A’vr[i].r] /*y’j=ΣiAi,j×x’iの計算
i←i+1 /*j列目における次の非ゼロ行に移動*/
while文終了
y’[j]←sum /*y’に結果を格納*/
for文終了
for文終了
上記図13及び図14を参照して上述したルーチンは、それらを実行するための特定のハードウェアセットアップに固有のものではなかった。GPUを含むハードウェアは、特定のパーティションに対してSpMV及びSpMTVを実行するために特定のカーネルを割り当てることによってより高速になる可能性がある。図15は、1つの実施形態にかかる図9の方法90において使用するためのパーティション化された圧縮表現処理のためのカーネル関数を起動するためのルーチンである。カーネル関数Fの配列は、図3を参照して上述したカーネルのリストを含むF=[f1T1R,f1W1R,f1B1R]のように作成され、カーネルは、0から2にインデックス付けされ、f1T1Rはインデックス0を有し、f1B1Rはインデックス2を有する(ステップ171)。さらなる実施形態において、他のカーネルは、配列の一部とすることができる。作成された全てのパーティションであるインデックスk=0からAp−1を有するパーティションについての反復処理ループ(ステップ172〜176)が開始される(ステップ172)。カーネル関数の1つは、kの値及びF配列におけるカーネルのインデックスに基づいてパーティションkについて選択される(ステップ173)。例えば、k=0の場合、0番目のインデックスを有するカーネルf1T1Rが選択され、k=1の場合、1のインデックスを有するカーネルf1W1Rが選択され、k=2の場合、2のインデックスを有するカーネルf1B1Rが選択される。3超のパーティションがある場合、2よりも大きいインデックスkを有する全てのパーティションについてf1B1Rが選択される。パーティション及びカーネルについての他のインデックスもまた使用可能であり、カーネル及びパーティションを一致させる他の方法も可能である。起動引数は、さらに図16を参照して記載されるように、選択されたカーネルについて設定される(ステップ174)。選択された関数は、各引数について起動され、さらに図17〜図22を参照して後述するように、パーティション化された圧縮符号化の処理に使用される(ステップ175)。反復処理ループは、次のパーティションに移動し(ステップ176)、全てのパーティションが処理されるまでループ(ステップ172〜176)を介したパーティションの処理が継続する。全てのパーティションが処理されると、ルーチン170は終了する。ルーチン170はまた、以下の擬似コードを使用して表現することができる−擬似コードは、行ベースのパーティションについて記載されているが、列ベースのパーティションについての擬似コードを準用して記載することができる。
for k=0からAp−1 /*Ap個の行ベースのパーティションにわたるループ*/
args←kernel_launch_args(F[k],Ao[k+1]−Ao[k]) /*F[k]:k番目のカーネル関数*/
F[k]<<<args>>>(y,x,Avc,A,A,Ao[k],Ao[k+1]) /*k番目のパーティションについてのk番目のカーネルを起動*/
for文終了
擬似コードは、PSCSRを参照して記載されているが、他のパーティション化された圧縮表現についての擬似コードを準用して表すことができる。図15を参照して上述したもの以外の特定のパーティションを処理するための特定のカーネル関数を選択する方法もまた可能である。
カーネルが選択された後に設定された起動引数は、パーティション化された圧縮表現の処理を実行するサーバ33などのシステムの実行時に入力される制約を提供する。GPU上でカーネルを起動するために、スレッドブロックの寸法及び可能であればそのようなスレッドブロックのグリッドの寸法などの起動引数を指定する必要がある。同じベンダからの異なるGPUは、ブロック又は各寸法に沿ったグリッドの最大サイズについての制約を課すことがあり、異なるCPUベンダは、異なる数のスレッドブロック及び/又はグリッド寸法をサポートすることがある。図16は、1つの実施形態にかかる図15のルーチン170において使用するための選択されたカーネル関数の起動引数を設定するためのルーチン180である。nを選択された行列に割り当てられた行列の部分、行又は列の数とする。選択されたカーネル関数がf1T1Rである場合(ステップ181)、nが起動するスレッドの最小数であるという制約が設定される。残りの引数が設定されて制約に基づいて返され(ステップ183)、ルーチン180を終了する。選択された関数がf1T1Rでない場合(ステップ181)、選択された関数がf1W1Rであるかどうかが判定される(ステップ184)。選択された関数がf1W1Rである場合(ステップ184)、ブロックサイズがワープサイズに等しく設定され且つnが起動するブロックの最小数であるという制約が設定される(185)。ルーチン180は、後述するステップ183に移動する。選択された関数がf1W1Rでない場合(ステップ181)、選択された関数がf1B1Rであるかどうかが判定される(ステップ186)。選択された関数がf1B1Rである場合(ステップ187)、ブロックサイズは、ブロック内のスレッド数に設定され、nが起動するブロックの最小数に設定される。ルーチン180は、後述するステップ183に移動する。選択された関数がf1B1Rでない場合(ステップ186)、カーネル関数は未知であり、ルーチン180は終了する。ルーチン180はまた、以下の擬似コードを使用して表現することができる−擬似コードは、SpMVを実行するためのカーネルの起動について記載されているが、SpMTVを実行するためのカーネルの起動についての擬似コードを準用して記載することができる。
関数 kernel_launch_args(f,n) /*f:SpMVカーネル;n:fに割り当てられた行数*/
(f=f1T1R)である場合 /*f1T1R:1スレッド1行カーネル*/
args.set_min_threads(n) /*n=起動されるスレッドの最小数*/
(f=f1W1R)である場合 /*f1W1R:1ワープ1行カーネル*/
args.set_block_size(WARPSIZE) /*BLOCKSIZE=WARPSIZEに設定*/
args.set_min_blocks(n) /*n=起動されるブロックの最小数*/
(f=f1B1R)である場合 /*f1B1R:1ブロック1行カーネル*/
args.set_block_size(BLOCKSIZE) /*BLOCKSIZE:ブロックにおけるスレッド数*/
args.set_min_blocks(n) /*n=起動されるブロックの最小数*/
それ以外の場合
エラー「中断:未知のカーネル関数」
if文終了
args.compute_satisfy_args() /*上記制約に基づいて残りの引数を設定*/
リターン args
固定されたスレッド−ブロック及びグリッド寸法を使用する代わりに、上記擬似コードは、3つのSpMVカーネルのそれぞれによって課される明示的な制約を形成することにより、制約を満たす問題を設定する起動引数に接近する。例えば、f1T1Rカーネルは、スレッドの総数がカーネルに割り当てられた行数に少なくとも等しくなければならないという単一の制約の充足を必要とするのみである。一方、f1W1Rカーネルは、(1)ブロック内のスレッド数がWARPSIZEと同じでなければならない、及び(2)カーネルに割り当てられた行があるのと少なくとも同数のスレッドブロックがなければならないという2つの制約を同時に充足することを要求する。ハードウェアによって課される制約を尊重しつつ、それらのカーネル起動制約を満たす方法は、GPU又はベンダ依存とすることができ、具体的な説明は提供されない。通常、最良の起動引数は、2の整数累乗である値を想定し、それゆえに、そのような制約充足問題の探索空間は、通常、非常に小さく且つ扱いやすい。例えば、Nvidia(登録商標)のFermi GPU上の実装は、f1T1Rカーネルについては8×8のスレッドブロックサイズを使用し、f1W1Rカーネルについては32×1のブロックサイズを使用し、f1B1Rカーネルについては512×1のブロックサイズを使用する。
選択されると、カーネル関数は、行列の割り当てられた部分を処理することができる。行列パーティション化の目的は、同じパーティションにおける行がカーネルが最適化された非ゼロエントリ数などのいくつかの共通の特徴を共有するように、単一のGPUカーネル関数について1つのパーティションを作成することであることを思い出されたい。上述したGPUベースのSpMVアルゴリズムは、任意数のパーティション及びカーネルを扱うことができる一方で、実験は、GTX480及び580を含むNvidia(登録商標)のFermiクラスのGPUにおいて最良のSpMV結果を生成する3つの異なるカーネル(f1T1R,f1W1R,f1B1R)を有することを示している。
WARPSIZE(本特許出願において「ワープサイズ」とも称される)及びBLOCKSIZE(本特許出願において「ブロックサイズ」とも称される)を、それぞれ、ワープ及びスレッドブロックにおけるスレッド数とする。GPUコアに対する行列の行(又は列)の分布を容易とするために、図17〜図22Bを参照して以下に記載されるルーチンにおいて以下のヘルパー関数を使用可能である。
・thread_id()は、現在のスレッドについてのグローバルに固有のスレッドidを返す
・warp_id()は、現在のスレッドについてのグローバルに固有のワープidを返す
・warp_thread_id()は、現在のスレッドについての(ワープ内でのみ固有の)ローカルスレッドidを返す
・block_thread_id()は、現在のスレッドについての(ブロック内でのみ固有の)ローカルスレッドidを返す
・sync_warp_threads()は、ワープ内の全てのスレッド間で同期する
・sync_block_threads()は、ブロック内の全てのスレッド間で同期する
全てのグローバル(又はローカル)スレッド(又はワープ)idは、ゼロで始まるスレッドのインデックスを有するゼロベースである。WARPSIZEは、ハードウェアに依存する:Nvidia(登録商標)については32個又はアドバンスト・マイクロ・デバイス(登録商標)については64個のGPU。一方、BLOCKSIZEは、通常はWARPSIZEの整数倍であるプログラム可能なパラメータとすることができる。共有されるキーワードは、同じブロック(又はBLOCKSIZE=WARPSIZEである場合にはワープ)内の全てのスレッド間で共有されるスレッド変数を宣言する。
図17は、1つの実施形態にかかる図15のルーチンにおいて使用するためのf1T1Rカーネル関数によってSpMVを実行するためのルーチン190を示すフロー図である。ルーチン190は、PSCSR又はPCSR符号化のいずれかでSpMVを実行するために使用することができる。カーネルにおける全ての起動されるスレッドについて反復処理ループが開始され、スレッドは並列に動作し、スレッドは並列に動作し、それゆえに、以下のステップは、1つのスレッドを参照して説明され、任意の他の起動スレッドは、並列にステップを行う(ステップ191)。thread_id()関数を使用することによってスレッドの固有idが取得され、Ao[k]に格納されたパーティションk(スレッドが割り当てられたパーティション)の第1の行のランクに等しい変数rminの値に追加され、加算結果は、そのスレッドに割り当てられた行のランクを識別する変数rによって表される(192)。rの値は、Ao[k+1]によって与えられる次のパーティションk+1の第1の行のランクに等しい変数rmaxと比較することができ、起動されるスレッド数がパーティションkにおける行数(rmax−rmin)に等しい場合には比較は任意である(193)。rがrmax未満である場合(ステップ193)、r番目のランク付けされた行のidiは、マッピング配列Aにおいて識別される(ステップ194)。rがrmax以上である場合、ルーチン190は、後述するステップ202に移動する。Ar配列におけるi行目についてのエントリが配置され、初期の圧縮表現がCSR又は構造化CSRであったかどうかに応じて、行列におけるその行についての第1の非ゼロエントリのAv又はAvc配列におけるインデックスを識別し、識別されたインデックスは、変数jとして設定される(ステップ195)。A配列における次の(i+1)エントリが配置され、変数jmaxとして設定される(ステップ196)。i行目が行列における最後の行でない限り、A配列における次のエントリは、(i+1)行目における第1の非ゼロエントリのAvc(又はAv)配列におけるインデックスである。i行目が行列における最後の行である場合、A配列における次のエントリは、Avc(又はAv)配列におけるエントリの総数である。以下のステップ199において記載される非ゼロ配列の値の乗算結果をともに加算する関数である総和計算部は、ゼロにおいて合計値を設定することによって初期化される(ステップ197)。jがjmax未満である場合(ステップ198)、SpMVが行われた符号化がPSCSR又はPCSRであるかどうかに応じて、乗算及び加算計算がjの値に対して実行される(ステップ199)。符号化がPSCSRである場合、計算は、以下の式にしたがって行われる。
sum←sum+Avc[j].v×x[Avc[j].c]
計算において、インデックスjを有するAvc配列の要素に格納された値はxの要素によって乗算され、そのインデックスは、j番目のインデックスを有するAvc要素に格納された列のインデックスであり、乗算結果は、ステップ199の前の反復中に実行された乗算結果の合計に加算される。あるいは、符号化がPCSRである場合、計算は、以下の式にしたがって行われる。
sum←sum+Av[j]×x[A[j]]
ここで、インデックスjを有するAv配列における値は、xの要素によって乗算され、そのインデックスは、j番目のインデックスを有するA配列における数であり、乗算結果は、ステップ199の前の反復中に実行された乗算結果の合計に加算される。いずれかの式に基づいて計算が終了すると、jの値に1が加算され、加算結果がjと設定され(ステップ200)、その行における次の列のエントリに処理を移動する。ルーチン190は、上述したステップ198に戻り、i行目における非ゼロ値が処理されるまでステップ199〜201を繰り返す。jがjmax以上である場合(ステップ198)、ループ198〜200における反復中に乗算結果を加算した合計は、密ベクトルyに格納される(ステップ201)。スレッドの処理の実行は停止され(ステップ202)、ルーチン190は、全ての起動したスレッドの処理の停止によって終了する。ルーチン190はまた、以下の擬似コードを使用して表現することができ−擬似コードは、PSCSR符号化に関して記載されるが、PCSR符号化についての擬似コードを準用して記載することができる。
カーネル f1T1R(y,x,Avc,A,A,rmin,rmax) /*1T1R:1スレッド1行SpMVカーネル*/
r←rmin+thread_id() /*このスレッドに割り当てられた頂点のランクを計算*/
(r<rmax)である場合 /*任意の正確に(rmax−rmin)スレッドが作成された場合*/
i←A[r] /*i:r番目にランク付けされた行のid*/
j←A[i] /*j:i行目における第1の非ゼロエントリのAvcにおけるインデックス*/
max←A[i+1] /*jmax:(i+1)行目における第1の非ゼロエントリのAvcにおけるインデックス*/
sum←0 /*総和積算器の初期化*/
while(j<jmax) /*行の終わりに到達したかどうかのテスト*/
sum←sum+Avc[j].v×x[Avc[j].c] /*yi=ΣjAi,j×xjの計算
j←j+1 /*i行目における次の非ゼロ列に移動*/
while文終了
y[i]←sum /*yに結果を格納*/
if文終了
図18は、1つの実施形態にかかる図15のルーチン170において使用するためのf1T1Rカーネル関数によってSpMTVを実行するためのルーチン210を示すフロー図である。ルーチン210は、図17を参照して上述したように、PSSCR又はPCSR符号化のいずれかにおいてSpMTVを実行するために使用することができ、全ての起動したスレッドは並列に動作する。反復処理ループは、カーネル内の全ての起動したスレッドについて開始され、スレッドは並列に動作する(ステップ211)。スレッドの固有idがthread_id()関数を使用して取得され、A’[k]に格納されたパーティションk(スレッドが割り当てられたパーティション)の第1の列のランクに等しい変数rminの値に加算され、加算結果は、変数rによって表される。rの値は、A’[k+1]によって与えられる次のパーティションk+1の第1の列のランクに等しい変数rmaxと比較することができ、起動されるスレッド数がパーティションkにおける列数(rmax−rmin)に等しい場合には比較は任意である。rがrmax未満である場合(ステップ213)、r番目のランク付けされた列のidjは、マッピング配列A’において識別される(ステップ214)。rがrmax以上である場合、ルーチン210は、後述するステップ222に移動する。A’配列におけるj列目についてのエントリが配置され、圧縮表現がCSC又は構造化CSCであったかどうかに応じて、行列におけるその列についての第1の非ゼロエントリのA’又はA’vr配列におけるインデックスを識別し、識別されたインデックスは、変数iとして設定される(ステップ215)。A’配列における次の(j+1)エントリが配置され、変数imaxとして設定される(ステップ216)。j列目が行列における最後の列でない限り、A’配列における次のエントリは、(j+1)列目における第1の非ゼロエントリのA’vr(又はA’)配列におけるインデックスである。j列目が行列における最後の列である場合、A’配列における次のエントリは、A’vr(又はA’)配列におけるエントリの総数である。以下のステップ219において記載される非ゼロ配列の値の乗算結果をともに加算する関数である総和計算部は、ゼロにおいて合計値を設定することによって初期化される(ステップ217)。iがimax未満である場合(ステップ218)、SpMTVが行われた符号化がPSCSC又はPCSCであるかどうかに応じて、乗算及び加算計算がiの値に対して実行される(ステップ219)。符号化がPSCSCである場合、計算は、以下の式にしたがって行われる。
sum←sum+A’vr[i].v×x’[A’vr[i].r]
ここで、インデックスiを有するA’vr配列の要素に格納された値は、x’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’vr要素に格納された行のインデックスであり、乗算結果は、ステップ219の前の反復中に実行された乗算結果の合計に加算される。あるいは、符号化がPCSRである場合、計算は、以下の式にしたがって行われる。
sum←sum+A’[i]×x’[A’[i]]
ここで、インデックスiを有するA’配列における値は、x’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’配列における数であり、乗算結果は、ステップ219の前の反復中に実行された乗算結果の合計に加算される。いずれかの式に基づいて計算が終了すると、iの値に1が加算され、加算結果がiと設定され(ステップ220)、その列における次の行のエントリに処理を移動する。ルーチン210は、上述したステップ218に戻り、j列目における非ゼロ値が処理されるまでステップ218〜220を繰り返す。iがimax以上である場合(ステップ218)、ループ218〜220における反復中に乗算結果を加算した合計は、密ベクトルy’に格納される(ステップ221)。スレッドの実行が停止される(ステップ222)。ルーチン210は、全ての起動した処理スレッドの実行が停止されると終了する。ルーチン210についての擬似コードは、準用する図17を参照して上記に示された擬似コードと同様に記載することができる。
非ゼロエントリ数が多い行列の部分は、各行などの各部分を処理するためにスレッドのワープを割り当てるf1W1Rカーネルによって処理されることから利益を得ることができる。図19A〜図19Bは、1つの実施形態にかかる図15のルーチン170において使用するためのf1W1Rカーネル関数によってSpMVを実行するためのルーチン230を示すフロー図である。ルーチン230は、PSSCR又はPCSR符号化のいずれかにおいてSpMVを行うために使用することができる。カーネルにおける全ての起動ワープについて反復処理ループが開始される(ステップ231)。起動ワープ及び起動ワープにおけるスレッドは、互いに並列に実行される。それゆえに、以下の説明がワープのいずれか又はスレッドのいずれかを参照する場合には、全ての他の起動ワープ又はスレッドは、スレッド又はワープのいずれかの実行が停止されるまで(以下のステップ251など)参照されるワープ又はスレッドと並列にルーチン230の同じステップを実行する。
ローカル共有は、同じワープにおける全てのスレッド間で開始され、スレッドがステップ240〜242において後述する計算結果を共有するのを可能とする(ステップ232)。ワープのいずれかの固有のグローバルidを取得するために関数warp__id()を使用し且つrminの値に固有のグローバルidを追加することによってワープのいずれかに割り当てられた行のランク(図17を参照して上述したようにランクが割り当てられるパーティションの第1の行のランク)が取得され、加算結果は、ワープに割り当てられた行のランクであり、変数rを使用して表される(ステップ233)。値rは、Ao[k+1]によって与えられる次のパーティションk+1の第1の行のランクに等しい変数であるrmaxと比較することができ、起動されるワープ数がパーティションkにおける行数(rmax−rmin)に等しい場合には比較は任意である。rがrmax未満である場合(ステップ234)、r番目のランク付けされた行のidiは、マッピング配列Aにおいて識別される(ステップ235)。rがrmax以上である場合(ステップ234)、ルーチン230は、後述するステップ251に移動する。関数warp_thread_id()を使用することなどによってワープにおけるスレッドのいずれかのローカルid(ワープ内のid)が取得され、変数tを使用して表される(ステップ236)。そのスレッドに割り当てられた第1の非ゼロエントリのインデックスは、A配列におけるi行目についてのエントリの値をtに追加することによって取得され、加算結果は、変数jを使用して表される(ステップ237)。A配列における次の(i+1)エントリが配置され、変数jmaxとして設定される(ステップ238)。以下のステップ241において記載される非ゼロ配列の値の乗算結果をともに加算する関数であるローカル総和計算部は、ゼロにおいて合計値を設定することによって初期化される(ステップ239)。jがjmax未満である場合(ステップ240)、SpMVが行われた符号化がPSCSR又はPCSRであるかに応じた動作で、値jに対して乗算及び加算計算が実行される(ステップ241)。符号化がPSCSRである場合、計算は、以下の式にしたがって行われる。
sum←sum+Avc[j].v×x[Avc[j].c]
計算において、インデックスjを有するAvc配列の要素に格納された値はxの要素によって乗算され、そのインデックスは、j番目のインデックスを有するAvc要素に格納された列のインデックスであり、乗算結果は、ステップ241の前の反復中に実行された乗算結果の合計に加算される。あるいは、符号化がPCSRである場合、計算は、以下の式にしたがって行われる。
sum←sum+Av[j]×x[A[j]]
ここで、インデックスjを有するAv配列における値は、xの要素によって乗算され、そのインデックスは、j番目のインデックスを有するA配列における数であり、乗算結果は、ステップ241の前の反復中に実行された乗算結果の合計に加算される。いずれかの式に基づいて計算が終了すると、jの値にワープサイズが加算され、加算結果がjと設定され(ステップ242)、スレッドが割り当てられたその行における次の列のエントリに処理を移動する。ルーチン230は、上述したステップ240に戻り、そのスレッドが割り当てられるi行目における非ゼロ値が処理されるまでステップ240〜242を繰り返す。jがjmax以上である場合(ステップ240)、ワープにおける全てのスレッドは、必要に応じて同期され、ワープにおける全てのスレッドがステップ240〜242のループの実行を終了するのを可能とし、同期は、並列に動作するワープにおけるスレッドに起因して1つの実施形態では必要なく、同期から出ることがない一方で、さらなる実施形態においては同期が行われる。
ワープのサイズ、ワープにおけるスレッド数は、整数除算を使用して2で割られ、除算結果は、変数tmaxを使用して表される(ステップ244)。tの場合、ワープにおけるスレッドのidは、tmax未満であり(ステップ245)、スレッドは、ともに組み合わせて、そのスレッドt及びスレッドidがt+tmaxである他のスレッドによって上記実行されるステップ240〜242における計算から得られる合計の削減を行う(ステップ246)。ステップ246は、半分の合計数を削減し、組み合わせた合計は、スレッドt(組み合わせを行ったスレッド)についての合計として設定され、ワープにおけるスレッドidがt+tmaxであるスレッドと前に関連付けられた合計は破棄される(ステップ246)。例えば、ワープにおいて32のスレッドがあり且つスレッドについてのtが0である場合、スレッドtは、ステップ240〜242のスレッドt自身の性能から得られた合計と、スレッドidが16であるスレッドによってステップ240〜242の性能から得られた合計とを組み合わせ、組み合わせた合計は、ステップ246のその後の反復についてのスレッドtの合計として設定され、16のidを有するスレッドは、もはや合計に関連付けられない。組み合わせに続いて、値tmaxは、整数除算を使用して半分にカットされ、tmaxとして設定される(ステップ247)。必要に応じて、合計の組み合わせを行ったスレッドは再度同期され(ステップ248)、ルーチン230は、上記ステップ245に戻る。ステップ245〜247のループの各後続反復中において、ループに参加しているスレッド数は、値tmaxの削減に起因して半分に削減される。
tがtmax以上である場合(ステップ245)、スレッドがワープにおける第1のスレッドであるかどうか(t=0及びr<rmax)が判定される(ステップ249)。スレッドが第1のスレッドでない場合、スレッドの実行は終了する(ステップ251)。スレッドが第1のスレッドである場合、ステップ245〜247において記載された削減から生じる総合計は、密ベクトルyに格納され(ステップ250)、そのスレッドの実行が終了する(ステップ251)。ルーチン230は、全ての起動ワープについての全てのスレッドの実行の終了によって終了する。ルーチン230はまた、以下の擬似コードを使用して表現することができる−擬似コードは、PSCSR符号化に関連して記載されるが、PCSRに関連する擬似コードを準用して記載することができる。
カーネル f1W1R(y,x,Avc,A,A,rmin,rmax) /*1W1R:1ワープ1行SpMVカーネル*/
共有 sum[WARPSIZE] /*sum:ワープにおけるスレッドによって共有されるローカル合計*/
r←rmin+warp_id() /*このワープに割り当てられた頂点のランクの計算*/
(r<rmax)である場合 /*任意の正確に(rmax−rmin)ワープが作成された場合*/
i←A[r] /*i:r番目にランク付けされた行のid*/
t←warp_thread_id() /*t:ワープにおけるローカルスレッドid*/
j←A[i]+t /*j:このスレッドに割り当てられた第1の非ゼロエントリのインデックス*/
max←A[i+1] /*jmax:(i+1)行目における第1の非ゼロエントリのインデックス*/
sum[t]←0 /*ローカル総和積算器の初期化*/
while(j<jmax) /*行の終わりに到達したかどうかのテスト*/
sum[t]←sum[t]+Avc[j].v×x[Avc[j].c] /*yi=ΣjAi,j×xjの計算
j←j+WARPSIZE /*このスレッドについての次の非ゼロエントリに移動*/
while文終了
if文終了
sync_warp_threads() /*任意にワープにおけるスレッドが常に同期している場合*/
max←WARPSIZE/2 /*tmax:ローカル合計まで追加したスレッド数*/
while(t<tmax) /*このスレッドが参加すべきであるかどうかのテスト*/
sum[t]←sum[t]+sum[t+tmax] /*2つのローカル合計を1つに削減*/
max←tmax/2 /*合計追加スレッド数を半分にカット*/
sync_warp_threads() /*任意にワープにおけるスレッドが常に同期している場合*/
while文終了
(t=0及びr<rmax)である場合 /*これはワープにおける第1のスレッド?*/
y[i]←sum[0] /*yに総合計を格納*/
if文終了
SpMTVは、ブロックを処理するためにスレッドのワープを割り当てるf1W1Rカーネルを使用して同様に行うことができる。図20A〜図20Bは、1つの実施形態にかかる図15のルーチン170において使用するためのf1W1RカーネルによってSpMTVを実行するためのルーチン260を示すフロー図である。ルーチン260は、PSCCR又はPCSR符号化のいずれかにおいてSpMVを行うために使用することができる。カーネルにおける全ての起動ワープについて反復処理ループが開始される(ステップ261)。起動ワープ及び起動ワープにおけるスレッドは、互いに並列に実行され、それゆえに、以下の説明がワープのいずれか又はスレッドのいずれかを参照する場合には、全ての他の起動ワープ又はスレッドは、スレッド又はワープのいずれかの実行が停止されるまで(以下のステップ281など)同時に並列にルーチン260の同じステップを実行する。
ローカル共有は、同じワープにおける全てのスレッド間で開始され、スレッドがステップ270〜272において後述する計算結果を共有するのを可能とする(ステップ262)。ワープのいずれかの固有のグローバルidを取得するために関数warp__id()を使用し且つrminの値に固有のグローバルidを追加することによってワープのいずれかに割り当てられた列のランク(図18を参照して上述したようにランクが割り当てられるパーティションの第1の列のランク)が取得され、加算結果は、ワープに割り当てられた列のランクであり、変数rを使用して表される(ステップ263)。値rは、A’[k+1]によって与えられる次のパーティションk+1の第1の列のランクに等しい変数であるrmaxと比較することができ、起動されるワープ数がパーティションkにおける列数(rmax−rmin)に等しい場合には比較は任意である。rがrmax未満である場合(ステップ264)、r番目のランク付けされた列のidjは、マッピング配列A’において識別される(ステップ265)。rがrmax以上である場合(ステップ264)、ルーチン260は、後述するステップ281に移動する。関数warp_thread_id()を使用することなどによってワープにおけるスレッドのいずれかのローカルid(ワープ内のid)が取得され、変数tを使用して表される(ステップ266)。そのスレッドに割り当てられた第1の非ゼロエントリのインデックスは、A’配列におけるj列目についてのエントリの値をtに追加することによって取得され、加算結果は、変数iを使用して表される(ステップ267)。A’配列における次の(j+1)エントリが配置され、変数imaxとして設定される(ステップ268)。j列目が行列における最後の列でない限り、A’配列における次のエントリは、(j+1)列目における第1の非ゼロエントリのA’vr(又はA’)配列におけるインデックスであり、j列目が行列における最後の列である場合、A’配列における次のエントリは、A’vr(又はA’)配列におけるエントリの総数である。以下のステップにおいて記載される非ゼロ配列の値の乗算結果をともに加算する関数であるローカル総和計算部は、ゼロにおいて合計値を設定することによって初期化される(ステップ269)。jがjmax未満である場合(ステップ270)、SpMTVが行われた符号化がPSCSC又はPCSCであるかに応じた動作で、値iに対して乗算及び加算計算が実行される(ステップ271)。符号化がPSCSCである場合、計算は、以下の式にしたがって行われる。
sum←sum+A’vr[i].v×x’[A’vr[i].r]
ここで、インデックスiを有するA’vr配列の要素に格納された値は、x’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’vr要素に格納された行のインデックスであり、乗算結果は、ステップ271の前の反復中に実行された乗算結果の合計に加算される。あるいは、符号化がPCSCである場合、計算は、以下の式にしたがって行われる。
sum←sum+A’[i]×x’[A’[i]]
ここで、インデックスiを有するA’配列における値は、x’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’配列における数であり、乗算結果は、ステップ271の前の反復中に実行された乗算結果の合計に加算される。いずれかの式に基づいて計算が終了すると、値iにワープサイズが加算され、加算結果がiと設定され(ステップ272)、そのスレッドに割り当てられたその列における次のエントリに処理を移動する。ルーチン260は、上述したステップ270に戻り、j列目における非ゼロ値の全てが処理されるまでステップ270〜272を繰り返す。iがimax以上である場合(ステップ270)、ワープにおける全てのスレッドは、必要に応じて同期し、ワープにおける全てのスレッドがステップ270〜272のループの実行を終了するのを可能とし、同期は、並列に動作するワープにおけるスレッドに起因して1つの実施形態では必要なく、同期から出ることがない一方で、さらなる実施形態においては同期が行われる。
ワープのサイズ、ワープにおけるスレッド数は、整数除算を使用して2で割られ、除算結果は、変数tmaxを使用して表される(ステップ274)。tの場合、ワープにおけるスレッドのidは、tmax未満であり、スレッドは、ともに組み合わせて、そのスレッドt及びスレッドidがt+tmaxである他のスレッドによって上記実行されるステップ270〜272における計算から得られる合計の削減を行う(ステップ276)。ステップ276は、上述したステップ246と同様に、半分の合計数を削減し、組み合わせた合計は、スレッドt(組み合わせを行ったスレッド)についての合計として設定され、ワープにおけるスレッドidがt+tmaxであるスレッドと前に関連付けられた合計は破棄される(ステップ276)。組み合わせに続いて、値tmaxは、整数除算を使用して半分にカットされ、tmaxとして設定される(ステップ277)。必要に応じて、合計の組み合わせを行ったスレッドは再度同期され(ステップ278)、ルーチン260は、上記ステップ275に戻る。ステップ275〜277のループの各後続反復中において、ループに参加しているスレッド数は、値tmaxの削減に起因して半分に削減される。
tがtmax以上である場合、スレッドがワープにおける第1のスレッドであるかどうか(t=0及びr<rmax)が判定される(ステップ279)。スレッドが第1のスレッドでない場合、スレッドの実行は終了する(ステップ281)。スレッドが第1のスレッドである場合、ステップ275〜277において記載された削減から生じる総合計は、密ベクトルy’に格納され(ステップ280)、そのスレッドの実行が終了する(ステップ281)。ルーチン260は、全ての起動ワープについての全てのスレッドの実行の終了によって終了する。ルーチン260についての擬似コードは、準用する図19A〜図19Bを参照して上記に示された擬似コードと同様である。
32又は64超の非ゼロエントリを有する行及び列の処理は、単一の行又は列を処理するためにスレッドのブロックを割り当てるf1B1Rカーネルを使用して最速に処理することができる。図21A〜図21Bは、1つの実施形態にかかる図15のルーチン170において使用するためのf1B1RカーネルによってSpMVを実行するためのルーチン290を示すフロー図である。ルーチン290は、PSCSR又はPCSR符号化のいずれかにおいてSpMVを行うために使用することができる。カーネルにおける全ての起動ブロックについて反復処理ループが開始される(ステップ291)。起動ブロック及び起動ブロックにおけるスレッドは、互いに並列に実行され、それゆえに、以下の説明がブロックのいずれか又はスレッドのいずれかを参照する場合には、全ての他の起動ブロック又はスレッドは、スレッド又はブロックのいずれかの実行が停止されるまで(以下のステップ311など)同時に並列にルーチン290のステップを実行する。
ローカル共有は、同じブロックにおける全てのスレッド間で開始され、スレッドがステップ300〜302において後述する計算結果を共有するのを可能とする(ステップ292)。ブロックのいずれかの固有のグローバルidを取得するために関数block__id()を使用し且つrminの値に固有のグローバルidを追加することによってブロックのいずれかに割り当てられた行のランク(図17を参照して上述したようにランクが割り当てられるパーティションの第1の行のランク)が取得され、加算結果は、ブロックに割り当てられた行のランクであり、変数rを使用して表される(ステップ293)。値rは、Ao[k+1]によって与えられる次のパーティションk+1の第1の行のランクに等しい変数であるrmaxと比較することができ、起動されるブロック数がパーティションkにおける行数(rmax−rmin)に等しい場合には比較は任意である。rがrmax未満である場合(ステップ294)、r番目のランク付けされた行のidiは、マッピング配列Aにおいて識別される(ステップ295)。rがrmax以上である場合(ステップ294)、ルーチン290は、後述するステップ311に移動し、スレッドの処理を停止する。関数block_thread_id()を使用することなどによってブロックにおけるスレッドのいずれかのローカルid(ブロック内のid)が取得され、変数tを使用して表される(ステップ296)。そのスレッドに割り当てられた第1の非ゼロエントリのインデックスは、A配列におけるi行目についてのエントリの値をtに追加することによって取得され、加算結果は、変数jを使用して表される(ステップ297)。A配列における次の(i+1)エントリが配置され、変数jmaxとして設定される(ステップ298)。以下のステップ301において記載される非ゼロ配列の値の乗算結果をともに加算する関数であるローカル総和計算部は、ゼロにおいて合計値を設定することによって初期化される(ステップ299)。jがjmax未満である場合(ステップ300)、SpMVが行われた符号化がPSCSR又はPCSRであるかに応じた動作で、値jに対して乗算及び加算計算が実行される(ステップ301)。符号化がPSCSRである場合、計算は、以下の式にしたがって行われる。
sum←sum+Avc[j].v×x[Avc[j].c]
計算において、インデックスjを有するAvc配列の要素に格納された値はxの要素によって乗算され、そのインデックスは、j番目のインデックスを有するAvc要素に格納された列のインデックスであり、乗算結果は、ステップ301の前の反復中に実行された乗算結果の合計に加算される。あるいは、符号化がPCSRである場合、計算は、以下の式にしたがって行われる。
sum←sum+Av[j]×x[v[j]]
ここで、インデックスjを有するAv配列における値は、xの要素によって乗算され、そのインデックスは、j番目のインデックスを有するA配列における数であり、乗算結果は、ステップ301の前の反復中に実行された乗算結果の合計に加算される。いずれかの式に基づいて計算が終了すると、jの値にブロックサイズが加算され、加算結果がjと設定され(ステップ302)、そのスレッドに割り当てられたその行における次のエントリに処理を移動する。ルーチン290は、上述したステップ300に戻り、そのスレッドに割り当てられるi行目における非ゼロ値が処理されるまでステップ300〜302を繰り返す。jがjmax以上である場合(ステップ300)、ブロックにおける全てのスレッドは同期し、ブロックにおける全てのスレッドがステップ300〜302のループの実行を終了するのを可能とする。
ブロックのサイズ、ブロックにおけるスレッド数は、整数除算を使用して2で割られ、除算結果は、変数tmaxを使用して表される(ステップ304)。tの場合、ブロックにおけるスレッドのidは、tmax未満であり、スレッドは、ともに組み合わせて、そのスレッドt及びスレッドidがt+tmaxである他のスレッドによって上記実行されるステップ300〜302における計算から得られる合計の削減を行う(ステップ306)。ステップ306は、半分の合計数を削減し、組み合わせた合計は、スレッドt(組み合わせを行ったスレッド)についての合計として設定され、ブロックにおけるスレッドidがt+tmaxであるスレッドと前に関連付けられた合計は破棄される(ステップ306)。組み合わせに続いて、値tmaxは、整数除算を使用して半分にカットされ、tmaxとして設定される(ステップ307)。合計の組み合わせを行ったスレッドは再度同期され(ステップ308)、ルーチン290は、上記ステップ305に戻る。ステップ305〜307のループの各後続反復中において、ループに参加しているスレッド数は、値tmaxの削減に起因して半分に削減される。
tがtmax以上である場合、スレッドがブロックにおける第1のスレッドであるかどうか(t=0及びr<rmax)が判定される(ステップ309)。スレッドが第1のスレッドでない場合、スレッドの実行は終了する(ステップ311)。スレッドが第1のスレッドである場合、ステップ305〜307において記載された削減から生じる総合計は、密ベクトルyに格納され(ステップ310)、そのスレッドの実行が終了する(ステップ311)。ルーチン290は、全てのスレッドの実行の終了によって終了する。ルーチン290はまた、以下の擬似コードを使用して表現することができる−擬似コードは、PSCSR符号化に関連して記載されるが、PCSRに関連する擬似コードを準用して記載することができる。
カーネル f1B1R(y,x,Avc,A,A,rmin,rmax) /*1B1R:1ブロック1行SpMVカーネル*/
共有 sum[BLOCKSIZE] /*sum:ブロックにおけるスレッドによって共有されるローカル合計*/
r←rmin+block_id() /*このブロックに割り当てられた頂点のランクの計算*/
(r<rmax)である場合 /*任意の正確に(rmax−rmin)ブロックが作成された場合*/
i←A[r] /*i:r番目にランク付けされた行のid*/
t←block_thread_id() /*t:ブロックにおけるローカルスレッドid*/
j←A[i]+t /*j:このスレッドに割り当てられた第1の非ゼロエントリのインデックス*/
max←A[i+1] /*jmax:(i+1)行目における第1の非ゼロエントリのインデックス*/
sum[t]←0 /*ローカル総和積算器の初期化*/
while(j<jmax) /*行の終わりに到達したかどうかのテスト*/
sum[t]←sum[t]+Avc[j].v×x[Avc[j].c] /*yi=ΣjAi,j×xjの計算
j←j+BLOCKSIZE /*このスレッドについての次の非ゼロエントリに移動*/
while文終了
if文終了
sync_block_threads() /*ブロックにおける全てのスレッドが同期している*/
max←BLOCKSIZE/2 /*tmax:ローカル合計まで追加したスレッド数*/
while(t<tmax) /*このスレッドが参加すべきであるかどうかのテスト*/
sum[t]←sum[t]+sum[t+tmax] /*2つのローカル合計を1つに削減*/
max←tmax/2 /*合計追加スレッド数を半分にカット*/
sync_block_threads() /*ブロックにおける全てのスレッドが同期している*/
while文終了
(t=0及びr<rmax)である場合 /*これはブロックにおける第1のスレッド?*/
y[i]←sum[0] /*yに総合計を格納*/
if文終了
同様に、SpMTVは、1B1Rカーネルを使用して計算することができる。図22A〜図22Bは、1つの実施形態にかかる図15のルーチン170において使用するためのf1B1Rカーネル関数によってSpMTVを実行するためのルーチン320を示すフロー図である。ルーチン320は、PSCSC又はPCSC符号化のいずれかにおいてSpMTVを行うために使用することができる。カーネルにおける全ての起動ブロックについて反復処理ループが開始される(ステップ321)。起動ブロック及び起動ブロックにおけるスレッドは、互いに並列に実行され、それゆえに、以下の説明がブロックのいずれか又はスレッドのいずれかを参照する場合には、全ての他の起動ブロック及びスレッドは、スレッド又はブロックのいずれかの実行が停止されるまで(以下のステップ341など)同時に並列にルーチン320のステップを実行する。
ローカル共有は、同じブロックにおける全てのスレッド間で開始され、スレッドがステップ330〜332において後述する計算結果を共有するのを可能とする(ステップ322)。ブロックのいずれかの固有のグローバルidを取得するために関数block__id()を使用し且つrminの値に固有のグローバルidを追加することによってブロックのいずれかに割り当てられた列のランク(図18を参照して上述したようにランクが割り当てられるパーティションの第1の列のランク)が取得され、加算結果は、ブロックに割り当てられた列のランクであり、変数rを使用して表される(ステップ323)。値rは、A’[k+1]によって与えられる次のパーティションk+1の第1の列のランクに等しい変数であるrmaxと比較することができ、起動されるブロック数がパーティションkにおける列数(rmax−rmin)に等しい場合には比較は任意である。rがrmax未満である場合(ステップ324)、r番目のランク付けされた列のidjは、マッピング配列A’において識別される(ステップ325)。rがrmax以上である場合(ステップ324)、ルーチン320は、後述するステップ341に移動する。関数block_thread_id()を使用することなどによってブロックにおけるスレッドのいずれかのローカルid(ブロック内のid)が取得され、変数tを使用して表される(ステップ326)。そのスレッドに割り当てられた第1の非ゼロエントリのインデックスは、A’配列におけるj列目についてのエントリの値をtに追加することによって取得され、加算結果は、変数iを使用して表される(ステップ327)。A’配列における次の(j+1)エントリが配置され、変数imaxとして設定される(ステップ328)。j列目が行列における最後の列でない限り、A’配列における次のエントリは、(j+1)列目における第1の非ゼロエントリのA’vr(又はA’)配列におけるインデックスであり、j列目が行列における最後の列である場合、A’配列における次のエントリは、A’r(又はA’)配列におけるエントリの総数である。以下のステップ331において記載される非ゼロ配列の値の乗算結果をともに加算する関数であるローカル総和計算部は、ゼロにおいて合計値を設定することによって初期化される(ステップ329)。iがimax未満である場合(ステップ330)、SpMTVが行われた符号化がPSCSC又はPCSCであるかに応じた動作で、値iに対して乗算及び加算計算が実行される(ステップ331)。符号化がPSCSCである場合、計算は、以下の式にしたがって行われる。
sum←sum+A’vr[i].v×x’[A’vr[i].r]
ここで、インデックスiを有するA’vr配列の要素に格納された値は、x’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’vr要素に格納された行のインデックスであり、乗算結果は、ステップ331の前の反復中に実行された乗算結果の合計に加算される。あるいは、符号化がPCSRである場合、計算は、以下の式にしたがって行われる。
sum←sum+A’[i]×x’[A’[i]]
ここで、インデックスiを有するA’配列における値は、x’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’配列における数であり、乗算結果は、ステップ331の前の反復中に実行された乗算結果の合計に加算される。いずれかの式に基づいて計算が終了すると、iの値にブロックサイズが加算され、加算結果がiと設定され(ステップ332)、その列における次のエントリに処理を移動する。ルーチン320は、上述したステップ330に戻り、j列目における非ゼロ値が処理されるまでステップ330〜332を繰り返す。iがimax以上である場合(ステップ330)、ブロックにおける全てのスレッドは同期し、ブロックにおける全てのスレッドがステップ330〜332のループの実行を終了するのを可能とする。
ブロックのサイズ、ブロックにおけるスレッド数は、整数除算を使用して2で割られ、除算結果は、変数tmaxを使用して表される(ステップ334)。tの場合、ブロックにおけるスレッドのidは、tmax未満であり、スレッドは、ともに組み合わせて、そのスレッドt及びスレッドidがt+tmaxである他のスレッドによって上記実行されるステップ330〜332における計算から得られる合計の削減を行う(ステップ336)。ステップ336は、半分の合計数を削減し、組み合わせた合計は、スレッドt(組み合わせを行ったスレッド)についての合計として設定され、上記のステップ246と同様に、ブロックにおけるスレッドidがt+tmaxであるスレッドと前に関連付けられた合計は破棄される(ステップ336)。組み合わせに続いて、値tmaxは、整数除算を使用して半分にカットされ、tmaxとして設定される(ステップ337)。合計の組み合わせを行ったスレッドは再度同期され(ステップ338)、ルーチン320は、上記ステップ335に戻る。ステップ335〜337のループの各後続反復中において、ループに参加しているスレッド数は、値tmaxの削減に起因して半分に削減される。
tがtmax以上である場合、スレッドがブロックにおける第1のスレッドであるかどうか(t=0及びr<rmax)が判定される(ステップ339)。スレッドが第1のスレッドでない場合、スレッドの実行は終了する(ステップ341)。スレッドが第1のスレッドである場合、ステップ335〜337において記載された削減から生じる総合計は、密ベクトルy’に格納され(ステップ340)、そのスレッドの実行が終了する(ステップ341)。ルーチン320は、全てのスレッドの実行の終了によって終了する。ルーチン320についての擬似コードは、準用する図21A〜図21Bを参照して上記に示された擬似コードと同様である。
図3及び図4を参照して上述したように、SpMV及びSpMTVの結果の一般的な用途の1つは、べき乗法を使用して実行されることができるページランク(登録商標)アルゴリズムなどのランク付けアルゴリズムである。図23は、1つの実施形態にかかる図6及び図9の方法60及び90において使用するためのべき乗法を実行するためのルーチン350を示すフロー図である。ルーチン350は、図6及び図9の方法において使用されるルーチンを参照して上述したSpMV及びSpMTVの全てのバリエーションの結果に適用するために使用することができる。d∈(0,1)を減衰係数とし、nをウェブページ数とする。Pをn×nの正方行列とする。
Figure 2016119084
ページjからページiへのリンクがある場合
それ以外
ここで、Ljは、ページjから出るリンクの数である。x及びyをサイズnの2つの密ベクトルとし、ε*を停止閾値とする。初期のページランクの確率分布がベクトルxについて設定される(ステップ351)。εの値がε*の値未満である場合に反復処理ループ(ステップ352〜356)が実行される(ステップ352)。値yは、以下の式に基づいて設定される。
Figure 2016119084
ここで、1は、(n×1)列ベクトルである(ステップ353)。値εは、以下の式にしたがって決定される。
Figure 2016119084
ここで、値εは、y及びxの差異に等しい(ステップ53)。最後に、xは、値yに等しくなるように設定される。ルーチン350は、ループの次の反復に移動し(ステップ356)、εが値ε*未満であるまでループ(352〜356)を介した処理を継続し、その後にルーチン350は終了する。ルーチン350はまた、以下の擬似コードを使用して表現することができる。
x←x0 /*初期のページランクの確率分布*/
ループ
Figure 2016119084
ε←|y−x|
x←y
ε<ε*まで
上述したシステム30及び方法60、90は、疎行列、すなわち、SCSR、SCSC、PCSR、PCSC、PSCSR及びPSCSCについての6つの新たな符号化の利点を活用し、従来技術よりも優れている。6つのうち、SCSR及びSCSCは、それらの構造化されていない相手方のCSR及びCSCと全く同じ空間の複雑性を有する。一方、残りの4つの新たな符号化は、僅かに高い空間的要件を有する。PCSR及びPSCSRについて、過剰な空間的オーバーヘッドは、非ゼロ行(すなわち、少なくとも1つの非ゼロエントリを有する行)の数において線形である。PCSC及びPSCSCについて、オーバーヘッドは、非ゼロ列(すなわち、少なくとも1つの非ゼロエントリを有する列)の数において線形である。PSCSR符号化の有効性を評価するために、フロリダ大学の疎行列コレクションにおいてみられる様々な疎行列に対してSpMVを実行する実験を行った。行列のパーティションを作成するために、実験における全てのデータセットについてAs=[1,32,1024,∞]が使用された。その結果、PSCSRの第1のパーティションは、少なくとも1つの非ゼロを有し且つ32未満の非ゼロエントリを有する行を含み、第2のパーティションは、32以上1024未満の非ゼロエントリを有する行を含み、第3のパーティションは、1024以上の非ゼロエントリを有する行を含んでいた。
表1は、(「#行」とラベル付けされている)行数、非ゼロ行の数(「#非ゼロ行」)、行列の非ゼロエントリの総数(#非ゼロ)、行あたりの非ゼロエントリの平均数(「平均#非ゼロ/行」)、(SCSRと同じである)CSR符号化のサイズ、PSCSR符号化のサイズ、及び、元のCSR符号化の割合としてのPSCSRの過剰な空間的オーバーヘッドを含む、実験において使用された疎行列の統計を示している。観察されるように、PSCSRの過剰な空間的オーバーヘッドは、行の非ゼロエントリの平均数が増加するのにともない低減する。オーバーヘッドは、行あたり2.1個の非ゼロエントリのみを有する行列「wiki−Talk」について最高である(19.06%)一方で、オーバーヘッドは、行あたり22.3個の非ゼロエントリの平均を有する行列「eu−2005」について2.19%まで低下する。表1における全ての行列にわたって平均化したPSCSRの過剰な空間的オーバーヘッドは5.45%である。
Figure 2016119084
λを通常は実装依存定数であるPSCSR(又はPSCSC)のA(又はA’)における要素のサイズに対するAvc(又はA’vc)における要素のサイズの比率とする。bをPSCSR(又はPSCSC)における行(又は列)の数に対する非ゼロエントリの数の比率とし、γをPSCSR(又はPSCSC)における行(又は列)の総数に対する非ゼロ行(又は列)の数の比率とする。δを通常のCSR(又はCSC)に対するPSCSR/PCSR(又はPSCSC/PCSC)の過剰な空間的オーバーヘッドとする。そのδは、
Figure 2016119084
によって与えられて示すことができる。実施形態の1つにおいて、Avcにおける要素は、Aにおける要素の2倍の大きさであり、それゆえにλ=2である。行列「wiki−Talk」について、γ=2,369,181/2,394,385=0.9895、b=5,021,410/2,394,385=2.097を有する。それゆえに、過剰な空間的オーバーヘッドは、δ=γ(λb+1)−1=0.9895×(2×2.097+1)−1=19.05%であり、実験において観察された実際のオーバーヘッドに近い。行列「eu−2005」について、γ=862,664/862,664=1、b=19,235,140/862,664=22.30である。それゆえに、δ=γ(λb+1)−1=1×(2×22.30+1)−1=2.19%である(経験的な数と同じである)。δ=γ(λb+1)−1という解析式は、表1に記録されたPSCSRの実際の過剰なオーバーヘッドと非常によく一致する(大抵の場合、検出可能なエラーなし)ことを確認することができる。この式は、(1+δ)の係数を乗じた通常のCSR(又はCSC)符号化のサイズとしてPSCSR/PCSR(又はPSCSC/PCSC)符号化のサイズを正確に予測することを可能とする。予測された過剰な空間的オーバーヘッドは、PSCSR/PCSR符号化を形成することが、特定のハードウェアセットアップについて有用であるかどうか、又は、不十分なハードウェアリソースがそのような符号化の形成又は使用を不可能とするかどうかを判定するために使用することができる。
図13を参照して始めて記載されたGPUベースのSpMVルーチンの性能が表1の全ての行列についてテストされた。比較のため、CPUベースのSpMV実装の結果が表3に含まれる。使用されたテストマシンは、それぞれが3.46GHzで動作する6個のコアを有する2つのIntel Xeon X5690プロセッサを有する。単一のX5690プロセッサのキャッシュサイズは12MBであり、同じボックス内のNvidia(登録商標)のGTX580 GPUのL2キャッシュサイズよりもかなり大きい。単一のCPUコアに対するSpMVにおけるGPUの高速化を測定するために、CPU実装は、単一スレッドを使用する。
Figure 2016119084
CPUベースの実装についての語は、実装が大規模グラフにおける最先端のSpMVにあるということである。テストマシンの単一コアを使用して、CPUは、Kwakら、「What is Twitter, a Social Network or a News Media?」、Proceedings of the 19th international conference on World Wide Web、2010年、591〜600ページによって導入された、ある4170万人のユーザ(すなわち、行)及び14.7億の接続(すなわち、非ゼロ)のTwitter(登録商標)フォロワーネットワークにおいて36秒でSpMVの反復を実行する。同じTwitter(登録商標)フォロワーネットワークは、ベンチマークデータセットとしての様々なSpMV実装によって使用されている。表2は、文献において見出された最良のCPUベースの大規模SpMVシステムのいくつかの実行時性能を比較する。最初の4つのシステムについての番号は、Gonzalezら、PowerGraph: Distributed Graph−Parallel Computation on Natural Graphs、OSDI、2012年、第12巻、第1号、2ページによって報告され、そのPowerGraphシステムは、それぞれが10ギガイーサネット(登録商標)を介して接続された23GBのRAMを有する2つのクワッドコアのIntel Xeon X5570プロセッサを有するAmazon EC2 cc1.4xlarge Linux(登録商標)インスタンスの64のノードクラスタにおいて3.6秒でSpMVの反復を実行することができる。HiperGraphと称される単一ノードシステムはまた、全12個のコアが使用された場合に3.6秒でSpMVの反復を終了することができる。コアあたりの高速化を測定するために、HiperGraphは、同じTwitter(登録商標)のフォロワーネットワーク上の完全なパスを終了するために36秒かかった単一スレッドモードで実行された。換言すれば、PowerGraphは、3.46GHz@の単一コアでHiperGraphに対して10倍の高速化を得るために2.93GHz@の512個のコアを使用する。SpMVに対するHiperGraphの単一コアの性能は競合的であり、それゆえに、後述する実験についてのベースラインとして使用された。
表3は、CPU(すなわち、1コアを有するHiperGraph)及び様々な疎行列に対するGPUベースのSpMVの性能を示している。GPUベースのSpMVの実行は、以下の3つの部分にわけられる:(a)PSCSR符号化によって必要とされるパーティションを生成するパーティション、(b)CPUからCPUに対してパーティション化された行列を転送する負荷、及び、(c)GPUにおいてSpMVの単一反復を実行するSpMV。(a)及び(b)の双方の部分は、そのコストがSpMVの複数の反復にわたって償却されることができるワンタイム動作である。表3の最後の列は、シーケンシャルCPUベースのSpMVに対するGPUベースの高速化を示しており、100回の反復にわたって積算される。行列が既にPSCSRフォーマットで符号化されている場合、パーティション時間は、表3において単にゼロとしなければならないことに留意されたい。高速化は、23.5×から38.3×まで及び、平均は31倍速である。
Figure 2016119084
実験は、さらに、上述した方法を実装するGPUベースのハードウェアが上述した従来技術を越えて実現することができるという符号化のさらなる利点を示す。大幅に高速化するのみならず、GPUはまた、CPUよりもはるかに安価である:GTX 580のGPUのコストは、現在$400未満であるが、単一のXeon X5690プロセッサは、約$1800のコストがかかる。X5690の6個全てのコアが使用された場合に我々のシーケンシャルなCPU SpMVの線形高速化を仮定すると、GPUは、X5690プロセッサにおける完全に並列化されたCPUの実装よりも31/6=5.2×高速化するであろう。ドルあたりの性能の観点から、GPUベースのSpMVは、上記使用されたテストマシンについて約23×以上の費用対効果がある。
GPUベースの解決策はまた、既存のシステムのGPUの追加又は更新が同じボックス内にCPUを追加又は更新するよりもはるかに容易であることから、より拡張性がある。コモディティハードウェアについて、CPUの最大数は、通常は2又は4である。一方、単一のコモディティGPUサーバは、8個のGPUまで保持することができ、それぞれが数千コアまで有することができる(例えば、Nvidia(登録商標)のTesla K40 GPUは2880のコアを有する)。それゆえに、ドルあたりの性能又はワットあたりのFLOPSのみならず、データセンタのために特に重要である1Uラック空間あたりの性能においても、GPUはCPUに優れている。
本発明は、その実施形態を参照して具体的に示されて説明されたが、当業者は、形態及び詳細における上述した及び他の変更が本発明の精神及び範囲から逸脱することなく行われ得ることを理解するであろう。

Claims (10)

  1. 構造化疎行列表現を取得するためのコンピュータ実装方法において、
    各部分が行及び列のうちの1つを含む行列の部分における1つ以上の順序で配置された1つ以上の非ゼロエントリを含む前記行列の構造化圧縮表現を取得することと、
    前記表現に含まれる複合配列であり、各要素が非ゼロエントリのいずれかとその非ゼロエントリを含む前記部分のいずれかのインデックスとを含む1つ以上の要素を含む前記複合配列を取得することと、
    前記順序のうちの1つ以上における第1である前記非ゼロエントリを含む各要素の前記複合配列におけるインデックスを含むインデックス配列であり、さらに前記行列における複数の前記非ゼロエントリを含む前記インデックス配列を取得することとを備える、方法。
  2. 各複合配列の要素インデックスが、列優先順序でその非ゼロエントリの行のインデックスを含み、前記順序が各列における前記非ゼロエントリの順序を含む、請求項1に記載の方法。
  3. さらに、
    各列を処理することによって複数の要素を含む密ベクトルによって前記行列の疎行列転置ベクトル乗算を実行することと、
    前記インデックス配列を使用してその列における前記非ゼロエントリを含む全ての前記複合配列要素からなる前記複合配列におけるインデックスを識別することと、
    そのエントリの前記インデックスを使用して前記密ベクトルの要素のいずれかによってその列における前記非ゼロエントリのそれぞれを乗算することと、
    その列における前記非ゼロエントリのそれぞれについての乗算結果を加算し、異なる密ベクトルの加算結果を格納することとを備える、請求項2に記載の方法。
  4. 各複合配列の要素インデックスが、行優先順序でその非ゼロエントリの列のインデックスを含み、前記順序が各行における前記非ゼロエントリの順序を含む、請求項1に記載の方法。
  5. さらに、
    各行を処理することによって複数の要素を含む密ベクトルによって前記行列の疎行列ベクトルの乗算を実行することと、
    前記インデックス配列を使用してその行における前記非ゼロエントリを含む全ての前記複合配列要素からなる前記複合配列におけるインデックスを識別することと、
    そのエントリの前記インデックスを使用して前記密ベクトルの要素のいずれかによってその行における前記非ゼロエントリのそれぞれを乗算することと、
    その行における前記非ゼロエントリのそれぞれについての乗算結果を加算し、異なる密ベクトルの加算結果を格納することとを備える、請求項4に記載の方法。
  6. さらに、
    前記構造化圧縮表現を処理することと、
    前記処理の結果を使用してランク付け分析を実行することとを備える、請求項1に記載の方法。
  7. 効率的な疎行列表現のためのコンピュータ実装方法において、
    1つ以上の非ゼロエントリを含む疎行列の圧縮表現を取得することであり、前記行列の行及び列のいずれかを含む前記行列の部分が前記行列におけるそれらの位置に基づいてインデックス付けされることと、
    前記行列の部分について複数のパーティションを定義することと、
    前記圧縮表現を使用して前記部分のそれぞれにおける複数の前記非ゼロエントリを取得することと、
    その部分における複数の前記非ゼロエントリに基づいて前記パーティションのいずれかを有する部分のそれぞれを関連付けることと、
    前記パーティションのそれぞれに関連する全ての部分であり、それらのインデックスの順序でリスト化された前記部分のリストを作成することと、
    前記リストを含むマッピング配列を含む前記行列のパーティション化された圧縮表現を作成することとを備える、方法。
  8. さらに、
    前記マッピング配列を処理することと、
    前記処理の結果に基づいてランク付け分析を実行することとを備える、請求項7に記載の方法。
  9. さらに、
    前記パーティションのそれぞれに関連する前記行列の部分における非ゼロエントリの数に基づいて前記パーティションをインデックス付けることと、
    前記パーティションのそれぞれについて、前記パーティションのインデックスに基づいて前記配列におけるそのパーティションに先行する全てのパーティションのサイズであり、前記パーティションに関連する複数の部分の数を含む前記サイズを識別し、そのパーティションについての前記リストを前記サイズに基づく位置において前記配列に挿入することとを備える、請求項7に記載の方法。
  10. 前記行列の行を含む前記各部分が、さらに、
    前記マッピング配列にリスト化されたパーティションのそれぞれの行のそれぞれをシーケンシャルに処理することによって複数の要素を含む密ベクトルによって前記行列の疎行列ベクトル乗算を実行することと、
    その行における前記非ゼロエントリのそれぞれに関連付けられた前記圧縮表現におけるインデックスを識別することと、
    その非ゼロエントリの前記インデックスを使用して前記密ベクトルの要素のいずれかによってその行における前記非ゼロエントリのそれぞれを乗算することと、
    その行における前記非ゼロエントリのそれぞれについての乗算結果を加算し、異なる密ベクトルにおける加算結果を格納することとを備える、請求項7に記載の方法。
JP2015240835A 2014-12-22 2015-12-10 効率的な疎行列表現及び処理のためのコンピュータ実装システム及び方法 Active JP6630558B2 (ja)

Applications Claiming Priority (2)

Application Number Priority Date Filing Date Title
US14/580,110 2014-12-22
US14/580,110 US9760538B2 (en) 2014-12-22 2014-12-22 Computer-implemented system and method for efficient sparse matrix representation and processing

Publications (3)

Publication Number Publication Date
JP2016119084A true JP2016119084A (ja) 2016-06-30
JP2016119084A5 JP2016119084A5 (ja) 2019-01-24
JP6630558B2 JP6630558B2 (ja) 2020-01-15

Family

ID=54849867

Family Applications (1)

Application Number Title Priority Date Filing Date
JP2015240835A Active JP6630558B2 (ja) 2014-12-22 2015-12-10 効率的な疎行列表現及び処理のためのコンピュータ実装システム及び方法

Country Status (3)

Country Link
US (2) US9760538B2 (ja)
EP (1) EP3037980A3 (ja)
JP (1) JP6630558B2 (ja)

Cited By (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
WO2018135599A3 (ja) * 2017-01-19 2018-09-13 新妻弘崇 遅延疎行列

Families Citing this family (40)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US9760538B2 (en) * 2014-12-22 2017-09-12 Palo Alto Research Center Incorporated Computer-implemented system and method for efficient sparse matrix representation and processing
US9606934B2 (en) * 2015-02-02 2017-03-28 International Business Machines Corporation Matrix ordering for cache efficiency in performing large sparse matrix operations
US10600151B2 (en) * 2015-06-05 2020-03-24 Apple Inc. Automatic determination of a region of influence
US10061832B2 (en) * 2016-11-28 2018-08-28 Oracle International Corporation Database tuple-encoding-aware data partitioning in a direct memory access engine
US10073815B2 (en) * 2016-05-31 2018-09-11 Palo Alto Research Cener Incorporated System and method for speeding up general matrix-matrix multiplication on the GPU
US10067910B2 (en) * 2016-07-01 2018-09-04 Palo Alto Research Center Incorporated System and method for GPU maximum register count optimization applied to general matrix-matrix multiplication
US10476896B2 (en) * 2016-09-13 2019-11-12 Accenture Global Solutions Limited Malicious threat detection through time series graph analysis
US10489063B2 (en) * 2016-12-19 2019-11-26 Intel Corporation Memory-to-memory instructions to accelerate sparse-matrix by dense-vector and sparse-vector by dense-vector multiplication
WO2018134740A2 (en) * 2017-01-22 2018-07-26 Gsi Technology Inc. Sparse matrix multiplication in associative memory device
US10146740B1 (en) * 2017-03-08 2018-12-04 Symantec Corporation Sparse data set processing
CN108664447B (zh) * 2017-03-31 2022-05-17 华为技术有限公司 一种矩阵与矢量的乘法运算方法及装置
US10346944B2 (en) * 2017-04-09 2019-07-09 Intel Corporation Machine learning sparse computation mechanism
US10409732B2 (en) 2017-05-31 2019-09-10 Nxp Usa, Inc. Sparse matrix accelerator
WO2019090325A1 (en) 2017-11-06 2019-05-09 Neuralmagic, Inc. Methods and systems for improved transforms in convolutional neural networks
US20190156214A1 (en) 2017-11-18 2019-05-23 Neuralmagic Inc. Systems and methods for exchange of data in distributed training of machine learning algorithms
CN110120251A (zh) * 2018-02-07 2019-08-13 北京第一视角科技有限公司 基于Spark的多维健康数据的统计分析方法及系统
US11995448B1 (en) * 2018-02-08 2024-05-28 Marvell Asia Pte Ltd Method and apparatus for performing machine learning operations in parallel on machine learning hardware
US10691772B2 (en) * 2018-04-20 2020-06-23 Advanced Micro Devices, Inc. High-performance sparse triangular solve on graphics processing units
US10929778B1 (en) 2018-05-22 2021-02-23 Marvell Asia Pte, Ltd. Address interleaving for machine learning
US10997510B1 (en) 2018-05-22 2021-05-04 Marvell Asia Pte, Ltd. Architecture to support tanh and sigmoid operations for inference acceleration in machine learning
US11016801B1 (en) 2018-05-22 2021-05-25 Marvell Asia Pte, Ltd. Architecture to support color scheme-based synchronization for machine learning
US10929779B1 (en) 2018-05-22 2021-02-23 Marvell Asia Pte, Ltd. Architecture to support synchronization between core and inference engine for machine learning
US11449363B2 (en) 2018-05-31 2022-09-20 Neuralmagic Inc. Systems and methods for improved neural network execution
US11216732B2 (en) 2018-05-31 2022-01-04 Neuralmagic Inc. Systems and methods for generation of sparse code for convolutional neural networks
WO2021061172A1 (en) * 2019-09-27 2021-04-01 Neuralmagic Inc. System and method of executing neural networks
US10832133B2 (en) 2018-05-31 2020-11-10 Neuralmagic Inc. System and method of executing neural networks
US10963787B2 (en) 2018-05-31 2021-03-30 Neuralmagic Inc. Systems and methods for generation of sparse code for convolutional neural networks
CN113190791A (zh) 2018-08-06 2021-07-30 华为技术有限公司 矩阵的处理方法、装置及逻辑电路
CN110147222B (zh) * 2018-09-18 2021-02-05 安徽寒武纪信息科技有限公司 运算装置及方法
WO2020072274A1 (en) 2018-10-01 2020-04-09 Neuralmagic Inc. Systems and methods for neural network pruning with accuracy preservation
US11132423B2 (en) 2018-10-31 2021-09-28 Hewlett Packard Enterprise Development Lp Partition matrices into sub-matrices that include nonzero elements
US12008475B2 (en) * 2018-11-14 2024-06-11 Nvidia Corporation Transposed sparse matrix multiply by dense matrix for neural network training
US11544559B2 (en) 2019-01-08 2023-01-03 Neuralmagic Inc. System and method for executing convolution in a neural network
US11195095B2 (en) 2019-08-08 2021-12-07 Neuralmagic Inc. System and method of accelerating execution of a neural network
EP4073667A1 (en) * 2020-01-15 2022-10-19 Google LLC Sparse matrix operations for deep learning
CN112416433B (zh) * 2020-11-24 2023-01-17 中科寒武纪科技股份有限公司 一种数据处理装置、数据处理方法及相关产品
US11556757B1 (en) 2020-12-10 2023-01-17 Neuralmagic Ltd. System and method of executing deep tensor columns in neural networks
CN113297537B (zh) * 2021-06-04 2022-10-25 中国科学院软件研究所 一种稀疏结构化三角方程组求解的高性能实现方法和装置
CN113377534A (zh) * 2021-06-08 2021-09-10 东南大学 一种基于csr格式的高性能稀疏矩阵向量乘法计算方法
US11960982B1 (en) 2021-10-21 2024-04-16 Neuralmagic, Inc. System and method of determining and executing deep tensor columns in neural networks

Citations (2)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
JP2008507789A (ja) * 2004-07-26 2008-03-13 ソースファイア インコーポレイテッド マルチパターン検索のための方法およびシステム
US20110307685A1 (en) * 2010-06-11 2011-12-15 Song William S Processor for Large Graph Algorithm Computations and Matrix Operations

Family Cites Families (7)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US8775495B2 (en) * 2006-02-13 2014-07-08 Indiana University Research And Technology Compression system and method for accelerating sparse matrix computations
US20080126467A1 (en) * 2006-09-26 2008-05-29 Anwar Ghuloum Technique for transposing nonsymmetric sparse matrices
US9317482B2 (en) * 2012-10-14 2016-04-19 Microsoft Technology Licensing, Llc Universal FPGA/ASIC matrix-vector multiplication architecture
US9760538B2 (en) * 2014-12-22 2017-09-12 Palo Alto Research Center Incorporated Computer-implemented system and method for efficient sparse matrix representation and processing
US20160259826A1 (en) * 2015-03-02 2016-09-08 International Business Machines Corporation Parallelized Hybrid Sparse Matrix Representations for Performing Personalized Content Ranking
US10067910B2 (en) * 2016-07-01 2018-09-04 Palo Alto Research Center Incorporated System and method for GPU maximum register count optimization applied to general matrix-matrix multiplication
US10180928B2 (en) * 2016-12-31 2019-01-15 Intel Corporation Heterogeneous hardware accelerator architecture for processing sparse matrix data with skewed non-zero distributions

Patent Citations (2)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
JP2008507789A (ja) * 2004-07-26 2008-03-13 ソースファイア インコーポレイテッド マルチパターン検索のための方法およびシステム
US20110307685A1 (en) * 2010-06-11 2011-12-15 Song William S Processor for Large Graph Algorithm Computations and Matrix Operations

Non-Patent Citations (2)

* Cited by examiner, † Cited by third party
Title
ARASH ASHARI ET. AL.: "Fast Sparse Matrix-Vector Multiplication on GPUs for Graph Applications", PROCEEDINGS OF THE INTERNATIONAL CONFERENCE FOR HIGH PERFORMANCE COMPUTING, NETWORKING, STORAGE AND, JPN6019017059, 16 November 2014 (2014-11-16), pages 781 - 792, XP032725190, ISSN: 0004032928, DOI: 10.1109/SC.2014.69 *
椋木大地ほか: "GPUにおける高速なCRS形式疎行列ベクトル積の実装", 情報処理学会研究報告 2012(平成24)年度▲6▼ [DVD−ROM], vol. Vol. 2013-HPC-138, No. 5, JPN6019017056, 15 April 2013 (2013-04-15), JP, pages 1 - 7, ISSN: 0004032927 *

Cited By (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
WO2018135599A3 (ja) * 2017-01-19 2018-09-13 新妻弘崇 遅延疎行列

Also Published As

Publication number Publication date
US10296556B2 (en) 2019-05-21
US20170371839A1 (en) 2017-12-28
US9760538B2 (en) 2017-09-12
EP3037980A2 (en) 2016-06-29
JP6630558B2 (ja) 2020-01-15
EP3037980A3 (en) 2016-07-06
US20160179750A1 (en) 2016-06-23

Similar Documents

Publication Publication Date Title
JP6630558B2 (ja) 効率的な疎行列表現及び処理のためのコンピュータ実装システム及び方法
Ashari et al. Fast sparse matrix-vector multiplication on GPUs for graph applications
Besta et al. Slimsell: A vectorizable graph representation for breadth-first search
Li et al. GPU-accelerated preconditioned iterative linear solvers
Dang et al. CUDA-enabled Sparse Matrix–Vector Multiplication on GPUs using atomic operations
Li et al. Walking in the cloud: Parallel simrank at scale
CN108170639B (zh) 基于分布式环境的张量cp分解实现方法
Ma et al. Optimizing sparse tensor times matrix on GPUs
Choi et al. Blocking optimization techniques for sparse tensor computation
Corral et al. Execution of a parallel edge-based Navier–Stokes solver on commodity graphics processor units
Elafrou et al. SparseX: A library for high-performance sparse matrix-vector multiplication on multicore platforms
Canny et al. Machine learning at the limit
Zheng et al. Semi-external memory sparse matrix multiplication for billion-node graphs
Lai et al. Accelerating Strassen-Winograd's matrix multiplication algorithm on GPUs
Martínez-del-Amor et al. Population Dynamics P systems on CUDA
Roy et al. Numa-caffe: Numa-aware deep learning neural networks
Loffeld et al. Considerations on the implementation and use of Anderson acceleration on distributed memory and GPU-based parallel computers
Burchard et al. ipug: Accelerating breadth-first graph traversals using manycore graphcore ipus
Nagasaka et al. Cache-aware sparse matrix formats for Kepler GPU
Chen et al. fgSpMSpV: A fine-grained parallel SpMSpV framework on HPC platforms
Demirci et al. Scaling sparse matrix-matrix multiplication in the accumulo database
JP7109576B2 (ja) グラフィックス処理ユニット上の高性能スパース三角解
Pan et al. G-slide: A gpu-based sub-linear deep learning engine via lsh sparsification
Busato et al. A dynamic approach for workload partitioning on GPU architectures
Feng et al. A segment‐based sparse matrix–vector multiplication on CUDA

Legal Events

Date Code Title Description
RD03 Notification of appointment of power of attorney

Free format text: JAPANESE INTERMEDIATE CODE: A7423

Effective date: 20151221

RD04 Notification of resignation of power of attorney

Free format text: JAPANESE INTERMEDIATE CODE: A7424

Effective date: 20151222

A521 Request for written amendment filed

Free format text: JAPANESE INTERMEDIATE CODE: A523

Effective date: 20181210

A621 Written request for application examination

Free format text: JAPANESE INTERMEDIATE CODE: A621

Effective date: 20181210

A871 Explanation of circumstances concerning accelerated examination

Free format text: JAPANESE INTERMEDIATE CODE: A871

Effective date: 20181210

A975 Report on accelerated examination

Free format text: JAPANESE INTERMEDIATE CODE: A971005

Effective date: 20190208

A977 Report on retrieval

Free format text: JAPANESE INTERMEDIATE CODE: A971007

Effective date: 20190417

A131 Notification of reasons for refusal

Free format text: JAPANESE INTERMEDIATE CODE: A131

Effective date: 20190516

A521 Request for written amendment filed

Free format text: JAPANESE INTERMEDIATE CODE: A523

Effective date: 20190809

TRDD Decision of grant or rejection written
A01 Written decision to grant a patent or to grant a registration (utility model)

Free format text: JAPANESE INTERMEDIATE CODE: A01

Effective date: 20191107

A61 First payment of annual fees (during grant procedure)

Free format text: JAPANESE INTERMEDIATE CODE: A61

Effective date: 20191209

R150 Certificate of patent or registration of utility model

Ref document number: 6630558

Country of ref document: JP

Free format text: JAPANESE INTERMEDIATE CODE: R150

R250 Receipt of annual fees

Free format text: JAPANESE INTERMEDIATE CODE: R250

R250 Receipt of annual fees

Free format text: JAPANESE INTERMEDIATE CODE: R250