//
//  FMMetalBrushKernels.metal
//  acorn
//

#include <metal_stdlib>
using namespace metal;

struct FMMetalBrushDab {
    float2 center;
    float2 size;
    float sinAngle;
    float cosAngle;
    float opacity;
    float padding;
    float2 sourceOffset;
};

struct FMMetalBrushBatchUniforms {
    uint2 origin;
    uint dabCount;
    uint colorStorageMode;
    float4 color;
};

enum FMMetalBrushColorStorageMode {
    FMMetalBrushColorStorageModeSRGB8 = 0,
    FMMetalBrushColorStorageModeLinearFloat = 1,
};

static inline float FMMetalBrushSRGBToLinearComponent(float c) {
    return (c <= 0.04045f) ? (c / 12.92f) : powr((c + 0.055f) / 1.055f, 2.4f);
}

static inline float FMMetalBrushLinearToSRGBComponent(float c) {
    return (c <= 0.0031308f) ? (c * 12.92f) : ((1.055f * powr(c, 1.0f / 2.4f)) - 0.055f);
}

static inline float3 FMMetalBrushSRGBToLinear(float3 color) {
    return float3(FMMetalBrushSRGBToLinearComponent(color.r),
                  FMMetalBrushSRGBToLinearComponent(color.g),
                  FMMetalBrushSRGBToLinearComponent(color.b));
}

static inline float3 FMMetalBrushLinearToSRGB(float3 color) {
    return float3(FMMetalBrushLinearToSRGBComponent(color.r),
                  FMMetalBrushLinearToSRGBComponent(color.g),
                  FMMetalBrushLinearToSRGBComponent(color.b));
}

static inline float3 FMMetalBrushPremultipliedSRGBToPremultipliedLinear(float4 color) {
    if (color.a <= 0.00001f) {
        return float3(0.0f);
    }
    float invAlpha = 1.0f / color.a;
    float3 straightSRGB = clamp(color.rgb * invAlpha, 0.0f, 1.0f);
    float3 straightLinear = FMMetalBrushSRGBToLinear(straightSRGB);
    return straightLinear * color.a;
}

static inline float4 FMMetalBrushPremultipliedLinearToPremultipliedSRGB(float3 rgb, float alpha) {
    if (alpha <= 0.00001f) {
        return float4(0.0f);
    }
    float invAlpha = 1.0f / alpha;
    float3 straightLinear = clamp(rgb * invAlpha, 0.0f, 1.0f);
    float3 straightSRGB = FMMetalBrushLinearToSRGB(straightLinear);
    return float4(clamp(straightSRGB * alpha, 0.0f, 1.0f), clamp(alpha, 0.0f, 1.0f));
}

static inline float4 FMMetalBrushTextureSampleToWorkingColor(float4 color, uint colorStorageMode) {
    if (colorStorageMode == FMMetalBrushColorStorageModeLinearFloat) {
        return color;
    }
    return float4(FMMetalBrushPremultipliedSRGBToPremultipliedLinear(color), color.a);
}

static inline float4 FMMetalBrushWorkingColorToTextureSample(float4 color, uint colorStorageMode) {
    if (colorStorageMode == FMMetalBrushColorStorageModeLinearFloat) {
        return color;
    }
    return FMMetalBrushPremultipliedLinearToPremultipliedSRGB(color.rgb, color.a);
}

static inline float3 FMMetalBrushSourceColorToWorkingColor(float3 color, uint colorStorageMode) {
    if (colorStorageMode == FMMetalBrushColorStorageModeLinearFloat) {
        return clamp(color, 0.0f, 1.0f);
    }
    return FMMetalBrushSRGBToLinear(clamp(color, 0.0f, 1.0f));
}

static inline float3 FMMetalBrushStraightRGB(float4 color) {
    if (color.a <= 0.00001f) {
        return float3(0.0f);
    }
    return clamp(color.rgb / color.a, 0.0f, 1.0f);
}

