"Locality is efficiency, Efficiency is power, Power is performance, Performance is King", Bill Dally

マルチスレッディングとは？

CPUとGPUのマルチスレッディングの違いを��ログにまとめていたけど例によって誰も興味なさそう
— arutema47 (@arutema47) 2021年8月16日

つぶやいたら読みたい方が多そうだったので完成させました。

マルチスレッディングとはメモリ遅延を隠蔽しスループットを上げるハードウェアのテクニックです。

ただCPUとGPUで使われ方がかなり異なるため、その違いについて考えてみる記事です。

（SIMDについて並列プログラミングの観点から触れるべきでしたが、~~時間無い~~マルチスレッディングに注目するため初版では省きました。）

本記事について

本記事はCPUとGPUにおけるマルチスレッディングの違いを出来るだけ平易に説明する。

そもそもメモリ遅延とはなにかから始めそれを減らすテクニックについて触れたあとにCPUのマルチスレッディングについて説明。

最後にGPUのマルチスレッディングについて説明し、実際のCUDAコードでどう活用されるか見る。

本記事の想定読者はハードウェアやCSに専門性がない方である。並列プログラミングやHWの専門性を持つ方にとっては教科書レベルの内容で退屈だと思う。

（追記：本記事はコンピュータアーキテクチャ的目線で書かれており、隣接分野から見ると言葉の定義が怪しいかもしれません。随時改定し語弊を減らしていきたいと思います）

tl;dr

（追記）

メモリ遅延を隠すためにCPUではマルチスレッディングが活用されてきた。一方でCPUの主目的はレイテンシを削減することであり、それら機能とのバランス上、スレッド数を大きく増やすことはできなかった。

対照的にGPUはCPU機能を削減し、スループットを最大化するアーキテクチャ。

GPUは大量のマルチスレッディングを行うことで非常に高いスループットが得られ、それは画像処理、科学計算、深層学習に好適。

メモリ遅延とは？

まずメモリ遅延(memory latency)とはなんだろう？

文字通りデータをメモリに読みに行く待機時間でコンピューティングには重要なもの。

Googleの神ことJeff Deanがよく言っている"Latency Numbers Every Programmer Should Know"にメモリ遅延の概算値が載っているのでみてみよう。

L1 cache reference                           0.5 ns
Branch mispredict                            5   ns
L2 cache reference                           7   ns                      14x L1 cache
Mutex lock/unlock                           25   ns
Main memory reference                      100   ns                      20x L2 cache, 200x L1 cache
Compress 1K bytes with Zippy             3,000   ns        3 us
Send 1K bytes over 1 Gbps network       10,000   ns       10 us
Read 4K randomly from SSD*             150,000   ns      150 us          ~1GB/sec SSD
Read 1 MB sequentially from memory     250,000   ns      250 us

https://gist.github.com/jboner/2841832

要するにCPU外からデータを読み込むとすごく（100-1000倍）くらい遅いよ、といっている。

このCPU内メモリ（キャッシュ）からデータを読める（読む）ことをキャッシュヒット、メインメモリ（DRAM)からデータを読まないといけないことをキャッシュミスと言う。

f:id:aru47:20210825211639p:plain

L1キャッシュに対しキャッシュミス（DRAM読み出し）は200倍遅いのでCPUの仕事イメージ的には図のようにほぼメモリ読み出しを待っている状態になる。これはストール(stall)と呼ばれ、CPUは何もしていない待機状態になってしまう。

メモリ遅延を減らすテクニック

このようなメモリ遅延（メモリ読み出しの待機時間）を減らし、CPUにより多くの仕事をしてもらうために色々な技術が考え出された。

キャッシュの階層化、大容量化
プリフェッチ、分岐予測、投機的実行
マルチスレッディング

キャッシュ

キャッシュの大容量化はメモリ遅延を減らす一番ストレートかつ効果的なアプローチ。 大きいは正義。

なのでCPUは世代が進むごとにキャッシュをどんどん大きくする。

半導体屋さん的にもキャッシュの性能向上・拡大は至上命題であり、最新CMOSプロセスがランチされる際には一番最初にキャッシュ性能（SRAMセル性能）が学会で報告される。今だと3nmプロセスですね。

プリフェッチ、分岐予測、投機的実行

f:id:aru47:20210825212711p:plain

またプリフェッチといって将来読まれるであろうデータを予め読み込む強力なテクニックがある。

図のように計算Aの次に必要となるデータBをデータAと同時に読み込むことで、計算Bを遅延なく開始することができる。（正確に予測できれば）メモリ遅延を大幅に減らしCPUの仕事効率を大きく改善する。

