diff --git a/piet-gpu/shader/drawtag.h b/piet-gpu/shader/drawtag.h index a9e8a1d..17105f6 100644 --- a/piet-gpu/shader/drawtag.h +++ b/piet-gpu/shader/drawtag.h @@ -26,9 +26,10 @@ DrawMonoid map_tag(uint tag_word) { case Element_FillImage: return DrawMonoid(1, 0); case Element_BeginClip: - return DrawMonoid(1, 1); + // TODO: endclip should be (0, 1), ie not generate a path. But for now + // we generate a dummy path. case Element_EndClip: - return DrawMonoid(0, 1); + return DrawMonoid(1, 1); default: return DrawMonoid(0, 0); } diff --git a/piet-gpu/shader/gen/draw_leaf.dxil b/piet-gpu/shader/gen/draw_leaf.dxil index f95e5bc..86b37e9 100644 Binary files a/piet-gpu/shader/gen/draw_leaf.dxil and b/piet-gpu/shader/gen/draw_leaf.dxil differ diff --git a/piet-gpu/shader/gen/draw_leaf.hlsl b/piet-gpu/shader/gen/draw_leaf.hlsl index 0dec2cd..e3cb387 100644 --- a/piet-gpu/shader/gen/draw_leaf.hlsl +++ b/piet-gpu/shader/gen/draw_leaf.hlsl @@ -154,14 +154,13 @@ struct Config static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u); static const DrawMonoid _418 = { 0u, 0u }; -static const DrawMonoid _443 = { 1u, 0u }; -static const DrawMonoid _445 = { 1u, 1u }; -static const DrawMonoid _447 = { 0u, 1u }; +static const DrawMonoid _442 = { 1u, 0u }; +static const DrawMonoid _444 = { 1u, 1u }; RWByteAddressBuffer _201 : register(u0, space0); ByteAddressBuffer _225 : register(t2, space0); -ByteAddressBuffer _1008 : register(t3, space0); -ByteAddressBuffer _1042 : register(t1, space0); +ByteAddressBuffer _1005 : register(t3, space0); +ByteAddressBuffer _1039 : register(t1, space0); static uint3 gl_WorkGroupID; static uint3 gl_LocalInvocationID; @@ -190,15 +189,12 @@ DrawMonoid map_tag(uint tag_word) case 5u: case 6u: { - return _443; + return _442; } case 9u: - { - return _445; - } case 10u: { - return _447; + return _444; } default: { @@ -293,9 +289,9 @@ void Annotated_Color_write(Alloc a, AnnotatedRef ref, uint flags, AnnoColor s) uint param_1 = ref.offset >> uint(2); uint param_2 = (flags << uint(16)) | 1u; write_mem(param, param_1, param_2); - AnnoColorRef _808 = { ref.offset + 4u }; + AnnoColorRef _805 = { ref.offset + 4u }; Alloc param_3 = a; - AnnoColorRef param_4 = _808; + AnnoColorRef param_4 = _805; AnnoColor param_5 = s; AnnoColor_write(param_3, param_4, param_5); } @@ -369,9 +365,9 @@ void Annotated_LinGradient_write(Alloc a, AnnotatedRef ref, uint flags, AnnoLinG uint param_1 = ref.offset >> uint(2); uint param_2 = (flags << uint(16)) | 2u; write_mem(param, param_1, param_2); - AnnoLinGradientRef _829 = { ref.offset + 4u }; + AnnoLinGradientRef _826 = { ref.offset + 4u }; Alloc param_3 = a; - AnnoLinGradientRef param_4 = _829; + AnnoLinGradientRef param_4 = _826; AnnoLinGradient param_5 = s; AnnoLinGradient_write(param_3, param_4, param_5); } @@ -433,9 +429,9 @@ void Annotated_Image_write(Alloc a, AnnotatedRef ref, uint flags, AnnoImage s) uint param_1 = ref.offset >> uint(2); uint param_2 = (flags << uint(16)) | 3u; write_mem(param, param_1, param_2); - AnnoImageRef _850 = { ref.offset + 4u }; + AnnoImageRef _847 = { ref.offset + 4u }; Alloc param_3 = a; - AnnoImageRef param_4 = _850; + AnnoImageRef param_4 = _847; AnnoImage param_5 = s; AnnoImage_write(param_3, param_4, param_5); } @@ -490,9 +486,9 @@ void Annotated_BeginClip_write(Alloc a, AnnotatedRef ref, uint flags, AnnoBeginC uint param_1 = ref.offset >> uint(2); uint param_2 = (flags << uint(16)) | 4u; write_mem(param, param_1, param_2); - AnnoBeginClipRef _871 = { ref.offset + 4u }; + AnnoBeginClipRef _868 = { ref.offset + 4u }; Alloc param_3 = a; - AnnoBeginClipRef param_4 = _871; + AnnoBeginClipRef param_4 = _868; AnnoBeginClip param_5 = s; AnnoBeginClip_write(param_3, param_4, param_5); } @@ -531,9 +527,9 @@ void Annotated_EndClip_write(Alloc a, AnnotatedRef ref, AnnoEndClip s) uint param_1 = ref.offset >> uint(2); uint param_2 = 5u; write_mem(param, param_1, param_2); - AnnoEndClipRef _889 = { ref.offset + 4u }; + AnnoEndClipRef _886 = { ref.offset + 4u }; Alloc param_3 = a; - AnnoEndClipRef param_4 = _889; + AnnoEndClipRef param_4 = _886; AnnoEndClip param_5 = s; AnnoEndClip_write(param_3, param_4, param_5); } @@ -541,8 +537,8 @@ void Annotated_EndClip_write(Alloc a, AnnotatedRef ref, AnnoEndClip s) void comp_main() { uint ix = gl_GlobalInvocationID.x * 8u; - ElementRef _907 = { ix * 36u }; - ElementRef ref = _907; + ElementRef _904 = { ix * 36u }; + ElementRef ref = _904; ElementRef param = ref; uint tag_word = Element_tag(param).tag; uint param_1 = tag_word; @@ -579,11 +575,11 @@ void comp_main() DrawMonoid row = tag_monoid_identity(); if (gl_WorkGroupID.x > 0u) { - DrawMonoid _1014; - _1014.path_ix = _1008.Load((gl_WorkGroupID.x - 1u) * 8 + 0); - _1014.clip_ix = _1008.Load((gl_WorkGroupID.x - 1u) * 8 + 4); - row.path_ix = _1014.path_ix; - row.clip_ix = _1014.clip_ix; + DrawMonoid _1011; + _1011.path_ix = _1005.Load((gl_WorkGroupID.x - 1u) * 8 + 0); + _1011.clip_ix = _1005.Load((gl_WorkGroupID.x - 1u) * 8 + 4); + row.path_ix = _1011.path_ix; + row.clip_ix = _1011.clip_ix; } if (gl_LocalInvocationID.x > 0u) { @@ -592,9 +588,9 @@ void comp_main() row = combine_tag_monoid(param_10, param_11); } uint out_ix = gl_GlobalInvocationID.x * 8u; - uint out_base = (_1042.Load(44) >> uint(2)) + (out_ix * 2u); - AnnotatedRef _1058 = { _1042.Load(32) + (out_ix * 40u) }; - AnnotatedRef out_ref = _1058; + uint out_base = (_1039.Load(44) >> uint(2)) + (out_ix * 2u); + AnnotatedRef _1055 = { _1039.Load(32) + (out_ix * 40u) }; + AnnotatedRef out_ref = _1055; float4 mat; float2 translate; AnnoColor anno_fill; @@ -621,7 +617,7 @@ void comp_main() tag_word = Element_tag(param_16).tag; if (((tag_word == 4u) || (tag_word == 5u)) || (tag_word == 6u)) { - uint bbox_offset = (_1042.Load(40) >> uint(2)) + (6u * (m.path_ix - 1u)); + uint bbox_offset = (_1039.Load(40) >> uint(2)) + (6u * (m.path_ix - 1u)); float bbox_l = float(_201.Load(bbox_offset * 4 + 8)) - 32768.0f; float bbox_t = float(_201.Load((bbox_offset + 1u) * 4 + 8)) - 32768.0f; float bbox_r = float(_201.Load((bbox_offset + 2u) * 4 + 8)) - 32768.0f; @@ -632,7 +628,7 @@ void comp_main() if ((linewidth >= 0.0f) || (tag_word == 5u)) { uint trans_ix = _201.Load((bbox_offset + 5u) * 4 + 8); - uint t = (_1042.Load(36) >> uint(2)) + (6u * trans_ix); + uint t = (_1039.Load(36) >> uint(2)) + (6u * trans_ix); mat = asfloat(uint4(_201.Load(t * 4 + 8), _201.Load((t + 1u) * 4 + 8), _201.Load((t + 2u) * 4 + 8), _201.Load((t + 3u) * 4 + 8))); if (tag_word == 5u) { @@ -653,9 +649,9 @@ void comp_main() anno_fill.bbox = bbox; anno_fill.linewidth = linewidth; anno_fill.rgba_color = fill.rgba_color; - Alloc _1261; - _1261.offset = _1042.Load(32); - param_18.offset = _1261.offset; + Alloc _1258; + _1258.offset = _1039.Load(32); + param_18.offset = _1258.offset; AnnotatedRef param_19 = out_ref; uint param_20 = fill_mode; AnnoColor param_21 = anno_fill; @@ -678,9 +674,9 @@ void comp_main() anno_lin.line_x = line_x; anno_lin.line_y = line_y; anno_lin.line_c = -((p0.x * line_x) + (p0.y * line_y)); - Alloc _1357; - _1357.offset = _1042.Load(32); - param_23.offset = _1357.offset; + Alloc _1354; + _1354.offset = _1039.Load(32); + param_23.offset = _1354.offset; AnnotatedRef param_24 = out_ref; uint param_25 = fill_mode; AnnoLinGradient param_26 = anno_lin; @@ -695,9 +691,9 @@ void comp_main() anno_img.linewidth = linewidth; anno_img.index = fill_img.index; anno_img.offset = fill_img.offset; - Alloc _1385; - _1385.offset = _1042.Load(32); - param_28.offset = _1385.offset; + Alloc _1382; + _1382.offset = _1039.Load(32); + param_28.offset = _1382.offset; AnnotatedRef param_29 = out_ref; uint param_30 = fill_mode; AnnoImage param_31 = anno_img; @@ -714,9 +710,9 @@ void comp_main() Clip begin_clip = Element_BeginClip_read(param_32); anno_begin_clip.bbox = begin_clip.bbox; anno_begin_clip.linewidth = 0.0f; - Alloc _1413; - _1413.offset = _1042.Load(32); - param_33.offset = _1413.offset; + Alloc _1410; + _1410.offset = _1039.Load(32); + param_33.offset = _1410.offset; AnnotatedRef param_34 = out_ref; uint param_35 = 0u; AnnoBeginClip param_36 = anno_begin_clip; @@ -729,9 +725,9 @@ void comp_main() ElementRef param_37 = this_ref; Clip end_clip = Element_EndClip_read(param_37); anno_end_clip.bbox = end_clip.bbox; - Alloc _1438; - _1438.offset = _1042.Load(32); - param_38.offset = _1438.offset; + Alloc _1435; + _1435.offset = _1039.Load(32); + param_38.offset = _1435.offset; AnnotatedRef param_39 = out_ref; AnnoEndClip param_40 = anno_end_clip; Annotated_EndClip_write(param_38, param_39, param_40); diff --git a/piet-gpu/shader/gen/draw_leaf.msl b/piet-gpu/shader/gen/draw_leaf.msl index f713186..e20fcb2 100644 --- a/piet-gpu/shader/gen/draw_leaf.msl +++ b/piet-gpu/shader/gen/draw_leaf.msl @@ -251,12 +251,9 @@ DrawMonoid map_tag(thread const uint& tag_word) return DrawMonoid{ 1u, 0u }; } case 9u: - { - return DrawMonoid{ 1u, 1u }; - } case 10u: { - return DrawMonoid{ 0u, 1u }; + return DrawMonoid{ 1u, 1u }; } default: { @@ -609,7 +606,7 @@ void Annotated_EndClip_write(thread const Alloc& a, thread const AnnotatedRef& r 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]]) +kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1039 [[buffer(1)]], const device SceneBuf& v_225 [[buffer(2)]], const device ParentBuf& _1005 [[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[512]; uint ix = gl_GlobalInvocationID.x * 8u; @@ -650,9 +647,9 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1 DrawMonoid row = tag_monoid_identity(); if (gl_WorkGroupID.x > 0u) { - uint _1011 = gl_WorkGroupID.x - 1u; - row.path_ix = _1008.parent[_1011].path_ix; - row.clip_ix = _1008.parent[_1011].clip_ix; + uint _1008 = gl_WorkGroupID.x - 1u; + row.path_ix = _1005.parent[_1008].path_ix; + row.clip_ix = _1005.parent[_1008].clip_ix; } if (gl_LocalInvocationID.x > 0u) { @@ -661,8 +658,8 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1 row = combine_tag_monoid(param_10, param_11); } 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) }; + uint out_base = (_1039.conf.drawmonoid_alloc.offset >> uint(2)) + (out_ix * 2u); + AnnotatedRef out_ref = AnnotatedRef{ _1039.conf.anno_alloc.offset + (out_ix * 40u) }; float4 mat; float2 translate; AnnoColor anno_fill; @@ -689,7 +686,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1 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)); + uint bbox_offset = (_1039.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; @@ -700,7 +697,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1 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); + uint t = (_1039.conf.trans_alloc.offset >> uint(2)) + (6u * trans_ix); mat = as_type(uint4(v_201.memory[t], v_201.memory[t + 1u], v_201.memory[t + 2u], v_201.memory[t + 3u])); if (tag_word == 5u) { @@ -721,7 +718,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1 anno_fill.bbox = bbox; anno_fill.linewidth = linewidth; anno_fill.rgba_color = fill.rgba_color; - param_18.offset = _1042.conf.anno_alloc.offset; + param_18.offset = _1039.conf.anno_alloc.offset; AnnotatedRef param_19 = out_ref; uint param_20 = fill_mode; AnnoColor param_21 = anno_fill; @@ -744,7 +741,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1 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; + param_23.offset = _1039.conf.anno_alloc.offset; AnnotatedRef param_24 = out_ref; uint param_25 = fill_mode; AnnoLinGradient param_26 = anno_lin; @@ -759,7 +756,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1 anno_img.linewidth = linewidth; anno_img.index = fill_img.index; anno_img.offset = fill_img.offset; - param_28.offset = _1042.conf.anno_alloc.offset; + param_28.offset = _1039.conf.anno_alloc.offset; AnnotatedRef param_29 = out_ref; uint param_30 = fill_mode; AnnoImage param_31 = anno_img; @@ -776,7 +773,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1 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; + param_33.offset = _1039.conf.anno_alloc.offset; AnnotatedRef param_34 = out_ref; uint param_35 = 0u; AnnoBeginClip param_36 = anno_begin_clip; @@ -789,7 +786,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1 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; + param_38.offset = _1039.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); diff --git a/piet-gpu/shader/gen/draw_leaf.spv b/piet-gpu/shader/gen/draw_leaf.spv index 8fade68..77ed9cf 100644 Binary files a/piet-gpu/shader/gen/draw_leaf.spv and b/piet-gpu/shader/gen/draw_leaf.spv differ diff --git a/piet-gpu/shader/gen/draw_reduce.dxil b/piet-gpu/shader/gen/draw_reduce.dxil index f1e48e1..c001e89 100644 Binary files a/piet-gpu/shader/gen/draw_reduce.dxil and b/piet-gpu/shader/gen/draw_reduce.dxil differ diff --git a/piet-gpu/shader/gen/draw_reduce.hlsl b/piet-gpu/shader/gen/draw_reduce.hlsl index 216d923..e56ec3d 100644 --- a/piet-gpu/shader/gen/draw_reduce.hlsl +++ b/piet-gpu/shader/gen/draw_reduce.hlsl @@ -44,15 +44,14 @@ struct Config uint pathseg_offset; }; -static const DrawMonoid _88 = { 1u, 0u }; -static const DrawMonoid _90 = { 1u, 1u }; -static const DrawMonoid _92 = { 0u, 1u }; -static const DrawMonoid _94 = { 0u, 0u }; +static const DrawMonoid _87 = { 1u, 0u }; +static const DrawMonoid _89 = { 1u, 1u }; +static const DrawMonoid _91 = { 0u, 0u }; ByteAddressBuffer _46 : register(t2, space0); -RWByteAddressBuffer _203 : register(u3, space0); -RWByteAddressBuffer _217 : register(u0, space0); -ByteAddressBuffer _223 : register(t1, space0); +RWByteAddressBuffer _200 : register(u3, space0); +RWByteAddressBuffer _214 : register(u0, space0); +ByteAddressBuffer _220 : register(t1, space0); static uint3 gl_WorkGroupID; static uint3 gl_LocalInvocationID; @@ -81,19 +80,16 @@ DrawMonoid map_tag(uint tag_word) case 5u: case 6u: { - return _88; + return _87; } case 9u: - { - return _90; - } case 10u: { - return _92; + return _89; } default: { - return _94; + return _91; } } } @@ -115,8 +111,8 @@ DrawMonoid combine_tag_monoid(DrawMonoid a, DrawMonoid b) void comp_main() { uint ix = gl_GlobalInvocationID.x * 8u; - ElementRef _110 = { ix * 36u }; - ElementRef ref = _110; + ElementRef _107 = { ix * 36u }; + ElementRef ref = _107; ElementRef param = ref; uint tag_word = Element_tag(param).tag; uint param_1 = tag_word; @@ -148,8 +144,8 @@ void comp_main() } if (gl_LocalInvocationID.x == 0u) { - _203.Store(gl_WorkGroupID.x * 8 + 0, agg.path_ix); - _203.Store(gl_WorkGroupID.x * 8 + 4, agg.clip_ix); + _200.Store(gl_WorkGroupID.x * 8 + 0, agg.path_ix); + _200.Store(gl_WorkGroupID.x * 8 + 4, agg.clip_ix); } } diff --git a/piet-gpu/shader/gen/draw_reduce.msl b/piet-gpu/shader/gen/draw_reduce.msl index 550cf8c..26a8793 100644 --- a/piet-gpu/shader/gen/draw_reduce.msl +++ b/piet-gpu/shader/gen/draw_reduce.msl @@ -98,12 +98,9 @@ DrawMonoid map_tag(thread const uint& tag_word) return DrawMonoid{ 1u, 0u }; } case 9u: - { - return DrawMonoid{ 1u, 1u }; - } case 10u: { - return DrawMonoid{ 0u, 1u }; + return DrawMonoid{ 1u, 1u }; } default: { @@ -127,7 +124,7 @@ DrawMonoid combine_tag_monoid(thread const DrawMonoid& a, thread const DrawMonoi return c; } -kernel void main0(const device SceneBuf& v_46 [[buffer(2)]], device OutBuf& _203 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) +kernel void main0(const device SceneBuf& v_46 [[buffer(2)]], device OutBuf& _200 [[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[512]; uint ix = gl_GlobalInvocationID.x * 8u; @@ -163,8 +160,8 @@ kernel void main0(const device SceneBuf& v_46 [[buffer(2)]], device OutBuf& _203 } if (gl_LocalInvocationID.x == 0u) { - _203.outbuf[gl_WorkGroupID.x].path_ix = agg.path_ix; - _203.outbuf[gl_WorkGroupID.x].clip_ix = agg.clip_ix; + _200.outbuf[gl_WorkGroupID.x].path_ix = agg.path_ix; + _200.outbuf[gl_WorkGroupID.x].clip_ix = agg.clip_ix; } } diff --git a/piet-gpu/shader/gen/draw_reduce.spv b/piet-gpu/shader/gen/draw_reduce.spv index 02ebc5d..73430c6 100644 Binary files a/piet-gpu/shader/gen/draw_reduce.spv and b/piet-gpu/shader/gen/draw_reduce.spv differ diff --git a/piet-gpu/src/encoder.rs b/piet-gpu/src/encoder.rs index 338c01c..fb32f26 100644 --- a/piet-gpu/src/encoder.rs +++ b/piet-gpu/src/encoder.rs @@ -62,6 +62,9 @@ const DRAWOBJ_PART_SIZE: usize = 4096; // Element struct in piet-gpu-types; that's pretty much going away. const ELEMENT_FILLCOLOR: u32 = 4; +const ELEMENT_FILLLINGRADIENT: u32 = 5; +const ELEMENT_BEGINCLIP: u32 = 9; +const ELEMENT_ENDCLIP: u32 = 10; #[repr(C)] #[derive(Clone, Copy, Debug, Default, Zeroable, Pod)] @@ -71,6 +74,24 @@ pub struct FillColor { padding: [u32; 7], } +#[repr(C)] +#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)] +pub struct FillLinGradient { + tag: u32, + index: u32, + p0: [f32; 2], + p1: [f32; 2], + padding: [u32; 3], +} + +#[repr(C)] +#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)] +pub struct Clip { + tag: u32, + bbox: [f32; 4], + padding: [u32; 4], +} + impl Encoder { pub fn new() -> Encoder { Encoder { @@ -116,6 +137,45 @@ impl Encoder { self.drawobj_stream.extend(bytemuck::bytes_of(&element)); } + /// Encode a fill linear gradient draw object. + /// + /// This should be encoded after a path. + pub fn fill_lin_gradient(&mut self, index: u32, p0: [f32; 2], p1: [f32; 2]) { + let element = FillLinGradient { + tag: ELEMENT_FILLLINGRADIENT, + index, + p0, + p1, + ..Default::default() + }; + self.drawobj_stream.extend(bytemuck::bytes_of(&element)); + } + + /// Start a clip and return a save point to be filled in later. + pub fn begin_clip(&mut self) -> usize { + let saved = self.drawobj_stream.len(); + let element = Clip { + tag: ELEMENT_BEGINCLIP, + ..Default::default() + }; + self.drawobj_stream.extend(bytemuck::bytes_of(&element)); + saved + } + + pub fn end_clip(&mut self, bbox: [f32; 4], save_point: usize) { + let element = Clip { + tag: ELEMENT_ENDCLIP, + bbox, + ..Default::default() + }; + self.drawobj_stream[save_point + 4..save_point + 20] + .clone_from_slice(bytemuck::bytes_of(&bbox)); + self.drawobj_stream.extend(bytemuck::bytes_of(&element)); + // This is a dummy path, and will go away with the new clip impl. + self.tag_stream.push(0x10); + self.n_path += 1; + } + /// Return a config for the element processing pipeline. /// /// This does not include further pipeline processing. Also returns the diff --git a/piet-gpu/src/render_ctx.rs b/piet-gpu/src/render_ctx.rs index 2c5ee73..31fbf9e 100644 --- a/piet-gpu/src/render_ctx.rs +++ b/piet-gpu/src/render_ctx.rs @@ -11,7 +11,7 @@ use piet::{ use piet_gpu_hal::BufWrite; use piet_gpu_types::encoder::{Encode, Encoder}; -use piet_gpu_types::scene::{Clip, Element, FillLinGradient, SetFillMode}; +use piet_gpu_types::scene::{Element, SetFillMode}; use crate::gradient::{LinearGradient, RampCache}; use crate::text::Font; @@ -64,8 +64,8 @@ struct State { } struct ClipElement { - /// Index of BeginClip element in element vec, for bbox fixup. - begin_ix: usize, + /// Byte offset of BeginClip element in element vec, for bbox fixup. + save_point: usize, bbox: Option, } @@ -242,21 +242,17 @@ impl RenderContext for PietGpuRenderContext { fn fill_even_odd(&mut self, _shape: impl Shape, _brush: &impl IntoBrush) {} fn clip(&mut self, shape: impl Shape) { - self.set_fill_mode(FillMode::Nonzero); + self.encode_linewidth(-1.0); let path = shape.path_elements(TOLERANCE); self.encode_path(path, true); - let begin_ix = self.elements.len(); - self.elements.push(Element::BeginClip(Clip { - bbox: Default::default(), - })); + let save_point = self.new_encoder.begin_clip(); if self.clip_stack.len() >= MAX_BLEND_STACK { panic!("Maximum clip/blend stack size {} exceeded", MAX_BLEND_STACK); } self.clip_stack.push(ClipElement { bbox: None, - begin_ix, + save_point, }); - self.path_count += 1; if let Some(tos) = self.state_stack.last_mut() { tos.n_clip += 1; } @@ -405,14 +401,7 @@ impl PietGpuRenderContext { let tos = self.clip_stack.pop().unwrap(); let bbox = tos.bbox.unwrap_or_default(); let bbox_f32_4 = rect_to_f32_4(bbox); - self.elements - .push(Element::EndClip(Clip { bbox: bbox_f32_4 })); - self.path_count += 1; - if let Element::BeginClip(begin_clip) = &mut self.elements[tos.begin_ix] { - begin_clip.bbox = bbox_f32_4; - } else { - unreachable!("expected BeginClip, not found"); - } + self.new_encoder.end_clip(bbox_f32_4, tos.save_point); if let Some(bbox) = tos.bbox { self.union_bbox(bbox); } @@ -468,13 +457,8 @@ impl PietGpuRenderContext { self.new_encoder.fill_color(*rgba_color); } PietGpuBrush::LinGradient(lin) => { - let fill_lin = FillLinGradient { - index: lin.ramp_id, - p0: lin.start, - p1: lin.end, - }; - self.elements.push(Element::FillLinGradient(fill_lin)); - self.path_count += 1; + self.new_encoder + .fill_lin_gradient(lin.ramp_id, lin.start, lin.end); } } } diff --git a/piet-gpu/src/text.rs b/piet-gpu/src/text.rs index d25320b..a47c614 100644 --- a/piet-gpu/src/text.rs +++ b/piet-gpu/src/text.rs @@ -12,8 +12,8 @@ use piet::{ use crate::encoder::GlyphEncoder; use crate::render_ctx::{self, FillMode}; -use crate::PietGpuRenderContext; use crate::stages::Transform; +use crate::PietGpuRenderContext; // This is very much a hack to get things working. // On Windows, can set this to "c:\\Windows\\Fonts\\seguiemj.ttf" to get color emoji @@ -51,7 +51,6 @@ struct Glyph { y: f32, } - struct TextRenderCtx<'a> { scaler: Scaler<'a>, }