From d4f77fd4cecc1a60975f0a20647ab775226d0655 Mon Sep 17 00:00:00 2001 From: Blaise Tine Date: Sat, 14 Sep 2024 02:24:56 -0700 Subject: [PATCH] fixed output merger synchronization bug --- ci/regression.sh.in | 16 ++--- hw/rtl/om/VX_om_arb.sv | 2 +- hw/rtl/om/VX_om_blend.sv | 68 ++++++++++---------- hw/rtl/om/VX_om_blend_func.sv | 28 ++++---- hw/rtl/om/VX_om_blend_minmax.sv | 58 ++++++++--------- hw/rtl/om/VX_om_blend_multadd.sv | 46 ++++++------- hw/rtl/om/VX_om_bus_if.sv | 8 +-- hw/rtl/om/VX_om_mem.sv | 4 +- hw/rtl/om/VX_om_pkg.sv | 25 ++++--- hw/rtl/om/VX_om_unit.sv | 60 ++++++++++------- hw/unit_tests/rop_unit/blending/blend_tb.sv | 28 ++++---- sim/simx/execute.cpp | 2 +- sim/simx/om_unit.cpp | 5 +- tests/regression/om/kernel.cpp | 33 ++++++++-- tests/regression/om/whitebox_128.png | Bin 390 -> 399 bytes tests/regression/om/whitebox_16.png | Bin 86 -> 98 bytes tests/regression/om/whitebox_32.png | Bin 104 -> 119 bytes tests/regression/om/whitebox_64.png | Bin 207 -> 0 bytes tests/regression/om/whitebox_8.png | Bin 79 -> 91 bytes tests/regression/om/whitebox_b_128.png | Bin 0 -> 405 bytes tests/regression/om/whitebox_b_32.png | Bin 167 -> 0 bytes tests/regression/om/whitebox_d_128.png | Bin 0 -> 27123 bytes tests/regression/om/whitebox_d_32.png | Bin 110 -> 0 bytes tests/regression/om/whitebox_db_128.png | Bin 0 -> 37363 bytes tests/regression/om/whitebox_db_32.png | Bin 195 -> 0 bytes 25 files changed, 207 insertions(+), 176 deletions(-) delete mode 100644 tests/regression/om/whitebox_64.png create mode 100644 tests/regression/om/whitebox_b_128.png delete mode 100644 tests/regression/om/whitebox_b_32.png create mode 100644 tests/regression/om/whitebox_d_128.png delete mode 100644 tests/regression/om/whitebox_d_32.png create mode 100644 tests/regression/om/whitebox_db_128.png delete mode 100644 tests/regression/om/whitebox_db_32.png diff --git a/ci/regression.sh.in b/ci/regression.sh.in index d3ab5a764..9f90f7ec8 100755 --- a/ci/regression.sh.in +++ b/ci/regression.sh.in @@ -157,14 +157,14 @@ om() { echo "begin render output tests..." -CONFIGS="-DEXT_OM_ENABLE" ./ci/blackbox.sh --driver=simx --app=om --args="-w32 -h32 -rwhitebox_32.png" -CONFIGS="-DEXT_OM_ENABLE" ./ci/blackbox.sh --driver=simx --app=om --args="-w32 -h32 -d -rwhitebox_d_32.png" -CONFIGS="-DEXT_OM_ENABLE" ./ci/blackbox.sh --driver=simx --app=om --args="-w32 -h32 -b -rwhitebox_b_32.png" -CONFIGS="-DEXT_OM_ENABLE" ./ci/blackbox.sh --driver=simx --app=om --args="-w32 -h32 -db -rwhitebox_db_32.png" -CONFIGS="-DEXT_OM_ENABLE" ./ci/blackbox.sh --driver=rtlsim --app=om --args="-w32 -h32 -rwhitebox_32.png" -CONFIGS="-DEXT_OM_ENABLE" ./ci/blackbox.sh --driver=rtlsim --app=om --args="-w32 -h32 -d -rwhitebox_d_32.png" -CONFIGS="-DEXT_OM_ENABLE" ./ci/blackbox.sh --driver=rtlsim --app=om --args="-w32 -h32 -b -rwhitebox_b_32.png" -CONFIGS="-DEXT_OM_ENABLE" ./ci/blackbox.sh --driver=rtlsim --app=om --args="-w32 -h32 -db -rwhitebox_db_32.png" +CONFIGS="-DEXT_OM_ENABLE" ./ci/blackbox.sh --driver=simx --app=om --args="-rwhitebox_128.png" +CONFIGS="-DEXT_OM_ENABLE" ./ci/blackbox.sh --driver=simx --app=om --args="-d -rwhitebox_d_128.png" +CONFIGS="-DEXT_OM_ENABLE" ./ci/blackbox.sh --driver=simx --app=om --args="-b -rwhitebox_b_128.png" +CONFIGS="-DEXT_OM_ENABLE" ./ci/blackbox.sh --driver=simx --app=om --args="-db -rwhitebox_db_128.png" +CONFIGS="-DEXT_OM_ENABLE" ./ci/blackbox.sh --driver=rtlsim --app=om --args="-rwhitebox_128.png" +CONFIGS="-DEXT_OM_ENABLE" ./ci/blackbox.sh --driver=rtlsim --app=om --args="-d -rwhitebox_d_128.png" +CONFIGS="-DEXT_OM_ENABLE" ./ci/blackbox.sh --driver=rtlsim --app=om --args="-b -rwhitebox_b_128.png" +CONFIGS="-DEXT_OM_ENABLE" ./ci/blackbox.sh --driver=rtlsim --app=om --args="-db -rwhitebox_db_128.png" CONFIGS="-DEXT_OM_ENABLE -DOCACHE_DISABLE" ./ci/blackbox.sh --driver=rtlsim --app=om --args="-rwhitebox_128.png" CONFIGS="-DEXT_OM_ENABLE -DOCACHE_NUM_BANKS=8" ./ci/blackbox.sh --driver=rtlsim --app=om --args="-rwhitebox_128.png" --perf=5 CONFIGS="-DEXT_OM_ENABLE -DNUM_OM_UNITS=2 -DL1_DISABLE -DLM_DISABLE -DOCACHE_DISABLE" ./ci/blackbox.sh --driver=simx --app=om --args="-rwhitebox_128.png" --cores=4 --warps=1 --threads=2 diff --git a/hw/rtl/om/VX_om_arb.sv b/hw/rtl/om/VX_om_arb.sv index e4fbed555..a75b52505 100644 --- a/hw/rtl/om/VX_om_arb.sv +++ b/hw/rtl/om/VX_om_arb.sv @@ -32,7 +32,7 @@ module VX_om_arb import VX_om_pkg::*; #( VX_om_bus_if.master bus_out_if [NUM_OUTPUTS] ); - localparam REQ_DATAW = `UUID_WIDTH + NUM_LANES * (1 + 2 * `VX_OM_DIM_BITS + $bits(rgba_t) + `VX_OM_DEPTH_BITS + 1); + localparam REQ_DATAW = `UUID_WIDTH + NUM_LANES * (1 + 2 * `VX_OM_DIM_BITS + $bits(om_color_t) + `VX_OM_DEPTH_BITS + 1); wire [NUM_INPUTS-1:0] req_valid_in; wire [NUM_INPUTS-1:0][REQ_DATAW-1:0] req_data_in; diff --git a/hw/rtl/om/VX_om_blend.sv b/hw/rtl/om/VX_om_blend.sv index a4fc1f107..a6b922f1c 100644 --- a/hw/rtl/om/VX_om_blend.sv +++ b/hw/rtl/om/VX_om_blend.sv @@ -36,11 +36,11 @@ module VX_om_blend import VX_om_pkg::*; #( input wire ready_out, // Input values - input rgba_t [NUM_LANES-1:0] src_color, - input rgba_t [NUM_LANES-1:0] dst_color, + input om_color_t [NUM_LANES-1:0] src_color, + input om_color_t [NUM_LANES-1:0] dst_color, // Output values - output rgba_t [NUM_LANES-1:0] color_out + output om_color_t [NUM_LANES-1:0] color_out ); `UNUSED_SPARAM (INSTANCE_ID) @@ -52,8 +52,8 @@ module VX_om_blend import VX_om_pkg::*; #( assign ready_in = ~stall; - rgba_t [NUM_LANES-1:0] src_factor; - rgba_t [NUM_LANES-1:0] dst_factor; + om_color_t [NUM_LANES-1:0] src_factor; + om_color_t [NUM_LANES-1:0] dst_factor; for (genvar i = 0; i < NUM_LANES; ++i) begin : g_blend_func_src VX_om_blend_func blend_func_src ( @@ -80,10 +80,10 @@ module VX_om_blend import VX_om_pkg::*; #( wire valid_s1, valid_s2; wire [TAG_WIDTH-1:0] tag_s1, tag_s2; - rgba_t [NUM_LANES-1:0] src_color_s1; - rgba_t [NUM_LANES-1:0] dst_color_s1; - rgba_t [NUM_LANES-1:0] src_factor_s1; - rgba_t [NUM_LANES-1:0] dst_factor_s1; + om_color_t [NUM_LANES-1:0] src_color_s1; + om_color_t [NUM_LANES-1:0] dst_color_s1; + om_color_t [NUM_LANES-1:0] src_factor_s1; + om_color_t [NUM_LANES-1:0] dst_factor_s1; VX_pipe_register #( .DATAW (1 + TAG_WIDTH + 32 * 4 * NUM_LANES), @@ -97,10 +97,10 @@ module VX_om_blend import VX_om_pkg::*; #( .data_out ({valid_s1, tag_s1, src_color_s1, dst_color_s1, src_factor_s1, dst_factor_s1}) ); - rgba_t [NUM_LANES-1:0] mult_add_color_s2; - rgba_t [NUM_LANES-1:0] min_color_s2; - rgba_t [NUM_LANES-1:0] max_color_s2; - rgba_t [NUM_LANES-1:0] logic_op_color_s2; + om_color_t [NUM_LANES-1:0] mult_add_color_s2; + om_color_t [NUM_LANES-1:0] min_color_s2; + om_color_t [NUM_LANES-1:0] max_color_s2; + om_color_t [NUM_LANES-1:0] logic_op_color_s2; for (genvar i = 0; i < NUM_LANES; ++i) begin : g_blend_multadd VX_om_blend_multadd #( @@ -158,7 +158,7 @@ module VX_om_blend import VX_om_pkg::*; #( .data_out ({valid_s2, tag_s2}) ); - rgba_t [NUM_LANES-1:0] color_out_s2; + om_color_t [NUM_LANES-1:0] color_out_s2; for (genvar i = 0; i < NUM_LANES; ++i) begin : g_blending always @(*) begin @@ -167,29 +167,29 @@ module VX_om_blend import VX_om_pkg::*; #( `VX_OM_BLEND_MODE_ADD, `VX_OM_BLEND_MODE_SUB, `VX_OM_BLEND_MODE_REV_SUB: begin - color_out_s2[i].r = mult_add_color_s2[i].r; - color_out_s2[i].g = mult_add_color_s2[i].g; - color_out_s2[i].b = mult_add_color_s2[i].b; + color_out_s2[i].argb[23:16] = mult_add_color_s2[i].argb[23:16]; + color_out_s2[i].argb[15:8] = mult_add_color_s2[i].argb[15:8]; + color_out_s2[i].argb[7:0] = mult_add_color_s2[i].argb[7:0]; end `VX_OM_BLEND_MODE_MIN: begin - color_out_s2[i].r = min_color_s2[i].r; - color_out_s2[i].g = min_color_s2[i].g; - color_out_s2[i].b = min_color_s2[i].b; + color_out_s2[i].argb[23:16] = min_color_s2[i].argb[23:16]; + color_out_s2[i].argb[15:8] = min_color_s2[i].argb[15:8]; + color_out_s2[i].argb[7:0] = min_color_s2[i].argb[7:0]; end `VX_OM_BLEND_MODE_MAX: begin - color_out_s2[i].r = max_color_s2[i].r; - color_out_s2[i].g = max_color_s2[i].g; - color_out_s2[i].b = max_color_s2[i].b; + color_out_s2[i].argb[23:16] = max_color_s2[i].argb[23:16]; + color_out_s2[i].argb[15:8] = max_color_s2[i].argb[15:8]; + color_out_s2[i].argb[7:0] = max_color_s2[i].argb[7:0]; end `VX_OM_BLEND_MODE_LOGICOP: begin - color_out_s2[i].r = logic_op_color_s2[i].r; - color_out_s2[i].g = logic_op_color_s2[i].g; - color_out_s2[i].b = logic_op_color_s2[i].b; + color_out_s2[i].argb[23:16] = logic_op_color_s2[i].argb[23:16]; + color_out_s2[i].argb[15:8] = logic_op_color_s2[i].argb[15:8]; + color_out_s2[i].argb[7:0] = logic_op_color_s2[i].argb[7:0]; end default: begin - color_out_s2[i].r = 'x; - color_out_s2[i].g = 'x; - color_out_s2[i].b = 'x; + color_out_s2[i].argb[23:16] = 'x; + color_out_s2[i].argb[15:8] = 'x; + color_out_s2[i].argb[7:0] = 'x; end endcase // Alpha Component @@ -197,19 +197,19 @@ module VX_om_blend import VX_om_pkg::*; #( `VX_OM_BLEND_MODE_ADD, `VX_OM_BLEND_MODE_SUB, `VX_OM_BLEND_MODE_REV_SUB: begin - color_out_s2[i].a = mult_add_color_s2[i].a; + color_out_s2[i].argb[31:24] = mult_add_color_s2[i].argb[31:24]; end `VX_OM_BLEND_MODE_MIN: begin - color_out_s2[i].a = min_color_s2[i].a; + color_out_s2[i].argb[31:24] = min_color_s2[i].argb[31:24]; end `VX_OM_BLEND_MODE_MAX: begin - color_out_s2[i].a = max_color_s2[i].a; + color_out_s2[i].argb[31:24] = max_color_s2[i].argb[31:24]; end `VX_OM_BLEND_MODE_LOGICOP: begin - color_out_s2[i].a = logic_op_color_s2[i].a; + color_out_s2[i].argb[31:24] = logic_op_color_s2[i].argb[31:24]; end default: begin - color_out_s2[i].a = 'x; + color_out_s2[i].argb[31:24] = 'x; end endcase end diff --git a/hw/rtl/om/VX_om_blend_func.sv b/hw/rtl/om/VX_om_blend_func.sv index 7bbc2d376..3bfe9b8e3 100644 --- a/hw/rtl/om/VX_om_blend_func.sv +++ b/hw/rtl/om/VX_om_blend_func.sv @@ -1,12 +1,12 @@ //!/bin/bash // Copyright © 2019-2023 -// +// // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. // You may obtain a copy of the License at // http://www.apache.org/licenses/LICENSE-2.0 -// +// // Unless required by applicable law or agreed to in writing, software // distributed under the License is distributed on an "AS IS" BASIS, // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. @@ -18,9 +18,9 @@ module VX_blend_func #( parameter INDEX = 0 ) ( - input wire [`VX_OM_BLEND_FUNC_BITS-1:0] func, - input wire [3:0][7:0] src, - input wire [3:0][7:0] dst, + input wire [`VX_OM_BLEND_FUNC_BITS-1:0] func, + input wire [3:0][7:0] src, + input wire [3:0][7:0] dst, input wire [3:0][7:0] cst, output wire [7:0] result ); @@ -57,7 +57,7 @@ module VX_blend_func #( end assign result = result_r; - + endmodule module VX_om_blend_func import VX_om_pkg::*; #( @@ -66,15 +66,15 @@ module VX_om_blend_func import VX_om_pkg::*; #( input wire [`VX_OM_BLEND_FUNC_BITS-1:0] func_rgb, input wire [`VX_OM_BLEND_FUNC_BITS-1:0] func_a, - input rgba_t src_color, - input rgba_t dst_color, - input rgba_t cst_color, + input om_color_t src_color, + input om_color_t dst_color, + input om_color_t cst_color, - output rgba_t factor_out + output om_color_t factor_out ); - VX_blend_func #(0) blend_func_b (func_rgb, src_color, dst_color, cst_color, factor_out.b); - VX_blend_func #(1) blend_func_g (func_rgb, src_color, dst_color, cst_color, factor_out.g); - VX_blend_func #(2) blend_func_r (func_rgb, src_color, dst_color, cst_color, factor_out.r); - VX_blend_func #(3) blend_func_a (func_a, src_color, dst_color, cst_color, factor_out.a); + VX_blend_func #(0) blend_func_b (func_rgb, src_color, dst_color, cst_color, factor_out.argb[7:0]); + VX_blend_func #(1) blend_func_g (func_rgb, src_color, dst_color, cst_color, factor_out.argb[15:8]); + VX_blend_func #(2) blend_func_r (func_rgb, src_color, dst_color, cst_color, factor_out.argb[23:16]); + VX_blend_func #(3) blend_func_a (func_a, src_color, dst_color, cst_color, factor_out.argb[31:24]); endmodule diff --git a/hw/rtl/om/VX_om_blend_minmax.sv b/hw/rtl/om/VX_om_blend_minmax.sv index 1ec10e40c..2bf3023d4 100644 --- a/hw/rtl/om/VX_om_blend_minmax.sv +++ b/hw/rtl/om/VX_om_blend_minmax.sv @@ -1,12 +1,12 @@ //!/bin/bash // Copyright © 2019-2023 -// +// // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. // You may obtain a copy of the License at // http://www.apache.org/licenses/LICENSE-2.0 -// +// // Unless required by applicable law or agreed to in writing, software // distributed under the License is distributed on an "AS IS" BASIS, // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. @@ -23,49 +23,49 @@ module VX_om_blend_minmax import VX_om_pkg::*; #( input wire enable, - input rgba_t src_color, - input rgba_t dst_color, + input om_color_t src_color, + input om_color_t dst_color, - output rgba_t min_out, - output rgba_t max_out + output om_color_t min_out, + output om_color_t max_out ); `UNUSED_VAR (reset) - rgba_t tmp_min; - rgba_t tmp_max; + om_color_t tmp_min; + om_color_t tmp_max; - always @(*) begin - if (src_color.r > dst_color.r) begin - tmp_max.r = src_color.r; - tmp_min.r = dst_color.r; + always @(*) begin + if (src_color.argb[23:16] > dst_color.argb[23:16]) begin + tmp_max.argb[23:16] = src_color.argb[23:16]; + tmp_min.argb[23:16] = dst_color.argb[23:16]; end else begin - tmp_max.r = dst_color.r; - tmp_min.r = src_color.r; + tmp_max.argb[23:16] = dst_color.argb[23:16]; + tmp_min.argb[23:16] = src_color.argb[23:16]; end - if (src_color.g > dst_color.g) begin - tmp_max.g = src_color.g; - tmp_min.g = dst_color.g; + if (src_color.argb[15:8] > dst_color.argb[15:8]) begin + tmp_max.argb[15:8] = src_color.argb[15:8]; + tmp_min.argb[15:8] = dst_color.argb[15:8]; end else begin - tmp_max.g = dst_color.g; - tmp_min.g = src_color.g; + tmp_max.argb[15:8] = dst_color.argb[15:8]; + tmp_min.argb[15:8] = src_color.argb[15:8]; end - if (src_color.b > dst_color.b) begin - tmp_max.b = src_color.b; - tmp_min.b = dst_color.b; + if (src_color.argb[7:0] > dst_color.argb[7:0]) begin + tmp_max.argb[7:0] = src_color.argb[7:0]; + tmp_min.argb[7:0] = dst_color.argb[7:0]; end else begin - tmp_max.b = dst_color.b; - tmp_min.b = src_color.b; + tmp_max.argb[7:0] = dst_color.argb[7:0]; + tmp_min.argb[7:0] = src_color.argb[7:0]; end - if (src_color.a > dst_color.a) begin - tmp_max.a = src_color.a; - tmp_min.a = dst_color.a; + if (src_color.argb[31:24] > dst_color.argb[31:24]) begin + tmp_max.argb[31:24] = src_color.argb[31:24]; + tmp_min.argb[31:24] = dst_color.argb[31:24]; end else begin - tmp_max.a = dst_color.a; - tmp_min.a = src_color.a; + tmp_max.argb[31:24] = dst_color.argb[31:24]; + tmp_min.argb[31:24] = src_color.argb[31:24]; end end diff --git a/hw/rtl/om/VX_om_blend_multadd.sv b/hw/rtl/om/VX_om_blend_multadd.sv index 8d1676a1d..971e25c35 100644 --- a/hw/rtl/om/VX_om_blend_multadd.sv +++ b/hw/rtl/om/VX_om_blend_multadd.sv @@ -1,12 +1,12 @@ //!/bin/bash // Copyright © 2019-2023 -// +// // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. // You may obtain a copy of the License at // http://www.apache.org/licenses/LICENSE-2.0 -// +// // Unless required by applicable law or agreed to in writing, software // distributed under the License is distributed on an "AS IS" BASIS, // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. @@ -26,13 +26,13 @@ module VX_om_blend_multadd import VX_om_pkg::*; #( input wire [`VX_OM_BLEND_MODE_BITS-1:0] mode_rgb, input wire [`VX_OM_BLEND_MODE_BITS-1:0] mode_a, - input rgba_t src_color, - input rgba_t dst_color, + input om_color_t src_color, + input om_color_t dst_color, - input rgba_t src_factor, - input rgba_t dst_factor, + input om_color_t src_factor, + input om_color_t dst_factor, - output rgba_t color_out + output om_color_t color_out ); `STATIC_ASSERT((LATENCY == 3), ("invalid parameter")) @@ -42,19 +42,19 @@ module VX_om_blend_multadd import VX_om_pkg::*; #( reg [15:0] prod_src_r, prod_src_g, prod_src_b, prod_src_a; reg [15:0] prod_dst_r, prod_dst_g, prod_dst_b, prod_dst_a; - reg [16:0] sum_r, sum_g, sum_b, sum_a; + reg [16:0] sum_r, sum_g, sum_b, sum_a; - always @(posedge clk) begin + always @(posedge clk) begin if (enable) begin - prod_src_r <= src_color.r * src_factor.r; - prod_src_g <= src_color.g * src_factor.g; - prod_src_b <= src_color.b * src_factor.b; - prod_src_a <= src_color.a * src_factor.a; + prod_src_r <= src_color.argb[23:16] * src_factor.argb[23:16]; + prod_src_g <= src_color.argb[15:8] * src_factor.argb[15:8]; + prod_src_b <= src_color.argb[7:0] * src_factor.argb[7:0]; + prod_src_a <= src_color.argb[31:24] * src_factor.argb[31:24]; - prod_dst_r <= dst_color.r * dst_factor.r; - prod_dst_g <= dst_color.g * dst_factor.g; - prod_dst_b <= dst_color.b * dst_factor.b; - prod_dst_a <= dst_color.a * dst_factor.a; + prod_dst_r <= dst_color.argb[23:16] * dst_factor.argb[23:16]; + prod_dst_g <= dst_color.argb[15:8] * dst_factor.argb[15:8]; + prod_dst_b <= dst_color.argb[7:0] * dst_factor.argb[7:0]; + prod_dst_a <= dst_color.argb[31:24] * dst_factor.argb[31:24]; case (mode_rgb) `VX_OM_BLEND_MODE_ADD: begin @@ -65,7 +65,7 @@ module VX_om_blend_multadd import VX_om_pkg::*; #( `VX_OM_BLEND_MODE_SUB: begin sum_r <= prod_src_r - prod_dst_r + 16'h80; sum_g <= prod_src_g - prod_dst_g + 16'h80; - sum_b <= prod_src_b - prod_dst_b + 16'h80; + sum_b <= prod_src_b - prod_dst_b + 16'h80; end `VX_OM_BLEND_MODE_REV_SUB: begin sum_r <= prod_dst_r - prod_src_r + 16'h80; @@ -126,11 +126,11 @@ module VX_om_blend_multadd import VX_om_pkg::*; #( // divide by 255 - rgba_t result; - assign result.r = 8'((clamp_r + (clamp_r >> 8)) >> 8); - assign result.g = 8'((clamp_g + (clamp_g >> 8)) >> 8); - assign result.b = 8'((clamp_b + (clamp_b >> 8)) >> 8); - assign result.a = 8'((clamp_a + (clamp_a >> 8)) >> 8); + om_color_t result; + assign result.argb[31:24] = 8'((clamp_a + (clamp_a >> 8)) >> 8); + assign result.argb[23:16] = 8'((clamp_r + (clamp_r >> 8)) >> 8); + assign result.argb[15:8] = 8'((clamp_g + (clamp_g >> 8)) >> 8); + assign result.argb[7:0] = 8'((clamp_b + (clamp_b >> 8)) >> 8); VX_pipe_register #( .DATAW (32) diff --git a/hw/rtl/om/VX_om_bus_if.sv b/hw/rtl/om/VX_om_bus_if.sv index e95f53fdd..7fcfc7fcd 100644 --- a/hw/rtl/om/VX_om_bus_if.sv +++ b/hw/rtl/om/VX_om_bus_if.sv @@ -1,12 +1,12 @@ //!/bin/bash // Copyright © 2019-2023 -// +// // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. // You may obtain a copy of the License at // http://www.apache.org/licenses/LICENSE-2.0 -// +// // Unless required by applicable law or agreed to in writing, software // distributed under the License is distributed on an "AS IS" BASIS, // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. @@ -20,10 +20,10 @@ interface VX_om_bus_if import VX_om_pkg::*; #( ) (); typedef struct packed { logic [`UUID_WIDTH-1:0] uuid; - logic [NUM_LANES-1:0] mask; + logic [NUM_LANES-1:0] mask; logic [NUM_LANES-1:0][`VX_OM_DIM_BITS-1:0] pos_x; logic [NUM_LANES-1:0][`VX_OM_DIM_BITS-1:0] pos_y; - rgba_t [NUM_LANES-1:0] color; + om_color_t [NUM_LANES-1:0] color; logic [NUM_LANES-1:0][`VX_OM_DEPTH_BITS-1:0] depth; logic [NUM_LANES-1:0] face; } req_data_t; diff --git a/hw/rtl/om/VX_om_mem.sv b/hw/rtl/om/VX_om_mem.sv index bb9c314ca..450e3a395 100644 --- a/hw/rtl/om/VX_om_mem.sv +++ b/hw/rtl/om/VX_om_mem.sv @@ -37,7 +37,7 @@ module VX_om_mem import VX_gpu_pkg::*; import VX_om_pkg::*; #( input wire req_rw, input wire [NUM_LANES-1:0][`VX_OM_DIM_BITS-1:0] req_pos_x, input wire [NUM_LANES-1:0][`VX_OM_DIM_BITS-1:0] req_pos_y, - input rgba_t [NUM_LANES-1:0] req_color, + input om_color_t [NUM_LANES-1:0] req_color, input wire [NUM_LANES-1:0][`VX_OM_DEPTH_BITS-1:0] req_depth, input wire [NUM_LANES-1:0][`VX_OM_STENCIL_BITS-1:0] req_stencil, input wire [NUM_LANES-1:0] req_face, @@ -48,7 +48,7 @@ module VX_om_mem import VX_gpu_pkg::*; import VX_om_pkg::*; #( // Response interface output wire rsp_valid, output wire [NUM_LANES-1:0] rsp_mask, - output rgba_t [NUM_LANES-1:0] rsp_color, + output om_color_t [NUM_LANES-1:0] rsp_color, output wire [NUM_LANES-1:0][`VX_OM_DEPTH_BITS-1:0] rsp_depth, output wire [NUM_LANES-1:0][`VX_OM_STENCIL_BITS-1:0] rsp_stencil, output wire [TAG_WIDTH-1:0] rsp_tag, diff --git a/hw/rtl/om/VX_om_pkg.sv b/hw/rtl/om/VX_om_pkg.sv index 18f8f22f9..15fc3c30c 100644 --- a/hw/rtl/om/VX_om_pkg.sv +++ b/hw/rtl/om/VX_om_pkg.sv @@ -1,12 +1,12 @@ //!/bin/bash // Copyright © 2019-2023 -// +// // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. // You may obtain a copy of the License at // http://www.apache.org/licenses/LICENSE-2.0 -// +// // Unless required by applicable law or agreed to in writing, software // distributed under the License is distributed on an "AS IS" BASIS, // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. @@ -21,11 +21,8 @@ package VX_om_pkg; typedef struct packed { - logic [7:0] a; - logic [7:0] r; - logic [7:0] g; - logic [7:0] b; -} rgba_t; + logic [31:0] argb; +} om_color_t; typedef struct packed { logic [`OM_ADDR_BITS-1:0] cbuf_addr; @@ -38,16 +35,16 @@ typedef struct packed { logic depth_enable; logic [`VX_OM_DEPTH_FUNC_BITS-1:0] depth_func; logic depth_writemask; - + logic [1:0] stencil_enable; - logic [1:0][`VX_OM_DEPTH_FUNC_BITS-1:0] stencil_func; + logic [1:0][`VX_OM_DEPTH_FUNC_BITS-1:0] stencil_func; logic [1:0][`VX_OM_STENCIL_OP_BITS-1:0] stencil_zpass; logic [1:0][`VX_OM_STENCIL_OP_BITS-1:0] stencil_zfail; logic [1:0][`VX_OM_STENCIL_OP_BITS-1:0] stencil_fail; logic [1:0][`VX_OM_STENCIL_BITS-1:0] stencil_ref; - logic [1:0][`VX_OM_STENCIL_BITS-1:0] stencil_mask; + logic [1:0][`VX_OM_STENCIL_BITS-1:0] stencil_mask; logic [1:0][`VX_OM_STENCIL_BITS-1:0] stencil_writemask; - + logic blend_enable; logic [`VX_OM_BLEND_MODE_BITS-1:0] blend_mode_rgb; logic [`VX_OM_BLEND_MODE_BITS-1:0] blend_mode_a; @@ -55,15 +52,15 @@ typedef struct packed { logic [`VX_OM_BLEND_FUNC_BITS-1:0] blend_src_a; logic [`VX_OM_BLEND_FUNC_BITS-1:0] blend_dst_rgb; logic [`VX_OM_BLEND_FUNC_BITS-1:0] blend_dst_a; - rgba_t blend_const; - + om_color_t blend_const; + logic [`VX_OM_LOGIC_OP_BITS-1:0] logic_op; } om_dcrs_t; typedef struct packed { logic [1:0] rt_idx; logic [`VX_OM_DIM_BITS-1:0] pos_x; - logic [`VX_OM_DIM_BITS-1:0] pos_y; + logic [`VX_OM_DIM_BITS-1:0] pos_y; logic [23:0] depth; logic [2:0] sample_idx; logic [7:0] sample_mask; diff --git a/hw/rtl/om/VX_om_unit.sv b/hw/rtl/om/VX_om_unit.sv index df7cf1598..1e6506a55 100644 --- a/hw/rtl/om/VX_om_unit.sv +++ b/hw/rtl/om/VX_om_unit.sv @@ -34,9 +34,9 @@ module VX_om_unit import VX_gpu_pkg::*; import VX_om_pkg::*; #( VX_dcr_bus_if.slave dcr_bus_if, VX_om_bus_if.slave om_bus_if ); - localparam MEM_TAG_WIDTH = `UUID_WIDTH + NUM_LANES * (`VX_OM_DIM_BITS + `VX_OM_DIM_BITS + 32 + `VX_OM_DEPTH_BITS + 1); - localparam DS_TAG_WIDTH = NUM_LANES * (`VX_OM_DIM_BITS + `VX_OM_DIM_BITS + 1 + 1 + 32); - localparam BLEND_TAG_WIDTH = NUM_LANES * (`VX_OM_DIM_BITS + `VX_OM_DIM_BITS + 1); + localparam MEM_TAG_WIDTH = `UUID_WIDTH + NUM_LANES * (`VX_OM_DIM_BITS + `VX_OM_DIM_BITS + 32 + `VX_OM_DEPTH_BITS + 1); + localparam DS_TAG_WIDTH = `UUID_WIDTH + NUM_LANES * (`VX_OM_DIM_BITS + `VX_OM_DIM_BITS + 1 + 1 + 32); + localparam BLEND_TAG_WIDTH = `UUID_WIDTH + NUM_LANES * (`VX_OM_DIM_BITS + `VX_OM_DIM_BITS + 1); // DCRs @@ -59,7 +59,7 @@ module VX_om_unit import VX_gpu_pkg::*; import VX_om_pkg::*; #( wire mem_req_rw, mem_req_rw_r; wire [NUM_LANES-1:0][`VX_OM_DIM_BITS-1:0] mem_req_pos_x, mem_req_pos_x_r; wire [NUM_LANES-1:0][`VX_OM_DIM_BITS-1:0] mem_req_pos_y, mem_req_pos_y_r; - rgba_t [NUM_LANES-1:0] mem_req_color, mem_req_color_r; + om_color_t [NUM_LANES-1:0] mem_req_color, mem_req_color_r; wire [NUM_LANES-1:0][`VX_OM_DEPTH_BITS-1:0] mem_req_depth, mem_req_depth_r; wire [NUM_LANES-1:0][`VX_OM_STENCIL_BITS-1:0] mem_req_stencil, mem_req_stencil_r; wire [NUM_LANES-1:0] mem_req_face, mem_req_face_r; @@ -68,7 +68,7 @@ module VX_om_unit import VX_gpu_pkg::*; import VX_om_pkg::*; #( wire mem_rsp_valid; wire [NUM_LANES-1:0] mem_rsp_mask; - rgba_t [NUM_LANES-1:0] mem_rsp_color; + om_color_t [NUM_LANES-1:0] mem_rsp_color; wire [NUM_LANES-1:0][`VX_OM_DEPTH_BITS-1:0] mem_rsp_depth; wire [NUM_LANES-1:0][`VX_OM_STENCIL_BITS-1:0] mem_rsp_stencil; wire [MEM_TAG_WIDTH-1:0] mem_rsp_tag; @@ -166,9 +166,9 @@ module VX_om_unit import VX_gpu_pkg::*; import VX_om_pkg::*; #( wire [BLEND_TAG_WIDTH-1:0] blend_tag_out; wire blend_ready_out; - rgba_t [NUM_LANES-1:0] blend_src_color; - rgba_t [NUM_LANES-1:0] blend_dst_color; - rgba_t [NUM_LANES-1:0] blend_color_out; + om_color_t [NUM_LANES-1:0] blend_src_color; + om_color_t [NUM_LANES-1:0] blend_dst_color; + om_color_t [NUM_LANES-1:0] blend_color_out; VX_om_blend #( .INSTANCE_ID ($sformatf("%s-blend", INSTANCE_ID)), @@ -224,30 +224,39 @@ module VX_om_unit import VX_gpu_pkg::*; import VX_om_pkg::*; #( wire [NUM_LANES-1:0][`VX_OM_DIM_BITS-1:0] ds_write_pos_x, ds_write_pos_y; wire [NUM_LANES-1:0] ds_write_face, ds_rsp_mask; - rgba_t [NUM_LANES-1:0] ds_write_color; + om_color_t [NUM_LANES-1:0] ds_write_color; + wire [`UUID_WIDTH-1:0] ds_write_uuid; wire [NUM_LANES-1:0][`VX_OM_DIM_BITS-1:0] blend_write_pos_x, blend_write_pos_y; wire [NUM_LANES-1:0] blend_rsp_mask; + wire [`UUID_WIDTH-1:0] blend_write_uuid; + + wire [MEM_TAG_WIDTH-1:0] def_mem_req_tag, ds_write_tag, blend_write_tag; wire pending_reads_full; - assign mem_req_tag = {om_bus_if.req_data.uuid, om_bus_if.req_data.pos_x, om_bus_if.req_data.pos_y, om_bus_if.req_data.color, om_bus_if.req_data.depth, om_bus_if.req_data.face}; + assign def_mem_req_tag = {om_bus_if.req_data.uuid, om_bus_if.req_data.pos_x, om_bus_if.req_data.pos_y, om_bus_if.req_data.color, om_bus_if.req_data.depth, om_bus_if.req_data.face}; assign {mem_rsp_uuid, mem_rsp_pos_x, mem_rsp_pos_y, blend_src_color, ds_depth_ref, ds_face} = mem_rsp_tag; - assign ds_tag_in = {mem_rsp_pos_x, mem_rsp_pos_y, mem_rsp_mask, ds_face, blend_src_color}; - assign {ds_write_pos_x, ds_write_pos_y, ds_rsp_mask, ds_write_face, ds_write_color} = ds_tag_out; + assign ds_tag_in = {mem_rsp_pos_x, mem_rsp_pos_y, mem_rsp_mask, ds_face, blend_src_color, mem_rsp_uuid}; + assign {ds_write_pos_x, ds_write_pos_y, ds_rsp_mask, ds_write_face, ds_write_color, ds_write_uuid} = ds_tag_out; + assign ds_write_tag = {ds_write_uuid, (MEM_TAG_WIDTH-`UUID_WIDTH)'(0)}; + + assign blend_tag_in = {mem_rsp_pos_x, mem_rsp_pos_y, mem_rsp_mask, mem_rsp_uuid}; + assign {blend_write_pos_x, blend_write_pos_y, blend_rsp_mask, blend_write_uuid} = blend_tag_out; + assign blend_write_tag = {blend_write_uuid, (MEM_TAG_WIDTH-`UUID_WIDTH)'(0)}; + + wire color_write = om_bus_if.req_valid && write_bypass; + + wire ds_blend_read = om_bus_if.req_valid && mem_readen && ~pending_reads_full; - assign blend_tag_in = {mem_rsp_pos_x, mem_rsp_pos_y, mem_rsp_mask}; - assign {blend_write_pos_x, blend_write_pos_y, blend_rsp_mask} = blend_tag_out; + wire ds_write = ds_color_writeen && ds_valid_out; - wire color_write = write_bypass && om_bus_if.req_valid; + wire blend_write = blend_writeen && blend_valid_out; - wire ds_blend_read = mem_readen && om_bus_if.req_valid && ~pending_reads_full; + wire ds_blend_write_any = ds_write || blend_write; - wire ds_blend_write = (ds_color_writeen && blend_writeen) ? (ds_valid_out && blend_valid_out) : - (ds_color_writeen ? ds_valid_out : - (blend_writeen ? blend_valid_out : - 1'b0)); + wire ds_blend_write_sync = (ds_color_writeen && blend_writeen) ? (ds_valid_out && blend_valid_out) : ds_blend_write_any; wire [NUM_LANES-1:0] ds_read_mask, ds_write_mask; wire [NUM_LANES-1:0] blend_read_mask, blend_write_mask; @@ -262,20 +271,21 @@ module VX_om_unit import VX_gpu_pkg::*; import VX_om_pkg::*; #( assign ds_color_write_mask[i] = ds_rsp_mask[i] && ds_pass_out[i]; end - assign mem_req_valid = ds_blend_write || ds_blend_read || color_write; + assign mem_req_valid = ds_blend_write_sync || ds_blend_read || color_write; assign mem_req_ds_mask = ds_valid_out ? ds_write_mask : ds_read_mask; assign mem_req_c_mask = write_bypass ? color_bypass_mask : (blend_valid_out ? blend_write_mask : (ds_valid_out ? ds_color_write_mask : blend_read_mask)); - assign mem_req_rw = ds_blend_write || write_bypass; + assign mem_req_rw = ds_blend_write_any || write_bypass; assign mem_req_face = ds_write_face; assign mem_req_pos_x = ds_valid_out ? ds_write_pos_x : (blend_valid_out ? blend_write_pos_x : om_bus_if.req_data.pos_x); assign mem_req_pos_y = ds_valid_out ? ds_write_pos_y : (blend_valid_out ? blend_write_pos_y : om_bus_if.req_data.pos_y); assign mem_req_color = blend_enable ? blend_color_out : (ds_enable ? ds_write_color : om_bus_if.req_data.color); assign mem_req_depth = ds_depth_out; assign mem_req_stencil = ds_stencil_out; + assign mem_req_tag = ds_valid_out ? ds_write_tag : (blend_valid_out ? blend_write_tag : def_mem_req_tag); + assign om_bus_if.req_ready = mem_req_ready && ~ds_blend_write_any && ~(mem_readen && pending_reads_full); assign ds_ready_out = mem_req_ready && (~blend_writeen || blend_valid_out); assign blend_ready_out = mem_req_ready && (~ds_color_writeen || ds_valid_out); - assign om_bus_if.req_ready = mem_req_ready && ~ds_blend_write && ~pending_reads_full; assign ds_valid_in = ds_enable && mem_rsp_valid && (~blend_enable || blend_ready_in); assign blend_valid_in = blend_enable && mem_rsp_valid && (~ds_enable || ds_ready_in); @@ -312,7 +322,7 @@ module VX_om_unit import VX_gpu_pkg::*; import VX_om_pkg::*; #( wire mem_req_valid_unqual_r; VX_elastic_buffer #( - .DATAW (1 + NUM_LANES * (1 + 1 + 2 * `VX_OM_DIM_BITS + $bits(rgba_t) + `VX_OM_DEPTH_BITS + `VX_OM_STENCIL_BITS + 1) + MEM_TAG_WIDTH), + .DATAW (1 + NUM_LANES * (1 + 1 + 2 * `VX_OM_DIM_BITS + $bits(om_color_t) + `VX_OM_DEPTH_BITS + `VX_OM_STENCIL_BITS + 1) + MEM_TAG_WIDTH), .OUT_REG (1) ) mem_req_buf ( .clk (clk), @@ -416,7 +426,7 @@ module VX_om_unit_top import VX_gpu_pkg::*; import VX_om_pkg::*; #( input wire [NUM_LANES-1:0] om_req_mask, input wire [NUM_LANES-1:0][`VX_OM_DIM_BITS-1:0] om_req_pos_x, input wire [NUM_LANES-1:0][`VX_OM_DIM_BITS-1:0] om_req_pos_y, - input rgba_t [NUM_LANES-1:0] om_req_color, + input om_color_t [NUM_LANES-1:0] om_req_color, input wire [NUM_LANES-1:0][`VX_OM_DEPTH_BITS-1:0] om_req_depth, input wire [NUM_LANES-1:0] om_req_face, output wire om_req_ready, diff --git a/hw/unit_tests/rop_unit/blending/blend_tb.sv b/hw/unit_tests/rop_unit/blending/blend_tb.sv index 7a3120cae..db0f09c03 100644 --- a/hw/unit_tests/rop_unit/blending/blend_tb.sv +++ b/hw/unit_tests/rop_unit/blending/blend_tb.sv @@ -9,16 +9,16 @@ module testbench(); reg reset; // reg valid_in; // reg ready_out; - om_dcrs_t dcrs; - rgba_t src_color; - rgba_t dst_color; - rgba_t color_out; - rgba_t src_factor; - rgba_t dst_factor; - + om_dcrs_t dcrs; + om_color_t src_color; + om_color_t dst_color; + om_color_t color_out; + om_color_t src_factor; + om_color_t dst_factor; + VX_om_blend #() dut ( .clk (clk), - .reset (reset), + .reset (reset), .valid_in (1), // only one needed .ready_in (), .ready_out (0), // only one needed @@ -34,25 +34,25 @@ module testbench(); end initial begin - $monitor ("%d: clk=%b rst=%b mode_rgb=%h mode_a=%h src_color=%p, dst_color=%p, src_factor=%p, dst_factor=%p, out_color=%p", + $monitor ("%d: clk=%b rst=%b mode_rgb=%h mode_a=%h src_color=%p, dst_color=%p, src_factor=%p, dst_factor=%p, out_color=%p", $time, clk, reset, dcrs.blend_mode_rgb, dcrs.blend_mode_a, src_color, dst_color, src_factor, dst_factor, out_color); #0 clk=0; reset=1; dcrs.blend_mode_rgb=0; dcrs.blend_mode_a=0; src_color=0; dst_color=0; src_factor=0; dst_factor=0; drcs.logic_op=0; - + #2 reset=0; dcrs.blend_src_rgb=`OM_BLEND_FUNC_ONE; dcrs.blend_src_a=`OM_BLEND_FUNC_ONE; dcrs.blend_dst_rgb=`OM_BLEND_FUNC_ZERO; dcrs.blend_dst_a=`OM_BLEND_FUNC_ZERO; - dcrs.blend_mode_rgb=`OM_BLEND_MODE_ADD; dcrs.blend_mode_a=`OM_BLEND_MODE_SUB; + dcrs.blend_mode_rgb=`OM_BLEND_MODE_ADD; dcrs.blend_mode_a=`OM_BLEND_MODE_SUB; drcs.blend_const=32'h0; - src_color='{8'hb4, 8'hef, 8'h4b, 8'h7b}; dst_color='{8'hc2, 8'hc4, 8'h26, 8'hf5}; + src_color='{8'hb4, 8'hef, 8'h4b, 8'h7b}; dst_color='{8'hc2, 8'hc4, 8'h26, 8'hf5}; #2 `check(color_out, '{8'hb4, 8'hef, 8'h4b, 8'h7b}); #2 dcrs.blend_src_rgb=`OM_BLEND_FUNC_SRC_RGB; dcrs.blend_src_a=`OM_BLEND_FUNC_SRC_A; - dcrs.blend_dst_rgb=`OM_BLEND_FUNC_SRC_RGB; dcrs.blend_dst_a=`OM_BLEND_FUNC_SRC_A; + dcrs.blend_dst_rgb=`OM_BLEND_FUNC_SRC_RGB; dcrs.blend_dst_a=`OM_BLEND_FUNC_SRC_A; #2 `check(color_out, '{8'h2b, 8'hff, 8'hc7, 8'hb0}); #2 dcrs.blend_mode_rgb=`OM_BLEND_MODE_LOGICOP; dcrs.blend_mode_a=`OM_BLEND_MODE_LOGICOP; dcrs.logic_op=`OM_LOGIC_OP_AND_INVERTED; #2 `check(color_out, '{8'h42, 8'h00, 8'h24, 8'h84}); - #1 $finish; + #1 $finish; end endmodule \ No newline at end of file diff --git a/sim/simx/execute.cpp b/sim/simx/execute.cpp index 48d2b29aa..f4dd92e2b 100644 --- a/sim/simx/execute.cpp +++ b/sim/simx/execute.cpp @@ -1486,7 +1486,7 @@ void Emulator::execute(const Instr &instr, uint32_t wid, instr_trace_t *trace) { auto f = (pos_face >> 0) & 0x1; auto x = (pos_face >> 1) & 0x7fff; auto y = (pos_face >> 16) & 0x7fff; - DT(2, "om-write: cid=" << std::dec << core_->id() << ", wid=" << wid << ", tid=" << t << ", x=" << x << ", y=" << y << ", backface=" << f << ", color=0x" << std::hex << color << ", depth=0x" << depth); + DT(2, "om-write: cid=" << std::dec << core_->id() << ", wid=" << wid << ", tid=" << t << ", x=" << x << ", y=" << y << ", backface=" << f << ", color=0x" << std::hex << color << ", depth=0x" << depth << std::dec << " (#" << trace->uuid << ")"); om_units_.at(trace_data->om_idx)->write(x, y, f, color, depth, trace_data); } } break; diff --git a/sim/simx/om_unit.cpp b/sim/simx/om_unit.cpp index e98ee7a20..cfb5e93ec 100644 --- a/sim/simx/om_unit.cpp +++ b/sim/simx/om_unit.cpp @@ -205,8 +205,9 @@ class OMUnit::Impl { } for (int i = 0, n = pending_reqs_.size(); i < n; ++i) { - if (pending_reqs_.contains(i)) + if (pending_reqs_.contains(i)) { perf_stats_.latency += pending_reqs_.at(i).count; + } } // check input trace @@ -282,7 +283,7 @@ class OMUnit::Impl { uint32_t count; }; - OMUnit* simobject_; + OMUnit* simobject_; const Arch& arch_; const DCRS& dcrs_; std::unordered_map csrs_; diff --git a/tests/regression/om/kernel.cpp b/tests/regression/om/kernel.cpp index 3835e42a7..6b7955f42 100644 --- a/tests/regression/om/kernel.cpp +++ b/tests/regression/om/kernel.cpp @@ -6,26 +6,42 @@ #include typedef struct { - uint32_t tile_height; - float alpha; + uint32_t tile_height; + int a_scale; + int r_scale; + int g_scale; + int b_scale; } tile_info_t; static tile_info_t g_tileinfo; +#define FXP_FRAC 16 + void kernel_body(kernel_arg_t* __UNIFORM__ arg) { + int a_scale = g_tileinfo.a_scale; + int r_scale = g_tileinfo.r_scale; + int g_scale = g_tileinfo.g_scale; + int b_scale = g_tileinfo.b_scale; + auto y_start = blockIdx.x * g_tileinfo.tile_height; auto y_end = std::min(y_start + g_tileinfo.tile_height, arg->dst_height); auto x_start = 0; auto x_end = arg->dst_width; - uint32_t alpha = arg->blend_enable ? static_cast(blockIdx.x * g_tileinfo.alpha) : 0xff; - uint32_t color = (alpha << 24) | (arg->color & 0x00ffffff); uint32_t backface = arg->backface; uint32_t depth = arg->depth; for (uint32_t y = y_start; y < y_end; ++y) { for (uint32_t x = x_start; x < x_end; ++x) { + + uint32_t alpha = arg->blend_enable ? ((y * a_scale) >> FXP_FRAC) : 0xff; + uint32_t red = (x * r_scale) >> FXP_FRAC; + uint32_t green = (y * g_scale) >> FXP_FRAC; + uint32_t blue = ((x + y) * b_scale) >> FXP_FRAC; + + uint32_t color = (alpha << 24) | (red << 16) | (green << 8) | blue; + vx_om(x, y, backface, color, depth); } } @@ -34,8 +50,15 @@ void kernel_body(kernel_arg_t* __UNIFORM__ arg) { int main() { auto __UNIFORM__ arg = (kernel_arg_t*)csr_read(VX_CSR_MSCRATCH); + int red = (arg->color >> 16) & 0xff; + int green = (arg->color >> 8) & 0xff; + int blue = arg->color & 0xff; + g_tileinfo.tile_height = (arg->dst_height + arg->num_tasks - 1) / arg->num_tasks; - g_tileinfo.alpha = 255.0f / g_tileinfo.tile_height; + g_tileinfo.a_scale = (255 << FXP_FRAC) / arg->dst_height; + g_tileinfo.r_scale = (red << FXP_FRAC) / arg->dst_width; + g_tileinfo.g_scale = (green << FXP_FRAC) / arg->dst_height; + g_tileinfo.b_scale = (blue << FXP_FRAC) / (arg->dst_width + arg->dst_height); return vx_spawn_threads(1, &arg->num_tasks, nullptr, (vx_kernel_func_cb)kernel_body, arg); } \ No newline at end of file diff --git a/tests/regression/om/whitebox_128.png b/tests/regression/om/whitebox_128.png index f62e87b60366e65656b5b1538c9ace4010462e66..28139c8febbf13a11d8fc713b2f7363a0a9db8fe 100644 GIT binary patch delta 84 zcmV-a0IUCo1CIlcBqvr$L_t(|obA$q4FWL$#IP&+x!D~0gT$7yLC${Pdv^9wvSAg9%(rfcot(%hJaNL4 di4#mjYs-JL{3(2V-FMGaV$qwNQW5`e{`|&f#|ED>hdw-HXI4$P(h$h(#=Bq^qkMf(3Ijvu WFNGyZ2kvfX00K`}KbLh*2~7Zx5E=~t delta 55 zcmYcao1kJO>*?YcV$qxY=l}oz^=w=W!m6{}c_mB}t~4A>VPMEU#Uiu(VQDY}5O})! KxvX^Am9Z diff --git a/tests/regression/om/whitebox_32.png b/tests/regression/om/whitebox_32.png index 97003211e645414e128b8f5ae2344534f3092fc2..96784870f4e4388449642aa5ba3f0ee3e9ec7778 100644 GIT binary patch delta 88 zcmc~OpP&+K=jq}Y64Cnh+(FI*3L>lr)j5_u7s+E{3qJPUfV+DBw95Ar=dFG*FMX*I h12Sk}W_TaY^@C9$)ZxRM!%WE_IZszVmvv4FO#q=|9uWWl delta 73 zcmXTVn4sdK@9E+g64CnhoFO9v1IOVFjrHf9+5EE>%-nW1`L`uI0}Ak%#O&L_Vzcgy QN)t%j)78&qol`;+0BOb*od5s; diff --git a/tests/regression/om/whitebox_64.png b/tests/regression/om/whitebox_64.png deleted file mode 100644 index 86d7c587ac25daf4291168b096e4685e986dd041..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 207 zcmeAS@N?(olHy`uVBq!ia0vp^4j|0I1|(Ny7TyC=(>z@qLn`LHz2GQ#z<}ZKhSEQl z%MHG)Xg%zopr E05SnU!~g&Q diff --git a/tests/regression/om/whitebox_8.png b/tests/regression/om/whitebox_8.png index 297dd6430ff92a7d24ea15fd166388e8b1e53de5..a60e8a43fe31be6e277c8d7beb3fbd614bcb73bd 100644 GIT binary patch delta 60 zcmeZ|o}gl`gTe~DWM4fftD1K delta 48 zcma#{pP-^G=IP=XA~7}j&;S4b>)C7?IOjk45XitFoyWpcs1o*^0SG)@{an^LB{Ts5 Dzz-2N diff --git a/tests/regression/om/whitebox_b_128.png b/tests/regression/om/whitebox_b_128.png new file mode 100644 index 0000000000000000000000000000000000000000..0fd7dfed056b68096a79a571a0fa767e4ccde512 GIT binary patch literal 405 zcmeAS@N?(olHy`uVBq!ia0vp^4Is?H1|$#LC7xzrV2tr}aSW-L^Y+q4-Ub5!mV@R+ zcT4z|@ifhysJB1iy3O~^X=hVUzO6iOox-W$$56=BA?z@R;Sq~~dV>Yykx?ORA@OaF grSbFk@_(K%^RP8^Jos056&TzMp00i_>zopr0GO+FRR910 literal 0 HcmV?d00001 diff --git a/tests/regression/om/whitebox_b_32.png b/tests/regression/om/whitebox_b_32.png deleted file mode 100644 index 3152f632295be1b33f428e2b6eb2ed796e902a62..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 167 zcmeAS@N?(olHy`uVBq!ia0vp^3LwnE1|*BCs=ffJJWm(Lkcif|mu~VlIPfqBvIm-M zSmSbArdu=lkE7Iy7bjDz1uKMq#!Nc(&6CZAV@--f%K=v)BL~Fr1u{&G6r3I?GApr^ z0T~Gs7zLT;@^lDPoMebem{u}@aficAGsy<-Khhc%OnVNnpV@43!SeaTmq42tJYD@< J);T3K0RS=VHogD= diff --git a/tests/regression/om/whitebox_d_128.png b/tests/regression/om/whitebox_d_128.png new file mode 100644 index 0000000000000000000000000000000000000000..f62c45e8dd0a1c1fbfc6bcc23cb2a27b14d1a42e GIT binary patch literal 27123 zcmYhj2Xs?M(7!ERt)#Ur%VXn)>(EUz#x`KUn7aurRMQO@Q0UbtfI$HRx+@Y05O4wk zhpHrmUJRiqV2Bk+AXJB53@B8Gsw4y)>NmXSeb4vT5#BTV>}Y>J=G?h+cSfIOcWoaW z7!#<|>4HD#(545C&UgQ~iGCi+`tZ0;hq@2ih+dABozMMs_KeEF|N_dZ_z|GvNe z|6ig1`)Wnw?^QxWYfP(7uV0rnK3Gt7=z_+-M!);9`hUm&F*08KG2_~QZ(;{MkDPsI z*2m#YR6d&X_I?tV_P_jt6O96r6W^@(4*AdL56zlR*fX>8$wG;|S;*Vh-GTk_?2lO) zffCDGy;{^>!0eZE-u_9>e!susQ;~ty(tc?W1?AUig1{ zx9WPrn?>)oczNjU|0r-&{`78(H6d@GzjJ=@$J=+#1+8EG^v-$BoEz^}LE_9?^cmRd z)i&~#ul)NLkXZZ-) zTE6nWgSY&1-oaCc60-<;LT242=xg=r_jmBE|HwP29ZG!XIomqxCP8Ta>NkS#J@l`C z2R}Qs=AGwS>#Vy33!A^X{0=_$-+Ko?Ike`T=Mw9zKL}!*WD@!< zft`Ma4nmt-@#_zTKQBavOoswO#}n9%XIQzztc?m$Xx+|MJbXHAc!p&Uh5M?jwAk#W z*}JXymg#WmIc8jQmF_lcBSVxM32fwaV9$|15pd0g2hCz)hYrwods{*Yp<7nAW-3+^ zdLGYbK34I%rtULn1!m=rmEDi$yIMlUR7`oS+OBA&I}O@+vvS1BPQ~-`mf)C*{wFG2 z(r}kS3^U71R`yvu@6i&5O~Ip&)a*-|o4|0hykcda$MgOzVe%BLc&wJ)&}<%q*I|}_ zS{W%?+7h-*#Y>OX;v26I47L^W86c#s@B0Lm{(--Spvbq{*>8ojQDJYl@> z4-1kb2S5p-^O0=A9&`|z(~xHoBCiPp;MFCm?lNk_f|SdVY+gg&YXFSegJt7|jeKbK zpGbD0A%8dkR^|!c6eNH^^U2Y-=!Y5et)>VZ{VBrQB&N(R(KwBq=%` z%rC^kxMp~1!kfp9%sDRK&oOM-1@!DRd3QJAd0T{iY_M!O7Rm>>eBKzg=mI(j{m_ID z+9Ld8f@QO@P(tWX3|n&nD+%pw!e?#~*x#k@G->04m0x1mrY1anEL2>;6#6)t{Ww^8 z5W}7|;axAHeXH>GbCvEjc`wH>ZxjCWSZH$*=WZ8;xh_0l5>tZZs~Gl66Mk$g^ty-} zwu`b{m-~Rp`(F$@(1e@D!h#D}v_+&0aM^a4yrk%44BOv?TOa&1J$&~0S-xX{WjHEX zlns);f&h6vjJay@cAcQ}QY>mHigTom6paXw*TUGGJ2#pK{(Q6b9bFLQ<~n_UE39u6 zdJgKmCyhL>k!qi+moI}Llh9uZ>lTF$LdT7KP$ShpMK7NRK>?wE6gD^tD+!%8^20&U zAqoo{sqS2-dO)X5(JQYN_N$S14uYIWEN-a!a-5|<>a^*4<%z=H7`YMzvmmS*>Tu6B3OdK3ijVR^`d3RQty!xj7AT*J!@G z7PfaJdRiL1udDKeQ>x!#lIznTlh7Rt+ddKuu z1~I}U%Ps7as=Q+w^csO1PO4=^nw!8#lYDGpBj}9<(sw4D^<9Uarbh4Qb$MPFVILVNr}uhjER!Gb+3P+%<2xouQ{+%%It3 zMz1YUlBC6T`G}s-U>)Xk5yf9i8-WjB&|Oi=e~X%zG8R<_J(+s%`2fDHn`)n=$g5SL ze3jGphOv-3=pb}9fUoYR`X?&#VihPM^qR56I#@~QVgNr#Xp*A*#n_nue!B`J)xqp; zs(Y1FZLQZnQk1idEe+t~sz63N#eILN>+<~?&G!C^o-~v9Kx3Y;K=?;lWSyaqN$8how*6Oh5Zd3EXAvT= zMnj>1(2;2NZDU?@C>;G2vlj?=j7v>3X?YgqP&6a&q9a(mP}mY(zR4ElVl-Rbm=_F% zm7fcUb77`QOt#3%X!cuUer6~PJAy|SiR?I+n?SxrzKCWg8}lo)e32-Nb=gu(UXt`% zH2a}3j~ohHe#J`*L@{Ydd9$Luh<@EZ&c@{Y8jzcy`R;Reu|0a4E8a(;{P9B7K8wjY zH6WAFpPZd;j}Af)LizQDs(&Vvi)ugtp$DAZXpfc670tmETK9J-Us3}$w8yfAsyk6D z%}}(_OgYZko=_fF12Wp9eUa*m*GgL}g7z-s?2Ax7qz2S%haEmw3#)3_Oc5V4d4{vE zLiv~)klqeQeXf>O)!Z!=L6ZL9>{cjWR|8J8$DD;~ae`)Rsd#5TZ8o!cDX^$d(UZ=+ zUsdDVf`xr3-BP^|d4JLKJ=3wNX7seyc@OD%LWt_0tdsW)P(tWm z9h+cA2caW+o<)efo*3X&oKwxzX_Ixze>!$d&(9fPlo`uHR9l?0G+n0+*C|(ZY@wd7 zG=Rs9eyi$BbYhB5jMB-YI<`a4!wj(2ghwrEVXV_lV6;vi)3M!p9%F!ACakciW${io z?Q7S`zjW-Zp5HY<12g7?sKxP4TkFaPKgGEz{h0x|_g%h@iR|JF^!!fgPZodtRM@vu z`ZEJE34M^rPQO40q0L$R`cvWGM(NKCC?M1+k==NKmA{*{4Iv7x+nL3~XTXLRSoT!7 zAGk`(%wC$ko5i=xfJ-ki7CS@aoeHg+PO-Bcze-ikhXR3c6r9Z8pgwW6s zwq`n35_+D%XFgN$v8G-zXeDN4cnG_nz<0HRis_j0OtszDO3xUyt!AZt2s@R)%UgkC zI{KfhP@&;DgIH^pIU($`1m2?+44Z~W|5dXqG&h0o&9Z+8`#gd7Zv~U5VZ}4G?1^T( zWbk^-a&!nIMN3=3mg#uunOgkh_3pveGD?4vd2+t6e^2SpKq#+}zCKZG>OS;br1U45 zC*%wNT1tNgLJ6VXQEb9KbP$@8%(DoQ*Mx!aszR#gjM~~DrC$`Am&|(&gi-sjEMM5j zhh`^5u?xxk;Xqh9MfmQ^(z8ZwJEcF#Tn>cU`*H3R0Z$}eHi~^ga$FSqFPYC92ut?k zhAATZiF6a#Pw7uGpEeLI`!G6Rr2Hvu7mVI-gCr?hk<6zIgt&dP=&{DVw^`;4v-g)t zkeg-po~NO;*?W@idC0m(16hYAK_*#eX&~!3-Sm)kl?JkoOo9TkF492OX}aq~lF zuyRv8ENFp6lhu?5PTTK#FDdF`X8Wz&y6UHC;j^FpbVPpzJ(o=0J5BkvJ;HtqAOUuOW4M`N`E(LKLjgNW7+kl{MC41SCH=(FxrI|Oyb*M*(;WP z-IPxp58+oZX}8E8?Q#?NE?D-7WviOsTi*~o{n5Wm z^^MhhTa3~c#Quon-!_E91F*v$wJ=Y^tBN>iltY5pok+f^A)Fh4qxPs}d77`pCnSl<=sxz4;NYxBHj z!oD{^zU%^-gg&amx~)J5q2smrpk~6qCqO>$0tJL7RbhiyU?ri`wfW&L&|w7@HWTh~ zQoY8sJpsy$D(u(VymJ@GS&qd`gm08Az09=z0m_6b>`iU1bb;B+ac;VRu@Zl0VpD+Z zT!sBun+JA*CChO`y2u_Y-2^rV$gWk`wc7ke7qG0r=w>2iytL7_I|F2X6?U~Ye|GYx zSI6}+=qb~Azc=x1>8gFJUd9GP`533KtARa_K?kAlOnh~^>ffT5QNd6`sEdI;iNQ)j z>rMP1A@WKLhPg3Vn6Bbjr~122+oV@$**+7u1jEX7)t2Why`s|&=;c`x&kY7g4Emd> z_6s^ute18J``yHE21CtQ?9fat9Ph*%I9IJfnQFhT6WjGt7#K;a z9SkR8FehCt9^i9B^tkZ^Th6K!ILit`;S(+IUjPjyL=}}8DRFl zZov~S2*1ZF*XKhfq2cjt`vY_kdf9?!5hAbc`A|S;WIVgsg0IboqYp6qf^cthsRPW~ zL8~$}o~>)a8|1^K`)EHWd|OH|X)XAB`OxM*&OI**H@L93S*)|l+<5j$ z3*Ipwdfmqj=SA5Dmz%(PtIUgMB`k7qAm45l=w3FL0j zd_96$-!bUvWAJ_+!}HFo_Vp$?y$NIz>K@FxjX?*Y`7wOZdDXwpB&Rlk0zy55+2Ao) zNoY|FZ_@-ijKRY5s(XW`_BLqiOv=Du_Gt|7ItEiNsV%F)WYqWo4}7IIVPC>8^f13fn8%UgJ%00yl#^uNsD9n zh$hfr4Cb6yi#KXE0^dAawoqoj-uv!>_)iA{<^0}|yG{DWM6ipS(bLc9eN>M>?j!6! z19^jG8A

==P<4bzOhRs;k zN4U4k(j23YSalIj9GurzK-&R>V*eKkAGCP8OQI8Ml4K+7mhu)%aqr|>O z@lBxY5y8Hy$H(-B^i4Rbw?_CwhANQuKoFdytKqjGnM%JhvItZ-@Q+f+DET_t0=vVY$G*_?fR+Q;Rwmt$p_0Zo-^=)?g zwkXnOWPb$mZzJGveeBRjE!^hB0eW#zkwc8^P9WbD0q5%Ds6J}hHm9#dk%NrP70BB~ zK<9c`)LSjy;IxsVb&4ElWOI%U&J29-zZ*Ao!y&hU%Qvwq>w6qMgG}C&sXT8PUF9tD z<#5O(^ifsT?KnCJ9Z%(hmeEztBA*Y30z#9jvcbo(lF;c?et0-^IF5zOgu8)D9ca?_ zSdt5!v-!ZUUPvvTId#EtS6*4wmB>y-cJublGUzofes2m0eBc&+h!>Kd$eH zo@~YYeHh=igsyT-#@2%JdYZ3m0DImM9fZCMsZOK|`e?>dMKuTMHZ=(f^fd@1uxf zChY<2_b`657S!y79hRwu4K*CBh#gE$3}ENO`2AXt-U&x7Q|*Hkv7Jc~z(`WuI*0igz26@Bi5-&iGB`go)1TJUNoCa$Ih13%QM?Z$v2DJD;v{%=@}JPl%+e z9GC0cLMEZ%p=|qny2>%{%j!If5P5BH3k8HmhO(R0`P#N{bUtQB3U{(p2Qcj*SB8eN zb=7%;ws2`4+QWsffh--!w7p!J70R^g{Jpl&W**Ls5QRw+do!_)%iK`*Np;?_E%cg) z8zMwml5`VT&t+aHBT1jOg$47mC{m;}lD2-#yN}E9q3q?&HE-|ejF6k;^!3oQzJYX= z(|JEvcwPiu<#cko5i$vN*RyVcbd}S2=PP_r1YPBHa;gyu2=&yn!GUy@(|H#uyp0h$ z1Y%)?>P~X1y>;3;oib3*K2>ICSQF|_)|mV z#gCzc(2N9@_zzYRx|qoi5=skEK1yI`GWqR~A?Y8?z9-x#UFsUM);L7znZTB2@^K$S z#@}eaD}1M1r7O%@>kvtMf0N0pe+-BJ#tyed;V~C3HjCsC`CbD1K9kq`7|#8Tqi&0` zV=gy=W+9R!9m?dtehi)e!J>Pj_=L;0%Is|#B4@lg{d%G^jwjz%?RCwvesjn@uKC(o z+1`ohS#0ngh~o*jRezFM)@crzgxXr!_KD~qv_Fn#5hAZf&7pu$dn@}kj@N7sM<-(T zZPk5DQx_YwWV7b}bhFae%2vnmg66RDjtVChHFsi%;iB+2iOYi*5CF4Jq(43fN0RE3;0EFP)am+Q6q2Bq7tg&R-&UoC)%9klRMsmMpDaLS3x^&-A~{eXec3+ zAHyb`MF*iH={$=Nc|92ouR6Qbl_o7cSot`H9ZTouM#HGHShil+I=f1jn6z5K%77TQ zFrBX)4W6^;UnhLsTv%ulNx`y14BL^;!$!l}GkA2ZDD32N6Q~<3JI1ix={#mM>^g%L zYeiWXmz(xY36>wlu(Rp>?r3Om7IW5%;w~=Rnm1`b#i3CK)rH)Sny+IZyVwIgUnt(K z;r#Jh)!xV`&2=G@&jLYEe7y=?G+eESRHK|6 z$S#NTzw1JHPfS{`W_Qut1k#LhN+7#e7oPUO=(TD}N6q%7;%#7*Z37u8YFZZ__P~s_ zZ%%ZFFW=TPcaXjzVQkrA^ekuIe{1r*+QQy4K<@tl%DYM5hhc2dVssGtrzRg%TliZ9 z$h{vx38A53Y|UbbbHdnXHF=K@VAz*nxt!RhOx zu&H6_S*G*47*D9B`fKXt$N(rI)LUT_!q7oz4&zyb$ZJ9Xyz1aom*}*bdZnMj<}uzY z07ivjSuNE@J~TT?VHX&G7yv74tGsts%1b*&P6Ub^*;~D(IM=q_nlLP0>8PF;on} z6#6)tZEjM=TG+E_-gPM2e^Gs(X{AvHZ$At3M)RK=Lz|&E_lR1!K*M;0XljxL7WPXt zKh_v}4aE&d)UpMdJKo^^*uoA(b5mnjFa(PZt10s}Td2WHin>_X{%CI9(KUP4P87P-^hY~{L!`YMXu#ym! z*c>E8UTeC;-0!fkoxnv>t!mU-1S+(Q{tGPV4n^N$aT{TqFH1v>T3VnoHJn{f;;*QX z>09J&1$-tk(kQA2%3k5@>m)w0JA{9WNo_^;XVOieMxg8y&Q>Myecj>dcNpDHq%4%S z7^9aI4G(7zlK7tP@bJ4?e=jq%QSH?z6$*vi&nOjQtXXyRL{KWk`Qx^#A1M_Ig-k;I z7;97=9fVeJew`3`bqj?8LIW6UT^*sVik~^vNWE5_QX$5cb3Qo~qN`)ZXU@_%z1DyA&fA4T=jvG0RxO_Ew2`9PlnOC6=ja+rg(?T%)O`rK7D|PpS>HqGNuX3H zh3Cx?_GXj{eF&L^K8j}D4xxk4@f1F2j_^07ROmw}AT%kO4L*dGgifdM!yiJ2Ls&RR zxGgR<-lR39R4AJLn!-DO2suAv@h8H^U8Pk`T1!fWqS>1iu6ziyf5y491q8V;)FkRr zDiqCrOyPka!jhkH!)%cq&#eZI_t5`|YL8Sz z3Z+7v{SN$QHK_SMc9^3U25Xq8h{lu(adsZ~{c4c@K8~8B+7lGfh*BZWNK)-;aN<48 znXMKFYPKqhH*M*)~N`;Uon1nxrQXv7E zgoXp#?m`D4IswTdL|)qk6c8E->?ZQH0*<;c+a%mpsm3xbi&7zA>yS4Pa7m(_37

    I$mlJO-{Dx`;8PN|TN^;OUlNvV*I=NVOd9ZH4tkV&Y!j&)PeL1?~? z4>GF$+LQ|Ep@2|N9UH7*C80$+-bN1{6f88VZthegbXsjng>>vw9q+1OidnS~R`6POHAw2KtM!Z>TuolnN~k;--adOoL&k)Iu`;u(0o`I^mR> zO$OCBY)ym7r_{1m20=9qhp0T^w3{9hl^fM)V1cxmBRj8h^#*o za{qDpK22nMU!mug+1si$PpB0Bqam`+OvohkNg~_+3LS)6w&q!c$g9yzC?GT^kyUTa zYtDqDuQ0n(xc_#kx6In95XF_q$h+t@7QYa-=Puv55M@~+OK8mtX2Qz<1Uz-&L$f#+ zB0Y)hgVy}aOc?eWkG>SyPhD;T=R@SaMAoG>ze3Ajin4!PHlNu`l0HjhX{~wWOxW@Y zFI9?S(%^eLw7u}Z{wJ=~EazuI?laA|DTG~|fu27N-tme2@e9>{!z|}yK_;P%A?)-F zbPyVw$gjUp{nyQMQ5F;s+8n}e%)rV&4caBMLhB|b@+DcYVFs4HP~Fe9(h7t2n^~D3 z!kQ-XxGcz+f%ccG?_aI-i9yibg&{07kq^m&ng!V5Keg~L4euF**DSvbVG)UZOctaU z;Hdx9vcEL<1A`z*--NKCiF{oaoS1<*FVx~^n(cwXd;NL)SBzJ4nkEL@zoy-|Ib0P?;t24v?_`vmS82Js7CxCp+iB+7f~#% z5g#@Pl1ec9W8wZws`reVH%OBAi9wKa0E?#!`+cK!JV;p-#j+alQzgi!33w*)zEPA1 z$+b~zbR)iV5Zpe1Nz+93GwCL9El945V&iG~0j&5~r1+(c_B|XVzm8%-jd-6yFs=kI zk#heVs`hofX>RrDy=_By@^saH-XI?&K>1UrZ;zQxZH1mEdhh!oJYl-(KWC8l5}<_8 zZZn(E3LS*nhwv;y5@1v-ESs*{o;pi!>$P7D$_g{H zhVYdM;Aw^aX{ztJ6McGd)F9`YS>q5MmH=xr@u)*B{Kx4g@S8!-GqWZkJSG8lWnzUx zE&JE$rhUr{a)+7q4B>YZpg}9lnXVT9>$E-juf|VtrNJ^N4{|5EeDh-2#d7rgVe&R? z#vglx{dBN2=Rqc+&tlo>a&!=?-;7`P2>+>I8IlJDgyzSx8|7H}heK#7aU_qxj5yDo)YV`-)a>RK5>lW1{%3WT+U3Df?90M6L9uqMbAPVwB5+Sa1~YpA3_8v0|TE zHdV7#DBep(xiN^5qNT~OWguSKrxs6r{l{Rdk9m!Ccye=L|0O_04TbOuo4}F%yoDcA@Z8g6<*n;dYfs71C+0;u)sRJS63Ld z63dzk8~M=e4^>#7I{aZ*SlL4OCd$%VOgkMQ(oV=3j4SY zpVt+ZtilZ~MD|qaCU7-CPO8JFbp^{xjBYMcK9aV(%=>eIBt@}x_>`^?w{lix5?}w; zV$Lvnf9Xl*-$w6w8d@8@CwtQQx6ylz2C@$Ir1Nj1_bd%$9q&o!-$w6M8pt}*lg_`5 z-itJlb-E|a?MCO{M(tHNJJu66bffccqxL48{n`_Xy3+Z#QF|257WRaUZgl=_)Sid4 zH$8!MrSos2xE;=R^n{vJPDaMvaQ0(Q2=7Yg-$rpKobB!j>0Rmk+bHgZGZOZ+8=Zd} zy_MmNbe-r%=if%}i*QD=9(H?M*2rvk`lNv^i$l*%owtdZ=Vhq&<9gZB3gwfWKFz=u z#i4^xnwbyEQ2odBvY8c12)PVwO&nGdYG&p$GgO@DRB!6E6ME$n1FL7|=~k$S!xZ{B zn!T)7HXGO!Gw&LY_U5YZBWLM7o%d@46K4Li723q(+?HzLWG6n-i3@sJVqk5}{FoJb z#p8ySYT0C``;pH3qk*+DbCVSo#9>i}nlj00`$Oj?MV}j3OEb4FY&|)gy=;9_{}4T+ z%--QG`L-*nSa;sc61xg5Qjb~3DVI`rFE%`x0 z6{|warnTe?ra;j{EWRXcKe$RiG;3$A%HDW3s3m_j1=vI6 zmj&!_;Rv(%*(#UBv#Krm#3>N|5R)#8>>VyQfkRfgG@iw`V`9VU5Ov)F*EG(7}YYIu@F#EFV-l3@@4Vu>^$@@f8 z$Qh5tS5^BMgLd4cEDB~>vHa9HxA#2P)r3urc-du0Bw)6;t{P+fG?J#HgJNROUpmuJw6&A8>3%{{R8S5{DVZ z;Xqjw!2;^@9(`cg7ChQdWEV>}fnNgU@(31OpZD(rleb_+e^Itq+C~|@mjdO+2u6yQ z_JJ*1@ltzjX^xQpKAX_kz*sFyx8emYGhLzpl77sYY5^A{Z#*9MUISw z5<*`Y*@On@AY=~WS%k=ILL|H@cB;el+F?cc+QHor(@o&2 zA}0m$X^~)QfYJTblpRjn2)*}bMUtY}AU-7$;u^f`AG$m1OGA;c|89}TM?h{9mv3KH z*6bvD#+$qY)A-{R!e3^QM@B#V3nClOW%Y~oV$ zOxhWXvb!qlpT?h#fasH$^0lxvc9o7WX@@Mz`l>81jhBpo%9F^81f;rfh)JBV$k|m{ z-87y$0&btcq#}`>>T(k}X_22+WewAKixDvSBvz~tDd{fTIFt9BMUtXUY5c+n74$pE&jGkS(7-X>LeUJ+g8m~2@a%F{HT7QhyDMhBs^Dtu58UFDc;RvStP zxdPal&R9vPSrtCBNX1l59inI_nDR*gt5=1m*M^GDm_i>%vzM8&Ie<;6!n<}s``42yXZD`X4=dMr-(={Bghzm@X1hBSM__5m1s|#*ep_Zj-3;K#Yn;COR1HkS-2k3P z!{Y${n1(+BcmfT7QaLaUcLMkz8Xg7k>onX7;8`?0pn6~$?gj9}b+B-?Q=O?OwD4RV zEF)ux9UwTRSKqv?7$XUVu*|G3`88_J^|d?ZC4D{n5ggCVeNkbceDJYw&N| z!Qs!bLyRbFDse0m*SY*Qlnt-JH?@OvpW~<)QPx!Y%DG$}$_zDlyLQld0Txvi#VOK8 zihkvCRVXvxdeSmb{*ic3XM)^xr*9{nO$MVUPv`Ao;CZpC{i063G(jez9dtGsj1EGb z417?m>c60q&rMK3=m$ER48}@A-3PvB!j@D^cb;@=o;J8jM z)U$yG{=x*7V2qAcQ<^w!wCx$4+^uH=4E))rH(nmsKSht*?5)*`Z>tdYE+H~@8kC=L z`4%Ry=TFf=s75Qkx8l?Y=qaFzRHnZp!BDDlYl4?(b6HhUxLKL5I*}gM-yM{>D zo9EP0`Fznwx_F*^U$yr$%grqy_oC)oM`x4y=-F)W&WPs;_f>z6S*~vZnS{Qlv&npP z5Skv(vj~yb_7+e;XuXvUiRWuuz|nlnzOTA3XzFHzHpr}eXJxhHd4m>kX)@X?RNpzR zbgMz@YgT@;GJQOMuLZQ3jC21~3(GY8-XPkU9~>Rx6u+?a6<4bp0W>O!~fzWPNv`XPwa-+>qyu7WVc* za(aKrB=kij>$VFWgn}CKL8FDgU67pG9|{O9j%0&(VI`rEhP+LG=&%b5M+RZfuH5XnAj$iM6lYj@+( zF{1FIbkj4eenE0mB+GBem-mNVyD@`i*BiZRkR(Y}8uAhSp}{W987+!`mo@_8<@ANp z{a^d$PCA=xf}SmUZ|`8fZLDhVX^>ZAp!~ekcaY8|o1lYG&tSfKtm^M!kQZa1gwO#x zn{0xWg!%^agM@k-lpQA4BbeWgfutswJyvy}cdFa>IcHb~mLS}?C31BcVG!zi`zj1#}pi>?MKO{`WhuNMR7(s9%%wd{=3O`xwqlBD;8 z`L8k1xd|4HRg2F#ZCmx;J_dQ7?yYrD+xt4;Rgx5AFO;4!@`>IVPhca9ASLyjj>cnY2#9%Ay#S)r6lqhkTQOJ}%s361l;0Z44XTgzp>!x6fnJ zCXwC8ecc<>qWk zS_9d#-sssxNoxep+oakDP|{it%KK`*Zvxq(-sm7SDS{8$r26|)(pnEn2z?jG*7U|o zLQ^C7%uOox(bP?fmP<)%AR80Gch!T6-k7pUwe`?Szg4s@l(Yu2t`WSv9yofVf3phN z8m?7D8%kOOSwIBuQ4fao!lN72>}<_Vpe-e>fh;(J_pb+&dtt>UwXCmZ+n{)JC}|C3 zq-be9*wPy>ZBmQ-zP{VYoGpEl&LNke=UZw}ti|(^guPRMY}paYdrF^1=a5U#K`5;j zACx5g9Rp;uj!;6#Mdy%9u#!--T6|`bz#dZlmT8>>luyD~y;?lIBUCKG6#6)t?H8bI z4r5bl@vdK?y{_0Ywr`l16n!4XTGry$V>jq%M&XBg?&_=dzSMIU2$_U-P|sZiItX>5p1b<0zYq1?1wsL#AE@Up0xJo1qn^7! z=n#R0^;LH-r@B?A_0cN_6xNPc5+J_8N}-7Q?c5)^j@#(}P1tKwD6-1pIM$u|?_Oh1_Knim3l? zIC|<+D5CzmV-$+0|1K3W34Krfcf-*^Xgc-Z5hAbcsZc;@y@d^-{<~B-Ivlf)sqUqk zTAxCZN%_vgYE%DRDqI?d_M@tAiB{T>LXk=N$-?wi`Fp9*W*E->O)V_cuoi_PlU!+G zM(V#ygM(k0QYdQ7w|z#T$RY;}f$~raMbRwe zFggfTY0OuDMxn?e`woE;LaU-#;$f^L6xEm?B=o*T`68NyHRi*HK+<8%{!F+-U1}W) zMHWfkCx$@IFIYTZ*y~X!vM7t9Syp3y>M-(o0wP^lk3x|}u8n4+8}pq*;Px+=G*4tl zy4(Z?TI9NDHlCLMf)$^M6mZ#S-!>Ncbu<|NlY6T}38C*eThkUR2~7>I8|3~+4^UeR`>i@W?=>}muR+iAf8LJ$);jo1uV%%)P} z`Z%|bioI3rs}^o?svDiyQ?K=ZR~M{oRb&q%n`UIasWiBkn%!3~-Z!!?5wNS5TGmG| z+8f!x2zW|HU%gi~viBokK`%9>kKWtH$a+S=2{MShXJl0Kyrq{~Or)KW4Tyk;WDpTX zMkUWHdr|o-9aRnkj6VD#C;4SiD@&l|mFUqLEKHKaqw#x4D>{bh2J0Chz<^B2psl4xQ_U^RG z-tp`~3*Iyz#y!AG=SA^$mybMlPy~r*_o%agS_&Ru<@YXMSF_h;m387-qZa)7WGJYh zW}>%$?$sZ|l&^%XuB)__Z`CuGPYZ!mdZaphS7Ch zwk;O< zjsW&q81GRF)^@_9OVz^qnmhad3Y@bmv+J&GRWuu>+0&Gqb%?2J8zzu`At5b z2?UY=@{l2fd;kH&1Ow)MAqh}I5+RHsiVPq_5==mBQ_QK!yuWa^>4G{fsg1ia42w{d zuNeG@iXj8FQ~^OjNf&6Rh;H1XQqct!aIbW`XV3n5&vWkZ9`3o%^M2g>%ZWsKWxg5! z0Ov>|Jvm?HPe6S~@G%6u=d0_B_3%;`_H#Mo*FS9J;rm_K!R60>ed~+0QA15%7WXd{ z)?U)*xW&FSno%sT|Ffc@rp#g^hgG$(Dh2iqvxl&KaH$(Ri1;QAo5&B@=0U`plgx@m z!d&VWamDgZ1f3TO^M1D&RV;rDfR$xrIZ0+W zR;t}%M6rAr04opTl|!){nPjfR$`ZF2j!-uMRuGcDh8Q8&a3VEZ{P36Gu|1WSZQ(Ms=Jq za`#$3v`P*-xHwK6OC|3#xo3?UIAe5U0acGa2hWZZN6{^rCOg-t-ZKc@Tan)&rG*#X zspu2QgEeZyY1AK(XZ-HypB>`^{gs-vL(}iKi6?b>z%DmLt1C%{93N0vKn(jOO@Cg~ z5ypPVCVpfS&**eUr!VcY#xC2V)u|Y@D=QE}DExJuQlnK3zz&UH(dq7JH3;xqjr(+R zM62YifbThtU({(_wCVr==QW*%A;S#-&dWO86|DvUzGXLFZWsa)T+>;d{uF~-gh1qtD^C?7{{2O4&}mIi(AV4{{%ZzxJ|=swsw+rnpG4}z zY33*J+&jdJGw9-Dax4I*SJeOjI5%fd z{8g3zZfgXbt21aKLh1nb-64MZn4EFRywdIr&%Z;geN6rg;Md>({LhX*SpDr$R=QKa zKVCdJk{*~YH>_7zkjwsqrLfTPLo{2@bn4Tc{7}63QM`C&B%K*aUrv`b(`Ea5b?Q;I z>m4gJ&dINjq}27Q1|Z$ZuZ*O-*Q-H*3@7)EB*%J{{Ep=t>*NN-Hd zwFG$PR-U86^m$F+V$*k@8L2cQUNf@TjNFH~4j!;gd9NuKnesj}SYZYinnAA_TxEaWVXQ4y7V&4v-$hh^yb-166QO z9m9P&>C@yFXqs=9;m`Ta9+!yVLecVec-&DLwEH+72Iye@S8al-vd=R0?vUP zn%DzXI0DW&#Osl!Kh69cyRfaqnP`|X9SyB9BhxlVGtxu6dr~D&e{wtgL6Nh*?MV?uWc_TLSQH)ed zkuL!}gRvG4@n(wrYPFHuVGFNwi0vsDKw{)RgBeQ>aUew=Uu|?_WxYe}PLX>6u+r!d zZ=}e}tBrM7dBh=pn<6^_u!4W*&=l2sED%9J{lYVz6{adY^TPy!iE61Kz6}O}1p`B*1ywyO!9Fg=c zhq*XaOskb=c9^Gv4s%hen1HA`z%L!QwRPPrtwFp?6xAYj)wsS1h||w2LykXs`>E$(ro~_j}9j!fh9m?cx!;aM?ws zUDVnI3UAZx!fqG12V7+rNp|r;G}ZmRkA}O;J%P|-oB4*$_^k`k)P#G^Y?X|g%RK>K zr7hg6^E|sa9Zi1Ra%QV4+*@J+UA5PBUSJpJqNx#goY^WLx0ZVX5inju^wciSMpJX9 z>Rp8DT_7^ghRzV@*~RE+6uAz4S05NkrJ1s7q3YgYMCSgVc)T@~W@XB@Oj)!@1sAH| zA~koX5qi!D&a(^z=C|tB)4G{y*Js-C{7H2S)vozgXpEhA4yEdass_Mg=Y2z|X`vdN zZ~2h)-!_y=ke~-h!z-(ycr?YMoaF;&$58SvR8;`rAphUEPz}ttBH(NvN@WPw1Av3D zeZxZ4yTC$W6g?E4t3s+$s=xGw?*wL=g?#qe}cgL z7F7c<)y1EWrKT-vu$+ClF8<3{D%qlv0j9VZzvSP7=`idA=i{;D-J+@h!1-(}H6m4C z&Jl1v8cStcR6YPWe?`(i-CXw@5uTkW%FzQDYhHO;gpuoC_V%Hh^#MVpE?JeRy3?Uf zp#RsJm0{DD>B15;%O%@fvS^43CaT~NHFum58Vhv-#)5#_s##BKW~NP_Y2z0J)g?lm zKxmAOcM7UbggQSwSMxqZ<`bb#z^81yO;8EK^8kp=KQG9i2z3HJaCQjtCPJNn51hS7 z(ditwih zZ;EJ45ojwmq=*MoM01KTQ^fKV;Z6~gQv|w{E_^3V5w$5|Ns5@Bg7-F}I7O^V5lJcH zgDF(EN}f5UYTC_HIVcZ1%s2kgd~8}Jze4k|-OSBt^;J5;y|FwmMVy{O{#Ek$F_Z_* zZY(TzgkO*41-C1ajjQD4V=BMhT$j@t0pqn;o}VJlPNC+vQRy7Idf)`crBQ^K;X6Scj z@E5Jd6Rm2r(;DV9GvoD{@%%zQ)m>D(HnoPv#PiO4s=lad06g)$FQ1w&s=-aIJ{8Z~ z@~PyaN(M-aXVeA!7gYxUI6Lymdr?&ZfP=n3<3%;FsWk%5_IxV4sPX~8*_}@f7gg^Q zt&w}sY{;jskJJVLT*1kF>ieh#JK6g!4{0vpcL^`P_qqgnpbakZpi4C4TdzwjcL}#k zOm>M!T*BoNnJ!W55=&fSx=Yyc&DSMXxkS?K+Gbt4JcDPMcH>lz6A6H|OqZAOAk%KF%dsM0 zye4?QOPo!l=6O&B$$7Mb^zbHp*~ZfhDyh#Mx>X-Yrqao>Y6Vm=B5hWGQk0e9(3fEv zS2E3-EZc4u6@n{Na3xeRLccSDzp#ua5Vwx9hDDi~4t=JB(R`>|0adV*AqJ96)hnP1 zb~3ypOs1w4Pz5`ggSRD9$qJ}~oy@_{CzBtQ1^{q&B;#EURKZS$*N4f}xB{wRCv)(2 z)F)Oz73^eqC5o!X3aEme%)v*KscR)v!A|DjlgZS#vgM|S-uu2^LsenB>Mk*TPjNpw z3OFnf3o-hRG!=-c`LQK?&2 zrK#>>s1xWHxn+?@1=CdUc2gl#1a$(&69F~aW(~8MnYunx=NE|T z(x47@GM#r4UOzz{>|{FcBWg;6I@rl{-i8iC8q~o~rsI7S`O#(o0A~l0Hx24wC)0T^ zQDYj^!A_?0cA_$rDFDFPP1KMEb+D7^e3Yom19h;I>3ou?&(rd?hb|skt+@qS#$LB* zbmQ%rK(7~%t<7#>;=8+BxZPs1TRh^%`!kX07PW4H=h*3PVRwsSe1CU~B)9k=jp~Z! znN6x@D0Z^j1>rXk_H&C1Y1C9KzuKgdhhisVU!@46zdX+^PN$K-SRUV`s)k}GW8dQ2 z&EW-ZaW0J-i{<4_Dt{<;GLC?OTJwCjIGaYzzfiqHu#<6Qo(OLu&U1^=Y1H?-LpSOJ z|F3M3Hq*|M=p5} zt119+_U2OKVKoqgoh+8O=Tg~WB=ND6#qw_SOb)BwIP7GI>*P|`5w!sTSAZ@_-;tJY zJhZLls5VA+%vRn1WcU)S{uJFR5c*mYw{ZmrP=_PnAoy0sP=_Pn>{iskP=_P;5Fb_4)e3bmoSFEfqQ2Iazj?@N zIfg<d<{Mm;upNmt_q)Cq(} zYG%5||EW=36x0cXCTskOM%7VJ2ZQ=F{#>J`D5w+g$B2NCHnRq^28c5LJc#1RhgJB{x(u zz~KbWEF}L8)d2ubdLemls44((CKOU5iVSEmfa58ovKuNN0G#YXYPg|#ziy5EDS_t{ zQrAtj0RUG}UPyg6TW)!%ujPn#kL-9H9Sg&E&FcR&$~x!N_a}(v(Zu)2jgPCg>xR5) zsn4U#!%l8b5VfPJ=N@_KXKMU)qv9(obUDg=%gJ>pCp@lpU9&>(IoUs&svlQ10H>YY zFq)bkSA*9q-zg_AA5A5Xt7L%pF?4J+`5#vu0N|_|OBj9272^zcJSRM^v~qL6qMk4_-H}d=eG@d4gEhSNL7MHi=5nM>wduAA>sF|FPXV z8l!I=E^L!%)`PO`L0R;s3LaF!H`UxGGqeHfu<Tn1J zhjEjr`XJQd5Xuu1CsETusKdVg7#=f;N)AFD_MtnGIf?uSp$_}NNk?TZHMf&Ncz*10Hszf&|P5Z@s;&Qxt1piV%2ZZi)@al2F0CQ#2E^3n`7ego7Agf836 zx1zW%f$C;Loj~Y46webd0S@W}LZ_p+A%U7^LY;sQMTO-FR5BCl1bpvDVI&;+XF?rJ z=!)W13FJk20RWu(1ZtcKbpjD^Rwhu{OsErxfU_onnjeEYws5Ocl%s;N0qWSoKX-~{ zU4=L6wLB_qkX7eZ_iJWk9lZ08Hi9jWW;MvR23d4L1<$MC1vU3|Gt_GaA89py0`EAi zPaNjkar%Wg&dQ^@^J>?l@J<|$%cJV^ss`Xl98b)nrt@m>QFteg$K+ATd6f)sIF2*( z$bVjS0DzO8N8au--VWMJ+!aoEp4LgSgN|wO86ow$~@R4pFwsU)os2xr{6Xm4`)c9Vb;>YmO?S=xH4s}a0FbzJkvmafBrK$$tw4ED< zW5k^r{3(27=jFqxWT{F9c;C+WCI3>@0RYab;pAPassO;LACBpFYT&2vk)2l#r?RCg z9{`*+C|SG)rSx!Xj3`G1qZdl)VN@!X?JNAIUdy1;GFjEBx~D@a_()hsUHaA`!j?g^ z%4A!aEP7T2J5}&mHTP~alnx5G9k&DM=P<1Dif{$E0 z5p9J|CR_=?v<67D~ZKm?51(eb2Uh>7gfE-qzw}$0*f(HsY?@}Umu^-*HklZChHwE`+_lvN1u|Iv(exAw2X85FN?9VKi9 zGz*dcMp=Z^|78{YOwIjwGc;fZUu-pAZdGfYmLF$OynZ2`vkIv0vf6b3K8ojY1yp@m z)c_#YKe2$CE~~)<@KHREDWHs5Hi z!Tu3cU9V~YPCK|^1T`Vo|0=xW;N>Hzq+TThyzgLC0sQr<0|1;=2>REnDgbcM1!$~S z1Fyn64qiEe%IZ}<061$#Q1c&+b*pg(VelBb3IoO!Kh7YG8hfSi+j=dVO1-jbkLq3m zb+~^O&Y)<0Ym%^K(=4xS^U9*VD!4}l_o}(oW~d74u<h?e# zcqf|2A;iB2>cBhEJTaS^5a&Mt??m&MY%19Ub>N+7&detN9;gHFM00vJdG|macqf`C zWK-iFr~~grvnQL%_COtYCz`X-S=a-0;GJlmlTBUdD*)gM%Co6&Z_8I6db;IZZJ6ws zpt`G}4#K%OgLHjAUU~j+Xuk^T1XKggAf4Os&NGI3hRI7gYJ3&c34|WT8KiSv4AmjW z|0=wrvpJPH@k{;*PzT=8c~uN~CqNx|N9TGZ z{U<;jct_`zF;q4I>cBfXuZf}NB~S-vkX@9ctWX7Ya0c1MvZo7g)N50z)Gw=!tL|sa zNG`mCGbmQynj&mdX_jBM`DM`w3^P~36Kd|in4u0cIJwoB2Jbj5KhB_7{X#5fO{Kcy zYS&bFCzi)erRw9V2H;36Pn=3k$JO9ecqf*}Or?_JDjDE#EF=H#Kdw3ez)7D<-s7qY z02~wm8jq`isqjuLd!|y^ag`4MPWDu4IIeo9!8@@$XDW4}qX2*_D4$AwCt9w1XlKjs zwQ;gzf$IJj!%r1QWO!4iaX&Q zoI8T+hEN?M{CV(>VE+)RUVwRn@Q&byA=I=$4d%f+f|n1Wk_DKg2=55SFZma!4ghdg z4I%FWRRsV}{Sazgpa$~b9lfIuI`9r23=w>GbtjZl>_@)K= z8#L>*roU0QEyv7@CA~TBKcgTn0)E`oeirhe`Ml=7Y@ju(d z2M&2A)u>pjncHpRghRfRiea6a`CFT~;E-Pdti?1gn>gZ-n^TS41}*%YO`LXMI;4?{ zm4CB|V-C3|)#%2`PMbLAke30lvdt!rI%H?6u?{OcY~rj#-k1`IysNP@in?OehRK1* z35^q?s4w<*ktXcnpy3{VWQIQb{nVPvp<6$;Yj}O#B{@A4k-NvbmfX~pvL5 zg)DWc4!4nu1yc+EO5-G(ScA{gLO$$L18$=`BM<@Oml_YTiPiWzE#w0()$7JU6bydC t=uVAS+r)l+oW`Uyj7S@H>)Ofhohwv&TKnYQHvBOsZC-I%(Y{$e{%_&s9NGW? literal 0 HcmV?d00001 diff --git a/tests/regression/om/whitebox_d_32.png b/tests/regression/om/whitebox_d_32.png deleted file mode 100644 index bc469bdfdb3d747a4cec7816796db68f9eaab493..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 110 zcmeAS@N?(olHy`uVBq!ia0vp^3LwnE1|*BCs=fdzQ%@Jikcif|7Z&m|7ziBMaJ&9x zNsqn(SI?8s+}ok$mg{%T6;^O)U|?io;Shi@>^W>gczK@G=Ra=)z4*}Q$iB} D3Va?_ diff --git a/tests/regression/om/whitebox_db_128.png b/tests/regression/om/whitebox_db_128.png new file mode 100644 index 0000000000000000000000000000000000000000..08160be33b7b72fbdc3b3876b85902dbc4eef455 GIT binary patch literal 37363 zcmZtNXH*jl)GvIHUP4fsbU6x2S9(X}C@4Y@q=zO_L&7Kmp?7i=6%_&Lp!8lPKqyJ9 zRHc^y85EE{fIc)B>2tpDw0Q55C5oq-Z&6#bui$UjtTsezIdE6@`hp zZ;H?2!<+n>_nxMwP=9}YbW%^tQ9BN2Qh5#f;j_aoAz3X<9DdV$78})sVOD-JmO@SY znskz&F{Kr!c`e%n8uRI8pA|-n6T{!2&(hqQBAF$h1V5&xejPX&p)sW%r+ls51RC`j zWbYJS7bk|ii9AblZ;EC%e8TgXn(>u*@<>BREpFlm#D40{*|Y4|OiZ3hh>IF|J&0g# z{rd1E|F!cbmLkUnaXmj$PL5oQep+zp0L5JWb?&TB%5T-UjUQcSf28~lU?zY4Bf9@l zmVDh<3#}E(?#*4w4OW38QaP$Q8VTD;!#IEleBf0Q84x2$lO2rl{#xPecHB$cU##JeEC#JA)gn0f(e0sz{B5!=zpu)D67nl@O^B zBb_DN7-RkMVeAN=ac*!DoRT7FAvF^+aH}|I1t9afNMgiD+2np>Y&hPNeUqn!2RsJf zN=0Z<3=&dsx=GM7u*7R1xgkQzAkP@n!tmbgb3A)I;33#4b)-cxOvuNvB|)nYl~)4> zhe;Xe=_P!JKJi+S-x#ywJ|;aAmyo;?g}uZ69{eoz64wPU)H(8C-vRTDeR%$eTv0m; zXU_$rlfH=l4q=Z6|G?bUj-riO`PhIJw;=`#$(s`0nennrTHzOi7?yZSVALbxFGi%+ zt4lG-Y^rR)vfCJgl4PGmcUrtQ6F%%>2*WtfB8+-KJYuA1* z7&?X5gu6rHJD6zhbpZ^9uw6JcgBZ(rq@kr1GeP~ra4Me6#LjBM`bBfyG%|#ZjQvZcGbXZcgOd@o!?0Ygvq%HT+3nhA%uxGlbqcQ zSSRC-(UCNKR|hS5VsdU2S9{og2%)19lAW8I>*SQ<$VggiR~Ni%7yirL z`j4CB{IPMQ=`MUvYSeHDR~Ov`pZr~K@cr=Pq2*fs598M82Keaj&B(I7KR=8c|9DHj z{_Pt1d|bI z1acyKHahIYG3B3uY*?3&Id>G^t;0c%5tX?4wcZ}*|J6lTFly}B< zscExq5`;0Qxp>Vm(J;t1zym(JjNTxIcPQ%93R;@kBzNBKP*Z1jB_hVAzs4(v>4rl# z0e0}t<@F6>NQb^Yt+1tq&2VRThng|_TjFR;_e;F}rW^cJhp|3;7%u$ThWN5kNk|i0 z>ukLQf7r)x%az<8M#UiwY|XQq`aZvY|1hfD^M+4#bcKCcKc0z?+BDQZz59E5JH@(< zuMKbvxg#|2!(xTxF!R}2t5mm`YbmHS7>wI!QbLa8&L*0*skQ+QAznf?KL%DvHZzmX z_!8YxuJNFyU@&QephOwTolA6T<7xw(LfnNmekiSwoM+xU<4fC%xb}iFg2C901|_`_ z=h;M)Hu*NlKEzWfZ>?s9)HpNhjID1k;ZhD64+cj!+LQ>h&U1;H?V@c^aEOP{>^f$H zbU35vOsi{e=8_C)2?B>Vc9kM#XTByXx9he;K_PZRo$DJLq`et^XIg!G3zuQYUJy9G z@mtA%MdHU5LUHSKk@=A_#phr4(r#^TBZ7}u%+hGCNOSyL<2dJHYaw{!>@|w`PcfB8`o{n*!qq(m~zW_6F9m%r9qf=8BfxjUY&-5!;M8} z(Ed+Rf3T75?;%&=we zEjYg0sp0Rp8DLC^EP#F9C2AB+;ih=s!ds-K4Th2;4t+!)t$Uy^w4}5YP8obn3Ol?j znzcSR&D;FP$GG&^TQqFFYwDuUpP8gvhlc3a_fF=BJV~Mc^0~t7eU0y|mTne5vGwbB z_JNd=e*t$j^|<6Wz9*3)<|))IpKHu^*Hq8qb+QDA&0p8pYf={f1w7O=~Nah=8MI`R`czRt66q@@1~xizW9;of+j)Duxmq5kx_%8W3PKZ`YR6d;;? zU14{5IQB2(*rXAs`8wz%L7+N?n(JfCym?jhEbYvc7}4ZwI=jKcx7Q$tCdD}A*Xt)E z1lIqfX8U+E&s}Z*hbA>8PDFgoV~={cdku1GQjW84GK~B8!I`)-I;V|vboh;Abi|C5 zbWRyb>4+G~>IfQ1>qrYOw>t0QbAf1n>1{L6%S=EorW!B^Ll+;6(h{?RH{uVE=| z{?I~wJeDAsGnzFK7wDMk2)Ll`b1C_l`x38Kmn-6n)T>!CnoC-!kH#JgB1f|(<4}%` zj({^Nmu6%h zSf}7Sr$eV_{^CAH6+v|@?yHtu4b^1DPc2WI3MO=>;_k27b5&G-Q!Hz-XeyZ0nT#vj z3Fp$OHgx=gI-ZO(TCEj)>=@nR^At@g@4siR)uG+V;jM9A4P0A%CzCUsGm@~Koth1} z2j1XQ!dRe5#r?kKc!G8Zhnb z5uzca25u~pWGK_fk%YtS;jqdzw-{5{asnV@7?h5XR<&& zMs*86`3cjiyi~}zq$Z<5C44AiMSgijiAff*76Gn%_%SJ=Zlbyco-D)g6_?&Hj;ncT zP}@FKu~NRgnu*C42^0a$JTOccR3EBa{K-BHUs+$kxThwlLGAqTnpN+MJtYPwVkQEa zdiXL)pcGKusZSPQ*thks8I?6`HK+?8Dp(0q_A@c#q9_r_%LC0ci>X3&w?El|VY})J z86`C|G^pP`EV4#SO$@~hiaLoxo*qt2of!Xs?v^LwFj{AQBcq{4usU_=!v^co)WmR% z`I8RVf111MXC^1+Vs3xC%tZP?WW6hj6L}UOhPhmSuZVHs!&lbZq8~*p1H=L@H{N^A zcj9qF+!)Ut^5E4YEEKiB>O5DJPfj)CXyG7}W51k9`~m*A5bNa-8z*aoek@ zD}a0CeQ~84L^-*5r#uKxFsSCS8!H_LR}LQQAf56ZaI340D}YC2usEh>u$)}I(;9^D zFsS7T8Cx0$HxA--C{C{SxWm=t704~}p14Gfayj|UPH_;H&{@TU7|$37_YWTF5N^9V z;Otf^S0Km8Ao1C<+H$hfPIC~pw6lh1b98(hj6X=yiMTyyk3+0vuR;!yX5yV?0o7!? z9dr<_q_cx(ZZu>B#2pOi9NnI?!MU!quR>0d2=O~R9YLled!tfio3+MmQ3X7i2SlCk zE4UTQGRYd_@~9%7w1clYBP$66D0x5VKH3xJQe9m^Tx}jz$2Pt#VGS4`8>7xlwTTd1 zom}JLC{aJ?9@+zDP>rr3E;TEv(;D73v$~GI9i>iB?TSRWPTz``MCqcSK(rk!s(QVK zSTdC-@(vY_dRDbwX;ksHi`9MoGOVFFUj5TJ&uD?YE;9c9x+$!p`Hgzn_|wq>2VIx= zvf$UOxl^A-e))BcyICH-VSO`|uU?^i8ys1AC70tO?v$un`Kq~p;<8*E)$87JIM$PWCv@NDmVjf+u4EDD_L>`YiFm#y^7Te$S5dMz@P$CMmj(G#D-SY*2?L6yX7`m zxiTgfVLfA;xLGk>2^j~O2t-wERFRZNoo#5{d@(7Qol2yg8wH+sl?((o@ArvC4Fgot6uf157+IEE^SlmEHPj`vUmdhJq9D+k&># z!k*WhdR7j|*pG5%a?o8*UzWUTh05;av;_fdYs2dkwYIjl)SaFRPJ)Ajb8L$|N)Ez! zqFH8Zt17#3X*&Wmd_&=hd|M4$>UPf}XN1F~Rcx)ilRSj>bYkhO^{?$tNfQ^KwKg=K z=&}vAakKUeVVO_U7cjNmwLN1qX%}19&~;+6CzkWOJe#~_l~~2)*8BJqqdm_#N90Z9 zzf_4;UcRmE5^J&WFH2Zjs=(>nQMKJ=TPBHHfu35;M^}U`lv8BMnG5xvTJ3r@+)L>n zGr*&*1Qn!xo_$hsZ)y+V8+u<@Nf{waPFpDV#JA~HbC0K&W`GAYiiowZL= z?#b?jY(mY1J7ohD$sr49Pg)_NgWE7I!Hf8-t)Z|75=#F zrdZz?OrtgpaY<{vQ7d8gPWzBf%^AX<%CQr1jpMq(cLTWPdxt|w)?AM0S^QZV|5~k> z*)FXlotib&bP8#gm30XxsqDS$2X>3DN5%XYrdOG#}P{Sz#($OW(&P@IW6z1#6F{@B5N1QGm zmBzL-l`wB+jHOe_Lv5!bkdCf#2s6i8x-YC{26Y)XY4W@ivY6acK9{ z5pwEQoU56xSzUQkCG*drjZ^2$!f!m2TURu)pR$Yd_l@S*DrS_DO%5enw7Q-l+5efw^Po;VJGL=glt49lN}y-8^ZbXkeExoUF_b_RYWgC&yXL zDIqZV+an{ex6DQ2X!Q|v z{a)lr19{8^lF<7nS7{aaR`xmm7JjgOe(awkONwnm{-fz+$n2iU$td{^g)3F~O7@5I z8~-rB*S~5iv?nNKO|eTr-itoT`JU@-LS-E7JT`|Hw-QOUV_xOq^c8?Td$mB+BpL` z0fak-aYMdOzMCm=p9Nc1SHO_MFDOW@C%rzQ=U|VF@#Zq)gzmWeGP)=!$ahyJF0f!b z>RvMh2-*r#*GUy82u}9SF?(DnPRQLI&1j%hCEs0|xWhst)D<$g3Tg;aS4oQ}BAh0y zVm7&)xF8RAC&nlxf93AlL~)jh+Y=5kSa&L_U_&0v3Xw4AXGd+k@Z1Uq^wV zTB?7pNDK_5d?vT<*N18u8Y;ns7C$b5^G6AK$gKQVxZds5Ex4V|o$FG9=P!&efyqZndJ$=JuW$id+1rqFl$k`-r+|JkVjmq!lkU=ocP)f`mr9?@ z#Qhs(ClPnFp;xpWzV+<$X1}pgmohx@=(pa>Ew(Mo&ysz{iqRL~PmVynrY)1LF9VYO z#+uOz@Z+POdfeORx9G3b8|7j->ESYL-rw_ji4Lto_W=vcO`d_?C;dd0)_dCc3z`OO z8Q*2TQyVDn#E{RkU&dG7PYndjF*kT>dd>TZuvTAfysV}^Tg!K^@6@-HJ~7Jk?6mmy z`;CDB0%O6m(c9loWNHo7#$M1gWJ~!j_?=o$A&Kc_I=zhd4onV&?qII-xC|)v6DwOi zwXymwdTarUw%@7ilpZlcwo_XCUSMS)OIfjoe)*iR7CXVFXQpImmptLvWSYVr2=0teQr6ezj5xryP8fek=(Fh)?KZ=jb%p4@l9 zsyG7EMRjOxV< zUQuX2%k^+=K!JJE*}>F1#9JzLCN&Wm93aR!8R>ArbI-f*h5oZd$6Yr;kx63*A5XY< zVXpp*M9i*_;K8Ja!)dP+Z~7eS;}h0W+`i2^@JTN2?z=4WBNlBzV5=G_u6YDr%cH|0fUUhK4rstRD@ z<;A+u)7R6jpE)UvS5`gG1i!oJMJ?Sy;>sRRMPR@KUjMZ*D z$5iW;?nT|{Yv3kiI^@N+sLHEA7%xxO*}5v~= zt17Ajp}jmGV0g=yN2>P%f;Z#~_ty~0`NGsT5Xy@I{Yq)n#9UToB5n(WKt z{;nFXYT38eds(~DkZH6pmU~2%P4!EkSkL8`+BvZni>9n$nQw$o-+k)o_EI}f>L`T0 zGzpOagwbxZdnXMh0KqgQ`%;o}0uZGolfXr;mqLm?mZMF$*qm31)L z>?t^|D_9-mY0@_Mn%9L>2nXH3pGz1g6-yCLDsh3;Kq&iWg8UV%GE$ht23GJIaN6OZ z>v)-jK~l3+#K{^iuo8%5pGy$ELQ|rgCc(i)-YCwN1jr1po-jv3OT8nlOPwQ$OEH53 zyhC6OuL`)qdmqf>@fF+*68Vau|nL#o37BFt7~!)(xk%(SfS zi@l()2}>zjDhBI+9jWPMx;#tr9_AW`?u6YEaj{UgAXhF+`(pJw)L;Q68O7lGuPrq~ zw#)ORyvdtG*YKsI*sm&uV_zBM)v*KE;daWP}zfjIuvuaZ!TX~u+?LxC0OMa1UDYbDV zi$CJIV|HAv)v6U_?ib0~X@)i>!pon?(qx;pS-R+XnA$Qj#{bT0+Ul7p8e#No>UrD) zKW|Qza%b63#lfWo&!$r2to#Bv<;o*v%ZhhP3!dvdi?dyEzdEg$@px%Qg0KRNy6)z{u*urX zgf+B3*^MLiFc|RIvOvCdIx&Os>WmS>RWRzRn=`{4>nan?V1l%Wb+ylO?8$h`fWN<2ozr`jlbe7`4SWIl+{c`fKn$3XhClCL#~QSr z;ZT+qk^$4czdlDuN9H7GzA$?M+1~d(HG5sjj1>NP;X3V1>pvWl(snXn>i3Fs5$VX> z1mzc~7m&?;^r_D4HD;uc&pX#?Qms-PhSHicV8-{wb4Tfv7YX+FQKwFQMqg)tF`fF% zV#D-`OsgtK>-U;-{4d&5EpKdGzarhL$JK?M@Jg31F>0QZk*1cs`_4>x@#;m)^eJmTz~qkOiH#dK5#9O^ zg*oubmsAjJrtuRMq+NZ|U@}c-3^d09D zoro$#cgjN;2Ti;F5@VN)#(6jEJ7y>5A69di%Dk06lQxkTbFuy`M;WG}U|fi@%X#N>96Zac18<=7G=%f3cd!Daqu~pFi*6 zH4LP9#y_6_0DhnqX;H-T#c)pFtmlF72tTon$B1O|%+KX}cuj+|JbNF#KY&D9m6o1Z zzBtb0Ti`r|jlhWIJsnIYZ~xrChn4L-!&CcC=mXeHd#yz{=_-b6nKzqdq+zxU~V$cIASIRTtiM2MK&PyKsk@Ap0wo}3fH z*+xW*P5orL_u~WnL!qQ0KhFIKpcq;_JVtO^Bp(NG5Bzs5&+7gDw|IQu9mG2-VZ!<)8Zly>1Lh)vIl;2SsCSILJo)i_0ue|f^4q)yd!CjMV{vQ!m|L_uCRzsGx<(f4Ezrzt11MdI7<4Br6|{$bo3$^HKknd&nyVJ~RNv!=ZB`be#x$T_VioGlvf?Xi6a zy5k?k?ebLdKVoHl;w7y9+Y77#9|S*A*C$?_CWtzT#P50Z-htfxL%9u}R{uvVtxvl| z(|aq&>iWUvRPDJ9Ye~$Lbx(!#|Qc>M8m^Vr~7COB4LlV)0miGwv7l)t5}) zzkPq^F*?Oay0m+e$Vr9;@2N^oVDihpo~`J@E=T^s^&3@FQ;R9E}Kd*NBGBWYw9<1C{cd($>6cW zf!Wj{b7lYeKcyTS#`y~TNsO*1|F@+6=c2 zdXp}C;Un1LJa*h52=++n=^)h;KH@5I(08DpH&3EQoYY7TGR8*Xqu7;sG`YbA@PiaV zhoY7+j?2dJ3S)$?_~V1vAJTD6QPMQ-nefKxE5Z0c_IPKCL(K(l1sl+t0TC)KI zZc_{cqBkYGS>lD6@IDs<88W!N=nqNa8YA-3E7h1v)<8C3$8DCOM!Zk5dncZS34gB+ zV9%jqh~ZnJn86V|it9z!im1Lk(OE|lRbF`e-pn-kkJF@Mp^VRexVM7NI& z_v3S!E^r6K3e+a_V|v^M=`^N;$%lIvR-is%7&AcaW*`xXjKA1$tQW<^B)YTX&oW`d zfd>qmJWpZNW4fz}(9^mUGst$H4O*a%Fmy`fNOZTv!&#wY>Kn!*E!tfML%4xxcWwM1 zZt5N}nhjc_(%T;O59_bITa`TcQbB)-Y85}vW= zy$AOE&M|nrGPq9HHKBiXO&49)2>qh#8q!}~GlKhi{5}%j(NzQw*~!>-zZvZyDY2%! zNhTi)L}Ht|8sI5=c01tb-zx@sGp@rp??cof6cOz#>9SV4N$x#{MPi$~>fyC}nmbZM za}=CC9oL*Xu8uU_Y1xq)o}0z_N1I4)9p^?~*bUh!q|D9W{OGb~<+vl#XLoO>aAs}} zhxv0^GVC}$^7JqG4*29ZXh8T2YeEi84k4F6_L2r`lAGkxW4=h*YdWoQ-9hYt*}seJa55|m-5srbY2-Z02AAQUdKT>XRS(UGQ)ece*QhS;&)qh`)#N$8C` zO~(s`t%gBX0g-T*W%Lgss^f|N>1FhaQBzAR+t1l_iF4sjbQ*(RHTp6=8;=Yyfpd1` z>R;Fi-YFQGo{4u1K)^*hI`n<$tfpXidM>^$q=4;smMFm;jthglHUao7oz)QAJKpGH zyF#kiB=?@~QMYFwN<>WRPR7@Umxe*;%?I$#&B!nDlHrm&)TP<;`n2Xf2YA%xaVc9g z-Pgo-h5wm8o5N4a1i$kR@RTM17%laLZL+Tv^0WmUoCmJ4XC$c_;IKp^`=z>K{dK(S{6 zPVAFFIQu?e!u}rcW?um8*hc~S00I!~{XhWwI^fDa14OfTK$qF$AvgAb1l$#@yq*|I zmmb|_xl)oKcdSAJg8Yph5%bu=O6V>{-M}kp;BuOy( z3Qhhf;D@C5gY!XP)y8+FoZ;M|#O*fgcEBjaQAlaUVv{62BkYW?(rx6jd*B@eR&1;( zA!l-j6H)Df?LcIRix6gQV3X88!{Ur@Zm;8N2@(ndpKi=6QHGs|5)a$W+9Asj2O)`N zrA?B@jDR!tb$bIBBFHWX%-mQ}(i?XkP84cGwL?}R&O$EBHJc>V49xlTGG^^cQ+qkr z&y7W;bL~!T&o(gYSH8@QB_cz@g*a!#oi7B11Q!m?j3zpUgb0byv5`;EUU1>?%y?p5 zdl%R7MywKdJ6k*CwPz_byCJbbYM(jljO_|D6Fu zKQ^Vco2%XLbC<>OO{#C0GD5t9sXYT#JdaoUHoDbwwZ-u@sx3^K4;F){B?Cu1lp%*N zv0rl&|AM13rjQ`P5VxD&*V!c26xO@Rxt`+KrbYr&N{H

    gRzT9=#cd(O7TYN?jmE0L|pWxjG1v7MOgzi~+%1_q zr~}z88H)eWdZS+t^KL%h4Cd^9*T@;q8%pAuvYrNRg*%EW`62?yp6kMB{A>M2m|e4X zGq|$5rh%N!8%|oC4x9!&!d*l$n85&Y^*Re0U)@m$3uzK+2J?33H7Ii~LrJ(PvuVgQ z+(A?Vr5r%!TNglMD>@orgH3kLVC(LR2H}g#a8kw;Y8tu`?kwtpstq8U(yh_!^(W~6 zS~Q!bCTA&055mnv<95e1zE4e0S^5M77?*dvfu-#ZXpBs?Pkr&(#28o7tx@jotVY6m zB>G3oUW@yL!52DQVxz;QLrXeblEl#A(pfYOAK3)UZ%Jtedv||oM9k2&QNZ;1X(;mW zwrG@pK88HF4xfThhi^2FW`7sKqTd>#$aCv|ntaWDh?&(76}8&c%UPD*%Cu0A#$E~{ zC$gsElAVSf0bf)Ymr_8#4>7GeQ}Og@*63ec<3z_b7JqRKnvQEi|Kgf1mN!mZjzBod zoomBS(U16QY+g`r!f_}L=Vayt`Jf!QBmxuzh^f^AirBX12A09Mw$0T2u@ynWgyV2r zh7-yO@}lSo&qn@nzc092sh*>R`j-FB#<5oUaL$S*qDUP@pYe4SE1W z+UeIe$JcAOaA;~Sse@ySM>3Q#6VaipkcsDlp| zw`BB2kzW(Mv*fcOWS|Y7i%$)PWY({4jxBC|%TcR7uI^q$S2bT3lVsjyi)KIb$M{`o zYi;E)3N+)p-(PKhL9Io-aD+0Oa3j!!&#yn%+(#`$y>OJSYAgZ~e5A!$nP1t5Sr`2` z0!Z2YOy<~jx~kdKfUASYi=a#dMQ^4zxP>%>R|kXrir%xRr;SYxPxPzRS5 zmt~HojtV%sG!6Yo`TdA2C}g`)=C7-1yi~~Otx0z^6uPU)nWMWJE>XIxx%b!Apy{q= z;jgP{pu3tkBC4W*xrZ%N4XO{*4SS*u!|dIBRt(^ic-`JsnZZ_a)+Mo6?$6uR$W&g6owLU*q| z34;BnHl}%IY~oAIzaAb;R!_`fpVTxo&rH*0O@s%E>CqE3?1JW7&4Q_kuQ6dB-b`6f zoMAqi!I}lrbXMcy5x~UyBoc;wQ{TlHpw6R4&852IS~FkHR}YgS09Q3VD`B z&X0kFgCZTIWj+#@ykflqT#xh^Q^ApD!o0-n7Qp)v6H!D}78+pb8Q$%@$uWFP4(o>!(^Nw(fu2*Q?h zzUCPmwHpHy4=Voh8fQAMLH*@5Xz^<+t1GNN!?B=X+ zS0Yj3H|VUUY@}tp(01+??pma`xZ{p<(2tRjajB}!a^tSf7CMaS)FC*`x#C(@*;XLT zv77j8tz-rH;|@89Mxe7A*Rhl_aP^>GC&J0l1-H3+eg*P4zAGM8lV3ru-5Cy|b^JCG z|6LV*8`rg(KT`grl9<{2Q(dcFqmXrJOlFLFG*u*mw9iJyC!?(W0AI8pOsUqql9<-K ztbW=i`_^AugQD9S+rPG^hHh)L|Js@s);FkaKOh_(1>2}lEF->}B8lkPJKc`Qp^{OM z5BfSxqE@k*nA+^Aj%|BOU>zK@9ii?|^@tGcoowPWP?adi4{Z)}sjaRyDtcSPT0eCl za?VfG?^y-9%BbXRIqT%qqR4lYlb>Z3y4tAqZ8hsC-PMesbp5_mtydYfzpZ4QncAS! zm~g*~W$V>M(dI~XY$cu5Y>x1ZP&=lGA`$kwx8o5g916Nhhc%rQ;?+dWW)F2*O)!Dg zaKvJST1uxiM^?Hv@pFEqD9B_zUgXH;uxwG%`}Rj0uxf=Smt&o4nYdlSR|yyeSqUgr zSk#cDM`dmBRXWX_cDGAyz>1Y$a!BXgTZyPjRr&|{AX@=U)j$oYe^lEB->lQZ*a#cc>LAL}X%9UzJ9;4DW*w<}MoCuo?8!&TaOHR+u`BtJ(g?uGs z6=WmeQeIO-LXE21oG!1azS7iI!})XNK<-?nXvMP{Ow|?3ncIoTpg;l6QAL{zHZ8V= zR&-V47!)obGAeH4V-sRqXgy<-SVuQC$8=M}T^U{pc?G!&%+_qwlG;a^Y_MHnb)1rR za2s%Y1(b_$(6>sgtu(2G&^z`5owX9Rq?A!U8(M2v180{l!Uo)1S(ZC;I4AN>)ua$D1fh1DL)}(Cu2t)>UqhDbjfs#m6souyI7&pabo)|@>XoHr#DMmnzO(Un_xRP z%Sr24O}TJ+$kx-9Wwv~+wmUqHNq|;&zmCpxcx+Nv z+l@%$6QC8}Z#dCqV_`#0=~?DHve_?~m-K%5F&(VhqA!Z9C0vuXd-!?)qtH9TN^%G} zvh>1dPkfbLF}K~z(sZz5t4Rgvoad5+>Q(IlB164|F$#lnBcm2ZJx|k3jd2sXm#ct8W>n zYUW`E1p6X61{9wt6ZeZbr15oX^~^#UGHKMIq3@@VE?HJ_@@Cd%7c11Nna5w2WKi3N z9-pFEW;w)V-)Ou6nE3{CT+3K|N!=ZyoTAud*~MKd-j~MKH5D+!GX&GAUx!|w(sOY{ z#+^4ay8%7$_2tM@RFo%f7cWR->ziIPD`(iIQ%8p?P7z!jo#RY>%{lHAH%pto97{W6 zLpL=yeT_L9il0b-%FswFu%AN4-S;))s3@+ME~BfO0*5J=xL{usj<(`l=^tsqX>OKN z&T*PI!f!yfz5yJwvg`81@Zw`>THS*J=6pJiZxvrE{1^eq`Mx_9NIlAMJC zN&5ZutN1FdR(3o7Qhu;ve&nyKvHk06lIgDI+Fw_rKzB6`e_c&GdyBvlKlpT>@(;y2 z$1dS8c{>@hyk~S$LRLwE3yO(nYpL zfrNT*bM>@rjnb7SdNr^TdB5 z9O$T~HW`-;p?{j5>{OOeB&EFfxJqlq64<*05&YoZ{O5m;9Q14wx{^zi`8VXdGZWc3 zf#AQWM)9$H_x@j0Q%6TNLPB&@GxQhLxX@9JJeM`+#fpo~4CDMI0@OCruM-r@SGQua zIaRp;Gk06YYeHT^)Ll}}35v}trx_Ad3X0v^iJmOj`Z@xGvXCu5b(GX|g5Y9r6JzQg#CRuBoyC-YOz@12yS8_Aq~L-`-$5z_uT^-D-so1%7oqu7C21U#ss^EGq%x9%@2WKK{v|@n;1>bb}q=) z9lWC>+>%5@`d}-n10vLBh(?h<@e}*&L;@busqO2rp zJ|PCkHv6)ncpXD^`1oS!0(fxrOONs@|2ZyuTXh>iMcGPVdItx{;``d6`16LE@V%v_ z1+ehwNKfyD>oeTPEv{|o->6#>c|FPlWYvA?P^?my8eI8%#sXM;w53OQ;rbjmz9qkX z@r7#|&WA2)=J(Y@O&7-(rE*;}aA&#<;H`8|!@niEZP^phV=UdJ4R1b5*Hhba+WyiL z&})32E^6GPOe7BWb3;!rhAe>nbWuadolC`qY=>_{QBkfE1_PUeN`n|7$0;w~E>QkH6o&ES`KQ%J`@aR-r8Cc+QsZx9&^#)1QS)hWZ0St? z^t1Ro7!RJFR&(vLMUAC`yy+M5UUXA4)r!{svG{hWAcIb7jOe6h_S1R~(W!Mt8+*Rx zJlp2N?jn^;M>P>y)A{kr_u1}4HXCj{ou9>fh#{?HZJJcbIX1(Er;F4KO1;=ox^7;) z;r;XXA-j#c1BL^iX0sC;8}>ZGoHZr~EoEH|E}^?0J;B#&GzH{~+-Hfs4SbEj)nF^Z znjVV*(&(hD171T{oon2?)D!%$_Dg}1llwf;sZq5Na0<2+-1sywK+>PocEB6zYI4Cn zmpsAeYex!t>CVp*&o^*2Lbro&3FduP8XyHuN;_bs+SRy}Ju^JPw6!e-Lb~(w#OG^0 z3ja1%Ht-JC3|vuZm*HAoo1*KOAcs$$Ytn%izW*(#3XFt^ltPPjG#0LxGS@7d3m0;f>I}U{^td!Hq#u#w3#iO(*;! z*PIvJ6Fgi46^>rfQB8iMNh5^XwHJ&Ulo%xCPx3j?bi-B5fDi_6);(4r7;t4+VudpU zQ5Q90xZ>|IBE=rVsY7m87&6%SnD8pE_`A#8?lX)tTc5xyy%On$5t|q=oF#04*~POl zU+D{N4BU$au%j=uFA;MXB)At@A>g8Nj7j`Hqn?BUJN2`hI|G6(`2>BtE!MqFe8J$t zoz4o~uTzMzi(g6bSos6o~u73Pk<63WWUOZoF&k)FoP`_Z?@eKW#Z z@__-Y1D`1wvD+%}yVSV832x*|;fz=_m2awmgVz<-nol?Tx@|JA3*&WEwVB4fO72n* z`bM}ZuQJnPvsH&x0hHG@Rt)`j_iph_2VwkqRb8e%@5Q@xVMO^Wj56rLXj>Kf*UN}C zuScP;TQ&2lFjlEii%I!z`dw;q-w-$9g+qGmxJs|;#TO3GV}0n#Xg<@1E`@Bp&*VC! z#-3@^U}~isBYu@tRm&d#PnV?|FEKUum2j)6WUGGZ@&A1JJY5;NdpWTlWUdRJ_6ojB z?e81lCgjqUQHW~0Diq~~U^N(6>+fF5j1;D+(4CR1H_u(_bRUsBB47JutfVSV6$+$F zqo{td{_f384`G^WpdM4Sx5Zu4xy&Th;gC{r?N^7dmS&`@$N<9dSrXT<$1rd&%=x{K z<^Ng$q==Gv3WB__%;Aey(Z2+HGn6_+CpP#05}WIEVpIN?*eL!bHl^+l!mf#65Q8_# z&dUzIr&IHClh}p>haeyfBO+lwXh#0Fyzh%u={N&JlnRxAlfPc8>E*lRBnb_h4MF$A zd_`Q$mCeYj%L~5$71LMIJ5h_+BX0EGo>J zeU%gzW-XGn-0$mC+EP@QJ)4!}66Qn~GV8wBvmIw)0d!ph=KiWtBV^3xCc%fphalgu z0Fh49O*69H^06;XCh`nyt{7ed;(isX9nlv%Nv=aCLy*&fkH~LRNi(wR@;BcAWDW{!~f^)eh^#tB&^TSHcU*EwCTE1Lg*Y~6V{R1f_B{mQ=Y`<4i0%f5^y z2}$;?7<@`dwy}@wTa-|#Bzv1AB-_Y3Ga_Wm76vm(3NbC13f=GXz3=aJ|Ni*Rg z9eW*R{}rccgvRV@DAC|=clpdami9Xe)=GbejuKoHjuZNhw%wJ!E1gP z=*R=0Bj|jfBVxCKjyNK8#1iPpmruWBNhL+M(pqn8&^Rln|d)Ufw#m)S zEc@-*+f?(gSl;>5>!%2)zG6$PYKtO^lw7ntW$aUlY}_qUVOs6&zqhH@VTrsQR+6U( zTz#pQv?DPJES<7k@|5OJw`6y289Yeqyxn_yFVy_w0*kI|sSu7$GT4yhUKcReg8OJ3>jrbidvzmv4em&b?kzHs_Lu+16kuIg+4 zEI4H~!~M~+zOv+!?sxR0())LlGI;CxsHZO29FVknXWhkjErtbqRLyw*u&kggxuu&x z-*)rtouqu;EI#V#3zi2wt(2_0PJtYGcx{O#Q{F}PV8Pjfq=*a72U5U|eBuq@wLA5~ z`uLGLmHl_S|Im-~it&z|3OsXMvF^zJ;_lD%Y1folzQ{Mp{dfPvy(AsJSw3pi1&ago z)@x_F{$BGy7z%2Csr;5Kc>wH4T=9qeq!7NpeAM`l#}6c&5kJ#440?p76{@?xQ$9zQ zJls7>zf-Jx@8JKhM;1gtk3{!^9%%*En-~b%*A`U|L<_*2EdUq1deC}92xEfIPwO8 zBjNx@To(b3Tmd+8=(|s7=9~OpX&x`045v6Bz#q<`Z+l#OrKP<1GA@7`>Hi)X`mATF zw_I|8JHt258}NrU;18YNa-jv@43oG3@%hYk5t#4OC;m}6OTmq z03MkGJQ5WGcw`R2BU}g``GDXNAq0J8G-3D5 zRK9_?sxwo=9&XZQtbKE@`#;$xO~%m=uY|k32fT0IsLnLeZ&GJ833O%;tD8_$=$`B@ z;|DZyI?#jta^0rdrEcx+GJ#p)RGUCYcEX3x!oR%TFC4oOl#MT}3sS?PnhY5?dZN0? zy8uVx#AgLkhrL_9sXl?W>>WAZvhmGz@oF@qCSyk99?ov^@&~-|jwp&hb-_Eso9efA zihVCfBAaKoRG~;@*Oh-LXAi}?Y`$@W@O!lN5$61_c?Gn1yDDqGxl|@$X|&EBr>?KU z`sm@}BaCI3Wan;QgN=5q&+@=93jC&qqXD<`2%b$s%}+ezh0oE9OaOdZh7) ziivRHm&_i@AuNvS*<;o9by#gJmdz;BUv?A>q&&FOg&wgzqJ~A-^Lu8izZ1SQpH;%_tdP_7ryb4Y||39x^?my8p7{-^-NDA$X0-*u#IbAI3boia%l5=HHGSrxhuJj!@)^fV0ZW;2>v|s5QSg_bKBlo`D zW8M+h<4*RHVUnT7ezK$AyCI&_wQx&@$9!*rxk&Gau?%MGZKpYzbSp8p>$gmFPFK22 zbZ*xhneCjeCo13`|AR*z^EMTK2cfO>KTPR25>on~ zn$rKvmJ0s=*Ys|betY}>YN!AGjo)@1PW`U|$q!wJ{E#B#hpI(Qd>XFMH;z8w(KHXR z;Le%ME;Dp9PFcT1Wd8FgTShoXJDa>U`k9BMH=~}?YT<5y*|;>zye<_ag-?AXA%n#> zgwyBdL}in|kFN6==*?)PR9M7XP}eRkFnfxBlfqX$qLaaTM_}kfvN^NK#iJCSHXVKS z6fUSCQAfXgU^Wtwkiz>t;*+6WhzO(i$}SPZ>pT)Nr_OzO!n1SiiVxJ3JjL*;kN({t zH;*P-P=9WVm2ZK5n`pOZh&qO^D=t{7p0AILmGXoNG~nxYW-DA zA7+bNMQlif;7n{W#RcnzFzawGgG8sKNR#x@x@m90Y+WF}F+Bw_Bxx-mTkkN_+^>jB zrf<1i*~@j&dTBo}6Wq^<6Q<9&rr68%(5SRY%pA8qantk*hYtJsV0 z!Gox5it_`mwawAYK}n`;NgHUA)Y2@jqB5u-FjL$m!t0U*N1~xA))&OqFxNH*aIZb- zHY0;$X*wB98PqN62=}~%q$J@lQO1;p!A3AkapiIV)%q?SXJnw3rkWw2LEWUbNXMb# z&NDA?d5IEoiLO_vKdFteLf#@ZBnhY?eMM?W;ZQ^J<9fMt-`B@$~AU0JE|blAw3iTgs>IoZk8ptjmHbTyN<88euuc=FqQMxx>pUGk~kbhJQB z+&(XM8+LLPzHfh=1~z!#0*9w?7guUNBXx{?cmIw$HfUcb$DT%#G2=gpF0a&o613hU zyO!~~)y6Gce!nVyQC{qg8x@U}PGIuel80L~AKG2qAs>HMLnj<+YGul4vSrzMCD7d(>G8+)HPhQnGi?h%~h@ zlfm?7_7BxAF;+_Q>&6$NH-Ihd!_F!$+s){0O2J- zXTfY?n-iaLZESOGQ;2)FjN(YfGa0sjd;d^l{$Zs$AcsVp+RclnHT^?FOsR7uL+Y>U zCS`d~F`m}=&rK>}<6i|=;?}>}jLyIRW-D~+m@2m9Hp#!{KS>r%7rsxek9Ca21c$hB z4Q*I&;M3bh3~@M}dZzp>-%axM`87%OSm9Wz*p={DjAMu=SMb-t4SY~LgJI^D;MT!; z_wm$=A)#FF+YdvA$SU^cs`9Eyd2<`nllkkCPFECT?N+N-Odd7AW;!2Y4fUZ$!!KKt zn`H~oK{_8|$`!m3v<5LHYb2)B0x_kgUl3E8gqYIoB*c_vVnDg;ah{}&HCwYvfUxpIW%AB*u zr6||v@(Hj>Hqk*sSLPx;gd?V+&yUZIq@!MowPIYO9fj7vLJKJ#T1eP3Xd!V$Hb+u! zK?})X5?V;|uc3wH1TCbJuh2sJ2|Nkg0xhIkq=hsHEu^+N!;$o81Sll{P%>HrpwtUM ziPiu>sT2W90|-!>1E3Uw0Hyft7L}b3@#PRx@>(I(&bhv(#%?#N>`X%uX>Wb6*~;$( z6MTwX>-t)3DWO}h(L(1J{WOA18XD&;S}2%IWQCm@A%x_C+AS%gbO z+2pzDWOF}fCIW~IHbiVwGhIl;RwTK4XpggPOm8H4 zTb~IDemMCXq^Jue#6voZ$mQB#(~9<|9PQ9SVxL)^Id&OnlEgQTAY4^@ z9mjTqP9vpb>$A4O$UK6a+!5sD4#+74K~53~a;gX96x~Sqy|u30Mw%bHIyD<7S|SG(2lDOZEACR;(jfDumeljW3AS5&w$A zxT!P%aFP&$lb!%hN_YY|$?N~aN!cOz%4!y6-FB@y2EQh`M)KO^C&8lG2k%pUx;wgK z&Ih=$sRUbx;Ga~BDB~V#)iX>r_%@QSPOb@(#vhEOxVweBW84Eg*&bH^R?ch|Y&bX# z;h%^Ad$tttNuS)T-0UtX2bokt_UG>855aLaMYoYl)qy4-+$K_@Ao}yVT1)v$^JLS1 zAKa!=X5ClaF@6E2Z0o^6f%x`nYh_$#V+#XUQ&R((0y@cHW@a*_*L~KVdNII)tt9wc zAiloZLm6Aw*ve4bwADb)2B6e7F*B7C?e^E5>bZKH&FGRuAl{g^OkSyENNn)*$9q*zxxu5jx<+YZ z1QzKc9vjF+VjP&4)ACGpRf~8D43u$h71oM$jHRw&+5kaB8Z9uyPI--SVb-}?Z;I0^ z;v-a3NZf}&kzSzeaPv}9xE&}2cqs27_bbXr?jMxf+}A1Zx%(-vxSvygaIa9Vb3dT; z!-ILxx%DY4+*bgQGE!!^%_*Fs?Z1v{*v6G)X9!m68mAlq=o+1Xh+LO7NyM zahJdYh*T$>Oqw5|SlU2pot3f#hO{;#_5@agBu1Ep$N%Ozwlk;+*n??YM6l`NlB|?5 zDu%fA_zAwDK$}3q^q{4+ZnIts%d$>|CWY{MTnNpdF(e?H(2845Y^s-=wOE3}a42hFVCY6Zn8a{L zKFPq~U;=Sa*&5r>(#Vpl)2v15`@AJ*Fh>IVRfsC2UI=&P^9))KCWs9hT4SqPnprk9 zx3wtkpZDbAMhr&Mavv}iQZ0Yl@fih31`~t^Wvpq>A&0b}nX5%<{rp{SXT)GMt@MF> zA=T<p+Tscdkrq^YH*P!eF6p_$K5h(K@90ab_~ZU&CfaqC(e!}o$)?T_bS_r zdNmKbN)`>!bN;Qma z^B?Atbi`sssZoHo=7ZJ(y8dSGa?oCemhG3);nXIN^^`Nl&FhkqRK+YrsgaA9=sJSL z1G=`e7ddEkp%3>P>)2?Mr+UIUXzkxm(Cr1bjVE=A)oMP=t%f2}C{je~fFhFoI24h( zks^{4QbbyXBGPH3h*S+lB>zab^?781F7o zQWH{3tlyM~#MX7xbJk8qPf)V|tn0TqFF0mIuf(oWJ>$J3j39rpOOx?ueot3f)o%E{|QtI0s$fsHqx! zhR<#+(89H7S1~qE`c55mD0ECc8x+jJ+ZdyDYU1O^JVOyH&8EdDaCdh+;w%R;q z9aE#Bg+xsV^5QUh@vRo$+eoKHYlwQuSPC^H^1uv5G|mI%kQx$T5kQUqc9tXIg+wiW zxRFnb))@7Qv2&`V3SZhNbeS6c?TP4wnzL@BeiLhjE4Zz9h}-f2w-sImZYu$CTQ9(E zy+Pbo!~eRi2jI3ycHp)OpM%@7Lflr>P|b-3{R%)zp0ES8g^oNA$jS_m zm5^(>gNFkkD}O*%KF-Wl163n6HP22ozkUVC>Ny~*swzxuq_x2Mi<)Y}+(e`^uK9Hb z>$aawO&nNnuvF%J$v2BTI$5*M5$lFAC0p&bV{#%o>4`$h8$8ym1H%6Yr@-^ zR(Yk3wbRMgiPF5%cXY=-zJ@i?+3OiRMtN6FK(|`G>42Y8k2Cp2KaI7>vA{0*=TSaI zjERo}gUa(WHC>1A9udI3P;Wfo>ipChu+}eD(%Az0*Ew`5-sTW-)n`r_o`ROWi_ zT(XIeGlTiv2?5-rS5FT_J5RchOZs24l3dSuB{LlrJ4&?!u_aj<_`J*e?w|nn;j1SH zY@M52$OZjPtOj1wp2_)&sH4;?fVSpetUd3Vy*np>ZGKg8V8eCGg-q=4V2wjhdnMZ{ zG99JHEd9sO0X3wqsk=l0T2t8L0}HM>F67?+I@TRD^pLz1<&RQhmo75wy)bZ1?o?ED zQtNw!tkGT>PC^Dp!t4N$)`tGCKY*5a=oG*WmLiq7m9gQwQU$q)=?8lhJTo$Q0 zVyj5xSo-RNt2B7T-h9XJ9;NBW3-yx&^#k0{LNa}2(9Hf7*cScLZ=bh!h3=HwK?%tr z#$Mzopsi zZSg+QJA4XKNCzDOZ3X%#KzxV+;zNQEA992Ekbfw|hYmx0=-$0DXE&D=8~+%VD^COt z3*T|M_n(uSa|-Hol@-Pnz}EVaYGr)6eb^2~hAK{~R(v1KP6463e1uLmhca zl!8?VV{eqruuZr|W>vvv0m{V3twU}0;I^8tqq3+rp{_jhit8!_CTJnip5dCAOZm42 zDD@xr4#n9U*rnB8XUd{lf4Af5P?S_5a6t}x_lN~ zK~rW96uIC!_WAMm@RMHFj-%q(ZhaHcY^I^Zkd?3m7s5d<707DI}-!Hv84nTNvI7eik_Bb>A!L!g-VZ z{pslp&jbYp$~>Wi3Ix4T$d^1JB=~AhKcc#L+^6?x;h!L|7!st zJp)}N9q1x0FGClJfOL_Hpo?UMbdk)EE>fRuACUv;A{~J)Qu8P1A~iu5=_DX7?&0Bn zQ*~n$*VmmQ1A3?;yt{ z4668zOc~mFZa^ybD$a=g8Sn#>m=Gi>`L1?=Nc_DSgVpSm=E_~x`Aq5C zd1_!#yikb8M!}Ve^4zG zLL`j_HU@}7zjtG3M-vZnEqvnoOljRIH`uvpT$E7?X(X!E-X+O|0jU9^)bGU@nquOi z4XKJxHA8s+mVj`K`T8I6qPt>wso`PX;TX?gG^fhj>|wlF3yU63R!5B~e@SkM{G9Yj z9GzcyJ5?;)F&q;f?8X^9XgzEqt0Tiy2U#T9u==o@{ewd$ays%%(3i5a`{ZIAw3r-<+&H9 z{tXKVqaFvoCD~s+fR}06)Wa$@OEcMi(*8`YBs~>3$b%}InjA)H3T-v;pcExUR&{pf>C3!gbZlm9{)Y3Dww zD77@)B8(caeu*<-SYiNQ3T-4BD*BMcR=-lwgLZwk3cuPxLc#QEJtivZjNs5<&LH9U zbm%$UF_aOTke~cT%KNo$CG@T0JLv(h^6D|BQKtoi`)vjZ%G3VmaLOn{w$Zg4O5K3C zZoV7pKcSC8v6)2K2!>6!oKv{C>{EWHxHSEAl!@Tw>GX4#E(-Zbzf&1Hfohdyn_GLY z^s9d0)v0`Lz`#QmIE?I#Yc>Jrxt>At*UmGfY^V<~WB{ zZ&hQn_0#d8OhXl^?LHVT*LtWTnL`z6zCU=FfSQ&#hgAi|C3O)L*CS9|akoKn4Zi}# zWd@3?qaPHP4x+eJKyevg1jSW>D6ZR}xE5Z4;<5n6wbw5>3|B_@yZeg-(Gp!E_{RN3 z0%)l&J^cNBLu@^a7?WcvP$>Bhg_1MOP$;>u&y^t$ph<*uw%wamtP)Cy7KZ^lJ84}Bp2wB*u0cY z@_eb8WCD82{-r?@Dn(#FN#F%Lx#U0RecKFGbl8Ju`MON;Gy5c2CK;$_Um7LJr*P~y z5C~)^7yReAZ68qbv*4W8Jh>}P#wbIcvKk#)e zO^qIG*sw>{CfPz@86}dqi+@5PYzJeQ^%P4 zhA{^BpXuvTE~pg2DQhUxjlTQE-mRgbeiQDh9oN@Tq%-l2Vhk(L6H$OtRM{OMT&H~# z7%vwvh+MiT=zH*;QvY!j0Is=rx7mMbSZN%4r!>&@tzbz6r%`{1F51tqdaOFt%r}IwqknCnOQ&F0gr*XHkk0rbr!TqU zT{-*CZRkZUs9UI0Ek0jj-0K$~fcqL>2gWdNfBH3xZguE$_&uq$ChC1%9m9BDIU}la zD%YB5_BGTWCu?ZLk^e;gG3EKMZcTK4$(^fWV|ink@GGZ9gRO0?O=ZDt)j<_W_PxgY zo2Q4*n93O{a!mhP)BFT#%kIpuwW*$=D#rvQk;cc^-jAG7J!7iauE0_J>!aqc_g3$Z zol>VGG810u}bYK9JY_ADkzn| zu4)?Of!bOgD@4dwxTvRXh$XQQqDWX&yE4bd)8-1w?ynrpxcl=3S98Z&5v~<2YGk`% zNzDJ+?oZQiSLIlEx?Mr}3t^<4`$qS#mX2k-rw07Nh$h%dSrSVjjYLDmD_O5B9MRQ^ zjA6C0^(aWuaCEI8uZ=ez(bb8JXC1Q5w#27b)*Zp=YN;^fKkca?Ums87N9R4bpCaav z=YVm%9L5@KX>ExQs;oSc`Sj_NgC!42QZ8OLXMGP*BxZ+thnrS|HYU25n7b8V@p@zkD7tL}rj75`+wq+wGL|m;G6TZI!L?#g$7(a40}q+fQAdlJ~}+ z^BdfED^AI93~-=c0=6Y-TWy8UtK2+-)o+w$7=Eh#l>BD=DSz98nW7XO2X_am%jE#p zdCMRxd~D^nBea8!G7M5rqMnlT#%uWF?#&dZs5-Q z3FlLC{&+e6&OQC26k`Vq2deA;q>(K5P%IUPRc#3(eTsP)mme##MR{%7EX#*$Z3zs0 zYL++!J#ChJ+1|sHSD$FI=(|NX(^hZu;Q?-kukI~jI` z*P&0u@~y0&oOEH)-L$x{DBkORdX@_E%W~xpit^L00NQ%c$6$F0azEvH=sRhiVWGT? zeTOZD59iBKaz4G1CEZ1rq{ZG=yp8b+^Wl{|RdbqP+Bac|l?Akwdw5EYa_!S=S%bUi z(lo!@Vz;SrVcxu+Ho>O}t9^r(SPgJnA@a>~l(|n$vTdcLk~FS7)NN{Lm>;i^_4;YT za^IXKO}#~)#Y=u$jxr0?pSV&|X_{0X(`{<_kNsTpmPGknIm;cz z{S|V1#Ytm)`SRp<-FI1ZQK69shOE(rNk93fbjt-FN6u@;UNx+n%vHt`2hRYxfO$`fCL&Zx8z(F1vZZILQIr z*3oM%EYA*qK3sO|ekpW#4(|U3YAb>-i+99U$)@XAo$UT*`FnEYq3&1oq~fzBNwIv2 zyqLrbJ_jVvyg1!Oyf(pt)vS}-pDRBpNABx>O>ak6A-Vgs`{i0nCm@zFQ9nMW!o37^K&&84Js+p$p9mhD^ z!ECXA%vwBlUG*AK;`@_OoMML(=XQ5~FQtQaUDtryOC%%XW7S9MsklgS$*kIIM85Bh zq1dAxDxAaJQ@xbmw5wpg7KAgV77IU8edEH#=W{}?5nI1khhmL8lsTn(ntLgWv>e21 ziDV2fwtl4gK=Y>~XX6^N`FndP&8S0_)3|55m$FPt*4+^`7S32$%=k$4+rfzMLBt-6nSz#pu_gl8c|NQ}o6JudN6dNwsP-3iZ_F;@{M%so^xWwHaNzd)|;A ze4q)FxC^;c3or0_W2^!l*n_jIb4)a~4>NuTx+Uv9cHw4@b*>3_Ax|ng(3hPR;FiK0 z-`+CbLcUbDz+iUKx)?QP@1Ncykoz&wX*$NZ`{AK5yLXm1CN|KLef@@VHh!+oM-8Xb zWWc!H>(Wief9MxB5O)_yweYU;reXuF*-LIzXXD4}BGs_kO-78WJ=)#mi4Q%(Z9KF5 zsm&J(y{R^k{F%=JyTt@{i>3f}tF#;JRy|_3c))JeUI4pgjo7U$uv=V+-8uqxtFs&I zRx@I^c))ITUdZr;3eEw!RSj}0^Q63mv|uqWFssYJto(slRai)if!m6U z@C9ad4w#jnMY(7(A26$6U{*2qm&}$e%J~3oc}AQEW@TnCY$jwu$@y|yfh6D|njZUb z;338f(3WIoO(wx~^sYTt_VrQLTr*b-%C#>K6buAB#M1p9u0Ev3f!p%T3C;xXm2HpJ zczuL5#6rh{GWX@4LYvf_C|Ir8ht$xBSbn4I^-RL@Xt6y_{k0;imqoM%W%f&nLY&l` zSi01szYnS55sCZ>*^-%rjnPzln#L6c)~i)%Zb%tfNVOPPx#f4kBT-Pl1L@Uu|K z?P|HDD1%A(w8D%QZqYs1Kr}f5=vK6>#wBn=@V$q+l5P49)-72j&EkT|fHH>^)S9vH3qSyw{_DIzMu|G4El$ ztk3i<%(F@B6zQiCFf-hziC;~buWDSilNb^|5w9D-OaYDLM$|L?!oJKYEjlE2!l1}F z4G+s~=EQVUU2ohaWNUDkoN^oo4euKyt^L=(??-ASodT@0l^NVOf<#jfB$^5!(d5ffrgz^k zY3skCeXmnn=)Q1#<|s41Z=6K=4-eCBO5!l>W(L!4fRUCt@3F(Q+k2RHJF5-TZUe}) z+Z&j6dm#?fZq6|6_VNl$yG^pgwA&0a?WP0MZmY<&8xf}60>xq4%^I0@i-BpkCU!C& zKe^vP|ATgtJmV}AbsQp6B_eC$UE!%7blCGRV`18Diyfxj=8K!{a@b+oZ4{>6bhTmHt(S>9LGIqyAwtqg+F;~jCuiW{NPp_veF-84 z+DQwH|3tc^V4sa#i~q;Cck7Gon(1@nO0I&vPez~qIR3HwV*Aas64u#Fpds~qR|M9} z{I;amhF46hKqD1HLFyz-vt|DM`RCY|;SJ)m&N6UJLR;XN9)e>6jP&!5)gN$7{@|EY zSHUqYfn(}?$mROWb&K+R@3|3r3_Q~6AGkU;O@wyHO6@99d3Q4&SKV3875z(Ri&C@q z)QB_&0ja-#-TzQAi2+iQ8zHO2yxni{*w>wpxQ4f)w5 zEjVR+m+24H`(L@y4k__5OuNAW8@vC0QHJ*t|5Br2pY0=yn(X2=wa1H=Y8My#)IKd9 zQR82%RJ*h2ujcgetJ?TS5jD+^1!|g$=hPY&r_~}B_0)cSY*G8Q7^B9%_*<<=QZ{BGm(@B&U}UYwZv{hqL7c5@H)WSjEUijUva4EQ!O zDJ@t(cGd3I3)my`5=6!CCBmLr+dWkCHr1=mA9vN-_@pu^?X>^crMfdN&=icsiSplX z2^(eQ_E62@l&;o(TvUtWE6t>|(mI3_GU3GgULFDUr&FPjYC?iOsAS@V*Z!pI1JWtZ~1POO%zU~+Jr=LJ%p}S zcYBMW@UPFC|Bb_1oyHYdt8;<1x}mSIRu_Y;)s;0rW)j+jwYuA|R)?O1wK~NsidQhs zuvS;|71rt^VXY2V-u#qF2U)AT2WxdCWUUSbYjqe`SgTuytQLN^eb5k#YktC%yWO-! z?wfCtG?;_1mR~ID3NyS=qV9-QxxY@36&R&*@W6>0)Z z))}oytl)dwiH5X>=*LW@+c{h00T``|n?uc~hQ!KWp~i1?NbY=yEn_m?mR!Ntwp(1G z#%|O}?o9u+<=R`{ZMM)fO1Y>hof4s`mJ+CWBqdH$FXe*f;gm2K(K@ZEoZ_Zwm}0Ie zlj5tXkrJh;m=dgsO7YZ`Pf36Ut@E0yDPxa*^Aw3FXOh=0ou^+m&pM=Q>?V#+AHB`2 zl5@`-}llJZ(Y}ndHWkz3MsG29WagMtog%p~W3DQ8Z|lDtOF{ZhxgcYTlVx5YzL2V3QOn##r3e62TEZB|Z{qgvRWtlv?g z2f6kvsv@#{9%lSrq8_nTufJ1?0Jp{+=@%`%4s5Lq)y(z?HvEc4VZ*N;+3>Rgwl=g1 z8-5JPhTpT+$87l#u;KR#HvG_|u;Et)8-5sTkZU!ou;Ir68-BP-*zjwPfDOMU*zhBb z!iHb#Yq8fDTiEbh-w0kI5a$Lxux+hR*leSkqbOh3n^X*D=jPLwTT!p67oz=yJlEG( z2vu`)9@zHQ3bv4_?I_B(^$wM`*|~-E?N+AO)brc_2^p^e*HrKA4wW5Jd>LD3l++4gW6t6=)o;5_WoK^JHh&1mDT`V}-CRHO z<#-be`Td5r7C$NvWw(Cj>+vTVjdYDe4V;?#W2lj}cVCXT!iwJIu} zTL7=MjR0QDhwnhCrva~dt^!^YL-1M^;I$0|ueAeSiyHyFmW$vuO9Zd20$vkB@Y-|0 zYYPZoYX!WvGXi+6RDU~?oIgY-L9700$F%pw;C*tZes1KB%&6qlRj1fEG4*dOL2n!7 zo61K@#;-cWhC#&Yj<(aRAlKHDHu;Ql=ZY^`Beulf*ulf84bhYebgJJ(~ zH~{wle#8DBu(jnzp(fb>dkFh~o+Q}+vzmeZKY!T&Qw@Oqza`lJ1GzTU=-LGPf6rn6 z&vOp;{{m)U|1Svk|7rtZ|8Eo7|10O%Zq{j{P_~|G8_YvkOMcdUh8htYAXX9-53Fe{f3eH>MeC?B8&QAMjcLg4Z4aUJGdiyjG0hwW*n@6fOj>*&}$(AMlzV zg4Z4bUh_imS|NhhrU0)=A$ZLZ!E63=Qz;8>Tn*&hNjS1<^VEPb?_izvDXm^D=`rzz zG&CqJ*1>Dm3R?17(wA(4Pc-U1Vc7(ul?H=#m$WK$%EyXE)1K+ovxGuei*az;T38E! z)(rq!MH59+Y2gok7GgZZP54wUsRR>Z)dUi(2c4~PPg-8HG;2+1QgS}e$dM+{Q)#geRuP?L3M+TPfrQ0DYisN)C~Mj3 zG;30(KTpaTERZJCgbHU1sbSF80-Y5|SRM4R#@4sAvV`btYf|Q*uGO|cno8p;{98y3 z{dt_vC`dApusmpLO?wrCWAW0-)uhaR9+it*Fqlk}D(o$!hW|Xtmk=~Rm9|jGrAZn7 z3`e&9K9z?#!`z1nWd4dlLwPX{%8MtUym$f1iw?|pF=5P&nA6PJn9ERKbc6ch6{s(o zGq+=WnRmhV5iwEB6&c|tbTyz5ISsaNhyx0d2~db+Bc{V(fC?)N286<3fE5e|phaOY z;1vu8tuZgx?ehzcwBae$^oSHqt=&set7}!f%1__XJHQ?qFIcA>+iqQ<^=R2PcGQL7AcH!XP-N zped(WHh%v@Qkh!i{t1w4q7Z1|xk&ytc9VZ#sO405f;Un!s~5;pvB<#iAD z>u9-XlkdTXA88Uc{7|9+q8L{Mulc_S=-LImhQ$G1%hdtA)`#G=Ilyav5KyKfxiQbc zz<{m+1g|v$UaLj$+8~10<^ZonBX}(V@S0KJT0mDXg4Y@VuazQrZ2-Y*bAZ=E5WE%- zcr76i@Y*nf*BV0~?(fuD59lh*W)h`FFP1ZIj;-8$0gvOpHfrLW{)5y1w7R%z_rms> zX^T-ED6_|)%<@)8p4nEK)*Cg##$5w!+>Nb{R1H5f?KG-~oYxEeB2UPTX&||=D9DWs zA-S>NNN(&U=iAUF0Na${&u$c-tkTCHFrAvac24Y{#R$c^FZAUCEn1-Y?F$c>Ra zAvcCv4OqcMLvC!ny0(g#y15yNZGhgG?YPb)rSi{HeFMh@bjI?E`wA5)j(I*0sUj+F zev8CbcRb<@8IPW%RR5{bZ*yGm$k<+qU7?!9he#Mb->4$WZ|+9YUci)}*ErWCrRq<) zew?E*I%8wSVufnDdr2bUxl|QVb#pP2Rt;-@iQ{mJ*Cx}--mPcFE~!pIXE@Xw>$KiZ z^h}i5!H}Or&X z8$x1Z=O8wwTLH1L<%!P|mF+cLu%y-cpPn3PCalrNjd1Ho7 z>v@HX^f|+7)~~SScd^R5>Sk?D)rmU&4%VfWw4+Xt9J8ykd4A%Nehcd$%=k4_IaZB0 z#q3kqw~yz#6h})x1!tAxuDf4A%thu_;$*kM%%%jD&{vn0u8X0@$1pz*Yl*O$z~RO{`ol zQ_ipf_~)pB{hVET#dEPL>X}GKfs(4=8Uh>S$FP;J8(C|enw=?)D_aO)vj@P2LI9g9 z0@!K*uo)tNtr-B@wle^>Jp{1X0bpZ70Gk~G*lGZ<$smBO0RUUBGXS>l0N4!e(>tGU zJMYTa>op#z?DrRV>(t|P&=vAyCw$Nhr(x8O`6%B}yXrG9j<-X6?7UAHkh3E|&Yrp) zghjvO=mJk*XXe1p3Iuc?saGBFbM|#57xfph7P$gqV>|luD8}9=fkCBG=~zruzW~^L5)r`m3;TJz;+7(Y;FM9dJ({e1;AGHtZU)! z-=p*u^f%dUI1Wk_2_DxL<~wdAEOuN?Sm?N(u+;Ix_qmTN-xoeEbD#gX#(nYQsQbLf zb?%EESG_NI-1xp^*D+Ys_zD4%A$ACmM8WOr@KnA1MN;oL$OKr_5MugChtq^bjU`yr zs9_gmz(m3w>;PEQKualckb7WJLx*XA4yOf+8sA`1W1ZcF0TT_28nLjbVIcg1gFFh0 z8n*PSbl9V?sG$Li8YOJn4Ac`Z^^)>al=rvs26B)a|8w6LLO)A~Rs6uyg~AW*n`et+ zpqjiyCFQ0V?vLYL;~>|=vW6G^Upm@RUG6R&JbqsX8z%$R6t*==Q)Kq<@QQPgo5^SQ zCDQlO(G+!gXciCSPK!7!IS$-BQtASovop{+tAoy2?_FtyTvgLXpnaP^y4BN`ETMHK zpI`1+;&AnuK_h!;jGG7}xIW>#es?H0e!NpJaxQVc8s1+&fbop65m8YXRwkM)#rWVJ z8+5YggY2VR|3OekJ9|1^4d~Arz<9)1iUcdDDif`je0*?E+S(xD=6VMbZvE;cXHVy= zvHev8;Q6dYBvopaiHl2-KG;`n9qhLGI(I13KYG*+yyl&*3Jt)}VHgY@dMbw~6IWsA z5L*vJhavgVcPMi(blBzvLx)@gtpn817+Vn|0AH?G;QNdgXK7q0IiktH*g6 zIbD?+$QYo8|2`#>pe&_K+*tDVq1D6MVPd}Rog>BQl#5^Q@IJPxxOdQ9zncE{`}@57 z5M{Gd$yPLZqS-;;_&$wS^E98;h?4bDlUB_Z`nmUacz>OCv>H(b+V@hkoxUR9_YV2p z`**xW9zgrTt?I2X3H~-LDoTUO_%}~t4&xd$JL&WD1@Dmi-xGMz9-#XItg@^y@&1-9 z!Ai=?_`xSWhjEP!ZS>9gE_cX7@B4X49svAet*Wf3{{irmRIXOW^C7_RIRL-mI{^Gz z0Qj}p0r1O5fZu5Vex8a!%6KsZ_*DV$+qeV3uN{D2oE-qa+|vO3ED_+R_)QrvgaE(i z0Q?s20Pt%C;J0H3z_0W)06!}P_$dPLlR|)BRYcn-scRqm%uWf|9y#$8dzmNf&DeE2 zg`6WNDs`KfO+wK;&)0^&7JvM7=;WItW|>)zS;rJ?6ba9C+nN3N z{g9B`I|#WIIYY>;9tpYGK*()K8A5I`NXV@dLT*Ax$c=Ca?FhDS^?Fto#x&HDC-|s} zGC}!`j~T8SbYC>TD?g>?V?Tu4(9RHYtIDEcLaliuRcn+9d2b@ku&;3)%)^2@{FL7x zdk(doB{`+#XBB2qFNa$3c&Y{~6AGcqhHb^QGj9k+^HX*|4jqa+OL9)jMY?SM&}B1H zT~{XLLzj*A8oF!?g3x9A3thIIv(RNLMY?PO&}B!8vH&l;5f+*V{q5D7sTOR9(@)?!mwa1?9+XdUl z&Njx1xpMx>{voAz3SHA+`*fR+Rqb0wl57M>vZ2wCWV2%X$%=_NpRm71NlBq=0(c*; z>E*M1Q$m75;YWK0!mpGcgkL=( z{Om#a)v|){vqprUA_zY&MEF&L@ayCU;n$1^KYI{2@A$_kEcJepew#aJss*5L6RK=!zgc1~Ey;f_&x6AijPa8OOwmwbkbGvwmY-;Q(eOf!>De-;gDYo>Vy1RHIiRbhWuL7GC7r$ zBOZ5Wwjk9utTl{!E;yRA)0I7hlQyN#`3LW=mhIgats`ffF|+^3pHE(%}k-KIZzx z`jLK{wh8@Z~3C%MM~S+J0d2v5tK+bJ?^9zzRIe{SVdV11`ntV5+tU5&fyeW z4cN9Xx?ZI0L6ps)$m4E$Z+%rg^S;riBks!x=yV8_t+A;i&lMzFXjG{~wJZ{zlRz$8vBXX?| zLG*OC#rb|2(g1|rumh+ONVSsdD})qMQ^pN*+pyn|@i*Ss;iQ>C}B zI{X>er2jhh-p)GL($o-U{mx%Oi@JUN{{K{T|4&IDa2&@62RR_S4QO}EYi04#BL;rs z11f$gXJ0`4pnj&fj=MIk{IYUnx1{A=L|!|uFUOZsGE(mb5!-6>2LqcqP}b&Z*-e&K z9(;6`>CO7gb4rfk-n{;f*F7I~FYL+oM1@NEfeqekJD@tY3rxSvtzi1qIKcF?!F$ct z4yGR~X8M(a=_itd>30B3KhG90{nQRH{W{>lrfCP$FX4CAs1^s)uR#u`-(fKQJX^u^ zYjS|;*9oTIL>HKTU&c&79c0Hkm0*)F#Y;%VEUaY@{43!2WTMrEVq@h{XH29&R`p|RQAy0GMqG(O~9c1-@f z7!^lsdB!vQ0QVzLtdwksDmP(6q1mlHdz>Kq9h$Rqr}1QGw>2A~KMT0{ z*X>x*4k4OEKOY$9ih=ceh@DsMkfMq7J5U~5#4Tm1%3{nn>*Dh>Gfr0sp#X`(C z??U*Vj9xLl$=T*0f`NtDZOsG@pa3Te>&uy)iW`_O!-ev7qFxC-#s0`aI07r&{#bl$ ztSkX1ch%#}pyDv*levU^T?*vKHra1Eh_3=m+-5V-ej{ZvoZMNjW)3LkF`v|x%MXis zrSyFJl7nytqTGl%+h_=t#o^@cxe{hnu@QT{d&}#S@;%7vr55EvNd60=QWPW}Z)7Wm zar0X|RdE4&d@p+r9&B$~s)2R8P*jkdmB%QB@YvL}AG%%^bE*RFYwyB+EmM$%GG=&A zWU?Q+V;UoL$Ie1`Oe=4!V2g?#V1w)7!RFbPC1s59W{}C-(-Gc;QDbbmVrE&iMUAEG zh9VI*xDifl9zma&vCL~nfV*4Z<=8YPNR1Y|uTN zD-Y)ZU6-G^;nE)bF2KWg9X0ussI2&Zx3o)K^VS z(=KwHyl!(QO6m32F=}~tdEJ-cz;+DOy4$UPGhRS?th>u<)qhe|M#WT0%zL?&54X28 zlroca4Iv`_MViKHHR;cCGC110stT$|bszIyYRzJJveTrLv83k+vEq-=Cajf4{S^+2 zqy4n1l-dwGv0ZIN7#>-Vn6jMIh7dmg0xhSb(xgvj3pnJ5&>riQe~WpKwxSGOVvmH9 zf;>WqSbeO&ZQQ7rvJ*IDMHNmB$`4`Q6VM+EZ|*@Un~+-wVev1~n%l-r`V_W|LzY*m zsRQ!AFmF?9E+d@SBcJo~#YJ>CLr*@VJC@%@apu-u7dL@8+qYC-DRB txJYeE_;DWj&3{BwoVRP8;m3^&tCYfjX^nkp`~G+XCf|t_-Yodo@IP|>JPH5+ literal 0 HcmV?d00001 diff --git a/tests/regression/om/whitebox_db_32.png b/tests/regression/om/whitebox_db_32.png deleted file mode 100644 index 3001396836c622889405d71b2494d9c23f07baa5..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 195 zcmeAS@N?(olHy`uVBq!ia0vp^3LwnE1|*BCs=ffJE>9Q7kcif!7d6G07#NrW``*6W z8GF2`m*=JZB^DMLmCU8n3LhvucxV`SHbTUr$|uBFjCG!z@YNpI#>vNBOp^m2JcUSA zgQf04q`DzekHJzqAiDNJq-KMq82UT2jx9QHH+j*ljW7OpO}%WfKtASlP*ymr|4E>$ O7(8A5T-G@yGywoXvqzKw