Bitonic sort

Kernels.metal

// ref :: http://t-pot.com/program/90_BitonicSort/index.html

#include <metal_stdlib>
using namespace metal;

kernel void bitonicSort(texture2d<float, access::read> inTexture [[texture(0)]],
                    texture2d<float, access::write> outTexture [[texture(1)]],
                    device float *stepno [[buffer(0)]],
                    device float *offset [[buffer(1)]],
                    device float *stage [[buffer(2)]],
                    uint2 gid [[thread_position_in_grid]])
{
    float4 dst;
    float w = inTexture.get_width();
    
    float2 elem2d =  float2(gid.x, gid.y);
    float elem1d = elem2d.x + elem2d.y * w;

    float csign = (fmod(elem1d, *stage) < *offset) ? 1.0 : -1.0;
    float cdir  = (fmod(floor(elem1d / *stepno), 2.0) <= 0.5) ? 1.0 : -1.0;

    float4 val0 = inTexture.read(gid);

    float adr1d = csign * *offset + elem1d;
    uint2 adr2d = uint2(fmod(adr1d, w), floor(adr1d / w));
    float4 val1 = inTexture.read(adr2d);
    
    float b0 = max(max(val0.r, val0.g), val0.b);
    float b1 = max(max(val1.r, val1.g), val1.b);
    
    float4 cmin = (b0<b1) ? val0: val1;
    float4 cmax = (b0<b1) ? val1: val0;
    
    dst = (csign==cdir) ? cmin : cmax;
    
	outTexture.write(dst.rgba, gid);
}