Source/ThirdParty/ANGLE/src/libANGLE/renderer/metal/shaders/blit.metal

//
// Copyright 2019 The ANGLE Project. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
// blit.metal: Implements blitting texture content to current frame buffer.

#include "common.h"

using namespace rx::mtl_shader;

// function_constant(0) is already used by common.h
constant bool kPremultiplyAlpha [[function_constant(1)]];
constant bool kUnmultiplyAlpha [[function_constant(2)]];
constant int kSourceTextureType [[function_constant(3)]];   // Source color/depth texture type.
constant int kSourceTexture2Type [[function_constant(4)]];  // Source stencil texture type.

constant bool kSourceTextureType2D      = kSourceTextureType == kTextureType2D;
constant bool kSourceTextureType2DArray = kSourceTextureType == kTextureType2DArray;
constant bool kSourceTextureType2DMS    = kSourceTextureType == kTextureType2DMultisample;
constant bool kSourceTextureTypeCube    = kSourceTextureType == kTextureTypeCube;
constant bool kSourceTextureType3D      = kSourceTextureType == kTextureType3D;

constant bool kSourceTexture2Type2D      = kSourceTexture2Type == kTextureType2D;
constant bool kSourceTexture2Type2DArray = kSourceTexture2Type == kTextureType2DArray;
constant bool kSourceTexture2Type2DMS    = kSourceTexture2Type == kTextureType2DMultisample;
constant bool kSourceTexture2TypeCube    = kSourceTexture2Type == kTextureTypeCube;

struct BlitParams
{
    // 0: lower left, 1: lower right, 2: upper left
    float2 srcTexCoords[3];
    int srcLevel;  // Source texture level.
    int srcLayer;  // Source texture layer.
    bool dstFlipViewportX;
    bool dstFlipViewportY;
    bool dstLuminance;  // destination texture is luminance. Unused by depth & stencil blitting.
    uint8_t padding[13];
};

struct BlitVSOut
{
    float4 position [[position]];
    float2 texCoords [[user(locn1)]];
};

vertex BlitVSOut blitVS(unsigned int vid [[vertex_id]], constant BlitParams &options [[buffer(0)]])
{
    BlitVSOut output;
    output.position  = float4(gCorners[vid], 0.0, 1.0);
    output.texCoords = options.srcTexCoords[vid];

    if (options.dstFlipViewportX)
    {
        output.position.x = -output.position.x;
    }
    if (!options.dstFlipViewportY)
    {
        // If viewport is not flipped, we have to flip Y in normalized device coordinates.
        // Since NDC has Y is opposite direction of viewport coodrinates.
        output.position.y = -output.position.y;
    }

    return output;
}

template <typename SrcTexture2d>
static uint2 getImageCoords(SrcTexture2d srcTexture, float2 texCoords)
{
    uint2 dimens(srcTexture.get_width(), srcTexture.get_height());
    uint2 coords = uint2(texCoords * float2(dimens));

    return coords;
}

template <typename T>
static inline vec<T, 4> blitSampleTextureMS(texture2d_ms<T> srcTexture, float2 texCoords)
{
    uint2 coords = getImageCoords(srcTexture, texCoords);
    return resolveTextureMS(srcTexture, coords);
}

template <typename T>
static inline vec<T, 4> blitSampleTexture3D(texture3d<T> srcTexture,
                                            sampler textureSampler,
                                            float2 texCoords,
                                            constant BlitParams &options)
{
    uint depth   = srcTexture.get_depth(options.srcLevel);
    float zCoord = (float(options.srcLayer) + 0.5) / float(depth);

    return srcTexture.sample(textureSampler, float3(texCoords, zCoord), level(options.srcLevel));
}

