title: What's wrong with Metal instrument simdgroup_load or simdgroup_store?
body: <p>OS: MacOS 12.2.1</p>
  <p>Hardwear: MacBook Pro 2020, M1</p>
  <p>Metal: 2.4</p>
  <p>Xcode： 13.2.1</p>
  <p>Here is my test computer kernel,which <strong>read input buffer with simdgroup_load adn write output buffer with simdgroup_store</strong></p>
  <pre><code>kernel void fun(
  const device half * Src                 [[ buffer(0) ]],
  constant uint4 &amp; SrcShape               [[ buffer(1) ]],
  device half * Dst                       [[ buffer(2) ]],
  constant uint4 &amp; DstShape               [[ buffer(3) ]],
  const device half * Weight              [[ buffer(4) ]],
  ushort3 threadgroup_position_in_grid    [[ threadgroup_position_in_grid ]],
  ushort3 thread_position_in_threadgroup  [[ thread_position_in_threadgroup ]],
  ushort3 threads_per_threadgroup         [[ threads_per_threadgroup ]],
  ushort3 thread_position_in_grid         [[ thread_position_in_grid ]])
  {
  
      const int SrcSlices = (int)SrcShape[0];
      const int SrcHeight = (int)SrcShape[1];
      const int SrcWidth  = (int)SrcShape[2];
      const int DstSlices = (int)DstShape[0];
      const int DstHeight = (int)DstShape[1];
      const int DstWidth  = (int)DstShape[2];
      const int Kernel_X = 3;
      const int KernelElemNum = 3 * 3;
      const int N_Pack = 8;
  
     // Test only 1 thread
     if(thread_position_in_grid.z != 0|| thread_position_in_grid.y != 0|| thread_position_in_grid.x * N_Pack != 0) return;
  
      simdgroup_half8x8 sgMatY;
      simdgroup_load(sgMatY, Src);
  
      simdgroup_store(sgMatY, Dst);
  
  }
  </code></pre>
  <p>It's a simple shader, <strong>however output buffer only save first 2 values from input buffer，the other 62 values are ALL ZERO.</strong></p>
  <p>Here is the result from Xcode Metal Capture</p>
  <p><img src="https://i.stack.imgur.com/fI1pU.png" alt="" /><img src="https://i.stack.imgur.com/nhDlZ.png" alt="" /></p>
  <p>How to debug or fix it？</p>