| #pragma clang diagnostic ignored "-Wmissing-prototypes" |
| #pragma clang diagnostic ignored "-Wmissing-braces" |
| |
| #include <metal_stdlib> |
| #include <simd/simd.h> |
| |
| using namespace metal; |
| |
| template<typename T, size_t Num> |
| struct spvUnsafeArray |
| { |
| T elements[Num ? Num : 1]; |
| |
| thread T& operator [] (size_t pos) thread |
| { |
| return elements[pos]; |
| } |
| constexpr const thread T& operator [] (size_t pos) const thread |
| { |
| return elements[pos]; |
| } |
| |
| device T& operator [] (size_t pos) device |
| { |
| return elements[pos]; |
| } |
| constexpr const device T& operator [] (size_t pos) const device |
| { |
| return elements[pos]; |
| } |
| |
| constexpr const constant T& operator [] (size_t pos) const constant |
| { |
| return elements[pos]; |
| } |
| |
| threadgroup T& operator [] (size_t pos) threadgroup |
| { |
| return elements[pos]; |
| } |
| constexpr const threadgroup T& operator [] (size_t pos) const threadgroup |
| { |
| return elements[pos]; |
| } |
| }; |
| |
| struct Alloc |
| { |
| uint offset; |
| }; |
| |
| struct ElementRef |
| { |
| uint offset; |
| }; |
| |
| struct FillColorRef |
| { |
| uint offset; |
| }; |
| |
| struct FillColor |
| { |
| uint rgba_color; |
| }; |
| |
| struct FillLinGradientRef |
| { |
| uint offset; |
| }; |
| |
| struct FillLinGradient |
| { |
| uint index; |
| float2 p0; |
| float2 p1; |
| }; |
| |
| struct FillImageRef |
| { |
| uint offset; |
| }; |
| |
| struct FillImage |
| { |
| uint index; |
| int2 offset; |
| }; |
| |
| struct ClipRef |
| { |
| uint offset; |
| }; |
| |
| struct Clip |
| { |
| float4 bbox; |
| }; |
| |
| struct ElementTag |
| { |
| uint tag; |
| uint flags; |
| }; |
| |
| struct DrawMonoid |
| { |
| uint path_ix; |
| uint clip_ix; |
| }; |
| |
| struct AnnoImageRef |
| { |
| uint offset; |
| }; |
| |
| struct AnnoImage |
| { |
| float4 bbox; |
| float linewidth; |
| uint index; |
| int2 offset; |
| }; |
| |
| struct AnnoColorRef |
| { |
| uint offset; |
| }; |
| |
| struct AnnoColor |
| { |
| float4 bbox; |
| float linewidth; |
| uint rgba_color; |
| }; |
| |
| struct AnnoLinGradientRef |
| { |
| uint offset; |
| }; |
| |
| struct AnnoLinGradient |
| { |
| float4 bbox; |
| float linewidth; |
| uint index; |
| float line_x; |
| float line_y; |
| float line_c; |
| }; |
| |
| struct AnnoBeginClipRef |
| { |
| uint offset; |
| }; |
| |
| struct AnnoBeginClip |
| { |
| float4 bbox; |
| float linewidth; |
| }; |
| |
| struct AnnoEndClipRef |
| { |
| uint offset; |
| }; |
| |
| struct AnnoEndClip |
| { |
| float4 bbox; |
| }; |
| |
| struct AnnotatedRef |
| { |
| uint offset; |
| }; |
| |
| struct Memory |
| { |
| uint mem_offset; |
| uint mem_error; |
| uint memory[1]; |
| }; |
| |
| struct SceneBuf |
| { |
| uint scene[1]; |
| }; |
| |
| struct DrawMonoid_1 |
| { |
| uint path_ix; |
| uint clip_ix; |
| }; |
| |
| struct ParentBuf |
| { |
| DrawMonoid_1 parent[1]; |
| }; |
| |
| struct Alloc_1 |
| { |
| uint offset; |
| }; |
| |
| struct Config |
| { |
| uint n_elements; |
| uint n_pathseg; |
| uint width_in_tiles; |
| uint height_in_tiles; |
| Alloc_1 tile_alloc; |
| Alloc_1 bin_alloc; |
| Alloc_1 ptcl_alloc; |
| Alloc_1 pathseg_alloc; |
| Alloc_1 anno_alloc; |
| Alloc_1 trans_alloc; |
| Alloc_1 bbox_alloc; |
| Alloc_1 drawmonoid_alloc; |
| uint n_trans; |
| uint n_path; |
| uint trans_offset; |
| uint linewidth_offset; |
| uint pathtag_offset; |
| uint pathseg_offset; |
| }; |
| |
| struct ConfigBuf |
| { |
| Config conf; |
| }; |
| |
| constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u); |
| |
| static inline __attribute__((always_inline)) |
| ElementTag Element_tag(thread const ElementRef& ref, const device SceneBuf& v_225) |
| { |
| uint tag_and_flags = v_225.scene[ref.offset >> uint(2)]; |
| return ElementTag{ tag_and_flags & 65535u, tag_and_flags >> uint(16) }; |
| } |
| |
| static inline __attribute__((always_inline)) |
| DrawMonoid map_tag(thread const uint& tag_word) |
| { |
| switch (tag_word) |
| { |
| case 4u: |
| case 5u: |
| case 6u: |
| { |
| return DrawMonoid{ 1u, 0u }; |
| } |
| case 9u: |
| case 10u: |
| { |
| return DrawMonoid{ 1u, 1u }; |
| } |
| default: |
| { |
| return DrawMonoid{ 0u, 0u }; |
| } |
| } |
| } |
| |
| static inline __attribute__((always_inline)) |
| ElementRef Element_index(thread const ElementRef& ref, thread const uint& index) |
| { |
| return ElementRef{ ref.offset + (index * 36u) }; |
| } |
| |
| static inline __attribute__((always_inline)) |
| DrawMonoid combine_tag_monoid(thread const DrawMonoid& a, thread const DrawMonoid& b) |
| { |
| DrawMonoid c; |
| c.path_ix = a.path_ix + b.path_ix; |
| c.clip_ix = a.clip_ix + b.clip_ix; |
| return c; |
| } |
| |
| static inline __attribute__((always_inline)) |
| DrawMonoid tag_monoid_identity() |
| { |
| return DrawMonoid{ 0u, 0u }; |
| } |
| |
| static inline __attribute__((always_inline)) |
| FillColor FillColor_read(thread const FillColorRef& ref, const device SceneBuf& v_225) |
| { |
| uint ix = ref.offset >> uint(2); |
| uint raw0 = v_225.scene[ix + 0u]; |
| FillColor s; |
| s.rgba_color = raw0; |
| return s; |
| } |
| |
| static inline __attribute__((always_inline)) |
| FillColor Element_FillColor_read(thread const ElementRef& ref, const device SceneBuf& v_225) |
| { |
| FillColorRef param = FillColorRef{ ref.offset + 4u }; |
| return FillColor_read(param, v_225); |
| } |
| |
| static inline __attribute__((always_inline)) |
| bool touch_mem(thread const Alloc& alloc, thread const uint& offset) |
| { |
| return true; |
| } |
| |
| static inline __attribute__((always_inline)) |
| void write_mem(thread const Alloc& alloc, thread const uint& offset, thread const uint& val, device Memory& v_201) |
| { |
| Alloc param = alloc; |
| uint param_1 = offset; |
| if (!touch_mem(param, param_1)) |
| { |
| return; |
| } |
| v_201.memory[offset] = val; |
| } |
| |
| static inline __attribute__((always_inline)) |
| void AnnoColor_write(thread const Alloc& a, thread const AnnoColorRef& ref, thread const AnnoColor& s, device Memory& v_201) |
| { |
| uint ix = ref.offset >> uint(2); |
| Alloc param = a; |
| uint param_1 = ix + 0u; |
| uint param_2 = as_type<uint>(s.bbox.x); |
| write_mem(param, param_1, param_2, v_201); |
| Alloc param_3 = a; |
| uint param_4 = ix + 1u; |
| uint param_5 = as_type<uint>(s.bbox.y); |
| write_mem(param_3, param_4, param_5, v_201); |
| Alloc param_6 = a; |
| uint param_7 = ix + 2u; |
| uint param_8 = as_type<uint>(s.bbox.z); |
| write_mem(param_6, param_7, param_8, v_201); |
| Alloc param_9 = a; |
| uint param_10 = ix + 3u; |
| uint param_11 = as_type<uint>(s.bbox.w); |
| write_mem(param_9, param_10, param_11, v_201); |
| Alloc param_12 = a; |
| uint param_13 = ix + 4u; |
| uint param_14 = as_type<uint>(s.linewidth); |
| write_mem(param_12, param_13, param_14, v_201); |
| Alloc param_15 = a; |
| uint param_16 = ix + 5u; |
| uint param_17 = s.rgba_color; |
| write_mem(param_15, param_16, param_17, v_201); |
| } |
| |
| static inline __attribute__((always_inline)) |
| void Annotated_Color_write(thread const Alloc& a, thread const AnnotatedRef& ref, thread const uint& flags, thread const AnnoColor& s, device Memory& v_201) |
| { |
| Alloc param = a; |
| uint param_1 = ref.offset >> uint(2); |
| uint param_2 = (flags << uint(16)) | 1u; |
| write_mem(param, param_1, param_2, v_201); |
| Alloc param_3 = a; |
| AnnoColorRef param_4 = AnnoColorRef{ ref.offset + 4u }; |
| AnnoColor param_5 = s; |
| AnnoColor_write(param_3, param_4, param_5, v_201); |
| } |
| |
| static inline __attribute__((always_inline)) |
| FillLinGradient FillLinGradient_read(thread const FillLinGradientRef& ref, const device SceneBuf& v_225) |
| { |
| uint ix = ref.offset >> uint(2); |
| uint raw0 = v_225.scene[ix + 0u]; |
| uint raw1 = v_225.scene[ix + 1u]; |
| uint raw2 = v_225.scene[ix + 2u]; |
| uint raw3 = v_225.scene[ix + 3u]; |
| uint raw4 = v_225.scene[ix + 4u]; |
| FillLinGradient s; |
| s.index = raw0; |
| s.p0 = float2(as_type<float>(raw1), as_type<float>(raw2)); |
| s.p1 = float2(as_type<float>(raw3), as_type<float>(raw4)); |
| return s; |
| } |
| |
| static inline __attribute__((always_inline)) |
| FillLinGradient Element_FillLinGradient_read(thread const ElementRef& ref, const device SceneBuf& v_225) |
| { |
| FillLinGradientRef param = FillLinGradientRef{ ref.offset + 4u }; |
| return FillLinGradient_read(param, v_225); |
| } |
| |
| static inline __attribute__((always_inline)) |
| void AnnoLinGradient_write(thread const Alloc& a, thread const AnnoLinGradientRef& ref, thread const AnnoLinGradient& s, device Memory& v_201) |
| { |
| uint ix = ref.offset >> uint(2); |
| Alloc param = a; |
| uint param_1 = ix + 0u; |
| uint param_2 = as_type<uint>(s.bbox.x); |
| write_mem(param, param_1, param_2, v_201); |
| Alloc param_3 = a; |
| uint param_4 = ix + 1u; |
| uint param_5 = as_type<uint>(s.bbox.y); |
| write_mem(param_3, param_4, param_5, v_201); |
| Alloc param_6 = a; |
| uint param_7 = ix + 2u; |
| uint param_8 = as_type<uint>(s.bbox.z); |
| write_mem(param_6, param_7, param_8, v_201); |
| Alloc param_9 = a; |
| uint param_10 = ix + 3u; |
| uint param_11 = as_type<uint>(s.bbox.w); |
| write_mem(param_9, param_10, param_11, v_201); |
| Alloc param_12 = a; |
| uint param_13 = ix + 4u; |
| uint param_14 = as_type<uint>(s.linewidth); |
| write_mem(param_12, param_13, param_14, v_201); |
| Alloc param_15 = a; |
| uint param_16 = ix + 5u; |
| uint param_17 = s.index; |
| write_mem(param_15, param_16, param_17, v_201); |
| Alloc param_18 = a; |
| uint param_19 = ix + 6u; |
| uint param_20 = as_type<uint>(s.line_x); |
| write_mem(param_18, param_19, param_20, v_201); |
| Alloc param_21 = a; |
| uint param_22 = ix + 7u; |
| uint param_23 = as_type<uint>(s.line_y); |
| write_mem(param_21, param_22, param_23, v_201); |
| Alloc param_24 = a; |
| uint param_25 = ix + 8u; |
| uint param_26 = as_type<uint>(s.line_c); |
| write_mem(param_24, param_25, param_26, v_201); |
| } |
| |
| static inline __attribute__((always_inline)) |
| void Annotated_LinGradient_write(thread const Alloc& a, thread const AnnotatedRef& ref, thread const uint& flags, thread const AnnoLinGradient& s, device Memory& v_201) |
| { |
| Alloc param = a; |
| uint param_1 = ref.offset >> uint(2); |
| uint param_2 = (flags << uint(16)) | 2u; |
| write_mem(param, param_1, param_2, v_201); |
| Alloc param_3 = a; |
| AnnoLinGradientRef param_4 = AnnoLinGradientRef{ ref.offset + 4u }; |
| AnnoLinGradient param_5 = s; |
| AnnoLinGradient_write(param_3, param_4, param_5, v_201); |
| } |
| |
| static inline __attribute__((always_inline)) |
| FillImage FillImage_read(thread const FillImageRef& ref, const device SceneBuf& v_225) |
| { |
| uint ix = ref.offset >> uint(2); |
| uint raw0 = v_225.scene[ix + 0u]; |
| uint raw1 = v_225.scene[ix + 1u]; |
| FillImage s; |
| s.index = raw0; |
| s.offset = int2(int(raw1 << uint(16)) >> 16, int(raw1) >> 16); |
| return s; |
| } |
| |
| static inline __attribute__((always_inline)) |
| FillImage Element_FillImage_read(thread const ElementRef& ref, const device SceneBuf& v_225) |
| { |
| FillImageRef param = FillImageRef{ ref.offset + 4u }; |
| return FillImage_read(param, v_225); |
| } |
| |
| static inline __attribute__((always_inline)) |
| void AnnoImage_write(thread const Alloc& a, thread const AnnoImageRef& ref, thread const AnnoImage& s, device Memory& v_201) |
| { |
| uint ix = ref.offset >> uint(2); |
| Alloc param = a; |
| uint param_1 = ix + 0u; |
| uint param_2 = as_type<uint>(s.bbox.x); |
| write_mem(param, param_1, param_2, v_201); |
| Alloc param_3 = a; |
| uint param_4 = ix + 1u; |
| uint param_5 = as_type<uint>(s.bbox.y); |
| write_mem(param_3, param_4, param_5, v_201); |
| Alloc param_6 = a; |
| uint param_7 = ix + 2u; |
| uint param_8 = as_type<uint>(s.bbox.z); |
| write_mem(param_6, param_7, param_8, v_201); |
| Alloc param_9 = a; |
| uint param_10 = ix + 3u; |
| uint param_11 = as_type<uint>(s.bbox.w); |
| write_mem(param_9, param_10, param_11, v_201); |
| Alloc param_12 = a; |
| uint param_13 = ix + 4u; |
| uint param_14 = as_type<uint>(s.linewidth); |
| write_mem(param_12, param_13, param_14, v_201); |
| Alloc param_15 = a; |
| uint param_16 = ix + 5u; |
| uint param_17 = s.index; |
| write_mem(param_15, param_16, param_17, v_201); |
| Alloc param_18 = a; |
| uint param_19 = ix + 6u; |
| uint param_20 = (uint(s.offset.x) & 65535u) | (uint(s.offset.y) << uint(16)); |
| write_mem(param_18, param_19, param_20, v_201); |
| } |
| |
| static inline __attribute__((always_inline)) |
| void Annotated_Image_write(thread const Alloc& a, thread const AnnotatedRef& ref, thread const uint& flags, thread const AnnoImage& s, device Memory& v_201) |
| { |
| Alloc param = a; |
| uint param_1 = ref.offset >> uint(2); |
| uint param_2 = (flags << uint(16)) | 3u; |
| write_mem(param, param_1, param_2, v_201); |
| Alloc param_3 = a; |
| AnnoImageRef param_4 = AnnoImageRef{ ref.offset + 4u }; |
| AnnoImage param_5 = s; |
| AnnoImage_write(param_3, param_4, param_5, v_201); |
| } |
| |
| static inline __attribute__((always_inline)) |
| Clip Clip_read(thread const ClipRef& ref, const device SceneBuf& v_225) |
| { |
| uint ix = ref.offset >> uint(2); |
| uint raw0 = v_225.scene[ix + 0u]; |
| uint raw1 = v_225.scene[ix + 1u]; |
| uint raw2 = v_225.scene[ix + 2u]; |
| uint raw3 = v_225.scene[ix + 3u]; |
| Clip s; |
| s.bbox = float4(as_type<float>(raw0), as_type<float>(raw1), as_type<float>(raw2), as_type<float>(raw3)); |
| return s; |
| } |
| |
| static inline __attribute__((always_inline)) |
| Clip Element_BeginClip_read(thread const ElementRef& ref, const device SceneBuf& v_225) |
| { |
| ClipRef param = ClipRef{ ref.offset + 4u }; |
| return Clip_read(param, v_225); |
| } |
| |
| static inline __attribute__((always_inline)) |
| void AnnoBeginClip_write(thread const Alloc& a, thread const AnnoBeginClipRef& ref, thread const AnnoBeginClip& s, device Memory& v_201) |
| { |
| uint ix = ref.offset >> uint(2); |
| Alloc param = a; |
| uint param_1 = ix + 0u; |
| uint param_2 = as_type<uint>(s.bbox.x); |
| write_mem(param, param_1, param_2, v_201); |
| Alloc param_3 = a; |
| uint param_4 = ix + 1u; |
| uint param_5 = as_type<uint>(s.bbox.y); |
| write_mem(param_3, param_4, param_5, v_201); |
| Alloc param_6 = a; |
| uint param_7 = ix + 2u; |
| uint param_8 = as_type<uint>(s.bbox.z); |
| write_mem(param_6, param_7, param_8, v_201); |
| Alloc param_9 = a; |
| uint param_10 = ix + 3u; |
| uint param_11 = as_type<uint>(s.bbox.w); |
| write_mem(param_9, param_10, param_11, v_201); |
| Alloc param_12 = a; |
| uint param_13 = ix + 4u; |
| uint param_14 = as_type<uint>(s.linewidth); |
| write_mem(param_12, param_13, param_14, v_201); |
| } |
| |
| static inline __attribute__((always_inline)) |
| void Annotated_BeginClip_write(thread const Alloc& a, thread const AnnotatedRef& ref, thread const uint& flags, thread const AnnoBeginClip& s, device Memory& v_201) |
| { |
| Alloc param = a; |
| uint param_1 = ref.offset >> uint(2); |
| uint param_2 = (flags << uint(16)) | 4u; |
| write_mem(param, param_1, param_2, v_201); |
| Alloc param_3 = a; |
| AnnoBeginClipRef param_4 = AnnoBeginClipRef{ ref.offset + 4u }; |
| AnnoBeginClip param_5 = s; |
| AnnoBeginClip_write(param_3, param_4, param_5, v_201); |
| } |
| |
| static inline __attribute__((always_inline)) |
| Clip Element_EndClip_read(thread const ElementRef& ref, const device SceneBuf& v_225) |
| { |
| ClipRef param = ClipRef{ ref.offset + 4u }; |
| return Clip_read(param, v_225); |
| } |
| |
| static inline __attribute__((always_inline)) |
| void AnnoEndClip_write(thread const Alloc& a, thread const AnnoEndClipRef& ref, thread const AnnoEndClip& s, device Memory& v_201) |
| { |
| uint ix = ref.offset >> uint(2); |
| Alloc param = a; |
| uint param_1 = ix + 0u; |
| uint param_2 = as_type<uint>(s.bbox.x); |
| write_mem(param, param_1, param_2, v_201); |
| Alloc param_3 = a; |
| uint param_4 = ix + 1u; |
| uint param_5 = as_type<uint>(s.bbox.y); |
| write_mem(param_3, param_4, param_5, v_201); |
| Alloc param_6 = a; |
| uint param_7 = ix + 2u; |
| uint param_8 = as_type<uint>(s.bbox.z); |
| write_mem(param_6, param_7, param_8, v_201); |
| Alloc param_9 = a; |
| uint param_10 = ix + 3u; |
| uint param_11 = as_type<uint>(s.bbox.w); |
| write_mem(param_9, param_10, param_11, v_201); |
| } |
| |
| static inline __attribute__((always_inline)) |
| void Annotated_EndClip_write(thread const Alloc& a, thread const AnnotatedRef& ref, thread const AnnoEndClip& s, device Memory& v_201) |
| { |
| Alloc param = a; |
| uint param_1 = ref.offset >> uint(2); |
| uint param_2 = 5u; |
| write_mem(param, param_1, param_2, v_201); |
| Alloc param_3 = a; |
| AnnoEndClipRef param_4 = AnnoEndClipRef{ ref.offset + 4u }; |
| AnnoEndClip param_5 = s; |
| AnnoEndClip_write(param_3, param_4, param_5, v_201); |
| } |
| |
| kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1038 [[buffer(1)]], const device SceneBuf& v_225 [[buffer(2)]], const device ParentBuf& _1004 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) |
| { |
| threadgroup DrawMonoid sh_scratch[256]; |
| uint ix = gl_GlobalInvocationID.x * 8u; |
| ElementRef ref = ElementRef{ ix * 36u }; |
| ElementRef param = ref; |
| uint tag_word = Element_tag(param, v_225).tag; |
| uint param_1 = tag_word; |
| DrawMonoid agg = map_tag(param_1); |
| spvUnsafeArray<DrawMonoid, 8> local; |
| local[0] = agg; |
| for (uint i = 1u; i < 8u; i++) |
| { |
| ElementRef param_2 = ref; |
| uint param_3 = i; |
| ElementRef param_4 = Element_index(param_2, param_3); |
| tag_word = Element_tag(param_4, v_225).tag; |
| uint param_5 = tag_word; |
| DrawMonoid param_6 = agg; |
| DrawMonoid param_7 = map_tag(param_5); |
| agg = combine_tag_monoid(param_6, param_7); |
| local[i] = agg; |
| } |
| sh_scratch[gl_LocalInvocationID.x] = agg; |
| for (uint i_1 = 0u; i_1 < 8u; i_1++) |
| { |
| threadgroup_barrier(mem_flags::mem_threadgroup); |
| if (gl_LocalInvocationID.x >= (1u << i_1)) |
| { |
| DrawMonoid other = sh_scratch[gl_LocalInvocationID.x - (1u << i_1)]; |
| DrawMonoid param_8 = other; |
| DrawMonoid param_9 = agg; |
| agg = combine_tag_monoid(param_8, param_9); |
| } |
| threadgroup_barrier(mem_flags::mem_threadgroup); |
| sh_scratch[gl_LocalInvocationID.x] = agg; |
| } |
| threadgroup_barrier(mem_flags::mem_threadgroup); |
| DrawMonoid row = tag_monoid_identity(); |
| if (gl_WorkGroupID.x > 0u) |
| { |
| uint _1007 = gl_WorkGroupID.x - 1u; |
| row.path_ix = _1004.parent[_1007].path_ix; |
| row.clip_ix = _1004.parent[_1007].clip_ix; |
| } |
| if (gl_LocalInvocationID.x > 0u) |
| { |
| DrawMonoid param_10 = row; |
| DrawMonoid param_11 = sh_scratch[gl_LocalInvocationID.x - 1u]; |
| row = combine_tag_monoid(param_10, param_11); |
| } |
| uint out_ix = gl_GlobalInvocationID.x * 8u; |
| uint out_base = (_1038.conf.drawmonoid_alloc.offset >> uint(2)) + (out_ix * 2u); |
| AnnotatedRef out_ref = AnnotatedRef{ _1038.conf.anno_alloc.offset + (out_ix * 40u) }; |
| float4 mat; |
| float2 translate; |
| AnnoColor anno_fill; |
| Alloc param_18; |
| AnnoLinGradient anno_lin; |
| Alloc param_23; |
| AnnoImage anno_img; |
| Alloc param_28; |
| AnnoBeginClip anno_begin_clip; |
| Alloc param_33; |
| AnnoEndClip anno_end_clip; |
| Alloc param_38; |
| for (uint i_2 = 0u; i_2 < 8u; i_2++) |
| { |
| DrawMonoid param_12 = row; |
| DrawMonoid param_13 = local[i_2]; |
| DrawMonoid m = combine_tag_monoid(param_12, param_13); |
| v_201.memory[out_base + (i_2 * 2u)] = m.path_ix; |
| v_201.memory[(out_base + (i_2 * 2u)) + 1u] = m.clip_ix; |
| ElementRef param_14 = ref; |
| uint param_15 = i_2; |
| ElementRef this_ref = Element_index(param_14, param_15); |
| ElementRef param_16 = this_ref; |
| tag_word = Element_tag(param_16, v_225).tag; |
| if (((tag_word == 4u) || (tag_word == 5u)) || (tag_word == 6u)) |
| { |
| uint bbox_offset = (_1038.conf.bbox_alloc.offset >> uint(2)) + (6u * (m.path_ix - 1u)); |
| float bbox_l = float(v_201.memory[bbox_offset]) - 32768.0; |
| float bbox_t = float(v_201.memory[bbox_offset + 1u]) - 32768.0; |
| float bbox_r = float(v_201.memory[bbox_offset + 2u]) - 32768.0; |
| float bbox_b = float(v_201.memory[bbox_offset + 3u]) - 32768.0; |
| float4 bbox = float4(bbox_l, bbox_t, bbox_r, bbox_b); |
| float linewidth = as_type<float>(v_201.memory[bbox_offset + 4u]); |
| uint fill_mode = uint(linewidth >= 0.0); |
| if ((linewidth >= 0.0) || (tag_word == 5u)) |
| { |
| uint trans_ix = v_201.memory[bbox_offset + 5u]; |
| uint t = (_1038.conf.trans_alloc.offset >> uint(2)) + (6u * trans_ix); |
| mat = as_type<float4>(uint4(v_201.memory[t], v_201.memory[t + 1u], v_201.memory[t + 2u], v_201.memory[t + 3u])); |
| if (tag_word == 5u) |
| { |
| translate = as_type<float2>(uint2(v_201.memory[t + 4u], v_201.memory[t + 5u])); |
| } |
| } |
| if (linewidth >= 0.0) |
| { |
| linewidth *= sqrt(abs((mat.x * mat.w) - (mat.y * mat.z))); |
| } |
| linewidth = fast::max(linewidth, 0.0); |
| switch (tag_word) |
| { |
| case 4u: |
| { |
| ElementRef param_17 = this_ref; |
| FillColor fill = Element_FillColor_read(param_17, v_225); |
| anno_fill.bbox = bbox; |
| anno_fill.linewidth = linewidth; |
| anno_fill.rgba_color = fill.rgba_color; |
| param_18.offset = _1038.conf.anno_alloc.offset; |
| AnnotatedRef param_19 = out_ref; |
| uint param_20 = fill_mode; |
| AnnoColor param_21 = anno_fill; |
| Annotated_Color_write(param_18, param_19, param_20, param_21, v_201); |
| break; |
| } |
| case 5u: |
| { |
| ElementRef param_22 = this_ref; |
| FillLinGradient lin = Element_FillLinGradient_read(param_22, v_225); |
| anno_lin.bbox = bbox; |
| anno_lin.linewidth = linewidth; |
| anno_lin.index = lin.index; |
| float2 p0 = ((mat.xy * lin.p0.x) + (mat.zw * lin.p0.y)) + translate; |
| float2 p1 = ((mat.xy * lin.p1.x) + (mat.zw * lin.p1.y)) + translate; |
| float2 dxy = p1 - p0; |
| float scale = 1.0 / ((dxy.x * dxy.x) + (dxy.y * dxy.y)); |
| float line_x = dxy.x * scale; |
| float line_y = dxy.y * scale; |
| anno_lin.line_x = line_x; |
| anno_lin.line_y = line_y; |
| anno_lin.line_c = -((p0.x * line_x) + (p0.y * line_y)); |
| param_23.offset = _1038.conf.anno_alloc.offset; |
| AnnotatedRef param_24 = out_ref; |
| uint param_25 = fill_mode; |
| AnnoLinGradient param_26 = anno_lin; |
| Annotated_LinGradient_write(param_23, param_24, param_25, param_26, v_201); |
| break; |
| } |
| case 6u: |
| { |
| ElementRef param_27 = this_ref; |
| FillImage fill_img = Element_FillImage_read(param_27, v_225); |
| anno_img.bbox = bbox; |
| anno_img.linewidth = linewidth; |
| anno_img.index = fill_img.index; |
| anno_img.offset = fill_img.offset; |
| param_28.offset = _1038.conf.anno_alloc.offset; |
| AnnotatedRef param_29 = out_ref; |
| uint param_30 = fill_mode; |
| AnnoImage param_31 = anno_img; |
| Annotated_Image_write(param_28, param_29, param_30, param_31, v_201); |
| break; |
| } |
| } |
| } |
| else |
| { |
| if (tag_word == 9u) |
| { |
| ElementRef param_32 = this_ref; |
| Clip begin_clip = Element_BeginClip_read(param_32, v_225); |
| anno_begin_clip.bbox = begin_clip.bbox; |
| anno_begin_clip.linewidth = 0.0; |
| param_33.offset = _1038.conf.anno_alloc.offset; |
| AnnotatedRef param_34 = out_ref; |
| uint param_35 = 0u; |
| AnnoBeginClip param_36 = anno_begin_clip; |
| Annotated_BeginClip_write(param_33, param_34, param_35, param_36, v_201); |
| } |
| else |
| { |
| if (tag_word == 10u) |
| { |
| ElementRef param_37 = this_ref; |
| Clip end_clip = Element_EndClip_read(param_37, v_225); |
| anno_end_clip.bbox = end_clip.bbox; |
| param_38.offset = _1038.conf.anno_alloc.offset; |
| AnnotatedRef param_39 = out_ref; |
| AnnoEndClip param_40 = anno_end_clip; |
| Annotated_EndClip_write(param_38, param_39, param_40, v_201); |
| } |
| } |
| } |
| out_ref.offset += 40u; |
| } |
| } |
| |