static inline float FMMetalBrushBlendColorDodge(float d, float s) {
    return (s >= 1.0f) ? 1.0f : min(1.0f, d / max(1.0f - s, 0.00001f));
}

static inline float FMMetalBrushBlendColorBurn(float d, float s) {
    return (s <= 0.0f) ? 0.0f : 1.0f - min(1.0f, (1.0f - d) / max(s, 0.00001f));
}

static inline float FMMetalBrushBlendSoftLight(float d, float s) {
    float g = (d <= 0.25f) ? (((16.0f * d - 12.0f) * d + 4.0f) * d) : sqrt(d);
    return (s <= 0.5f) ? (d - ((1.0f - 2.0f * s) * d * (1.0f - d))) : (d + ((2.0f * s - 1.0f) * (g - d)));
}

static inline float FMMetalBrushLuminosity(float3 color) {
    return dot(color, float3(0.3f, 0.59f, 0.11f));
}

static inline float FMMetalBrushSaturation(float3 color) {
    return max(color.r, max(color.g, color.b)) - min(color.r, min(color.g, color.b));
}

static inline float3 FMMetalBrushClipColor(float3 color) {
    float l = FMMetalBrushLuminosity(color);
    float n = min(color.r, min(color.g, color.b));
    float x = max(color.r, max(color.g, color.b));
    if (n < 0.0f) {
        color = l + (((color - l) * l) / max(l - n, 0.00001f));
    }
    if (x > 1.0f) {
        color = l + (((color - l) * (1.0f - l)) / max(x - l, 0.00001f));
    }
    return clamp(color, 0.0f, 1.0f);
}

static inline float3 FMMetalBrushSetLuminosity(float3 color, float luminosity) {
    return FMMetalBrushClipColor(color + (luminosity - FMMetalBrushLuminosity(color)));
}

static inline float3 FMMetalBrushSetSaturation(float3 color, float saturation) {
    float r = color.r;
    float g = color.g;
    float b = color.b;
    float minValue = min(r, min(g, b));
    float maxValue = max(r, max(g, b));
    float range = maxValue - minValue;
    if (range <= 0.00001f) {
        return float3(0.0f);
    }
    float3 result = (color - minValue) * (saturation / range);
    return clamp(result, 0.0f, saturation);
}

static inline float4 FMMetalBrushCompositeSeparable(float4 D, float4 S, float3 blendColor) {
    float alpha = S.a + D.a - (S.a * D.a);
    float3 rgb = (S.rgb * (1.0f - D.a)) + (D.rgb * (1.0f - S.a)) + (clamp(blendColor, 0.0f, 1.0f) * S.a * D.a);
    return float4(clamp(rgb, 0.0f, alpha), clamp(alpha, 0.0f, 1.0f));
}

static inline float4 FMMetalBrushCompositeNonSeparable(float4 D, float4 S, float3 blendColor) {
    return FMMetalBrushCompositeSeparable(D, S, blendColor);
}

static inline float4 FMMetalBrushNormalBlend(float4 D, float4 S) {
    return S + (D * (1.0f - S.a));
}

static inline float4 FMMetalBrushMultiplyBlend(float4 D, float4 S) {
    float3 d = FMMetalBrushStraightRGB(D);
    float3 s = FMMetalBrushStraightRGB(S);
    return FMMetalBrushCompositeSeparable(D, S, d * s);
}

static inline float4 FMMetalBrushScreenBlend(float4 D, float4 S) {
    float3 d = FMMetalBrushStraightRGB(D);
    float3 s = FMMetalBrushStraightRGB(S);
    return FMMetalBrushCompositeSeparable(D, S, d + s - (d * s));
}

static inline float4 FMMetalBrushOverlayBlend(float4 D, float4 S) {
    float3 d = FMMetalBrushStraightRGB(D);
    float3 s = FMMetalBrushStraightRGB(S);
    float3 blend = select(1.0f - (2.0f * (1.0f - d) * (1.0f - s)), 2.0f * d * s, d <= 0.5f);
    return FMMetalBrushCompositeSeparable(D, S, blend);
}

