diff options
48 files changed, 2047 insertions, 69 deletions
diff --git a/intern/cycles/blender/CMakeLists.txt b/intern/cycles/blender/CMakeLists.txt index 1f05b1aa234..fe7d0b89bb0 100644 --- a/intern/cycles/blender/CMakeLists.txt +++ b/intern/cycles/blender/CMakeLists.txt @@ -40,6 +40,7 @@ set(SRC object_cull.cpp output_driver.cpp particles.cpp + pointcloud.cpp curves.cpp logging.cpp python.cpp diff --git a/intern/cycles/blender/addon/ui.py b/intern/cycles/blender/addon/ui.py index fd86d75a301..e4fbc898070 100644 --- a/intern/cycles/blender/addon/ui.py +++ b/intern/cycles/blender/addon/ui.py @@ -1020,7 +1020,7 @@ class CYCLES_OBJECT_PT_motion_blur(CyclesButtonsPanel, Panel): def poll(cls, context): ob = context.object if CyclesButtonsPanel.poll(context) and ob: - if ob.type in {'MESH', 'CURVE', 'CURVE', 'SURFACE', 'FONT', 'META', 'CAMERA'}: + if ob.type in {'MESH', 'CURVE', 'CURVE', 'SURFACE', 'FONT', 'META', 'CAMERA', 'HAIR', 'POINTCLOUD'}: return True if ob.instance_type == 'COLLECTION' and ob.instance_collection: return True diff --git a/intern/cycles/blender/geometry.cpp b/intern/cycles/blender/geometry.cpp index a5f36a3a17d..78c803b7adb 100644 --- a/intern/cycles/blender/geometry.cpp +++ b/intern/cycles/blender/geometry.cpp @@ -19,6 +19,7 @@ #include "scene/hair.h" #include "scene/mesh.h" #include "scene/object.h" +#include "scene/pointcloud.h" #include "scene/volume.h" #include "blender/sync.h" @@ -39,6 +40,10 @@ static Geometry::Type determine_geom_type(BObjectInfo &b_ob_info, bool use_parti return Geometry::HAIR; } + if (b_ob_info.object_data.is_a(&RNA_PointCloud)) { + return Geometry::POINTCLOUD; + } + if (b_ob_info.object_data.is_a(&RNA_Volume) || (b_ob_info.object_data == b_ob_info.real_object.data() && object_fluid_gas_domain_find(b_ob_info.real_object))) { @@ -111,6 +116,9 @@ Geometry *BlenderSync::sync_geometry(BL::Depsgraph &b_depsgraph, else if (geom_type == Geometry::VOLUME) { geom = scene->create_node<Volume>(); } + else if (geom_type == Geometry::POINTCLOUD) { + geom = scene->create_node<PointCloud>(); + } else { geom = scene->create_node<Mesh>(); } @@ -170,6 +178,10 @@ Geometry *BlenderSync::sync_geometry(BL::Depsgraph &b_depsgraph, Volume *volume = static_cast<Volume *>(geom); sync_volume(b_ob_info, volume); } + else if (geom_type == Geometry::POINTCLOUD) { + PointCloud *pointcloud = static_cast<PointCloud *>(geom); + sync_pointcloud(pointcloud, b_ob_info); + } else { Mesh *mesh = static_cast<Mesh *>(geom); sync_mesh(b_depsgraph, b_ob_info, mesh); @@ -231,6 +243,10 @@ void BlenderSync::sync_geometry_motion(BL::Depsgraph &b_depsgraph, object_fluid_gas_domain_find(b_ob_info.real_object)) { /* No volume motion blur support yet. */ } + else if (b_ob_info.object_data.is_a(&RNA_PointCloud)) { + PointCloud *pointcloud = static_cast<PointCloud *>(geom); + sync_pointcloud_motion(pointcloud, b_ob_info, motion_step); + } else { Mesh *mesh = static_cast<Mesh *>(geom); sync_mesh_motion(b_depsgraph, b_ob_info, mesh, motion_step); diff --git a/intern/cycles/blender/object.cpp b/intern/cycles/blender/object.cpp index d6c52f7faf9..86314d3b196 100644 --- a/intern/cycles/blender/object.cpp +++ b/intern/cycles/blender/object.cpp @@ -72,7 +72,8 @@ bool BlenderSync::object_is_geometry(BObjectInfo &b_ob_info) BL::Object::type_enum type = b_ob_info.iter_object.type(); - if (type == BL::Object::type_VOLUME || type == BL::Object::type_HAIR) { + if (type == BL::Object::type_VOLUME || type == BL::Object::type_HAIR || + type == BL::Object::type_POINTCLOUD) { /* Will be exported attached to mesh. */ return true; } @@ -206,7 +207,7 @@ Object *BlenderSync::sync_object(BL::Depsgraph &b_depsgraph, return NULL; } - /* only interested in object that we can create meshes from */ + /* only interested in object that we can create geometry from */ if (!object_is_geometry(b_ob_info)) { return NULL; } diff --git a/intern/cycles/blender/pointcloud.cpp b/intern/cycles/blender/pointcloud.cpp new file mode 100644 index 00000000000..a9e616a468f --- /dev/null +++ b/intern/cycles/blender/pointcloud.cpp @@ -0,0 +1,253 @@ +/* + * Copyright 2011-2013 Blender Foundation + * + * 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. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "scene/pointcloud.h" +#include "scene/attribute.h" +#include "scene/scene.h" + +#include "blender/sync.h" +#include "blender/util.h" + +#include "util/foreach.h" +#include "util/hash.h" + +CCL_NAMESPACE_BEGIN + +template<typename TypeInCycles, typename GetValueAtIndex> +static void fill_generic_attribute(BL::PointCloud &b_pointcloud, + TypeInCycles *data, + const GetValueAtIndex &get_value_at_index) +{ + const int num_points = b_pointcloud.points.length(); + for (int i = 0; i < num_points; i++) { + data[i] = get_value_at_index(i); + } +} + +static void copy_attributes(PointCloud *pointcloud, BL::PointCloud b_pointcloud) +{ + AttributeSet &attributes = pointcloud->attributes; + for (BL::Attribute &b_attribute : b_pointcloud.attributes) { + const ustring name{b_attribute.name().c_str()}; + + if (attributes.find(name)) { + continue; + } + + const AttributeElement element = ATTR_ELEMENT_VERTEX; + const BL::Attribute::data_type_enum b_data_type = b_attribute.data_type(); + switch (b_data_type) { + case BL::Attribute::data_type_FLOAT: { + BL::FloatAttribute b_float_attribute{b_attribute}; + Attribute *attr = attributes.add(name, TypeFloat, element); + float *data = attr->data_float(); + fill_generic_attribute( + b_pointcloud, data, [&](int i) { return b_float_attribute.data[i].value(); }); + break; + } + case BL::Attribute::data_type_BOOLEAN: { + BL::BoolAttribute b_bool_attribute{b_attribute}; + Attribute *attr = attributes.add(name, TypeFloat, element); + float *data = attr->data_float(); + fill_generic_attribute( + b_pointcloud, data, [&](int i) { return (float)b_bool_attribute.data[i].value(); }); + break; + } + case BL::Attribute::data_type_INT: { + BL::IntAttribute b_int_attribute{b_attribute}; + Attribute *attr = attributes.add(name, TypeFloat, element); + float *data = attr->data_float(); + fill_generic_attribute( + b_pointcloud, data, [&](int i) { return (float)b_int_attribute.data[i].value(); }); + break; + } + case BL::Attribute::data_type_FLOAT_VECTOR: { + BL::FloatVectorAttribute b_vector_attribute{b_attribute}; + Attribute *attr = attributes.add(name, TypeVector, element); + float3 *data = attr->data_float3(); + fill_generic_attribute(b_pointcloud, data, [&](int i) { + BL::Array<float, 3> v = b_vector_attribute.data[i].vector(); + return make_float3(v[0], v[1], v[2]); + }); + break; + } + case BL::Attribute::data_type_FLOAT_COLOR: { + BL::FloatColorAttribute b_color_attribute{b_attribute}; + Attribute *attr = attributes.add(name, TypeRGBA, element); + float4 *data = attr->data_float4(); + fill_generic_attribute(b_pointcloud, data, [&](int i) { + BL::Array<float, 4> v = b_color_attribute.data[i].color(); + return make_float4(v[0], v[1], v[2], v[3]); + }); + break; + } + case BL::Attribute::data_type_FLOAT2: { + BL::Float2Attribute b_float2_attribute{b_attribute}; + Attribute *attr = attributes.add(name, TypeFloat2, element); + float2 *data = attr->data_float2(); + fill_generic_attribute(b_pointcloud, data, [&](int i) { + BL::Array<float, 2> v = b_float2_attribute.data[i].vector(); + return make_float2(v[0], v[1]); + }); + break; + } + default: + /* Not supported. */ + break; + } + } +} + +static void export_pointcloud(Scene *scene, PointCloud *pointcloud, BL::PointCloud b_pointcloud) +{ + /* TODO: optimize so we can straight memcpy arrays from Blender? */ + + /* Add requested attributes. */ + Attribute *attr_random = NULL; + if (pointcloud->need_attribute(scene, ATTR_STD_POINT_RANDOM)) { + attr_random = pointcloud->attributes.add(ATTR_STD_POINT_RANDOM); + } + + /* Reserve memory. */ + const int num_points = b_pointcloud.points.length(); + pointcloud->reserve(num_points); + + /* Export points. */ + BL::PointCloud::points_iterator b_point_iter; + for (b_pointcloud.points.begin(b_point_iter); b_point_iter != b_pointcloud.points.end(); + ++b_point_iter) { + BL::Point b_point = *b_point_iter; + const float3 co = get_float3(b_point.co()); + const float radius = b_point.radius(); + pointcloud->add_point(co, radius); + + /* Random number per point. */ + if (attr_random != NULL) { + attr_random->add(hash_uint2_to_float(b_point.index(), 0)); + } + } + + /* Export attributes */ + copy_attributes(pointcloud, b_pointcloud); +} + +static void export_pointcloud_motion(PointCloud *pointcloud, + BL::PointCloud b_pointcloud, + int motion_step) +{ + /* Find or add attribute. */ + Attribute *attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); + bool new_attribute = false; + + if (!attr_mP) { + attr_mP = pointcloud->attributes.add(ATTR_STD_MOTION_VERTEX_POSITION); + new_attribute = true; + } + + /* Export motion points. */ + const int num_points = pointcloud->num_points(); + float3 *mP = attr_mP->data_float3() + motion_step * num_points; + bool have_motion = false; + int num_motion_points = 0; + const array<float3> &pointcloud_points = pointcloud->get_points(); + + BL::PointCloud::points_iterator b_point_iter; + for (b_pointcloud.points.begin(b_point_iter); b_point_iter != b_pointcloud.points.end(); + ++b_point_iter) { + BL::Point b_point = *b_point_iter; + + if (num_motion_points < num_points) { + float3 P = get_float3(b_point.co()); + P.w = b_point.radius(); + mP[num_motion_points] = P; + have_motion = have_motion || (P != pointcloud_points[num_motion_points]); + num_motion_points++; + } + } + + /* In case of new attribute, we verify if there really was any motion. */ + if (new_attribute) { + if (num_motion_points != num_points || !have_motion) { + pointcloud->attributes.remove(ATTR_STD_MOTION_VERTEX_POSITION); + } + else if (motion_step > 0) { + /* Motion, fill up previous steps that we might have skipped because + * they had no motion, but we need them anyway now. */ + for (int step = 0; step < motion_step; step++) { + pointcloud->copy_center_to_motion_step(step); + } + } + } + + /* Export attributes */ + copy_attributes(pointcloud, b_pointcloud); +} + +void BlenderSync::sync_pointcloud(PointCloud *pointcloud, BObjectInfo &b_ob_info) +{ + size_t old_numpoints = pointcloud->num_points(); + + array<Node *> used_shaders = pointcloud->get_used_shaders(); + + PointCloud new_pointcloud; + new_pointcloud.set_used_shaders(used_shaders); + + /* TODO: add option to filter out points in the view layer. */ + BL::PointCloud b_pointcloud(b_ob_info.object_data); + export_pointcloud(scene, &new_pointcloud, b_pointcloud); + + /* update original sockets */ + for (const SocketType &socket : new_pointcloud.type->inputs) { + /* Those sockets are updated in sync_object, so do not modify them. */ + if (socket.name == "use_motion_blur" || socket.name == "motion_steps" || + socket.name == "used_shaders") { + continue; + } + pointcloud->set_value(socket, new_pointcloud, socket); + } + + pointcloud->attributes.clear(); + foreach (Attribute &attr, new_pointcloud.attributes.attributes) { + pointcloud->attributes.attributes.push_back(std::move(attr)); + } + + /* tag update */ + const bool rebuild = (pointcloud && old_numpoints != pointcloud->num_points()); + pointcloud->tag_update(scene, rebuild); +} + +void BlenderSync::sync_pointcloud_motion(PointCloud *pointcloud, + BObjectInfo &b_ob_info, + int motion_step) +{ + /* Skip if nothing exported. */ + if (pointcloud->num_points() == 0) { + return; + } + + /* Export deformed coordinates. */ + if (ccl::BKE_object_is_deform_modified(b_ob_info, b_scene, preview)) { + /* PointCloud object. */ + BL::PointCloud b_pointcloud(b_ob_info.object_data); + export_pointcloud_motion(pointcloud, b_pointcloud, motion_step); + } + else { + /* No deformation on this frame, copy coordinates if other frames did have it. */ + pointcloud->copy_center_to_motion_step(motion_step); + } +} + +CCL_NAMESPACE_END diff --git a/intern/cycles/blender/sync.h b/intern/cycles/blender/sync.h index 2db4125ba1c..d074f90bb1b 100644 --- a/intern/cycles/blender/sync.h +++ b/intern/cycles/blender/sync.h @@ -169,12 +169,16 @@ class BlenderSync { Hair *hair, BL::Mesh &b_mesh, BObjectInfo &b_ob_info, bool motion, int motion_step = 0); bool object_has_particle_hair(BL::Object b_ob); + /* Point Cloud */ + void sync_pointcloud(PointCloud *pointcloud, BObjectInfo &b_ob_info); + void sync_pointcloud_motion(PointCloud *pointcloud, BObjectInfo &b_ob_info, int motion_step = 0); + /* Camera */ void sync_camera_motion( BL::RenderSettings &b_render, BL::Object &b_ob, int width, int height, float motion_time); /* Geometry */ - Geometry *sync_geometry(BL::Depsgraph &b_depsgrpah, + Geometry *sync_geometry(BL::Depsgraph &b_depsgraph, BObjectInfo &b_ob_info, bool object_updated, bool use_particle_hair, diff --git a/intern/cycles/bvh/build.cpp b/intern/cycles/bvh/build.cpp index 3ce268dfb25..91198e4e2a2 100644 --- a/intern/cycles/bvh/build.cpp +++ b/intern/cycles/bvh/build.cpp @@ -26,6 +26,7 @@ #include "scene/hair.h" #include "scene/mesh.h" #include "scene/object.h" +#include "scene/pointcloud.h" #include "scene/scene.h" #include "util/algorithm.h" @@ -113,9 +114,9 @@ void BVHBuild::add_reference_triangles(BoundBox &root, else { /* Motion triangles, trace optimized case: we split triangle * primitives into separate nodes for each of the time steps. - * This way we minimize overlap of neighbor curve primitives. + * This way we minimize overlap of neighbor triangle primitives. */ - const int num_bvh_steps = params.num_motion_curve_steps * 2 + 1; + const int num_bvh_steps = params.num_motion_triangle_steps * 2 + 1; const float num_bvh_steps_inv_1 = 1.0f / (num_bvh_steps - 1); const size_t num_verts = mesh->verts.size(); const size_t num_steps = mesh->motion_steps; @@ -269,6 +270,101 @@ void BVHBuild::add_reference_curves(BoundBox &root, BoundBox ¢er, Hair *hair } } +void BVHBuild::add_reference_points(BoundBox &root, + BoundBox ¢er, + PointCloud *pointcloud, + int i) +{ + const Attribute *point_attr_mP = NULL; + if (pointcloud->has_motion_blur()) { + point_attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); + } + + const float3 *points_data = &pointcloud->points[0]; + const float *radius_data = &pointcloud->radius[0]; + const size_t num_points = pointcloud->num_points(); + const float3 *motion_data = (point_attr_mP) ? point_attr_mP->data_float3() : NULL; + const size_t num_steps = pointcloud->get_motion_steps(); + + if (point_attr_mP == NULL) { + /* Really simple logic for static points. */ + for (uint j = 0; j < num_points; j++) { + const PointCloud::Point point = pointcloud->get_point(j); + BoundBox bounds = BoundBox::empty; + point.bounds_grow(points_data, radius_data, bounds); + if (bounds.valid()) { + references.push_back(BVHReference(bounds, j, i, PRIMITIVE_POINT)); + root.grow(bounds); + center.grow(bounds.center2()); + } + } + } + else if (params.num_motion_point_steps == 0 || params.use_spatial_split) { + /* Simple case of motion points: single node for the whole + * shutter time. Lowest memory usage but less optimal + * rendering. + */ + /* TODO(sergey): Support motion steps for spatially split BVH. */ + for (uint j = 0; j < num_points; j++) { + const PointCloud::Point point = pointcloud->get_point(j); + BoundBox bounds = BoundBox::empty; + point.bounds_grow(points_data, radius_data, bounds); + for (size_t step = 0; step < num_steps - 1; step++) { + point.bounds_grow(motion_data + step * num_points, radius_data, bounds); + } + if (bounds.valid()) { + references.push_back(BVHReference(bounds, j, i, PRIMITIVE_MOTION_POINT)); + root.grow(bounds); + center.grow(bounds.center2()); + } + } + } + else { + /* Motion points, trace optimized case: we split point + * primitives into separate nodes for each of the time steps. + * This way we minimize overlap of neighbor point primitives. + */ + const int num_bvh_steps = params.num_motion_point_steps * 2 + 1; + const float num_bvh_steps_inv_1 = 1.0f / (num_bvh_steps - 1); + + for (uint j = 0; j < num_points; j++) { + const PointCloud::Point point = pointcloud->get_point(j); + const size_t num_steps = pointcloud->get_motion_steps(); + const float3 *point_steps = point_attr_mP->data_float3(); + + /* Calculate bounding box of the previous time step. + * Will be reused later to avoid duplicated work on + * calculating BVH time step boundbox. + */ + float4 prev_key = point.motion_key( + points_data, radius_data, point_steps, num_points, num_steps, 0.0f, j); + BoundBox prev_bounds = BoundBox::empty; + point.bounds_grow(prev_key, prev_bounds); + /* Create all primitive time steps, */ + for (int bvh_step = 1; bvh_step < num_bvh_steps; ++bvh_step) { + const float curr_time = (float)(bvh_step)*num_bvh_steps_inv_1; + float4 curr_key = point.motion_key( + points_data, radius_data, point_steps, num_points, num_steps, curr_time, j); + BoundBox curr_bounds = BoundBox::empty; + point.bounds_grow(curr_key, curr_bounds); + BoundBox bounds = prev_bounds; + bounds.grow(curr_bounds); + if (bounds.valid()) { + const float prev_time = (float)(bvh_step - 1) * num_bvh_steps_inv_1; + references.push_back( + BVHReference(bounds, j, i, PRIMITIVE_MOTION_POINT, prev_time, curr_time)); + root.grow(bounds); + center.grow(bounds.center2()); + } + /* Current time boundbox becomes previous one for the + * next time step. + */ + prev_bounds = curr_bounds; + } + } + } +} + void BVHBuild::add_reference_geometry(BoundBox &root, BoundBox ¢er, Geometry *geom, @@ -282,6 +378,10 @@ void BVHBuild::add_reference_geometry(BoundBox &root, Hair *hair = static_cast<Hair *>(geom); add_reference_curves(root, center, hair, object_index); } + else if (geom->geometry_type == Geometry::POINTCLOUD) { + PointCloud *pointcloud = static_cast<PointCloud *>(geom); + add_reference_points(root, center, pointcloud, object_index); + } } void BVHBuild::add_reference_object(BoundBox &root, BoundBox ¢er, Object *ob, int i) @@ -311,6 +411,10 @@ static size_t count_primitives(Geometry *geom) Hair *hair = static_cast<Hair *>(geom); return count_curve_segments(hair); } + else if (geom->geometry_type == Geometry::POINTCLOUD) { + PointCloud *pointcloud = static_cast<PointCloud *>(geom); + return pointcloud->num_points(); + } return 0; } @@ -328,8 +432,9 @@ void BVHBuild::add_references(BVHRange &root) if (!ob->get_geometry()->is_instanced()) { num_alloc_references += count_primitives(ob->get_geometry()); } - else + else { num_alloc_references++; + } } else { num_alloc_references += count_primitives(ob->get_geometry()); @@ -394,7 +499,7 @@ BVHNode *BVHBuild::run() spatial_min_overlap = root.bounds().safe_area() * params.spatial_split_alpha; spatial_free_index = 0; - need_prim_time = params.num_motion_curve_steps > 0 || params.num_motion_triangle_steps > 0; + need_prim_time = params.use_motion_steps(); /* init progress updates */ double build_start_time; @@ -535,7 +640,8 @@ bool BVHBuild::range_within_max_leaf_size(const BVHRange &range, const vector<BVHReference> &references) const { size_t size = range.size(); - size_t max_leaf_size = max(params.max_triangle_leaf_size, params.max_curve_leaf_size); + size_t max_leaf_size = max(max(params.max_triangle_leaf_size, params.max_curve_leaf_size), + params.max_point_leaf_size); if (size > max_leaf_size) return false; @@ -544,6 +650,8 @@ bool BVHBuild::range_within_max_leaf_size(const BVHRange &range, size_t num_motion_triangles = 0; size_t num_curves = 0; size_t num_motion_curves = 0; + size_t num_points = 0; + size_t num_motion_points = 0; for (int i = 0; i < size; i++) { const BVHReference &ref = references[range.start() + i]; @@ -564,12 +672,22 @@ bool BVHBuild::range_within_max_leaf_size(const BVHRange &range, num_triangles++; } } + else if (ref.prim_type() & PRIMITIVE_ALL_POINT) { + if (ref.prim_type() & PRIMITIVE_ALL_MOTION) { + num_motion_points++; + } + else { + num_points++; + } + } } return (num_triangles <= params.max_triangle_leaf_size) && (num_motion_triangles <= params.max_motion_triangle_leaf_size) && (num_curves <= params.max_curve_leaf_size) && - (num_motion_curves <= params.max_motion_curve_leaf_size); + (num_motion_curves <= params.max_motion_curve_leaf_size) && + (num_points <= params.max_point_leaf_size) && + (num_motion_points <= params.max_motion_point_leaf_size); } /* multithreaded binning builder */ diff --git a/intern/cycles/bvh/build.h b/intern/cycles/bvh/build.h index 06b318f1ee0..5b9bb59d9f8 100644 --- a/intern/cycles/bvh/build.h +++ b/intern/cycles/bvh/build.h @@ -39,6 +39,7 @@ class Geometry; class Hair; class Mesh; class Object; +class PointCloud; class Progress; /* BVH Builder */ @@ -68,6 +69,7 @@ class BVHBuild { /* Adding references. */ void add_reference_triangles(BoundBox &root, BoundBox ¢er, Mesh *mesh, int i); void add_reference_curves(BoundBox &root, BoundBox ¢er, Hair *hair, int i); + void add_reference_points(BoundBox &root, BoundBox ¢er, PointCloud *pointcloud, int i); void add_reference_geometry(BoundBox &root, BoundBox ¢er, Geometry *geom, int i); void add_reference_object(BoundBox &root, BoundBox ¢er, Object *ob, int i); void add_references(BVHRange &root); diff --git a/intern/cycles/bvh/bvh2.cpp b/intern/cycles/bvh/bvh2.cpp index 04290602145..744e7fa9898 100644 --- a/intern/cycles/bvh/bvh2.cpp +++ b/intern/cycles/bvh/bvh2.cpp @@ -20,6 +20,7 @@ #include "scene/hair.h" #include "scene/mesh.h" #include "scene/object.h" +#include "scene/pointcloud.h" #include "bvh/build.h" #include "bvh/node.h" @@ -409,6 +410,30 @@ void BVH2::refit_primitives(int start, int end, BoundBox &bbox, uint &visibility } } } + else if (pack.prim_type[prim] & PRIMITIVE_ALL_POINT) { + /* Points. */ + const PointCloud *pointcloud = static_cast<const PointCloud *>(ob->get_geometry()); + int prim_offset = (params.top_level) ? pointcloud->prim_offset : 0; + const float3 *points = &pointcloud->points[0]; + const float *radius = &pointcloud->radius[0]; + PointCloud::Point point = pointcloud->get_point(pidx - prim_offset); + + point.bounds_grow(points, radius, bbox); + + /* Motion points. */ + if (pointcloud->get_use_motion_blur()) { + Attribute *attr = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); + + if (attr) { + size_t pointcloud_size = pointcloud->points.size(); + size_t steps = pointcloud->get_motion_steps() - 1; + float3 *point_steps = attr->data_float3(); + + for (size_t i = 0; i < steps; i++) + point.bounds_grow(point_steps + i * pointcloud_size, radius, bbox); + } + } + } else { /* Triangles. */ const Mesh *mesh = static_cast<const Mesh *>(ob->get_geometry()); @@ -505,7 +530,8 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size) pack.leaf_nodes.resize(leaf_nodes_size); pack.object_node.resize(objects.size()); - if (params.num_motion_curve_steps > 0 || params.num_motion_triangle_steps > 0) { + if (params.num_motion_curve_steps > 0 || params.num_motion_triangle_steps > 0 || + params.num_motion_point_steps > 0) { pack.prim_time.resize(prim_index_size); } diff --git a/intern/cycles/bvh/embree.cpp b/intern/cycles/bvh/embree.cpp index b54b38f2798..eab193f45cb 100644 --- a/intern/cycles/bvh/embree.cpp +++ b/intern/cycles/bvh/embree.cpp @@ -45,6 +45,7 @@ # include "scene/hair.h" # include "scene/mesh.h" # include "scene/object.h" +# include "scene/pointcloud.h" # include "util/foreach.h" # include "util/log.h" @@ -245,7 +246,7 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args) } } -static void rtc_filter_func_thick_curve(const RTCFilterFunctionNArguments *args) +static void rtc_filter_func_backface_cull(const RTCFilterFunctionNArguments *args) { const RTCRay *ray = (RTCRay *)args->ray; RTCHit *hit = (RTCHit *)args->hit; @@ -258,7 +259,7 @@ static void rtc_filter_func_thick_curve(const RTCFilterFunctionNArguments *args) } } -static void rtc_filter_occluded_func_thick_curve(const RTCFilterFunctionNArguments *args) +static void rtc_filter_occluded_func_backface_cull(const RTCFilterFunctionNArguments *args) { const RTCRay *ray = (RTCRay *)args->ray; RTCHit *hit = (RTCHit *)args->hit; @@ -410,6 +411,12 @@ void BVHEmbree::add_object(Object *ob, int i) add_curves(ob, hair, i); } } + else if (geom->geometry_type == Geometry::POINTCLOUD) { + PointCloud *pointcloud = static_cast<PointCloud *>(geom); + if (pointcloud->num_points() > 0) { + add_points(ob, pointcloud, i); + } + } } void BVHEmbree::add_instance(Object *ob, int i) @@ -624,6 +631,89 @@ void BVHEmbree::set_curve_vertex_buffer(RTCGeometry geom_id, const Hair *hair, c } } +void BVHEmbree::set_point_vertex_buffer(RTCGeometry geom_id, + const PointCloud *pointcloud, + const bool update) +{ + const Attribute *attr_mP = NULL; + size_t num_motion_steps = 1; + if (pointcloud->has_motion_blur()) { + attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); + if (attr_mP) { + num_motion_steps = pointcloud->get_motion_steps(); + } + } + + const size_t num_points = pointcloud->num_points(); + + /* Copy the point data to Embree */ + const int t_mid = (num_motion_steps - 1) / 2; + const float *radius = pointcloud->get_radius().data(); + for (int t = 0; t < num_motion_steps; ++t) { + const float3 *verts; + if (t == t_mid || attr_mP == NULL) { + verts = pointcloud->get_points().data(); + } + else { + int t_ = (t > t_mid) ? (t - 1) : t; + verts = &attr_mP->data_float3()[t_ * num_points]; + } + + float4 *rtc_verts = (update) ? (float4 *)rtcGetGeometryBufferData( + geom_id, RTC_BUFFER_TYPE_VERTEX, t) : + (float4 *)rtcSetNewGeometryBuffer(geom_id, + RTC_BUFFER_TYPE_VERTEX, + t, + RTC_FORMAT_FLOAT4, + sizeof(float) * 4, + num_points); + + assert(rtc_verts); + if (rtc_verts) { + for (size_t j = 0; j < num_points; ++j) { + rtc_verts[j] = float3_to_float4(verts[j]); + rtc_verts[j].w = radius[j]; + } + } + + if (update) { + rtcUpdateGeometryBuffer(geom_id, RTC_BUFFER_TYPE_VERTEX, t); + } + } +} + +void BVHEmbree::add_points(const Object *ob, const PointCloud *pointcloud, int i) +{ + size_t prim_offset = pointcloud->prim_offset; + + const Attribute *attr_mP = NULL; + size_t num_motion_steps = 1; + if (pointcloud->has_motion_blur()) { + attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); + if (attr_mP) { + num_motion_steps = pointcloud->get_motion_steps(); + } + } + + enum RTCGeometryType type = RTC_GEOMETRY_TYPE_SPHERE_POINT; + + RTCGeometry geom_id = rtcNewGeometry(rtc_device, type); + + rtcSetGeometryBuildQuality(geom_id, build_quality); + rtcSetGeometryTimeStepCount(geom_id, num_motion_steps); + + set_point_vertex_buffer(geom_id, pointcloud, false); + + rtcSetGeometryUserData(geom_id, (void *)prim_offset); + rtcSetGeometryIntersectFilterFunction(geom_id, rtc_filter_func_backface_cull); + rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func_backface_cull); + rtcSetGeometryMask(geom_id, ob->visibility_for_tracing()); + + rtcCommitGeometry(geom_id); + rtcAttachGeometryByID(scene, geom_id, i * 2); + rtcReleaseGeometry(geom_id); +} + void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i) { size_t prim_offset = hair->curve_segment_offset; @@ -678,8 +768,8 @@ void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i) rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func); } else { - rtcSetGeometryIntersectFilterFunction(geom_id, rtc_filter_func_thick_curve); - rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func_thick_curve); + rtcSetGeometryIntersectFilterFunction(geom_id, rtc_filter_func_backface_cull); + rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func_backface_cull); } rtcSetGeometryMask(geom_id, ob->visibility_for_tracing()); @@ -716,6 +806,14 @@ void BVHEmbree::refit(Progress &progress) rtcCommitGeometry(geom); } } + else if (geom->geometry_type == Geometry::POINTCLOUD) { + PointCloud *pointcloud = static_cast<PointCloud *>(geom); + if (pointcloud->num_points() > 0) { + RTCGeometry geom = rtcGetGeometry(scene, geom_id); + set_point_vertex_buffer(geom, pointcloud, true); + rtcCommitGeometry(geom); + } + } } geom_id += 2; } diff --git a/intern/cycles/bvh/embree.h b/intern/cycles/bvh/embree.h index 746ca97b504..3b30b2bbcf7 100644 --- a/intern/cycles/bvh/embree.h +++ b/intern/cycles/bvh/embree.h @@ -33,6 +33,7 @@ CCL_NAMESPACE_BEGIN class Hair; class Mesh; +class PointCloud; class BVHEmbree : public BVH { public: @@ -51,11 +52,15 @@ class BVHEmbree : public BVH { void add_object(Object *ob, int i); void add_instance(Object *ob, int i); void add_curves(const Object *ob, const Hair *hair, int i); + void add_points(const Object *ob, const PointCloud *pointcloud, int i); void add_triangles(const Object *ob, const Mesh *mesh, int i); private: void set_tri_vertex_buffer(RTCGeometry geom_id, const Mesh *mesh, const bool update); void set_curve_vertex_buffer(RTCGeometry geom_id, const Hair *hair, const bool update); + void set_point_vertex_buffer(RTCGeometry geom_id, + const PointCloud *pointcloud, + const bool update); RTCDevice rtc_device; enum RTCBuildQuality build_quality; diff --git a/intern/cycles/bvh/params.h b/intern/cycles/bvh/params.h index 8f185a2640f..16edf2e88e4 100644 --- a/intern/cycles/bvh/params.h +++ b/intern/cycles/bvh/params.h @@ -83,6 +83,8 @@ class BVHParams { int max_motion_triangle_leaf_size; int max_curve_leaf_size; int max_motion_curve_leaf_size; + int max_point_leaf_size; + int max_motion_point_leaf_size; /* object or mesh level bvh */ bool top_level; @@ -98,13 +100,13 @@ class BVHParams { /* Split time range to this number of steps and create leaf node for each * of this time steps. * - * Speeds up rendering of motion curve primitives in the cost of higher - * memory usage. + * Speeds up rendering of motion primitives in the cost of higher memory usage. */ - int num_motion_curve_steps; /* Same as above, but for triangle primitives. */ int num_motion_triangle_steps; + int num_motion_curve_steps; + int num_motion_point_steps; /* Same as in SceneParams. */ int bvh_type; @@ -132,6 +134,8 @@ class BVHParams { max_motion_triangle_leaf_size = 8; max_curve_leaf_size = 1; max_motion_curve_leaf_size = 4; + max_point_leaf_size = 8; + max_motion_point_leaf_size = 8; top_level = false; bvh_layout = BVH_LAYOUT_BVH2; @@ -139,6 +143,7 @@ class BVHParams { num_motion_curve_steps = 0; num_motion_triangle_steps = 0; + num_motion_point_steps = 0; bvh_type = 0; @@ -166,6 +171,12 @@ class BVHParams { return (size <= min_leaf_size || level >= MAX_DEPTH); } + bool use_motion_steps() + { + return num_motion_curve_steps > 0 || num_motion_triangle_steps > 0 || + num_motion_point_steps > 0; + } + /* Gets best matching BVH. * * If the requested layout is supported by the device, it will be used. diff --git a/intern/cycles/bvh/split.cpp b/intern/cycles/bvh/split.cpp index 102c50e2979..34d12de97c0 100644 --- a/intern/cycles/bvh/split.cpp +++ b/intern/cycles/bvh/split.cpp @@ -23,6 +23,7 @@ #include "scene/hair.h" #include "scene/mesh.h" #include "scene/object.h" +#include "scene/pointcloud.h" #include "util/algorithm.h" @@ -426,6 +427,32 @@ void BVHSpatialSplit::split_curve_primitive(const Hair *hair, } } +void BVHSpatialSplit::split_point_primitive(const PointCloud *pointcloud, + const Transform *tfm, + int prim_index, + int dim, + float pos, + BoundBox &left_bounds, + BoundBox &right_bounds) +{ + /* No real splitting support for points, assume they are small enough for it + * not to matter. */ + float3 point = pointcloud->get_points()[prim_index]; + + if (tfm != NULL) { + point = transform_point(tfm, point); + } + point = get_unaligned_point(point); + + if (point[dim] <= pos) { + left_bounds.grow(point); + } + + if (point[dim] >= pos) { + right_bounds.grow(point); + } +} + void BVHSpatialSplit::split_triangle_reference(const BVHReference &ref, const Mesh *mesh, int dim, @@ -453,6 +480,16 @@ void BVHSpatialSplit::split_curve_reference(const BVHReference &ref, right_bounds); } +void BVHSpatialSplit::split_point_reference(const BVHReference &ref, + const PointCloud *pointcloud, + int dim, + float pos, + BoundBox &left_bounds, + BoundBox &right_bounds) +{ + split_point_primitive(pointcloud, NULL, ref.prim_index(), dim, pos, left_bounds, right_bounds); +} + void BVHSpatialSplit::split_object_reference( const Object *object, int dim, float pos, BoundBox &left_bounds, BoundBox &right_bounds) { @@ -475,6 +512,13 @@ void BVHSpatialSplit::split_object_reference( } } } + else if (geom->geometry_type == Geometry::POINTCLOUD) { + PointCloud *pointcloud = static_cast<PointCloud *>(geom); + for (int point_idx = 0; point_idx < pointcloud->num_points(); ++point_idx) { + split_point_primitive( + pointcloud, &object->get_tfm(), point_idx, dim, pos, left_bounds, right_bounds); + } + } } void BVHSpatialSplit::split_reference(const BVHBuild &builder, @@ -499,6 +543,10 @@ void BVHSpatialSplit::split_reference(const BVHBuild &builder, Hair *hair = static_cast<Hair *>(ob->get_geometry()); split_curve_reference(ref, hair, dim, pos, left_bounds, right_bounds); } + else if (ref.prim_type() & PRIMITIVE_ALL_POINT) { + PointCloud *pointcloud = static_cast<PointCloud *>(ob->get_geometry()); + split_point_reference(ref, pointcloud, dim, pos, left_bounds, right_bounds); + } else { split_object_reference(ob, dim, pos, left_bounds, right_bounds); } diff --git a/intern/cycles/bvh/split.h b/intern/cycles/bvh/split.h index 2650a500ea9..92953c40040 100644 --- a/intern/cycles/bvh/split.h +++ b/intern/cycles/bvh/split.h @@ -26,6 +26,7 @@ CCL_NAMESPACE_BEGIN class BVHBuild; class Hair; class Mesh; +class PointCloud; struct Transform; /* Object Split */ @@ -123,6 +124,13 @@ class BVHSpatialSplit { float pos, BoundBox &left_bounds, BoundBox &right_bounds); + void split_point_primitive(const PointCloud *pointcloud, + const Transform *tfm, + int prim_index, + int dim, + float pos, + BoundBox &left_bounds, + BoundBox &right_bounds); /* Lower-level functions which calculates boundaries of left and right nodes * needed for spatial split. @@ -141,6 +149,12 @@ class BVHSpatialSplit { float pos, BoundBox &left_bounds, BoundBox &right_bounds); + void split_point_reference(const BVHReference &ref, + const PointCloud *pointcloud, + int dim, + float pos, + BoundBox &left_bounds, + BoundBox &right_bounds); void split_object_reference( const Object *object, int dim, float pos, BoundBox &left_bounds, BoundBox &right_bounds); diff --git a/intern/cycles/device/optix/device_impl.cpp b/intern/cycles/device/optix/device_impl.cpp index 11786f4a6e2..38cc3330ebd 100644 --- a/intern/cycles/device/optix/device_impl.cpp +++ b/intern/cycles/device/optix/device_impl.cpp @@ -28,6 +28,7 @@ # include "scene/mesh.h" # include "scene/object.h" # include "scene/pass.h" +# include "scene/pointcloud.h" # include "scene/scene.h" # include "util/debug.h" @@ -242,6 +243,9 @@ bool OptiXDevice::load_kernels(const uint kernel_features) else pipeline_options.usesPrimitiveTypeFlags |= OPTIX_PRIMITIVE_TYPE_FLAGS_CUSTOM; } + if (kernel_features & KERNEL_FEATURE_POINTCLOUD) { + pipeline_options.usesPrimitiveTypeFlags |= OPTIX_PRIMITIVE_TYPE_FLAGS_CUSTOM; + } /* Keep track of whether motion blur is enabled, so to enable/disable motion in BVH builds * This is necessary since objects may be reported to have motion if the Vector pass is @@ -372,6 +376,18 @@ bool OptiXDevice::load_kernels(const uint kernel_features) } } + /* Pointclouds */ + if (kernel_features & KERNEL_FEATURE_POINTCLOUD) { + group_descs[PG_HITD_POINTCLOUD] = group_descs[PG_HITD]; + group_descs[PG_HITD_POINTCLOUD].kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP; + group_descs[PG_HITD_POINTCLOUD].hitgroup.moduleIS = optix_module; + group_descs[PG_HITD_POINTCLOUD].hitgroup.entryFunctionNameIS = "__intersection__point"; + group_descs[PG_HITS_POINTCLOUD] = group_descs[PG_HITS]; + group_descs[PG_HITS_POINTCLOUD].kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP; + group_descs[PG_HITS_POINTCLOUD].hitgroup.moduleIS = optix_module; + group_descs[PG_HITS_POINTCLOUD].hitgroup.entryFunctionNameIS = "__intersection__point"; + } + if (kernel_features & (KERNEL_FEATURE_SUBSURFACE | KERNEL_FEATURE_NODE_RAYTRACE)) { /* Add hit group for local intersections. */ group_descs[PG_HITL].kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP; @@ -419,6 +435,10 @@ bool OptiXDevice::load_kernels(const uint kernel_features) stack_size[PG_HITD_MOTION].cssIS + stack_size[PG_HITD_MOTION].cssAH); trace_css = std::max(trace_css, stack_size[PG_HITS_MOTION].cssIS + stack_size[PG_HITS_MOTION].cssAH); + trace_css = std::max( + trace_css, stack_size[PG_HITD_POINTCLOUD].cssIS + stack_size[PG_HITD_POINTCLOUD].cssAH); + trace_css = std::max( + trace_css, stack_size[PG_HITS_POINTCLOUD].cssIS + stack_size[PG_HITS_POINTCLOUD].cssAH); OptixPipelineLinkOptions link_options = {}; link_options.maxTraceDepth = 1; @@ -444,6 +464,10 @@ bool OptiXDevice::load_kernels(const uint kernel_features) pipeline_groups.push_back(groups[PG_HITD_MOTION]); pipeline_groups.push_back(groups[PG_HITS_MOTION]); } + if (kernel_features & KERNEL_FEATURE_POINTCLOUD) { + pipeline_groups.push_back(groups[PG_HITD_POINTCLOUD]); + pipeline_groups.push_back(groups[PG_HITS_POINTCLOUD]); + } pipeline_groups.push_back(groups[PG_CALL_SVM_AO]); pipeline_groups.push_back(groups[PG_CALL_SVM_BEVEL]); @@ -483,6 +507,10 @@ bool OptiXDevice::load_kernels(const uint kernel_features) pipeline_groups.push_back(groups[PG_HITD_MOTION]); pipeline_groups.push_back(groups[PG_HITS_MOTION]); } + if (kernel_features & KERNEL_FEATURE_POINTCLOUD) { + pipeline_groups.push_back(groups[PG_HITD_POINTCLOUD]); + pipeline_groups.push_back(groups[PG_HITS_POINTCLOUD]); + } optix_assert(optixPipelineCreate(context, &pipeline_options, @@ -1378,6 +1406,86 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit) progress.set_error("Failed to build OptiX acceleration structure"); } } + else if (geom->geometry_type == Geometry::POINTCLOUD) { + /* Build BLAS for points primitives. */ + PointCloud *const pointcloud = static_cast<PointCloud *const>(geom); + const size_t num_points = pointcloud->num_points(); + if (num_points == 0) { + return; + } + + size_t num_motion_steps = 1; + Attribute *motion_points = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); + if (motion_blur && pointcloud->get_use_motion_blur() && motion_points) { + num_motion_steps = pointcloud->get_motion_steps(); + } + + device_vector<OptixAabb> aabb_data(this, "optix temp aabb data", MEM_READ_ONLY); + aabb_data.alloc(num_points * num_motion_steps); + + /* Get AABBs for each motion step. */ + for (size_t step = 0; step < num_motion_steps; ++step) { + /* The center step for motion vertices is not stored in the attribute. */ + const float3 *points = pointcloud->get_points().data(); + const float *radius = pointcloud->get_radius().data(); + size_t center_step = (num_motion_steps - 1) / 2; + if (step != center_step) { + size_t attr_offset = (step > center_step) ? step - 1 : step; + /* Technically this is a float4 array, but sizeof(float3) == sizeof(float4). */ + points = motion_points->data_float3() + attr_offset * num_points; + } + + for (size_t i = 0; i < num_points; ++i) { + const PointCloud::Point point = pointcloud->get_point(i); + BoundBox bounds = BoundBox::empty; + point.bounds_grow(points, radius, bounds); + + const size_t index = step * num_points + i; + aabb_data[index].minX = bounds.min.x; + aabb_data[index].minY = bounds.min.y; + aabb_data[index].minZ = bounds.min.z; + aabb_data[index].maxX = bounds.max.x; + aabb_data[index].maxY = bounds.max.y; + aabb_data[index].maxZ = bounds.max.z; + } + } + + /* Upload AABB data to GPU. */ + aabb_data.copy_to_device(); + + vector<device_ptr> aabb_ptrs; + aabb_ptrs.reserve(num_motion_steps); + for (size_t step = 0; step < num_motion_steps; ++step) { + aabb_ptrs.push_back(aabb_data.device_pointer + step * num_points * sizeof(OptixAabb)); + } + + /* Disable visibility test any-hit program, since it is already checked during + * intersection. Those trace calls that require anyhit can force it with a ray flag. + * For those, force a single any-hit call, so shadow record-all behavior works correctly. */ + unsigned int build_flags = OPTIX_GEOMETRY_FLAG_DISABLE_ANYHIT | + OPTIX_GEOMETRY_FLAG_REQUIRE_SINGLE_ANYHIT_CALL; + OptixBuildInput build_input = {}; + build_input.type = OPTIX_BUILD_INPUT_TYPE_CUSTOM_PRIMITIVES; +# if OPTIX_ABI_VERSION < 23 + build_input.aabbArray.aabbBuffers = (CUdeviceptr *)aabb_ptrs.data(); + build_input.aabbArray.numPrimitives = num_points; + build_input.aabbArray.strideInBytes = sizeof(OptixAabb); + build_input.aabbArray.flags = &build_flags; + build_input.aabbArray.numSbtRecords = 1; + build_input.aabbArray.primitiveIndexOffset = pointcloud->prim_offset; +# else + build_input.customPrimitiveArray.aabbBuffers = (CUdeviceptr *)aabb_ptrs.data(); + build_input.customPrimitiveArray.numPrimitives = num_points; + build_input.customPrimitiveArray.strideInBytes = sizeof(OptixAabb); + build_input.customPrimitiveArray.flags = &build_flags; + build_input.customPrimitiveArray.numSbtRecords = 1; + build_input.customPrimitiveArray.primitiveIndexOffset = pointcloud->prim_offset; +# endif + + if (!build_optix_bvh(bvh_optix, operation, build_input, num_motion_steps)) { + progress.set_error("Failed to build OptiX acceleration structure"); + } + } } else { unsigned int num_instances = 0; @@ -1461,12 +1569,22 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit) instance.sbtOffset = PG_HITD_MOTION - PG_HITD; } } + else if (ob->get_geometry()->geometry_type == Geometry::POINTCLOUD) { + /* Use the hit group that has an intersection program for point clouds. */ + instance.sbtOffset = PG_HITD_POINTCLOUD - PG_HITD; + + /* Also skip point clouds in local trace calls. */ + instance.visibilityMask |= 4; + } + # if OPTIX_ABI_VERSION < 55 /* Cannot disable any-hit program for thick curves, since it needs to filter out end-caps. */ else # endif { - /* Can disable __anyhit__kernel_optix_visibility_test by default. + /* Can disable __anyhit__kernel_optix_visibility_test by default (except for thick curves, + * since it needs to filter out end-caps there). + * It is enabled where necessary (visibility mask exceeds 8 bits or the other any-hit * programs like __anyhit__kernel_optix_shadow_all_hit) via OPTIX_RAY_FLAG_ENFORCE_ANYHIT. */ diff --git a/intern/cycles/device/optix/device_impl.h b/intern/cycles/device/optix/device_impl.h index cf4afb46280..25073c60e69 100644 --- a/intern/cycles/device/optix/device_impl.h +++ b/intern/cycles/device/optix/device_impl.h @@ -44,6 +44,8 @@ enum { PG_HITV, /* __VOLUME__ hit group. */ PG_HITD_MOTION, PG_HITS_MOTION, + PG_HITD_POINTCLOUD, + PG_HITS_POINTCLOUD, PG_CALL_SVM_AO, PG_CALL_SVM_BEVEL, NUM_PROGRAM_GROUPS @@ -52,9 +54,9 @@ enum { static const int MISS_PROGRAM_GROUP_OFFSET = PG_MISS; static const int NUM_MIS_PROGRAM_GROUPS = 1; static const int HIT_PROGAM_GROUP_OFFSET = PG_HITD; -static const int NUM_HIT_PROGRAM_GROUPS = 6; +static const int NUM_HIT_PROGRAM_GROUPS = 8; static const int CALLABLE_PROGRAM_GROUPS_BASE = PG_CALL_SVM_AO; -static const int NUM_CALLABLE_PROGRAM_GROUPS = 3; +static const int NUM_CALLABLE_PROGRAM_GROUPS = 2; /* List of OptiX pipelines. */ enum { PIP_SHADE_RAYTRACE, PIP_INTERSECT, NUM_PIPELINES }; diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt index 674eb702814..51158b86c5a 100644 --- a/intern/cycles/kernel/CMakeLists.txt +++ b/intern/cycles/kernel/CMakeLists.txt @@ -179,11 +179,14 @@ set(SRC_KERNEL_GEOM_HEADERS geom/curve.h geom/curve_intersect.h geom/motion_curve.h + geom/motion_point.h geom/motion_triangle.h geom/motion_triangle_intersect.h geom/motion_triangle_shader.h geom/object.h geom/patch.h + geom/point.h + geom/point_intersect.h geom/primitive.h geom/shader_data.h geom/subd_triangle.h diff --git a/intern/cycles/kernel/bvh/bvh.h b/intern/cycles/kernel/bvh/bvh.h index 37cfac25060..67804fb1d0d 100644 --- a/intern/cycles/kernel/bvh/bvh.h +++ b/intern/cycles/kernel/bvh/bvh.h @@ -49,24 +49,24 @@ CCL_NAMESPACE_BEGIN # include "kernel/bvh/nodes.h" # define BVH_FUNCTION_NAME bvh_intersect -# define BVH_FUNCTION_FEATURES 0 +# define BVH_FUNCTION_FEATURES BVH_POINTCLOUD # include "kernel/bvh/traversal.h" # if defined(__HAIR__) # define BVH_FUNCTION_NAME bvh_intersect_hair -# define BVH_FUNCTION_FEATURES BVH_HAIR +# define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_POINTCLOUD # include "kernel/bvh/traversal.h" # endif # if defined(__OBJECT_MOTION__) # define BVH_FUNCTION_NAME bvh_intersect_motion -# define BVH_FUNCTION_FEATURES BVH_MOTION +# define BVH_FUNCTION_FEATURES BVH_MOTION | BVH_POINTCLOUD # include "kernel/bvh/traversal.h" # endif # if defined(__HAIR__) && defined(__OBJECT_MOTION__) # define BVH_FUNCTION_NAME bvh_intersect_hair_motion -# define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_MOTION +# define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_MOTION | BVH_POINTCLOUD # include "kernel/bvh/traversal.h" # endif @@ -102,26 +102,27 @@ CCL_NAMESPACE_BEGIN # if defined(__SHADOW_RECORD_ALL__) # define BVH_FUNCTION_NAME bvh_intersect_shadow_all -# define BVH_FUNCTION_FEATURES 0 +# define BVH_FUNCTION_FEATURES BVH_POINTCLOUD # include "kernel/bvh/shadow_all.h" # if defined(__HAIR__) # define BVH_FUNCTION_NAME bvh_intersect_shadow_all_hair -# define BVH_FUNCTION_FEATURES BVH_HAIR +# define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_POINTCLOUD # include "kernel/bvh/shadow_all.h" # endif # if defined(__OBJECT_MOTION__) # define BVH_FUNCTION_NAME bvh_intersect_shadow_all_motion -# define BVH_FUNCTION_FEATURES BVH_MOTION +# define BVH_FUNCTION_FEATURES BVH_MOTION | BVH_POINTCLOUD # include "kernel/bvh/shadow_all.h" # endif # if defined(__HAIR__) && defined(__OBJECT_MOTION__) # define BVH_FUNCTION_NAME bvh_intersect_shadow_all_hair_motion -# define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_MOTION +# define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_MOTION | BVH_POINTCLOUD # include "kernel/bvh/shadow_all.h" # endif + # endif /* __SHADOW_RECORD_ALL__ */ /* Record all intersections - Volume BVH traversal. */ diff --git a/intern/cycles/kernel/bvh/shadow_all.h b/intern/cycles/kernel/bvh/shadow_all.h index 049c6a03fe0..caca85aac1a 100644 --- a/intern/cycles/kernel/bvh/shadow_all.h +++ b/intern/cycles/kernel/bvh/shadow_all.h @@ -28,6 +28,7 @@ * without new features slowing things down. * * BVH_HAIR: hair curve rendering + * BVH_POINTCLOUD: point cloud rendering * BVH_MOTION: motion blur rendering */ @@ -199,6 +200,34 @@ ccl_device_inline break; } #endif +#if BVH_FEATURE(BVH_POINTCLOUD) + case PRIMITIVE_POINT: + case PRIMITIVE_MOTION_POINT: { + if ((type & PRIMITIVE_ALL_MOTION) && kernel_data.bvh.use_bvh_steps) { + const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr); + if (ray->time < prim_time.x || ray->time > prim_time.y) { + hit = false; + break; + } + } + + const int point_object = (object == OBJECT_NONE) ? + kernel_tex_fetch(__prim_object, prim_addr) : + object; + const int point_prim = kernel_tex_fetch(__prim_index, prim_addr); + const int point_type = kernel_tex_fetch(__prim_type, prim_addr); + hit = point_intersect(kg, + &isect, + P, + dir, + t_max_current, + point_object, + point_prim, + ray->time, + point_type); + break; + } +#endif /* BVH_FEATURE(BVH_POINTCLOUD) */ default: { hit = false; break; diff --git a/intern/cycles/kernel/bvh/traversal.h b/intern/cycles/kernel/bvh/traversal.h index 1c17ebf767f..180f19d11c5 100644 --- a/intern/cycles/kernel/bvh/traversal.h +++ b/intern/cycles/kernel/bvh/traversal.h @@ -28,6 +28,7 @@ * without new features slowing things down. * * BVH_HAIR: hair curve rendering + * BVH_POINTCLOUD: point cloud rendering * BVH_MOTION: motion blur rendering */ @@ -188,6 +189,33 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg, break; } #endif /* BVH_FEATURE(BVH_HAIR) */ +#if BVH_FEATURE(BVH_POINTCLOUD) + case PRIMITIVE_POINT: + case PRIMITIVE_MOTION_POINT: { + for (; prim_addr < prim_addr2; prim_addr++) { + if ((type & PRIMITIVE_ALL_MOTION) && kernel_data.bvh.use_bvh_steps) { + const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr); + if (ray->time < prim_time.x || ray->time > prim_time.y) { + continue; + } + } + + const int point_object = (object == OBJECT_NONE) ? + kernel_tex_fetch(__prim_object, prim_addr) : + object; + const int point_prim = kernel_tex_fetch(__prim_index, prim_addr); + const int point_type = kernel_tex_fetch(__prim_type, prim_addr); + const bool hit = point_intersect( + kg, isect, P, dir, isect->t, point_object, point_prim, ray->time, point_type); + if (hit) { + /* shadow ray early termination */ + if (visibility & PATH_RAY_SHADOW_OPAQUE) + return true; + } + } + break; + } +#endif /* BVH_FEATURE(BVH_POINTCLOUD) */ } } else { diff --git a/intern/cycles/kernel/bvh/types.h b/intern/cycles/kernel/bvh/types.h index 6039e707fc3..f16f43333f8 100644 --- a/intern/cycles/kernel/bvh/types.h +++ b/intern/cycles/kernel/bvh/types.h @@ -34,6 +34,7 @@ CCL_NAMESPACE_BEGIN #define BVH_MOTION 1 #define BVH_HAIR 2 +#define BVH_POINTCLOUD 4 #define BVH_NAME_JOIN(x, y) x##_##y #define BVH_NAME_EVAL(x, y) BVH_NAME_JOIN(x, y) diff --git a/intern/cycles/kernel/bvh/util.h b/intern/cycles/kernel/bvh/util.h index 26ba136dd79..57593e42a88 100644 --- a/intern/cycles/kernel/bvh/util.h +++ b/intern/cycles/kernel/bvh/util.h @@ -118,14 +118,16 @@ ccl_device_forceinline int intersection_get_shader_flags(KernelGlobals kg, { int shader = 0; -#ifdef __HAIR__ - if (type & PRIMITIVE_ALL_TRIANGLE) -#endif - { + if (type & PRIMITIVE_ALL_TRIANGLE) { shader = kernel_tex_fetch(__tri_shader, prim); } +#ifdef __POINTCLOUD__ + else if (type & PRIMITIVE_ALL_POINT) { + shader = kernel_tex_fetch(__points_shader, prim); + } +#endif #ifdef __HAIR__ - else { + else if (type & PRIMITIVE_ALL_CURVE) { shader = kernel_tex_fetch(__curves, prim).shader_id; } #endif @@ -139,14 +141,16 @@ ccl_device_forceinline int intersection_get_shader_from_isect_prim(KernelGlobals { int shader = 0; -#ifdef __HAIR__ - if (isect_type & PRIMITIVE_ALL_TRIANGLE) -#endif - { + if (isect_type & PRIMITIVE_ALL_TRIANGLE) { shader = kernel_tex_fetch(__tri_shader, prim); } +#ifdef __POINTCLOUD__ + else if (isect_type & PRIMITIVE_ALL_POINT) { + shader = kernel_tex_fetch(__points_shader, prim); + } +#endif #ifdef __HAIR__ - else { + else if (isect_type & PRIMITIVE_ALL_CURVE) { shader = kernel_tex_fetch(__curves, prim).shader_id; } #endif diff --git a/intern/cycles/kernel/device/optix/kernel.cu b/intern/cycles/kernel/device/optix/kernel.cu index c9a48315544..c639dc87f35 100644 --- a/intern/cycles/kernel/device/optix/kernel.cu +++ b/intern/cycles/kernel/device/optix/kernel.cu @@ -97,9 +97,9 @@ extern "C" __global__ void __miss__kernel_optix_miss() extern "C" __global__ void __anyhit__kernel_optix_local_hit() { -#ifdef __HAIR__ +#if defined(__HAIR__) || defined(__POINTCLOUD__) if (!optixIsTriangleHit()) { - /* Ignore curves. */ + /* Ignore curves and points. */ return optixIgnoreIntersection(); } #endif @@ -194,7 +194,7 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() type = kernel_tex_fetch(__objects, object).primitive_type; } # ifdef __HAIR__ - else { + else if (optixGetHitKind() & PRIMITIVE_ALL_CURVE) { u = __uint_as_float(optixGetAttribute_0()); v = __uint_as_float(optixGetAttribute_1()); @@ -210,6 +210,11 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() # endif } # endif + else { + type = kernel_tex_fetch(__objects, object).primitive_type; + u = 0.0f; + v = 0.0f; + } # ifndef __TRANSPARENT_SHADOWS__ /* No transparent shadows support compiled in, make opaque. */ @@ -291,7 +296,7 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() extern "C" __global__ void __anyhit__kernel_optix_volume_test() { -#ifdef __HAIR__ +#if defined(__HAIR__) || defined(__POINTCLOUD__) if (!optixIsTriangleHit()) { /* Ignore curves. */ return optixIgnoreIntersection(); @@ -315,7 +320,7 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test() { #ifdef __HAIR__ # if OPTIX_ABI_VERSION < 55 - if (!optixIsTriangleHit()) { + if (optixGetHitKind() & PRIMITIVE_ALL_CURVE) { /* Filter out curve endcaps. */ const float u = __uint_as_float(optixGetAttribute_0()); if (u == 0.0f || u == 1.0f) { @@ -354,13 +359,19 @@ extern "C" __global__ void __closesthit__kernel_optix_hit() optixSetPayload_3(prim); optixSetPayload_5(kernel_tex_fetch(__objects, object).primitive_type); } - else { + else if (optixGetHitKind() & PRIMITIVE_ALL_CURVE) { const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim); optixSetPayload_1(optixGetAttribute_0()); /* Same as 'optixGetCurveParameter()' */ optixSetPayload_2(optixGetAttribute_1()); optixSetPayload_3(segment.prim); optixSetPayload_5(segment.type); } + else { + optixSetPayload_1(0); + optixSetPayload_2(0); + optixSetPayload_3(prim); + optixSetPayload_5(kernel_tex_fetch(__objects, object).primitive_type); + } } #ifdef __HAIR__ @@ -411,4 +422,45 @@ extern "C" __global__ void __intersection__curve_ribbon() optix_intersection_curve(prim, type); } } + +#endif + +#ifdef __POINTCLOUD__ +extern "C" __global__ void __intersection__point() +{ + const int prim = optixGetPrimitiveIndex(); + const int object = get_object_id(); + const int type = kernel_tex_fetch(__objects, object).primitive_type; + +# ifdef __VISIBILITY_FLAG__ + const uint visibility = optixGetPayload_4(); + if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) { + return; + } +# endif + + float3 P = optixGetObjectRayOrigin(); + float3 dir = optixGetObjectRayDirection(); + + /* The direction is not normalized by default, the point intersection routine expects that. */ + float len; + dir = normalize_len(dir, &len); + +# ifdef __OBJECT_MOTION__ + const float time = optixGetRayTime(); +# else + const float time = 0.0f; +# endif + + Intersection isect; + isect.t = optixGetRayTmax(); + /* Transform maximum distance into object space. */ + if (isect.t != FLT_MAX) { + isect.t *= len; + } + + if (point_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) { + optixReportIntersection(isect.t / len, type & PRIMITIVE_ALL); + } +} #endif diff --git a/intern/cycles/kernel/geom/geom.h b/intern/cycles/kernel/geom/geom.h index 9d023375a35..efc296c8d39 100644 --- a/intern/cycles/kernel/geom/geom.h +++ b/intern/cycles/kernel/geom/geom.h @@ -29,6 +29,9 @@ #include "kernel/geom/motion_triangle_intersect.h" #include "kernel/geom/motion_triangle_shader.h" #include "kernel/geom/motion_curve.h" +#include "kernel/geom/motion_point.h" +#include "kernel/geom/point.h" +#include "kernel/geom/point_intersect.h" #include "kernel/geom/curve.h" #include "kernel/geom/curve_intersect.h" #include "kernel/geom/volume.h" diff --git a/intern/cycles/kernel/geom/motion_point.h b/intern/cycles/kernel/geom/motion_point.h new file mode 100644 index 00000000000..d7ffc80c045 --- /dev/null +++ b/intern/cycles/kernel/geom/motion_point.h @@ -0,0 +1,74 @@ +/* + * 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. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +CCL_NAMESPACE_BEGIN + +/* Motion Point Primitive + * + * These are stored as regular points, plus extra positions and radii at times + * other than the frame center. Computing the point at a given ray time is + * a matter of interpolation of the two steps between which the ray time lies. + * + * The extra points are stored as ATTR_STD_MOTION_VERTEX_POSITION. + */ + +#ifdef __POINTCLOUD__ + +ccl_device_inline float4 +motion_point_for_step(KernelGlobals kg, int offset, int numkeys, int numsteps, int step, int prim) +{ + if (step == numsteps) { + /* center step: regular key location */ + return kernel_tex_fetch(__points, prim); + } + else { + /* center step is not stored in this array */ + if (step > numsteps) + step--; + + offset += step * numkeys; + + return kernel_tex_fetch(__attributes_float4, offset + prim); + } +} + +/* return 2 point key locations */ +ccl_device_inline float4 motion_point(KernelGlobals kg, int object, int prim, float time) +{ + /* get motion info */ + int numsteps, numkeys; + object_motion_info(kg, object, &numsteps, NULL, &numkeys); + + /* figure out which steps we need to fetch and their interpolation factor */ + int maxstep = numsteps * 2; + int step = min((int)(time * maxstep), maxstep - 1); + float t = time * maxstep - step; + + /* find attribute */ + int offset = intersection_find_attribute(kg, object, ATTR_STD_MOTION_VERTEX_POSITION); + kernel_assert(offset != ATTR_STD_NOT_FOUND); + + /* fetch key coordinates */ + float4 point = motion_point_for_step(kg, offset, numkeys, numsteps, step, prim); + float4 next_point = motion_point_for_step(kg, offset, numkeys, numsteps, step + 1, prim); + + /* interpolate between steps */ + return (1.0f - t) * point + t * next_point; +} + +#endif + +CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/geom/point.h b/intern/cycles/kernel/geom/point.h new file mode 100644 index 00000000000..021135b76fb --- /dev/null +++ b/intern/cycles/kernel/geom/point.h @@ -0,0 +1,124 @@ +/* + * 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. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +CCL_NAMESPACE_BEGIN + +/* Point Primitive + * + * Point primitive for rendering point clouds. + */ + +#ifdef __POINTCLOUD__ + +/* Reading attributes on various point elements */ + +ccl_device float point_attribute_float(KernelGlobals kg, + ccl_private const ShaderData *sd, + const AttributeDescriptor desc, + ccl_private float *dx, + ccl_private float *dy) +{ +# ifdef __RAY_DIFFERENTIALS__ + if (dx) + *dx = 0.0f; + if (dy) + *dy = 0.0f; +# endif + + if (desc.element == ATTR_ELEMENT_VERTEX) { + return kernel_tex_fetch(__attributes_float, desc.offset + sd->prim); + } + else { + return 0.0f; + } +} + +ccl_device float2 point_attribute_float2( + KernelGlobals kg, const ShaderData *sd, const AttributeDescriptor desc, float2 *dx, float2 *dy) +{ +# ifdef __RAY_DIFFERENTIALS__ + if (dx) + *dx = make_float2(0.0f, 0.0f); + if (dy) + *dy = make_float2(0.0f, 0.0f); +# endif + + if (desc.element == ATTR_ELEMENT_VERTEX) { + return kernel_tex_fetch(__attributes_float2, desc.offset + sd->prim); + } + else { + return make_float2(0.0f, 0.0f); + } +} + +ccl_device float3 point_attribute_float3( + KernelGlobals kg, const ShaderData *sd, const AttributeDescriptor desc, float3 *dx, float3 *dy) +{ +# ifdef __RAY_DIFFERENTIALS__ + if (dx) + *dx = make_float3(0.0f, 0.0f, 0.0f); + if (dy) + *dy = make_float3(0.0f, 0.0f, 0.0f); +# endif + + if (desc.element == ATTR_ELEMENT_VERTEX) { + return float4_to_float3(kernel_tex_fetch(__attributes_float4, desc.offset + sd->prim)); + } + else { + return make_float3(0.0f, 0.0f, 0.0f); + } +} + +ccl_device float4 point_attribute_float4( + KernelGlobals kg, const ShaderData *sd, const AttributeDescriptor desc, float4 *dx, float4 *dy) +{ +# ifdef __RAY_DIFFERENTIALS__ + if (dx) + *dx = make_float4(0.0f, 0.0f, 0.0f, 0.0f); + if (dy) + *dy = make_float4(0.0f, 0.0f, 0.0f, 0.0f); +# endif + + if (desc.element == ATTR_ELEMENT_VERTEX) { + return kernel_tex_fetch(__attributes_float4, desc.offset + sd->prim); + } + else { + return make_float4(0.0f, 0.0f, 0.0f, 0.0f); + } +} + +/* Point radius */ + +ccl_device float point_radius(KernelGlobals kg, ccl_private const ShaderData *sd) +{ + if (sd->type & PRIMITIVE_ALL_POINT) { + return kernel_tex_fetch(__points, sd->prim).w; + } + + return 0.0f; +} + +/* Point location for motion pass, linear interpolation between keys and + * ignoring radius because we do the same for the motion keys */ + +ccl_device float3 point_motion_center_location(KernelGlobals kg, ccl_private const ShaderData *sd) +{ + return float4_to_float3(kernel_tex_fetch(__points, sd->prim)); +} + +#endif /* __POINTCLOUD__ */ + +CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/geom/point_intersect.h b/intern/cycles/kernel/geom/point_intersect.h new file mode 100644 index 00000000000..da0b31f9a8c --- /dev/null +++ b/intern/cycles/kernel/geom/point_intersect.h @@ -0,0 +1,133 @@ +/* + * 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. + * See the License for the specific language governing permissions and + * limitations under the License. + * + * Based on Embree code, copyright 2009-2020 Intel Corporation. + */ + +#pragma once + +CCL_NAMESPACE_BEGIN + +/* Point primitive intersection functions. */ + +#ifdef __POINTCLOUD__ + +ccl_device_forceinline bool point_intersect_test( + const float4 point, const float3 P, const float3 dir, const float tmax, float *t) +{ + const float3 center = float4_to_float3(point); + const float radius = point.w; + + const float rd2 = 1.0f / dot(dir, dir); + + const float3 c0 = center - P; + const float projC0 = dot(c0, dir) * rd2; + const float3 perp = c0 - projC0 * dir; + const float l2 = dot(perp, perp); + const float r2 = radius * radius; + if (!(l2 <= r2)) { + return false; + } + + const float td = sqrt((r2 - l2) * rd2); + const float t_front = projC0 - td; + const bool valid_front = (0.0f <= t_front) & (t_front <= tmax); + + /* Always backface culling for now. */ +# if 0 + const float t_back = projC0 + td; + const bool valid_back = (0.0f <= t_back) & (t_back <= tmax); + + /* check if there is a first hit */ + const bool valid_first = valid_front | valid_back; + if (!valid_first) { + return false; + } + + *t = (valid_front) ? t_front : t_back; + return true; +# else + if (!valid_front) { + return false; + } + *t = t_front; + return true; +# endif +} + +ccl_device_forceinline bool point_intersect(KernelGlobals kg, + ccl_private Intersection *isect, + const float3 P, + const float3 dir, + const float tmax, + const int object, + const int prim, + const float time, + const int type) +{ + const float4 point = (type & PRIMITIVE_ALL_MOTION) ? motion_point(kg, object, prim, time) : + kernel_tex_fetch(__points, prim); + + if (!point_intersect_test(point, P, dir, tmax, &isect->t)) { + return false; + } + + isect->prim = prim; + isect->object = object; + isect->type = type; + isect->u = 0.0f; + isect->v = 0.0f; + return true; +} + +ccl_device_inline void point_shader_setup(KernelGlobals kg, + ccl_private ShaderData *sd, + ccl_private const Intersection *isect, + const Ray *ray) +{ + sd->shader = kernel_tex_fetch(__points_shader, isect->prim); + sd->P = ray->P + ray->D * isect->t; + + /* Texture coordinates, zero for now. */ +# ifdef __UV__ + sd->u = isect->u; + sd->v = isect->v; +# endif + + /* Computer point center for normal. */ + float3 center = float4_to_float3((isect->type & PRIMITIVE_ALL_MOTION) ? + motion_point(kg, sd->object, sd->prim, sd->time) : + kernel_tex_fetch(__points, sd->prim)); + + if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) { + const Transform tfm = object_get_transform(kg, sd); + +# ifndef __KERNEL_OPTIX__ + center = transform_point(&tfm, center); +# endif + } + + /* Normal */ + sd->Ng = normalize(sd->P - center); + sd->N = sd->Ng; + +# ifdef __DPDU__ + /* dPdu/dPdv */ + sd->dPdu = make_float3(0.0f, 0.0f, 0.0f); + sd->dPdv = make_float3(0.0f, 0.0f, 0.0f); +# endif +} + +#endif + +CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/geom/primitive.h b/intern/cycles/kernel/geom/primitive.h index 6d7b550d82f..9416385638c 100644 --- a/intern/cycles/kernel/geom/primitive.h +++ b/intern/cycles/kernel/geom/primitive.h @@ -48,6 +48,11 @@ ccl_device_inline float primitive_surface_attribute_float(KernelGlobals kg, return curve_attribute_float(kg, sd, desc, dx, dy); } #endif +#ifdef __POINTCLOUD__ + else if (sd->type & PRIMITIVE_ALL_POINT) { + return point_attribute_float(kg, sd, desc, dx, dy); + } +#endif else { if (dx) *dx = 0.0f; @@ -74,6 +79,11 @@ ccl_device_inline float2 primitive_surface_attribute_float2(KernelGlobals kg, return curve_attribute_float2(kg, sd, desc, dx, dy); } #endif +#ifdef __POINTCLOUD__ + else if (sd->type & PRIMITIVE_ALL_POINT) { + return point_attribute_float2(kg, sd, desc, dx, dy); + } +#endif else { if (dx) *dx = make_float2(0.0f, 0.0f); @@ -100,6 +110,11 @@ ccl_device_inline float3 primitive_surface_attribute_float3(KernelGlobals kg, return curve_attribute_float3(kg, sd, desc, dx, dy); } #endif +#ifdef __POINTCLOUD__ + else if (sd->type & PRIMITIVE_ALL_POINT) { + return point_attribute_float3(kg, sd, desc, dx, dy); + } +#endif else { if (dx) *dx = make_float3(0.0f, 0.0f, 0.0f); @@ -126,6 +141,11 @@ ccl_device_forceinline float4 primitive_surface_attribute_float4(KernelGlobals k return curve_attribute_float4(kg, sd, desc, dx, dy); } #endif +#ifdef __POINTCLOUD__ + else if (sd->type & PRIMITIVE_ALL_POINT) { + return point_attribute_float4(kg, sd, desc, dx, dy); + } +#endif else { if (dx) *dx = make_float4(0.0f, 0.0f, 0.0f, 0.0f); @@ -225,8 +245,8 @@ ccl_device bool primitive_ptex(KernelGlobals kg, ccl_device float3 primitive_tangent(KernelGlobals kg, ccl_private ShaderData *sd) { -#ifdef __HAIR__ - if (sd->type & PRIMITIVE_ALL_CURVE) +#if defined(__HAIR__) || defined(__POINTCLOUD__) + if (sd->type & (PRIMITIVE_ALL_CURVE | PRIMITIVE_ALL_POINT)) # ifdef __DPDU__ return normalize(sd->dPdu); # else @@ -261,10 +281,21 @@ ccl_device_inline float4 primitive_motion_vector(KernelGlobals kg, /* center position */ float3 center; -#ifdef __HAIR__ - bool is_curve_primitive = sd->type & PRIMITIVE_ALL_CURVE; - if (is_curve_primitive) { - center = curve_motion_center_location(kg, sd); +#if defined(__HAIR__) || defined(__POINTCLOUD__) + bool is_curve_or_point = sd->type & (PRIMITIVE_ALL_CURVE | PRIMITIVE_ALL_POINT); + if (is_curve_or_point) { + center = make_float3(0.0f, 0.0f, 0.0f); + + if (sd->type & PRIMITIVE_ALL_CURVE) { +# if defined(__HAIR__) + center = curve_motion_center_location(kg, sd); +# endif + } + else if (sd->type & PRIMITIVE_ALL_POINT) { +# if defined(__POINTCLOUD__) + center = point_motion_center_location(kg, sd); +# endif + } if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) { object_position_transform(kg, sd, ¢er); @@ -272,7 +303,9 @@ ccl_device_inline float4 primitive_motion_vector(KernelGlobals kg, } else #endif + { center = sd->P; + } float3 motion_pre = center, motion_post = center; @@ -284,8 +317,8 @@ ccl_device_inline float4 primitive_motion_vector(KernelGlobals kg, int numverts, numkeys; object_motion_info(kg, sd->object, NULL, &numverts, &numkeys); -#ifdef __HAIR__ - if (is_curve_primitive) { +#if defined(__HAIR__) || defined(__POINTCLOUD__) + if (is_curve_or_point) { motion_pre = float4_to_float3(curve_attribute_float4(kg, sd, desc, NULL, NULL)); desc.offset += numkeys; motion_post = float4_to_float3(curve_attribute_float4(kg, sd, desc, NULL, NULL)); diff --git a/intern/cycles/kernel/geom/shader_data.h b/intern/cycles/kernel/geom/shader_data.h index 46bda2b656c..d3932545ef6 100644 --- a/intern/cycles/kernel/geom/shader_data.h +++ b/intern/cycles/kernel/geom/shader_data.h @@ -75,6 +75,13 @@ ccl_device_inline void shader_setup_from_ray(KernelGlobals kg, } else #endif +#ifdef __POINTCLOUD__ + if (sd->type & PRIMITIVE_ALL_POINT) { + /* point */ + point_shader_setup(kg, sd, isect, ray); + } + else +#endif if (sd->type & PRIMITIVE_TRIANGLE) { /* static triangle */ float3 Ng = triangle_normal(kg, sd); diff --git a/intern/cycles/kernel/osl/services.cpp b/intern/cycles/kernel/osl/services.cpp index 389c1e2b746..4007005dee7 100644 --- a/intern/cycles/kernel/osl/services.cpp +++ b/intern/cycles/kernel/osl/services.cpp @@ -28,6 +28,7 @@ #include "scene/colorspace.h" #include "scene/mesh.h" #include "scene/object.h" +#include "scene/pointcloud.h" #include "scene/scene.h" #include "kernel/osl/closures.h" @@ -113,6 +114,8 @@ ustring OSLRenderServices::u_curve_thickness("geom:curve_thickness"); ustring OSLRenderServices::u_curve_length("geom:curve_length"); ustring OSLRenderServices::u_curve_tangent_normal("geom:curve_tangent_normal"); ustring OSLRenderServices::u_curve_random("geom:curve_random"); +ustring OSLRenderServices::u_is_point("geom:is_point"); +ustring OSLRenderServices::u_point_radius("geom:point_radius"); ustring OSLRenderServices::u_normal_map_normal("geom:normal_map_normal"); ustring OSLRenderServices::u_path_ray_length("path:ray_length"); ustring OSLRenderServices::u_path_ray_depth("path:ray_depth"); @@ -994,6 +997,15 @@ bool OSLRenderServices::get_object_standard_attribute(const KernelGlobalsCPU *kg float3 f = curve_tangent_normal(kg, sd); return set_attribute_float3(f, type, derivatives, val); } + /* point attributes */ + else if (name == u_is_point) { + float f = (sd->type & PRIMITIVE_ALL_POINT) != 0; + return set_attribute_float(f, type, derivatives, val); + } + else if (name == u_point_radius) { + float f = point_radius(kg, sd); + return set_attribute_float(f, type, derivatives, val); + } else if (name == u_normal_map_normal) { if (sd->type & PRIMITIVE_ALL_TRIANGLE) { float3 f = triangle_smooth_normal_unnormalized(kg, sd, sd->Ng, sd->prim, sd->u, sd->v); diff --git a/intern/cycles/kernel/osl/services.h b/intern/cycles/kernel/osl/services.h index d9f57c642ad..9526c92b8fb 100644 --- a/intern/cycles/kernel/osl/services.h +++ b/intern/cycles/kernel/osl/services.h @@ -297,6 +297,8 @@ class OSLRenderServices : public OSL::RendererServices { static ustring u_curve_length; static ustring u_curve_tangent_normal; static ustring u_curve_random; + static ustring u_is_point; + static ustring u_point_radius; static ustring u_normal_map_normal; static ustring u_path_ray_length; static ustring u_path_ray_depth; diff --git a/intern/cycles/kernel/svm/wireframe.h b/intern/cycles/kernel/svm/wireframe.h index 530a9601bce..645dc59f22e 100644 --- a/intern/cycles/kernel/svm/wireframe.h +++ b/intern/cycles/kernel/svm/wireframe.h @@ -42,7 +42,7 @@ ccl_device_inline float wireframe(KernelGlobals kg, int pixel_size, ccl_private float3 *P) { -#ifdef __HAIR__ +#if defined(__HAIR__) || defined(__POINTCLOUD__) if (sd->prim != PRIM_NONE && sd->type & PRIMITIVE_ALL_TRIANGLE) #else if (sd->prim != PRIM_NONE) diff --git a/intern/cycles/kernel/textures.h b/intern/cycles/kernel/textures.h index 58edb239007..06074a1eca1 100644 --- a/intern/cycles/kernel/textures.h +++ b/intern/cycles/kernel/textures.h @@ -55,6 +55,10 @@ KERNEL_TEX(KernelCurveSegment, __curve_segments) /* patches */ KERNEL_TEX(uint, __patches) +/* pointclouds */ +KERNEL_TEX(float4, __points) +KERNEL_TEX(uint, __points_shader) + /* attributes */ KERNEL_TEX(uint4, __attributes_map) KERNEL_TEX(float, __attributes_float) diff --git a/intern/cycles/kernel/types.h b/intern/cycles/kernel/types.h index 34f909a06d9..855cc97edbf 100644 --- a/intern/cycles/kernel/types.h +++ b/intern/cycles/kernel/types.h @@ -86,6 +86,7 @@ CCL_NAMESPACE_BEGIN #define __AO__ #define __PASSES__ #define __HAIR__ +#define __POINTCLOUD__ #define __SVM__ #define __EMISSION__ #define __HOLDOUT__ @@ -125,6 +126,9 @@ CCL_NAMESPACE_BEGIN # if !(__KERNEL_FEATURES & KERNEL_FEATURE_HAIR) # undef __HAIR__ # endif +# if !(__KERNEL_FEATURES & KERNEL_FEATURE_POINTCLOUD) +# undef __POINTCLOUD__ +# endif # if !(__KERNEL_FEATURES & KERNEL_FEATURE_VOLUME) # undef __VOLUME__ # endif @@ -538,19 +542,22 @@ typedef enum PrimitiveType { PRIMITIVE_MOTION_CURVE_THICK = (1 << 3), PRIMITIVE_CURVE_RIBBON = (1 << 4), PRIMITIVE_MOTION_CURVE_RIBBON = (1 << 5), - PRIMITIVE_VOLUME = (1 << 6), - PRIMITIVE_LAMP = (1 << 7), + PRIMITIVE_POINT = (1 << 6), + PRIMITIVE_MOTION_POINT = (1 << 7), + PRIMITIVE_VOLUME = (1 << 8), + PRIMITIVE_LAMP = (1 << 9), PRIMITIVE_ALL_TRIANGLE = (PRIMITIVE_TRIANGLE | PRIMITIVE_MOTION_TRIANGLE), PRIMITIVE_ALL_CURVE = (PRIMITIVE_CURVE_THICK | PRIMITIVE_MOTION_CURVE_THICK | PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON), + PRIMITIVE_ALL_POINT = (PRIMITIVE_POINT | PRIMITIVE_MOTION_POINT), PRIMITIVE_ALL_VOLUME = (PRIMITIVE_VOLUME), PRIMITIVE_ALL_MOTION = (PRIMITIVE_MOTION_TRIANGLE | PRIMITIVE_MOTION_CURVE_THICK | - PRIMITIVE_MOTION_CURVE_RIBBON), + PRIMITIVE_MOTION_CURVE_RIBBON | PRIMITIVE_MOTION_POINT), PRIMITIVE_ALL = (PRIMITIVE_ALL_TRIANGLE | PRIMITIVE_ALL_CURVE | PRIMITIVE_ALL_VOLUME | - PRIMITIVE_LAMP), + PRIMITIVE_LAMP | PRIMITIVE_ALL_POINT), - PRIMITIVE_NUM = 8, + PRIMITIVE_NUM = 10, } PrimitiveType; #define PRIMITIVE_PACK_SEGMENT(type, segment) ((segment << PRIMITIVE_NUM) | (type)) @@ -605,6 +612,7 @@ typedef enum AttributeStandard { ATTR_STD_CURVE_INTERCEPT, ATTR_STD_CURVE_LENGTH, ATTR_STD_CURVE_RANDOM, + ATTR_STD_POINT_RANDOM, ATTR_STD_PTEX_FACE_ID, ATTR_STD_PTEX_UV, ATTR_STD_VOLUME_DENSITY, @@ -1604,6 +1612,9 @@ enum KernelFeatureFlag : unsigned int { KERNEL_FEATURE_AO_PASS = (1U << 25U), KERNEL_FEATURE_AO_ADDITIVE = (1U << 26U), KERNEL_FEATURE_AO = (KERNEL_FEATURE_AO_PASS | KERNEL_FEATURE_AO_ADDITIVE), + + /* Point clouds. */ + KERNEL_FEATURE_POINTCLOUD = (1U << 27U), }; /* Shader node feature mask, to specialize shader evaluation for kernels. */ diff --git a/intern/cycles/scene/CMakeLists.txt b/intern/cycles/scene/CMakeLists.txt index a3fde99306b..6f8b97a4998 100644 --- a/intern/cycles/scene/CMakeLists.txt +++ b/intern/cycles/scene/CMakeLists.txt @@ -40,6 +40,7 @@ set(SRC mesh_displace.cpp mesh_subdivision.cpp procedural.cpp + pointcloud.cpp object.cpp osl.cpp particles.cpp @@ -81,6 +82,7 @@ set(SRC_HEADERS particles.h pass.h procedural.h + pointcloud.h curves.h scene.h shader.h diff --git a/intern/cycles/scene/alembic.cpp b/intern/cycles/scene/alembic.cpp index 39b5f467736..71fa1863b8b 100644 --- a/intern/cycles/scene/alembic.cpp +++ b/intern/cycles/scene/alembic.cpp @@ -21,6 +21,7 @@ #include "scene/curves.h" #include "scene/mesh.h" #include "scene/object.h" +#include "scene/pointcloud.h" #include "scene/scene.h" #include "scene/shader.h" @@ -99,6 +100,9 @@ void CachedData::clear() triangles.clear(); uv_loops.clear(); vertices.clear(); + points.clear(); + radiuses.clear(); + points_shader.clear(); for (CachedAttribute &attr : attributes) { attr.data.clear(); @@ -146,6 +150,9 @@ bool CachedData::is_constant() const CHECK_IF_CONSTANT(triangles) CHECK_IF_CONSTANT(uv_loops) CHECK_IF_CONSTANT(vertices) + CHECK_IF_CONSTANT(points) + CHECK_IF_CONSTANT(radiuses) + CHECK_IF_CONSTANT(points_shader) for (const CachedAttribute &attr : attributes) { if (!attr.data.is_constant()) { @@ -185,6 +192,9 @@ void CachedData::invalidate_last_loaded_time(bool attributes_only) triangles.invalidate_last_loaded_time(); uv_loops.invalidate_last_loaded_time(); vertices.invalidate_last_loaded_time(); + points.invalidate_last_loaded_time(); + radiuses.invalidate_last_loaded_time(); + points_shader.invalidate_last_loaded_time(); } void CachedData::set_time_sampling(TimeSampling time_sampling) @@ -206,6 +216,9 @@ void CachedData::set_time_sampling(TimeSampling time_sampling) triangles.set_time_sampling(time_sampling); uv_loops.set_time_sampling(time_sampling); vertices.set_time_sampling(time_sampling); + points.set_time_sampling(time_sampling); + radiuses.set_time_sampling(time_sampling); + points_shader.set_time_sampling(time_sampling); for (CachedAttribute &attr : attributes) { attr.data.set_time_sampling(time_sampling); @@ -233,6 +246,9 @@ size_t CachedData::memory_used() const mem_used += triangles.memory_used(); mem_used += uv_loops.memory_used(); mem_used += vertices.memory_used(); + mem_used += points.memory_used(); + mem_used += radiuses.memory_used(); + mem_used += points_shader.memory_used(); for (const CachedAttribute &attr : attributes) { mem_used += attr.data.memory_used(); @@ -901,6 +917,9 @@ void AlembicProcedural::generate(Scene *scene, Progress &progress) else if (object->schema_type == AlembicObject::CURVES) { read_curves(object, frame_time); } + else if (object->schema_type == AlembicObject::POINTS) { + read_points(object, frame_time); + } else if (object->schema_type == AlembicObject::SUBD) { read_subd(object, frame_time); } @@ -970,6 +989,9 @@ void AlembicProcedural::load_objects(Progress &progress) if (abc_object->schema_type == AlembicObject::CURVES) { geometry = scene_->create_node<Hair>(); } + else if (abc_object->schema_type == AlembicObject::POINTS) { + geometry = scene_->create_node<PointCloud>(); + } else if (abc_object->schema_type == AlembicObject::POLY_MESH || abc_object->schema_type == AlembicObject::SUBD) { geometry = scene_->create_node<Mesh>(); @@ -1202,6 +1224,45 @@ void AlembicProcedural::read_curves(AlembicObject *abc_object, Abc::chrono_t fra hair->tag_update(scene_, rebuild); } +void AlembicProcedural::read_points(AlembicObject *abc_object, Abc::chrono_t frame_time) +{ + CachedData &cached_data = abc_object->get_cached_data(); + + /* update sockets */ + + Object *object = abc_object->get_object(); + cached_data.transforms.copy_to_socket(frame_time, object, object->get_tfm_socket()); + + if (object->is_modified()) { + object->tag_update(scene_); + } + + /* Only update sockets for the original Geometry. */ + if (abc_object->instance_of) { + return; + } + + PointCloud *point_cloud = static_cast<PointCloud *>(object->get_geometry()); + + /* Make sure shader ids are also updated. */ + if (point_cloud->used_shaders_is_modified()) { + point_cloud->tag_shader_modified(); + } + + cached_data.points.copy_to_socket(frame_time, point_cloud, point_cloud->get_points_socket()); + cached_data.radiuses.copy_to_socket(frame_time, point_cloud, point_cloud->get_radius_socket()); + cached_data.points_shader.copy_to_socket( + frame_time, point_cloud, point_cloud->get_shader_socket()); + + /* update attributes */ + + update_attributes(point_cloud->attributes, cached_data, frame_time); + + const bool rebuild = (point_cloud->points_is_modified() || point_cloud->radius_is_modified() || + point_cloud->shader_is_modified()); + point_cloud->tag_update(scene_, rebuild); +} + void AlembicProcedural::walk_hierarchy( IObject parent, const ObjectHeader &header, @@ -1314,7 +1375,23 @@ void AlembicProcedural::walk_hierarchy( // ignore the face set, it will be read along with the data } else if (IPoints::matches(header)) { - // unsupported for now + IPoints points(parent, header.getName()); + + unordered_map<std::string, AlembicObject *>::const_iterator iter; + iter = object_map.find(points.getFullName()); + + if (iter != object_map.end()) { + AlembicObject *abc_object = iter->second; + abc_object->iobject = points; + abc_object->schema_type = AlembicObject::POINTS; + + if (matrix_samples_data.samples) { + abc_object->xform_samples = *matrix_samples_data.samples; + abc_object->xform_time_sampling = matrix_samples_data.time_sampling; + } + } + + next_object = points; } else if (INuPatch::matches(header)) { // unsupported for now @@ -1391,6 +1468,14 @@ void AlembicProcedural::build_caches(Progress &progress) object->load_data_in_cache(object->get_cached_data(), this, schema, progress); } } + else if (object->schema_type == AlembicObject::POINTS) { + if (!object->has_data_loaded() || default_radius_is_modified() || + object->radius_scale_is_modified()) { + IPoints points(object->iobject, Alembic::Abc::kWrapExisting); + IPointsSchema schema = points.getSchema(); + object->load_data_in_cache(object->get_cached_data(), this, schema, progress); + } + } else if (object->schema_type == AlembicObject::SUBD) { if (!object->has_data_loaded()) { ISubD subd_mesh(object->iobject, Alembic::Abc::kWrapExisting); diff --git a/intern/cycles/scene/alembic.h b/intern/cycles/scene/alembic.h index 77aafd0ab32..3a4d37da3ff 100644 --- a/intern/cycles/scene/alembic.h +++ b/intern/cycles/scene/alembic.h @@ -327,6 +327,11 @@ struct CachedData { DataStore<array<int>> curve_first_key; DataStore<array<int>> curve_shader; + /* point data */ + DataStore<array<float3>> points; + DataStore<array<float>> radiuses; + DataStore<array<int>> points_shader; + struct CachedAttribute { AttributeStandard std; AttributeElement element; @@ -414,6 +419,7 @@ class AlembicObject : public Node { POLY_MESH, SUBD, CURVES, + POINTS, }; bool need_shader_update = true; @@ -550,6 +556,10 @@ class AlembicProcedural : public Procedural { * Object Nodes in the Cycles scene if none exist yet. */ void read_curves(AlembicObject *abc_object, Alembic::AbcGeom::Abc::chrono_t frame_time); + /* Read the data for an IPoints at the specified frame_time. Creates corresponding Geometry and + * Object Nodes in the Cycles scene if none exist yet. */ + void read_points(AlembicObject *abc_object, Alembic::AbcGeom::Abc::chrono_t frame_time); + /* Read the data for an ISubD at the specified frame_time. Creates corresponding Geometry and * Object Nodes in the Cycles scene if none exist yet. */ void read_subd(AlembicObject *abc_object, Alembic::AbcGeom::Abc::chrono_t frame_time); diff --git a/intern/cycles/scene/alembic_read.cpp b/intern/cycles/scene/alembic_read.cpp index 35f4854127a..07874b7528e 100644 --- a/intern/cycles/scene/alembic_read.cpp +++ b/intern/cycles/scene/alembic_read.cpp @@ -609,6 +609,55 @@ void read_geometry_data(AlembicProcedural *proc, read_data_loop(proc, cached_data, data, read_curves_data, progress); } +/* Points Geometries. */ + +static void read_points_data(CachedData &cached_data, const PointsSchemaData &data, chrono_t time) +{ + const ISampleSelector iss = ISampleSelector(time); + + const P3fArraySamplePtr position = data.positions.getValue(iss); + FloatArraySamplePtr radiuses; + + array<float3> a_positions; + array<float> a_radius; + array<int> a_shader; + a_positions.reserve(position->size()); + a_radius.reserve(position->size()); + a_shader.reserve(position->size()); + + if (data.radiuses.valid()) { + IFloatGeomParam::Sample wsample = data.radiuses.getExpandedValue(iss); + radiuses = wsample.getVals(); + } + + const bool do_radius = (radiuses != nullptr) && (radiuses->size() > 1); + float radius = (radiuses && radiuses->size() == 1) ? (*radiuses)[0] : data.default_radius; + + int offset = 0; + for (size_t i = 0; i < position->size(); i++) { + const V3f &f = position->get()[offset + i]; + a_positions.push_back_slow(make_float3_from_yup(f)); + + if (do_radius) { + radius = (*radiuses)[offset + i]; + a_radius.push_back_slow(radius); + } + + a_shader.push_back_slow((int)0); + } + + cached_data.points.add_data(a_positions, time); + cached_data.radiuses.add_data(a_radius, time); + cached_data.points_shader.add_data(a_shader, time); +} + +void read_geometry_data(AlembicProcedural *proc, + CachedData &cached_data, + const PointsSchemaData &data, + Progress &progress) +{ + read_data_loop(proc, cached_data, data, read_points_data, progress); +} /* Attributes conversions. */ /* Type traits for converting between Alembic and Cycles types. diff --git a/intern/cycles/scene/alembic_read.h b/intern/cycles/scene/alembic_read.h index 6b656b59481..e95750ce691 100644 --- a/intern/cycles/scene/alembic_read.h +++ b/intern/cycles/scene/alembic_read.h @@ -122,6 +122,26 @@ void read_geometry_data(AlembicProcedural *proc, const CurvesSchemaData &data, Progress &progress); +/* Data of a ICurvesSchema that we need to read. */ +struct PointsSchemaData { + Alembic::AbcGeom::TimeSamplingPtr time_sampling; + size_t num_samples; + + float default_radius; + float radius_scale; + + Alembic::AbcGeom::IP3fArrayProperty positions; + Alembic::AbcGeom::IInt32ArrayProperty num_points; + Alembic::AbcGeom::IFloatGeomParam radiuses; + // Those are unsupported for now. + Alembic::AbcGeom::IV3fArrayProperty velocities; +}; + +void read_geometry_data(AlembicProcedural *proc, + CachedData &cached_data, + const PointsSchemaData &data, + Progress &progress); + void read_attributes(AlembicProcedural *proc, CachedData &cache, const Alembic::AbcGeom::ICompoundProperty &arb_geom_params, diff --git a/intern/cycles/scene/attribute.cpp b/intern/cycles/scene/attribute.cpp index 6d15f3325f7..f1021a26e8a 100644 --- a/intern/cycles/scene/attribute.cpp +++ b/intern/cycles/scene/attribute.cpp @@ -18,6 +18,7 @@ #include "scene/hair.h" #include "scene/image.h" #include "scene/mesh.h" +#include "scene/pointcloud.h" #include "util/foreach.h" #include "util/log.h" @@ -205,6 +206,10 @@ size_t Attribute::element_size(Geometry *geom, AttributePrimitive prim) const size -= mesh->get_num_subd_verts(); } } + else if (geom->geometry_type == Geometry::POINTCLOUD) { + PointCloud *pointcloud = static_cast<PointCloud *>(geom); + size = pointcloud->num_points(); + } break; case ATTR_ELEMENT_VERTEX_MOTION: if (geom->geometry_type == Geometry::MESH) { @@ -215,6 +220,10 @@ size_t Attribute::element_size(Geometry *geom, AttributePrimitive prim) const size -= mesh->get_num_subd_verts() * (mesh->get_motion_steps() - 1); } } + else if (geom->geometry_type == Geometry::POINTCLOUD) { + PointCloud *pointcloud = static_cast<PointCloud *>(geom); + size = pointcloud->num_points() * (pointcloud->get_motion_steps() - 1); + } break; case ATTR_ELEMENT_FACE: if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) { @@ -346,6 +355,8 @@ const char *Attribute::standard_name(AttributeStandard std) return "curve_length"; case ATTR_STD_CURVE_RANDOM: return "curve_random"; + case ATTR_STD_POINT_RANDOM: + return "point_random"; case ATTR_STD_PTEX_FACE_ID: return "ptex_face_id"; case ATTR_STD_PTEX_UV: @@ -555,6 +566,28 @@ Attribute *AttributeSet::add(AttributeStandard std, ustring name) break; } } + else if (geometry->geometry_type == Geometry::POINTCLOUD) { + switch (std) { + case ATTR_STD_UV: + attr = add(name, TypeFloat2, ATTR_ELEMENT_VERTEX); + break; + case ATTR_STD_GENERATED: + attr = add(name, TypeDesc::TypePoint, ATTR_ELEMENT_VERTEX); + break; + case ATTR_STD_MOTION_VERTEX_POSITION: + attr = add(name, TypeDesc::TypeFloat4, ATTR_ELEMENT_VERTEX_MOTION); + break; + case ATTR_STD_POINT_RANDOM: + attr = add(name, TypeDesc::TypeFloat, ATTR_ELEMENT_VERTEX); + break; + case ATTR_STD_GENERATED_TRANSFORM: + attr = add(name, TypeDesc::TypeMatrix, ATTR_ELEMENT_MESH); + break; + default: + assert(0); + break; + } + } else if (geometry->geometry_type == Geometry::VOLUME) { switch (std) { case ATTR_STD_VERTEX_NORMAL: diff --git a/intern/cycles/scene/geometry.cpp b/intern/cycles/scene/geometry.cpp index ca330c5544d..558b0e13b0f 100644 --- a/intern/cycles/scene/geometry.cpp +++ b/intern/cycles/scene/geometry.cpp @@ -26,6 +26,7 @@ #include "scene/light.h" #include "scene/mesh.h" #include "scene/object.h" +#include "scene/pointcloud.h" #include "scene/scene.h" #include "scene/shader.h" #include "scene/shader_nodes.h" @@ -240,6 +241,7 @@ void Geometry::compute_bvh( params->use_bvh_unaligned_nodes; bparams.num_motion_triangle_steps = params->num_bvh_time_steps; bparams.num_motion_curve_steps = params->num_bvh_time_steps; + bparams.num_motion_point_steps = params->num_bvh_time_steps; bparams.bvh_type = params->bvh_type; bparams.curve_subdivisions = params->curve_subdivisions(); @@ -723,6 +725,12 @@ void GeometryManager::update_attribute_element_offset(Geometry *geom, else if (element == ATTR_ELEMENT_CURVE_KEY_MOTION) offset -= hair->curve_key_offset; } + else if (geom->is_pointcloud()) { + if (element == ATTR_ELEMENT_VERTEX) + offset -= geom->prim_offset; + else if (element == ATTR_ELEMENT_VERTEX_MOTION) + offset -= geom->prim_offset; + } } else { /* attribute not found */ @@ -1006,6 +1014,8 @@ void GeometryManager::mesh_calc_offset(Scene *scene, BVHLayout bvh_layout) size_t curve_key_size = 0; size_t curve_segment_size = 0; + size_t point_size = 0; + size_t patch_size = 0; size_t face_size = 0; size_t corner_size = 0; @@ -1054,6 +1064,14 @@ void GeometryManager::mesh_calc_offset(Scene *scene, BVHLayout bvh_layout) curve_key_size += hair->get_curve_keys().size(); curve_segment_size += hair->num_segments(); } + else if (geom->is_pointcloud()) { + PointCloud *pointcloud = static_cast<PointCloud *>(geom); + + prim_offset_changed = (pointcloud->prim_offset != point_size); + + pointcloud->prim_offset = point_size; + point_size += pointcloud->num_points(); + } if (prim_offset_changed) { /* Need to rebuild BVH in OptiX, since refit only allows modified mesh data there */ @@ -1079,6 +1097,8 @@ void GeometryManager::device_update_mesh(Device *, size_t curve_size = 0; size_t curve_segment_size = 0; + size_t point_size = 0; + size_t patch_size = 0; foreach (Geometry *geom, scene->geometry) { @@ -1106,6 +1126,10 @@ void GeometryManager::device_update_mesh(Device *, curve_size += hair->num_curves(); curve_segment_size += hair->num_segments(); } + else if (geom->is_pointcloud()) { + PointCloud *pointcloud = static_cast<PointCloud *>(geom); + point_size += pointcloud->num_points(); + } } /* Fill in all the arrays. */ @@ -1201,6 +1225,26 @@ void GeometryManager::device_update_mesh(Device *, dscene->curve_segments.copy_to_device_if_modified(); } + if (point_size != 0) { + progress.set_status("Updating Mesh", "Copying Point clouds to device"); + + float4 *points = dscene->points.alloc(point_size); + uint *points_shader = dscene->points_shader.alloc(point_size); + + foreach (Geometry *geom, scene->geometry) { + if (geom->is_pointcloud()) { + PointCloud *pointcloud = static_cast<PointCloud *>(geom); + pointcloud->pack( + scene, &points[pointcloud->prim_offset], &points_shader[pointcloud->prim_offset]); + if (progress.get_cancel()) + return; + } + } + + dscene->points.copy_to_device(); + dscene->points_shader.copy_to_device(); + } + if (patch_size != 0 && dscene->patches.need_realloc()) { progress.set_status("Updating Mesh", "Copying Patches to device"); @@ -1242,6 +1286,7 @@ void GeometryManager::device_update_bvh(Device *device, scene->params.use_bvh_unaligned_nodes; bparams.num_motion_triangle_steps = scene->params.num_bvh_time_steps; bparams.num_motion_curve_steps = scene->params.num_bvh_time_steps; + bparams.num_motion_point_steps = scene->params.num_bvh_time_steps; bparams.bvh_type = scene->params.bvh_type; bparams.curve_subdivisions = scene->params.curve_subdivisions(); @@ -1323,6 +1368,7 @@ void GeometryManager::device_update_bvh(Device *device, enum { DEVICE_CURVE_DATA_MODIFIED = (1 << 0), DEVICE_MESH_DATA_MODIFIED = (1 << 1), + DEVICE_POINT_DATA_MODIFIED = (1 << 2), ATTR_FLOAT_MODIFIED = (1 << 2), ATTR_FLOAT2_MODIFIED = (1 << 3), @@ -1332,17 +1378,20 @@ enum { CURVE_DATA_NEED_REALLOC = (1 << 7), MESH_DATA_NEED_REALLOC = (1 << 8), + POINT_DATA_NEED_REALLOC = (1 << 9), + + ATTR_FLOAT_NEEDS_REALLOC = (1 << 10), + ATTR_FLOAT2_NEEDS_REALLOC = (1 << 11), + ATTR_FLOAT3_NEEDS_REALLOC = (1 << 12), + ATTR_FLOAT4_NEEDS_REALLOC = (1 << 13), - ATTR_FLOAT_NEEDS_REALLOC = (1 << 9), - ATTR_FLOAT2_NEEDS_REALLOC = (1 << 10), - ATTR_FLOAT3_NEEDS_REALLOC = (1 << 11), - ATTR_FLOAT4_NEEDS_REALLOC = (1 << 12), - ATTR_UCHAR4_NEEDS_REALLOC = (1 << 13), + ATTR_UCHAR4_NEEDS_REALLOC = (1 << 14), ATTRS_NEED_REALLOC = (ATTR_FLOAT_NEEDS_REALLOC | ATTR_FLOAT2_NEEDS_REALLOC | ATTR_FLOAT3_NEEDS_REALLOC | ATTR_FLOAT4_NEEDS_REALLOC | ATTR_UCHAR4_NEEDS_REALLOC), DEVICE_MESH_DATA_NEEDS_REALLOC = (MESH_DATA_NEED_REALLOC | ATTRS_NEED_REALLOC), + DEVICE_POINT_DATA_NEEDS_REALLOC = (POINT_DATA_NEED_REALLOC | ATTRS_NEED_REALLOC), DEVICE_CURVE_DATA_NEEDS_REALLOC = (CURVE_DATA_NEED_REALLOC | ATTRS_NEED_REALLOC), }; @@ -1529,6 +1578,17 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro device_update_flags |= DEVICE_MESH_DATA_MODIFIED; } } + + if (geom->is_pointcloud()) { + PointCloud *pointcloud = static_cast<PointCloud *>(geom); + + if (pointcloud->need_update_rebuild) { + device_update_flags |= DEVICE_POINT_DATA_NEEDS_REALLOC; + } + else if (pointcloud->is_modified()) { + device_update_flags |= DEVICE_POINT_DATA_MODIFIED; + } + } } if (update_flags & (MESH_ADDED | MESH_REMOVED)) { @@ -1539,10 +1599,15 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro device_update_flags |= DEVICE_CURVE_DATA_NEEDS_REALLOC; } + if (update_flags & (POINT_ADDED | POINT_REMOVED)) { + device_update_flags |= DEVICE_POINT_DATA_NEEDS_REALLOC; + } + /* tag the device arrays for reallocation or modification */ DeviceScene *dscene = &scene->dscene; - if (device_update_flags & (DEVICE_MESH_DATA_NEEDS_REALLOC | DEVICE_CURVE_DATA_NEEDS_REALLOC)) { + if (device_update_flags & (DEVICE_MESH_DATA_NEEDS_REALLOC | DEVICE_CURVE_DATA_NEEDS_REALLOC | + DEVICE_POINT_DATA_NEEDS_REALLOC)) { delete scene->bvh; scene->bvh = nullptr; @@ -1570,6 +1635,11 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro dscene->curve_keys.tag_realloc(); dscene->curve_segments.tag_realloc(); } + + if (device_update_flags & DEVICE_POINT_DATA_NEEDS_REALLOC) { + dscene->points.tag_realloc(); + dscene->points_shader.tag_realloc(); + } } if ((update_flags & VISIBILITY_MODIFIED) != 0) { @@ -1630,6 +1700,11 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro dscene->curve_segments.tag_modified(); } + if (device_update_flags & DEVICE_POINT_DATA_MODIFIED) { + dscene->points.tag_modified(); + dscene->points_shader.tag_modified(); + } + need_flags_update = false; } @@ -2064,6 +2139,8 @@ void GeometryManager::device_update(Device *device, dscene->curves.clear_modified(); dscene->curve_keys.clear_modified(); dscene->curve_segments.clear_modified(); + dscene->points.clear_modified(); + dscene->points_shader.clear_modified(); dscene->patches.clear_modified(); dscene->attributes_map.clear_modified(); dscene->attributes_float.clear_modified(); @@ -2092,6 +2169,8 @@ void GeometryManager::device_free(Device *device, DeviceScene *dscene, bool forc dscene->curves.free_if_need_realloc(force_free); dscene->curve_keys.free_if_need_realloc(force_free); dscene->curve_segments.free_if_need_realloc(force_free); + dscene->points.free_if_need_realloc(force_free); + dscene->points_shader.free_if_need_realloc(force_free); dscene->patches.free_if_need_realloc(force_free); dscene->attributes_map.free_if_need_realloc(force_free); dscene->attributes_float.free_if_need_realloc(force_free); diff --git a/intern/cycles/scene/geometry.h b/intern/cycles/scene/geometry.h index 91799d7fde8..b02387c3020 100644 --- a/intern/cycles/scene/geometry.h +++ b/intern/cycles/scene/geometry.h @@ -55,6 +55,7 @@ class Geometry : public Node { MESH, HAIR, VOLUME, + POINTCLOUD, }; Type geometry_type; @@ -155,6 +156,11 @@ class Geometry : public Node { return geometry_type == HAIR; } + bool is_pointcloud() const + { + return geometry_type == POINTCLOUD; + } + bool is_volume() const { return geometry_type == VOLUME; @@ -181,12 +187,14 @@ class GeometryManager { MESH_REMOVED = (1 << 5), HAIR_ADDED = (1 << 6), HAIR_REMOVED = (1 << 7), + POINT_ADDED = (1 << 12), + POINT_REMOVED = (1 << 13), SHADER_ATTRIBUTE_MODIFIED = (1 << 8), SHADER_DISPLACEMENT_MODIFIED = (1 << 9), - GEOMETRY_ADDED = MESH_ADDED | HAIR_ADDED, - GEOMETRY_REMOVED = MESH_REMOVED | HAIR_REMOVED, + GEOMETRY_ADDED = MESH_ADDED | HAIR_ADDED | POINT_ADDED, + GEOMETRY_REMOVED = MESH_REMOVED | HAIR_REMOVED | POINT_REMOVED, TRANSFORM_MODIFIED = (1 << 10), diff --git a/intern/cycles/scene/object.cpp b/intern/cycles/scene/object.cpp index 05f6f6ce197..f8110b20d6e 100644 --- a/intern/cycles/scene/object.cpp +++ b/intern/cycles/scene/object.cpp @@ -23,6 +23,7 @@ #include "scene/light.h" #include "scene/mesh.h" #include "scene/particles.h" +#include "scene/pointcloud.h" #include "scene/scene.h" #include "scene/stats.h" #include "scene/volume.h" @@ -69,6 +70,7 @@ struct UpdateObjectTransformState { /* Flags which will be synchronized to Integrator. */ bool have_motion; bool have_curves; + // bool have_points; /* ** Scheduling queue. ** */ Scene *scene; @@ -435,7 +437,7 @@ void ObjectManager::device_update_object_transform(UpdateObjectTransformState *s state->have_motion = true; } - if (geom->geometry_type == Geometry::MESH) { + if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::POINTCLOUD) { /* TODO: why only mesh? */ Mesh *mesh = static_cast<Mesh *>(geom); if (mesh->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION)) { @@ -491,6 +493,8 @@ void ObjectManager::device_update_object_transform(UpdateObjectTransformState *s kobject.dupli_generated[2] = ob->dupli_generated[2]; kobject.numkeys = (geom->geometry_type == Geometry::HAIR) ? static_cast<Hair *>(geom)->get_curve_keys().size() : + (geom->geometry_type == Geometry::POINTCLOUD) ? + static_cast<PointCloud *>(geom)->num_points() : 0; kobject.dupli_uv[0] = ob->dupli_uv[0]; kobject.dupli_uv[1] = ob->dupli_uv[1]; diff --git a/intern/cycles/scene/pointcloud.cpp b/intern/cycles/scene/pointcloud.cpp new file mode 100644 index 00000000000..4f88fe9db3d --- /dev/null +++ b/intern/cycles/scene/pointcloud.cpp @@ -0,0 +1,304 @@ +/* + * Copyright 2011-2020 Blender Foundation + * + * 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. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "bvh/bvh.h" + +#include "scene/pointcloud.h" +#include "scene/scene.h" + +CCL_NAMESPACE_BEGIN + +/* PointCloud Point */ + +void PointCloud::Point::bounds_grow(const float3 *points, + const float *radius, + BoundBox &bounds) const +{ + bounds.grow(points[index], radius[index]); +} + +void PointCloud::Point::bounds_grow(const float3 *points, + const float *radius, + const Transform &aligned_space, + BoundBox &bounds) const +{ + float3 P = transform_point(&aligned_space, points[index]); + bounds.grow(P, radius[index]); +} + +void PointCloud::Point::bounds_grow(const float4 &point, BoundBox &bounds) const +{ + bounds.grow(float4_to_float3(point), point.w); +} + +float4 PointCloud::Point::motion_key(const float3 *points, + const float *radius, + const float3 *point_steps, + size_t num_points, + size_t num_steps, + float time, + size_t p) const +{ + /* Figure out which steps we need to fetch and their + * interpolation factor. */ + const size_t max_step = num_steps - 1; + const size_t step = min((int)(time * max_step), max_step - 1); + const float t = time * max_step - step; + /* Fetch vertex coordinates. */ + const float4 curr_key = point_for_step( + points, radius, point_steps, num_points, num_steps, step, p); + const float4 next_key = point_for_step( + points, radius, point_steps, num_points, num_steps, step + 1, p); + /* Interpolate between steps. */ + return (1.0f - t) * curr_key + t * next_key; +} + +float4 PointCloud::Point::point_for_step(const float3 *points, + const float *radius, + const float3 *point_steps, + size_t num_points, + size_t num_steps, + size_t step, + size_t p) const +{ + const size_t center_step = ((num_steps - 1) / 2); + if (step == center_step) { + /* Center step: regular key location. */ + return make_float4(points[p].x, points[p].y, points[p].z, radius[p]); + } + else { + /* Center step is not stored in this array. */ + if (step > center_step) { + step--; + } + const size_t offset = step * num_points; + return make_float4(point_steps[offset + p].x, + point_steps[offset + p].y, + point_steps[offset + p].z, + radius[offset + p]); + } +} + +/* PointCloud */ + +NODE_DEFINE(PointCloud) +{ + NodeType *type = NodeType::add( + "pointcloud", create, NodeType::NONE, Geometry::get_node_base_type()); + + SOCKET_POINT_ARRAY(points, "Points", array<float3>()); + SOCKET_FLOAT_ARRAY(radius, "Radius", array<float>()); + SOCKET_INT_ARRAY(shader, "Shader", array<int>()); + + return type; +} + +PointCloud::PointCloud() : Geometry(node_type, Geometry::POINTCLOUD) +{ +} + +PointCloud::~PointCloud() +{ +} + +void PointCloud::resize(int numpoints) +{ + points.resize(numpoints); + radius.resize(numpoints); + shader.resize(numpoints); + attributes.resize(); + + tag_points_modified(); + tag_radius_modified(); + tag_shader_modified(); +} + +void PointCloud::reserve(int numpoints) +{ + points.reserve(numpoints); + radius.reserve(numpoints); + shader.reserve(numpoints); + attributes.resize(true); +} + +void PointCloud::clear(const bool preserve_shaders) +{ + Geometry::clear(preserve_shaders); + + points.clear(); + radius.clear(); + shader.clear(); + attributes.clear(); + + tag_points_modified(); + tag_radius_modified(); + tag_shader_modified(); +} + +void PointCloud::add_point(float3 co, float r, int shader_index) +{ + points.push_back_reserved(co); + radius.push_back_reserved(r); + shader.push_back_reserved(shader_index); + + tag_points_modified(); + tag_radius_modified(); + tag_shader_modified(); +} + +void PointCloud::copy_center_to_motion_step(const int motion_step) +{ + Attribute *attr_mP = attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); + if (attr_mP) { + float3 *points_data = points.data(); + size_t numpoints = points.size(); + memcpy( + attr_mP->data_float3() + motion_step * numpoints, points_data, sizeof(float3) * numpoints); + } +} + +void PointCloud::get_uv_tiles(ustring map, unordered_set<int> &tiles) +{ + Attribute *attr; + + if (map.empty()) { + attr = attributes.find(ATTR_STD_UV); + } + else { + attr = attributes.find(map); + } + + if (attr) { + attr->get_uv_tiles(this, ATTR_PRIM_GEOMETRY, tiles); + } +} + +void PointCloud::compute_bounds() +{ + BoundBox bnds = BoundBox::empty; + size_t numpoints = points.size(); + + if (numpoints > 0) { + for (size_t i = 0; i < numpoints; i++) { + bnds.grow(points[i], radius[i]); + } + + Attribute *attr = attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); + if (use_motion_blur && attr) { + size_t steps_size = points.size() * (motion_steps - 1); + float3 *point_steps = attr->data_float3(); + + for (size_t i = 0; i < steps_size; i++) + bnds.grow(point_steps[i]); + } + + if (!bnds.valid()) { + bnds = BoundBox::empty; + + /* skip nan or inf coordinates */ + for (size_t i = 0; i < numpoints; i++) + bnds.grow_safe(points[i], radius[i]); + + if (use_motion_blur && attr) { + size_t steps_size = points.size() * (motion_steps - 1); + float3 *point_steps = attr->data_float3(); + + for (size_t i = 0; i < steps_size; i++) + bnds.grow_safe(point_steps[i]); + } + } + } + + if (!bnds.valid()) { + /* empty mesh */ + bnds.grow(make_float3(0.0f, 0.0f, 0.0f)); + } + + bounds = bnds; +} + +void PointCloud::apply_transform(const Transform &tfm, const bool apply_to_motion) +{ + /* compute uniform scale */ + float3 c0 = transform_get_column(&tfm, 0); + float3 c1 = transform_get_column(&tfm, 1); + float3 c2 = transform_get_column(&tfm, 2); + float scalar = powf(fabsf(dot(cross(c0, c1), c2)), 1.0f / 3.0f); + + /* apply transform to curve keys */ + for (size_t i = 0; i < points.size(); i++) { + float3 co = transform_point(&tfm, points[i]); + float r = radius[i] * scalar; + + /* scale for curve radius is only correct for uniform scale + */ + points[i] = co; + radius[i] = r; + } + + if (apply_to_motion) { + Attribute *attr = attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); + + if (attr) { + /* apply transform to motion curve keys */ + size_t steps_size = points.size() * (motion_steps - 1); + float4 *point_steps = attr->data_float4(); + + for (size_t i = 0; i < steps_size; i++) { + float3 co = transform_point(&tfm, float4_to_float3(point_steps[i])); + float radius = point_steps[i].w * scalar; + + /* scale for curve radius is only correct for uniform + * scale */ + point_steps[i] = float3_to_float4(co); + point_steps[i].w = radius; + } + } + } +} + +void PointCloud::pack(Scene *scene, float4 *packed_points, uint *packed_shader) +{ + size_t numpoints = points.size(); + float3 *points_data = points.data(); + float *radius_data = radius.data(); + int *shader_data = shader.data(); + + for (size_t i = 0; i < numpoints; i++) { + packed_points[i] = make_float4( + points_data[i].x, points_data[i].y, points_data[i].z, radius_data[i]); + } + + uint shader_id = 0; + uint last_shader = -1; + for (size_t i = 0; i < numpoints; i++) { + if (last_shader != shader_data[i]) { + last_shader = shader_data[i]; + Shader *shader = (last_shader < used_shaders.size()) ? + static_cast<Shader *>(used_shaders[last_shader]) : + scene->default_surface; + shader_id = scene->shader_manager->get_shader_id(shader); + } + packed_shader[i] = shader_id; + } +} + +PrimitiveType PointCloud::primitive_type() const +{ + return has_motion_blur() ? PRIMITIVE_MOTION_POINT : PRIMITIVE_POINT; +} + +CCL_NAMESPACE_END diff --git a/intern/cycles/scene/pointcloud.h b/intern/cycles/scene/pointcloud.h new file mode 100644 index 00000000000..08987ee6c44 --- /dev/null +++ b/intern/cycles/scene/pointcloud.h @@ -0,0 +1,114 @@ +/* + * Copyright 2011-2020 Blender Foundation + * + * 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. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#ifndef __POINTCLOUD_H__ +# define __POINTCLOUD_H__ + +# include "scene/geometry.h" + +CCL_NAMESPACE_BEGIN + +class PointCloud : public Geometry { + public: + NODE_DECLARE + + /* PointCloud Point */ + struct Point { + int index; + + void bounds_grow(const float3 *points, const float *radius, BoundBox &bounds) const; + void bounds_grow(const float3 *points, + const float *radius, + const Transform &aligned_space, + BoundBox &bounds) const; + void bounds_grow(const float4 &point, BoundBox &bounds) const; + + float4 motion_key(const float3 *points, + const float *radius, + const float3 *point_steps, + size_t num_points, + size_t num_steps, + float time, + size_t p) const; + float4 point_for_step(const float3 *points, + const float *radius, + const float3 *point_steps, + size_t num_points, + size_t num_steps, + size_t step, + size_t p) const; + }; + + NODE_SOCKET_API_ARRAY(array<float3>, points) + NODE_SOCKET_API_ARRAY(array<float>, radius) + NODE_SOCKET_API_ARRAY(array<int>, shader) + + /* Constructor/Destructor */ + PointCloud(); + ~PointCloud(); + + /* Geometry */ + void clear(const bool preserver_shaders = false) override; + + void resize(int numpoints); + void reserve(int numpoints); + void add_point(float3 loc, float radius, int shader = 0); + + void copy_center_to_motion_step(const int motion_step); + + void compute_bounds() override; + void apply_transform(const Transform &tfm, const bool apply_to_motion) override; + + /* Points */ + Point get_point(int i) const + { + Point point = {i}; + return point; + } + + size_t num_points() const + { + return points.size(); + } + + size_t num_attributes() const + { + return 1; + } + + /* UDIM */ + void get_uv_tiles(ustring map, unordered_set<int> &tiles) override; + + PrimitiveType primitive_type() const override; + + /* BVH */ + void pack(Scene *scene, float4 *packed_points, uint *packed_shader); + + private: + friend class BVH2; + friend class BVHBuild; + friend class BVHSpatialSplit; + friend class DiagSplit; + friend class EdgeDice; + friend class GeometryManager; + friend class ObjectManager; +}; + +CCL_NAMESPACE_END + +#endif /* __POINTCLOUD_H__ */ diff --git a/intern/cycles/scene/scene.cpp b/intern/cycles/scene/scene.cpp index 452b5215836..1963ebbbb19 100644 --- a/intern/cycles/scene/scene.cpp +++ b/intern/cycles/scene/scene.cpp @@ -30,6 +30,7 @@ #include "scene/object.h" #include "scene/osl.h" #include "scene/particles.h" +#include "scene/pointcloud.h" #include "scene/procedural.h" #include "scene/scene.h" #include "scene/shader.h" @@ -64,6 +65,8 @@ DeviceScene::DeviceScene(Device *device) curve_keys(device, "__curve_keys", MEM_GLOBAL), curve_segments(device, "__curve_segments", MEM_GLOBAL), patches(device, "__patches", MEM_GLOBAL), + points(device, "__points", MEM_GLOBAL), + points_shader(device, "__points_shader", MEM_GLOBAL), objects(device, "__objects", MEM_GLOBAL), object_motion_pass(device, "__object_motion_pass", MEM_GLOBAL), object_motion(device, "__object_motion", MEM_GLOBAL), @@ -523,6 +526,9 @@ void Scene::update_kernel_features() else if (geom->is_hair()) { kernel_features |= KERNEL_FEATURE_HAIR; } + else if (geom->is_pointcloud()) { + kernel_features |= KERNEL_FEATURE_POINTCLOUD; + } } if (bake_manager->get_baking()) { @@ -575,6 +581,7 @@ static void log_kernel_features(const uint features) VLOG(2) << "Use Path Tracing " << string_from_bool(features & KERNEL_FEATURE_PATH_TRACING) << "\n"; VLOG(2) << "Use Hair " << string_from_bool(features & KERNEL_FEATURE_HAIR) << "\n"; + VLOG(2) << "Use Pointclouds " << string_from_bool(features & KERNEL_FEATURE_POINTCLOUD) << "\n"; VLOG(2) << "Use Object Motion " << string_from_bool(features & KERNEL_FEATURE_OBJECT_MOTION) << "\n"; VLOG(2) << "Use Camera Motion " << string_from_bool(features & KERNEL_FEATURE_CAMERA_MOTION) @@ -757,6 +764,15 @@ template<> Volume *Scene::create_node<Volume>() return node; } +template<> PointCloud *Scene::create_node<PointCloud>() +{ + PointCloud *node = new PointCloud(); + node->set_owner(this); + geometry.push_back(node); + geometry_manager->tag_update(this, GeometryManager::POINT_ADDED); + return node; +} + template<> Object *Scene::create_node<Object>() { Object *node = new Object(); @@ -844,6 +860,12 @@ template<> void Scene::delete_node_impl(Volume *node) geometry_manager->tag_update(this, GeometryManager::MESH_REMOVED); } +template<> void Scene::delete_node_impl(PointCloud *node) +{ + delete_node_from_array(geometry, static_cast<Geometry *>(node)); + geometry_manager->tag_update(this, GeometryManager::POINT_REMOVED); +} + template<> void Scene::delete_node_impl(Geometry *node) { uint flag; diff --git a/intern/cycles/scene/scene.h b/intern/cycles/scene/scene.h index f8f672a079a..ec935b41be6 100644 --- a/intern/cycles/scene/scene.h +++ b/intern/cycles/scene/scene.h @@ -54,6 +54,7 @@ class Object; class ObjectManager; class ParticleSystemManager; class ParticleSystem; +class PointCloud; class Procedural; class ProceduralManager; class CurveSystemManager; @@ -94,6 +95,10 @@ class DeviceScene { device_vector<uint> patches; + /* pointcloud */ + device_vector<float4> points; + device_vector<uint> points_shader; + /* objects */ device_vector<KernelObject> objects; device_vector<Transform> object_motion_pass; @@ -365,6 +370,8 @@ template<> Hair *Scene::create_node<Hair>(); template<> Volume *Scene::create_node<Volume>(); +template<> PointCloud *Scene::create_node<PointCloud>(); + template<> ParticleSystem *Scene::create_node<ParticleSystem>(); template<> Shader *Scene::create_node<Shader>(); @@ -379,6 +386,8 @@ template<> void Scene::delete_node_impl(Mesh *node); template<> void Scene::delete_node_impl(Volume *node); +template<> void Scene::delete_node_impl(PointCloud *node); + template<> void Scene::delete_node_impl(Hair *node); template<> void Scene::delete_node_impl(Geometry *node); diff --git a/tests/python/CMakeLists.txt b/tests/python/CMakeLists.txt index a2c2d44fe20..7f221e1c4eb 100644 --- a/tests/python/CMakeLists.txt +++ b/tests/python/CMakeLists.txt @@ -649,6 +649,7 @@ if(WITH_CYCLES OR WITH_OPENGL_RENDER_TESTS) integrator light mesh + pointcloud shader shadow_catcher sss |