2021-12-03 03:41:41 +11:00
|
|
|
#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];
|
|
|
|
}
|
|
|
|
};
|
|
|
|
|
2021-12-03 10:07:33 +11:00
|
|
|
struct Alloc
|
|
|
|
{
|
|
|
|
uint offset;
|
|
|
|
};
|
|
|
|
|
2021-12-03 03:41:41 +11:00
|
|
|
struct ElementRef
|
|
|
|
{
|
|
|
|
uint offset;
|
|
|
|
};
|
|
|
|
|
2021-12-03 10:07:33 +11:00
|
|
|
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;
|
|
|
|
};
|
|
|
|
|
2021-12-03 03:41:41 +11:00
|
|
|
struct ElementTag
|
|
|
|
{
|
|
|
|
uint tag;
|
|
|
|
uint flags;
|
|
|
|
};
|
|
|
|
|
|
|
|
struct DrawMonoid
|
|
|
|
{
|
|
|
|
uint path_ix;
|
|
|
|
uint clip_ix;
|
|
|
|
};
|
|
|
|
|
2021-12-03 10:07:33 +11:00
|
|
|
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];
|
|
|
|
};
|
|
|
|
|
2021-12-03 03:41:41 +11:00
|
|
|
struct SceneBuf
|
|
|
|
{
|
|
|
|
uint scene[1];
|
|
|
|
};
|
|
|
|
|
|
|
|
struct DrawMonoid_1
|
|
|
|
{
|
|
|
|
uint path_ix;
|
|
|
|
uint clip_ix;
|
|
|
|
};
|
|
|
|
|
|
|
|
struct ParentBuf
|
|
|
|
{
|
|
|
|
DrawMonoid_1 parent[1];
|
|
|
|
};
|
|
|
|
|
2021-12-03 10:07:33 +11:00
|
|
|
struct Alloc_1
|
2021-12-03 03:41:41 +11:00
|
|
|
{
|
|
|
|
uint offset;
|
|
|
|
};
|
|
|
|
|
|
|
|
struct Config
|
|
|
|
{
|
|
|
|
uint n_elements;
|
|
|
|
uint n_pathseg;
|
|
|
|
uint width_in_tiles;
|
|
|
|
uint height_in_tiles;
|
2021-12-03 10:07:33 +11:00
|
|
|
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;
|
2021-12-03 03:41:41 +11:00
|
|
|
uint n_trans;
|
2021-12-03 10:07:33 +11:00
|
|
|
uint n_path;
|
2021-12-03 03:41:41 +11:00
|
|
|
uint trans_offset;
|
|
|
|
uint linewidth_offset;
|
2021-12-03 10:07:33 +11:00
|
|
|
uint pathtag_offset;
|
2021-12-03 03:41:41 +11:00
|
|
|
uint pathseg_offset;
|
|
|
|
};
|
|
|
|
|
|
|
|
struct ConfigBuf
|
|
|
|
{
|
|
|
|
Config conf;
|
|
|
|
};
|
|
|
|
|
|
|
|
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(512u, 1u, 1u);
|
|
|
|
|
|
|
|
static inline __attribute__((always_inline))
|
2021-12-03 10:07:33 +11:00
|
|
|
ElementTag Element_tag(thread const ElementRef& ref, const device SceneBuf& v_225)
|
2021-12-03 03:41:41 +11:00
|
|
|
{
|
2021-12-03 10:07:33 +11:00
|
|
|
uint tag_and_flags = v_225.scene[ref.offset >> uint(2)];
|
2021-12-03 03:41:41 +11:00
|
|
|
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:
|
|
|
|
{
|
|
|
|
return DrawMonoid{ 1u, 1u };
|
|
|
|
}
|
|
|
|
case 10u:
|
|
|
|
{
|
|
|
|
return DrawMonoid{ 0u, 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 };
|
|
|
|
}
|
|
|
|
|
2021-12-03 10:07:33 +11:00
|
|
|
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& _1042 [[buffer(1)]], const device SceneBuf& v_225 [[buffer(2)]], const device ParentBuf& _1008 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
|
2021-12-03 03:41:41 +11:00
|
|
|
{
|
|
|
|
threadgroup DrawMonoid sh_scratch[512];
|
|
|
|
uint ix = gl_GlobalInvocationID.x * 8u;
|
|
|
|
ElementRef ref = ElementRef{ ix * 36u };
|
|
|
|
ElementRef param = ref;
|
2021-12-03 10:07:33 +11:00
|
|
|
uint tag_word = Element_tag(param, v_225).tag;
|
2021-12-03 03:41:41 +11:00
|
|
|
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);
|
2021-12-03 10:07:33 +11:00
|
|
|
tag_word = Element_tag(param_4, v_225).tag;
|
2021-12-03 03:41:41 +11:00
|
|
|
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 < 9u; 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)
|
|
|
|
{
|
2021-12-03 10:07:33 +11:00
|
|
|
uint _1011 = gl_WorkGroupID.x - 1u;
|
|
|
|
row.path_ix = _1008.parent[_1011].path_ix;
|
|
|
|
row.clip_ix = _1008.parent[_1011].clip_ix;
|
2021-12-03 03:41:41 +11:00
|
|
|
}
|
|
|
|
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);
|
|
|
|
}
|
2021-12-03 10:07:33 +11:00
|
|
|
uint out_ix = gl_GlobalInvocationID.x * 8u;
|
|
|
|
uint out_base = (_1042.conf.drawmonoid_alloc.offset >> uint(2)) + (out_ix * 2u);
|
|
|
|
AnnotatedRef out_ref = AnnotatedRef{ _1042.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;
|
2021-12-03 03:41:41 +11:00
|
|
|
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);
|
2021-12-03 10:07:33 +11:00
|
|
|
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 = (_1042.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 = (_1042.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 = _1042.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 = _1042.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 = _1042.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 = _1042.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 = _1042.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;
|
2021-12-03 03:41:41 +11:00
|
|
|
}
|
|
|
|
}
|
|
|
|
|