static inline float4 FMMetalBrushDarkenBlend(float4 D, float4 S) {
    return FMMetalBrushCompositeSeparable(D, S, min(FMMetalBrushStraightRGB(D), FMMetalBrushStraightRGB(S)));
}

static inline float4 FMMetalBrushLightenBlend(float4 D, float4 S) {
    return FMMetalBrushCompositeSeparable(D, S, max(FMMetalBrushStraightRGB(D), FMMetalBrushStraightRGB(S)));
}

static inline float4 FMMetalBrushColorDodgeBlend(float4 D, float4 S) {
    float3 d = FMMetalBrushStraightRGB(D);
    float3 s = FMMetalBrushStraightRGB(S);
    float3 blend = float3(FMMetalBrushBlendColorDodge(d.r, s.r),
                          FMMetalBrushBlendColorDodge(d.g, s.g),
                          FMMetalBrushBlendColorDodge(d.b, s.b));
    return FMMetalBrushCompositeSeparable(D, S, blend);
}

static inline float4 FMMetalBrushColorBurnBlend(float4 D, float4 S) {
    float3 d = FMMetalBrushStraightRGB(D);
    float3 s = FMMetalBrushStraightRGB(S);
    float3 blend = float3(FMMetalBrushBlendColorBurn(d.r, s.r),
                          FMMetalBrushBlendColorBurn(d.g, s.g),
                          FMMetalBrushBlendColorBurn(d.b, s.b));
    return FMMetalBrushCompositeSeparable(D, S, blend);
}

static inline float4 FMMetalBrushSoftLightBlend(float4 D, float4 S) {
    float3 d = FMMetalBrushStraightRGB(D);
    float3 s = FMMetalBrushStraightRGB(S);
    float3 blend = float3(FMMetalBrushBlendSoftLight(d.r, s.r),
                          FMMetalBrushBlendSoftLight(d.g, s.g),
                          FMMetalBrushBlendSoftLight(d.b, s.b));
    return FMMetalBrushCompositeSeparable(D, S, blend);
}

static inline float4 FMMetalBrushHardLightBlend(float4 D, float4 S) {
    float3 d = FMMetalBrushStraightRGB(D);
    float3 s = FMMetalBrushStraightRGB(S);
    float3 blend = select(1.0f - (2.0f * (1.0f - d) * (1.0f - s)), 2.0f * d * s, s <= 0.5f);
    return FMMetalBrushCompositeSeparable(D, S, blend);
}

static inline float4 FMMetalBrushDifferenceBlend(float4 D, float4 S) {
    return FMMetalBrushCompositeSeparable(D, S, abs(FMMetalBrushStraightRGB(D) - FMMetalBrushStraightRGB(S)));
}

static inline float4 FMMetalBrushExclusionBlend(float4 D, float4 S) {
    float3 d = FMMetalBrushStraightRGB(D);
    float3 s = FMMetalBrushStraightRGB(S);
    return FMMetalBrushCompositeSeparable(D, S, d + s - (2.0f * d * s));
}

static inline float4 FMMetalBrushHueBlend(float4 D, float4 S) {
    float3 d = FMMetalBrushStraightRGB(D);
    float3 s = FMMetalBrushStraightRGB(S);
    float3 blend = FMMetalBrushSetLuminosity(FMMetalBrushSetSaturation(s, FMMetalBrushSaturation(d)), FMMetalBrushLuminosity(d));
    return FMMetalBrushCompositeNonSeparable(D, S, blend);
}