// clang-format off
#define BLIT_COLOR_FS_PARAMS(TYPE)                                                               \
    BlitVSOut input [[stage_in]],                                                                \
    texture2d<TYPE> srcTexture2d [[texture(0), function_constant(kSourceTextureType2D)]],        \
    texture2d_array<TYPE> srcTexture2dArray                                                      \
    [[texture(0), function_constant(kSourceTextureType2DArray)]],                                \
    texture2d_ms<TYPE> srcTexture2dMS [[texture(0), function_constant(kSourceTextureType2DMS)]], \
    texturecube<TYPE> srcTextureCube [[texture(0), function_constant(kSourceTextureTypeCube)]],  \
    texture3d<TYPE> srcTexture3d [[texture(0), function_constant(kSourceTextureType3D)]],        \
    sampler textureSampler [[sampler(0)]],                                                       \
    constant BlitParams &options [[buffer(0)]]
// clang-format on

#define FORWARD_BLIT_COLOR_FS_PARAMS                                                      \
    input, srcTexture2d, srcTexture2dArray, srcTexture2dMS, srcTextureCube, srcTexture3d, \
        textureSampler, options

template <typename T>
static inline vec<T, 4> blitReadTexture(BLIT_COLOR_FS_PARAMS(T))
{
    vec<T, 4> output;

    switch (kSourceTextureType)
    {
        case kTextureType2D:
            output = srcTexture2d.sample(textureSampler, input.texCoords, level(options.srcLevel));
            break;
        case kTextureType2DArray:
            output = srcTexture2dArray.sample(textureSampler, input.texCoords, options.srcLayer,
                                              level(options.srcLevel));
            break;
        case kTextureType2DMultisample:
            output = blitSampleTextureMS(srcTexture2dMS, input.texCoords);
            break;
        case kTextureTypeCube:
            output = srcTextureCube.sample(textureSampler,
                                           cubeTexcoords(input.texCoords, options.srcLayer),
                                           level(options.srcLevel));
            break;
        case kTextureType3D:
            output = blitSampleTexture3D(srcTexture3d, textureSampler, input.texCoords, options);
            break;
    }

    if (kPremultiplyAlpha)
    {
        output.xyz *= output.a;
    }
    else if (kUnmultiplyAlpha)
    {
        if (output.a != 0.0)
        {
            output.xyz /= output.a;
        }
    }

    if (options.dstLuminance)
    {
        output.g = output.b = output.r;
    }

    return output;
}

template <typename T>
static inline MultipleColorOutputs<T> blitFS(BLIT_COLOR_FS_PARAMS(T))
{
    vec<T, 4> output = blitReadTexture(FORWARD_BLIT_COLOR_FS_PARAMS);

    return toMultipleColorOutputs(output);
}

fragment MultipleColorOutputs<float> blitFloatFS(BLIT_COLOR_FS_PARAMS(float))
{
    return blitFS(FORWARD_BLIT_COLOR_FS_PARAMS);
}
fragment MultipleColorOutputs<int> blitIntFS(BLIT_COLOR_FS_PARAMS(int))
{
    return blitFS(FORWARD_BLIT_COLOR_FS_PARAMS);
}
fragment MultipleColorOutputs<uint> blitUIntFS(BLIT_COLOR_FS_PARAMS(uint))
{
    return blitFS(FORWARD_BLIT_COLOR_FS_PARAMS);
}

fragment MultipleColorOutputs<uint> copyTextureFloatToUIntFS(BLIT_COLOR_FS_PARAMS(float))
{
    float4 inputColor = blitReadTexture<>(FORWARD_BLIT_COLOR_FS_PARAMS);
    uint4 output = uint4(inputColor * float4(255.0));

    return toMultipleColorOutputs(output);
}

// Depth & stencil blitting.
struct FragmentDepthOut
{
    float depth [[depth(any)]];
};

static inline float sampleDepth(
    texture2d<float> srcTexture2d [[function_constant(kSourceTextureType2D)]],
    texture2d_array<float> srcTexture2dArray [[function_constant(kSourceTextureType2DArray)]],
    texture2d_ms<float> srcTexture2dMS [[function_constant(kSourceTextureType2DMS)]],
    texturecube<float> srcTextureCube [[function_constant(kSourceTextureTypeCube)]],
    float2 texCoords,
    constant BlitParams &options)
{
    float4 output;

    constexpr sampler textureSampler(mag_filter::nearest, min_filter::nearest);

    switch (kSourceTextureType)
    {
        case kTextureType2D:
            output = srcTexture2d.sample(textureSampler, texCoords, level(options.srcLevel));
            break;
        case kTextureType2DArray:
            output = srcTexture2dArray.sample(textureSampler, texCoords, options.srcLayer,
                                              level(options.srcLevel));
            break;
        case kTextureType2DMultisample:
            // Always use sample 0 for depth resolve:
            output = srcTexture2dMS.read(getImageCoords(srcTexture2dMS, texCoords), 0);
            break;
        case kTextureTypeCube:
            output =
                srcTextureCube.sample(textureSampler, cubeTexcoords(texCoords, options.srcLayer),
                                      level(options.srcLevel));
            break;
    }

    return output.r;
}

