blender
/
blender
130
397
Fork 571
Code Issues 6.2k Pull Requests 755 Projects 19 Wiki Activity

Cycles: pointcloud rendering 

This add support for rendering of the point cloud object in Blender, as a native
geometry type in Cycles that is more memory and time efficient than instancing
sphere meshes. This can be useful for rendering sand, water splashes, particles,
motion graphics, etc.

Points are currently always rendered as spheres, with backface culling. More
shapes are likely to be added later, but this is the most important one and can
be customized with shaders.

For CPU rendering the Embree primitive is used, for GPU there is our own
intersection code. Motion blur is suppored. Volumes inside points are not
currently supported.

Implemented with help from:
* Kévin Dietrich: Alembic procedural integration
* Patrick Mourse: OptiX integration
* Josh Whelchel: update for cycles-x changes

Ref T92573

Differential Revision: https://developer.blender.org/D9887
This commit is contained in:
Brecht Van Lommel 2021-12-01 17:30:46 +01:00
parent 2229179faa
commit 35b1e9fc3a
Notes: blender-bot 2023-02-14 09:43:37 +01:00 
Referenced by issue #94142, Cycles Metal crash with simultaneous viewport and final render
Referenced by issue #92573, Cycles: point cloud geometry improvements
48 changed files with 2047 additions and 69 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
intern/cycles/blender/CMakeLists.txt
View File

@ -40,6 +40,7 @@ set(SRC
object_cull.cpp
output_driver.cpp
particles.cpp
pointcloud.cpp
curves.cpp
logging.cpp
python.cpp

2
intern/cycles/blender/addon/ui.py
View File

@ -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

16
intern/cycles/blender/geometry.cpp
View File

@ -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);

5
intern/cycles/blender/object.cpp
View File

@ -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;
}

253
intern/cycles/blender/pointcloud.cpp Normal file
View File

@ -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

6
intern/cycles/blender/sync.h
View File

@ -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,

130
intern/cycles/bvh/build.cpp
View File

@ -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 &center, Hair *hair
}
}
void BVHBuild::add_reference_points(BoundBox &root,
BoundBox &center,
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 &center,
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 &center, 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 */

2
intern/cycles/bvh/build.h
View File

@ -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 &center, Mesh *mesh, int i);
void add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair, int i);
void add_reference_points(BoundBox &root, BoundBox &center, PointCloud *pointcloud, int i);
void add_reference_geometry(BoundBox &root, BoundBox &center, Geometry *geom, int i);
void add_reference_object(BoundBox &root, BoundBox &center, Object *ob, int i);
void add_references(BVHRange &root);

28
intern/cycles/bvh/bvh2.cpp
View File

@ -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);
}

106
intern/cycles/bvh/embree.cpp
View File

@ -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;
}

5
intern/cycles/bvh/embree.h
View File

@ -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;

17
intern/cycles/bvh/params.h
View File

@ -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.

48
intern/cycles/bvh/split.cpp
View File

@ -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);
}

14
intern/cycles/bvh/split.h
View File

@ -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);

120
intern/cycles/device/optix/device_impl.cpp
View File

@ -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,
@ -1374,6 +1402,86 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
build_input.triangleArray.numSbtRecords = 1;
build_input.triangleArray.primitiveIndexOffset = mesh->prim_offset;
if (!build_optix_bvh(bvh_optix, operation, build_input, num_motion_steps)) {
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");
}
@ -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.
*/

6
intern/cycles/device/optix/device_impl.h
View File

@ -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 };

3
intern/cycles/kernel/CMakeLists.txt
View File

@ -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

17
intern/cycles/kernel/bvh/bvh.h
View File

@ -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. */

29
intern/cycles/kernel/bvh/shadow_all.h
View File

@ -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;

28
intern/cycles/kernel/bvh/traversal.h
View File

@ -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 {

1
intern/cycles/kernel/bvh/types.h
View File

@ -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)

24
intern/cycles/kernel/bvh/util.h
View File

@ -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

64
intern/cycles/kernel/device/optix/kernel.cu
View File

@ -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

3
intern/cycles/kernel/geom/geom.h
View File

@ -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"

74
intern/cycles/kernel/geom/motion_point.h Normal file
View File

@ -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

124
intern/cycles/kernel/geom/point.h Normal file
View File

@ -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

133
intern/cycles/kernel/geom/point_intersect.h Normal file
View File

@ -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

49
intern/cycles/kernel/geom/primitive.h
View File

@ -47,6 +47,11 @@ ccl_device_inline float primitive_surface_attribute_float(KernelGlobals kg,
else if (sd->type & PRIMITIVE_ALL_CURVE) {
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)
@ -73,6 +78,11 @@ ccl_device_inline float2 primitive_surface_attribute_float2(KernelGlobals kg,
else if (sd->type & PRIMITIVE_ALL_CURVE) {
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)
@ -99,6 +109,11 @@ ccl_device_inline float3 primitive_surface_attribute_float3(KernelGlobals kg,
else if (sd->type & PRIMITIVE_ALL_CURVE) {
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)
@ -125,6 +140,11 @@ ccl_device_forceinline float4 primitive_surface_attribute_float4(KernelGlobals k
else if (sd->type & PRIMITIVE_ALL_CURVE) {
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)
@ -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, &center);
@ -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));

7
intern/cycles/kernel/geom/shader_data.h
View File

@ -74,6 +74,13 @@ ccl_device_inline void shader_setup_from_ray(KernelGlobals kg,
curve_shader_setup(kg, sd, ray->P, ray->D, isect->t, isect->object, isect->prim);
}
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 */

12
intern/cycles/kernel/osl/services.cpp
View File

@ -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);

2
intern/cycles/kernel/osl/services.h
View File

@ -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;

2
intern/cycles/kernel/svm/wireframe.h
View File

@ -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)

4
intern/cycles/kernel/textures.h
View File

@ -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)

21
intern/cycles/kernel/types.h
View File

@ -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. */

2
intern/cycles/scene/CMakeLists.txt
View File

@ -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

87
intern/cycles/scene/alembic.cpp
View File

@ -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);

10
intern/cycles/scene/alembic.h
View File

@ -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);

49
intern/cycles/scene/alembic_read.cpp
View File

@ -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.

20
intern/cycles/scene/alembic_read.h
View File

@ -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,

33
intern/cycles/scene/attribute.cpp
View File

@ -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:

91
intern/cycles/scene/geometry.cpp
View File

@ -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 << 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_FLOAT_NEEDS_REALLOC = (1 << 10),
ATTR_FLOAT2_NEEDS_REALLOC = (1 << 11),
ATTR_FLOAT3_NEEDS_REALLOC = (1 << 12),
ATTR_FLOAT4_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);

12
intern/cycles/scene/geometry.h
View File

@ -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),

6
intern/cycles/scene/object.cpp
View File

@ -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];

304
intern/cycles/scene/pointcloud.cpp Normal file
View File

@ -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

114
intern/cycles/scene/pointcloud.h Normal file
View File

@ -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__ */

22
intern/cycles/scene/scene.cpp
View File

@ -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;

9
intern/cycles/scene/scene.h
View File

@ -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);

1
tests/python/CMakeLists.txt
View File

@ -649,6 +649,7 @@ if(WITH_CYCLES OR WITH_OPENGL_RENDER_TESTS)
integrator
light
mesh
pointcloud
shader
shadow_catcher
sss