static inline float4 FMMetalBrushSaturationBlend(float4 D, float4 S) {
    float3 d = FMMetalBrushStraightRGB(D);
    float3 s = FMMetalBrushStraightRGB(S);
    float3 blend = FMMetalBrushSetLuminosity(FMMetalBrushSetSaturation(d, FMMetalBrushSaturation(s)), FMMetalBrushLuminosity(d));
    return FMMetalBrushCompositeNonSeparable(D, S, blend);
}

static inline float4 FMMetalBrushColorBlend(float4 D, float4 S) {
    return FMMetalBrushCompositeNonSeparable(D, S, FMMetalBrushSetLuminosity(FMMetalBrushStraightRGB(S), FMMetalBrushLuminosity(FMMetalBrushStraightRGB(D))));
}

static inline float4 FMMetalBrushLuminosityBlend(float4 D, float4 S) {
    return FMMetalBrushCompositeNonSeparable(D, S, FMMetalBrushSetLuminosity(FMMetalBrushStraightRGB(D), FMMetalBrushLuminosity(FMMetalBrushStraightRGB(S))));
}

static inline float4 FMMetalBrushClearBlend(float4 D, float4 S) {
    return float4(0.0f);
}

static inline float4 FMMetalBrushCopyBlend(float4 D, float4 S) {
    return S;
}

static inline float4 FMMetalBrushSourceInBlend(float4 D, float4 S) {
    return S * D.a;
}

static inline float4 FMMetalBrushSourceOutBlend(float4 D, float4 S) {
    return S * (1.0f - D.a);
}

static inline float4 FMMetalBrushSourceAtopBlend(float4 D, float4 S) {
    return (S * D.a) + (D * (1.0f - S.a));
}

static inline float4 FMMetalBrushDestinationOverBlend(float4 D, float4 S) {
    return (S * (1.0f - D.a)) + D;
}

static inline float4 FMMetalBrushDestinationInBlend(float4 D, float4 S) {
    return D * S.a;
}

static inline float4 FMMetalBrushDestinationOutBlend(float4 D, float4 S) {
    return D * (1.0f - S.a);
}

static inline float4 FMMetalBrushDestinationAtopBlend(float4 D, float4 S) {
    return (S * (1.0f - D.a)) + (D * S.a);
}

static inline float4 FMMetalBrushXORBlend(float4 D, float4 S) {
    return (S * (1.0f - D.a)) + (D * (1.0f - S.a));
}

static inline float4 FMMetalBrushPlusDarkerBlend(float4 D, float4 S) {
    float3 blend = clamp(FMMetalBrushStraightRGB(D) + FMMetalBrushStraightRGB(S) - 1.0f, 0.0f, 1.0f);
    return FMMetalBrushCompositeSeparable(D, S, blend);
}

static inline float4 FMMetalBrushPlusLighterBlend(float4 D, float4 S) {
    float3 blend = clamp(FMMetalBrushStraightRGB(D) + FMMetalBrushStraightRGB(S), 0.0f, 1.0f);
    return FMMetalBrushCompositeSeparable(D, S, blend);
}

static inline float4 FMMetalBrushDivideBlend(float4 D, float4 S) {
    return FMMetalBrushCompositeSeparable(D, S, clamp(FMMetalBrushStraightRGB(D) / max(FMMetalBrushStraightRGB(S), float3(0.00001f)), 0.0f, 1.0f));
}

static inline float4 FMMetalBrushLighterColorBlend(float4 D, float4 S) {
    float3 d = FMMetalBrushStraightRGB(D);
    float3 s = FMMetalBrushStraightRGB(S);
    return FMMetalBrushCompositeSeparable(D, S, (FMMetalBrushLuminosity(s) > FMMetalBrushLuminosity(d)) ? s : d);
}

static inline float4 FMMetalBrushDarkerColorBlend(float4 D, float4 S) {
    float3 d = FMMetalBrushStraightRGB(D);
    float3 s = FMMetalBrushStraightRGB(S);
    return FMMetalBrushCompositeSeparable(D, S, (FMMetalBrushLuminosity(s) < FMMetalBrushLuminosity(d)) ? s : d);
}

