title: GPU Memory bandwidth theoretical vs practical
body: <p>As part of an algorithm profiling running on GPU I feel that I'm hitting the memory bandwidth.</p>
  
  <p>I have several complex kernels performing some complicated operations (sparse matrix multiplications, reduction etc) and some very simple ones and it seems that all (significant ones) hit ~79GB/s bandwidth wall when I calculate the total data read/written for each one of them, regardless the complexity of them, while the theoretical GPU bandwidth is 112GB/s (nVidia GTX 960)</p>
  
  <p>The data set is very large operating on vectors of ~10,000,000 float entries so I get good measurements/statistics from <code>clGetEventProfilingInfo</code> between <code>COMMAND_START</code> and <code>COMMAND_END</code>. All the data remains in GPU memory during algorithm run so there virtually no host/device memory transfer (also it is not measured by profiling counters)</p>
  
  <p>Even for a very simple kernel (see below) that solves <code>x=x+alpha*b</code> where x and b are huge vectors of ~10,000,000 entries, I don't get close to the theoretical bandwidth (112GB/s) but rather is running on ~70% of the maximum (~79GB/s) </p>
  
  <pre><code>__kernel void add_vectors(int N,__global float *x,__global float const *b,float factor)
  {
      int gid = get_global_id(0);
      if(gid &lt; N)
          x[gid]+=b[gid]*factor;
  }
  </code></pre>
  
  <p>I calculate data transfer for this particular kernel per run as N * (2 + 1) * 4:</p>
  
  <ul>
  <li>N - size of vector = ~10,000,000</li>
  <li>2 loads and 1 store per vector entry</li>
  <li>4 for sizeof float</li>
  </ul>
  
  <p>I expected that for such a simple kernel I need to get much closer to the bandwidth limits, what do I miss?</p>
  
  <p>P.S.: I get similar numbers from CUDA implementation of the same algorithm</p>