Teknik Pemrograman Parallel: Loop Unrolling

Loop unrolling, atau juga biasa disebut loop unwinding, adalah salah satu teknik mendasar dalam pemrograman parallel, dan biasanya efektif untuk mempercepat program. Pada umumnya, loop unrolling dilakukan dengan cara men-disassemble sebuah loop sehingga menjadi deretan instruksi yang sama (dengan index pointer ke memory yang berbeda).

Walaupun dengan serial processor, loop unrolling juga bisa mengoptimisasi eksekusi program, karena instruksi untuk mengontrol loop tersebut akan berkurang, seperti instruksi branching atau pointer arithmetic.

Sebagai info tambahan, gcc menjalankan optimisasi loop unrolling terhadap sebuah program C secara implicit, dengan memberi opsi -O, -O1, atau -O3 sewaktu mengompile program tersebut. Opsi -O2 tidak menjalankan optimisasi ini.

Loop unrolling juga menstimulasi software pipelining, dan dengan menggunakan pipelining, overhead yang diakibatkan oleh data read-write bisa dikurangi.

Contoh kode:

1
2
3

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

Setelah dilakukan loop unrolling:

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];
  }

Pada device yang mempunyai processor lebih dari satu (atau parallel processors), workload di dalam loop tersebut akan dioper ke masing-masing processor dan dijalankan secara parallel.

Yang perlu diingat waktu membagi workload dari sebuah loop, adalah data dependency. Kalau kita tidak teliti sewaktu membagi workload untuk masing-masing thread, maka kita bisa dihadapkan dengan hasil eksekusi yang tidak sesuai/salah, atau optimisasi yang kurang efektif. Di artikel ini, tidak ada data dependency di dalam contoh kodenya, jadi tidak akan dibahas lebih lanjut disini.

Implementasi OpenMP:

1
2
3
4

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

Implementasi 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];
}

Implementasi 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];
}