static inline float4 FMMetalBrushSubtractBlend(float4 D, float4 S) {
    return FMMetalBrushCompositeSeparable(D, S, clamp(FMMetalBrushStraightRGB(D) - FMMetalBrushStraightRGB(S), 0.0f, 1.0f));
}

static inline float4 FMMetalBrushPinLightBlend(float4 D, float4 S) {
    float3 d = FMMetalBrushStraightRGB(D);
    float3 s = FMMetalBrushStraightRGB(S);
    float3 blend = select(max(d, (2.0f * s) - 1.0f), min(d, 2.0f * s), s <= 0.5f);
    return FMMetalBrushCompositeSeparable(D, S, blend);
}

static inline float4 FMMetalBrushVividLightBlend(float4 D, float4 S) {
    float3 d = FMMetalBrushStraightRGB(D);
    float3 s = FMMetalBrushStraightRGB(S);
    float3 burn = float3(FMMetalBrushBlendColorBurn(d.r, 2.0f * s.r),
                         FMMetalBrushBlendColorBurn(d.g, 2.0f * s.g),
                         FMMetalBrushBlendColorBurn(d.b, 2.0f * s.b));
    float3 dodge = float3(FMMetalBrushBlendColorDodge(d.r, (2.0f * s.r) - 1.0f),
                          FMMetalBrushBlendColorDodge(d.g, (2.0f * s.g) - 1.0f),
                          FMMetalBrushBlendColorDodge(d.b, (2.0f * s.b) - 1.0f));
    return FMMetalBrushCompositeSeparable(D, S, select(dodge, burn, s <= 0.5f));
}

static inline float4 FMMetalBrushLinearLightBlend(float4 D, float4 S) {
    return FMMetalBrushCompositeSeparable(D, S, clamp(FMMetalBrushStraightRGB(D) + (2.0f * FMMetalBrushStraightRGB(S)) - 1.0f, 0.0f, 1.0f));
}

static inline float4 FMMetalBrushHardMixBlend(float4 D, float4 S) {
    float3 vivid = FMMetalBrushStraightRGB(FMMetalBrushVividLightBlend(float4(FMMetalBrushStraightRGB(D), 1.0f), float4(FMMetalBrushStraightRGB(S), 1.0f)));
    return FMMetalBrushCompositeSeparable(D, S, select(float3(1.0f), float3(0.0f), vivid < 0.5f));
}

static inline float4 FMMetalBrushMaskBlend(float4 D, float4 S) {
    return D * clamp(S.g, 0.0f, 1.0f);
}

static inline float4 FMMetalBrushCopyOpaqueOnlyBlend(float4 D, float4 S) {
    return (S.a > 0.0f) ? S : D;
}

static inline float4 FMMetalBrushReadSourceTexture(texture2d<float, access::read> sourceTexture, int2 coord) {
    if (coord.x < 0 || coord.y < 0 || coord.x >= int(sourceTexture.get_width()) || coord.y >= int(sourceTexture.get_height())) {
        return float4(0.0f);
    }
    return sourceTexture.read(uint2(coord));
}

static inline float4 FMMetalBrushBilinearSourceSample(texture2d<float, access::read> sourceTexture, float2 pixelCoord) {
    float2 base = floor(pixelCoord);
    float2 fraction = pixelCoord - base;
    int2 p00 = int2(base);
    int2 p10 = p00 + int2(1, 0);
    int2 p01 = p00 + int2(0, 1);
    int2 p11 = p00 + int2(1, 1);
    float4 c00 = FMMetalBrushReadSourceTexture(sourceTexture, p00);
    float4 c10 = FMMetalBrushReadSourceTexture(sourceTexture, p10);
    float4 c01 = FMMetalBrushReadSourceTexture(sourceTexture, p01);
    float4 c11 = FMMetalBrushReadSourceTexture(sourceTexture, p11);
    return mix(mix(c00, c10, fraction.x), mix(c01, c11, fraction.x), fraction.y);
}

