1 GPUのPTXコードを用いた
ランダムアドレスシフトの厳密評価
広島大学 藤田 徹
研究概要
GPUのアセンブリ言語であるPTXコードを用いて，メモリアクセスの高速 化手法であるランダムアドレスシフトの厳密な性能評価を実施
◦
命令をPTXコードで記述することによって，ランダムアドレスシフトの実行ク ロックサイクル数を計測GPU（Graphics Processing Unit）
◦
グラフィックス処理に適したハードウェア◦
内部に多数のコアを搭載し，並列演算能力に優れる◦
NVIDIA社が提供するGPUの統合開発環境CUDAを利用NVIDIA社のGPUには複数のストリーミングマルチプロセッサ（SM）が 搭載されている
• 複数のコアとシェアードメモリから構成
• シェアードメモリは同じSM内のコアのみアクセス可能 GPU
SM
グローバルメモリ
シェアードメモリ
コア コア
SM シェアードメモリ
SM シェアードメモリ
・・・・
…
コア コアコア コア
…
コア コアコア コア
…
コア コアGPUアーキテクチャ
各コアにはスレッドが割り当てられる
処理はスレッドを32個毎にまとめたワープ単位で実行される
今回は1つのSMのみに注目して計測を行う
シェアードメモリとバンクコンフリクト
（例）バンク数4の場合
バンクコンフリクト シェアードメモリは1ワードずつ32個のバンクによって構成されている 複数のスレッドが同じバンクの異なるアドレスに同時にアクセス
アクセスが逐次化され，スループットが低下
バンク0 バンク1 バンク2 バンク3
0 1 2 3
4 5 6 7
8 9 10 11
12 13 14 15
スレッド0 スレッド1 スレッド2 スレッド3
衝突は発生しない
同時にアクセス
シェアードメモリとバンクコンフリクト
（例）バンク数4の場合
シェアードメモリは1ワードずつ32個のバンクから構成されている 複数のスレッドが同じバンクの異なるアドレスに同時にアクセス
アクセスが逐次化され，スループットが低下
バンク0 バンク1 バンク2 バンク3
0] 1 2 3
4 5 6 7
8 9 10 11
12 13 14 15
衝突が発生
アクセスが逐次化 バンクコンフリクト
スレッド0 スレッド1 スレッド2 スレッド3
ランダムアドレスシフト
シェアードメモリに格納する要素をシフト数にしたがって行方向にサイ クリックシフト
論文名：The Random Address Shift to Reduce the Memory Access Congestion on the Discrete Memory Machine 著者：Koji Nakano, Susumu Matsumae, and Yasuaki Ito
出典：International Symposium on Computing and Networking, pp. 95–103, Dec. 2013
3 1 3 2 バンク0 バンク1 バンク2 バンク3 シフト数
0 1 2 3
4 5 6 7
8 9 10 11
12 13 14 15
バンク0 バンク1 バンク2 バンク3
1 2 3 0
7 4 5 6
9 10 11 8
14 15 12 13
アクセスの衝突回数：4 アクセスの衝突回数：2
アドレスシフトを行うことによってバンクコンフリクトが低減
スレッド0 スレッド1 スレッド2 スレッド3 スレッド0 スレッド1 スレッド2 スレッド3
2 シフト数の決定方法
Random Address Shift
各行のシフト数は独立した乱数に
よって決定されるRandom Address Permute-Shift
ランダムに置換した数列に従って
シフトする各行のシフト数はすべて異なる
3 1 3 2バンク0 バンク1 バンク2 バンク3
1 2 3 0
7 4 5 6
9 10 11 8
14 15 12 13
シフト数
2 1 3 0
バンク0 バンク1 バンク2 バンク3
2 3 0 1
7 4 5 6
9 10 11 8
12 13 14 15
シフト数
アクセス時間の計測
GPU：GeForce GTX 680
◦
1SM当たりのコア数：192◦
動作クロック： 1006MHz次の４つのアクセスパターンについてワープ数を1から32まで変化させて 1000回計測を行い，平均時間を算出
◦
Contiguous◦
Stride◦
Diagonal◦
Randomアクセスに必要な実行クロックサイクル数を計測することで 厳密な性能評価を実施
asm volatile (
"mov.u32%%a0, %3;¥n¥t"
"mov.u32 %%a1, sh_mem;¥n¥t"
"bar.sync0;¥n¥t"
"mov.u32%0, %%clock;¥n¥t"
"shl.b32%%a2, %%a0, 3;¥n¥t"
"add.u32%%a2, %%a1, %%a2;¥n¥t"
"ld.volatile.shared.u64%2, [%%a2];¥n¥t"
"mov.u32%1, %%clock;¥n¥t"
:"=r"(begin), "=r"(end), "=l"(temp) :"r"(AID[threadIdx.x]) );
PTXコードを用いたクロックサイクル 数の計測方法
GPUのアセンブリ言語
◦
インラインアセンブラを用いてソースコード内に記述できる測定したい命令をクロックカウンタの取得命令で挟み，その差分を 求めることで命令の実行クロックサイクル数が計測できる
クロックレジスタの値を取得
クロックカウンタの値が 格納されている この範囲のクロック
サイクル数を計測
z
計測範囲
アドレスシフトなし
シェアードメモリアクセス アクセスアドレスに対応したシフト数の取得
開始クロック取得
終了クロック取得
開始クロック取得
終了クロック取得 アドレスシフトあり
𝑥 ← 𝑀[𝑎]
𝑥 ← 𝑀[𝑠𝑎]
𝑠𝑎 ← 𝑠ℎ𝑖𝑓𝑡(𝑎)
計測するアクセスパターン
1.
Contiguous Access◦
各スレッドは行方向にアクセスバンク0 バンク1 バンク2 バンク3
0 1 2 3
4 5 6 7
8 9 10 11
12 13 14 15
• Random Shift
• Permute Shift 3 1 3 2
バンク0 バンク1 バンク2 バンク3
1 2 3 0
7 4 5 6
9 10 11 8
14 15 12 13
2 1 3 0
バンク0 バンク1 バンク2 バンク3
2 3 0 1
7 4 5 6
9 10 11 8
12 13 14 15
アクセスは衝突しない
アクセスは衝突しない アクセスは衝突しない
計測するアクセスパターン
2.
Stride Access◦
各スレッドは列方向にアクセスバンク0 バンク1 バンク2 バンク3
0 1 2 3
4 5 6 7
8 9 10 11
12 13 14 15
• Random Shift
• Permute Shift 3 1 3 2
バンク0 バンク1 バンク2 バンク3
1 2 3 0
7 4 5 6
9 10 11 8
14 15 12 13
2 1 3 0
バンク0 バンク1 バンク2 バンク3
2 3 0 1
7 4 5 6
9 10 11 8
12 13 14 15
一部のスレッドでアクセスが衝突
アクセスは衝突しない 全てのスレッドでアクセスが衝突
3 計測するアクセスパターン
3.
Diagonal Access◦
各スレッドは斜め方向にアクセスバンク0 バンク1 バンク2 バンク3
0 1 2 3
4 5 6 7
8 9 10 11
12 13 14 15
• Random Shift
• Permute Shift 3 1 3 2
バンク0 バンク1 バンク2 バンク3
1 2 3 0
7 4 5 6
9 10 11 8
14 15 12 13
2 1 3 0
バンク0 バンク1 バンク2 バンク3
2 3 0 1
7 4 5 6
9 10 11 8
12 13 14 15
一部のスレッドでアクセスが衝突
アクセスは衝突しない
一部のスレッドでアクセスが衝突
計測するアクセスパターン
4.
Random Access◦
各スレッドはランダムにアクセスバンク0 バンク1 バンク2 バンク3
0 1 2 3
4 5 6 7
8 9 10 11
12 13 14 15
• Random Shift
• Permute Shift 3 1 3 2
バンク0 バンク1 バンク2 バンク3
1 2 3 0
7 4 5 6
9 10 11 8
14 15 12 13
2 1 3 0
バンク0 バンク1 バンク2 バンク3
2 3 0 1
7 4 5 6
9 10 11 8
12 13 14 15
一部のスレッドでアクセスが衝突
一部のスレッドでアクセスが衝突
一部のスレッドでアクセスが衝突
0 200 400 600 800 1000 1200 1400 1600 1800 2000
1 35 7 911 13 15 17 19 21 23 25 27 29 31
クロックサイクル
ワープ数
Random Read
計測結果（read命令）
0 200 400 600 800 1000 1200 1400 1600 1800 2000
13 57 911 13 15 17 19 21 23 25 27 29 31
クロックサイクル
ワープ数
Contiguous Read
0 200 400 600 800 1000 1200 1400 1600 1800 2000
1 35 7911 13 15 17 19 21 23 25 27 29 31
クロックサイクル
ワープ数
Stride Read
0 200 400 600 800 1000 1200 1400 1600 1800 2000
1 35 7911 13 15 17 19 21 23 25 27 29 31
クロックサイクル
ワープ数
Diagonal Read
アドレスシフトなし Random Shift Permute Shift x2.29
x0.36
x0.18 x0.29
アドレスシフトを行わない場合に 対する高速化率
0 500 1000 1500 2000 2500
1357 9 11 13 15 17 19 21 23 25 27 29 31
クロックサイクル
ワープ数
Random Write
計測結果（write命令）
アドレスシフトなし Random Shift Permute Shift
0 500 1000 1500 2000 2500
135 79 11 13 15 17 19 21 23 25 27 29 31
クロックサイクル
ワープ数
Contiguous Write
0 500 1000 1500 2000 2500
1357 9 11 13 15 17 19 21 23 25 27 29 31
クロックサイクル
ワープ数
Stride Write
0 500 1000 1500 2000 2500
13 579 11 13 15 17 19 21 23 25 27 29 31
クロックサイクル
ワープ数
Diagonal Write
x0.51x2.32
x0.25 x0.43
アドレスシフトを行わない場合に 対する高速化率
計測範囲の変更
アドレスシフトありの場合の計測範囲をメモリアクセスのみに変更する
開始クロック取得
終了クロック取得
アドレスシフトあり アドレスシフトあり
開始クロック取得
終了クロック取得
アドレスシフトがハードウェア化された場合を想定
アドレスシフトの計算時間を無視できる 𝑥 ← 𝑀[𝑠𝑎]
𝑠𝑎 ← 𝑠ℎ𝑖𝑓𝑡(𝑎)
𝑠𝑎 ← 𝑠ℎ𝑖𝑓𝑡(𝑎)
𝑥 ← 𝑀[𝑠𝑎]
計測結果（read命令）
アドレスシフトなし Random Shift Permute Shift
0 10 20 30 40 50 60 70 80 90
13579 11 13 15 17 19 21 23 25 27 29 31
クロックサイクル
ワープ数
Contiguous Read
0 500 1000 1500 2000
13579 11 13 15 17 19 21 23 25 27 29 31
クロックサイクル
ワープ数
Stride Read
0 50 100 150 200 250
135 7911 13 15 17 19 21 23 25 27 29 31
クロックサイクル
ワープ数
Diagonal Read
x22.94
x0.36
アドレスシフトを行わない場合に 対する高速化率
x1
0 50 100 150 200 250
13579 11 13 15 17 19 21 23 25 27 29 31
クロックサイクル
ワープ数
Random Read
x1
4 計測結果（write命令）
アドレスシフトなし Random Shift Permute Shift
0 20 40 60 80 100
13579 11 13 15 17 19 21 23 25 27 29 31
クロックサイクル
ワープ数
Contiguous Write
0 500 1000 1500 2000
13579 11 13 15 17 19 21 23 25 27 29 31
クロックサイクル
ワープ数
Stride Write
0 50 100 150 200 250
13579 11 13 15 17 19 21 23 25 27 29 31
クロックサイクル
ワープ数
Diagonal Write
x22
x0.37
アドレスシフトを行わない場合に 対する高速化率
x1
0 50 100 150 200 250
135 79 11 13 15 17 19 21 23 25 27 29 31
クロックサイクル
ワープ数
Random Write
x1
まとめ
GPUのアセンブリ言語であるPTXコードを用いて，ランダムアドレスシフトの 厳密な性能評価を行った
アドレスシフトを求める時間を含めた場合
◦
strideアクセスで約2.3倍の高速化となり，その他のアクセスパターンで はアドレスシフトを行わない場合が高速となったアドレスシフトを求める時間を含めない場合
