CUDA PagedLockedHostMemoryの性能について
マルチポスト元→http://exth.net/~tgbt/wordpress/2010/05/18/3135/
なんだか久々にCUDAネタ。twitterでぶつぶつ言いながらやった実験なので、こちらに投稿しておく。（後であちらにも投げるつもりだけど。）
メチャクチャ長くなったけど仕方がない。
Well-KnownなCUDA高速化技術の１つにPage-Lockedなメモリを使うっていうのがある。というわけで、PinnedMemoryに関する実験をやってみた。
一応真面目にやってはいるんだけど、まだ理解が追いついていない部分もあるので、参考にする際は自己責任で。むしろ情報交換しましょう。間違いへの突っ込みも歓迎。
実験に使ったプログラムのフルソースコードなどは末尾。
とりあえずProgrammingGuideを読んでみる
3.0のProgrammingGuideだと3.2.5項。
以下要約：
- 3.2.5 Page-Locked Host Memory
- 確保しすぎるとOSの制御を阻害するとか、副作用もあるので気をつけようね！
- 3.2.5.1 Portable Memory
- 複数スレッド（OS側のスレッド）で使うときはcudaHostAllocPortableを設定すると良いよ！
- 3.2.5.2 Write-Combining Memory
- cudaHostAllocWriteCombinedを設定するとOSのキャッシュ管理から切り離されるよ！
- host-device間の転送速度が上がるよ！
- hostでWCメモリを読むと遅いよ！だからhostからは書き込むだけにした方が良いよ！（device側の読み書き速度がどうなるのかは書かれてないなあ）
- 3.2.5.3 Mapped Memory
- cudaHostAllocMappedを設定すると、hostとdeviceで同じメモリ空間を使えるよ！
- host-device間でメモリコピーを明示的に書かなくても良くなるし、計算と転送のオーバーラップも簡単だよ！
- ただし、read-after-write/write-after-read/write-after-writeの問題があるのでstreamsかeventsでタイミング制御をする必要はあるよ！
- atomic演算は注意してね！
だいたいこんなかんじ。間違ってたらごめん。
考えてみる
なんとなくわかった。
- WriteCombinedにすれば転送速度が上がる。
OK、OK。ただしOSのキャッシュ制御の都合もあるから、CPUでもGPUでもいじりまくるようなデータは気をつけろということだな。
- Mappedにすればアドレスの一元化が可能になる。memcpyとか不要。通信と演算のオーバーラップとか楽勝。
これは便利かもしれない。でも、それってつまり必要に応じてデータがコピーされるってことだよね？アクセスパターンとか次第で酷い性能になるよね？？？
テストしてみる
まずはサンプルに入っているbandwidthTestを走らせる。
>/usr/local/cudasdk/3.0/C/bin/linux/release/bandwidthTest --memory=pageable [bandwidthTest] /usr/local/cudasdk/3.0/C/bin/linux/release/bandwidthTest Starting... Running on... Device 0: GeForce GTX 480 Quick Mode Host to Device Bandwidth, 1 Device(s), Paged memory Transfer Size (Bytes) Bandwidth(MB/s) 33554432 5395.7 Device to Host Bandwidth, 1 Device(s), Paged memory Transfer Size (Bytes) Bandwidth(MB/s) 33554432 4587.9 Device to Device Bandwidth, 1 Device(s) Transfer Size (Bytes) Bandwidth(MB/s) 33554432 111025.4 [bandwidthTest] - Test results: PASSED Pressto Quit... >/usr/local/cudasdk/3.0/C/bin/linux/release/bandwidthTest --memory=pinned [bandwidthTest] /usr/local/cudasdk/3.0/C/bin/linux/release/bandwidthTest Starting... Running on... Device 0: GeForce GTX 480 Quick Mode Host to Device Bandwidth, 1 Device(s), Pinned memory, Write-Combined Memory Enabled Transfer Size (Bytes) Bandwidth(MB/s) 33554432 5815.4 Device to Host Bandwidth, 1 Device(s), Pinned memory, Write-Combined Memory Enabled Transfer Size (Bytes) Bandwidth(MB/s) 33554432 6105.8 Device to Device Bandwidth, 1 Device(s) Transfer Size (Bytes) Bandwidth(MB/s) 33554432 110967.5 [bandwidthTest] - Test results: PASSED Press
- -
to Quit...
- -
H2Dで10%、D2Hで30%くらい高速化されてる。これはすごい。
ソースを眺めたところ、どうやら本当にcudaHostAllocにcudaHostAllocWriteCombinedを指定しているだけっぽい。
コードを書いてテストしてみる
以下の比較をしてみよう：
- 普通にMalloc/cudaMalloc
- cudaHostAlloc with cudaHostAllocWriteCombined
- HostAllocMapped
それぞれについて、
を比較。
ついでに、カーネルをCPUにやらせた版も作成して様子見。
GPUの計算は、いずれもgridDim=(1,1,1) blockDim=(128,1,1)という適当な設定。これでも十分に傾向くらいは掴めると思いたい。
カーネルは以下の通り：
// simpleテスト用(part1)
__global__ void kernel1
(float *data1, float *data2, int size)
{
int i;
int begin, end, step;
begin = threadIdx.x;
end = size;
step = blockDim.x;
for(i=begin; iカーネル以外は添付ファイルにて。
基本的にカーネルを一発実行するだけのプログラム。
まずは普通にMalloc/cudaMallocした場合：
./normal 0 50000000 size = 50000000Byte(48828KByte, 47MByte), 12500000 elements device max=2, n=0 part1(simple) memcpy : 34878.208160 nsec, 2695.092580 MByte/sec all : 1128581.047058 nsec part2(indirect) memcpy : 34505.214691 nsec, 2724.225913 MByte/sec all : 49117.088318 nsec part3(indirect random) memcpy : 34498.271942 nsec, 2724.774161 MByte/sec all : 49160.003662 nsec ./normal_wk 0 50000000 size = 50000000Byte(48828KByte, 47MByte), 12500000 elements device max=2, n=0 part1(simple) memcpy+kernel : 77812.156677 nsec, 1208.037459 MByte/sec all : 1167029.857635 nsec part1(CPU) CPUkernel : 24100.780487 nsec memcpy+CPUkernel : 43626.399994 nsec, 2154.658647 MByte/sec all : 48449.993134 nsec part2(indirect) memcpy+kernel : 115774.589539 nsec, 811.922550 MByte/sec all : 130471.944809 nsec part2(CPU) CPUkernel : 28672.933578 nsec memcpy+CPUkernel : 48061.695099 nsec, 1955.819490 MByte/sec all : 62689.065933 nsec part3(indirect random) memcpy+kernel : 184142.211914 nsec, 510.475024 MByte/sec all : 198773.860931 nsec part3(CPU) CPUkernel : 126466.035843 nsec memcpy+CPUkernel : 145819.137573 nsec, 644.634179 MByte/sec all : 160428.047180 nsec
- _wkがカーネルも実行している版。流石に間接参照やランダム参照が遅い。でも実はもっと遅くなるんじゃないかと思っていた。
- これが基準値ということで。
- （CPUがkernelを実行した場合にMByte/secを表示する意味は無さそうだなぁ。）
WriteCombined設定をした場合：
./wc 0 50000000 size = 50000000Byte(48828KByte, 47MByte), 12500000 elements device max=2, n=0 part1(simple) memcpy : 16025.951385 nsec, 5865.486406 MByte/sec all : 1158394.098282 nsec part2(indirect) memcpy : 16022.495270 nsec, 5866.751615 MByte/sec all : 82098.960876 nsec part3(indirect random) memcpy : 16023.904800 nsec, 5866.235551 MByte/sec all : 82008.123398 nsec ./wc_wk 0 50000000 size = 50000000Byte(48828KByte, 47MByte), 12500000 elements device max=2, n=0 part1(simple) memcpy+kernel : 58642.913818 nsec, 1602.921715 MByte/sec all : 1202525.138855 nsec part1(CPU) CPUkernel : 1574666.976929 nsec memcpy+CPUkernel : 1590638.549805 nsec, 59.095764 MByte/sec all : 1648941.993713 nsec part2(indirect) memcpy+kernel : 97171.073914 nsec, 967.366071 MByte/sec all : 163262.128830 nsec part2(CPU) CPUkernel : 1573044.061661 nsec memcpy+CPUkernel : 1589031.250000 nsec, 59.155539 MByte/sec all : 1657025.098801 nsec part3(indirect random) memcpy+kernel : 165483.520508 nsec, 568.032392 MByte/sec all : 231626.033783 nsec part3(CPU) CPUkernel : 1745186.090469 nsec memcpy+CPUkernel : 1761167.358398 nsec, 53.373690 MByte/sec all : 1827242.136002 nsec
- カーネル非実行のmemcpyをみてわかるように、転送速度がとても速くなった。ただしmalloc等の処理に手間取るようで、allの値はむしろ増加。
- CPUkernelの値が増加。ProgrammingGuideにあったhostから叩くデータは良くない的な記述が反映されている模様。注意が必要だなあ。
つづいてMapped：
./mapped 0 50000000 size = 50000000Byte(48828KByte, 47MByte), 12500000 elements device max=2, n=0 part1(simple) memcpy : 2.528000 nsec, 37183545.414364 MByte/sec all : 1177767.992020 nsec part2(indirect) memcpy : 2.336000 nsec, 40239724.436970 MByte/sec all : 101921.796799 nsec part3(indirect random) memcpy : 2.336000 nsec, 40239724.436970 MByte/sec all : 101827.144623 nsec ./mapped_wk 0 50000000 size = 50000000Byte(48828KByte, 47MByte), 12500000 elements device max=2, n=0 part1(simple) memcpy+kernel : 141869.567871 nsec, 662.580435 MByte/sec all : 1318675.041199 nsec part1(CPU) CPUkernel : 17521.142960 nsec memcpy+CPUkernel : 17529.504776 nsec, 5362.387654 MByte/sec all : 109644.174576 nsec part2(indirect) memcpy+kernel : 158290.374756 nsec, 593.845331 MByte/sec all : 260251.045227 nsec part2(CPU) CPUkernel : 20175.933838 nsec memcpy+CPUkernel : 20186.048508 nsec, 4656.681567 MByte/sec all : 123075.962067 nsec part3(indirect random) memcpy+kernel : 340836.029053 nsec, 275.792440 MByte/sec all : 442532.062531 nsec part3(CPU) CPUkernel : 152827.024460 nsec memcpy+CPUkernel : 152834.014893 nsec, 615.046330 MByte/sec all : 255613.803864 nsec
- カーネル非実行：明示的なmallocが消えたのでmemcpyの値が吹っ飛んだ。allの値がさらに増えていることから、やはりオーバーヘッドは小さくないのだろうか。
- カーネル実行：困ったことに、何も速くなっている気がしない。
最後に、2つのパラメタを両方設定してみた：
./mpwc 0 50000000 size = 50000000Byte(48828KByte, 47MByte), 12500000 elements device max=2, n=0 part1(simple) memcpy : 2.528000 nsec, 37183545.414364 MByte/sec all : 1180523.872375 nsec part2(indirect) memcpy : 2.336000 nsec, 40239724.436970 MByte/sec all : 106031.894684 nsec part3(indirect random) memcpy : 2.336000 nsec, 40239724.436970 MByte/sec all : 105930.805206 nsec ./mpwc_wk 0 50000000 size = 50000000Byte(48828KByte, 47MByte), 12500000 elements device max=2, n=0 part1(simple) memcpy+kernel : 141808.486938 nsec, 662.865827 MByte/sec all : 1324396.133423 nsec part1(CPU) CPUkernel : 1578817.129135 nsec memcpy+CPUkernel : 1578768.432617 nsec, 59.540081 MByte/sec all : 1676369.905472 nsec part2(indirect) memcpy+kernel : 158298.950195 nsec, 593.813161 MByte/sec all : 264221.906662 nsec part2(CPU) CPUkernel : 1572522.163391 nsec memcpy+CPUkernel : 1572480.102539 nsec, 59.778181 MByte/sec all : 1678436.040878 nsec part3(indirect random) memcpy+kernel : 340668.151855 nsec, 275.928347 MByte/sec all : 446609.020233 nsec part3(CPU) CPUkernel : 1745092.153549 nsec memcpy+CPUkernel : 1745044.189453 nsec, 53.866831 MByte/sec all : 1851984.977722 nsec
- まぁ、memcpyが要らないプログラムにmemcpy高速化をしているようなものなので、どうにもならんわな。
結論
というわけで、とりあえず書いて走らせてはみたんだけど、今ひとつだ。
- WriteCombined設定で通信速度が向上
- ただしhostから叩く際の速度低下があり得るので対象データに注意
- Mappedが使い物にならん（コーディングが楽なのは事実）
という感じ。
もちろん、使い方が悪い可能性は否定しないので、今後追加調査を要すると言うことで。特に、通信と演算のオーバーラップがまだ試せていない。Mappedの本懐はそちらかなあと思っている。あとはPortable設定で複数GPU実験だな！
