Qiita Teams that are logged in
You are not logged in to any team

Log in to Qiita Team
Community
OrganizationAdvent CalendarQiitadon (β)
Service
Qiita JobsQiita ZineQiita Blog
Help us understand the problem. What is going on with this article?

cuBLAS と cuBLAS-XT の調査(その2)。行列の積演算にて。転置の影響。

More than 3 years have passed since last update.

cuBLSA を用いて GPU で 行列の積演算を行わせるにあたり、行列の積演算関数の引き数には、入力行列の転置の指定があり、転置によって計算時間にどう影響あるのか調査しました。
調査結果を公開します。

背景

深層学習の理解を深めようと自前で実装を行っていると、行列の積演算 の所要時間を少しでも短くしようと感じます。 
BLASでの計算時間は、いかにメモリアクセスの時間を短縮するかの工夫(キャッシュメモリに当てる、メモリを連続的にアクセスする)の世界になっていると想像し、入力行列の転置によっても、所要時間に影響あるかと予想し、今回測定しました。

行列の積演算関数 gemm()

参照先「cuda tool kit documentation cublasgemm()」

gemm_sample.c
cublasSgemm(    
        handle,
        CUBLAS_OP_N, //行列A 転置有無
        CUBLAS_OP_N, //行列B 転置有無
        num,    // 行列Aの行数
        num,    // 行列Bの列数
        num,    // 行列Aの列数(=行列Bの行数)
        &alpha, // 行列の積に掛ける値(なければ1)
        devA,   // 行列A
        num,    // 行列Aの行数
        devB,   // 行列B
        num,    // 行列Bの行数
        &beta,  // 行列Cに掛けるスカラ値(なければ0)
        devC,   // 行列Cの初期値 兼 出力先
        num // 行列Cの行数
);

↑の 行列A と 行列B の転置有無の引数を、

意味
CUBLAS_OP_N 転置なし
CUBLAS_OP_T 転置あり

の 値を指定することで、転置の有無を指定します。

転置による所要時間

  • 一辺 num の 正方行列(下図)で、 num を増やしていった際の 所要時間で比較する。

行列の積.PNG

cublas_Bt.png

↑を num 4096 以下 で拡大したのが↓

cublas_AtBt__u4096.png

↑を num 2048 以下 で拡大したのが↓

cublas_AtBt__u2048.png

結果

  • numの値に関わらず「Bのみ転置」が 所要時間が最も短い。
  • 「転置なし」は num が 2000 以下のあたりでは、「Bのみ転置」と同等。numが4096を超えたあたりから ~3%位の差が生じる。

(「転置なし」と「Bのみ転置」間は予想より差が少なかった。「Aのみ転置」と「転置なし」と同様の差が、「転置なし」と「Bのみ転置」間にもあるように予想してました)

→「Bのみ転置」を用いるのが良さそう。

cuBLAS-XT も含めた比較

cublas_vs_cublasXt_all.png

  • 所要時間の短い方から、「cublas(Bのみ転置)」「cublasXt(Bのみ転置)」「cublas(転置なし)」「cublasXt(転置なし)」
  • 「Bのみ転置」にしても、cuBLAS-XT より cuBLA のが わずかに所要時間が短い。(グラフィックカードが1枚の場合)

実装の話

BLASで扱う行列は Column優先(列優先)のメモリ配置。

col_row_01.PNG

画像などでよく用いる Row とは向きが異なるので注意が必要。

col優先での B転置!? メモリにどう配置するん?

Row 優先表記で↓の配置。

cor_bt_mem.PNG

特徴

  • MとPの共通項がどちらも 行数 になる。 (どっちがどっちでの悩みが一つ減る)
  • MとP、どちらをAにするかで、結果の転置を調整できる。

これはこれで便利そう。 (うまくできてるもんですね)

検算したの?

むしろ、結果から逆算しました...。

入力:
--- M ---
mem:[ 1.0000 ,2.0000 ,3.0000 ,4.0000 ,5.0000 ,6.0000 ,7.0000 ,8.0000 ,9.0000 ]
1.0000 2.0000 3.0000 
4.0000 5.0000 6.0000 
7.0000 8.0000 9.0000 

--- P ---
mem:[ 0.0001 ,0.0010 ,0.0100 ,0.1000 ,1.0000 ,10.0000 ,100.0000 ,1000.0000 ,10000.0000 ]
  0.0001    0.0010     0.0100 
  0.1000    1.0000    10.0000 
100.0000 1000.0000 10000.0000 

結果:
--- M×Pt ---
mem:[ 700.4001 ,800.5002 ,900.6003 ,7004.0010 ,8005.0020 ,9006.0030 ,70040.0100 ,80050.0200 ,90060.0300 ]
  700.4001   800.5002   900.6003 
 7004.0010  8005.0020  9006.0030 
70040.0100 80050.0200 90060.0300

--- P×Mt ---
mem:[700.4001 ,7004.0010 ,70040.0100 ,800.5002 ,8005.0020 ,80050.0200 ,900.6003 ,9006.0030 ,90060.0300 ]
700.4001 7004.0010 70040.0100 
800.5002 8005.0020 80050.0200 
900.6003 9006.0030 90060.0300 
t-tkd3a
機械学習 について実装できる位の理解を目指します。学ぶ過程の資料・成果を公開していきます。 また Linux での 開発環境・ツール類についても忘備録かねて記載していきます。
Why not register and get more from Qiita?
  1. We will deliver articles that match you
    By following users and tags, you can catch up information on technical fields that you are interested in as a whole
  2. you can read useful information later efficiently
    By "stocking" the articles you like, you can search right away