title: OpenCL slow -- not sure why
body: <p>I'm teaching myself OpenCL by trying to optimize the mpeg4dst reference audio encoder. I achieved a 3x speedup by using vector instructions on CPU but I figured the GPU could probably do better.</p>
  
  <p>I'm focusing on computing auto-correlation vectors in OpenCL as my first area of improvement. The CPU code is:</p>
  
  <pre><code>for (int i = 0; i &lt; NrOfChannels; i++) {
      for (int shift = 0; shift &lt;= PredOrder[ChannelFilter[i]]; shift++)
          vDSP_dotpr(Signal[i] + shift, 1, Signal[i], 1, &amp;out, NrOfChannelBits - shift);
  }
  NrOfChannels = 6
  PredOrder = 129
  NrOfChannelBits = 150528.
  </code></pre>
  
  <p>On my test file, this function take approximately 188ms to complete.</p>
  
  <p>Here's my OpenCL method:</p>
  
  <pre><code>kernel void calculateAutocorrelation(size_t offset,
                                   global const float *input,
                                   global float *output,
                                   size_t size) {
  size_t index = get_global_id(0);
  size_t end = size - index;
  float sum = 0.0;
  
  for (size_t i = 0; i &lt; end; i++)
      sum += input[i + offset] * input[i + offset + index];
  
  output[index] = sum;
  }
  </code></pre>
  
  <p>This is how it is called:</p>
  
  <pre><code>gcl_memcpy(gpu_signal_in, Signal, sizeof(float) * NrOfChannels * MAXCHBITS);
  
  for (int i = 0; i &lt; NrOfChannels; i++) {
      size_t sz = PredOrder[ChannelFilter[i]] + 1;
      cl_ndrange range = { 1, { 0, 0, 0 }, { sz, 0, 0}, { 0, 0, 0 } };
  
      calculateAutocorrelation_kernel(&amp;range, i * MAXCHBITS, (cl_float *)gpu_signal_in, (cl_float *)gpu_out, NrOfChannelBits);
      gcl_memcpy(out, gpu_out, sizeof(float) * sz);
  }
  </code></pre>
  
  <p>According to Instruments, my OpenCL implementation seems to take about 13ms, with about 54ms of memory copy overhead (gcl_memcpy).</p>
  
  <p>When I use a much larger test file, 1 minute of 2-channel music vs, 1 second of 6-channel, while the measured performance of the OpenCL code seems to be the same, the CPU usage falls to about 50% and the whole program takes about 2x longer to run.</p>
  
  <p>I can't find a cause for this in Instruments and I haven't read anything yet that suggests that I should expect very heavy overhead switching in and out of OpenCL.</p>