title: metal compute function limitations
body: <p>I experienced that <code>MTLBuffer</code>s with computionally intensive shader functions tend to stop calculating before all <code>threadgroups</code> are done. When I use a <code>MTLComputePipelineState</code> and <code>MTLComputeCommandEncoder</code> to blur an image with very large blur radii the resulting image half way processed and one can actually see half finished <code>threadgroups</code>. I did not narrow it down to the exact amount of blur radius, but 16 pixels works fine, 32 is already too much and not even half the groups are computed.</p>
  <p>So are there any limitations on how long a shader function call should take to finish or anything like that? I just finished most of the documentation about how to use the Metal framework and I cannot recall stumbling upon any such statements.</p>
  <h2>EDIT</h2>
  <p>Since in my case the problem was not a simple timeout but some internal error I'm going to add some code.</p>
  <p>The most expensive part of is the block-matching algorithm that finds matching blocks in two images (i.e consecutive frames in a movie)</p>
  <pre><code>//Exhaustive Search Block-matching algorithm
  kernel void naiveMotion(
      texture2d&lt;float,access::read&gt;   inputImage1   [[ texture(0) ]],
      texture2d&lt;float,access::read&gt;   inputImage2   [[ texture(1) ]],
      texture2d&lt;float,access::write&gt;  outputImage  [[ texture(2) ]],
  uint2 gid                                    [[ thread_position_in_grid ]]
  )
  {
      //area to search for matches
      float searchSize = 10.0;
      int searchRadius = searchSize/2;
  
      //window size to search in
      int kernelSize = 6;
      int kernelRadius = kernelSize/2;
  
      //this will store the motion direction
      float2 vector = float2(0.0,0.0);
      float2 maxVector = float2(searchSize,searchSize/2);
      float maxVectorLength = length(maxVector);
  
      //maximum error caused by noise
      float error = kernelSize*kernelSize*(10.0/255.0);
  
  
      for (int y = -searchRadius; y &lt; searchRadius; ++y)
      {
          for (int x = 0; x &lt; searchSize; ++x)
          {
              float diff = 0;
          
              for (int b = - kernelRadius; b &lt; kernelRadius; ++b)
              {
                  for (int a = - kernelRadius; a &lt; kernelRadius; ++a)
                  {
                      uint2 textureIndex(gid.x + x + a, gid.y + y + b);
                      float4 targetColor = inputImage2.read(textureIndex).rgba;
                      float4 referenceColor = inputImage1.read(gid).rgba;
                      float targetGray = 0.299*targetColor.r + 0.587*targetColor.g + 0.114*targetColor.b;
                      float referenceGray = 0.299*referenceColor.r + 0.587*referenceColor.g + 0.114*referenceColor.b;
                      diff = diff + abs(targetGray - referenceGray);
                  }
              }
          
              if ( error &gt; diff )
              {
                  error = diff;
                  //vertical motion is rather irrelevant but negative values can't be stored so just take the absolute value
                  vector = float2(x, abs(y));
              }
          }
      }
  
      float intensity = length(vector)/maxVectorLength;
      outputImage.write(float4(normalize(vector), intensity, 1),gid);
  }
  </code></pre>
  <p>I am using that shader on a 960x540px image. With a <code>searchSize</code> of 9 and <code>kernelSize</code> of 8 the shader runs over the whole image. Changing the searchSize to 10 and the shader will stop early with an error code 1.</p>