金属テクスチャ内の最小値と最大値を見つける

Question

私はMTLTextureに16ビット符号なし整数（MTLPixelFormatR16Uint）を含んでいます。値は約7000から20000の範囲で、0が 'nodata'値として使用されます。そのため、以下のコードでは省略されています。最小値と最大値を探して、0〜255の間でこれらの値を再スケーリングできるようにしたいと思います。最終的には、データのヒストグラムに最小値と最大値を設定することになりますが（これにはいくつかの異常値があります）、今は単純に最小値/最大値を抽出しています。

私はGPUからCPUにデータを読み込み、最小/最大値を引き出すことができますが、GPUでこのタスクを実行することをお勧めします。

コマンドエンコーダがスレッドグループ当たり16×16スレッドをディスパッチさ

最初の試みは、スレッドグループの数は、テクスチャサイズに基づいている（例えば、幅= textureWidth/16、高さ= 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); 
     } 
    } 
} 

これより最小値と最大値が得られますが、同じデータセットでは最小値と最大値が異なる値を返すことがよくあります。これは、複数のスレッドが実行されている場合、単一スレッドからの最小値と最大値です。

第二の試行前の試行で

ビル、私はそれぞれのスレッドから個々の最小/最大値を格納しています。この時、すべて256（16×16）。

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

これは、最小値/最大値の256セットを含む配列を返します。これらから、私は最小値の最低値を見つけることができたと思いますが、これは貧弱なアプローチのようです。正しい方向にポインタを感謝します、ありがとう！

Answer 1

金属シェーディング言語には、メモリ位置の既存の値と値を比較し、等しい場合はその位置の値を置き換えるために使用できるアトミックの比較およびスワップ関数があります。これらを使用すると、アトミック置き換える比較-と-IF-のセットを作成することができます[大きい|縮小] -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)); 
} 

これらを適用するには、あなたは、最大の1人のインターリーブ分が含まれているバッファを作成しますスレッドグループごとのペア。次に、カーネル関数では、テクスチャから読み取られ、条件付きで最小値と最大値を書き込む：テクスチャ全体にわたる

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

最後に、このバッファを介して削減し、最終分を収集するために別のカーネルを実行して、最大値：

もちろん

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

、あなただけのデバイス（〜256）の総許さスレッド実行幅にわたって減らすことができますので、あなたはのサイズを小さく各1と、複数のパスの削減を行う必要があるかもしれません最大スレッド実行幅のファクタで操作されるデータ。

免責事項：これは最善の方法ではないかもしれませんが、限られたOS X実装のテストでは正しいと思われます。 Intel Iris Proの256x256テクスチャ上の単純なCPUインプリメンテーションよりもわずかに速かったが、（ディスパッチオーバーヘッドのために）Nvidia GT 750Mでは大幅に遅い。