実際にデータBのメモリ遅延はなくしたわけではありませんが、プロセッサから見るとあたかもメモリ遅延がなくなったように見えるためメモリ遅延を隠蔽する(hiding memory latency)といったりする。

f:id:aru47:20210826215131p:plain

イメージとしてプリフェッチは必要となるデータを先読みする命令を追加し、予めキャッシュに格納する。

プリフェッチには次に必要となるデータと命令を正確に予測する必要がある。このような命令予測技術を分岐予測や投機的実行といってCPU分野で大きく発展してきた。

余談だが投機的実行はGoogleが発見したx86 CPUの脆弱性であるSpectreやMeltdownで話題になった機能。ものすごく平易に言うと投機的実行のプリフェッチ時に本来権限がないメモリ領域にアクセスできてしまい、悪用すると機密データへアプリケーションがアクセスできてしまう。結局Spectreを完全に抑えるには投機的実行を一部無効化するしかなく、OSアップデートで一部アプリケーションが大きく低速化したのが話題になった。ハードウェア脆弱性はコンピューティングの本質と向き合うものが多く、調べると面白いです（DRAMの脆弱性をつくrowhammerとか）

CPUのマルチスレッディング

(追記:はてブにて演算器を共用する同時マルチスレッディングとメモリ隠蔽するマルチスレッディングの話が混ざっているという指摘がありました。仰る通りなので補足・改定します。）

それではようやく本題のマルチスレッディングについて説明する。

CPUコア（物理的に存在するコア）を2倍に見せる技術がマルチスレッディング。

f:id:aru47:20210825215230p:plain

図の通り自分のCPU(Ryzen4800)は物理的に8コアですが、マルチスレッディングにより16コア（論理プレセッサ数）とコンピュータは認識する。

f:id:aru47:20210825220500p:plain

図のようにマルチスレッディングを使うCPUコアは2つに見える（コアAとコアB）。

コアAが計算、データ読み込みと動作するときにキャッシュミスが発生したとします。そうすると前述したように大きなメモリ遅延が発生し、CPUはストールし待機状態に入ってしまう。

そこでマルチスレッドCPUはプロセススイッチと呼ばれる動作が入り、コアBの命令動作に移行します。コアBが計算している間、コアAはデータを待っているため結果的に1つの物理CPUコアを有効活用できる。

また同様にコアBでもキャッシュミスが生じると今度はコアAにプロセススイッチし、データが届き次第計算を開始する。

このようにCPUのマルチスレッディングは生じてしまったメモリ遅延の裏側で別の命令・計算を行うことで、あたかも2つのコアがあるかのように見せかける技術だ。

上記がマルチスレッディングの重要な意義であり、それはCPUもGPUも変わらない。

またマルチスレッディングによりレイテンシは変わらないが、スループットを大幅に向上することが可能だ。

（補足：上記例では並列プログラミングの観点からマルチスレッディングの単純な例を書きました。実際のCPUでは同時マルチスレッディング（SMT)というもっと複雑な技術が使用されています。上記ではメモリ遅延を隠蔽する効果でしたが、SMTではCPU内のある処理で使用していない演算器を別スレッドでも利用することで更にハードウェア利用効率を高めています。こちらはメモリ遅延を緩和する効果はありませんが、タスクが演算律速である場合は高速化に寄与します。）

ja.wikipedia.org

レイテンシとスループット

f:id:aru47:20210826215835p:plain

レイテンシとスループットという概念は重要でそれぞれどう改善されるか理解しておこう。

レイテンシとはある命令が実行されてから完了するまでにかかった時間。例えば命令Aが終了するまでにかかる時間をレイテンシとして評価する。これはマルチスレッディングによっては改善しないが、分岐予測で上手くメモリ遅延を隠蔽できたら改善できる。

一方でスループットとは1秒間に実行した命令の量（データ量）を表す。例えば図のように計算が1bitで全部で1秒かかったとすると4bit/sのように表す。

スループットはマルチスレッディングを活用することで大きく改善でき、更にスレッド数を増やしても改善できることがわかる。だが後述するようにCPUは分岐予測、SMTに加え多種多様の命令を実現するためパイプラインが複雑でスレッド数をこれ以上増やしてもあまりリターンが得られない事が知られている。

では分岐予測や煩雑なx86命令などスループット改善に邪魔な機能()を全て取っ払ってマルチスレッド数を増やしまくったら最速プロセッサが作れる！という発想で発展してきたのがGPUだ。

GPUのマルチスレッディング

GPUマルチスレッディングの基本思想

