Metal implementation pt 1: autogen resources

scripts/code_generation_hashes/Metal_default_shaders.json

+{
+  "src/libANGLE/renderer/metal/shaders/compiled/mtl_default_shaders.inc":
+    "a60a4682f4aa7cea2d35e551af938e3b",
+  "src/libANGLE/renderer/metal/shaders/gen_mtl_internal_shaders.py":
+    "5d8832978de07470b2d6dbf18104ba76",
+  "src/libANGLE/renderer/metal/shaders/master_source.metal":
+    "fbe6f4bfb49a48ae87791a4cff5fab0a",
+  "src/libANGLE/renderer/metal/shaders/mtl_default_shaders_src_autogen.inc":
+    "e3546ea73c88d12f8d59b2b6a286b4bf"
+}
\ No newline at end of file

scripts/code_generation_hashes/Metal_format_table.json

+{
+  "src/libANGLE/renderer/metal/gen_mtl_format_table.py":
+    "5dd54a352213e303ba517fa105bca8fb",
+  "src/libANGLE/renderer/metal/mtl_format_map.json":
+    "c6f5b6dda11e456cfbcaeec53eb46fa0",
+  "src/libANGLE/renderer/metal/mtl_format_table_autogen.mm":
+    "9ccbbb1c6d4f84ea2f8a22c126040ec9"
+}
\ No newline at end of file

scripts/run_code_generation.py

@@ -112,7 +112,11 @@ generators = {
     'Emulated HLSL functions':
         'src/compiler/translator/gen_emulated_builtin_function_tables.py',
     'Static builtins':
-        'src/compiler/translator/gen_builtin_symbols.py'
+        'src/compiler/translator/gen_builtin_symbols.py',
+    'Metal format table':
+        'src/libANGLE/renderer/metal/gen_mtl_format_table.py',
+    'Metal default shaders':
+        'src/libANGLE/renderer/metal/shaders/gen_mtl_internal_shaders.py',
 }
 
 

src/libANGLE/renderer/metal/mtl_format_map.json

