JP6630558B2 - 効率的な疎行列表現及び処理のためのコンピュータ実装システム及び方法 - Google Patents
効率的な疎行列表現及び処理のためのコンピュータ実装システム及び方法 Download PDFInfo
- Publication number
- JP6630558B2 JP6630558B2 JP2015240835A JP2015240835A JP6630558B2 JP 6630558 B2 JP6630558 B2 JP 6630558B2 JP 2015240835 A JP2015240835 A JP 2015240835A JP 2015240835 A JP2015240835 A JP 2015240835A JP 6630558 B2 JP6630558 B2 JP 6630558B2
- Authority
- JP
- Japan
- Prior art keywords
- zero
- index
- row
- column
- matrix
- 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.)
- Active
Links
- 239000011159 matrix material Substances 0.000 title claims description 254
- 238000012545 processing Methods 0.000 title claims description 139
- 238000000034 method Methods 0.000 title claims description 113
- 238000005192 partition Methods 0.000 claims description 198
- 230000006870 function Effects 0.000 claims description 84
- 239000013598 vector Substances 0.000 claims description 62
- 238000013507 mapping Methods 0.000 claims description 38
- 230000008569 process Effects 0.000 claims description 34
- 239000002131 composite material Substances 0.000 claims description 18
- 230000001360 synchronised effect Effects 0.000 claims description 13
- 230000014509 gene expression Effects 0.000 claims description 7
- 230000004044 response Effects 0.000 claims description 3
- 238000004364 calculation method Methods 0.000 description 73
- 241001409283 Spartina mottle virus Species 0.000 description 72
- 238000010586 diagram Methods 0.000 description 36
- 238000004422 calculation algorithm Methods 0.000 description 26
- 238000003491 array Methods 0.000 description 15
- 238000012360 testing method Methods 0.000 description 15
- 238000000638 solvent extraction Methods 0.000 description 13
- 230000015654 memory Effects 0.000 description 12
- 238000002474 experimental method Methods 0.000 description 10
- 230000009467 reduction Effects 0.000 description 9
- 230000008901 benefit Effects 0.000 description 7
- 230000004913 activation Effects 0.000 description 5
- 230000001186 cumulative effect Effects 0.000 description 4
- 238000001193 catalytic steam reforming Methods 0.000 description 3
- 230000001133 acceleration Effects 0.000 description 2
- 230000001174 ascending effect Effects 0.000 description 2
- 230000009286 beneficial effect Effects 0.000 description 2
- 230000006835 compression Effects 0.000 description 2
- 238000007906 compression Methods 0.000 description 2
- 230000007423 decrease Effects 0.000 description 2
- 230000001419 dependent effect Effects 0.000 description 2
- 230000006872 improvement Effects 0.000 description 2
- 239000002699 waste material Substances 0.000 description 2
- 235000019227 E-number Nutrition 0.000 description 1
- 239000004243 E-number Substances 0.000 description 1
- 230000002411 adverse Effects 0.000 description 1
- 238000013459 approach Methods 0.000 description 1
- 230000015556 catabolic process Effects 0.000 description 1
- 230000001413 cellular effect Effects 0.000 description 1
- 230000008859 change Effects 0.000 description 1
- 230000002860 competitive effect Effects 0.000 description 1
- 238000004590 computer program Methods 0.000 description 1
- 238000006731 degradation reaction Methods 0.000 description 1
- 230000000694 effects Effects 0.000 description 1
- 230000000977 initiatory effect Effects 0.000 description 1
- 238000012423 maintenance Methods 0.000 description 1
- 238000012986 modification Methods 0.000 description 1
- 230000004048 modification Effects 0.000 description 1
- 238000012546 transfer Methods 0.000 description 1
- 230000007704 transition Effects 0.000 description 1
Images
Classifications
-
- G—PHYSICS
- G06—COMPUTING; CALCULATING OR COUNTING
- G06F—ELECTRIC DIGITAL DATA PROCESSING
- G06F17/00—Digital computing or data processing equipment or methods, specially adapted for specific functions
- G06F17/10—Complex mathematical operations
- G06F17/16—Matrix 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)
Description
・値配列Avは、行優先順序でAにおける全ての非ゼロエントリの値を保持し、
・列配列Acは、行優先順序でAにおける全ての非ゼロエントリの列を保持し、
・行インデックス配列Arは、Aにおける各行の第1の非ゼロエントリのAvにおけるインデックスを保持し、Av配列におけるエントリの総数であるAr配列における最後のエントリを有する。
sum←sum+Av[j]×x[Ac[j]]
が実行される(ステップ16)。計算において、インデックスjを有するAv配列における値はxの要素によって乗算され、そのインデックスは、j番目のインデックスを有するAc配列の数である。乗算結果は、ステップ16の前の反復中に実行された乗算結果の合計に加算される。上記ステップ14において設定されたように、本方法におけるステップ16の第1の反復中において合計はゼロである。計算が終了すると、jの値に1が加算され、加算結果はjに設定され(ステップ17)、その行における次の列のエントリに処理を移動する。本方法は、上述したステップ15に戻り、i行目における非ゼロ値が処理されるまでステップ15〜17を繰り返す。jがjmax以上である場合(ステップ15)、ループ15〜17における反復中に乗算結果を加算した合計は、密ベクトルyに格納される(ステップ18)。反復処理ループは、次の行に移動し(ステップ19)、全ての行が処理されるまでループ(ステップ11〜19)を介して処理を継続し、方法10は終了する。
for i=0からm−1 /*疎行列Aのm行に対するループ*/
j←Ar[i] /*j:i行目における第1の非ゼロエントリのAvにおけるインデックス*/
jmax←Ar[i+1] /*jmax:(i+1)行目における第1の非ゼロエントリのAvにおけるインデックス*/
sum←0 /*総和積算器を初期化*/
while (j<jmax) /*行の終わりに到達したかどうかのテスト/*
sum←sum+Av[j]×x[Ac[j]] /*yi=ΣjAi,j×xjを計算*/
j←j+1 /*i行目における次の非ゼロの列に移動*/
while文終了
y[i]←sum /*yに結果を格納*/
for文終了
sum←sum+Av[j]×x[Ac[j]]
ここで、それぞれインデックスj、j及びAc[j]を有するAv、Ac及びxの3つの配列がアクセスされる。Av及びAcは、双方とも、通常はSpMVにおいてAr及びxのサイズよりもはるかに大きいe個の要素を有することに留意されたい。CPU及びGPUの双方を含む現代のプロセッサにおけるこのサイズ差の意味は、上記総和行がこれらの配列のサイズに応じて2つ又は3つのキャッシュミスを引き起こす可能性が最も高いということである。3つの配列のうち、x(入力密ベクトルを格納する配列)は、通常最小であり、それゆえに、それは、プロセッサのL2キャッシュに適合する最善の機会を有する。もちろん、行列Aが十分に小さい場合、全てが適合し、自明なケースである。しかしながら、非自明のSpMVの問題について、Av又はAcのいずれもL2に収まることを仮定するべきではない。換言すれば、Av及びAcの双方にアクセスすることは、2つの別個のキャッシュミスをトリガすることができ、SpMVの性能に悪影響を与えることがある。キャッシュミスは、SpMVの速度を大幅に低減することができ、大抵の場合にメモリ参照の局所性をほとんど示さない。
・値配列A’vは、列優先順序でAにおける全ての非ゼロエントリの値を保持し、
・列配列A’rは、列優先順序でAにおける全ての非ゼロエントリの行を保持し、
・行インデックス配列A’cは、Aにおける各列の第1の非ゼロエントリのA’vにおけるインデックスを保持し、A’v配列におけるエントリの総数であるA’c配列における最後のエントリを有する。
sum←sum+A’v[i]×x’[A’r[i]]
が実行される(ステップ26)。計算において、インデックスiを有するA’v配列における値はx’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’r配列の数である。乗算結果は、ステップ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’c[j] /*i:j列目における第1の非ゼロエントリのA’vにおけるインデックス*/
imax←A’c[j+1] /*imax:(j+1)列目における第1の非ゼロエントリのA’vにおけるインデックス*/
sum←0 /*総和積算器を初期化*/
while (i<imax) /*列の終わりに到達したかどうかのテスト/*
sum←sum+A’v[i]×x’[A’r[i]] /*y’j=ΣiAi,j×x’iを計算*/
i←i+1 /*j列目における次の非ゼロの行に移動*/
while文終了
y’[j]←sum /*y’に結果を格納*/
for文終了
・As[0]=1
・k=0,1,2,・・・,p−1について、As[k]<As[k+1]
・As[p]=∞
・∀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)
1.第1のパーティションが単一の非ゼロエントリを有する行のみを含み、
2.第2のパーティションが複数の非ゼロエントリを有する行を含む
ように、Aについて2つのパーティションを作成したい場合、Ap=A’p =2及びAs=A’s=[1,2,∞]である。他の配列は、以下のとおりである。
・PCSR及びPSCSRについて、Am=[0,3,1,2]及びAo=[0,2,4]
・PCSC及びPSCSCについて、A’m=[1,2,0,3]及びA’o=[0,2,4]
2.行又は列を処理するためにワープと称されるユニット内の全てのスレッドを割り当てる1ワープ1行(1W1R)カーネル
それ自体では、いずれのカーネルも、行列32の全ての部分について理想的ではない。上述したように、総ランタイムは、最大数の非ゼロエントリを有する行(又は列)に対応する最も遅いスレッドに依存することから、(以下ではf1T1Rカーネルとも称されることができる)1T1Rカーネルは負荷不均衡を被る。行(又は列)における非ゼロエントリ数が32(又は64)未満のワープ(又はウェーブフロント)サイズである場合、(以下ではf1W1Rカーネルとも称されることができる)1W1Rカーネルは、ハードウェアリソースの浪費を被る。単一の疎行列は、少数の行及び多くの非ゼロエントリを有する行の双方を有することができることに留意されたい。それゆえに、いずれのタイプのカーネルにコミットすることも、めったに問題を解決しない。さらにまた、実験は、同じ行に対して動作するように単一のワープにおけるものよりも多くのスレッドを有する利益があることを示す。これは、第3のタイプを追加する必要性を引き起こす。
3.行又は列を処理するためにスレッドのブロック(32又は64より多い)を割り当てる1ブロック1行(1B1R)カーネル
1W1Rカーネルと同様に、(以下ではf1B1Rカーネルとも称されることができる)1B1Rカーネルはまた、スレッドの数が行列のその部分における非ゼロエントリの数よりも大きい場合にハードウェアリソースの浪費を被ることがある。しかしながら、行列の特定のパーティションと各カーネルを相関させることにより、計算モジュールは、さらに図15を参照して以下に記載されるように、特定の行又は列を処理するための最良のカーネルを使用することができる。
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←Ar[i] /*j:i行目における第1の非ゼロエントリのAvcにおけるインデックス*/
jmax←Ar[i+1] /*jmax:(i+1)行目における第1の非ゼロエントリのAvcにおけるインデックス*/
sum←0 /*総和積算器を初期化*/
while (j<jmax) /*行の終わりに到達したかどうかをテスト*/
while文終了
y[i]←sum /*yに結果を格納*/
for文終了
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’c[j] /*i:j列目における第1の非ゼロエントリのA’vrにおけるインデックス*/
imax←A’c[j+1] /*imax:(j+1)列目における第1の非ゼロエントリのA’vrにおけるインデックス*/
sum←0 /*総和積算器を初期化*/
while (i<imax) /*列の終わりに到達したかどうかをテスト*/
while文終了
y’[j]←sum /*y’に結果を格納*/
for文終了
for i=0からm―1 /*疎行列Aのm行にわたるループ*/
j←Ar[i] /*j:i行目における第1の非ゼロエントリのインデックス*/
jmax←Ar[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文終了
for j=0からn―1 /*疎行列Aのn列にわたるループ*/
i←A’c[j] /*i:j列目における第1の非ゼロエントリのインデックス*/
imax←A’c[j+1] /*imax:(j+1)列目における第1の非ゼロエントリのインデックス*/
e〜j←imax−i /*e〜j:j列目における非ゼロエントリ数*/
(e〜j>0)の場合
k←lower_bound(A’s,e〜j) /*パーティションk s.t.A’s[k]≦e〜j<A’s[k+1]をみつける*/
partition[k].push_back(j) /*パーティションkの終わりに列idjを追加(すなわち、Ck)*/
if文終了
for文終了
Am.size←0 /*Am:(順序維持置換)マッピング配列*/
for k=0からp―1 /*p個のパーティションにわたるループ*/
Ao[k]←Am.size() /*Ao[k]:全ての前のパーティションの積算サイズ<k*/
Aminsert(partition[k]) /*Amの終わりにパーティションkを挿入*/
partition[k]を削除 /*パーティションkによって使用される空きメモリ(必要に応じて)*/
for文終了
Ao[p]←Am.size() /*Ao[p]:全てのパーティションの総サイズ*/
sum←sum+Avc[j].v×x[Avc[j].c]
sum←sum+Av[j]×x[Ac[j]]
ここで、インデックスjを有するAv配列における値は、xの要素によって乗算され、そのインデックスは、j番目のインデックスを有するAc配列における数であり、乗算結果は、ステップ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←Am[r] /*i:r番目にランク付けされた行のid*/
j←Ar[i] /*j:i行目における第1の非ゼロエントリのAvcにおけるインデックス*/
jmax←Ar[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を実行するための擬似コードを使用して表現することができる。
sum←sum+A’vr[i].v×x’[A’vr[i].r]
ここで、インデックスiを有するA’vr配列の要素に格納された値は、x’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’vr要素に格納された行のインデックスであり、乗算結果は、ステップ158の前の反復中に実行された乗算結果の合計に加算される。あるいは、符号化がPCSRである場合、計算は、以下の式にしたがって行われる。
sum←sum+A’v[i]×x’[A’r[i]]
ここで、インデックスiを有するA’v配列における値は、x’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’r配列における数であり、乗算結果は、ステップ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’p―1 /*A’p個の列ベースのパーティションにわたるループ*/
for r=A’o[k]からA’o[k+1]−1 /*k番目のパーティションにおける列にわたるループ*/
j←A’m[r] /*j:r番目にランク付けされた列のid*/
i←A’c[j] /*i:j列目における第1の非ゼロエントリのA’vrにおけるインデックス*/
imax←A’c[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文終了
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,Ar,Am,Ao[k],Ao[k+1]) /*k番目のパーティションについてのk番目のカーネルを起動*/
for文終了
擬似コードは、PSCSRを参照して記載されているが、他のパーティション化された圧縮表現についての擬似コードを準用して表すことができる。図15を参照して上述したもの以外の特定のパーティションを処理するための特定のカーネル関数を選択する方法もまた可能である。
関数 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
・thread_id()は、現在のスレッドについてのグローバルに固有のスレッドidを返す
・warp_id()は、現在のスレッドについてのグローバルに固有のワープidを返す
・warp_thread_id()は、現在のスレッドについての(ワープ内でのみ固有の)ローカルスレッドidを返す
・block_thread_id()は、現在のスレッドについての(ブロック内でのみ固有の)ローカルスレッドidを返す
・sync_warp_threads()は、ワープ内の全てのスレッド間で同期する
・sync_block_threads()は、ブロック内の全てのスレッド間で同期する
sum←sum+Avc[j].v×x[Avc[j].c]
sum←sum+Av[j]×x[Ac[j]]
ここで、インデックスjを有するAv配列における値は、xの要素によって乗算され、そのインデックスは、j番目のインデックスを有するAc配列における数であり、乗算結果は、ステップ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,Ar,Am,rmin,rmax) /*1T1R:1スレッド1行SpMVカーネル*/
r←rmin+thread_id() /*このスレッドに割り当てられた頂点のランクを計算*/
(r<rmax)である場合 /*任意の正確に(rmax−rmin)スレッドが作成された場合*/
i←Am[r] /*i:r番目にランク付けされた行のid*/
j←Ar[i] /*j:i行目における第1の非ゼロエントリのAvcにおけるインデックス*/
jmax←Ar[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文終了
sum←sum+A’vr[i].v×x’[A’vr[i].r]
ここで、インデックスiを有するA’vr配列の要素に格納された値は、x’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’vr要素に格納された行のインデックスであり、乗算結果は、ステップ219の前の反復中に実行された乗算結果の合計に加算される。あるいは、符号化がPCSRである場合、計算は、以下の式にしたがって行われる。
sum←sum+A’v[i]×x’[A’r[i]]
ここで、インデックスiを有するA’v配列における値は、x’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’r配列における数であり、乗算結果は、ステップ219の前の反復中に実行された乗算結果の合計に加算される。いずれかの式に基づいて計算が終了すると、iの値に1が加算され、加算結果がiと設定され(ステップ220)、その列における次の行のエントリに処理を移動する。ルーチン210は、上述したステップ218に戻り、j列目における非ゼロ値が処理されるまでステップ218〜220を繰り返す。iがimax以上である場合(ステップ218)、ループ218〜220における反復中に乗算結果を加算した合計は、密ベクトルy’に格納される(ステップ221)。スレッドの実行が停止される(ステップ222)。ルーチン210は、全ての起動した処理スレッドの実行が停止されると終了する。ルーチン210についての擬似コードは、準用する図17を参照して上記に示された擬似コードと同様に記載することができる。
sum←sum+Avc[j].v×x[Avc[j].c]
sum←sum+Av[j]×x[Ac[j]]
ここで、インデックスjを有するAv配列における値は、xの要素によって乗算され、そのインデックスは、j番目のインデックスを有するAc配列における数であり、乗算結果は、ステップ241の前の反復中に実行された乗算結果の合計に加算される。いずれかの式に基づいて計算が終了すると、jの値にワープサイズが加算され、加算結果がjと設定され(ステップ242)、スレッドが割り当てられたその行における次の列のエントリに処理を移動する。ルーチン230は、上述したステップ240に戻り、そのスレッドが割り当てられるi行目における非ゼロ値が処理されるまでステップ240〜242を繰り返す。jがjmax以上である場合(ステップ240)、ワープにおける全てのスレッドは、必要に応じて同期され、ワープにおける全てのスレッドがステップ240〜242のループの実行を終了するのを可能とし、同期は、並列に動作するワープにおけるスレッドに起因して1つの実施形態では必要なく、同期から出ることがない一方で、さらなる実施形態においては同期が行われる。
カーネル f1W1R(y,x,Avc,Ar,Am,rmin,rmax) /*1W1R:1ワープ1行SpMVカーネル*/
共有 sum[WARPSIZE] /*sum:ワープにおけるスレッドによって共有されるローカル合計*/
r←rmin+warp_id() /*このワープに割り当てられた頂点のランクの計算*/
(r<rmax)である場合 /*任意の正確に(rmax−rmin)ワープが作成された場合*/
i←Am[r] /*i:r番目にランク付けされた行のid*/
t←warp_thread_id() /*t:ワープにおけるローカルスレッドid*/
j←Ar[i]+t /*j:このスレッドに割り当てられた第1の非ゼロエントリのインデックス*/
jmax←Ar[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() /*任意にワープにおけるスレッドが常に同期している場合*/
tmax←WARPSIZE/2 /*tmax:ローカル合計まで追加したスレッド数*/
while(t<tmax) /*このスレッドが参加すべきであるかどうかのテスト*/
sum[t]←sum[t]+sum[t+tmax] /*2つのローカル合計を1つに削減*/
tmax←tmax/2 /*合計追加スレッド数を半分にカット*/
sync_warp_threads() /*任意にワープにおけるスレッドが常に同期している場合*/
while文終了
(t=0及びr<rmax)である場合 /*これはワープにおける第1のスレッド?*/
y[i]←sum[0] /*yに総合計を格納*/
if文終了
sum←sum+A’vr[i].v×x’[A’vr[i].r]
ここで、インデックスiを有するA’vr配列の要素に格納された値は、x’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’vr要素に格納された行のインデックスであり、乗算結果は、ステップ271の前の反復中に実行された乗算結果の合計に加算される。あるいは、符号化がPCSCである場合、計算は、以下の式にしたがって行われる。
sum←sum+A’v[i]×x’[A’r[i]]
ここで、インデックスiを有するA’v配列における値は、x’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’r配列における数であり、乗算結果は、ステップ271の前の反復中に実行された乗算結果の合計に加算される。いずれかの式に基づいて計算が終了すると、値iにワープサイズが加算され、加算結果がiと設定され(ステップ272)、そのスレッドに割り当てられたその列における次のエントリに処理を移動する。ルーチン260は、上述したステップ270に戻り、j列目における非ゼロ値の全てが処理されるまでステップ270〜272を繰り返す。iがimax以上である場合(ステップ270)、ワープにおける全てのスレッドは、必要に応じて同期し、ワープにおける全てのスレッドがステップ270〜272のループの実行を終了するのを可能とし、同期は、並列に動作するワープにおけるスレッドに起因して1つの実施形態では必要なく、同期から出ることがない一方で、さらなる実施形態においては同期が行われる。
sum←sum+Avc[j].v×x[Avc[j].c]
sum←sum+Av[j]×x[v[j]]
ここで、インデックスjを有するAv配列における値は、xの要素によって乗算され、そのインデックスは、j番目のインデックスを有するAc配列における数であり、乗算結果は、ステップ301の前の反復中に実行された乗算結果の合計に加算される。いずれかの式に基づいて計算が終了すると、jの値にブロックサイズが加算され、加算結果がjと設定され(ステップ302)、そのスレッドに割り当てられたその行における次のエントリに処理を移動する。ルーチン290は、上述したステップ300に戻り、そのスレッドに割り当てられるi行目における非ゼロ値が処理されるまでステップ300〜302を繰り返す。jがjmax以上である場合(ステップ300)、ブロックにおける全てのスレッドは同期し、ブロックにおける全てのスレッドがステップ300〜302のループの実行を終了するのを可能とする。
カーネル f1B1R(y,x,Avc,Ar,Am,rmin,rmax) /*1B1R:1ブロック1行SpMVカーネル*/
共有 sum[BLOCKSIZE] /*sum:ブロックにおけるスレッドによって共有されるローカル合計*/
r←rmin+block_id() /*このブロックに割り当てられた頂点のランクの計算*/
(r<rmax)である場合 /*任意の正確に(rmax−rmin)ブロックが作成された場合*/
i←Am[r] /*i:r番目にランク付けされた行のid*/
t←block_thread_id() /*t:ブロックにおけるローカルスレッドid*/
j←Ar[i]+t /*j:このスレッドに割り当てられた第1の非ゼロエントリのインデックス*/
jmax←Ar[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() /*ブロックにおける全てのスレッドが同期している*/
tmax←BLOCKSIZE/2 /*tmax:ローカル合計まで追加したスレッド数*/
while(t<tmax) /*このスレッドが参加すべきであるかどうかのテスト*/
sum[t]←sum[t]+sum[t+tmax] /*2つのローカル合計を1つに削減*/
tmax←tmax/2 /*合計追加スレッド数を半分にカット*/
sync_block_threads() /*ブロックにおける全てのスレッドが同期している*/
while文終了
(t=0及びr<rmax)である場合 /*これはブロックにおける第1のスレッド?*/
y[i]←sum[0] /*yに総合計を格納*/
if文終了
sum←sum+A’vr[i].v×x’[A’vr[i].r]
ここで、インデックスiを有するA’vr配列の要素に格納された値は、x’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’vr要素に格納された行のインデックスであり、乗算結果は、ステップ331の前の反復中に実行された乗算結果の合計に加算される。あるいは、符号化がPCSRである場合、計算は、以下の式にしたがって行われる。
sum←sum+A’v[i]×x’[A’r[i]]
ここで、インデックスiを有するA’v配列における値は、x’の要素によって乗算され、そのインデックスは、i番目のインデックスを有するA’r配列における数であり、乗算結果は、ステップ331の前の反復中に実行された乗算結果の合計に加算される。いずれかの式に基づいて計算が終了すると、iの値にブロックサイズが加算され、加算結果がiと設定され(ステップ332)、その列における次のエントリに処理を移動する。ルーチン320は、上述したステップ330に戻り、j列目における非ゼロ値が処理されるまでステップ330〜332を繰り返す。iがimax以上である場合(ステップ330)、ブロックにおける全てのスレッドは同期し、ブロックにおける全てのスレッドがステップ330〜332のループの実行を終了するのを可能とする。
それ以外
ここで、Ljは、ページjから出るリンクの数である。x及びyをサイズnの2つの密ベクトルとし、ε*を停止閾値とする。初期のページランクの確率分布がベクトルxについて設定される(ステップ351)。εの値がε*の値未満である場合に反復処理ループ(ステップ352〜356)が実行される(ステップ352)。値yは、以下の式に基づいて設定される。
x←x0 /*初期のページランクの確率分布*/
ループ
x←y
ε<ε*まで
Claims (18)
- 構造化疎行列表現の取得及び利用のためのコンピュータ実装方法において、
各部分が行及び列のうちの1つを含む疎行列の部分における1つ以上の順序で配置された1つ以上の非ゼロエントリを含む前記疎行列の構造化圧縮表現を、1つ以上のキャッシュ及び1つ以上のプロセッサを備えた少なくとも一つのサーバによって取得するステップを含み、
前記取得するステップは、
前記構造化疎行列表現に含まれる複合配列であって、前記非ゼロエントリと前記非ゼロエントリを含む部分の一つにおけるインデックスとの一つをそれぞれが有する1つ以上の要素を含む複合配列を取得するステップと、
前記順序の1つ以上における最初である前記非ゼロエントリを含む前記要素それぞれについての前記複合配列におけるインデックスを有するインデックス配列であって、前記行列における前記非ゼロエントリの数をさらに有するインデックス配列を取得するステップとを有し、
方法は、さらに、前記プロセッサのそれぞれが前記1つ以上のキャッシュのうちの一つにおいて前記複合配列の前記1つ以上の要素にアクセスすることを含む疎行列解法(SLA)を前記構造化圧縮表現上で前記少なくとも一つのサーバによって実行するステップであって、前記要素における前記非ゼロエントリと前記要素における前記インデックスとを取得する、ステップを含み、
前記複合配列の前記1つ以上の要素にアクセスする間に前記キャッシュから前記部分のうちの一つにおける前記非ゼロエントリと前記インデックスとを取得するキャッシュミスの数が削減され、
各複合配列の要素インデックスが、列優先順序でその非ゼロエントリの行のインデックスを含み、前記順序が各列における前記非ゼロエントリの順序を含み、
方法は、さらに、前記行列の各列を処理することによって複数の要素を含む密ベクトルによって前記行列の疎行列転置ベクトル乗算を実行するステップと、
前記インデックス配列を使用してその列における前記非ゼロエントリを含む全ての前記複合配列の要素の前記複合配列におけるインデックスを識別するステップと、
前記非ゼロエントリの前記インデックスを使用して前記密ベクトルの要素のいずれかによってその列における前記非ゼロエントリのそれぞれを乗算するステップと、
その列における前記非ゼロエントリのそれぞれについての乗算結果を加算し、異なる密ベクトルの加算結果を格納するステップとを含む、方法。 - 各複合配列の要素インデックスが、行優先順序でその非ゼロエントリの列のインデックスを含み、前記順序が各行における前記非ゼロエントリの順序を含む、請求項1に記載の方法。
- さらに、
前記行列の各行を処理することによって複数の要素を含む密ベクトルによって前記行列の疎行列ベクトルの乗算を実行するステップと、
前記インデックス配列を使用してその行における前記非ゼロエントリを含む全ての前記複合配列要素の前記複合配列におけるインデックスを識別するステップと、
前記非ゼロエントリの前記インデックスを使用して前記密ベクトルの要素のいずれかによってその行における前記非ゼロエントリのそれぞれを乗算するステップと、
その行における前記非ゼロエントリのそれぞれについての乗算結果を加算し、異なる密ベクトルの加算結果を格納するステップとを備える、請求項2に記載の方法。 - さらに、
前記構造化圧縮表現を処理することと、
前記処理の結果を使用してランク付け分析を実行することとを備える、請求項1に記載の方法。 - 効率的な疎行列表現及び使用のためのコンピュータ実装方法において、
疎行列の圧縮表現をそれぞれが1つ以上のグラフィックス処理ユニット(GPU)のプロセッサコアを備えた1つ以上のストリーミングプロセッサを備える少なくとも一つのサーバによって取得するステップであって、前記疎行列が前記疎行列の複数の部分の1つ以上において1つ以上の非ゼロエントリを含み、前記疎行列の部分が前記疎行列におけるそれらの位置に基づいてインデックス付けされ、前記部分が前記疎行列の行及び列の一つを含む、ステップと、
前記疎行列の部分について複数のパーティションを前記少なくとも一つのサーバによって定義するステップと、
前記圧縮表現を使用して前記部分のそれぞれにおける前記非ゼロエントリの数を前記少なくとも一つのサーバによって取得するステップと、
その部分における前記非ゼロエントリの数に基づいて前記パーティションのいずれかを有する部分のそれぞれを前記少なくとも一つのサーバによって関連付けるステップであって、前記パーティションのそれぞれに関連する部分が、残りのパーティションに関連する範囲と異なる範囲内にある前記非ゼロエントリの数を有する、ステップと、
前記パーティションのそれぞれに関連する全ての部分であり、それらのインデックスの順序でリスト化された前記部分のリストを前記少なくとも一つのサーバによって作成するステップと、
前記リストを含むマッピング配列を含む前記疎行列のパーティション化された圧縮表現を前記少なくとも一つのサーバによって作成するステップと、
前記パーティションのそれぞれに関連する前記部分の前記パーティションに関連する前記範囲に基づく処理のための前記ストリーミングプロセッサの1つ以上における所定の数の前記GPUのプロセッサコアを前記少なくとも一つのサーバによって割り当てるステップと、
前記パーティションのそれぞれに関連する前記部分のそれぞれを、割り当てられた
前記所定の数の前記GPUのプロセッサコアによって処理するステップとを含む、方法。 - 前記範囲の一つが1から31の前記非ゼロエントリを含み、前記範囲の第2の一つが31から1024の前記非ゼロエントリを含み、前記範囲の第3の一つが1024以上の前記非ゼロエントリの数を含む、請求項5に記載の方法。
- さらに、
前記パーティションのそれぞれに関連する前記疎行列の部分における非ゼロエントリの数に基づいて前記パーティションをインデックス付けるステップと、
前記パーティションのそれぞれについて、前記パーティションのインデックスに基づいて前記マッピング配列におけるそのパーティションに先行する全てのパーティションのサイズであり、前記パーティションに関連する複数の部分の数を含む前記サイズを識別するステップであって、そのパーティションについての前記リストを前記サイズに基づく位置において前記マッピング配列に挿入するステップとを含む、請求項5に記載の方法。 - 前記疎行列の列の一つを含む前記各部分が、さらに、
前記マッピング配列にリスト化されたパーティションのそれぞれの列のそれぞれを処理することによって複数の要素を含む密ベクトルによって前記疎行列の疎行列ベクトル転置乗算を実行するステップと、
その列における前記非ゼロエントリのそれぞれに関連付けられた前記圧縮表現におけるインデックスを識別するステップと、
その非ゼロエントリの前記インデックスを使用して前記密ベクトルの要素のいずれかによってその列における前記非ゼロエントリのそれぞれを乗算するステップと、
その列における前記非ゼロエントリのそれぞれについての乗算結果を加算し、異なる密ベクトルにおける加算結果を格納するステップとを備える、請求項7に記載の方法。 - 前記疎行列の行の一つを含む前記各部分が、さらに、
前記マッピング配列にリスト化されたパーティションのそれぞれの行のそれぞれをシーケンシャルに処理することによって複数の要素を含む密ベクトルによって前記疎行列の疎行列ベクトル乗算を実行するステップと、
その行における前記非ゼロエントリのそれぞれに関連付けられた前記圧縮表現におけるインデックスを識別するステップと、
その非ゼロエントリの前記インデックスを使用して前記密ベクトルの要素のいずれかによってその行における前記非ゼロエントリのそれぞれを乗算するステップと、
その行における前記非ゼロエントリのそれぞれについての乗算結果を加算し、異なる密ベクトルにおける加算結果を格納するステップとを備える、請求項5に記載の方法。 - 前記パーティション化された圧縮表現は、さらに、前記疎行列の前記部分のそれぞれにおける前記非ゼロエントリの最初の一つに関連するインデックスを含み、前記部分のそれぞれの前記処理は、割り当てられた前記GPUのプロセッサコアのそれぞれによって、前記サーバに含まれる1つ以上のキャッシュの一つにシーケンシャルにアクセスするステップを含み、前記インデックスの配列における複数のインデックスは、前記マッピング配列の前記リストにおいて、前記インデックスの配列における複数のインデックスに関連する前記部分の順序でリスト化されており、シーケンシャルに行われる前記アクセスは、前記複数のインデックスへの非シーケンシャルなアクセスと比較されるときに、各キャッシュにおける前記インデックスの配列へのアクセスを削減することによる前記処理の間に、キャッシュミスの数を削減する、請求項5に記載の方法。
- さらに、
前記パーティションのそれぞれに関連する前記部分を処理するための複数のカーネル関数から一つを選択するステップを含み、
前記カーネル関数が、前記部分の一つを処理する前記GPUの1つのプロセッサコアによって実行される一つの処理スレッドを割り当てるカーネル関数(f1T1Rカーネル関数)と、前記部分の一つを処理する前記GPUの1つより多くのプロセッサコアによって実行される処理スレッドのワープを割り当てるカーネル関数(f1W1Rカーネル関数)と、前記処理スレッドの前記ワープを実行する前記GPUのプロセッサコアの数より多い前記所定の数の前記GPUのプロセッサコアであって、前記部分の一つを処理するプロセッサコアによって実行される処理スレッドのブロックを割り当てるカーネル関数(f1B1Rカーネル関数)とのうちの1つ以上を含み、前記f1W1Rカーネル関数が選択される前記部分のそれぞれは、前記f1T1Rカーネル関数が選択される前記部分のそれぞれのものより多い前記非ゼロエントリを含み、前記f1B1Rカーネル関数が選択される前記部分のそれぞれは、前記f1W1Rカーネル関数が選択される前記部分のそれぞれのものより多い前記非ゼロエントリを含む、請求項5に記載の方法。 - 前記部分のそれぞれは、前記行の一つを含み、方法は、さらに、
複数の要素を含む密ベクトルによる前記疎行列における前記行の一つの疎行列ベクトル乗算を前記f1T1Rカーネル関数によって割り当てられる処理スレッドによって実行するステップを含み、前記実行するステップは、
前記マッピング配列を用いて、前記処理スレッドに割り当てられた前記行を識別するステップと、
前記行における前記非ゼロエントリのそれぞれに関連する前記圧縮表現においてインデックスを識別するステップと、
前記行における前記非ゼロエントリのそれぞれを、前記非ゼロエントリの前記インデックスを用いて前記密ベクトルの前記要素の一つと乗算するステップと、
前記行における前記非ゼロエントリのそれぞれについての乗算結果を加算し、異なる密ベクトルの加算結果を格納するステップとを有する、請求項11に記載の方法。 - 前記部分のそれぞれは、前記列のうちの一つを含み、前記方法は、さらに、
複数の要素を含む密ベクトルによる前記疎行列の前記列のうちの一つの疎行列ベクトル転置乗算を、前記f1T1Rカーネル関数により割り当てられた処理スレッドによって実行するステップと、
前記処理スレッドに割り当てられた前記列を前記マッピング配列を用いて識別するステップと、
前記列において、前記非ゼロエントリのそれぞれに関連する前記圧縮表現でインデックスを識別するステップと、
前記非ゼロエントリの前記インデックスを用いて、前記列における前記非ゼロエントリのそれぞれを前記密ベクトルの前記要素のうちの一つと乗算するステップと、
前記列における前記非ゼロエントリのそれぞれについて、前記乗算の結果を加算し、異なる密ベクトルの加算結果を格納するステップと、
を含む、請求項11に記載の方法。 - 前記部分のそれぞれは、前記行のうちの一つを含み、前記方法は、さらに、
前記疎行列の前記行のうちの一つの疎行列ベクトル乗算を複数の要素を含む密ベクトルにより、前記f1W1Rカーネル関数により割り当てられた処理スレッドのワープによって実行するステップと、
前記ワープに割り当てられた前記行を前記マッピング配列を用いて識別するステップと、
前記ワープにおける前記処理スレッドのそれぞれに割り当てられた前記行における前記非ゼロエントリの前記圧縮表現でインデックスを識別し、前記非ゼロエントリの前記インデックスを用いて、前記処理スレッドに割り当てられた前記非ゼロエントリのそれぞれを前記密ベクトルの前記要素のうちの一つと乗算するステップと、
前記処理スレッドに割り当てられた前記非ゼロエントリのそれぞれについて、前記乗算結果を合計するステップと、
前記ワープにおける前記処理スレッドの全てについて前記合計を組み合わせて、組み合わせた合計を異なる密ベクトルで格納するステップと、
を含む、請求項11に記載の方法。 - 前記部分のそれぞれは、前記列のうちの一つを含み、前記方法は、さらに、
複数の要素を含む密ベクトルによる前記疎行列の前記列のうちの一つの疎行列ベクトル転置乗算を、前記f1W1Rカーネル関数により割り当てられた処理スレッドのワープによって実行するステップと、
前記ワープに割り当てられた前記列を前記マッピング配列を用いて識別するステップと、
前記ワープにおける前記処理スレッドのそれぞれに割り当てられた前記列における前記非ゼロエントリの前記圧縮表現でインデックスを識別し、前記非ゼロエントリの前記インデックスを用いて、前記処理スレッドに割り当てられた前記非ゼロエントリのそれぞれを前記密ベクトルの前記要素のうちの一つと乗算するステップと、
前記処理スレッドに割り当てられた前記非ゼロエントリのそれぞれについて、前記乗算結果を合計するステップと、
前記ワープにおける前記処理スレッドの全てについて前記合計を組み合わせて、組み合わせた合計を異なる密ベクトルで格納するステップと、
を含む、請求項11に記載の方法。 - 前記部分のそれぞれは、前記行のうちの一つを含み、前記方法は、さらに、
前記疎行列の前記行のうちの一つの疎行列ベクトル乗算を複数の要素を含む密ベクトルにより、前記f1B1Rカーネル関数により割り当てられた処理スレッドのブロックによって実行するステップと、
前記ブロックに割り当てられた前記行を前記マッピング配列を用いて識別するステップと、
前記ブロックにおける前記処理スレッドのそれぞれに割り当てられた前記行における前記非ゼロエントリの前記圧縮表現でインデックスを識別し、前記非ゼロエントリの前記インデックスを用いて、前記処理スレッドに割り当てられた前記非ゼロエントリのそれぞれを前記密ベクトルの前記要素のうちの一つと乗算するステップと、
前記処理スレッドに割り当てられた前記非ゼロエントリのそれぞれについて、前記乗算結果を合計するステップと、
前記合計の完了に応じて、前記ブロックにおける前記処理スレッドを同期するステップと、
前記ブロックのうちの同期された処理スレッド全てについて前記合計を組み合わせて、組み合わせた合計を異なる密ベクトルで格納するステップと、
を含む、請求項11に記載の方法。 - 前記部分のそれぞれは、前記列のうちの一つを含み、前記方法は、さらに、
前記列のうちの一つの疎行列ベクトル乗算を複数の要素を含む密ベクトルにより、前記f1B1Rカーネル関数により割り当てられた処理スレッドのブロックによって実行するステップと、
前記ブロックに割り当てられた前記列を前記マッピング配列を用いて識別するステップと、
前記ブロックにおける前記処理スレッドのそれぞれに割り当てられた前記列における前記非ゼロエントリの前記圧縮表現でインデックスを識別し、前記非ゼロエントリの前記インデックスを用いて、前記処理スレッドに割り当てられた前記非ゼロエントリのそれぞれを前記密ベクトルの前記要素のうちの一つと乗算するステップと、
前記処理スレッドに割り当てられた前記非ゼロエントリのそれぞれについて、前記乗算結果を合計するステップと、
前記合計の完了に応じて、前記ブロックにおける前記処理スレッドを同期するステップと、
前記ブロックのうちの同期された処理スレッド全てについて前記合計を組み合わせて、組み合わせた合計を異なる密ベクトルで格納するステップと、
を含む、請求項11に記載の方法。 - 前記圧縮表現は、圧縮疎行符号化、圧縮疎列符号化、構造化圧縮疎行符号化及び構造化圧縮疎列符号化の一つを含み、方法は、さらに、
式δ=γ/(λb+1)を用いて、パーティション化された前記圧縮表現により要求される過剰な空間的オーバーヘッドを予測するステップを含み、
δが前記過剰な空間的オーバーヘッドであり、γがパーティション化された前記圧縮表現における前記部分の合計数に対する前記非ゼロエントリを含む前記部分の数の割合であり、bがパーティション化された前記圧縮表現における行(又は列)の数に対する前記非ゼロエントリの数の割合であり、λがパーティション化された前記圧縮表現のインデックス配列における要素のサイズに対する複合配列における要素のサイズの割合である、請求項5に記載の方法。
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 JP2016119084A (ja) | 2016-06-30 |
JP2016119084A5 JP2016119084A5 (ja) | 2019-01-24 |
JP6630558B2 true 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) |
Families Citing this family (45)
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 |
JP2018116561A (ja) * | 2017-01-19 | 2018-07-26 | 弘崇 新妻 | Delayed Sparse Matrix |
US10489480B2 (en) | 2017-01-22 | 2019-11-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 |
US10902318B2 (en) | 2017-11-06 | 2021-01-26 | 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的多维健康数据的统计分析方法及系统 |
US12112175B1 (en) | 2018-02-08 | 2024-10-08 | Marvell Asia Pte Ltd | Method and apparatus for performing machine learning operations in parallel on machine learning hardware |
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 |
US10970080B2 (en) | 2018-02-08 | 2021-04-06 | Marvell Asia Pte, Ltd. | Systems and methods for programmable hardware architecture for machine learning |
US10691772B2 (en) * | 2018-04-20 | 2020-06-23 | Advanced Micro Devices, Inc. | High-performance sparse triangular solve on graphics processing units |
US10929779B1 (en) | 2018-05-22 | 2021-02-23 | Marvell Asia Pte, Ltd. | Architecture to support synchronization between core and inference engine for machine learning |
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 |
US11449363B2 (en) | 2018-05-31 | 2022-09-20 | Neuralmagic Inc. | Systems and methods for improved neural network execution |
WO2021061172A1 (en) * | 2019-09-27 | 2021-04-01 | 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 |
US10832133B2 (en) | 2018-05-31 | 2020-11-10 | Neuralmagic Inc. | System and method of executing neural networks |
US11216732B2 (en) | 2018-05-31 | 2022-01-04 | Neuralmagic Inc. | Systems and methods for generation of sparse code for convolutional neural networks |
WO2020029018A1 (zh) | 2018-08-06 | 2020-02-13 | 华为技术有限公司 | 矩阵的处理方法、装置及逻辑电路 |
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 |
US11221848B2 (en) * | 2019-09-25 | 2022-01-11 | Intel Corporation | Sharing register file usage between fused processing resources |
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 |
US20240143199A1 (en) * | 2022-11-01 | 2024-05-02 | Advanced Micro Devices, Inc. | Sparse Matrix Operations Using Processing-in-Memory |
Family Cites Families (9)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
US7539681B2 (en) * | 2004-07-26 | 2009-05-26 | Sourcefire, Inc. | Methods and systems for multi-pattern searching |
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 |
WO2011156247A2 (en) | 2010-06-11 | 2011-12-15 | Massachusetts Institute Of Technology | Processor for large graph algorithm computations and matrix operations |
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 |
-
2014
- 2014-12-22 US US14/580,110 patent/US9760538B2/en active Active
-
2015
- 2015-12-10 JP JP2015240835A patent/JP6630558B2/ja active Active
- 2015-12-14 EP EP15199963.8A patent/EP3037980A3/en not_active Withdrawn
-
2017
- 2017-09-07 US US15/698,547 patent/US10296556B2/en active Active
Also Published As
Publication number | Publication date |
---|---|
US20160179750A1 (en) | 2016-06-23 |
JP2016119084A (ja) | 2016-06-30 |
US10296556B2 (en) | 2019-05-21 |
US9760538B2 (en) | 2017-09-12 |
US20170371839A1 (en) | 2017-12-28 |
EP3037980A3 (en) | 2016-07-06 |
EP3037980A2 (en) | 2016-06-29 |
Similar Documents
Publication | Publication Date | Title |
---|---|---|
JP6630558B2 (ja) | 効率的な疎行列表現及び処理のためのコンピュータ実装システム及び方法 | |
Ashari et al. | Fast sparse matrix-vector multiplication on GPUs for graph applications | |
Kronbichler et al. | Multigrid for matrix-free high-order finite element computations on graphics processors | |
Mittal et al. | A survey of deep learning on cpus: opportunities and co-optimizations | |
Benson et al. | Direct QR factorizations for tall-and-skinny matrices in MapReduce architectures | |
Li et al. | GPU-accelerated preconditioned iterative linear solvers | |
Choi et al. | Blocking optimization techniques for sparse tensor computation | |
Morozov et al. | Block-parallel data analysis with DIY2 | |
Li et al. | Walking in the cloud: Parallel simrank at scale | |
Ma et al. | Optimizing sparse tensor times matrix on GPUs | |
Canny et al. | Machine learning at the limit | |
Lee et al. | Fast matrix-vector multiplications for large-scale logistic regression on shared-memory systems | |
Vannieuwenhoven et al. | Computing the gradient in optimization algorithms for the CP decomposition in constant memory through tensor blocking | |
Angiulli et al. | GPU strategies for distance-based outlier detection | |
Zheng et al. | Semi-external memory sparse matrix multiplication for billion-node graphs | |
Roy et al. | Numa-caffe: Numa-aware deep learning neural networks | |
Martínez-del-Amor et al. | Population Dynamics P systems on CUDA | |
Burchard et al. | ipug: Accelerating breadth-first graph traversals using manycore graphcore ipus | |
Wang et al. | A survey of statistical methods and computing for big data | |
Sivkov et al. | DBCSR: A blocked sparse tensor algebra library | |
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 | |
Pan et al. | G-slide: A gpu-based sub-linear deep learning engine via lsh sparsification | |
Ranbirsingh et al. | Distributed neural networks using tensorflow over multicore and many-core systems | |
Busato et al. | A dynamic approach for workload partitioning on GPU architectures |
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 |