title: Worse performance when using a bigger multiple work group size of CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE
body: <p>I am writing a matrix multiplication program in OpenCL. In order to utilize local memory, I divided the original matrix into <code>block_size * block_size</code> sub-matrix, where <code>block_size</code> is equals to a square root of <code>work group size</code> and will be calculated in one work group.</p>
  <p>My GPU is <em>RX580</em>, and <code>CL_KERNEL_WORK_GROUP_SIZE</code> returns <code>256</code>, so a square root is <code>16</code>, while <code>CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE</code> returns <code>64</code>. And <strong>after some profiling, I found out that settting the <code>block_size</code> to <code>8*8</code> performs 4 times faster than setting it to <code>16*16</code>, which seems unintuitive to me.</strong> Any explanation to this?</p>
  <p>EDIT: CodeXL screenshot</p>
  <p>8*8 work group
  <a href="https://i.stack.imgur.com/9VyX5.png" rel="nofollow noreferrer"><img src="https://i.stack.imgur.com/9VyX5.png" alt="enter image description here" /></a>
  <a href="https://i.stack.imgur.com/v148L.png" rel="nofollow noreferrer"><img src="https://i.stack.imgur.com/v148L.png" alt="enter image description here" /></a></p>
  <p>16*16 work group
  <a href="https://i.stack.imgur.com/subEe.png" rel="nofollow noreferrer"><img src="https://i.stack.imgur.com/subEe.png" alt="enter image description here" /></a>
  <a href="https://i.stack.imgur.com/gnq3z.png" rel="nofollow noreferrer"><img src="https://i.stack.imgur.com/gnq3z.png" alt="enter image description here" /></a></p>
  <p>I noticed something unusual. How can <code>SALUBusy%</code> exceed 100% in the 8*8 work group case?</p>
  <p>EDIT2: Add my kernel code, which might be useful</p>
  <pre class="lang-c prettyprint-override"><code>size_t getIndex(int row, int col, int width)
  {
      return row*width+col;
  }
  
  
  __kernel void matrixMul(__global const float* restrict a, 
      __global const float* restrict b,  //assume tranposed
      __local float* restrict a_local,
      __local float* restrict b_local,
      __global float* restrict result) 
  {
      const int row=get_global_id(0);
      const int col=get_global_id(1);
  
      const int limit=get_global_size(0);
  
  
      const int blockSize=get_local_size(0);
      const int gidRow=get_local_id(0);
      const int gidCol=get_local_id(1);
  
      
      float sum=0;
      for(int blockIndex=0; blockIndex &lt; (limit/blockSize); ++blockIndex)
      {
          /*copy block -&gt; local memory*/
          a_local[getIndex(gidRow, gidCol, blockSize)]=a[getIndex(row, col, limit)];
          b_local[getIndex(gidRow, gidCol, blockSize)]=b[getIndex(row, col, limit)];
          barrier(CLK_LOCAL_MEM_FENCE);
  
          /*block multiply*/
          __local float* restrict a_local_row=&amp;a_local[getIndex(gidRow, 0, blockSize)];
          __local float* restrict b_local_row=&amp;b_local[getIndex(gidCol, 0, blockSize)];
          for(int i=0; i&lt;blockSize; ++i)
          {
              sum+= (*a_local_row) * (*b_local_row);
              ++a_local_row;
              ++b_local_row;
          }
      }
  
      result[getIndex(row, col, limit)]=sum;
  }
  </code></pre>