+{
+    "image": {
+        "map": {
+            "NONE": "MTLPixelFormatInvalid",
+            "A8_UNORM": "MTLPixelFormatA8Unorm",
+            "R8_UNORM": "MTLPixelFormatR8Unorm",
+            "R8G8_UNORM": "MTLPixelFormatRG8Unorm",
+            "R16_UNORM": "MTLPixelFormatR16Unorm",
+            "R16_FLOAT": "MTLPixelFormatR16Float",
+            "R16G16_FLOAT": "MTLPixelFormatRG16Float",
+            "R16G16B16A16_FLOAT": "MTLPixelFormatRGBA16Float",
+            "R32_FLOAT": "MTLPixelFormatR32Float",
+            "R32G32_FLOAT": "MTLPixelFormatRG32Float",
+            "R32G32B32A32_FLOAT": "MTLPixelFormatRGBA32Float",
+            "R8G8B8A8_UNORM": "MTLPixelFormatRGBA8Unorm",
+            "R8G8B8A8_UNORM_SRGB": "MTLPixelFormatRGBA8Unorm_sRGB",
+            "B8G8R8A8_UNORM": "MTLPixelFormatBGRA8Unorm",
+            "B8G8R8A8_UNORM_SRGB": "MTLPixelFormatBGRA8Unorm_sRGB",
+            "D32_FLOAT": "MTLPixelFormatDepth32Float",
+            "S8_UINT": "MTLPixelFormatStencil8",
+            "D32_FLOAT_S8X24_UINT": "MTLPixelFormatDepth32Float_Stencil8"
+        },
+        "map_ios": {
+            "R5G6B5_UNORM": "MTLPixelFormatB5G6R5Unorm",
+            "R5G5B5A1_UNORM": "MTLPixelFormatBGR5A1Unorm",
+            "R4G4B4A4_UNORM": "MTLPixelFormatABGR4Unorm",
+            "PVRTC1_RGB_4BPP_UNORM_BLOCK": "MTLPixelFormatPVRTC_RGB_4BPP",
+            "PVRTC1_RGB_2BPP_UNORM_BLOCK": "MTLPixelFormatPVRTC_RGB_2BPP",
+            "PVRTC1_RGBA_4BPP_UNORM_BLOCK": "MTLPixelFormatPVRTC_RGBA_4BPP",
+            "PVRTC1_RGBA_2BPP_UNORM_BLOCK": "MTLPixelFormatPVRTC_RGBA_2BPP",
+            "PVRTC1_RGB_2BPP_UNORM_SRGB_BLOCK": "MTLPixelFormatPVRTC_RGB_2BPP_sRGB",
+            "PVRTC1_RGB_4BPP_UNORM_SRGB_BLOCK": "MTLPixelFormatPVRTC_RGB_4BPP_sRGB",
+            "PVRTC1_RGBA_2BPP_UNORM_SRGB_BLOCK": "MTLPixelFormatPVRTC_RGBA_2BPP_sRGB",
+            "PVRTC1_RGBA_4BPP_UNORM_SRGB_BLOCK": "MTLPixelFormatPVRTC_RGBA_4BPP_sRGB",
+            "ETC1_R8G8B8_UNORM_BLOCK": "MTLPixelFormatETC2_RGB8",
+            "ETC2_R8G8B8_UNORM_BLOCK": "MTLPixelFormatETC2_RGB8",
+            "ETC2_R8G8B8_SRGB_BLOCK": "MTLPixelFormatETC2_RGB8_sRGB",
+            "ETC2_R8G8B8A1_UNORM_BLOCK": "MTLPixelFormatETC2_RGB8A1",
+            "ETC2_R8G8B8A1_SRGB_BLOCK": "MTLPixelFormatETC2_RGB8A1_sRGB",
+            "ETC2_R8G8B8A8_UNORM_BLOCK": "MTLPixelFormatEAC_RGBA8",
+            "ETC2_R8G8B8A8_SRGB_BLOCK": "MTLPixelFormatEAC_RGBA8_sRGB",
+            "EAC_R11_UNORM_BLOCK": "MTLPixelFormatEAC_R11Unorm",
+            "EAC_R11_SNORM_BLOCK": "MTLPixelFormatEAC_R11Snorm",
+            "EAC_R11G11_UNORM_BLOCK": "MTLPixelFormatEAC_R11Unorm",
+            "EAC_R11G11_SNORM_BLOCK": "MTLPixelFormatEAC_RG11Snorm"
+        },
+        "map_mac": {
+            "D16_UNORM": "MTLPixelFormatDepth16Unorm",
+            "D24_UNORM_S8_UINT": "MTLPixelFormatDepth24Unorm_Stencil8",
+            "BC1_RGB_UNORM_BLOCK": "MTLPixelFormatBC1_RGBA",
+            "BC1_RGB_UNORM_SRGB_BLOCK": "MTLPixelFormatBC1_RGBA_sRGB",
+            "BC1_RGBA_UNORM_BLOCK": "MTLPixelFormatBC1_RGBA",
+            "BC1_RGBA_UNORM_SRGB_BLOCK": "MTLPixelFormatBC1_RGBA_sRGB",
+            "BC2_RGBA_UNORM_BLOCK": "MTLPixelFormatBC2_RGBA",
+            "BC2_RGBA_UNORM_SRGB_BLOCK": "MTLPixelFormatBC2_RGBA_sRGB",
+            "BC3_RGBA_UNORM_BLOCK": "MTLPixelFormatBC3_RGBA",
+            "BC3_RGBA_UNORM_SRGB_BLOCK": "MTLPixelFormatBC3_RGBA_sRGB"
+        },
+        "override": {
+            "L8_UNORM": "R8G8B8A8_UNORM",
+            "L8A8_UNORM": "R8G8B8A8_UNORM",
+            "R8G8B8_UNORM": "R8G8B8A8_UNORM",
+            "R8G8B8_UNORM_SRGB": "R8G8B8A8_UNORM_SRGB",
+            "R32G32B32_FLOAT": "R32G32B32A32_FLOAT",
+            "R16G16B16_FLOAT": "R16G16B16A16_FLOAT",
+            "A16_FLOAT": "R16G16B16A16_FLOAT",
+            "L16_FLOAT": "R16G16B16A16_FLOAT",
+            "L16A16_FLOAT": "R16G16B16A16_FLOAT",
+            "A32_FLOAT": "R32G32B32A32_FLOAT",
+            "L32_FLOAT": "R32G32B32A32_FLOAT",
+            "L32A32_FLOAT": "R32G32B32A32_FLOAT",
+            "D24_UNORM_X8_UINT": "D32_FLOAT"
+        },
+        "override_ios": {
+            "D24_UNORM_S8_UINT": "D32_FLOAT_S8X24_UINT",
+            "D16_UNORM": "D32_FLOAT"
+        },
+        "override_mac": {
+            "R5G6B5_UNORM": "R8G8B8A8_UNORM",
+            "R5G5B5A1_UNORM": "R8G8B8A8_UNORM",
+            "R4G4B4A4_UNORM": "R8G8B8A8_UNORM"
+        },
+        "fallbacks_mac": {
+            "B8G8R8A8_UNORM": "R8G8B8A8_UNORM",
+            "D24_UNORM_S8_UINT": "D32_FLOAT_S8X24_UINT",
+            "R8_UNORM": "NONE",
+            "R16_FLOAT": "NONE",
+            "R32_FLOAT": "NONE",
+            "R16G16B16A16_FLOAT": "NONE",
+            "R32G32B32A32_FLOAT": "NONE"
+        }
+    },
+    "vertex": {
+        "map": {
+            "NONE": "MTLVertexFormatInvalid",
+            "R8_UNORM": "MTLVertexFormatUCharNormalized",
+            "R8_SNORM": "MTLVertexFormatCharNormalized",
+            "R8_UINT": "MTLVertexFormatUChar",
+            "R8_SINT": "MTLVertexFormatChar",
+            "R8_USCALED": "MTLVertexFormatUChar",
+            "R8_SSCALED": "MTLVertexFormatChar",
+            "R8G8_UNORM": "MTLVertexFormatUChar2Normalized",
+            "R8G8_SNORM": "MTLVertexFormatChar2Normalized",
+            "R8G8_UINT": "MTLVertexFormatUChar2",
+            "R8G8_SINT": "MTLVertexFormatChar2",
+            "R8G8_USCALED": "MTLVertexFormatUChar2",
+            "R8G8_SSCALED": "MTLVertexFormatChar2",
+            "R8G8B8_UNORM": "MTLVertexFormatUChar3Normalized",
+            "R8G8B8_SNORM": "MTLVertexFormatChar3Normalized",
+            "R8G8B8_UINT": "MTLVertexFormatUChar3",
+            "R8G8B8_SINT": "MTLVertexFormatChar3",
+            "R8G8B8_USCALED": "MTLVertexFormatUChar3",
+            "R8G8B8_SSCALED": "MTLVertexFormatChar3",
+            "R8G8B8A8_UNORM": "MTLVertexFormatUChar4Normalized",
+            "R8G8B8A8_SNORM": "MTLVertexFormatChar4Normalized",
+            "R8G8B8A8_UINT": "MTLVertexFormatUChar4",
+            "R8G8B8A8_SINT": "MTLVertexFormatChar4",
+            "R8G8B8A8_USCALED": "MTLVertexFormatUChar4",
+            "R8G8B8A8_SSCALED": "MTLVertexFormatChar4",
+
+            "R16_UNORM": "MTLVertexFormatUShortNormalized",
+            "R16_SNORM": "MTLVertexFormatShortNormalized",
+            "R16_UINT": "MTLVertexFormatUShort",
+            "R16_SINT": "MTLVertexFormatShort",
+            "R16_USCALED": "MTLVertexFormatUShort",
+            "R16_SSCALED": "MTLVertexFormatShort",
+            "R16G16_UNORM": "MTLVertexFormatUShort2Normalized",
+            "R16G16_SNORM": "MTLVertexFormatShort2Normalized",
+            "R16G16_UINT": "MTLVertexFormatUShort2",
+            "R16G16_SINT": "MTLVertexFormatShort2",
+            "R16G16_USCALED": "MTLVertexFormatUShort2",
+            "R16G16_SSCALED": "MTLVertexFormatShort2",
+            "R16G16B16_UNORM": "MTLVertexFormatUShort3Normalized",
+            "R16G16B16_SNORM": "MTLVertexFormatShort3Normalized",
+            "R16G16B16_UINT": "MTLVertexFormatUShort3",
+            "R16G16B16_SINT": "MTLVertexFormatShort3",
+            "R16G16B16_USCALED": "MTLVertexFormatUShort3",
+            "R16G16B16_SSCALED": "MTLVertexFormatShort3",
+            "R16G16B16A16_UNORM": "MTLVertexFormatUShort4Normalized",
+            "R16G16B16A16_SNORM": "MTLVertexFormatShort4Normalized",
+            "R16G16B16A16_UINT": "MTLVertexFormatUShort4",
+            "R16G16B16A16_SINT": "MTLVertexFormatShort4",
+            "R16G16B16A16_USCALED": "MTLVertexFormatUShort4",
+            "R16G16B16A16_SSCALED": "MTLVertexFormatShort4",
+
+            "R32_FLOAT": "MTLVertexFormatFloat",
+            "R32G32_FLOAT": "MTLVertexFormatFloat2",
+            "R32G32B32_FLOAT": "MTLVertexFormatFloat3",
+            "R32G32B32A32_FLOAT": "MTLVertexFormatFloat4"
+        },
+        "override": {
+            "R32_FIXED": "R32_FLOAT",
+            "R32G32_FIXED": "R32G32_FLOAT",
+            "R32G32B32_FIXED": "R32G32B32_FLOAT",
+            "R32G32B32A32_FIXED": "R32G32B32A32_FLOAT"
+        },
+        "override_tightly_packed": {
+            "R16_UNORM": "R32_FLOAT",
+            "R16_SNORM": "R32_FLOAT",
+            "R16_UINT": "R32_FLOAT",
+            "R16_SINT": "R32_FLOAT",
+            "R16_USCALED": "R32_FLOAT",
+            "R16_SSCALED": "R32_FLOAT",
+            "R16G16B16_UNORM": "R32G32B32_FLOAT",
+            "R16G16B16_SNORM": "R32G32B32_FLOAT",
+            "R16G16B16_UINT": "R32G32B32_FLOAT",
+            "R16G16B16_SINT": "R32G32B32_FLOAT",
+            "R16G16B16_USCALED": "R32G32B32_FLOAT",
+            "R16G16B16_SSCALED": "R32G32B32_FLOAT",
+            "R8_UNORM": "R32_FLOAT",
+            "R8_SNORM": "R32_FLOAT",
+            "R8_UINT": "R32_FLOAT",
+            "R8_SINT": "R32_FLOAT",
+            "R8_USCALED": "R32_FLOAT",
+            "R8_SSCALED": "R32_FLOAT",
+            "R8G8_UNORM": "R32G32_FLOAT",
+            "R8G8_SNORM": "R32G32_FLOAT",
+            "R8G8_UINT": "R32G32_FLOAT",
+            "R8G8_SINT": "R32G32_FLOAT",
+            "R8G8_USCALED": "R32G32_FLOAT",
+            "R8G8_SSCALED": "R32G32_FLOAT",
+            "R8G8B8_UNORM": "R32G32B32_FLOAT",
+            "R8G8B8_SNORM": "R32G32B32_FLOAT",
+            "R8G8B8_UINT": "R32G32B32_FLOAT",
+            "R8G8B8_SINT": "R32G32B32_FLOAT",
+            "R8G8B8_USCALED": "R32G32B32_FLOAT",
+            "R8G8B8_SSCALED": "R32G32B32_FLOAT"
+        }
+    }
+}
\ No newline at end of file

