CVPR/LIVE
34
1#include "scene.h"2#include "aabb.h"3#include "cuda_utils.h"4#include "filter.h"5#include "shape.h"6#include <numeric>7#include <algorithm>8#include <cstring>9#include <chrono>10#include <cstddef>11 12size_t align(size_t s) {13 auto a = alignof(std::max_align_t);14 return ((s + a - 1) / a) * a;15}16 17template <typename T>18void allocate(bool use_gpu, T **p) {19 if (use_gpu) {20#ifdef __NVCC__21 checkCuda(cudaMallocManaged(p, sizeof(T)));22#else23 throw std::runtime_error("diffvg not compiled with GPU");24 assert(false);25#endif26 } else {27 *p = (T*)malloc(sizeof(T));28 }29}30 31template <typename T>32void allocate(bool use_gpu, size_t size, T **p) {33 if (use_gpu) {34#ifdef __NVCC__35 checkCuda(cudaMallocManaged(p, size * sizeof(T)));36#else37 throw std::runtime_error("diffvg not compiled with GPU");38 assert(false);39#endif40 } else {41 *p = (T*)malloc(size * sizeof(T));42 }43}44 45void copy_and_init_shapes(Scene &scene,46 const std::vector<const Shape *> &shape_list) {47 for (int shape_id = 0; shape_id < scene.num_shapes; shape_id++) {48 switch (shape_list[shape_id]->type) {49 case ShapeType::Circle: {50 Circle *p = (Circle *)scene.shapes[shape_id].ptr;51 const Circle *p_ = (const Circle*)(shape_list[shape_id]->ptr);52 *p = *p_;53 Circle *d_p = (Circle *)scene.d_shapes[shape_id].ptr;54 d_p->radius = 0;55 d_p->center = Vector2f{0, 0};56 break;57 } case ShapeType::Ellipse: {58 Ellipse *p = (Ellipse *)scene.shapes[shape_id].ptr;59 const Ellipse *p_ = (const Ellipse*)(shape_list[shape_id]->ptr);60 *p = *p_;61 Ellipse *d_p = (Ellipse *)scene.d_shapes[shape_id].ptr;62 d_p->radius = Vector2f{0, 0};63 d_p->center = Vector2f{0, 0};64 break;65 } case ShapeType::Path: {66 Path *p = (Path *)scene.shapes[shape_id].ptr;67 const Path *p_ = (const Path*)(shape_list[shape_id]->ptr);68 p->num_points = p_->num_points;69 p->num_base_points = p_->num_base_points;70 for (int i = 0; i < p_->num_base_points; i++) {71 p->num_control_points[i] = p_->num_control_points[i];72 }73 for (int i = 0; i < 2 * p_->num_points; i++) {74 p->points[i] = p_->points[i];75 }76 p->is_closed = p_->is_closed;77 p->use_distance_approx = p_->use_distance_approx;78 Path *d_p = (Path *)scene.d_shapes[shape_id].ptr;79 d_p->num_points = p_->num_points;80 d_p->num_base_points = p_->num_base_points;81 for (int i = 0; i < 2 * p_->num_points; i++) {82 d_p->points[i] = 0;83 }84 d_p->is_closed = p_->is_closed;85 if (p_->thickness != nullptr) {86 for (int i = 0; i < p_->num_points; i++) {87 p->thickness[i] = p_->thickness[i];88 d_p->thickness[i] = 0;89 }90 }91 d_p->use_distance_approx = p_->use_distance_approx;92 break;93 } case ShapeType::Rect: {94 Rect *p = (Rect *)scene.shapes[shape_id].ptr;95 const Rect *p_ = (const Rect*)(shape_list[shape_id]->ptr);96 *p = *p_;97 Rect *d_p = (Rect *)scene.d_shapes[shape_id].ptr;98 d_p->p_min = Vector2f{0, 0};99 d_p->p_max = Vector2f{0, 0};100 break;101 } default: {102 assert(false);103 break;104 }105 }106 scene.shapes[shape_id].type = shape_list[shape_id]->type;107 scene.shapes[shape_id].stroke_width = shape_list[shape_id]->stroke_width;108 scene.d_shapes[shape_id].type = shape_list[shape_id]->type;109 scene.d_shapes[shape_id].stroke_width = 0;110 }111}112 113std::vector<float>114compute_shape_length(const std::vector<const Shape *> &shape_list) {115 int num_shapes = (int)shape_list.size();116 std::vector<float> shape_length_list(num_shapes, 0.f);117 for (int shape_id = 0; shape_id < num_shapes; shape_id++) {118 auto shape_length = 0.f;119 switch (shape_list[shape_id]->type) {120 case ShapeType::Circle: {121 const Circle *p_ = (const Circle*)(shape_list[shape_id]->ptr);122 shape_length += float(2.f * M_PI) * p_->radius;123 break;124 } case ShapeType::Ellipse: {125 const Ellipse *p_ = (const Ellipse*)(shape_list[shape_id]->ptr);126 // https://en.wikipedia.org/wiki/Ellipse#Circumference127 // Ramanujan's ellipse circumference approximation128 auto a = p_->radius.x;129 auto b = p_->radius.y;130 shape_length += float(M_PI) * (3 * (a + b) - sqrt((3 * a + b) * (a + 3 * b)));131 break;132 } case ShapeType::Path: {133 const Path *p_ = (const Path*)(shape_list[shape_id]->ptr);134 auto length = 0.f;135 auto point_id = 0;136 for (int i = 0; i < p_->num_base_points; i++) {137 if (p_->num_control_points[i] == 0) {138 // Straight line139 auto i0 = point_id;140 assert(i0 < p_->num_points);141 auto i1 = (i0 + 1) % p_->num_points;142 point_id += 1;143 auto p0 = Vector2f{p_->points[2 * i0], p_->points[2 * i0 + 1]};144 auto p1 = Vector2f{p_->points[2 * i1], p_->points[2 * i1 + 1]};145 length += distance(p1, p0);146 } else if (p_->num_control_points[i] == 1) {147 // Quadratic Bezier curve148 auto i0 = point_id;149 auto i1 = i0 + 1;150 auto i2 = (i0 + 2) % p_->num_points;151 point_id += 2;152 auto p0 = Vector2f{p_->points[2 * i0], p_->points[2 * i0 + 1]};153 auto p1 = Vector2f{p_->points[2 * i1], p_->points[2 * i1 + 1]};154 auto p2 = Vector2f{p_->points[2 * i2], p_->points[2 * i2 + 1]};155 auto eval = [&](float t) -> Vector2f {156 auto tt = 1 - t;157 return (tt*tt)*p0 + (2*tt*t)*p1 + (t*t)*p2;158 };159 // We use 3-point samples to approximate the length160 auto v0 = p0;161 auto v1 = eval(0.5f);162 auto v2 = p2;163 length += distance(v1, v0) + distance(v1, v2);164 } else if (p_->num_control_points[i] == 2) {165 // Cubic Bezier curve166 auto i0 = point_id;167 auto i1 = i0 + 1;168 auto i2 = i0 + 2;169 auto i3 = (i0 + 3) % p_->num_points;170 point_id += 3;171 auto p0 = Vector2f{p_->points[2 * i0], p_->points[2 * i0 + 1]};172 auto p1 = Vector2f{p_->points[2 * i1], p_->points[2 * i1 + 1]};173 auto p2 = Vector2f{p_->points[2 * i2], p_->points[2 * i2 + 1]};174 auto p3 = Vector2f{p_->points[2 * i3], p_->points[2 * i3 + 1]};175 auto eval = [&](float t) -> Vector2f {176 auto tt = 1 - t;177 return (tt*tt*tt)*p0 + (3*tt*tt*t)*p1 + (3*tt*t*t)*p2 + (t*t*t)*p3;178 };179 // We use 4-point samples to approximate the length180 auto v0 = p0;181 auto v1 = eval(1.f/3.f);182 auto v2 = eval(2.f/3.f);183 auto v3 = p3;184 length += distance(v1, v0) + distance(v1, v2) + distance(v2, v3);185 } else {186 assert(false);187 }188 }189 assert(isfinite(length));190 shape_length += length;191 break;192 } case ShapeType::Rect: {193 const Rect *p_ = (const Rect*)(shape_list[shape_id]->ptr);194 shape_length += 2 * (p_->p_max.x - p_->p_min.x + p_->p_max.y - p_->p_min.y);195 break;196 } default: {197 assert(false);198 break;199 }200 }201 assert(isfinite(shape_length));202 shape_length_list[shape_id] = shape_length;203 }204 return shape_length_list;205}206 207void build_shape_cdfs(Scene &scene,208 const std::vector<const ShapeGroup *> &shape_group_list,209 const std::vector<float> &shape_length_list) {210 int sample_id = 0;211 for (int shape_group_id = 0; shape_group_id < (int)shape_group_list.size(); shape_group_id++) {212 const ShapeGroup *shape_group = shape_group_list[shape_group_id];213 for (int i = 0; i < shape_group->num_shapes; i++) {214 int shape_id = shape_group->shape_ids[i];215 float length = shape_length_list[shape_id];216 scene.sample_shape_id[sample_id] = shape_id;217 if (sample_id == 0) {218 scene.sample_shapes_cdf[sample_id] = length;219 } else {220 scene.sample_shapes_cdf[sample_id] = length +221 scene.sample_shapes_cdf[sample_id - 1];222 }223 assert(isfinite(length));224 scene.sample_shapes_pmf[sample_id] = length;225 scene.sample_group_id[sample_id] = shape_group_id;226 sample_id++;227 }228 }229 assert(sample_id == scene.num_total_shapes);230 auto normalization = scene.sample_shapes_cdf[scene.num_total_shapes - 1];231 if (normalization <= 0) {232 char buf[256];233 sprintf(buf, "The total length of the shape boundaries in the scene is equal or less than 0. Length = %f", normalization);234 throw std::runtime_error(buf);235 }236 if (!isfinite(normalization)) {237 char buf[256];238 sprintf(buf, "The total length of the shape boundaries in the scene is not a number. Length = %f", normalization);239 throw std::runtime_error(buf);240 }241 assert(normalization > 0);242 for (int sample_id = 0; sample_id < scene.num_total_shapes; sample_id++) {243 scene.sample_shapes_cdf[sample_id] /= normalization;244 scene.sample_shapes_pmf[sample_id] /= normalization;245 }246}247 248void build_path_cdfs(Scene &scene,249 const std::vector<const Shape *> &shape_list,250 const std::vector<float> &shape_length_list) {251 for (int shape_id = 0; shape_id < scene.num_shapes; shape_id++) {252 if (shape_list[shape_id]->type == ShapeType::Path) {253 const Path &path = shape_list[shape_id]->as_path();254 float *pmf = scene.path_length_pmf[shape_id];255 float *cdf = scene.path_length_cdf[shape_id];256 int *point_id_map = scene.path_point_id_map[shape_id];257 auto path_length = shape_length_list[shape_id];258 auto inv_length = 1.f / path_length;259 auto point_id = 0;260 for (int i = 0; i < path.num_base_points; i++) {261 point_id_map[i] = point_id;262 if (path.num_control_points[i] == 0) {263 // Straight line264 auto i0 = point_id;265 auto i1 = (i0 + 1) % path.num_points;266 point_id += 1;267 auto p0 = Vector2f{path.points[2 * i0], path.points[2 * i0 + 1]};268 auto p1 = Vector2f{path.points[2 * i1], path.points[2 * i1 + 1]};269 auto d = distance(p0, p1) * inv_length;270 pmf[i] = d;271 if (i == 0) {272 cdf[i] = d;273 } else {274 cdf[i] = d + cdf[i - 1];275 }276 } else if (path.num_control_points[i] == 1) {277 // Quadratic Bezier curve278 auto i0 = point_id;279 auto i1 = i0 + 1;280 auto i2 = (i0 + 2) % path.num_points;281 point_id += 2;282 auto p0 = Vector2f{path.points[2 * i0], path.points[2 * i0 + 1]};283 auto p1 = Vector2f{path.points[2 * i1], path.points[2 * i1 + 1]};284 auto p2 = Vector2f{path.points[2 * i2], path.points[2 * i2 + 1]};285 auto eval = [&](float t) -> Vector2f {286 auto tt = 1 - t;287 return (tt*tt)*p0 + (2*tt*t)*p1 + (t*t)*p2;288 };289 // We use 3-point samples to approximate the length290 auto v0 = p0;291 auto v1 = eval(0.5f);292 auto v2 = p2;293 auto d = (distance(v0, v1) + distance(v1, v2)) * inv_length;294 pmf[i] = d;295 if (i == 0) {296 cdf[i] = d;297 } else {298 cdf[i] = d + cdf[i - 1];299 }300 } else if (path.num_control_points[i] == 2) {301 // Cubic Bezier curve302 auto i0 = point_id;303 auto i1 = point_id + 1;304 auto i2 = point_id + 2;305 auto i3 = (point_id + 3) % path.num_points;306 point_id += 3;307 auto p0 = Vector2f{path.points[2 * i0], path.points[2 * i0 + 1]};308 auto p1 = Vector2f{path.points[2 * i1], path.points[2 * i1 + 1]};309 auto p2 = Vector2f{path.points[2 * i2], path.points[2 * i2 + 1]};310 auto p3 = Vector2f{path.points[2 * i3], path.points[2 * i3 + 1]};311 auto eval = [&](float t) -> Vector2f {312 auto tt = 1 - t;313 return (tt*tt*tt)*p0 + (3*tt*tt*t)*p1 + (3*tt*t*t)*p2 + (t*t*t)*p3;314 };315 // We use 4-point samples to approximate the length316 auto v0 = p0;317 auto v1 = eval(1.f/3.f);318 auto v2 = eval(2.f/3.f);319 auto v3 = p3;320 auto d = (distance(v1, v0) + distance(v1, v2) + distance(v2, v3)) * inv_length;321 pmf[i] = d;322 if (i == 0) {323 cdf[i] = d;324 } else {325 cdf[i] = d + cdf[i - 1];326 }327 } else {328 assert(false);329 }330 }331 }332 }333}334 335void copy_and_init_shape_groups(Scene &scene,336 const std::vector<const ShapeGroup *> &shape_group_list) {337 for (int group_id = 0; group_id < scene.num_shape_groups; group_id++) {338 const ShapeGroup *shape_group = shape_group_list[group_id];339 auto copy_and_init_color = [&](const ColorType &color_type, void *color_ptr, void *target_ptr, void *d_target_ptr) {340 switch (color_type) {341 case ColorType::Constant: {342 Constant *c = (Constant*)target_ptr;343 Constant *d_c = (Constant*)d_target_ptr;344 const Constant *c_ = (const Constant*)color_ptr;345 *c = *c_;346 d_c->color = Vector4{0, 0, 0, 0};347 break;348 } case ColorType::LinearGradient: {349 LinearGradient *c = (LinearGradient*)target_ptr;350 LinearGradient *d_c = (LinearGradient*)d_target_ptr;351 const LinearGradient *c_ = (const LinearGradient*)color_ptr;352 c->begin = c_->begin;353 c->end = c_->end;354 c->num_stops = c_->num_stops;355 for (int i = 0; i < c_->num_stops; i++) {356 c->stop_offsets[i] = c_->stop_offsets[i];357 }358 for (int i = 0; i < 4 * c_->num_stops; i++) {359 c->stop_colors[i] = c_->stop_colors[i];360 }361 d_c->begin = Vector2f{0, 0};362 d_c->end = Vector2f{0, 0};363 d_c->num_stops = c_->num_stops;364 for (int i = 0; i < c_->num_stops; i++) {365 d_c->stop_offsets[i] = 0;366 }367 for (int i = 0; i < 4 * c_->num_stops; i++) {368 d_c->stop_colors[i] = 0;369 }370 break;371 } case ColorType::RadialGradient: {372 RadialGradient *c = (RadialGradient*)target_ptr;373 RadialGradient *d_c = (RadialGradient*)d_target_ptr;374 const RadialGradient *c_ = (const RadialGradient*)color_ptr;375 c->center = c_->center;376 c->radius = c_->radius;377 c->num_stops = c_->num_stops;378 for (int i = 0; i < c_->num_stops; i++) {379 c->stop_offsets[i] = c_->stop_offsets[i];380 }381 for (int i = 0; i < 4 * c_->num_stops; i++) {382 c->stop_colors[i] = c_->stop_colors[i];383 }384 d_c->center = Vector2f{0, 0};385 d_c->radius = Vector2f{0, 0};386 d_c->num_stops = c_->num_stops;387 for (int i = 0; i < c_->num_stops; i++) {388 d_c->stop_offsets[i] = 0;389 }390 for (int i = 0; i < 4 * c_->num_stops; i++) {391 d_c->stop_colors[i] = 0;392 }393 break;394 } default: {395 assert(false);396 }397 }398 };399 for (int i = 0; i < shape_group->num_shapes; i++) {400 scene.shape_groups[group_id].shape_ids[i] = shape_group->shape_ids[i];401 }402 scene.shape_groups[group_id].num_shapes = shape_group->num_shapes;403 scene.shape_groups[group_id].use_even_odd_rule = shape_group->use_even_odd_rule;404 scene.shape_groups[group_id].canvas_to_shape = shape_group->canvas_to_shape;405 scene.shape_groups[group_id].shape_to_canvas = shape_group->shape_to_canvas;406 scene.d_shape_groups[group_id].shape_ids = nullptr;407 scene.d_shape_groups[group_id].num_shapes = shape_group->num_shapes;408 scene.d_shape_groups[group_id].use_even_odd_rule = shape_group->use_even_odd_rule;409 scene.d_shape_groups[group_id].canvas_to_shape = Matrix3x3f{};410 scene.d_shape_groups[group_id].shape_to_canvas = Matrix3x3f{};411 412 scene.shape_groups[group_id].fill_color_type = shape_group->fill_color_type;413 scene.d_shape_groups[group_id].fill_color_type = shape_group->fill_color_type;414 if (shape_group->fill_color != nullptr) {415 copy_and_init_color(shape_group->fill_color_type,416 shape_group->fill_color,417 scene.shape_groups[group_id].fill_color,418 scene.d_shape_groups[group_id].fill_color);419 }420 scene.shape_groups[group_id].stroke_color_type = shape_group->stroke_color_type;421 scene.d_shape_groups[group_id].stroke_color_type = shape_group->stroke_color_type;422 if (shape_group->stroke_color != nullptr) {423 copy_and_init_color(shape_group->stroke_color_type,424 shape_group->stroke_color,425 scene.shape_groups[group_id].stroke_color,426 scene.d_shape_groups[group_id].stroke_color);427 }428 }429}430 431DEVICE uint32_t morton2D(const Vector2f &p, int canvas_width, int canvas_height) {432 auto scene_bounds = Vector2f{canvas_width, canvas_height};433 auto pp = p / scene_bounds;434 TVector2<uint32_t> pp_i{pp.x * 1023, pp.y * 1023};435 return (expand_bits(pp_i.x) << 1u) |436 (expand_bits(pp_i.y) << 0u);437}438 439template <bool sort>440void build_bvh(const Scene &scene, BVHNode *nodes, int num_primitives) {441 auto bvh_size = 2 * num_primitives - 1;442 if (bvh_size > 1) {443 if (sort) {444 // Sort by Morton code445 std::sort(nodes, nodes + num_primitives,446 [&] (const BVHNode &n0, const BVHNode &n1) {447 auto p0 = 0.5f * (n0.box.p_min + n0.box.p_max);448 auto p1 = 0.5f * (n1.box.p_min + n1.box.p_max);449 auto m0 = morton2D(p0, scene.canvas_width, scene.canvas_height);450 auto m1 = morton2D(p1, scene.canvas_width, scene.canvas_height);451 return m0 < m1;452 });453 }454 for (int i = num_primitives; i < bvh_size; i++) {455 nodes[i] = BVHNode{-1, -1, AABB{}, 0.f};456 }457 int prev_beg = 0;458 int prev_end = num_primitives;459 // For handling odd number of nodes at a level460 int leftover = prev_end % 2 == 0 ? -1 : prev_end - 1;461 while (prev_end - prev_beg >= 1 || leftover != -1) {462 int length = (prev_end - prev_beg) / 2;463 if ((prev_end - prev_beg) % 2 == 1 && leftover != -1 &&464 leftover != prev_end - 1) {465 length += 1;466 }467 for (int i = 0; i < length; i++) {468 BVHNode node;469 node.child0 = prev_beg + 2 * i;470 node.child1 = prev_beg + 2 * i + 1;471 if (node.child1 >= prev_end) {472 assert(leftover != -1);473 node.child1 = leftover;474 leftover = -1;475 }476 AABB child0_box = nodes[node.child0].box;477 AABB child1_box = nodes[node.child1].box;478 node.box = merge(child0_box, child1_box);479 node.max_radius = std::max(nodes[node.child0].max_radius,480 nodes[node.child1].max_radius);481 nodes[prev_end + i] = node;482 }483 if (length == 1 && leftover == -1) {484 break;485 }486 prev_beg = prev_end;487 prev_end = prev_beg + length;488 if (length % 2 == 1 && leftover == -1) {489 leftover = prev_end - 1;490 }491 }492 }493 assert(nodes[2 * num_primitives - 2].child0 != -1);494}495 496void compute_bounding_boxes(Scene &scene,497 const std::vector<const Shape *> &shape_list,498 const std::vector<const ShapeGroup *> &shape_group_list) {499 for (int shape_id = 0; shape_id < scene.num_shapes; shape_id++) {500 switch (shape_list[shape_id]->type) {501 case ShapeType::Circle: {502 const Circle *p = (const Circle*)(shape_list[shape_id]->ptr);503 scene.shapes_bbox[shape_id] = AABB{p->center - p->radius,504 p->center + p->radius};505 break;506 } case ShapeType::Ellipse: {507 const Ellipse *p = (const Ellipse*)(shape_list[shape_id]->ptr);508 scene.shapes_bbox[shape_id] = AABB{p->center - p->radius,509 p->center + p->radius};510 break;511 } case ShapeType::Path: {512 const Path *p = (const Path*)(shape_list[shape_id]->ptr);513 AABB box;514 if (p->num_points > 0) {515 box = AABB{Vector2f{p->points[0], p->points[1]},516 Vector2f{p->points[0], p->points[1]}};517 }518 for (int i = 1; i < p->num_points; i++) {519 box = merge(box, Vector2f{p->points[2 * i], p->points[2 * i + 1]});520 }521 scene.shapes_bbox[shape_id] = box;522 std::vector<AABB> boxes(p->num_base_points);523 std::vector<float> thickness(p->num_base_points);524 std::vector<int> first_point_id(p->num_base_points);525 auto r = shape_list[shape_id]->stroke_width;526 auto point_id = 0;527 for (int i = 0; i < p->num_base_points; i++) {528 first_point_id[i] = point_id;529 if (p->num_control_points[i] == 0) {530 // Straight line531 auto i0 = point_id;532 auto i1 = (i0 + 1) % p->num_points;533 point_id += 1;534 auto p0 = Vector2f{p->points[2 * i0], p->points[2 * i0 + 1]};535 auto p1 = Vector2f{p->points[2 * i1], p->points[2 * i1 + 1]};536 boxes[i] = AABB();537 boxes[i] = merge(boxes[i], p0);538 boxes[i] = merge(boxes[i], p1);539 auto r0 = r;540 auto r1 = r;541 // override radius if path has thickness542 if (p->thickness != nullptr) {543 r0 = p->thickness[i0];544 r1 = p->thickness[i1];545 }546 thickness[i] = max(r0, r1);547 } else if (p->num_control_points[i] == 1) {548 // Quadratic Bezier curve549 auto i0 = point_id;550 auto i1 = i0 + 1;551 auto i2 = (i0 + 2) % p->num_points;552 point_id += 2;553 auto p0 = Vector2f{p->points[2 * i0], p->points[2 * i0 + 1]};554 auto p1 = Vector2f{p->points[2 * i1], p->points[2 * i1 + 1]};555 auto p2 = Vector2f{p->points[2 * i2], p->points[2 * i2 + 1]};556 boxes[i] = AABB();557 boxes[i] = merge(boxes[i], p0);558 boxes[i] = merge(boxes[i], p1);559 boxes[i] = merge(boxes[i], p2);560 auto r0 = r;561 auto r1 = r;562 auto r2 = r;563 // override radius if path has thickness564 if (p->thickness != nullptr) {565 r0 = p->thickness[i0];566 r1 = p->thickness[i1];567 r2 = p->thickness[i2];568 }569 thickness[i] = max(max(r0, r1), r2);570 } else if (p->num_control_points[i] == 2) {571 // Cubic Bezier curve572 auto i0 = point_id;573 auto i1 = i0 + 1;574 auto i2 = i0 + 2;575 auto i3 = (i0 + 3) % p->num_points;576 point_id += 3;577 auto p0 = Vector2f{p->points[2 * i0], p->points[2 * i0 + 1]};578 auto p1 = Vector2f{p->points[2 * i1], p->points[2 * i1 + 1]};579 auto p2 = Vector2f{p->points[2 * i2], p->points[2 * i2 + 1]};580 auto p3 = Vector2f{p->points[2 * i3], p->points[2 * i3 + 1]};581 boxes[i] = AABB();582 boxes[i] = merge(boxes[i], p0);583 boxes[i] = merge(boxes[i], p1);584 boxes[i] = merge(boxes[i], p2);585 boxes[i] = merge(boxes[i], p3);586 auto r0 = r;587 auto r1 = r;588 auto r2 = r;589 auto r3 = r;590 // override radius if path has thickness591 if (p->thickness != nullptr) {592 r0 = p->thickness[i0];593 r1 = p->thickness[i1];594 r2 = p->thickness[i2];595 r3 = p->thickness[i3];596 }597 thickness[i] = max(max(max(r0, r1), r2), r3);598 } else {599 assert(false);600 }601 }602 // Sort the boxes by y603 std::vector<int> idx(boxes.size());604 std::iota(idx.begin(), idx.end(), 0);605 std::sort(idx.begin(), idx.end(), [&](int i0, int i1) {606 const AABB &b0 = boxes[i0];607 const AABB &b1 = boxes[i1];608 auto b0y = 0.5f * (b0.p_min.y + b0.p_max.y);609 auto b1y = 0.5f * (b1.p_min.y + b1.p_max.y);610 return b0y < b1y;611 });612 BVHNode *nodes = scene.path_bvhs[shape_id];613 for (int i = 0; i < (int)idx.size(); i++) {614 nodes[i] = BVHNode{idx[i],615 -(first_point_id[idx[i]]+1),616 boxes[idx[i]],617 thickness[idx[i]]};618 }619 build_bvh<false /*sort*/>(scene, nodes, boxes.size());620 break;621 } case ShapeType::Rect: {622 const Rect *p = (const Rect*)(shape_list[shape_id]->ptr);623 scene.shapes_bbox[shape_id] = AABB{p->p_min, p->p_max};624 break;625 } default: {626 assert(false);627 break;628 }629 }630 }631 632 for (int shape_group_id = 0; shape_group_id < (int)shape_group_list.size(); shape_group_id++) {633 const ShapeGroup *shape_group = shape_group_list[shape_group_id];634 // Build a BVH for each shape group635 BVHNode *nodes = scene.shape_groups_bvh_nodes[shape_group_id];636 for (int i = 0; i < shape_group->num_shapes; i++) {637 auto shape_id = shape_group->shape_ids[i];638 auto r = shape_group->stroke_color == nullptr ? 0 : shape_list[shape_id]->stroke_width;639 nodes[i] = BVHNode{shape_id,640 -1,641 scene.shapes_bbox[shape_id],642 r};643 }644 build_bvh<true /*sort*/>(scene, nodes, shape_group->num_shapes);645 }646 647 BVHNode *nodes = scene.bvh_nodes;648 for (int shape_group_id = 0; shape_group_id < (int)shape_group_list.size(); shape_group_id++) {649 const ShapeGroup *shape_group = shape_group_list[shape_group_id];650 auto max_radius = shape_list[shape_group->shape_ids[0]]->stroke_width;651 if (shape_list[shape_group->shape_ids[0]]->type == ShapeType::Path) {652 const Path *p = (const Path*)(shape_list[shape_group->shape_ids[0]]->ptr);653 if (p->thickness != nullptr) {654 const BVHNode *nodes = scene.path_bvhs[shape_group->shape_ids[0]];655 max_radius = nodes[0].max_radius;656 }657 }658 for (int i = 1; i < shape_group->num_shapes; i++) {659 auto shape_id = shape_group->shape_ids[i];660 auto shape = shape_list[shape_id];661 auto r = shape->stroke_width;662 if (shape->type == ShapeType::Path) {663 const Path *p = (const Path*)(shape_list[shape_id]->ptr);664 if (p->thickness != nullptr) {665 const BVHNode *nodes = scene.path_bvhs[shape_id];666 r = nodes[0].max_radius;667 }668 }669 max_radius = std::max(max_radius, r);670 }671 // Fetch group bbox from BVH672 auto bbox = scene.shape_groups_bvh_nodes[shape_group_id][2 * shape_group->num_shapes - 2].box;673 // Transform box from local to world space674 nodes[shape_group_id].child0 = shape_group_id;675 nodes[shape_group_id].child1 = -1;676 nodes[shape_group_id].box = transform(shape_group->shape_to_canvas, bbox);677 if (shape_group->stroke_color == nullptr) {678 nodes[shape_group_id].max_radius = 0;679 } else {680 nodes[shape_group_id].max_radius = max_radius;681 }682 }683 build_bvh<true /*sort*/>(scene, nodes, shape_group_list.size());684}685 686template <bool alloc_mode>687size_t allocate_buffers(Scene &scene,688 const std::vector<const Shape *> &shape_list,689 const std::vector<const ShapeGroup *> &shape_group_list) {690 auto num_shapes = shape_list.size();691 auto num_shape_groups = shape_group_list.size();692 693 size_t buffer_size = 0;694 if (alloc_mode) scene.shapes = (Shape*)&scene.buffer[buffer_size];695 buffer_size += align(sizeof(Shape) * num_shapes);696 if (alloc_mode) scene.d_shapes = (Shape*)&scene.buffer[buffer_size];697 buffer_size += align(sizeof(Shape) * num_shapes); 698 if (alloc_mode) scene.shape_groups = (ShapeGroup*)&scene.buffer[buffer_size];699 buffer_size += align(sizeof(ShapeGroup) * num_shape_groups);700 if (alloc_mode) scene.d_shape_groups = (ShapeGroup*)&scene.buffer[buffer_size];701 buffer_size += align(sizeof(ShapeGroup) * num_shape_groups);702 if (alloc_mode) scene.sample_shapes_cdf = (float*)&scene.buffer[buffer_size];703 buffer_size += align(sizeof(float) * scene.num_total_shapes);704 if (alloc_mode) scene.sample_shapes_pmf = (float*)&scene.buffer[buffer_size];705 buffer_size += align(sizeof(float) * scene.num_total_shapes);706 if (alloc_mode) scene.sample_shape_id = (int*)&scene.buffer[buffer_size];707 buffer_size += align(sizeof(int) * scene.num_total_shapes);708 if (alloc_mode) scene.sample_group_id = (int*)&scene.buffer[buffer_size];709 buffer_size += align(sizeof(int) * scene.num_total_shapes);710 if (alloc_mode) scene.shapes_length = (float*)&scene.buffer[buffer_size];711 buffer_size += align(sizeof(float) * num_shapes);712 if (alloc_mode) scene.path_length_cdf = (float**)&scene.buffer[buffer_size];713 buffer_size += align(sizeof(float*) * num_shapes);714 if (alloc_mode) scene.path_length_pmf = (float**)&scene.buffer[buffer_size];715 buffer_size += align(sizeof(float*) * num_shapes);716 if (alloc_mode) scene.path_point_id_map = (int**)&scene.buffer[buffer_size];717 buffer_size += align(sizeof(int*) * num_shapes);718 if (alloc_mode) scene.filter = (Filter*)&scene.buffer[buffer_size];719 buffer_size += align(sizeof(Filter));720 if (alloc_mode) scene.d_filter = (DFilter*)&scene.buffer[buffer_size];721 buffer_size += align(sizeof(DFilter));722 if (alloc_mode) scene.shapes_bbox = (AABB*)&scene.buffer[buffer_size];723 buffer_size += align(sizeof(AABB) * num_shapes);724 if (alloc_mode) scene.path_bvhs = (BVHNode**)&scene.buffer[buffer_size];725 buffer_size += align(sizeof(BVHNode*) * num_shapes);726 if (alloc_mode) scene.shape_groups_bvh_nodes = (BVHNode**)&scene.buffer[buffer_size];727 buffer_size += align(sizeof(BVHNode*) * num_shape_groups);728 if (alloc_mode) scene.bvh_nodes = (BVHNode*)&scene.buffer[buffer_size];729 buffer_size += align(sizeof(BVHNode) * (2 * num_shape_groups - 1));730 731 if (alloc_mode) {732 for (int i = 0; i < num_shapes; i++) {733 scene.path_length_cdf[i] = nullptr;734 scene.path_length_pmf[i] = nullptr;735 scene.path_point_id_map[i] = nullptr;736 scene.path_bvhs[i] = nullptr;737 }738 }739 740 for (int shape_id = 0; shape_id < scene.num_shapes; shape_id++) {741 switch (shape_list[shape_id]->type) {742 case ShapeType::Circle: {743 if (alloc_mode) scene.shapes[shape_id].ptr = (Circle*)&scene.buffer[buffer_size];744 buffer_size += align(sizeof(Circle)); // scene.shapes[shape_id].ptr745 if (alloc_mode) scene.d_shapes[shape_id].ptr = (Circle*)&scene.buffer[buffer_size];746 buffer_size += align(sizeof(Circle)); // scene.d_shapes[shape_id].ptr747 break;748 } case ShapeType::Ellipse: {749 if (alloc_mode) scene.shapes[shape_id].ptr = (Ellipse*)&scene.buffer[buffer_size];750 buffer_size += align(sizeof(Ellipse)); // scene.shapes[shape_id].ptr751 if (alloc_mode) scene.d_shapes[shape_id].ptr = (Ellipse*)&scene.buffer[buffer_size];752 buffer_size += align(sizeof(Ellipse)); // scene.d_shapes[shape_id].ptr753 break;754 } case ShapeType::Path: {755 if (alloc_mode) scene.shapes[shape_id].ptr = (Path*)&scene.buffer[buffer_size];756 buffer_size += align(sizeof(Path)); // scene.shapes[shape_id].ptr757 if (alloc_mode) scene.d_shapes[shape_id].ptr = (Path*)&scene.buffer[buffer_size];758 buffer_size += align(sizeof(Path)); // scene.d_shapes[shape_id].ptr759 760 const Path *p_ = (const Path*)(shape_list[shape_id]->ptr);761 Path *p = nullptr, *d_p = nullptr;762 if (alloc_mode) p = (Path*)scene.shapes[shape_id].ptr;763 if (alloc_mode) d_p = (Path*)scene.d_shapes[shape_id].ptr; 764 if (alloc_mode) p->num_control_points = (int*)&scene.buffer[buffer_size];765 buffer_size += align(sizeof(int) * p_->num_base_points); // p->num_control_points766 if (alloc_mode) p->points = (float*)&scene.buffer[buffer_size];767 buffer_size += align(sizeof(float) * (2 * p_->num_points)); // p->points768 if (alloc_mode) d_p->points = (float*)&scene.buffer[buffer_size];769 buffer_size += align(sizeof(float) * (2 * p_->num_points)); // d_p->points770 if (p_->thickness != nullptr) {771 if (alloc_mode) p->thickness = (float*)&scene.buffer[buffer_size];772 buffer_size += align(sizeof(float) * p_->num_points); // p->thickness773 if (alloc_mode) d_p->thickness = (float*)&scene.buffer[buffer_size];774 buffer_size += align(sizeof(float) * p_->num_points); // d_p->thickness775 } else {776 if (alloc_mode) p->thickness = nullptr;777 if (alloc_mode) d_p->thickness = nullptr;778 }779 if (alloc_mode) scene.path_length_pmf[shape_id] = (float*)&scene.buffer[buffer_size];780 buffer_size += align(sizeof(float) * p_->num_base_points); // scene.path_length_pmf781 if (alloc_mode) scene.path_length_cdf[shape_id] = (float*)&scene.buffer[buffer_size];782 buffer_size += align(sizeof(float) * p_->num_base_points); // scene.path_length_cdf783 if (alloc_mode) scene.path_point_id_map[shape_id] = (int*)&scene.buffer[buffer_size];784 buffer_size += align(sizeof(int) * p_->num_base_points); // scene.path_point_id_map785 if (alloc_mode) scene.path_bvhs[shape_id] = (BVHNode*)&scene.buffer[buffer_size];786 buffer_size += align(sizeof(BVHNode) * (2 * p_->num_base_points - 1));787 break;788 } case ShapeType::Rect: {789 if (alloc_mode) scene.shapes[shape_id].ptr = (Ellipse*)&scene.buffer[buffer_size];790 buffer_size += align(sizeof(Rect)); // scene.shapes[shape_id].ptr791 if (alloc_mode) scene.d_shapes[shape_id].ptr = (Ellipse*)&scene.buffer[buffer_size];792 buffer_size += align(sizeof(Rect)); // scene.d_shapes[shape_id].ptr793 break;794 } default: {795 assert(false);796 break;797 }798 }799 }800 801 for (int group_id = 0; group_id < scene.num_shape_groups; group_id++) {802 const ShapeGroup *shape_group = shape_group_list[group_id];803 if (shape_group->fill_color != nullptr) {804 switch (shape_group->fill_color_type) {805 case ColorType::Constant: {806 if (alloc_mode) scene.shape_groups[group_id].fill_color = (Constant*)&scene.buffer[buffer_size];807 buffer_size += align(sizeof(Constant)); // color808 if (alloc_mode) scene.d_shape_groups[group_id].fill_color = (Constant*)&scene.buffer[buffer_size];809 buffer_size += align(sizeof(Constant)); // d_color810 break;811 } case ColorType::LinearGradient: {812 if (alloc_mode) scene.shape_groups[group_id].fill_color = (LinearGradient*)&scene.buffer[buffer_size];813 buffer_size += align(sizeof(LinearGradient)); // color814 if (alloc_mode) scene.d_shape_groups[group_id].fill_color = (LinearGradient*)&scene.buffer[buffer_size];815 buffer_size += align(sizeof(LinearGradient)); // d_color816 817 const LinearGradient *c_ = (const LinearGradient *)shape_group->fill_color;818 LinearGradient *c = nullptr, *d_c = nullptr;819 if (alloc_mode) c = (LinearGradient *)scene.shape_groups[group_id].fill_color;820 if (alloc_mode) d_c = (LinearGradient *)scene.d_shape_groups[group_id].fill_color;821 if (alloc_mode) c->stop_offsets = (float*)&scene.buffer[buffer_size];822 buffer_size += align(sizeof(float) * c_->num_stops); // c->stop_offsets823 if (alloc_mode) c->stop_colors = (float*)&scene.buffer[buffer_size];824 buffer_size += align(sizeof(float) * 4 * c_->num_stops); // c->stop_colors825 if (alloc_mode) d_c->stop_offsets = (float*)&scene.buffer[buffer_size];826 buffer_size += align(sizeof(float) * c_->num_stops); // d_c->stop_offsets827 if (alloc_mode) d_c->stop_colors = (float*)&scene.buffer[buffer_size];828 buffer_size += align(sizeof(float) * 4 * c_->num_stops); // d_c->stop_colors829 break;830 } case ColorType::RadialGradient: {831 if (alloc_mode) scene.shape_groups[group_id].fill_color = (RadialGradient*)&scene.buffer[buffer_size];832 buffer_size += align(sizeof(RadialGradient)); // color833 if (alloc_mode) scene.d_shape_groups[group_id].fill_color = (RadialGradient*)&scene.buffer[buffer_size];834 buffer_size += align(sizeof(RadialGradient)); // d_color835 836 const RadialGradient *c_ = (const RadialGradient *)shape_group->fill_color;837 RadialGradient *c = nullptr, *d_c = nullptr;838 if (alloc_mode) c = (RadialGradient *)scene.shape_groups[group_id].fill_color;839 if (alloc_mode) d_c = (RadialGradient *)scene.d_shape_groups[group_id].fill_color;840 if (alloc_mode) c->stop_offsets = (float*)&scene.buffer[buffer_size];841 buffer_size += align(sizeof(float) * c_->num_stops); // c->stop_offsets842 if (alloc_mode) c->stop_colors = (float*)&scene.buffer[buffer_size];843 buffer_size += align(sizeof(float) * 4 * c_->num_stops); // c->stop_colors844 if (alloc_mode) d_c->stop_offsets = (float*)&scene.buffer[buffer_size];845 buffer_size += align(sizeof(float) * c_->num_stops); // d_c->stop_offsets846 if (alloc_mode) d_c->stop_colors = (float*)&scene.buffer[buffer_size];847 buffer_size += align(sizeof(float) * 4 * c_->num_stops); // d_c->stop_colors848 break;849 } default: {850 assert(false);851 }852 }853 } else {854 if (alloc_mode) scene.shape_groups[group_id].fill_color = nullptr;855 if (alloc_mode) scene.d_shape_groups[group_id].fill_color = nullptr;856 }857 if (shape_group->stroke_color != nullptr) {858 switch (shape_group->stroke_color_type) {859 case ColorType::Constant: {860 if (alloc_mode) scene.shape_groups[group_id].stroke_color = (Constant*)&scene.buffer[buffer_size];861 buffer_size += align(sizeof(Constant)); // color862 if (alloc_mode) scene.d_shape_groups[group_id].stroke_color = (Constant*)&scene.buffer[buffer_size];863 buffer_size += align(sizeof(Constant)); // d_color864 break;865 } case ColorType::LinearGradient: {866 if (alloc_mode) scene.shape_groups[group_id].stroke_color = (LinearGradient*)&scene.buffer[buffer_size];867 buffer_size += align(sizeof(LinearGradient)); // color868 if (alloc_mode) scene.shape_groups[group_id].stroke_color = (LinearGradient*)&scene.buffer[buffer_size];869 buffer_size += align(sizeof(LinearGradient)); // d_color870 871 const LinearGradient *c_ = (const LinearGradient *)shape_group->stroke_color;872 LinearGradient *c = nullptr, *d_c = nullptr;873 if (alloc_mode) c = (LinearGradient *)scene.shape_groups[group_id].stroke_color;874 if (alloc_mode) d_c = (LinearGradient *)scene.d_shape_groups[group_id].stroke_color;875 if (alloc_mode) c->stop_offsets = (float*)&scene.buffer[buffer_size];876 buffer_size += align(sizeof(float) * c_->num_stops); // c->stop_offsets877 if (alloc_mode) c->stop_colors = (float*)&scene.buffer[buffer_size];878 buffer_size += align(sizeof(float) * 4 * c_->num_stops); // c->stop_colors879 if (alloc_mode) d_c->stop_offsets = (float*)&scene.buffer[buffer_size];880 buffer_size += align(sizeof(float) * c_->num_stops); // d_c->stop_offsets881 if (alloc_mode) d_c->stop_colors = (float*)&scene.buffer[buffer_size];882 buffer_size += align(sizeof(float) * 4 * c_->num_stops); // d_c->stop_colors883 break;884 } case ColorType::RadialGradient: {885 if (alloc_mode) scene.shape_groups[group_id].stroke_color = (RadialGradient*)&scene.buffer[buffer_size];886 buffer_size += align(sizeof(RadialGradient)); // color887 if (alloc_mode) scene.shape_groups[group_id].stroke_color = (RadialGradient*)&scene.buffer[buffer_size];888 buffer_size += align(sizeof(RadialGradient)); // d_color889 890 const RadialGradient *c_ = (const RadialGradient *)shape_group->stroke_color;891 RadialGradient *c = nullptr, *d_c = nullptr;892 if (alloc_mode) c = (RadialGradient *)scene.shape_groups[group_id].stroke_color;893 if (alloc_mode) d_c = (RadialGradient *)scene.d_shape_groups[group_id].stroke_color;894 if (alloc_mode) c->stop_offsets = (float*)&scene.buffer[buffer_size];895 buffer_size += align(sizeof(float) * c_->num_stops); // c->stop_offsets896 if (alloc_mode) c->stop_colors = (float*)&scene.buffer[buffer_size];897 buffer_size += align(sizeof(float) * 4 * c_->num_stops); // c->stop_colors898 if (alloc_mode) d_c->stop_offsets = (float*)&scene.buffer[buffer_size];899 buffer_size += align(sizeof(float) * c_->num_stops); // d_c->stop_offsets900 if (alloc_mode) d_c->stop_colors = (float*)&scene.buffer[buffer_size];901 buffer_size += align(sizeof(float) * 4 * c_->num_stops); // d_c->stop_colors902 break;903 } default: {904 assert(false);905 }906 }907 } else {908 if (alloc_mode) scene.shape_groups[group_id].stroke_color = nullptr;909 if (alloc_mode) scene.d_shape_groups[group_id].stroke_color = nullptr;910 }911 if (alloc_mode) scene.shape_groups[group_id].shape_ids = (int*)&scene.buffer[buffer_size];912 buffer_size += align(sizeof(int) * shape_group->num_shapes); // shape_group->shape_ids913 if (alloc_mode) scene.shape_groups_bvh_nodes[group_id] = (BVHNode*)&scene.buffer[buffer_size];914 buffer_size += align(sizeof(BVHNode) * (2 * shape_group->num_shapes - 1)); // scene.shape_groups_bvh_nodes[group_id]915 }916 return buffer_size;917}918 919Scene::Scene(int canvas_width,920 int canvas_height,921 const std::vector<const Shape *> &shape_list,922 const std::vector<const ShapeGroup *> &shape_group_list,923 const Filter &filter,924 bool use_gpu,925 int gpu_index)926 : canvas_width(canvas_width),927 canvas_height(canvas_height),928 num_shapes(shape_list.size()),929 num_shape_groups(shape_group_list.size()),930 use_gpu(use_gpu),931 gpu_index(gpu_index) {932 if (num_shapes == 0) {933 return;934 }935 // Shape group may reuse some of the shapes,936 // record the total number of shapes.937 int num_total_shapes = 0;938 for (const ShapeGroup *sg : shape_group_list) {939 num_total_shapes += sg->num_shapes;940 }941 this->num_total_shapes = num_total_shapes;942 943 // Memory initialization944#ifdef __NVCC__945 int old_device_id = -1;946#endif947 if (use_gpu) {948#ifdef __NVCC__949 checkCuda(cudaGetDevice(&old_device_id));950 if (gpu_index != -1) {951 checkCuda(cudaSetDevice(gpu_index));952 }953#else954 throw std::runtime_error("diffvg not compiled with GPU");955 assert(false);956#endif957 }958 959 size_t buffer_size = allocate_buffers<false /*alloc_mode*/>(*this, shape_list, shape_group_list);960 // Allocate a huge buffer for everything961 allocate<uint8_t>(use_gpu, buffer_size, &buffer);962 // memset(buffer, 111, buffer_size);963 // Actually distribute the buffer964 allocate_buffers<true /*alloc_mode*/>(*this, shape_list, shape_group_list);965 copy_and_init_shapes(*this, shape_list);966 copy_and_init_shape_groups(*this, shape_group_list);967 968 std::vector<float> shape_length_list = compute_shape_length(shape_list);969 // Copy shape_length970 if (use_gpu) {971#ifdef __NVCC__972 checkCuda(cudaMemcpy(this->shapes_length, &shape_length_list[0], num_shapes * sizeof(float), cudaMemcpyHostToDevice));973#else974 throw std::runtime_error("diffvg not compiled with GPU");975 assert(false);976#endif977 } else {978 memcpy(this->shapes_length, &shape_length_list[0], num_shapes * sizeof(float));979 }980 build_shape_cdfs(*this, shape_group_list, shape_length_list);981 build_path_cdfs(*this, shape_list, shape_length_list);982 compute_bounding_boxes(*this, shape_list, shape_group_list);983 984 // Filter initialization985 *(this->filter) = filter;986 this->d_filter->radius = 0;987 988 if (use_gpu) {989#ifdef __NVCC__990 if (old_device_id != -1) {991 checkCuda(cudaSetDevice(old_device_id));992 }993#else994 throw std::runtime_error("diffvg not compiled with GPU");995 assert(false);996#endif997 }998}999 1000Scene::~Scene() {1001 if (num_shapes == 0) {1002 return;1003 }1004 if (use_gpu) {1005#ifdef __NVCC__1006 int old_device_id = -1;1007 checkCuda(cudaGetDevice(&old_device_id));1008 if (gpu_index != -1) {1009 checkCuda(cudaSetDevice(gpu_index));1010 }1011 1012 checkCuda(cudaFree(buffer));1013 1014 checkCuda(cudaSetDevice(old_device_id));1015#else1016 // Don't throw because C++ don't want a destructor to throw.1017 std::cerr << "diffvg not compiled with GPU";1018 exit(1);1019#endif1020 } else {1021 free(buffer);1022 }1023}1024 1025Shape Scene::get_d_shape(int shape_id) const {1026 return d_shapes[shape_id];1027}1028 1029ShapeGroup Scene::get_d_shape_group(int group_id) const {1030 return d_shape_groups[group_id];1031}1032 1033float Scene::get_d_filter_radius() const {1034 return d_filter->radius;1035}1036