opencl - OpenCL削減を理解するのに助けが必要

Question

私は次のメモを読んでいます: http://developer.amd.com/resources/documentation-articles/articles-whitepapers/opencl-optimization-case-study-simple-reductions/

次のカーネルはデータのチャンクを削減する必要がありますが、その一部が理解できません。

while (global_index < length) ....  global_index += get_global_size(0)

シーケンシャルに配置されたグローバル ストレージからデータを読み取る方が賢明であると私は信じていました。つまり、k、k+1、k+2 でのデータの読み取りは、k+1000、k+2000、k+3000 の読み取りよりも高速です。これは、 global_index += get_global_size(0) と言うときに彼らがしていることではありませんか?

__kernel
void reduce(__global float* buffer,
            __local float* scratch,
            __const int length,
            __global float* result) {

  int global_index = get_global_id(0);
  float accumulator = INFINITY;
  // Loop sequentially over chunks of input vector
  while (global_index < length) {
    float element = buffer[global_index];
    accumulator = (accumulator < element) ? accumulator : element;
    global_index += get_global_size(0);
  }

  // Perform parallel reduction
  int local_index = get_local_id(0);
  scratch[local_index] = accumulator;
  barrier(CLK_LOCAL_MEM_FENCE);
  for(int offset = get_local_size(0) / 2;
      offset > 0;
      offset = offset / 2) {
    if (local_index < offset) {
      float other = scratch[local_index + offset];
      float mine = scratch[local_index];
      scratch[local_index] = (mine < other) ? mine : other;
    }
    barrier(CLK_LOCAL_MEM_FENCE);
  }
  if (local_index == 0) {
    result[get_group_id(0)] = scratch[0];
  }
}