src/libANGLE/renderer/metal/mtl_format_utils.h

+//
+// Copyright 2019 The ANGLE Project Authors. All rights reserved.
+// Use of this source code is governed by a BSD-style license that can be
+// found in the LICENSE file.
+//
+// mtl_format_utils.h:
+//      Declares Format conversion utilities classes that convert from angle formats
+//      to respective MTLPixelFormat and MTLVertexFormat.
+//
+
+#ifndef LIBANGLE_RENDERER_METAL_MTL_FORMAT_UTILS_H_
+#define LIBANGLE_RENDERER_METAL_MTL_FORMAT_UTILS_H_
+
+#import <Metal/Metal.h>
+
+#include "common/angleutils.h"
+#include "libANGLE/Caps.h"
+#include "libANGLE/formatutils.h"
+#include "libANGLE/renderer/copyvertex.h"
+
+namespace rx
+{
+class RendererMtl;
+
+namespace mtl
+{
+
+struct FormatBase
+{
+    const angle::Format &actualAngleFormat() const;
+    const angle::Format &intendedAngleFormat() const;
+
+    angle::FormatID actualFormatId   = angle::FormatID::NONE;
+    angle::FormatID intendedFormatId = angle::FormatID::NONE;
+};
+
+// Pixel format
+struct Format : public FormatBase
+{
+    Format() = default;
+
+    const gl::InternalFormat &intendedInternalFormat() const;
+
+    bool valid() const { return metalFormat != MTLPixelFormatInvalid; }
+
+    static bool FormatRenderable(MTLPixelFormat format);
+    static bool FormatCPUReadable(MTLPixelFormat format);
+
+    MTLPixelFormat metalFormat = MTLPixelFormatInvalid;
+
+  private:
+    void init(const RendererMtl *renderer, angle::FormatID intendedFormatId);
+
+    friend class FormatTable;
+};
+
+// Vertex format
+struct VertexFormat : public FormatBase
+{
+    VertexFormat() = default;
+
+    MTLVertexFormat metalFormat = MTLVertexFormatInvalid;
+
+    VertexCopyFunction vertexLoadFunction = nullptr;
+
+  private:
+    void init(angle::FormatID angleFormatId, bool tightlyPacked = false);
+
+    friend class FormatTable;
+};
+
+class FormatTable final : angle::NonCopyable
+{
+  public:
+    FormatTable()  = default;
+    ~FormatTable() = default;
+
+    angle::Result initialize(const RendererMtl *renderer);
+
+    void generateTextureCaps(const RendererMtl *renderer,
+                             gl::TextureCapsMap *capsMapOut,
+                             std::vector<GLenum> *compressedFormatsOut) const;
+
+    const Format &getPixelFormat(angle::FormatID angleFormatId) const;
+
+    // tightlyPacked means this format will be used in a tightly packed vertex buffer.
+    // In that case, it's easier to just convert everything to float to ensure
+    // Metal alignment requirements between 2 elements inside the buffer will be met regardless
+    // of how many components each element has.
+    const VertexFormat &getVertexFormat(angle::FormatID angleFormatId, bool tightlyPacked) const;
+
+  private:
+    std::array<Format, angle::kNumANGLEFormats> mPixelFormatTable;
+    // One for tightly packed buffers, one for general cases.
+    std::array<VertexFormat, angle::kNumANGLEFormats> mVertexFormatTables[2];
+};
+
+}  // namespace mtl
+}  // namespace rx
+
+#endif /* LIBANGLE_RENDERER_METAL_MTL_FORMAT_UTILS_H_ */

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"
+
+
+struct BlitParams
+{
+    // 0: lower left, 1: lower right, 2: upper left, 3: upper right
+    float2 srcTexCoords[4];
+    int srcLevel;
+    bool srcLuminance; // source texture is luminance texture
+    bool dstFlipY;
+    bool dstLuminance; // destination texture is luminance;
+};
+
+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[gTexcoordsIndices[vid]];
+
+    if (options.dstFlipY)
+    {
+        output.position = -output.position;
+    }
+
+    return output;
+}
+
+float4 blitSampleTexture(texture2d<float> srcTexture,
+                     float2 texCoords,
+                     constant BlitParams &options)
+{
+    constexpr sampler textureSampler(mag_filter::linear,
+                                     min_filter::linear);
+    float4 output = srcTexture.sample(textureSampler, texCoords, level(options.srcLevel));
+
+    if (options.srcLuminance)
+    {
+        output.gb = float2(output.r, output.r);
+    }
+
+    return output;
+}
+
+float4 blitOutput(float4 color, constant BlitParams &options)
+{
+    float4 ret = color;
+
+    if (options.dstLuminance)
+    {
+        ret.r = ret.g = ret.b = (color.r * 0.3) + (color.g * 0.59) + (color.b * 0.11);
+    }
+
+    return ret;
+}
+
+fragment float4 blitFS(BlitVSOut input [[stage_in]],
+                       texture2d<float> srcTexture [[texture(0)]],
+                       constant BlitParams &options [[buffer(0)]])
+{
+    return blitOutput(blitSampleTexture(srcTexture, input.texCoords, options), options);
+}
+
+fragment float4 blitPremultiplyAlphaFS(BlitVSOut input [[stage_in]],
+                                       texture2d<float> srcTexture [[texture(0)]],
+                                       constant BlitParams &options [[buffer(0)]])
+{
+    float4 output = blitSampleTexture(srcTexture, input.texCoords, options);
+    output.xyz *= output.a;
+    return blitOutput(output, options);
+}
+
+fragment float4 blitUnmultiplyAlphaFS(BlitVSOut input [[stage_in]],
+                                      texture2d<float> srcTexture [[texture(0)]],
+                                      constant BlitParams &options [[buffer(0)]])
+{
+    float4 output = blitSampleTexture(srcTexture, input.texCoords, options);
+    if (output.a != 0.0)
+    {
+        output.xyz *= 1.0 / output.a;
+    }
+    return blitOutput(output, options);
+}

