Ricerca del valore minimo e massimo all'interno di una trama metallica

Question

Ho uno MTLTexture contenente numeri interi senza segno a 16 bit (MTLPixelFormatR16Uint). I valori vanno da circa 7000 a 20000, con 0 usato come valore 'nodata', motivo per cui viene saltato nel codice sottostante. Mi piacerebbe trovare i valori minimi e massimi in modo da poter ridimensionare questi valori tra 0-255. In definitiva cercherò di basare i valori minimi e massimi su un istogramma dei dati (ha alcuni valori anomali), ma per ora sono bloccato semplicemente estraendo il minimo/massimo.

Posso leggere i dati dalla GPU alla CPU e tirare fuori i valori min/max ma preferirei eseguire questa operazione sulla GPU.

Primo tentativo

L'encoder comando viene inviato con 16x16 filetti per gruppo filo, il numero di gruppi di thread si basa sulla dimensione trama (per esempio, larghezza = textureWidth/16, altezza = textureHeight/16).

typedef struct { 
    atomic_uint min; 
    atomic_uint max; 
} BandMinMax; 

kernel void minMax(texture2d<ushort, access::read> band1 [[texture(0)]], 
        device BandMinMax &out [[buffer(0)]], 
        uint2 gid [[thread_position_in_grid]]) 
{ 
    ushort value = band1.read(gid).r; 

    if (value != 0) { 
     uint currentMin = atomic_load_explicit(&out.min, memory_order_relaxed); 
     uint currentMax = atomic_load_explicit(&out.max, memory_order_relaxed); 

     if (value > currentMax) { 
      atomic_store_explicit(&out.max, value, memory_order_relaxed); 
     } 
     if (value < currentMin) { 
      atomic_store_explicit(&out.min, value, memory_order_relaxed); 
     } 
    } 
} 

Da questo ottengo un valore minimo e massimo, ma per lo stesso insieme di dati il ​​min e max sarà spesso restituire valori diversi. Abbastanza sicuro questo è il minimo e il massimo da un singolo thread quando ci sono più thread in esecuzione.

Secondo tentativo

Basandosi sul tentativo precedente, questa volta sto memorizzando i singoli valori min/max da ogni thread, tutti 256 (16x16).

kernel void minMax(texture2d<ushort, access::read> 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(&out[tid].min, memory_order_relaxed); 
     uint currentMax = atomic_load_explicit(&out[tid].max, memory_order_relaxed); 

     if (value > currentMax) { 
      atomic_store_explicit(&out[tid].max, value, memory_order_relaxed); 
     } 
     if (value < currentMin) { 
      atomic_store_explicit(&out[tid].min, value, memory_order_relaxed); 
     } 
    } 
} 

Questo restituisce un array contenente 256 serie di valori min/max. Da questi suppongo di aver trovato il minimo dei valori minimi, ma questo mi sembra un approccio inadeguato. Gradirei un puntatore nella giusta direzione, grazie!

Answer 1

Il Metal Shading Language ha funzioni di confronto e scambio atomiche che è possibile utilizzare per confrontare il valore esistente in una posizione di memoria con un valore e sostituire il valore in quella posizione se non sono uguali. Con questi, è possibile creare un set di atomica confronto-and-replace-se- [maggiore] | meno operazioni -than:

static void atomic_uint_exchange_if_less_than(volatile device atomic_uint *current, uint candidate) 
{ 
    uint val; 
    do { 
     val = *((device uint *)current); 
    } while ((candidate < val || val == 0) && !atomic_compare_exchange_weak_explicit(current, 
                        &val, 
                        candidate, 
                        memory_order_relaxed, 
                        memory_order_relaxed)); 
} 

static void atomic_uint_exchange_if_greater_than(volatile device atomic_uint *current, uint candidate) 
{ 
    uint val; 
    do { 
     val = *((device uint *)current); 
    } while (candidate > val && !atomic_compare_exchange_weak_explicit(current, 
                     &val, 
                     candidate, 
                     memory_order_relaxed, 
                     memory_order_relaxed)); 
} 

Per applicare questi, si potrebbe creare un buffer che contiene un min interlacciata, max coppia per threadgroup. Poi, in funzione del kernel, leggere la trama e condizionatamente scrivere i valori minimi e massimi:

kernel void min_max_per_threadgroup(texture2d<ushort, access::read> texture [[texture(0)]], 
            device uint *mapBuffer [[buffer(0)]], 
            uint2 tpig [[thread_position_in_grid]], 
            uint2 tgpig [[threadgroup_position_in_grid]], 
            uint2 tgpg [[threadgroups_per_grid]]) 
{ 
    ushort val = texture.read(tpig).r; 

    device atomic_uint *atomicBuffer = (device atomic_uint *)mapBuffer; 

    atomic_uint_exchange_if_less_than(atomicBuffer + ((tgpig[1] * tgpg[0] + tgpig[0]) * 2), 
             val); 

    atomic_uint_exchange_if_greater_than(atomicBuffer + ((tgpig[1] * tgpg[0] + tgpig[0]) * 2) + 1, 
             val); 
} 

Infine, eseguire un kernel separato per ridurre su questo buffer e raccogliere il minimo finale, valori massimi su tutta la trama :

kernel void min_max_reduce(constant uint *mapBuffer [[buffer(0)]], 
          device uint *reduceBuffer [[buffer(1)]], 
          uint2 tpig [[thread_position_in_grid]]) 
{ 
    uint minv = mapBuffer[tpig[0] * 2]; 
    uint maxv = mapBuffer[tpig[0] * 2 + 1]; 

    device atomic_uint *atomicBuffer = (device atomic_uint *)reduceBuffer; 

    atomic_uint_exchange_if_less_than(atomicBuffer, minv); 

    atomic_uint_exchange_if_greater_than(atomicBuffer + 1, maxv); 
} 

Naturalmente, è possibile ridurre solo su totale consentito la larghezza esecuzione filettatura del dispositivo (~ 256), quindi potrebbe essere necessario fare la riduzione in più passaggi, con ognuno ridurre le dimensioni del dati da utilizzare per un fattore della larghezza massima di esecuzione del filo.

Disclaimer: Questa potrebbe non essere la tecnica migliore, ma sembra essere corretta nei miei test limitati di un'implementazione di OS X. È stato leggermente più veloce di un'implementazione della CPU naive su una trama 256x256 su Intel Iris Pro, ma sostanzialmente più lento su una Nvidia GT 750M (a causa del dispatch overhead).