title: Kernel attributes in the two-stage reduction example in OpenCL proposed by AMD
body: <p>I have some problems understanding the two-stage reduction algorithm described <a href="http://developer.amd.com/resources/documentation-articles/articles-whitepapers/opencl-optimization-case-study-simple-reductions/" rel="nofollow">here</a>.</p>
  
  <pre><code>__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 &lt; length) {
      float element = buffer[global_index];
      accumulator = (accumulator &lt; 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 &gt; 0;
        offset = offset / 2) {
      if (local_index &lt; offset) {
        float other = scratch[local_index + offset];
        float mine = scratch[local_index];
        scratch[local_index] = (mine &lt; other) ? mine : other;
      }
      barrier(CLK_LOCAL_MEM_FENCE);
    }
    if (local_index == 0) {
      result[get_group_id(0)] = scratch[0];
    }
  }
  </code></pre>
  
  <p>I understand the basic idea, but I am not sure about the while-loop. As far as I inderstand, the attribute <em>length</em> specifies the number of elements in the buffer, i.e. how many elements do I want to process at all. But <em>get_global_size</em> returns the global number of work-items. Aren't <em>length</em> and <em>get_global_size</em> equal then? This would mean that the while-loop condition wil be satisfied only once. Shouldn't we use <em>get_local_size</em> instead of <em>get_global_size</em>?</p>