#define FM_METAL_BRUSH_KERNEL(kernelName, blendFunction) \
kernel void kernelName(texture2d<float, access::read_write> destinationTexture [[texture(0)]], \
                       texture2d<float, access::sample> maskTexture [[texture(1)]], \
                       constant FMMetalBrushBatchUniforms &uniforms [[buffer(0)]], \
                       constant FMMetalBrushDab *dabs [[buffer(1)]], \
                       uint2 gid [[thread_position_in_grid]]) { \
    uint2 pixel = uniforms.origin + gid; \
    if (pixel.x >= destinationTexture.get_width() || pixel.y >= destinationTexture.get_height()) { \
        return; \
    } \
    float4 destinationSample = destinationTexture.read(pixel); \
    float4 dst = FMMetalBrushTextureSampleToWorkingColor(destinationSample, uniforms.colorStorageMode); \
    float3 srcColorLinear = FMMetalBrushSourceColorToWorkingColor(uniforms.color.rgb, uniforms.colorStorageMode); \
    constexpr sampler maskSampler(coord::normalized, address::clamp_to_zero, filter::linear); \
    float2 pixelCenter = float2(pixel) + 0.5f; \
    for (uint i = 0; i < uniforms.dabCount; i++) { \
        FMMetalBrushDab dab = dabs[i]; \
        if (dab.opacity <= 0.0f || dab.size.x <= 0.0f || dab.size.y <= 0.0f) { \
            continue; \
        } \
        float2 relative = pixelCenter - dab.center; \
        float2 local = float2(relative.x * dab.cosAngle + relative.y * dab.sinAngle, \
                              -relative.x * dab.sinAngle + relative.y * dab.cosAngle); \
        float2 uv = (local / dab.size) + 0.5f; \
        if (uv.x < 0.0f || uv.x > 1.0f || uv.y < 0.0f || uv.y > 1.0f) { \
            continue; \
        } \
        float coverage = maskTexture.sample(maskSampler, uv).r; \
        float srcAlpha = clamp(uniforms.color.a * dab.opacity * coverage, 0.0f, 1.0f); \
        if (srcAlpha <= 0.00001f) { \
            continue; \
        } \
        float4 src = float4(srcColorLinear * srcAlpha, srcAlpha); \
        dst = clamp(blendFunction(dst, src), 0.0f, 1.0f); \
    } \
    destinationTexture.write(FMMetalBrushWorkingColorToTextureSample(dst, uniforms.colorStorageMode), pixel); \
}