fragment FragmentDepthOut blitDepthFS(BlitVSOut input [[stage_in]],
                                      texture2d<float> srcTexture2d
                                      [[texture(0), function_constant(kSourceTextureType2D)]],
                                      texture2d_array<float> srcTexture2dArray
                                      [[texture(0), function_constant(kSourceTextureType2DArray)]],
                                      texture2d_ms<float> srcTexture2dMS
                                      [[texture(0), function_constant(kSourceTextureType2DMS)]],
                                      texturecube<float> srcTextureCube
                                      [[texture(0), function_constant(kSourceTextureTypeCube)]],
                                      constant BlitParams &options [[buffer(0)]])
{
    FragmentDepthOut re;

    re.depth = sampleDepth(srcTexture2d, srcTexture2dArray, srcTexture2dMS, srcTextureCube,
                           input.texCoords, options);

    return re;
}

static inline uint32_t sampleStencil(
    texture2d<uint32_t> srcTexture2d [[function_constant(kSourceTexture2Type2D)]],
    texture2d_array<uint32_t> srcTexture2dArray [[function_constant(kSourceTexture2Type2DArray)]],
    texture2d_ms<uint32_t> srcTexture2dMS [[function_constant(kSourceTexture2Type2DMS)]],
    texturecube<uint32_t> srcTextureCube [[function_constant(kSourceTexture2TypeCube)]],
    float2 texCoords,
    int srcLevel,
    int srcLayer)
{
    uint4 output;
    constexpr sampler textureSampler(mag_filter::nearest, min_filter::nearest);

    switch (kSourceTexture2Type)
    {
        case kTextureType2D:
            output = srcTexture2d.sample(textureSampler, texCoords, level(srcLevel));
            break;
        case kTextureType2DArray:
            output = srcTexture2dArray.sample(textureSampler, texCoords, srcLayer, level(srcLevel));
            break;
        case kTextureType2DMultisample:
            // Always use sample 0 for stencil resolve:
            output = srcTexture2dMS.read(getImageCoords(srcTexture2dMS, texCoords), 0);
            break;
        case kTextureTypeCube:
            output = srcTextureCube.sample(textureSampler, cubeTexcoords(texCoords, srcLayer),
                                           level(srcLevel));
            break;
    }

    return output.r;
}

// Write stencil to a buffer
struct BlitStencilToBufferParams
{
    float2 srcStartTexCoords;
    float2 srcTexCoordSteps;
    int srcLevel;
    int srcLayer;

    uint2 dstSize;
    uint dstBufferRowPitch;
    // Is multisample resolve needed?
    bool resolveMS;
};

