Finding the minimum and maximum value within a Metal texture : <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>