src/libANGLE/renderer/metal/shaders/clear.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.
+//
+// clear.metal: Implements viewport clearing.
+
+#include "common.h"
+
+struct ClearParams
+{
+    float4 clearColor;
+    float clearDepth;
+};
+
+vertex float4 clearVS(unsigned int vid [[ vertex_id ]],
+                      constant ClearParams &clearParams [[buffer(0)]])
+{
+    return float4(gCorners[vid], clearParams.clearDepth, 1.0);
+}
+
+fragment float4 clearFS(constant ClearParams &clearParams [[buffer(0)]])
+{
+    return clearParams.clearColor;
+}

src/libANGLE/renderer/metal/shaders/common.h

+//
+// 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.
+//
+// common.h: Common header for other metal source code.
+
+#ifndef LIBANGLE_RENDERER_METAL_SHADERS_COMMON_H_
+#define LIBANGLE_RENDERER_METAL_SHADERS_COMMON_H_
+
+#ifndef SKIP_STD_HEADERS
+#    include <simd/simd.h>
+#    include <metal_stdlib>
+#endif
+
+using namespace metal;
+
+// Full screen quad's vertices
+constant float2 gCorners[6] = {
+    float2(-1.0f, 1.0f), float2(1.0f, -1.0f), float2(-1.0f, -1.0f),
+    float2(-1.0f, 1.0f), float2(1.0f, 1.0f),  float2(1.0f, -1.0f),
+};
+
+// Full screen quad's texcoords indices:
+// 0: lower left, 1: lower right, 2: upper left, 3: upper right
+constant int gTexcoordsIndices[6] = {2, 1, 0, 2, 3, 1};
+
+fragment float4 dummyFS()
+{
+    return float4(0, 0, 0, 0);
+}
+#endif /* LIBANGLE_RENDERER_METAL_SHADERS_COMMON_H_ */

