| // 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 |
| # 22 "./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 dstFlipViewportY; |
| 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.dstFlipViewportY) |
| { |
| |
| |
| output.position.y = -output.position.y; |
| } |
| |
| 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; |
| } |
| |
| 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 |
| |
| |
| |
| |
| |
| |
| |
| |
| constant bool kSourceBufferAligned[[function_constant(0)]]; |
| constant bool kSourceIndexIsU8[[function_constant(1)]]; |
| constant bool kSourceIndexIsU16[[function_constant(2)]]; |
| constant bool kSourceIndexIsU32[[function_constant(3)]]; |
| constant bool kSourceBufferUnaligned = !kSourceBufferAligned; |
| constant bool kUseSourceBufferU8 = kSourceIndexIsU8 || kSourceBufferUnaligned; |
| constant bool kUseSourceBufferU16 = kSourceIndexIsU16 && kSourceBufferAligned; |
| constant bool kUseSourceBufferU32 = kSourceIndexIsU32 && kSourceBufferAligned; |
| |
| struct IndexConversionParams |
| { |
| uint32_t srcOffset; |
| uint32_t indexCount; |
| }; |
| |
| |
| |
| inline ushort getIndexAligned(constant ushort *inputAligned, uint offset, uint idx) |
| { |
| return inputAligned[offset / 2 + idx]; |
| } |
| inline uint getIndexAligned(constant uint *inputAligned, uint offset, uint idx) |
| { |
| return inputAligned[offset / 4 + idx]; |
| } |
| inline uchar getIndexAligned(constant uchar *input, uint offset, uint idx) |
| { |
| return input[offset + idx]; |
| } |
| inline ushort getIndexUnalignedU16(constant uchar *input, uint offset, uint idx) |
| { |
| ushort inputLo = input[offset + 2 * idx]; |
| ushort inputHi = input[offset + 2 * idx + 1]; |
| |
| return inputLo | (inputHi << 8); |
| } |
| inline uint getIndexUnalignedU32(constant uchar *input, uint offset, uint idx) |
| { |
| uint input0 = input[offset + 4 * idx]; |
| uint input1 = input[offset + 4 * idx + 1]; |
| uint input2 = input[offset + 4 * idx + 2]; |
| uint input3 = input[offset + 4 * idx + 3]; |
| |
| return input0 | (input1 << 8) | (input2 << 16) | (input3 << 24); |
| } |
| |
| 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] = getIndexAligned(input, options.srcOffset, idx); |
| } |
| |
| kernel void convertIndexU16( |
| uint idx[[thread_position_in_grid]], |
| constant IndexConversionParams &options[[buffer(0)]], |
| constant uchar *input[[ buffer(1), function_constant(kSourceBufferUnaligned) ]], |
| constant ushort *inputAligned[[ buffer(1), function_constant(kSourceBufferAligned) ]], |
| device ushort *output[[buffer(2)]]) |
| { |
| if (idx >= options.indexCount) { return; }; |
| |
| ushort value; |
| if (kSourceBufferAligned) |
| { |
| value = getIndexAligned(inputAligned, options.srcOffset, idx); |
| } |
| else |
| { |
| value = getIndexUnalignedU16(input, options.srcOffset, idx); |
| } |
| output[idx] = value; |
| } |
| |
| kernel void convertIndexU32( |
| uint idx[[thread_position_in_grid]], |
| constant IndexConversionParams &options[[buffer(0)]], |
| constant uchar *input[[ buffer(1), function_constant(kSourceBufferUnaligned) ]], |
| constant uint *inputAligned[[ buffer(1), function_constant(kSourceBufferAligned) ]], |
| device uint *output[[buffer(2)]]) |
| { |
| if (idx >= options.indexCount) { return; }; |
| |
| uint value; |
| if (kSourceBufferAligned) |
| { |
| value = getIndexAligned(inputAligned, options.srcOffset, idx); |
| } |
| else |
| { |
| value = getIndexUnalignedU32(input, options.srcOffset, idx); |
| } |
| output[idx] = value; |
| } |
| |
| struct TriFanArrayParams |
| { |
| uint firstVertex; |
| uint vertexCountFrom3rd; |
| }; |
| kernel void genTriFanIndicesFromArray(uint idx[[thread_position_in_grid]], |
| constant TriFanArrayParams &options[[buffer(0)]], |
| device uint *output[[buffer(2)]]) |
| { |
| if (idx >= options.vertexCountFrom3rd) { return; }; |
| |
| uint vertexIdx = options.firstVertex + 2 + idx; |
| |
| output[3 * idx] = options.firstVertex; |
| output[3 * idx + 1] = vertexIdx - 1; |
| output[3 * idx + 2] = vertexIdx; |
| } |
| |
| inline uint getIndexU32(uint offset, |
| uint idx, |
| constant uchar *inputU8[[function_constant(kUseSourceBufferU8)]], |
| constant ushort *inputU16[[function_constant(kUseSourceBufferU16)]], |
| constant uint *inputU32[[function_constant(kUseSourceBufferU32)]]) |
| { |
| if (kUseSourceBufferU8) |
| { |
| if (kSourceIndexIsU16) |
| { |
| return getIndexUnalignedU16(inputU8, offset, idx); |
| } |
| else if (kSourceIndexIsU32) |
| { |
| return getIndexUnalignedU32(inputU8, offset, idx); |
| } |
| return getIndexAligned(inputU8, offset, idx); |
| } |
| else if (kUseSourceBufferU16) |
| { |
| return getIndexAligned(inputU16, offset, idx); |
| } |
| else if (kUseSourceBufferU32) |
| { |
| return getIndexAligned(inputU32, offset, idx); |
| } |
| return 0; |
| } |
| |
| |
| |
| kernel void genTriFanIndicesFromElements( |
| uint idx[[thread_position_in_grid]], |
| constant IndexConversionParams &options[[buffer(0)]], |
| constant uchar *inputU8[[ buffer(1), function_constant(kUseSourceBufferU8) ]], |
| constant ushort *inputU16[[ buffer(1), function_constant(kUseSourceBufferU16) ]], |
| constant uint *inputU32[[ buffer(1), function_constant(kUseSourceBufferU32) ]], |
| device uint *output[[buffer(2)]]) |
| { |
| if (idx >= options.indexCount) { return; }; |
| |
| uint elemIdx = 2 + idx; |
| |
| output[3 * idx] = getIndexU32(options.srcOffset, 0, inputU8, inputU16, inputU32); |
| output[3 * idx + 1] = getIndexU32(options.srcOffset, elemIdx - 1, inputU8, inputU16, inputU32); |
| output[3 * idx + 2] = getIndexU32(options.srcOffset, elemIdx, inputU8, inputU16, inputU32); |
| } |
| # 12 "master_source.metal" 2 |
| |
| )"; |