kernel void blitStencilToBufferCS(ushort2 gIndices [[thread_position_in_grid]],
                                  texture2d<uint32_t> srcTexture2d
                                  [[texture(1), function_constant(kSourceTexture2Type2D)]],
                                  texture2d_array<uint32_t> srcTexture2dArray
                                  [[texture(1), function_constant(kSourceTexture2Type2DArray)]],
                                  texture2d_ms<uint32_t> srcTexture2dMS
                                  [[texture(1), function_constant(kSourceTexture2Type2DMS)]],
                                  texturecube<uint32_t> srcTextureCube
                                  [[texture(1), function_constant(kSourceTexture2TypeCube)]],
                                  constant BlitStencilToBufferParams &options [[buffer(0)]],
                                  device uchar *buffer [[buffer(1)]])
{
    if (gIndices.x >= options.dstSize.x || gIndices.y >= options.dstSize.y)
    {
        return;
    }

    float2 srcTexCoords = options.srcStartTexCoords + float2(gIndices) * options.srcTexCoordSteps;

    if (kSourceTexture2Type == kTextureType2DMultisample && !options.resolveMS)
    {
        uint samples      = srcTexture2dMS.get_num_samples();
        uint2 imageCoords = getImageCoords(srcTexture2dMS, srcTexCoords);
        uint bufferOffset = options.dstBufferRowPitch * gIndices.y + samples * gIndices.x;

        for (uint sample = 0; sample < samples; ++sample)
        {
            uint stencilPerSample         = srcTexture2dMS.read(imageCoords, sample).r;
            buffer[bufferOffset + sample] = static_cast<uchar>(stencilPerSample);
        }
    }
    else
    {
        uint32_t stencil =
            sampleStencil(srcTexture2d, srcTexture2dArray, srcTexture2dMS, srcTextureCube,
                          srcTexCoords, options.srcLevel, options.srcLayer);

        buffer[options.dstBufferRowPitch * gIndices.y + gIndices.x] = static_cast<uchar>(stencil);
    }
}

// Fragment's stencil output is only available since Metal 2.1
@@#if __METAL_VERSION__ >= 210

struct FragmentStencilOut
{
    uint32_t stencil [[stencil]];
};

struct FragmentDepthStencilOut
{
    float depth [[depth(any)]];
    uint32_t stencil [[stencil]];
};

fragment FragmentStencilOut blitStencilFS(
    BlitVSOut input [[stage_in]],
    texture2d<uint32_t> srcTexture2d [[texture(1), function_constant(kSourceTexture2Type2D)]],
    texture2d_array<uint32_t> srcTexture2dArray
    [[texture(1), function_constant(kSourceTexture2Type2DArray)]],
    texture2d_ms<uint32_t> srcTexture2dMS
    [[texture(1), function_constant(kSourceTexture2Type2DMS)]],
    texturecube<uint32_t> srcTextureCube [[texture(1), function_constant(kSourceTexture2TypeCube)]],
    constant BlitParams &options [[buffer(0)]])
{
    FragmentStencilOut re;

    re.stencil = sampleStencil(srcTexture2d, srcTexture2dArray, srcTexture2dMS, srcTextureCube,
                               input.texCoords, options.srcLevel, options.srcLayer);

    return re;
}

fragment FragmentDepthStencilOut blitDepthStencilFS(
    BlitVSOut input [[stage_in]],
    // Source depth texture
    texture2d<float> srcDepthTexture2d [[texture(0), function_constant(kSourceTextureType2D)]],
    texture2d_array<float> srcDepthTexture2dArray
    [[texture(0), function_constant(kSourceTextureType2DArray)]],
    texture2d_ms<float> srcDepthTexture2dMS
    [[texture(0), function_constant(kSourceTextureType2DMS)]],
    texturecube<float> srcDepthTextureCube
    [[texture(0), function_constant(kSourceTextureTypeCube)]],

    // Source stencil texture
    texture2d<uint32_t> srcStencilTexture2d
    [[texture(1), function_constant(kSourceTexture2Type2D)]],
    texture2d_array<uint32_t> srcStencilTexture2dArray
    [[texture(1), function_constant(kSourceTexture2Type2DArray)]],
    texture2d_ms<uint32_t> srcStencilTexture2dMS
    [[texture(1), function_constant(kSourceTexture2Type2DMS)]],
    texturecube<uint32_t> srcStencilTextureCube
    [[texture(1), function_constant(kSourceTexture2TypeCube)]],

    constant BlitParams &options [[buffer(0)]])
{
    FragmentDepthStencilOut re;

    re.depth = sampleDepth(srcDepthTexture2d, srcDepthTexture2dArray, srcDepthTexture2dMS,
                           srcDepthTextureCube, input.texCoords, options);
    re.stencil =
        sampleStencil(srcStencilTexture2d, srcStencilTexture2dArray, srcStencilTexture2dMS,
                      srcStencilTextureCube, input.texCoords, options.srcLevel, options.srcLayer);
    return re;
}
@@#endif  // __METAL_VERSION__ >= 210