src/libANGLE/renderer/metal/shaders/gen_indices.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.
+//
+
+#include "common.h"
+
+struct IndexConversionParams
+{
+    uint32_t srcOffset;  // offset in bytes
+    uint32_t indexCount;
+};
+
+#define ANGLE_IDX_CONVERSION_GUARD(IDX, OPTS) \
+    if (IDX >= OPTS.indexCount)               \
+    {                                         \
+        return;                               \
+    }
+
+kernel void convertIndexU8ToU16(uint idx[[thread_position_in_grid]],
+                                constant IndexConversionParams &options[[buffer(0)]],
+                                constant uchar *input[[buffer(1)]],
+                                device ushort *output[[buffer(2)]])
+{
+    ANGLE_IDX_CONVERSION_GUARD(idx, options);
+    output[idx] = input[options.srcOffset + idx];
+}
+
+kernel void convertIndexU16Unaligned(uint idx[[thread_position_in_grid]],
+                                     constant IndexConversionParams &options[[buffer(0)]],
+                                     constant uchar *input[[buffer(1)]],
+                                     device ushort *output[[buffer(2)]])
+{
+    ANGLE_IDX_CONVERSION_GUARD(idx, options);
+    ushort inputLo = input[options.srcOffset + 2 * idx];
+    ushort inputHi = input[options.srcOffset + 2 * idx + 1];
+    // Little endian conversion:
+    ushort value = inputLo | (inputHi << 8);
+    output[idx] = value;
+}
+
+kernel void convertIndexU16Aligned(uint idx[[thread_position_in_grid]],
+                                   constant IndexConversionParams &options[[buffer(0)]],
+                                   constant ushort *input[[buffer(1)]],
+                                   device ushort *output[[buffer(2)]])
+{
+    ANGLE_IDX_CONVERSION_GUARD(idx, options);
+    ushort value = input[options.srcOffset / 2 + idx];
+    output[idx] = value;
+}
+
+kernel void convertIndexU32Unaligned(uint idx[[thread_position_in_grid]],
+                                     constant IndexConversionParams &options[[buffer(0)]],
+                                     constant uchar *input[[buffer(1)]],
+                                     device uint *output[[buffer(2)]])
+{
+    ANGLE_IDX_CONVERSION_GUARD(idx, options);
+    uint input0 = input[options.srcOffset + 4 * idx];
+    uint input1 = input[options.srcOffset + 4 * idx + 1];
+    uint input2 = input[options.srcOffset + 4 * idx + 2];
+    uint input3 = input[options.srcOffset + 4 * idx + 3];
+    // Little endian conversion:
+    uint value = input0 | (input1 << 8) | (input2 << 16) | (input3 << 24);
+    output[idx] = value;
+}
+
+kernel void convertIndexU32Aligned(uint idx[[thread_position_in_grid]],
+                                   constant IndexConversionParams &options[[buffer(0)]],
+                                   constant uint *input[[buffer(1)]],
+                                   device uint *output[[buffer(2)]])
+{
+    ANGLE_IDX_CONVERSION_GUARD(idx, options);
+    uint value = input[options.srcOffset / 4 + idx];
+    output[idx] = value;
+}

