VoltaのL1キャッシュとL2キャッシュの改善

Pascalの場合は、64KBの共有メモリ(Shared Memory)と24KBの1次キャッシュが設けられていた。これに対して、Voltaでは共有メモリと1次キャッシュを一体化して、128KBのメモリを設けた。この一体化されたメモリは、共有メモリとL1キャッシュに分けて使うことができ、64KBを共有メモリとした場合は残りの64KBはL1キャッシュとして使える。

そして、L1キャッシュは共有メモリと同じメモリアレイとなったことで、L1キャッシュの動作がほぼ共有メモリ並みに速くなった。

  • VoltaではL1キャッシュと共有メモリが128KBのメモリアレイを分割して使用する構造になった

    VoltaではL1キャッシュと共有メモリが128KBのメモリアレイを分割して使用する構造になった。結果として、L1キャッシュの容量の大きい分割も共有メモリの大きい分割も可能となった。また、L1キャッシュがほぼ共有メモリ並みの高速アクセスができるようになった

このメモリをL1キャッシュとして使う場合は、遅延の短いノンブロッキングキャッシュとして使える。メモリバンド幅はPascalの場合の4倍以上であり、容量も5倍以上と大きくなっている。なお、L1キャッシュと共有メモリの分割は、cudaFuncSetAttribute関数を呼び出して、共有メモリに割り当てるメモリを%で指定する。ただし、この割り当て%はヒントで、実際に割り当てられるメモリはドライバ依存と書かれている。また、共有メモリに割り当てることができる最大容量は96KBとなっている。

  • VoltaのL1キャッシュはノンブロッキングで、前世代と比べるとアクセス遅延が短く、4倍以上のバンド幅で最大では5倍以上の容量が持てる

    VoltaのL1キャッシュはノンブロッキングで、前世代と比べるとアクセス遅延が短く、4倍以上のバンド幅で最大では5倍以上の容量が持てる。共有メモリに割り当てるメモリ量は%で指定できるが、最大は96KBである

キャッシュは頻繁に使われるデータをハードウェアが自動的に取り込んでくれるので、ソースコードを変更する必要がなく、簡単に使えるというメリットがある。 一方、共有メモリは、格納するデータをソフトウェアでGPUメモリから読んで共有メモリにストアする必要があるが、キャッシュのように、知らないうちに勝手に追い出されてしまうことは無く、安定した性能が得られる。

次の図の右の棒グラフはL1キャッシュを使用した時の性能を共有メモリを使用した時の性能と比較したもので、Pascalの場合は70%程度の性能しか出なかったが、Voltaでは93%と共有メモリに迫る性能が得られている。ただし、スレッド間協調が必要な時や高速のアトミックが必要な時は、共有メモリを使う方が良い。また、キャッシュは知らない間に競合で追い出されてしまい、次のアクセスに長い時間が掛かるなど同じ処理をしても性能がばらつくことがあり得る。このため、安定した性能が必要な場合は、共有メモリを使うべきである。

  • PascalではL1キャッシュを使用すると、共有メモリを使う場合より30%性能が下がったが、Voltaでは、7%の性能低下と共有メモリに近い性能が得られるようになった

    PascalではL1キャッシュを使用すると、共有メモリを使う場合より30%性能が下がったが、Voltaでは、7%の性能低下と共有メモリに近い性能が得られるようになった。ただし、スレッド間協調が必要な時や高速のAtomicが必要な場合など、共有メモリを使う方が良いケースは残っている

L2キャッシュはPascalの4MBからVoltaでは6MBに増強されただけでなく、次の図に示すように、コンフリクトのない場合のアトミックアクセスのスループットが最大2倍に改善されている。一方、大容量のアレイのランダムアクセスを行う場合のスループットはほとんど改善されていない。

  • V100とP100 GPUのL2キャッシュのAtomicアクセス性能の比較

    V100とP100 GPUのL2キャッシュのAtomicアクセス性能の比較。連続アクセスでアレイサイズが大きい場合は、V100のスループットはP100の2倍になっている

Voltaのスレッドスケジューリングの改善

PascalのSMではワープスケジューラが2個であったが、Voltaでは4個に増えている。これに伴い、1個のワープスケジューラが分担するFP32ユニットは32個から16個に半減した。

  • Voltaではワープスケジューラの数が2個から4個に増加し、その結果、1個のワープスケジューラが担当するFP32ユニットは16個に減少した

    Voltaではワープスケジューラの数が2個から4個に増加し、その結果、1個のワープスケジューラが担当するFP32ユニットは16個に減少した

そして、Pascalでは1個のスケジューラに2個のディスパッチャがあり、各ディスパッチャが16CUDAコアを担当していた。スケジューラは1サイクルに1個のワープを選んで、それをどちらかのディスパッチャに渡す。各ディスパッチャは2サイクルに1回、ワープの命令を16CUDAコアに投入し、16CUDAコアは2サイクルを掛けて32スレッドを実行するという構造になっていた。

  • Pascalでは1個のスケジューラに2個のディスパッチャが付き、各ディスパッチャが16CUDAコアを担当していた

    Pascalでは1個のスケジューラに2個のディスパッチャが付き、各ディスパッチャが16CUDAコアを担当していた。そして、ディスパッチャは2サイクルに1回、命令を16CUDAコアに発行し、16CUDAコアはその命令を2回実行して32スレッドを処理していた

これがVoltaでは、1個のスケジューラに1個のディスパッチャとなり、1個のディスパッチャが16個のFP32ユニットとINTユニットを担当する。そして、ディスパッチャは1サイクルに1回、16個のFP32ユニット、あるいはINT32ユニットに命令を投入する。投入された命令は2サイクルを掛けて32スレッドの処理を行うという構造に変わった。

スケジューラとディスパッチャの切れ目が変わっただけと思うかも知れないが、この構造では、ディスパッチャが投入する連続した2命令が、例えばFP32命令とINT32命令と異なっていれば、それらの命令を別ユニットに発行して同時実行できるようになった。

  • Voltaではディスパッチャは毎サイクル命令を発行するので、FP32命令とINT32命令が連続している場合は、FP32命令とINT32命令を並行して実行できるようになった

    Voltaではディスパッチャは毎サイクル命令を発行するので、FP32命令とINT32命令が連続している場合は、FP32命令とINT32命令を並行して実行できるようになった

これによって、命令実行の並列度が改善されて、スループットが向上している。

(次回は3月1日に掲載します)