title: What exactly does thread_position_in_grid mean in a metal compute kernel shader?
body: <p>Although this is an inappropriate use of a compute shader I was doing some experiments to determine if I could use one to produce the general UV gradient where one channel of the image goes linearly from 0-1 across the x axis and the other channel goes from 0-1 across the y axis of the image. However I became confused when I generated this image by varying the b value of a texture by the thread_position_in_grid.x value divided by the image width. I edited the pixel of the texture at the thread_position_in_grid position:</p>
  
  <p><a href="https://i.stack.imgur.com/XZuAC.png" rel="nofollow noreferrer"><img src="https://i.stack.imgur.com/XZuAC.png" alt="enter image description here"></a></p>
  
  <p>Yes it was a gradient but it certainly did not appear to be the gradient from 0-1 that I wanted. I dropped it into an image editor and sure enough it was not linear. (The part added below shows what a linear gradient from 0-1 would look like)</p>
  
  <p><a href="https://i.stack.imgur.com/3tKde.png" rel="nofollow noreferrer"><img src="https://i.stack.imgur.com/3tKde.png" alt="enter image description here"></a></p>
  
  <p>It would appear that I do not understand what exactly the thread_position_in_grid value means. I know it has something to do with the threads per thread-groups and thread execution width but I dont exactly understand what. I suppose my end goal is to know whether it would be possible to generate the gradient below in a compute shader however I dont understand what is going on.</p>
  
  <p><a href="https://i.stack.imgur.com/CvQHU.png" rel="nofollow noreferrer"><img src="https://i.stack.imgur.com/CvQHU.png" alt="enter image description here"></a></p>
  
  <p>For reference I was working with a 100x100 texture with the following thread settings. Really I dont know why I use these values but this is what I saw recommended somewhere so I am sticking with them. I would love to be able to generalize this problem to other texture sizes as well including rectangles.</p>
  
  <pre><code>let w = greenPipeline.threadExecutionWidth
  let h = greenPipeline.maxTotalThreadsPerThreadgroup / w
  let threadsPerThreadgroup = MTLSizeMake(w, h, 1)
  let threadgroupsPerGrid = MTLSize(width: (texture.width + w - 1) / w,
                                    height: (texture.height + h - 1) / h,
                                    depth: 1)
  encoder.dispatchThreadgroups(threadgroupsPerGrid, threadsPerThreadgroup: threadsPerThreadgroup)
  </code></pre>
  
  <p>And my shader looks like this:</p>
  
  <pre><code>kernel void green(texture2d&lt;float, access::write&gt; outputTexture [[texture(0)]],
                               uint2 position [[thread_position_in_grid]])
  {
      if (position.x &gt;= outputTexture.get_width() || position.y &gt;= outputTexture.get_height()) {
          return;
      }
      outputTexture.write(float4(position.x / 100.0, 0.0, 0.0, 0.0), position);
  }
  </code></pre>
  
  <p>Two things about this shader confuse me because I cannot explain them:</p>
  
  <ol>
  <li>I am using position as the coordinate to write to on the texture so
  it bothers me that position doesnt work to generate the gradient.</li>
  <li>You cannot reaplace <code>position.x / 100.0</code> value with <code>position.x / outputTexture.getWidth()</code> even though it should also be 100. Doing so causes a black image. Yet when I made a shader that colored everything with outputTexture.getWidth() as its value it did indeed shade everything to a value equivalent to 100 (or more accurately 101 because of rounding)</li>
  <li>It is ok to use position to check if the kernel is within bounds but not to create the UV gradient.</li>
  </ol>
  
  <p>What is going on?</p>