src/libANGLE/renderer/metal/shaders/gen_mtl_internal_shaders.py

+#!/usr/bin/python
+# Copyright 2019 The ANGLE Project Authors. All rights reserved.
+# Use of this source code is governed by a BSD-style license that can be
+# found in the LICENSE file.
+#
+# gen_mtl_internal_shaders.py:
+#   Code generation for Metal backend's default shaders.
+#   NOTE: don't run this script directly. Run scripts/run_code_generation.py.
+
+import os
+import sys
+import json
+from datetime import datetime
+
+template_header_boilerplate = """// GENERATED FILE - DO NOT EDIT.
+// Generated by {script_name}
+//
+// Copyright {copyright_year} The ANGLE Project Authors. All rights reserved.
+// Use of this source code is governed by a BSD-style license that can be
+// found in the LICENSE file.
+//
+"""
+
+
+def main():
+    # auto_script parameters.
+    if len(sys.argv) > 1:
+        inputs = ['master_source.metal']
+        outputs = ['compiled/mtl_default_shaders.inc', 'mtl_default_shaders_src_autogen.inc']
+
+        if sys.argv[1] == 'inputs':
+            print ','.join(inputs)
+        elif sys.argv[1] == 'outputs':
+            print ','.join(outputs)
+        else:
+            print('Invalid script parameters')
+            return 1
+        return 0
+
+    os.chdir(sys.path[0])
+
+    print('Compiling macos version of default shaders ...')
+    os.system(
+        'xcrun -sdk macosx metal master_source.metal -mmacosx-version-min=10.13 -c -o compiled/default.air'
+    )
+    os.system('xcrun -sdk macosx metallib compiled/default.air -o compiled/default.metallib')
+
+    print('Compiling ios version of default shaders ...')
+    os.system(
+        'xcrun -sdk iphoneos metal master_source.metal -mios-version-min=8.0 -c -o compiled/default.ios.air'
+    )
+    os.system(
+        'xcrun -sdk iphoneos metallib compiled/default.ios.air -o compiled/default.ios.metallib')
+
+    print('Compiling ios simulator version of default shaders ...')
+    os.system(
+        'xcrun -sdk iphonesimulator metal master_source.metal -c -o compiled/default.ios_sim.air')
+    os.system(
+        'xcrun -sdk iphonesimulator metallib compiled/default.ios_sim.air -o compiled/default.ios_sim.metallib'
+    )
+
+    boilerplate_code = template_header_boilerplate.format(
+        script_name=sys.argv[0], copyright_year=datetime.today().year)
+
+    os.system("echo \"{0}\" > compiled/mtl_default_shaders.inc".format(boilerplate_code))
+    os.system(
+        'echo "// Compiled binary for Metal default shaders.\n\n" >> compiled/mtl_default_shaders.inc'
+    )
+    os.system('echo "#include <TargetConditionals.h>\n\n" >> compiled/mtl_default_shaders.inc')
+
+    # Mac version
+    os.system('echo "#if TARGET_OS_OSX\n" >> compiled/mtl_default_shaders.inc')
+
+    os.system('echo "constexpr" >> compiled/mtl_default_shaders.inc')
+    os.system('xxd -i compiled/default.metallib >> compiled/mtl_default_shaders.inc')
+
+    # iOS simulator version
+    os.system(
+        'echo "\n#elif TARGET_OS_SIMULATOR  // TARGET_OS_OSX\n" >> compiled/mtl_default_shaders.inc'
+    )
+
+    os.system(
+        'echo "#define compiled_default_metallib     compiled_default_ios_sim_metallib" >> compiled/mtl_default_shaders.inc'
+    )
+    os.system(
+        'echo "#define compiled_default_metallib_len compiled_default_ios_sim_metallib_len\n" >> compiled/mtl_default_shaders.inc'
+    )
+    os.system('echo "constexpr" >> compiled/mtl_default_shaders.inc')
+    os.system('xxd -i compiled/default.ios_sim.metallib >> compiled/mtl_default_shaders.inc')
+
+    # iOS version
+    os.system(
+        'echo "\n#elif TARGET_OS_IOS  // TARGET_OS_OSX\n" >> compiled/mtl_default_shaders.inc')
+
+    os.system(
+        'echo "#define compiled_default_metallib     compiled_default_ios_metallib" >> compiled/mtl_default_shaders.inc'
+    )
+    os.system(
+        'echo "#define compiled_default_metallib_len compiled_default_ios_metallib_len\n" >> compiled/mtl_default_shaders.inc'
+    )
+    os.system('echo "constexpr" >> compiled/mtl_default_shaders.inc')
+    os.system('xxd -i compiled/default.ios.metallib >> compiled/mtl_default_shaders.inc')
+
+    os.system('echo "#endif  // TARGET_OS_OSX\n" >> compiled/mtl_default_shaders.inc')
+
+    # Write full source string for debug purpose
+    os.system("echo \"{0}\" > mtl_default_shaders_src_autogen.inc".format(boilerplate_code))
+    os.system(
+        'echo "// C++ string version of Metal default shaders for debug purpose.\n\n" >> mtl_default_shaders_src_autogen.inc'
+    )
+    os.system(
+        'echo "\n\nconstexpr char default_metallib_src[] = R\\"(" >> mtl_default_shaders_src_autogen.inc'
+    )
+    os.system('echo "#include <metal_stdlib>" >> mtl_default_shaders_src_autogen.inc')
+    os.system('echo "#include <simd/simd.h>" >> mtl_default_shaders_src_autogen.inc')
+    os.system(
+        'clang -xc++ -E -DSKIP_STD_HEADERS master_source.metal >> mtl_default_shaders_src_autogen.inc'
+    )
+    os.system('echo ")\\";" >> mtl_default_shaders_src_autogen.inc')
+
+
+if __name__ == '__main__':
+    sys.exit(main())

