title: Finding the minimum and maximum value within a Metal texture
body: <p>I have a <code>MTLTexture</code> containing 16bit unsigned integers (<code>MTLPixelFormatR16Uint</code>). The values range from about 7000 to 20000, with 0 being used as a 'nodata' value, which is why it is skipped in the code below. I'd like to find the minimum and maximum values so I can rescale these values between 0-255. Ultimately I'll be looking to base the minimum and maximum values on a histogram of the data (it has some outliers), but for now I'm stuck on simply extracting the min/max.</p>
  
  <p><em>I can read the data from the GPU to CPU and pull the min/max values out but would prefer to perform this task on the GPU.</em></p>
  
  <p><strong>First attempt</strong></p>
  
  <p>The command encoder is dispatched with 16x16 threads per thread group, the number of thread groups is based on the texture size (eg; width = textureWidth / 16, height = textureHeight / 16).</p>
  
  <pre><code>typedef struct {
      atomic_uint min;
      atomic_uint max;
  } BandMinMax;
  
  kernel void minMax(texture2d&lt;ushort, access::read&gt; band1 [[texture(0)]],
                     device BandMinMax &amp;out [[buffer(0)]],
                     uint2 gid [[thread_position_in_grid]])
  {
      ushort value = band1.read(gid).r;
  
      if (value != 0) {
          uint currentMin = atomic_load_explicit(&amp;out.min, memory_order_relaxed);
          uint currentMax = atomic_load_explicit(&amp;out.max, memory_order_relaxed);
  
          if (value &gt; currentMax) {
              atomic_store_explicit(&amp;out.max, value, memory_order_relaxed);
          }
          if (value &lt; currentMin) {
              atomic_store_explicit(&amp;out.min, value, memory_order_relaxed);
          }
      }
  }
  </code></pre>
  
  <p>From this I get a minimum and maximum value, but for the same dataset the min and max will often return different values. Fairly certain this is the min and max from a single thread when there are multiple threads running.</p>
  
  <p><strong>Second attempt</strong></p>
  
  <p>Building on the previous attempt, this time I'm storing the individual min/max values from each thread, all 256 (16x16).</p>
  
  <pre><code>kernel void minMax(texture2d&lt;ushort, access::read&gt; band1 [[texture(0)]],
                     device BandMinMax *out [[buffer(0)]],
                     uint2 gid [[thread_position_in_grid]],
                     uint tid [[ thread_index_in_threadgroup ]])
  {
  
      ushort value = band1.read(gid).r;
  
      if (value != 0) {
          uint currentMin = atomic_load_explicit(&amp;out[tid].min, memory_order_relaxed);
          uint currentMax = atomic_load_explicit(&amp;out[tid].max, memory_order_relaxed);
  
          if (value &gt; currentMax) {
              atomic_store_explicit(&amp;out[tid].max, value, memory_order_relaxed);
          }
          if (value &lt; currentMin) {
              atomic_store_explicit(&amp;out[tid].min, value, memory_order_relaxed);
          }
      }
  }
  </code></pre>
  
  <p>This returns an array containing 256 sets of min/max values. From these I guess I could find the lowest of the minimum values, but this seems like a poor approach. Would appreciate a pointer in the right direction, thanks!</p>