FM_METAL_BRUSH_KERNEL(FMMetalBrushNormalBlendKernel, FMMetalBrushNormalBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushMultiplyBlendKernel, FMMetalBrushMultiplyBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushScreenBlendKernel, FMMetalBrushScreenBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushOverlayBlendKernel, FMMetalBrushOverlayBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushDarkenBlendKernel, FMMetalBrushDarkenBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushLightenBlendKernel, FMMetalBrushLightenBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushColorDodgeBlendKernel, FMMetalBrushColorDodgeBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushColorBurnBlendKernel, FMMetalBrushColorBurnBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushSoftLightBlendKernel, FMMetalBrushSoftLightBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushHardLightBlendKernel, FMMetalBrushHardLightBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushDifferenceBlendKernel, FMMetalBrushDifferenceBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushExclusionBlendKernel, FMMetalBrushExclusionBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushHueBlendKernel, FMMetalBrushHueBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushSaturationBlendKernel, FMMetalBrushSaturationBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushColorBlendKernel, FMMetalBrushColorBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushLuminosityBlendKernel, FMMetalBrushLuminosityBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushClearBlendKernel, FMMetalBrushClearBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushCopyBlendKernel, FMMetalBrushCopyBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushSourceInBlendKernel, FMMetalBrushSourceInBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushSourceOutBlendKernel, FMMetalBrushSourceOutBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushSourceAtopBlendKernel, FMMetalBrushSourceAtopBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushDestinationOverBlendKernel, FMMetalBrushDestinationOverBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushDestinationInBlendKernel, FMMetalBrushDestinationInBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushDestinationOutBlendKernel, FMMetalBrushDestinationOutBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushDestinationAtopBlendKernel, FMMetalBrushDestinationAtopBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushXORBlendKernel, FMMetalBrushXORBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushPlusDarkerBlendKernel, FMMetalBrushPlusDarkerBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushPlusLighterBlendKernel, FMMetalBrushPlusLighterBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushDivideBlendKernel, FMMetalBrushDivideBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushLighterColorBlendKernel, FMMetalBrushLighterColorBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushDarkerColorBlendKernel, FMMetalBrushDarkerColorBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushSubtractBlendKernel, FMMetalBrushSubtractBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushHardMixBlendKernel, FMMetalBrushHardMixBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushMaskBlendKernel, FMMetalBrushMaskBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushPinLightBlendKernel, FMMetalBrushPinLightBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushVividLightBlendKernel, FMMetalBrushVividLightBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushLinearLightBlendKernel, FMMetalBrushLinearLightBlend)
FM_METAL_BRUSH_KERNEL(FMMetalBrushCopyOpaqueOnlyBlendKernel, FMMetalBrushCopyOpaqueOnlyBlend)

kernel void FMMetalBrushSmudgeKernel(texture2d<float, access::read_write> destinationTexture [[texture(0)]],
                                     texture2d<float, access::sample> maskTexture [[texture(1)]],
                                     texture2d<float, access::read> sourceTexture [[texture(2)]],
                                     constant FMMetalBrushBatchUniforms &uniforms [[buffer(0)]],
                                     constant FMMetalBrushDab *dabs [[buffer(1)]],
                                     uint2 gid [[thread_position_in_grid]]) {
    uint2 pixel = uniforms.origin + gid;
    if (pixel.x >= destinationTexture.get_width() || pixel.y >= destinationTexture.get_height()) {
        return;
    }

    float4 dst = FMMetalBrushTextureSampleToWorkingColor(destinationTexture.read(pixel), uniforms.colorStorageMode);
    constexpr sampler maskSampler(coord::normalized, address::clamp_to_zero, filter::linear);
    float2 pixelCenter = float2(pixel) + 0.5f;

    for (uint i = 0; i < uniforms.dabCount; i++) {
        FMMetalBrushDab dab = dabs[i];
        if (dab.opacity <= 0.0f || dab.size.x <= 0.0f || dab.size.y <= 0.0f) {
            continue;
        }

        float2 relative = pixelCenter - dab.center;
        float2 local = float2(relative.x * dab.cosAngle + relative.y * dab.sinAngle,
                              -relative.x * dab.sinAngle + relative.y * dab.cosAngle);
        float2 uv = (local / dab.size) + 0.5f;
        if (uv.x < 0.0f || uv.x > 1.0f || uv.y < 0.0f || uv.y > 1.0f) {
            continue;
        }

        float coverage = clamp(maskTexture.sample(maskSampler, uv).r * dab.opacity, 0.0f, 1.0f);
        if (coverage <= 0.00001f) {
            continue;
        }

        float2 sourcePixelCoord = (pixelCenter - dab.sourceOffset) - 0.5f;
        float4 sourceSample = FMMetalBrushTextureSampleToWorkingColor(FMMetalBrushBilinearSourceSample(sourceTexture, sourcePixelCoord), uniforms.colorStorageMode);
        dst = mix(dst, sourceSample, coverage);
    }

    destinationTexture.write(FMMetalBrushWorkingColorToTextureSample(dst, uniforms.colorStorageMode), pixel);
}