src/libANGLE/renderer/metal/shaders/master_source.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.
+//
+// master_source.metal:
+//    Includes all other metal source code in one file.
+
+#include "clear.metal"
+#include "blit.metal"
+#include "gen_indices.metal"

src/libANGLE/renderer/metal/shaders/mtl_default_shaders_src_autogen.inc

+// GENERATED FILE - DO NOT EDIT.
+// Generated by gen_mtl_internal_shaders.py
+//
+// Copyright 2019 The ANGLE Project Authors. All rights reserved.
+// Use of this source code is governed by a BSD-style license that can be
+// found in the LICENSE file.
+//
+
+// C++ string version of Metal default shaders for debug purpose.
+
+
+
+
+constexpr char default_metallib_src[] = R"(
+#include <metal_stdlib>
+#include <simd/simd.h>
+# 1 "master_source.metal"
+# 1 "<built-in>" 1
+# 1 "<built-in>" 3
+# 374 "<built-in>" 3
+# 1 "<command line>" 1
+# 1 "<built-in>" 2
+# 1 "master_source.metal" 2
+
+
+
+
+
+
+
+
+# 1 "./clear.metal" 1
+
+
+
+
+
+
+
+# 1 "./common.h" 1
+# 16 "./common.h"
+using namespace metal;
+
+
+constant float2 gCorners[6] = {
+    float2(-1.0f, 1.0f), float2(1.0f, -1.0f), float2(-1.0f, -1.0f),
+    float2(-1.0f, 1.0f), float2(1.0f, 1.0f), float2(1.0f, -1.0f),
+};
+
+
+
+constant int gTexcoordsIndices[6] = {2, 1, 0, 2, 3, 1};
+
+fragment float4 dummyFS()
+{
+    return float4(0, 0, 0, 0);
+}
+# 9 "./clear.metal" 2
+
+struct ClearParams
+{
+    float4 clearColor;
+    float clearDepth;
+};
+
+vertex float4 clearVS(unsigned int vid [[ vertex_id ]],
+                      constant ClearParams &clearParams [[buffer(0)]])
+{
+    return float4(gCorners[vid], clearParams.clearDepth, 1.0);
+}
+
+fragment float4 clearFS(constant ClearParams &clearParams [[buffer(0)]])
+{
+    return clearParams.clearColor;
+}
+# 10 "master_source.metal" 2
+# 1 "./blit.metal" 1
+# 11 "./blit.metal"
+struct BlitParams
+{
+
+    float2 srcTexCoords[4];
+    int srcLevel;
+    bool srcLuminance;
+    bool dstFlipY;
+    bool dstLuminance;
+};
+
+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[gTexcoordsIndices[vid]];
+
+    if (options.dstFlipY)
+    {
+        output.position = -output.position;
+    }
+
+    return output;
+}
+
+float4 blitSampleTexture(texture2d<float> srcTexture,
+                     float2 texCoords,
+                     constant BlitParams &options)
+{
+    constexpr sampler textureSampler(mag_filter::linear,
+                                     min_filter::linear);
+    float4 output = srcTexture.sample(textureSampler, texCoords, level(options.srcLevel));
+
+    if (options.srcLuminance)
+    {
+        output.gb = float2(output.r, output.r);
+    }
+
+    return output;
+}
+
+float4 blitOutput(float4 color, constant BlitParams &options)
+{
+    float4 ret = color;
+
+    if (options.dstLuminance)
+    {
+        ret.r = ret.g = ret.b = (color.r * 0.3) + (color.g * 0.59) + (color.b * 0.11);
+    }
+
+    return ret;
+}
+
+fragment float4 blitFS(BlitVSOut input [[stage_in]],
+                       texture2d<float> srcTexture [[texture(0)]],
+                       constant BlitParams &options [[buffer(0)]])
+{
+    return blitOutput(blitSampleTexture(srcTexture, input.texCoords, options), options);
+}
+
+fragment float4 blitPremultiplyAlphaFS(BlitVSOut input [[stage_in]],
+                                       texture2d<float> srcTexture [[texture(0)]],
+                                       constant BlitParams &options [[buffer(0)]])
+{
+    float4 output = blitSampleTexture(srcTexture, input.texCoords, options);
+    output.xyz *= output.a;
+    return blitOutput(output, options);
+}
+
+fragment float4 blitUnmultiplyAlphaFS(BlitVSOut input [[stage_in]],
+                                      texture2d<float> srcTexture [[texture(0)]],
+                                      constant BlitParams &options [[buffer(0)]])
+{
+    float4 output = blitSampleTexture(srcTexture, input.texCoords, options);
+    if (output.a != 0.0)
+    {
+        output.xyz *= 1.0 / output.a;
+    }
+    return blitOutput(output, options);
+}
+# 11 "master_source.metal" 2
+# 1 "./gen_indices.metal" 1
+
+
+
+
+
+
+
+
+struct IndexConversionParams
+{
+    uint32_t srcOffset;
+    uint32_t indexCount;
+};
+
+
+
+
+
+
+
+kernel void convertIndexU8ToU16(uint idx[[thread_position_in_grid]],
+                                constant IndexConversionParams &options[[buffer(0)]],
+                                constant uchar *input[[buffer(1)]],
+                                device ushort *output[[buffer(2)]])
+{
+    if (idx >= options.indexCount) { return; };
+    output[idx] = input[options.srcOffset + idx];
+}
+
+kernel void convertIndexU16Unaligned(uint idx[[thread_position_in_grid]],
+                                     constant IndexConversionParams &options[[buffer(0)]],
+                                     constant uchar *input[[buffer(1)]],
+                                     device ushort *output[[buffer(2)]])
+{
+    if (idx >= options.indexCount) { return; };
+    ushort inputLo = input[options.srcOffset + 2 * idx];
+    ushort inputHi = input[options.srcOffset + 2 * idx + 1];
+
+    ushort value = inputLo | (inputHi << 8);
+    output[idx] = value;
+}
+
+kernel void convertIndexU16Aligned(uint idx[[thread_position_in_grid]],
+                                   constant IndexConversionParams &options[[buffer(0)]],
+                                   constant ushort *input[[buffer(1)]],
+                                   device ushort *output[[buffer(2)]])
+{
+    if (idx >= options.indexCount) { return; };
+    ushort value = input[options.srcOffset / 2 + idx];
+    output[idx] = value;
+}
+
+kernel void convertIndexU32Unaligned(uint idx[[thread_position_in_grid]],
+                                     constant IndexConversionParams &options[[buffer(0)]],
+                                     constant uchar *input[[buffer(1)]],
+                                     device uint *output[[buffer(2)]])
+{
+    if (idx >= options.indexCount) { return; };
+    uint input0 = input[options.srcOffset + 4 * idx];
+    uint input1 = input[options.srcOffset + 4 * idx + 1];
+    uint input2 = input[options.srcOffset + 4 * idx + 2];
+    uint input3 = input[options.srcOffset + 4 * idx + 3];
+
+    uint value = input0 | (input1 << 8) | (input2 << 16) | (input3 << 24);
+    output[idx] = value;
+}
+
+kernel void convertIndexU32Aligned(uint idx[[thread_position_in_grid]],
+                                   constant IndexConversionParams &options[[buffer(0)]],
+                                   constant uint *input[[buffer(1)]],
+                                   device uint *output[[buffer(2)]])
+{
+    if (idx >= options.indexCount) { return; };
+    uint value = input[options.srcOffset / 4 + idx];
+    output[idx] = value;
+}
+# 12 "master_source.metal" 2
+
+)";