高速化手法: ループ展開

ループ展開（英：Loop Unrolling/Loop Unwinding）は、高速化の基本でもあり、効果が高い時は少なくないです。 Wikipediaによると、

ループ展開の目的は、毎回の繰り返しごとに発生する「ループの終了」条件のテストを減少させる（もしくはなくす）事によって、実行速度を向上させることである。ループは、ループ自体を制御するためのオーバーヘッドがなくなるように、独立した命令ブロックの連続に書き換えることができる。

したがって、シリアルプロセッサでも、ループ展開によってプログラムの実行時間を減らすことができます。 最近、大体のプロセッサはパイプライン化が可能なので、ループ展開によってデータ読み書きのレイテンシー を減らすことにもつながります。つまり、ループ展開は（ソフトウェア）パイプラインを促します。

サンプルコード:

1
2
3

  for (int i = 0; i < N_DATA; i++) {
      a[i] = alpha * b[i] + c[i];
  }

サンプルコードにループ展開をあてると:

1
2
3
4
5
6

  for (int i = 0; i < N_DATA; i+=4) {
      a[i+0] = alpha * b[i+0] + c[i+0];
      a[i+1] = alpha * b[i+1] + c[i+1];
      a[i+2] = alpha * b[i+2] + c[i+2];
      a[i+3] = alpha * b[i+3] + c[i+3];
  }

2つ以上のプロセッサが備わっているデバイス（並列プロセッサ）には、ループ内の各イテレーションを 各プロセッサにあてることで、同時実行可能になり、実行時間をより低くすることができます。

しかし、並列実行可能な環境に、ループを展開してワークロードを分割する時に、一つ注意点を頭に入れなければ ならないことがあります。それは、データ依存性のことです。間違ってループ展開をしてしまうと、不正解の実行結果に導かれたり、 最大限の高速化を得られず、だったりします。 ここではサンプルコードにはデータ依存性がないので、後々時間があればデータ依存性について書きます。

OpenMP実装:

1
2
3
4

#pragma omp parallel for
  for (int i = 0; i < N_DATA; i++) {
      a[i] = alpha * b[i] + c[i];
  }

CUDA実装:

1
2
3
4
5

__global__
void saxpy(float *a, float *b, float *c, float alpha) {
     int i = threadIdx.x + blockDim.x * blockIdx.x;
     a[i] = alpha * b[i] + c[i];
}

OpenCL実装:

1
2
3
4
5

__kernel
void saxpy(__global float *a, __global float *b, __global float *c, float alpha) {
     int i = get_global_id(0);
     a[i] = alpha * b[i] + c[i];
}