JP2016119084A - 効率的な疎行列表現及び処理のためのコンピュータ実装システム及び方法 - Google Patents
効率的な疎行列表現及び処理のためのコンピュータ実装システム及び方法 Download PDFInfo
- 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
Links
- 239000011159 matrix material Substances 0.000 title claims abstract description 241
- 238000000034 method Methods 0.000 title claims abstract description 104
- 238000012545 processing Methods 0.000 title claims abstract description 102
- 238000005192 partition Methods 0.000 claims description 192
- 239000013598 vector Substances 0.000 claims description 34
- 238000013507 mapping Methods 0.000 claims description 29
- 239000002131 composite material Substances 0.000 claims description 16
- 230000003247 decreasing effect Effects 0.000 abstract 1
- 238000004364 calculation method Methods 0.000 description 77
- 230000006870 function Effects 0.000 description 66
- 238000010586 diagram Methods 0.000 description 38
- 230000008569 process Effects 0.000 description 37
- 238000004422 calculation algorithm Methods 0.000 description 26
- 238000003491 array Methods 0.000 description 15
- 238000012360 testing method Methods 0.000 description 15
- 230000009467 reduction Effects 0.000 description 14
- 238000000638 solvent extraction Methods 0.000 description 12
- 230000001360 synchronised effect Effects 0.000 description 12
- 238000002474 experimental method Methods 0.000 description 10
- 230000014509 gene expression Effects 0.000 description 10
- 230000008901 benefit Effects 0.000 description 7
- 230000001133 acceleration Effects 0.000 description 4
- 230000006835 compression Effects 0.000 description 3
- 238000007906 compression Methods 0.000 description 3
- 230000001419 dependent effect Effects 0.000 description 3
- 230000001174 ascending effect Effects 0.000 description 2
- 230000009286 beneficial effect Effects 0.000 description 2
- 230000006872 improvement Effects 0.000 description 2
- 238000012423 maintenance Methods 0.000 description 2
- 230000004044 response Effects 0.000 description 2
- 235000019227 E-number Nutrition 0.000 description 1
- 239000004243 E-number Substances 0.000 description 1
- 230000004913 activation Effects 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
- 230000001186 cumulative effect Effects 0.000 description 1
- 230000007423 decrease Effects 0.000 description 1
- 238000006731 degradation reaction Methods 0.000 description 1
- 230000000694 effects Effects 0.000 description 1
- 239000000203 mixture Substances 0.000 description 1
- 238000012986 modification Methods 0.000 description 1
- 230000004048 modification Effects 0.000 description 1
- 238000006467 substitution reaction Methods 0.000 description 1
- 238000012546 transfer Methods 0.000 description 1
- 230000007704 transition Effects 0.000 description 1
- 239000002699 waste material Substances 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)
Abstract
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 (10)
- 構造化疎行列表現を取得するためのコンピュータ実装方法において、
各部分が行及び列のうちの1つを含む行列の部分における1つ以上の順序で配置された1つ以上の非ゼロエントリを含む前記行列の構造化圧縮表現を取得することと、
前記表現に含まれる複合配列であり、各要素が非ゼロエントリのいずれかとその非ゼロエントリを含む前記部分のいずれかのインデックスとを含む1つ以上の要素を含む前記複合配列を取得することと、
前記順序のうちの1つ以上における第1である前記非ゼロエントリを含む各要素の前記複合配列におけるインデックスを含むインデックス配列であり、さらに前記行列における複数の前記非ゼロエントリを含む前記インデックス配列を取得することとを備える、方法。 - 各複合配列の要素インデックスが、列優先順序でその非ゼロエントリの行のインデックスを含み、前記順序が各列における前記非ゼロエントリの順序を含む、請求項1に記載の方法。
- さらに、
各列を処理することによって複数の要素を含む密ベクトルによって前記行列の疎行列転置ベクトル乗算を実行することと、
前記インデックス配列を使用してその列における前記非ゼロエントリを含む全ての前記複合配列要素からなる前記複合配列におけるインデックスを識別することと、
そのエントリの前記インデックスを使用して前記密ベクトルの要素のいずれかによってその列における前記非ゼロエントリのそれぞれを乗算することと、
その列における前記非ゼロエントリのそれぞれについての乗算結果を加算し、異なる密ベクトルの加算結果を格納することとを備える、請求項2に記載の方法。 - 各複合配列の要素インデックスが、行優先順序でその非ゼロエントリの列のインデックスを含み、前記順序が各行における前記非ゼロエントリの順序を含む、請求項1に記載の方法。
- さらに、
各行を処理することによって複数の要素を含む密ベクトルによって前記行列の疎行列ベクトルの乗算を実行することと、
前記インデックス配列を使用してその行における前記非ゼロエントリを含む全ての前記複合配列要素からなる前記複合配列におけるインデックスを識別することと、
そのエントリの前記インデックスを使用して前記密ベクトルの要素のいずれかによってその行における前記非ゼロエントリのそれぞれを乗算することと、
その行における前記非ゼロエントリのそれぞれについての乗算結果を加算し、異なる密ベクトルの加算結果を格納することとを備える、請求項4に記載の方法。 - さらに、
前記構造化圧縮表現を処理することと、
前記処理の結果を使用してランク付け分析を実行することとを備える、請求項1に記載の方法。 - 効率的な疎行列表現のためのコンピュータ実装方法において、
1つ以上の非ゼロエントリを含む疎行列の圧縮表現を取得することであり、前記行列の行及び列のいずれかを含む前記行列の部分が前記行列におけるそれらの位置に基づいてインデックス付けされることと、
前記行列の部分について複数のパーティションを定義することと、
前記圧縮表現を使用して前記部分のそれぞれにおける複数の前記非ゼロエントリを取得することと、
その部分における複数の前記非ゼロエントリに基づいて前記パーティションのいずれかを有する部分のそれぞれを関連付けることと、
前記パーティションのそれぞれに関連する全ての部分であり、それらのインデックスの順序でリスト化された前記部分のリストを作成することと、
前記リストを含むマッピング配列を含む前記行列のパーティション化された圧縮表現を作成することとを備える、方法。 - さらに、
前記マッピング配列を処理することと、
前記処理の結果に基づいてランク付け分析を実行することとを備える、請求項7に記載の方法。 - さらに、
前記パーティションのそれぞれに関連する前記行列の部分における非ゼロエントリの数に基づいて前記パーティションをインデックス付けることと、
前記パーティションのそれぞれについて、前記パーティションのインデックスに基づいて前記配列におけるそのパーティションに先行する全てのパーティションのサイズであり、前記パーティションに関連する複数の部分の数を含む前記サイズを識別し、そのパーティションについての前記リストを前記サイズに基づく位置において前記配列に挿入することとを備える、請求項7に記載の方法。 - 前記行列の行を含む前記各部分が、さらに、
前記マッピング配列にリスト化されたパーティションのそれぞれの行のそれぞれをシーケンシャルに処理することによって複数の要素を含む密ベクトルによって前記行列の疎行列ベクトル乗算を実行することと、
その行における前記非ゼロエントリのそれぞれに関連付けられた前記圧縮表現におけるインデックスを識別することと、
その非ゼロエントリの前記インデックスを使用して前記密ベクトルの要素のいずれかによってその行における前記非ゼロエントリのそれぞれを乗算することと、
その行における前記非ゼロエントリのそれぞれについての乗算結果を加算し、異なる密ベクトルにおける加算結果を格納することとを備える、請求項7に記載の方法。
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)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
WO2018135599A3 (ja) * | 2017-01-19 | 2018-09-13 | 新妻弘崇 | 遅延疎行列 |
Families Citing this family (40)
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)
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)
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 |
-
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
Patent Citations (2)
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)
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)
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 |