CPUのジョブ(OS等）は基本的に逐次的であり、ジョブのレイテンシを最適化することを目標に設計されている。そのため非常に複雑な分岐予測機を持っており、マルチスレッド数も2と小さく抑えられている。

一方でGPUはそもそもレイテンシは全く気にせず、スループットを最大化するシステム思想。根本からCPUの設計思想と異なることを意識する必要がある。

前のCPUでは各コアにジョブを分割していたが、GPU（CUDA)ではスレッドで命令を分割するマルチスレッドプログラミングを行う。

f:id:aru47:20210826222617p:plain

（参考：CS15-418)

ここではスレッドが4つあるGPUコアを考える。まずスレッド0が計算を始めるがストールが発生するとする。するとメモリ遅延が生じ、プロセススイッチによりスレッド1に切り替える。スレッド1も同様にストールが発生すると次はスレッド2に。。とプロセススイッチを繰り返す。

最後にスレッド3もストールすると今度はスレッド0のデータ読み出しが完了したためスレッド0にスイッチし計算を完了させる。そしてデータが到着したスレッドに順々に切り替え、計算を行っていく。

マルチスレッディング動作自体はCPUとほぼ変わらない（ストールしたらプロセススイッチを行う）。しかしCPUと違いGPUの計算機構成が単純であるため多くのスレッドを展開可能なことが特徴である（GPUはキャッシュはあるが分岐予測は持たない）。CUDAでは256スレッドまでプログラマは定義することができ、2スレッドしかないCPUと大きく異る。

ここで注目するべきことは

プロセッサの（計算器）使用率は非常に高い（ほぼ100%）である
メモリ遅延が発生しているにも関わらず、その影響をほぼ受けていない。

の2点である。

この例ではストールが発生してもすぐにデータがキャッシュに格納されているスレッドがあると仮定したため、ほぼプロセッサは100%計算を実施することができた。これは非常に計算機の利用効率が高く、プロセッサの理論スループット近く出せていることを示している。繰り返しになるが、マルチスレッド数が多ければ多いほどシステムスループットは改善でき、GPUはスループットを最大化するように設計されている。

また1点目が達成できている理由は、マルチスレッディングによって上手くメモリ遅延を隠蔽できているからだ。スレッド0がストールしている間に他スレッドが計算を行うことでプロセッサの計算機は動き続けることができる。

このスループット最大化処理方法はOSなどのシステムを走らせるには向かないが、科学計算、画像、ディープラーニングといった大量の計算が要求されるアプリケーションと相性が良い。

上記の考察によりGPUのハードウェア構造も見えてきたのではないだろうか？ GPUのようなスループット最大化システム（throughput oriented systems)は複数スレッドを一度に展開し多数のデータフェッチを実施し、データの準備ができたスレッドから片っ端に計算する。そのためデータフェッチ機能は大量に必要で、プロセススイッチ機構などもCPUより複雑になる。

GPUのコア構造

f:id:aru47:20210827211617p:plain

それでは単純化したGPUのコア(SM core)をみてみよう。

マルチスレッド管理をプログラマにやらせると煩雑すぎて死んでしまう。

そのためスレッド管理はGPUではコア内にスケジューラをもたせており、ハードウェアがスレッドのスケジュールを管理してくれる。このマルチスレッドスケジューラを備えている点がGPUの固有な点であるため取り上げた。

命令セットプールにはそのコアに割り当てられた命令がプールされている。そしてWarpスケジューラが上記の例で示したように、スレッドがストールしたらデータがフェッチできているスレッドに切り替えるというスケジューリングを行う（実際のスレッド管理はブロック毎になっているが基本概念は同じ）。

CUDAコードを読んでみよう

それでは最後にCUDAコードを読んでみてGPUのマルチスレッディングについて理解を深めよう。

developer.nvidia.com

CUDAのHello World的なコードを読んで見る。

#include <stdio.h>

__global__
void saxpy(int n, float a, float *x, float *y)
{
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  if (i < n) y[i] = a*x[i] + y[i];
}

int main(void)
{
  //ベクトル定義
  int N = 1<<20;  // 100万
  float *x, *y, *d_x, *d_y;
  x = (float*)malloc(N*sizeof(float));
  y = (float*)malloc(N*sizeof(float));

  cudaMalloc(&d_x, N*sizeof(float)); 
  cudaMalloc(&d_y, N*sizeof(float));

  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  //GPUへ配列をコピー
  cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);
  cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice);

  // 計算の実行
  saxpy<<<(N+255)/256, 256>>>(N, 2.0f, d_x, d_y);
  
  // 出力をコピー、メモリ開放
  cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost);

  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = max(maxError, abs(y[i]-4.0f));
  printf("Max error: %f\n", maxError);

  cudaFree(d_x);
  cudaFree(d_y);
  free(x);
  free(y);
}