Den minimalen und maximalen Wert innerhalb einer Metall Textur finden

Question

Ich habe eine MTLTexture enthält 16bit vorzeichenlose Ganzzahlen (MTLPixelFormatR16Uint). Die Werte reichen von etwa 7000 bis 20000, wobei 0 als "nodata" -Wert verwendet wird, weshalb sie im folgenden Code übersprungen wird. Ich möchte die minimalen und maximalen Werte finden, damit ich diese Werte zwischen 0 und 255 skalieren kann. Letztendlich werde ich versuchen, die Minimal- und Maximalwerte auf einem Histogramm der Daten zu basieren (es hat einige Ausreißer), aber im Moment bleibe ich dran, einfach die Min/Max-Werte zu extrahieren.

Ich kann die Daten von der GPU zur CPU lesen und die Min/Max-Werte herausziehen, würde aber diese Aufgabe lieber auf der GPU durchführen.

Erster Versuch

Der Befehlsgeber mit 16x16 Fäden pro Thread-Gruppe versandt wird, wird die Anzahl der Fadengruppen auf der Grundlage der Texturgröße (zB; width = textureWidth/16, Höhe = textureHeight/16).

Daraus erhalte ich einen minimalen und maximalen Wert, aber für den gleichen Datensatz geben Min und Max oft unterschiedliche Werte zurück. Ziemlich sicher, das ist die Min und Max von einem einzelnen Thread, wenn mehrere Threads ausgeführt werden.

Zweiter Versuch

Aufbauend auf dem vorherigen Versuch, dieses Mal habe ich von jedem Thread die einzelnen Min-/Max-Werte bin speichern, werden alle 256 (16x16).

Dies gibt ein Array mit 256 Min/Max-Werten zurück. Von diesen könnte ich den niedrigsten der Mindestwerte finden, aber das scheint ein schlechter Ansatz zu sein. Würde einen Zeiger in die richtige Richtung schätzen, danke!

Answer 1

Die Metal Shading Language verfügt über atomare Compare-and-Swap-Funktionen, mit denen Sie den vorhandenen Wert an einem Speicherort mit einem Wert vergleichen und den Wert an diesem Speicherort ersetzen können, wenn sie nicht gleich sind. Mit diesen können Sie eine Reihe von Atom vergleichen und Ersetzen-wenn- [mehr | weniger] erstellen -than Operationen:

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)); 
} 

diese anwenden zu können, Sie könnten einen Puffer erstellen, die eine verschachtelte min enthält, max Paar pro Threadgruppe. Dann wurde in der Kernel-Funktion, lesen Sie von der Textur und bedingt die Min- und Max-Werte schreiben:

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); 
} 

schließlich einen separaten Kernel läuft diese Puffer zu reduzieren über und die letzte min, max Werte über die gesamte Textur sammeln :

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); 
} 

natürlich können Sie nur über die gesamte erlaubte Thread-Ausführungs Breite der Vorrichtung (~ 256) reduzieren, so dass Sie kann die Reduktion in mehreren Durchgängen, mit jedem tun müssen, um die Größe des Reduktions Daten, die mit einem Faktor der maximalen Thread-Ausführungsbreite bearbeitet werden sollen.

Haftungsausschluss: Dies ist möglicherweise nicht die beste Technik, aber es scheint in meinen begrenzten Tests einer OS X-Implementierung korrekt zu sein. Es war marginal schneller als eine naive CPU-Implementierung auf einer 256x256 Textur auf Intel Iris Pro, aber wesentlich langsamer auf einer Nvidia GT 750M (wegen des Versand-Overheads).