ilusha/blender
1
0
Fork 0
mirror of https://github.com/blender/blender synced 2026-09-29 04:37:17 +03:00
Code Issues Projects Releases Packages Wiki Activity Actions

Refactor: Cycles: Hair/Point positions include motion, kernel attribute

Browse source 
Motion is now part of the position attribute. On the kernel side, a
combined position + radius attribute is now stored, replacing the
previous motion only attribute.

Legacy motion attributes are now removed.

Pull Request: https://projects.blender.org/blender/blender/pulls/158728
This commit is contained in:
Brecht Van Lommel 2026-04-27 02:05:46 +02:00
parent 0baa98866c
commit fc9917352b
32 changed files with 769 additions and 1015 deletions
Download patch file Download diff file Expand all files Collapse all files

106
intern/cycles/blender/curves.cpp
View file

@ -464,48 +464,52 @@ static void export_hair_motion_validate_attribute(Hair *hair,
const int num_motion_keys, const int num_motion_keys,
bool have_motion) bool have_motion)
{ {
Attribute *attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
Attribute *attr_R = hair->attributes.find(ATTR_STD_RADIUS);
const int num_keys = hair->num_keys(); const int num_keys = hair->num_keys();
if (num_motion_keys != num_keys || !have_motion) { if (num_motion_keys != num_keys || !have_motion) {
/* No motion or hair "topology" changed, remove attributes again. */ /* No motion or hair "topology" changed, remove motion steps. */
if (num_motion_keys != num_keys) { if (num_motion_keys != num_keys) {
LOG_DEBUG << "Hair topology changed, removing motion attribute."; LOG_DEBUG << "Hair topology changed, removing motion attribute.";
} }
hair->attributes.remove(ATTR_STD_MOTION_VERTEX_POSITION); attr_P->remove_motion();
attr_R->remove_motion();
} }
else if (motion_step > 0) { else if (motion_step > 0) {
/* Motion, fill up previous steps that we might have skipped because /* Motion, fill up previous steps that we might have skipped because
* they had no motion, but we need them anyway now. */ * they had no motion, but we need them anyway now. */
for (int step = 0; step < motion_step; step++) { for (int step = 1; step <= motion_step; step++) {
float4 *mP = attr_mP->data_for_write<float4>() + step * num_keys; packed_float3 *mP = attr_P->data_for_write<packed_float3>(step);
std::copy_n(hair->get_position(), num_keys, mP);
for (int key = 0; key < num_keys; key++) { float *mR = attr_R->data_for_write<float>(step);
mP[key] = make_float4(hair->get_position()[key]); std::copy_n(hair->get_radius(), num_keys, mR);
mP[key].w = hair->get_radius()[key];
}
} }
} }
} }
static void ExportCurveSegmentsMotion(Hair *hair, ParticleCurveData *CData, const int motion_step) static void ExportCurveSegmentsMotion(Hair *hair, ParticleCurveData *CData, const int motion_step)
{ {
/* find attribute */ /* Set motion steps on position and radius attributes. */
Attribute *attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
Attribute *attr_R = hair->attributes.find(ATTR_STD_RADIUS);
bool new_attribute = false; bool new_attribute = false;
/* add new attribute if it doesn't exist already */ if (!attr_P->has_motion()) {
if (!attr_mP) { attr_P->add_motion(hair);
attr_mP = hair->attributes.add(ATTR_STD_MOTION_VERTEX_POSITION); attr_R->add_motion(hair);
new_attribute = true; new_attribute = true;
} }
/* export motion vectors for curve keys */ /* export motion vectors for curve keys */
const size_t numkeys = hair->num_keys(); const int attr_step = motion_step + 1;
float4 *mP = attr_mP->data_for_write<float4>() + motion_step * numkeys; packed_float3 *mP = attr_P->data_for_write<packed_float3>(attr_step);
float *mR = attr_R->data_for_write<float>(attr_step);
bool have_motion = false; bool have_motion = false;
int i = 0; int i = 0;
int num_curves = 0; int num_curves = 0;
const int num_keys = hair->num_keys();
for (int sys = 0; sys < CData->psys_firstcurve.size(); sys++) { for (int sys = 0; sys < CData->psys_firstcurve.size(); sys++) {
for (int curve = CData->psys_firstcurve[sys]; for (int curve = CData->psys_firstcurve[sys];
@ -515,7 +519,7 @@ static void ExportCurveSegmentsMotion(Hair *hair, ParticleCurveData *CData, cons
/* Curve lengths may not match! Curves can be clipped. */ /* Curve lengths may not match! Curves can be clipped. */
const int curve_key_end = (num_curves + 1 < (int)hair->get_curve_first_key().size() ? const int curve_key_end = (num_curves + 1 < (int)hair->get_curve_first_key().size() ?
hair->get_curve_first_key()[num_curves + 1] : hair->get_curve_first_key()[num_curves + 1] :
(int)hair->num_keys()); num_keys);
const int num_center_curve_keys = curve_key_end - hair->get_curve_first_key()[num_curves]; const int num_center_curve_keys = curve_key_end - hair->get_curve_first_key()[num_curves];
const int is_num_keys_different = CData->curve_keynum[curve] - num_center_curve_keys; const int is_num_keys_different = CData->curve_keynum[curve] - num_center_curve_keys;
@ -524,15 +528,17 @@ static void ExportCurveSegmentsMotion(Hair *hair, ParticleCurveData *CData, cons
curvekey < CData->curve_firstkey[curve] + CData->curve_keynum[curve]; curvekey < CData->curve_firstkey[curve] + CData->curve_keynum[curve];
curvekey++) curvekey++)
{ {
if (i < hair->num_keys()) { if (i < num_keys) {
mP[i] = CurveSegmentMotionCV(CData, sys, curve, curvekey); const float4 cv = CurveSegmentMotionCV(CData, sys, curve, curvekey);
mP[i] = make_float3(cv);
mR[i] = cv.w;
if (!have_motion) { if (!have_motion) {
/* unlike mesh coordinates, these tend to be slightly different /* unlike mesh coordinates, these tend to be slightly different
* between frames due to particle transforms into/out of object * between frames due to particle transforms into/out of object
* space, so we use an epsilon to detect actual changes */ * space, so we use an epsilon to detect actual changes */
float4 curve_key = make_float4(hair->get_position()[i]); float4 curve_key = make_float4(hair->get_position()[i]);
curve_key.w = hair->get_radius()[i]; curve_key.w = hair->get_radius()[i];
if (len_squared(mP[i] - curve_key) > 1e-5f * 1e-5f) { if (len_squared(cv - curve_key) > 1e-5f * 1e-5f) {
have_motion = true; have_motion = true;
} }
} }
@ -547,7 +553,9 @@ static void ExportCurveSegmentsMotion(Hair *hair, ParticleCurveData *CData, cons
0.0f; 0.0f;
for (int step_index = 0; step_index < num_center_curve_keys; ++step_index) { for (int step_index = 0; step_index < num_center_curve_keys; ++step_index) {
const float step = step_index * step_size; const float step = step_index * step_size;
mP[i] = LerpCurveSegmentMotionCV(CData, sys, curve, step); const float4 cv = LerpCurveSegmentMotionCV(CData, sys, curve, step);
mP[i] = make_float3(cv);
mR[i] = cv.w;
i++; i++;
} }
have_motion = true; have_motion = true;
@ -717,19 +725,17 @@ static void attr_create_motion_from_velocity(Hair *hair,
/* Override motion steps to fixed number. */ /* Override motion steps to fixed number. */
hair->set_motion_steps(3); hair->set_motion_steps(3);
/* Find or add attribute */ /* Set motion steps on position attribute. Radius doesn't change for
* velocity-based motion. */
Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
attr_P->add_motion(hair);
const packed_float3 *P = hair->get_position(); const packed_float3 *P = hair->get_position();
Attribute *attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
if (!attr_mP) {
attr_mP = hair->attributes.add(ATTR_STD_MOTION_VERTEX_POSITION);
}
/* Only export previous and next frame, we don't have any in between data. */ /* Only export previous and next frame, we don't have any in between data. */
const float motion_times[2] = {-1.0f, 1.0f}; const float motion_times[2] = {-1.0f, 1.0f};
for (int step = 0; step < 2; step++) { for (int step = 1; step <= 2; step++) {
const float relative_time = motion_times[step] * 0.5f * motion_scale; const float relative_time = motion_times[step - 1] * 0.5f * motion_scale;
packed_float3 *mP = attr_mP->data_for_write<packed_float3>() + step * num_curve_keys; packed_float3 *mP = attr_P->data_for_write<packed_float3>(step);
for (int i = 0; i < num_curve_keys; i++) { for (int i = 0; i < num_curve_keys; i++) {
mP[i] = float3(P[i]) + make_float3(src[i][0], src[i][1], src[i][2]) * relative_time; mP[i] = float3(P[i]) + make_float3(src[i][0], src[i][1], src[i][2]) * relative_time;
@ -965,6 +971,9 @@ static void export_hair_curves(Scene *scene,
} }
} }
hair->tag_position_modified();
hair->tag_radius_modified();
attr_create_generic(scene, hair, b_curves, need_motion, motion_scale); attr_create_generic(scene, hair, b_curves, need_motion, motion_scale);
} }
@ -972,19 +981,23 @@ static void export_hair_curves_motion(Hair *hair,
const blender::bke::CurvesGeometry &b_curves, const blender::bke::CurvesGeometry &b_curves,
const int motion_step) const int motion_step)
{ {
/* Find or add attribute. */ /* Set motion steps on position and radius attributes. */
Attribute *attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
Attribute *attr_R = hair->attributes.find(ATTR_STD_RADIUS);
bool new_attribute = false; bool new_attribute = false;
if (!attr_mP) { if (!attr_P->has_motion()) {
attr_mP = hair->attributes.add(ATTR_STD_MOTION_VERTEX_POSITION); attr_P->add_motion(hair);
attr_R->add_motion(hair);
new_attribute = true; new_attribute = true;
} }
/* Export motion keys. */ /* Export motion keys. */
const size_t num_keys = hair->num_keys(); const size_t num_keys = hair->num_keys();
const size_t num_curves = hair->num_curves(); const size_t num_curves = hair->num_curves();
float4 *mP = attr_mP->data_for_write<float4>() + motion_step * num_keys; const int attr_step = motion_step + 1;
packed_float3 *mP = attr_P->data_for_write<packed_float3>(attr_step);
float *mR = attr_R->data_for_write<float>(attr_step);
bool have_motion = false; bool have_motion = false;
int num_motion_keys = 0; int num_motion_keys = 0;
int curve_index = 0; int curve_index = 0;
@ -1009,15 +1022,17 @@ static void export_hair_curves_motion(Hair *hair,
const int point = points[i]; const int point = points[i];
if (point < num_keys) { if (point < num_keys) {
mP[num_motion_keys] = curve_point_as_float4(b_positions, b_radius, point); const float4 cv = curve_point_as_float4(b_positions, b_radius, point);
mP[num_motion_keys] = make_float3(cv);
mR[num_motion_keys] = cv.w;
num_motion_keys++; num_motion_keys++;
if (!have_motion) { if (!have_motion) {
/* TODO: use epsilon for comparison? Was needed for particles due to /* TODO: use epsilon for comparison? Was needed for particles due to
* transform, but ideally should not happen anymore. */ * transform, but ideally should not happen anymore. */
float4 curve_key = make_float4(hair->get_position()[i]); float4 curve_key = make_float4(hair->get_position()[point]);
curve_key.w = hair->get_radius()[i]; curve_key.w = hair->get_radius()[point];
have_motion = !(mP[i] == curve_key); have_motion = !(cv == curve_key);
} }
} }
} }
@ -1028,8 +1043,10 @@ static void export_hair_curves_motion(Hair *hair,
const float step_size = curve.num_keys > 1 ? 1.0f / (curve.num_keys - 1) : 0.0f; const float step_size = curve.num_keys > 1 ? 1.0f / (curve.num_keys - 1) : 0.0f;
for (int i = 0; i < curve.num_keys; i++) { for (int i = 0; i < curve.num_keys; i++) {
const float step = i * step_size; const float step = i * step_size;
mP[num_motion_keys] = interpolate_curve_points( const float4 cv = interpolate_curve_points(
b_positions, b_radius, points.start(), points.size(), step); b_positions, b_radius, points.start(), points.size(), step);
mP[num_motion_keys] = make_float3(cv);
mR[num_motion_keys] = cv.w;
num_motion_keys++; num_motion_keys++;
} }
have_motion = true; have_motion = true;
@ -1100,12 +1117,13 @@ void BlenderSync::sync_hair(BObjectInfo &b_ob_info, Hair *hair)
{ {
new_hair.set_motion_steps(2); new_hair.set_motion_steps(2);
Attribute *attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
Attribute *new_attr_mP = new_hair.attributes.add(ATTR_STD_MOTION_VERTEX_POSITION); Attribute *new_attr_P = new_hair.attributes.find(ATTR_STD_POSITION);
if (attr_mP) { if (attr_P->has_motion()) {
new_attr_mP->set_data_from(std::move(*attr_mP)); new_attr_P->take_motion_from(*attr_P);
} }
else { else {
new_attr_P->add_motion(&new_hair);
new_hair.copy_center_to_motion_step(0); new_hair.copy_center_to_motion_step(0);
} }
} }

63
intern/cycles/blender/pointcloud.cpp
View file

@ -28,26 +28,25 @@ static void attr_create_motion_from_velocity(PointCloud *pointcloud,
/* Override motion steps to fixed number. */ /* Override motion steps to fixed number. */
pointcloud->set_motion_steps(3); pointcloud->set_motion_steps(3);
/* Find or add attribute */ /* Set motion steps on position and radius attributes. */
Attribute *attr_P = pointcloud->attributes.find(ATTR_STD_POSITION);
Attribute *attr_R = pointcloud->attributes.find(ATTR_STD_RADIUS);
attr_P->add_motion(pointcloud);
attr_R->add_motion(pointcloud);
const packed_float3 *P = pointcloud->get_position(); const packed_float3 *P = pointcloud->get_position();
const float *radius = pointcloud->get_radius(); const float *radius = pointcloud->get_radius();
Attribute *attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
if (!attr_mP) {
attr_mP = pointcloud->attributes.add(ATTR_STD_MOTION_VERTEX_POSITION);
}
/* Only export previous and next frame, we don't have any in between data. */ /* Only export previous and next frame, we don't have any in between data. */
const float motion_times[2] = {-1.0f, 1.0f}; const float motion_times[2] = {-1.0f, 1.0f};
for (int step = 0; step < 2; step++) { for (int step = 1; step <= 2; step++) {
const float relative_time = motion_times[step] * 0.5f * motion_scale; const float relative_time = motion_times[step - 1] * 0.5f * motion_scale;
float4 *mP = attr_mP->data_for_write<float4>() + step * num_points; packed_float3 *mP = attr_P->data_for_write<packed_float3>(step);
float *mR = attr_R->data_for_write<float>(step);
for (int i = 0; i < num_points; i++) { for (int i = 0; i < num_points; i++) {
const float3 Pi = float3(P[i]) + mP[i] = float3(P[i]) +
make_float3(b_attribute[i][0], b_attribute[i][1], b_attribute[i][2]) * make_float3(b_attribute[i][0], b_attribute[i][1], b_attribute[i][2]) * relative_time;
relative_time; mR[i] = radius[i];
mP[i] = make_float4(Pi, radius[i]);
} }
} }
} }
@ -141,6 +140,8 @@ static void export_pointcloud(Scene *scene,
else { else {
std::fill(radius, radius + b_positions.size(), 0.01f); std::fill(radius, radius + b_positions.size(), 0.01f);
} }
pointcloud->tag_position_modified();
pointcloud->tag_radius_modified();
int *shader = pointcloud->get_shader().data(); int *shader = pointcloud->get_shader().data();
std::fill(shader, shader + b_positions.size(), 0); std::fill(shader, shader + b_positions.size(), 0);
@ -160,22 +161,23 @@ static void export_pointcloud_motion(PointCloud *pointcloud,
const blender::PointCloud &b_pointcloud, const blender::PointCloud &b_pointcloud,
const int motion_step) const int motion_step)
{ {
/* Find or add attribute. */ /* Set motion steps on position and radius attributes. */
Attribute *attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); Attribute *attr_P = pointcloud->attributes.find(ATTR_STD_POSITION);
Attribute *attr_R = pointcloud->attributes.find(ATTR_STD_RADIUS);
bool new_attribute = false; bool new_attribute = false;
if (!attr_mP) { if (!attr_P->has_motion()) {
attr_mP = pointcloud->attributes.add(ATTR_STD_MOTION_VERTEX_POSITION); attr_P->add_motion(pointcloud);
attr_R->add_motion(pointcloud);
new_attribute = true; new_attribute = true;
} }
const int num_points = pointcloud->num_points(); const int num_points = pointcloud->num_points();
/* Point cloud attributes are stored as float4 with the radius in the w element. const int attr_step = motion_step + 1;
* This is explicit now as float3 is no longer interchangeable with float4 as it packed_float3 *mP = attr_P->data_for_write<packed_float3>(attr_step);
* is packed now. */ float *mR = attr_R->data_for_write<float>(attr_step);
float4 *mP = attr_mP->data_for_write<float4>() + motion_step * num_points;
bool have_motion = false; bool have_motion = false;
const packed_float3 *pointcloud_positions = pointcloud->get_position(); const packed_float3 *pointcloud_points = pointcloud->get_position();
const blender::Span<blender::float3> b_positions = b_pointcloud.positions(); const blender::Span<blender::float3> b_positions = b_pointcloud.positions();
const blender::VArraySpan b_radius = *b_pointcloud.attributes().lookup<float>( const blender::VArraySpan b_radius = *b_pointcloud.attributes().lookup<float>(
@ -184,14 +186,16 @@ static void export_pointcloud_motion(PointCloud *pointcloud,
for (int i = 0; i < std::min<int>(num_points, b_positions.size()); i++) { for (int i = 0; i < std::min<int>(num_points, b_positions.size()); i++) {
const float3 P = make_float3(b_positions[i][0], b_positions[i][1], b_positions[i][2]); const float3 P = make_float3(b_positions[i][0], b_positions[i][1], b_positions[i][2]);
const float radius = b_radius.is_empty() ? 0.01f : b_radius[i]; const float radius = b_radius.is_empty() ? 0.01f : b_radius[i];
mP[i] = make_float4(P, radius); mP[i] = P;
have_motion = have_motion || (P != pointcloud_positions[i]); mR[i] = radius;
have_motion = have_motion || (P != pointcloud_points[i]);
} }
/* In case of new attribute, we verify if there really was any motion. */ /* In case of new attribute, we verify if there really was any motion. */
if (new_attribute) { if (new_attribute) {
if (b_positions.size() != num_points || !have_motion) { if (b_positions.size() != num_points || !have_motion) {
pointcloud->attributes.remove(ATTR_STD_MOTION_VERTEX_POSITION); attr_P->remove_motion();
attr_R->remove_motion();
} }
else if (motion_step > 0) { else if (motion_step > 0) {
/* Motion, fill up previous steps that we might have skipped because /* Motion, fill up previous steps that we might have skipped because
@ -230,12 +234,13 @@ void BlenderSync::sync_pointcloud(PointCloud *pointcloud, BObjectInfo &b_ob_info
{ {
new_pointcloud.set_motion_steps(2); new_pointcloud.set_motion_steps(2);
Attribute *attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); Attribute *attr_P = pointcloud->attributes.find(ATTR_STD_POSITION);
Attribute *new_attr_mP = new_pointcloud.attributes.add(ATTR_STD_MOTION_VERTEX_POSITION); Attribute *new_attr_P = new_pointcloud.attributes.find(ATTR_STD_POSITION);
if (attr_mP) { if (attr_P->has_motion()) {
new_attr_mP->set_data_from(std::move(*attr_mP)); new_attr_P->take_motion_from(*attr_P);
} }
else { else {
new_attr_P->add_motion(&new_pointcloud);
new_pointcloud.copy_center_to_motion_step(0); new_pointcloud.copy_center_to_motion_step(0);
} }
} }

73
intern/cycles/bvh/build.cpp
View file

@ -141,10 +141,9 @@ void BVHBuild::add_reference_curves(BoundBox &root,
Hair *hair, Hair *hair,
const int object_index) const int object_index)
{ {
const Attribute *curve_attr_mP = nullptr; const Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
if (hair->has_motion_blur()) { const Attribute *attr_R = hair->attributes.find(ATTR_STD_RADIUS);
curve_attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); const bool has_motion = attr_P->has_motion();
}
const PrimitiveType primitive_type = hair->primitive_type(); const PrimitiveType primitive_type = hair->primitive_type();
@ -153,7 +152,7 @@ void BVHBuild::add_reference_curves(BoundBox &root,
const Hair::Curve curve = hair->get_curve(j); const Hair::Curve curve = hair->get_curve(j);
const float *curve_radius = hair->get_radius(); const float *curve_radius = hair->get_radius();
for (int k = 0; k < curve.num_keys - 1; k++) { for (int k = 0; k < curve.num_keys - 1; k++) {
if (curve_attr_mP == nullptr) { if (!has_motion) {
/* Really simple logic for static hair. */ /* Really simple logic for static hair. */
BoundBox bounds = BoundBox::empty; BoundBox bounds = BoundBox::empty;
curve.bounds_grow(k, hair->get_position(), curve_radius, bounds); curve.bounds_grow(k, hair->get_position(), curve_radius, bounds);
@ -165,18 +164,15 @@ void BVHBuild::add_reference_curves(BoundBox &root,
} }
} }
else if (params.num_motion_curve_steps == 0 || params.use_spatial_split) { else if (params.num_motion_curve_steps == 0 || params.use_spatial_split) {
/* Simple case of motion curves: single node for the while /* Simple case of motion curves: single node for the whole
* shutter time. Lowest memory usage but less optimal * shutter time. Lowest memory usage but less optimal
* rendering. * rendering.
*/ */
/* TODO(sergey): Support motion steps for spatially split BVH. */ /* TODO(sergey): Support motion steps for spatially split BVH. */
BoundBox bounds = BoundBox::empty; BoundBox bounds = BoundBox::empty;
curve.bounds_grow(k, hair->get_position(), curve_radius, bounds); for (int attr_step = 0; attr_step < attr_P->num_motion_steps(); attr_step++) {
const size_t num_keys = hair->num_keys(); curve.bounds_grow(
const size_t num_steps = hair->get_motion_steps(); k, attr_P->data<packed_float3>(attr_step), attr_R->data<float>(attr_step), bounds);
const float4 *key_steps = curve_attr_mP->data<float4>();
for (size_t step = 0; step < num_steps - 1; step++) {
curve.bounds_grow(k, key_steps + step * num_keys, bounds);
} }
if (bounds.valid()) { if (bounds.valid()) {
const int packed_type = PRIMITIVE_PACK_SEGMENT(primitive_type, k); const int packed_type = PRIMITIVE_PACK_SEGMENT(primitive_type, k);
@ -193,42 +189,21 @@ void BVHBuild::add_reference_curves(BoundBox &root,
const int num_bvh_steps = params.num_motion_curve_steps * 2 + 1; const int num_bvh_steps = params.num_motion_curve_steps * 2 + 1;
const float num_bvh_steps_inv_1 = 1.0f / (num_bvh_steps - 1); const float num_bvh_steps_inv_1 = 1.0f / (num_bvh_steps - 1);
const size_t num_steps = hair->get_motion_steps(); const size_t num_steps = hair->get_motion_steps();
const packed_float3 *curve_keys = hair->get_position();
const float4 *key_steps = curve_attr_mP->data<float4>();
const size_t num_keys = hair->num_keys();
/* Calculate bounding box of the previous time step. /* Calculate bounding box of the previous time step.
* Will be reused later to avoid duplicated work on * Will be reused later to avoid duplicated work on
* calculating BVH time step boundbox. * calculating BVH time step boundbox.
*/ */
float4 prev_keys[4]; float4 prev_keys[4];
curve.cardinal_motion_keys(curve_keys, curve.cardinal_motion_keys(
curve_radius, attr_P, attr_R, num_steps, 0.0f, k - 1, k, k + 1, k + 2, prev_keys);
key_steps,
num_keys,
num_steps,
0.0f,
k - 1,
k,
k + 1,
k + 2,
prev_keys);
BoundBox prev_bounds = BoundBox::empty; BoundBox prev_bounds = BoundBox::empty;
curve.bounds_grow(prev_keys, prev_bounds); curve.bounds_grow(prev_keys, prev_bounds);
/* Create all primitive time steps, */ /* Create all primitive time steps, */
for (int bvh_step = 1; bvh_step < num_bvh_steps; ++bvh_step) { for (int bvh_step = 1; bvh_step < num_bvh_steps; ++bvh_step) {
const float curr_time = (float)(bvh_step)*num_bvh_steps_inv_1; const float curr_time = (float)(bvh_step)*num_bvh_steps_inv_1;
float4 curr_keys[4]; float4 curr_keys[4];
curve.cardinal_motion_keys(curve_keys, curve.cardinal_motion_keys(
curve_radius, attr_P, attr_R, num_steps, curr_time, k - 1, k, k + 1, k + 2, curr_keys);
key_steps,
num_keys,
num_steps,
curr_time,
k - 1,
k,
k + 1,
k + 2,
curr_keys);
BoundBox curr_bounds = BoundBox::empty; BoundBox curr_bounds = BoundBox::empty;
curve.bounds_grow(curr_keys, curr_bounds); curve.bounds_grow(curr_keys, curr_bounds);
BoundBox bounds = prev_bounds; BoundBox bounds = prev_bounds;
@ -256,18 +231,16 @@ void BVHBuild::add_reference_points(BoundBox &root,
PointCloud *pointcloud, PointCloud *pointcloud,
const int i) const int i)
{ {
const Attribute *point_attr_mP = nullptr; const Attribute *attr_P = pointcloud->attributes.find(ATTR_STD_POSITION);
if (pointcloud->has_motion_blur()) { const Attribute *attr_R = pointcloud->attributes.find(ATTR_STD_RADIUS);
point_attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); const bool has_motion = attr_P->has_motion();
}
const packed_float3 *points_data = pointcloud->get_position(); const packed_float3 *points_data = pointcloud->get_position();
const float *radius_data = pointcloud->get_radius(); const float *radius_data = pointcloud->get_radius();
const size_t num_points = pointcloud->num_points(); const size_t num_points = pointcloud->num_points();
const float4 *motion_data = (point_attr_mP) ? point_attr_mP->data<float4>() : nullptr;
const size_t num_steps = pointcloud->get_motion_steps(); const size_t num_steps = pointcloud->get_motion_steps();
if (point_attr_mP == nullptr) { if (!has_motion) {
/* Really simple logic for static points. */ /* Really simple logic for static points. */
for (uint j = 0; j < num_points; j++) { for (uint j = 0; j < num_points; j++) {
const PointCloud::Point point = pointcloud->get_point(j); const PointCloud::Point point = pointcloud->get_point(j);
@ -289,9 +262,10 @@ void BVHBuild::add_reference_points(BoundBox &root,
for (uint j = 0; j < num_points; j++) { for (uint j = 0; j < num_points; j++) {
const PointCloud::Point point = pointcloud->get_point(j); const PointCloud::Point point = pointcloud->get_point(j);
BoundBox bounds = BoundBox::empty; BoundBox bounds = BoundBox::empty;
point.bounds_grow(points_data, radius_data, bounds); for (int attr_step = 0; attr_step < attr_P->num_motion_steps(); attr_step++) {
for (size_t step = 0; step < num_steps - 1; step++) { const packed_float3 *mP = attr_P->data<packed_float3>(attr_step);
point.bounds_grow(motion_data[step * num_points + j], bounds); const float *mR = attr_R->data<float>(attr_step);
point.bounds_grow(make_float4(float3(mP[j]), mR[j]), bounds);
} }
if (bounds.valid()) { if (bounds.valid()) {
references.push_back(BVHReference(bounds, j, i, PRIMITIVE_MOTION_POINT)); references.push_back(BVHReference(bounds, j, i, PRIMITIVE_MOTION_POINT));
@ -307,8 +281,7 @@ void BVHBuild::add_reference_points(BoundBox &root,
*/ */
const int num_bvh_steps = params.num_motion_point_steps * 2 + 1; 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); const float num_bvh_steps_inv_1 = 1.0f / (num_bvh_steps - 1);
const size_t num_steps = pointcloud->get_motion_steps(); const Attribute *attr_R_motion = attr_R->has_motion() ? attr_R : nullptr;
const float4 *point_steps = point_attr_mP->data<float4>();
for (uint j = 0; j < num_points; j++) { for (uint j = 0; j < num_points; j++) {
const PointCloud::Point point = pointcloud->get_point(j); const PointCloud::Point point = pointcloud->get_point(j);
@ -318,14 +291,14 @@ void BVHBuild::add_reference_points(BoundBox &root,
* calculating BVH time step boundbox. * calculating BVH time step boundbox.
*/ */
const float4 prev_key = point.motion_key( const float4 prev_key = point.motion_key(
points_data, radius_data, point_steps, num_points, num_steps, 0.0f, j); radius_data, attr_P, attr_R_motion, num_steps, 0.0f, j);
BoundBox prev_bounds = BoundBox::empty; BoundBox prev_bounds = BoundBox::empty;
point.bounds_grow(prev_key, prev_bounds); point.bounds_grow(prev_key, prev_bounds);
/* Create all primitive time steps, */ /* Create all primitive time steps, */
for (int bvh_step = 1; bvh_step < num_bvh_steps; ++bvh_step) { for (int bvh_step = 1; bvh_step < num_bvh_steps; ++bvh_step) {
const float curr_time = (float)(bvh_step)*num_bvh_steps_inv_1; const float curr_time = (float)(bvh_step)*num_bvh_steps_inv_1;
const float4 curr_key = point.motion_key( const float4 curr_key = point.motion_key(
points_data, radius_data, point_steps, num_points, num_steps, curr_time, j); radius_data, attr_P, attr_R_motion, num_steps, curr_time, j);
BoundBox curr_bounds = BoundBox::empty; BoundBox curr_bounds = BoundBox::empty;
point.bounds_grow(curr_key, curr_bounds); point.bounds_grow(curr_key, curr_bounds);
BoundBox bounds = prev_bounds; BoundBox bounds = prev_bounds;

30
intern/cycles/bvh/bvh2.cpp
View file

@ -380,15 +380,13 @@ void BVH2::refit_primitives(const int start, const int end, BoundBox &bbox, uint
/* Motion curves. */ /* Motion curves. */
if (hair->get_use_motion_blur()) { if (hair->get_use_motion_blur()) {
Attribute *attr = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); const Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
const Attribute *attr_R = hair->attributes.find(ATTR_STD_RADIUS);
if (attr) { if (attr_P->has_motion()) {
const size_t hair_size = hair->num_keys(); for (int attr_step = 1; attr_step < attr_P->num_motion_steps(); attr_step++) {
const size_t steps = hair->get_motion_steps() - 1; curve.bounds_grow(
const packed_float3 *key_steps = attr->data<packed_float3>(); k, attr_P->data<packed_float3>(attr_step), attr_R->data<float>(attr_step), bbox);
for (size_t i = 0; i < steps; i++) {
curve.bounds_grow(k, key_steps + i * hair_size, hair->get_radius(), bbox);
} }
} }
} }
@ -405,17 +403,13 @@ void BVH2::refit_primitives(const int start, const int end, BoundBox &bbox, uint
/* Motion points. */ /* Motion points. */
if (pointcloud->get_use_motion_blur()) { if (pointcloud->get_use_motion_blur()) {
Attribute *attr = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); const Attribute *attr_P = pointcloud->attributes.find(ATTR_STD_POSITION);
if (attr) { if (attr_P->has_motion()) {
const size_t pointcloud_size = pointcloud->num_points(); const Attribute *attr_R = pointcloud->attributes.find(ATTR_STD_RADIUS);
const size_t steps = pointcloud->get_motion_steps() - 1; for (int attr_step = 1; attr_step < attr_P->num_motion_steps(); attr_step++) {
const float4 *point_steps = attr->data<float4>(); const float3 P = attr_P->data<packed_float3>(attr_step)[point.index];
const float r = attr_R->data<float>(attr_step)[point.index];
for (size_t i = 0; i < steps; i++) {
const size_t idx = i * pointcloud_size + point.index;
const float3 P = make_float3(point_steps[idx]);
const float r = point_steps[idx].w;
bbox.grow(P, r); bbox.grow(P, r);
} }
} }

82
intern/cycles/bvh/embree.cpp
View file

@ -488,13 +488,11 @@ void pack_motion_verts(const size_t num_curves,
void BVHEmbree::set_curve_vertex_buffer(RTCGeometry geom_id, const Hair *hair, const bool update) void BVHEmbree::set_curve_vertex_buffer(RTCGeometry geom_id, const Hair *hair, const bool update)
{ {
const Attribute *attr_mP = nullptr; const Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
const Attribute *attr_R = hair->attributes.find(ATTR_STD_RADIUS);
size_t num_motion_steps = 1; size_t num_motion_steps = 1;
if (hair->has_motion_blur()) { if (hair->has_motion_blur() && attr_P->has_motion()) {
attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); num_motion_steps = hair->get_motion_steps();
if (attr_mP) {
num_motion_steps = hair->get_motion_steps();
}
} }
const size_t num_curves = hair->num_curves(); const size_t num_curves = hair->num_curves();
@ -509,13 +507,7 @@ void BVHEmbree::set_curve_vertex_buffer(RTCGeometry geom_id, const Hair *hair, c
num_keys_embree += num_curves * 2; num_keys_embree += num_curves * 2;
/* Copy the CV data to Embree */ /* Copy the CV data to Embree */
const int t_mid = (num_motion_steps - 1) / 2;
const float *curve_radius = hair->get_radius();
for (int t = 0; t < num_motion_steps; ++t) { for (int t = 0; t < num_motion_steps; ++t) {
// As float4 and float3 are no longer interchangeable the 2 types need to be
// handled separately. Attributes are float4s where the radius is stored in w and
// the middle motion vector is from the mesh points which are stored float3s with
// the radius stored in another array.
float4 *rtc_verts = nullptr; float4 *rtc_verts = nullptr;
if (update) { if (update) {
rtc_verts = (float4 *)rtcGetGeometryBufferData(geom_id, RTC_BUFFER_TYPE_VERTEX, t); rtc_verts = (float4 *)rtcGetGeometryBufferData(geom_id, RTC_BUFFER_TYPE_VERTEX, t);
@ -543,17 +535,13 @@ void BVHEmbree::set_curve_vertex_buffer(RTCGeometry geom_id, const Hair *hair, c
assert(rtc_verts); assert(rtc_verts);
if (rtc_verts) { if (rtc_verts) {
const size_t num_curves = hair->num_curves(); const size_t num_curves = hair->num_curves();
if (t == t_mid || attr_mP == nullptr) { pack_motion_verts<packed_float3>(
const packed_float3 *verts = hair->get_position(); num_curves,
pack_motion_verts<packed_float3>( hair,
num_curves, hair, verts, curve_radius, rtc_verts, hair->curve_shape); attr_P->data_at_time_step<packed_float3>(t, num_motion_steps),
} attr_R->data_at_time_step<float>(t, num_motion_steps),
else { rtc_verts,
const int t_ = (t > t_mid) ? (t - 1) : t; hair->curve_shape);
const float4 *verts = &attr_mP->data<float4>()[t_ * num_keys];
pack_motion_verts<float4>(
num_curves, hair, verts, curve_radius, rtc_verts, hair->curve_shape);
}
} }
if (update) { if (update) {
@ -566,25 +554,17 @@ void BVHEmbree::set_point_vertex_buffer(RTCGeometry geom_id,
const PointCloud *pointcloud, const PointCloud *pointcloud,
const bool update) const bool update)
{ {
const Attribute *attr_mP = nullptr; const Attribute *attr_P = pointcloud->attributes.find(ATTR_STD_POSITION);
const Attribute *attr_R = pointcloud->attributes.find(ATTR_STD_RADIUS);
size_t num_motion_steps = 1; size_t num_motion_steps = 1;
if (pointcloud->has_motion_blur()) { if (pointcloud->has_motion_blur() && attr_P->has_motion()) {
attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); num_motion_steps = pointcloud->get_motion_steps();
if (attr_mP) {
num_motion_steps = pointcloud->get_motion_steps();
}
} }
const size_t num_points = pointcloud->num_points(); const size_t num_points = pointcloud->num_points();
/* Copy the point data to Embree */ /* Copy the point data to Embree. */
const int t_mid = (num_motion_steps - 1) / 2;
const float *radius = pointcloud->get_radius();
for (int t = 0; t < num_motion_steps; ++t) { for (int t = 0; t < num_motion_steps; ++t) {
// As float4 and float3 are no longer interchangeable the 2 types need to be
// handled separately. Attributes are float4s where the radius is stored in w and
// the middle motion vector is from the mesh points which are stored float3s with
// the radius stored in another array.
float4 *rtc_verts = nullptr; float4 *rtc_verts = nullptr;
if (update) { if (update) {
@ -612,18 +592,10 @@ void BVHEmbree::set_point_vertex_buffer(RTCGeometry geom_id,
assert(rtc_verts); assert(rtc_verts);
if (rtc_verts) { if (rtc_verts) {
if (t == t_mid || attr_mP == nullptr) { const packed_float3 *verts = attr_P->data_at_time_step<packed_float3>(t, num_motion_steps);
/* Pack the motion points into a float4 as [x y z radius]. */ const float *radius = attr_R->data_at_time_step<float>(t, num_motion_steps);
const packed_float3 *verts = pointcloud->get_position(); for (size_t j = 0; j < num_points; ++j) {
for (size_t j = 0; j < num_points; ++j) { rtc_verts[j] = make_float4(float3(verts[j]), radius[j]);
rtc_verts[j] = make_float4(float3(verts[j]), radius[j]);
}
}
else {
/* Motion blur is already packed as [x y z radius]. */
const int t_ = (t > t_mid) ? (t - 1) : t;
const float4 *verts = &attr_mP->data<float4>()[t_ * num_points];
std::copy_n(verts, num_points, rtc_verts);
} }
} }
@ -637,11 +609,10 @@ void BVHEmbree::add_points(const Object *ob, const PointCloud *pointcloud, const
{ {
const size_t prim_offset = pointcloud->prim_offset; const size_t prim_offset = pointcloud->prim_offset;
const Attribute *attr_mP = nullptr;
size_t num_motion_steps = 1; size_t num_motion_steps = 1;
if (pointcloud->has_motion_blur()) { if (pointcloud->has_motion_blur()) {
attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); const Attribute *attr_P = pointcloud->attributes.find(ATTR_STD_POSITION);
if (attr_mP) { if (attr_P->has_motion()) {
num_motion_steps = pointcloud->get_motion_steps(); num_motion_steps = pointcloud->get_motion_steps();
} }
} }
@ -668,13 +639,10 @@ void BVHEmbree::add_curves(const Object *ob, const Hair *hair, const int i)
{ {
const size_t prim_offset = hair->curve_segment_offset; const size_t prim_offset = hair->curve_segment_offset;
const Attribute *attr_mP = nullptr; const Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
size_t num_motion_steps = 1; size_t num_motion_steps = 1;
if (hair->has_motion_blur()) { if (hair->has_motion_blur() && attr_P->has_motion()) {
attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); num_motion_steps = hair->get_motion_steps();
if (attr_mP) {
num_motion_steps = hair->get_motion_steps();
}
} }
assert(num_motion_steps <= RTC_MAX_TIME_STEP_COUNT); assert(num_motion_steps <= RTC_MAX_TIME_STEP_COUNT);

75
intern/cycles/device/hiprt/device_impl.cpp
View file

@ -537,13 +537,11 @@ hiprtGeometryBuildInput HIPRTDevice::prepare_curve_blas(BVHHIPRT *bvh, Hair *hai
const PrimitiveType primitive_type = hair->primitive_type(); const PrimitiveType primitive_type = hair->primitive_type();
const size_t num_curves = hair->num_curves(); const size_t num_curves = hair->num_curves();
const size_t num_segments = hair->num_segments(); const size_t num_segments = hair->num_segments();
const Attribute *curve_attr_mP = nullptr; const Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
const Attribute *attr_R = hair->attributes.find(ATTR_STD_RADIUS);
const bool has_motion = use_motion_blur && hair->has_motion_blur() && attr_P->has_motion();
if (use_motion_blur && hair->has_motion_blur()) { if (!has_motion || bvh->params.num_motion_curve_steps == 0) {
curve_attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
}
if (curve_attr_mP == nullptr || bvh->params.num_motion_curve_steps == 0) {
bvh->custom_prim_info.resize(num_segments); bvh->custom_prim_info.resize(num_segments);
bvh->custom_primitive_bound.alloc(num_segments); bvh->custom_primitive_bound.alloc(num_segments);
} }
@ -563,7 +561,7 @@ hiprtGeometryBuildInput HIPRTDevice::prepare_curve_blas(BVHHIPRT *bvh, Hair *hai
const float *curve_radius = hair->get_radius(); const float *curve_radius = hair->get_radius();
int first_key = curve.first_key; int first_key = curve.first_key;
for (int k = 0; k < curve.num_keys - 1; k++) { for (int k = 0; k < curve.num_keys - 1; k++) {
if (curve_attr_mP == nullptr) { if (!has_motion) {
float3 current_keys[4]; float3 current_keys[4];
current_keys[0] = curve_keys[max(first_key + k - 1, first_key)]; current_keys[0] = curve_keys[max(first_key + k - 1, first_key)];
current_keys[1] = curve_keys[first_key + k]; current_keys[1] = curve_keys[first_key + k];
@ -592,14 +590,12 @@ hiprtGeometryBuildInput HIPRTDevice::prepare_curve_blas(BVHHIPRT *bvh, Hair *hai
} }
else { else {
const size_t num_steps = hair->get_motion_steps(); const size_t num_steps = hair->get_motion_steps();
const float4 *key_steps = curve_attr_mP->data<float4>();
const size_t num_keys = hair->num_keys();
if (bvh->params.num_motion_curve_steps == 0 || bvh->params.use_spatial_split) { if (bvh->params.num_motion_curve_steps == 0 || bvh->params.use_spatial_split) {
BoundBox bounds = BoundBox::empty; BoundBox bounds = BoundBox::empty;
curve.bounds_grow(k, hair->get_position(), curve_radius, bounds); for (int attr_step = 0; attr_step < attr_P->num_motion_steps(); attr_step++) {
for (size_t step = 0; step < num_steps - 1; step++) { curve.bounds_grow(
curve.bounds_grow(k, key_steps + step * num_keys, bounds); k, attr_P->data<packed_float3>(attr_step), attr_R->data<float>(attr_step), bounds);
} }
if (bounds.valid()) { if (bounds.valid()) {
int type = PRIMITIVE_PACK_SEGMENT(primitive_type, k); int type = PRIMITIVE_PACK_SEGMENT(primitive_type, k);
@ -615,34 +611,16 @@ hiprtGeometryBuildInput HIPRTDevice::prepare_curve_blas(BVHHIPRT *bvh, Hair *hai
const float num_bvh_steps_inv_1 = 1.0f / (num_bvh_steps - 1); const float num_bvh_steps_inv_1 = 1.0f / (num_bvh_steps - 1);
float4 prev_keys[4]; float4 prev_keys[4];
curve.cardinal_motion_keys(curve_keys, curve.cardinal_motion_keys(
curve_radius, attr_P, attr_R, num_steps, 0.0f, k - 1, k, k + 1, k + 2, prev_keys);
key_steps,
num_keys,
num_steps,
0.0f,
k - 1,
k,
k + 1,
k + 2,
prev_keys);
BoundBox prev_bounds = BoundBox::empty; BoundBox prev_bounds = BoundBox::empty;
curve.bounds_grow(prev_keys, prev_bounds); curve.bounds_grow(prev_keys, prev_bounds);
for (int bvh_step = 1; bvh_step < num_bvh_steps; ++bvh_step) { for (int bvh_step = 1; bvh_step < num_bvh_steps; ++bvh_step) {
const float curr_time = (float)(bvh_step)*num_bvh_steps_inv_1; const float curr_time = (float)(bvh_step)*num_bvh_steps_inv_1;
float4 curr_keys[4]; float4 curr_keys[4];
curve.cardinal_motion_keys(curve_keys, curve.cardinal_motion_keys(
curve_radius, attr_P, attr_R, num_steps, curr_time, k - 1, k, k + 1, k + 2, curr_keys);
key_steps,
num_keys,
num_steps,
curr_time,
k - 1,
k,
k + 1,
k + 2,
curr_keys);
BoundBox curr_bounds = BoundBox::empty; BoundBox curr_bounds = BoundBox::empty;
curve.bounds_grow(curr_keys, curr_bounds); curve.bounds_grow(curr_keys, curr_bounds);
BoundBox bounds = prev_bounds; BoundBox bounds = prev_bounds;
@ -688,21 +666,26 @@ hiprtGeometryBuildInput HIPRTDevice::prepare_point_blas(BVHHIPRT *bvh, PointClou
{ {
hiprtGeometryBuildInput geom_input; hiprtGeometryBuildInput geom_input;
const Attribute *point_attr_mP = nullptr; const Attribute *attr_P = nullptr;
const Attribute *attr_R = nullptr;
if (use_motion_blur && pointcloud->has_motion_blur()) { if (use_motion_blur && pointcloud->has_motion_blur()) {
point_attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); attr_P = pointcloud->attributes.find(ATTR_STD_POSITION);
attr_R = pointcloud->attributes.find(ATTR_STD_RADIUS);
if (!attr_P->has_motion()) {
attr_P = nullptr;
attr_R = nullptr;
}
} }
const packed_float3 *points_data = pointcloud->get_position(); const packed_float3 *points_data = pointcloud->get_position();
const float *radius_data = pointcloud->get_radius(); const float *radius_data = pointcloud->get_radius();
const size_t num_points = pointcloud->num_points(); const size_t num_points = pointcloud->num_points();
const float4 *motion_data = (point_attr_mP) ? point_attr_mP->data<float4>() : nullptr; const bool has_motion_radius = attr_R && attr_R->has_motion();
const size_t num_steps = pointcloud->get_motion_steps();
int num_bounds = 0; int num_bounds = 0;
float sum_area = 0.0f; float sum_area = 0.0f;
if (point_attr_mP == nullptr) { if (attr_P == nullptr) {
bvh->custom_prim_info.resize(num_points); bvh->custom_prim_info.resize(num_points);
bvh->custom_primitive_bound.alloc(num_points); bvh->custom_primitive_bound.alloc(num_points);
for (uint j = 0; j < num_points; j++) { for (uint j = 0; j < num_points; j++) {
@ -725,9 +708,9 @@ hiprtGeometryBuildInput HIPRTDevice::prepare_point_blas(BVHHIPRT *bvh, PointClou
for (uint j = 0; j < num_points; j++) { for (uint j = 0; j < num_points; j++) {
const PointCloud::Point point = pointcloud->get_point(j); const PointCloud::Point point = pointcloud->get_point(j);
BoundBox bounds = BoundBox::empty; BoundBox bounds = BoundBox::empty;
point.bounds_grow(points_data, radius_data, bounds); for (int attr_step = 0; attr_step < attr_P->num_motion_steps(); attr_step++) {
for (size_t step = 0; step < num_steps - 1; step++) { point.bounds_grow(
point.bounds_grow(motion_data[step * num_points + j], bounds); attr_P->data<packed_float3>(attr_step), attr_R->data<float>(attr_step), bounds);
} }
if (bounds.valid()) { if (bounds.valid()) {
bvh->custom_primitive_bound[num_bounds] = bounds; bvh->custom_primitive_bound[num_bounds] = bounds;
@ -746,20 +729,20 @@ hiprtGeometryBuildInput HIPRTDevice::prepare_point_blas(BVHHIPRT *bvh, PointClou
bvh->custom_primitive_bound.alloc(num_points * num_bvh_steps); bvh->custom_primitive_bound.alloc(num_points * num_bvh_steps);
bvh->prims_time.resize(num_points * num_bvh_steps); bvh->prims_time.resize(num_points * num_bvh_steps);
const Attribute *attr_R_motion = has_motion_radius ? attr_R : nullptr;
for (uint j = 0; j < num_points; j++) { for (uint j = 0; j < num_points; j++) {
const PointCloud::Point point = pointcloud->get_point(j); const PointCloud::Point point = pointcloud->get_point(j);
const size_t num_steps = pointcloud->get_motion_steps(); const size_t num_steps = pointcloud->get_motion_steps();
const float4 *point_steps = point_attr_mP->data<float4>();
float4 prev_key = point.motion_key( float4 prev_key = point.motion_key(radius_data, attr_P, attr_R_motion, num_steps, 0.0f, j);
points_data, radius_data, point_steps, num_points, num_steps, 0.0f, j);
BoundBox prev_bounds = BoundBox::empty; BoundBox prev_bounds = BoundBox::empty;
point.bounds_grow(prev_key, prev_bounds); point.bounds_grow(prev_key, prev_bounds);
for (int bvh_step = 1; bvh_step < num_bvh_steps; ++bvh_step) { for (int bvh_step = 1; bvh_step < num_bvh_steps; ++bvh_step) {
const float curr_time = (float)(bvh_step)*num_bvh_steps_inv_1; const float curr_time = (float)(bvh_step)*num_bvh_steps_inv_1;
float4 curr_key = point.motion_key( float4 curr_key = point.motion_key(
points_data, radius_data, point_steps, num_points, num_steps, curr_time, j); radius_data, attr_P, attr_R_motion, num_steps, curr_time, j);
BoundBox curr_bounds = BoundBox::empty; BoundBox curr_bounds = BoundBox::empty;
point.bounds_grow(curr_key, curr_bounds); point.bounds_grow(curr_key, curr_bounds);
BoundBox bounds = prev_bounds; BoundBox bounds = prev_bounds;

74
intern/cycles/device/metal/bvh.mm
View file

@ -401,8 +401,9 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
const bool use_fast_trace_bvh = (params.bvh_type == BVH_TYPE_STATIC) || !support_refit_blas(); const bool use_fast_trace_bvh = (params.bvh_type == BVH_TYPE_STATIC) || !support_refit_blas();
size_t num_motion_steps = 1; size_t num_motion_steps = 1;
Attribute *motion_keys = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); const Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
if (motion_blur && hair->get_use_motion_blur() && motion_keys) { const Attribute *attr_R = hair->attributes.find(ATTR_STD_RADIUS);
if (motion_blur && hair->get_use_motion_blur() && attr_P->has_motion()) {
num_motion_steps = hair->get_motion_steps(); num_motion_steps = hair->get_motion_steps();
} }
@ -417,10 +418,9 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
uint64_t numKeys = hair->num_keys(); uint64_t numKeys = hair->num_keys();
uint64_t numCurves = hair->num_curves(); uint64_t numCurves = hair->num_curves();
const float *radiuses = hair->get_radius();
/* Gather the curve geometry. */ /* Gather the curve geometry. */
std::vector<float3> cpData; std::vector<packed_float3> cpData;
std::vector<int> idxData; std::vector<int> idxData;
std::vector<float> radiusData; std::vector<float> radiusData;
cpData.reserve(numKeys); cpData.reserve(numKeys);
@ -428,14 +428,9 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
std::vector<int> step_offsets; std::vector<int> step_offsets;
for (size_t step = 0; step < num_motion_steps; ++step) { for (size_t step = 0; step < num_motion_steps; ++step) {
const packed_float3 *keys = attr_P->data_at_time_step<packed_float3>(step,
/* The center step for motion vertices is not stored in the attribute. */ num_motion_steps);
const packed_float3 *keys = hair->get_position(); const float *radii = attr_R->data_at_time_step<float>(step, num_motion_steps);
size_t center_step = (num_motion_steps - 1) / 2;
if (step != center_step) {
size_t attr_offset = (step > center_step) ? step - 1 : step;
keys = motion_keys->data<packed_float3>() + attr_offset * numKeys;
}
step_offsets.push_back(cpData.size()); step_offsets.push_back(cpData.size());
@ -446,20 +441,20 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
uint64_t idxBase = cpData.size(); uint64_t idxBase = cpData.size();
if (hair->curve_shape != CURVE_THICK_LINEAR) { if (hair->curve_shape != CURVE_THICK_LINEAR) {
cpData.push_back(keys[firstKey]); cpData.push_back(keys[firstKey]);
radiusData.push_back(radiuses[firstKey]); radiusData.push_back(radii[firstKey]);
} }
for (int s = 0; s < segCount; ++s) { for (int s = 0; s < segCount; ++s) {
if (step == 0) { if (step == 0) {
idxData.push_back(idxBase + s); idxData.push_back(idxBase + s);
} }
cpData.push_back(keys[firstKey + s]); cpData.push_back(keys[firstKey + s]);
radiusData.push_back(radiuses[firstKey + s]); radiusData.push_back(radii[firstKey + s]);
} }
cpData.push_back(keys[firstKey + curve.num_keys - 1]); cpData.push_back(keys[firstKey + curve.num_keys - 1]);
radiusData.push_back(radiuses[firstKey + curve.num_keys - 1]); radiusData.push_back(radii[firstKey + curve.num_keys - 1]);
if (hair->curve_shape != CURVE_THICK_LINEAR) { if (hair->curve_shape != CURVE_THICK_LINEAR) {
cpData.push_back(keys[firstKey + curve.num_keys - 1]); cpData.push_back(keys[firstKey + curve.num_keys - 1]);
radiusData.push_back(radiuses[firstKey + curve.num_keys - 1]); radiusData.push_back(radii[firstKey + curve.num_keys - 1]);
} }
} }
} }
@ -470,7 +465,7 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
options:MTLResourceStorageModeShared]; options:MTLResourceStorageModeShared];
cpBuffer = [mtl_device newBufferWithBytes:cpData.data() cpBuffer = [mtl_device newBufferWithBytes:cpData.data()
length:cpData.size() * sizeof(float3) length:cpData.size() * sizeof(packed_float3)
options:MTLResourceStorageModeShared]; options:MTLResourceStorageModeShared];
radiusBuffer = [mtl_device newBufferWithBytes:radiusData.data() radiusBuffer = [mtl_device newBufferWithBytes:radiusData.data()
@ -485,7 +480,7 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
for (size_t step = 0; step < num_motion_steps; ++step) { for (size_t step = 0; step < num_motion_steps; ++step) {
MTLMotionKeyframeData *k = [MTLMotionKeyframeData data]; MTLMotionKeyframeData *k = [MTLMotionKeyframeData data];
k.buffer = cpBuffer; k.buffer = cpBuffer;
k.offset = step_offsets[step] * sizeof(float3); k.offset = step_offsets[step] * sizeof(packed_float3);
cp_ptrs.push_back(k); cp_ptrs.push_back(k);
k = [MTLMotionKeyframeData data]; k = [MTLMotionKeyframeData data];
@ -501,7 +496,7 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
/* controlPointCount should specify the *per-step* control point count. */ /* controlPointCount should specify the *per-step* control point count. */
geomDescCrv.controlPointCount = cpData.size() / num_motion_steps; geomDescCrv.controlPointCount = cpData.size() / num_motion_steps;
geomDescCrv.controlPointStride = sizeof(float3); geomDescCrv.controlPointStride = sizeof(packed_float3);
geomDescCrv.controlPointFormat = MTLAttributeFormatFloat3; geomDescCrv.controlPointFormat = MTLAttributeFormatFloat3;
geomDescCrv.radiusStride = sizeof(float); geomDescCrv.radiusStride = sizeof(float);
geomDescCrv.radiusFormat = MTLAttributeFormatFloat; geomDescCrv.radiusFormat = MTLAttributeFormatFloat;
@ -749,14 +744,13 @@ bool BVHMetal::build_BLAS_pointcloud(Progress &progress,
} }
const size_t num_points = pointcloud->num_points(); const size_t num_points = pointcloud->num_points();
const packed_float3 *points = pointcloud->get_position();
const float *radius = pointcloud->get_radius();
const bool use_fast_trace_bvh = (params.bvh_type == BVH_TYPE_STATIC) || !support_refit_blas(); const bool use_fast_trace_bvh = (params.bvh_type == BVH_TYPE_STATIC) || !support_refit_blas();
size_t num_motion_steps = 1; size_t num_motion_steps = 1;
Attribute *motion_keys = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); Attribute *attr_P = pointcloud->attributes.find(ATTR_STD_POSITION);
if (motion_blur && pointcloud->get_use_motion_blur() && motion_keys) { Attribute *attr_R = pointcloud->attributes.find(ATTR_STD_RADIUS);
if (motion_blur && pointcloud->get_use_motion_blur() && attr_P->has_motion()) {
num_motion_steps = pointcloud->get_motion_steps(); num_motion_steps = pointcloud->get_motion_steps();
} }
@ -769,33 +763,19 @@ bool BVHMetal::build_BLAS_pointcloud(Progress &progress,
MTLAxisAlignedBoundingBox *aabb_data = (MTLAxisAlignedBoundingBox *)[aabbBuf contents]; MTLAxisAlignedBoundingBox *aabb_data = (MTLAxisAlignedBoundingBox *)[aabbBuf contents];
/* Get AABBs for each motion step */ /* Get AABBs for each motion step */
size_t center_step = (num_motion_steps - 1) / 2;
for (size_t step = 0; step < num_motion_steps; ++step) { for (size_t step = 0; step < num_motion_steps; ++step) {
if (step == center_step) { const packed_float3 *step_points = attr_P->data_at_time_step<packed_float3>(
/* The center step for motion vertices is not stored in the attribute */ step, num_motion_steps);
for (size_t j = 0; j < num_points; ++j) { const float *step_radius = attr_R->data_at_time_step<float>(step, num_motion_steps);
const PointCloud::Point point = pointcloud->get_point(j);
BoundBox bounds = BoundBox::empty;
point.bounds_grow(points, radius, bounds);
const size_t index = step * num_points + j; for (size_t j = 0; j < num_points; ++j) {
aabb_data[index].min = (MTLPackedFloat3 &)bounds.min; const PointCloud::Point point = pointcloud->get_point(j);
aabb_data[index].max = (MTLPackedFloat3 &)bounds.max; BoundBox bounds = BoundBox::empty;
} point.bounds_grow(step_points, step_radius, bounds);
}
else {
size_t attr_offset = (step > center_step) ? step - 1 : step;
const float4 *motion_points = motion_keys->data<float4>() + attr_offset * num_points;
for (size_t j = 0; j < num_points; ++j) { const size_t index = step * num_points + j;
const PointCloud::Point point = pointcloud->get_point(j); aabb_data[index].min = (MTLPackedFloat3 &)bounds.min;
BoundBox bounds = BoundBox::empty; aabb_data[index].max = (MTLPackedFloat3 &)bounds.max;
point.bounds_grow(motion_points[j], bounds);
const size_t index = step * num_points + j;
aabb_data[index].min = (MTLPackedFloat3 &)bounds.min;
aabb_data[index].max = (MTLPackedFloat3 &)bounds.max;
}
} }
} }

110
intern/cycles/device/optix/device_impl.cpp
View file

@ -1334,8 +1334,9 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
const size_t num_segments = hair->num_segments(); const size_t num_segments = hair->num_segments();
size_t num_motion_steps = 1; size_t num_motion_steps = 1;
Attribute *motion_keys = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); const Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
if (pipeline_options.usesMotionBlur && hair->get_use_motion_blur() && motion_keys) { const Attribute *attr_R = hair->attributes.find(ATTR_STD_RADIUS);
if (pipeline_options.usesMotionBlur && hair->get_use_motion_blur() && attr_P->has_motion()) {
num_motion_steps = hair->get_motion_steps(); num_motion_steps = hair->get_motion_steps();
} }
@ -1360,19 +1361,9 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
/* Get AABBs for each motion step. */ /* Get AABBs for each motion step. */
for (size_t step = 0; step < num_motion_steps; ++step) { for (size_t step = 0; step < num_motion_steps; ++step) {
/* The center step for motion vertices is not stored in the attribute. */ const packed_float3 *keys = attr_P->data_at_time_step<packed_float3>(step,
const size_t center_step = (num_motion_steps - 1) / 2; num_motion_steps);
const packed_float3 *keys_center = nullptr; const float *curve_radius_step = attr_R->data_at_time_step<float>(step, num_motion_steps);
const packed_float3 *keys_motion = nullptr;
if (step == center_step) {
keys_center = hair->get_position();
}
else {
const size_t attr_offset = (step > center_step) ? step - 1 : step;
keys_motion = motion_keys->data<packed_float3>() + attr_offset * hair->num_keys();
}
const packed_float3 *keys = keys_center ? keys_center : keys_motion;
const auto get_key = [keys](size_t idx) -> float3 { return float3(keys[idx]); };
if (hair->curve_shape == CURVE_THICK || hair->curve_shape == CURVE_THICK_LINEAR) { if (hair->curve_shape == CURVE_THICK || hair->curve_shape == CURVE_THICK_LINEAR) {
for (size_t curve_index = 0, segment_index = 0, vertex_index = step * num_vertices; for (size_t curve_index = 0, segment_index = 0, vertex_index = step * num_vertices;
@ -1380,7 +1371,6 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
++curve_index) ++curve_index)
{ {
const Hair::Curve curve = hair->get_curve(curve_index); const Hair::Curve curve = hair->get_curve(curve_index);
const float *curve_radius = hair->get_radius();
if (hair->curve_shape == CURVE_THICK_LINEAR) { if (hair->curve_shape == CURVE_THICK_LINEAR) {
const int first_key_index = curve.first_key; const int first_key_index = curve.first_key;
@ -1389,37 +1379,32 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
if (step == 0) { if (step == 0) {
index_data[segment_index++] = vertex_index; index_data[segment_index++] = vertex_index;
} }
const float3 key = get_key(first_key_index + k); vertex_data[vertex_index++] = make_float4(float3(keys[first_key_index + k]),
vertex_data[vertex_index++] = make_float4(key, curve_radius[first_key_index + k]); curve_radius_step[first_key_index + k]);
} }
const int last_key_index = first_key_index + curve.num_keys - 1; const int last_key_index = first_key_index + curve.num_keys - 1;
{ vertex_data[vertex_index++] = make_float4(float3(keys[last_key_index]),
const float3 key = get_key(last_key_index); curve_radius_step[last_key_index]);
vertex_data[vertex_index++] = make_float4(key, curve_radius[last_key_index]);
}
} }
else { else {
const int first_key_index = curve.first_key; const int first_key_index = curve.first_key;
{ vertex_data[vertex_index++] = make_float4(float3(keys[first_key_index]),
const float3 key = get_key(first_key_index); curve_radius_step[first_key_index]);
vertex_data[vertex_index++] = make_float4(key, curve_radius[first_key_index]);
}
for (int k = 0; k < curve.num_segments(); ++k) { for (int k = 0; k < curve.num_segments(); ++k) {
if (step == 0) { if (step == 0) {
index_data[segment_index++] = vertex_index - 1; index_data[segment_index++] = vertex_index - 1;
} }
const float3 key = get_key(first_key_index + k); vertex_data[vertex_index++] = make_float4(float3(keys[first_key_index + k]),
vertex_data[vertex_index++] = make_float4(key, curve_radius[first_key_index + k]); curve_radius_step[first_key_index + k]);
} }
const int last_key_index = first_key_index + curve.num_keys - 1; const int last_key_index = first_key_index + curve.num_keys - 1;
{ vertex_data[vertex_index++] = make_float4(float3(keys[last_key_index]),
const float3 key = get_key(last_key_index); curve_radius_step[last_key_index]);
vertex_data[vertex_index++] = make_float4(key, curve_radius[last_key_index]); vertex_data[vertex_index++] = make_float4(float3(keys[last_key_index]),
vertex_data[vertex_index++] = make_float4(key, curve_radius[last_key_index]); curve_radius_step[last_key_index]);
}
} }
} }
} }
@ -1429,7 +1414,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
for (int segment = 0; segment < curve.num_segments(); ++segment, ++i) { for (int segment = 0; segment < curve.num_segments(); ++segment, ++i) {
BoundBox bounds = BoundBox::empty; BoundBox bounds = BoundBox::empty;
curve.bounds_grow(segment, keys, hair->get_radius(), bounds); curve.bounds_grow(segment, keys, curve_radius_step, bounds);
const size_t index = step * num_segments + i; const size_t index = step * num_segments + i;
aabb_data[index].minX = bounds.min.x; aabb_data[index].minX = bounds.min.x;
@ -1574,8 +1559,11 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
} }
size_t num_motion_steps = 1; size_t num_motion_steps = 1;
Attribute *motion_points = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); Attribute *attr_P = pointcloud->attributes.find(ATTR_STD_POSITION);
if (pipeline_options.usesMotionBlur && pointcloud->get_use_motion_blur() && motion_points) { Attribute *attr_R = pointcloud->attributes.find(ATTR_STD_RADIUS);
if (pipeline_options.usesMotionBlur && pointcloud->get_use_motion_blur() &&
attr_P->has_motion())
{
num_motion_steps = pointcloud->get_motion_steps(); num_motion_steps = pointcloud->get_motion_steps();
} }
@ -1584,44 +1572,22 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
/* Get AABBs for each motion step. */ /* Get AABBs for each motion step. */
for (size_t step = 0; step < num_motion_steps; ++step) { for (size_t step = 0; step < num_motion_steps; ++step) {
/* The center step for motion vertices is not stored in the attribute. */ const packed_float3 *points = attr_P->data_at_time_step<packed_float3>(step,
size_t center_step = (num_motion_steps - 1) / 2; num_motion_steps);
const float *radius = attr_R->data_at_time_step<float>(step, num_motion_steps);
if (step == center_step) { for (size_t i = 0; i < num_points; ++i) {
const packed_float3 *points = pointcloud->get_position(); const PointCloud::Point point = pointcloud->get_point(i);
const float *radius = pointcloud->get_radius(); BoundBox bounds = BoundBox::empty;
point.bounds_grow(points, radius, bounds);
for (size_t i = 0; i < num_points; ++i) { const size_t index = step * num_points + i;
const PointCloud::Point point = pointcloud->get_point(i); aabb_data[index].minX = bounds.min.x;
BoundBox bounds = BoundBox::empty; aabb_data[index].minY = bounds.min.y;
point.bounds_grow(points, radius, bounds); aabb_data[index].minZ = bounds.min.z;
aabb_data[index].maxX = bounds.max.x;
const size_t index = step * num_points + i; aabb_data[index].maxY = bounds.max.y;
aabb_data[index].minX = bounds.min.x; aabb_data[index].maxZ = bounds.max.z;
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;
}
}
else {
size_t attr_offset = (step > center_step) ? step - 1 : step;
const float4 *points = motion_points->data<float4>() + 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[i], 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;
}
} }
} }

2
intern/cycles/kernel/data_arrays.h
View file

@ -37,11 +37,9 @@ KERNEL_DATA_ARRAY(packed_uint3, tri_vindex)
/* curves */ /* curves */
KERNEL_DATA_ARRAY(KernelCurve, curves) KERNEL_DATA_ARRAY(KernelCurve, curves)
KERNEL_DATA_ARRAY(float4, curve_keys)
KERNEL_DATA_ARRAY(KernelCurveSegment, curve_segments) KERNEL_DATA_ARRAY(KernelCurveSegment, curve_segments)
/* pointclouds */ /* pointclouds */
KERNEL_DATA_ARRAY(float4, points)
KERNEL_DATA_ARRAY(uint, points_shader) KERNEL_DATA_ARRAY(uint, points_shader)
/* attributes */ /* attributes */

18
intern/cycles/kernel/device/metal/bvh.h
View file

@ -75,11 +75,12 @@ ccl_device_forceinline bool curve_ribbon_accept(KernelGlobals kg,
/* We can ignore motion blur here because we don't need the positions, and it doesn't affect the /* We can ignore motion blur here because we don't need the positions, and it doesn't affect the
* radius. */ * radius. */
const int position_offset = kernel_data_fetch(objects, object).position_offset;
float radius[4]; float radius[4];
radius[0] = kernel_data_fetch(curve_keys, ka).w; radius[0] = kernel_data_fetch(attributes_float4, position_offset + ka).w;
radius[1] = kernel_data_fetch(curve_keys, k0).w; radius[1] = kernel_data_fetch(attributes_float4, position_offset + k0).w;
radius[2] = kernel_data_fetch(curve_keys, k1).w; radius[2] = kernel_data_fetch(attributes_float4, position_offset + k1).w;
radius[3] = kernel_data_fetch(curve_keys, kb).w; radius[3] = kernel_data_fetch(attributes_float4, position_offset + kb).w;
const float r = metal::catmull_rom(u, radius[0], radius[1], radius[2], radius[3]); const float r = metal::catmull_rom(u, radius[0], radius[1], radius[2], radius[3]);
/* MPJ TODO: Can we ignore motion and/or object transforms here? Depends on scaling? */ /* MPJ TODO: Can we ignore motion and/or object transforms here? Depends on scaling? */
@ -124,10 +125,11 @@ ccl_device_forceinline float curve_ribbon_v(KernelGlobals kg,
float4 curve[4]; float4 curve[4];
if (!is_motion) { if (!is_motion) {
curve[0] = kernel_data_fetch(curve_keys, ka); const int position_offset = kernel_data_fetch(objects, object).position_offset;
curve[1] = kernel_data_fetch(curve_keys, k0); curve[0] = kernel_data_fetch(attributes_float4, position_offset + ka);
curve[2] = kernel_data_fetch(curve_keys, k1); curve[1] = kernel_data_fetch(attributes_float4, position_offset + k0);
curve[3] = kernel_data_fetch(curve_keys, kb); curve[2] = kernel_data_fetch(attributes_float4, position_offset + k1);
curve[3] = kernel_data_fetch(attributes_float4, position_offset + kb);
} }
else { else {
motion_curve_keys(kg, object, time, ka, k0, k1, kb, curve); motion_curve_keys(kg, object, time, ka, k0, k1, kb, curve);

10
intern/cycles/kernel/geom/curve.h
View file

@ -108,8 +108,9 @@ ccl_device float curve_thickness(KernelGlobals kg, const ccl_private ShaderData
else else
# endif # endif
{ {
P_curve[0] = kernel_data_fetch(curve_keys, k0); const int position_offset = kernel_data_fetch(objects, sd->object).position_offset;
P_curve[1] = kernel_data_fetch(curve_keys, k1); P_curve[0] = kernel_data_fetch(attributes_float4, position_offset + k0);
P_curve[1] = kernel_data_fetch(attributes_float4, position_offset + k1);
} }
float r = 2.0f * ((P_curve[1].w - P_curve[0].w) * sd->u + P_curve[0].w); float r = 2.0f * ((P_curve[1].w - P_curve[0].w) * sd->u + P_curve[0].w);
@ -146,8 +147,9 @@ ccl_device float3 curve_motion_center_location(KernelGlobals kg, const ccl_priva
float4 P_curve[2]; float4 P_curve[2];
P_curve[0] = kernel_data_fetch(curve_keys, k0); const int position_offset = kernel_data_fetch(objects, sd->object).position_offset;
P_curve[1] = kernel_data_fetch(curve_keys, k1); P_curve[0] = kernel_data_fetch(attributes_float4, position_offset + k0);
P_curve[1] = kernel_data_fetch(attributes_float4, position_offset + k1);
return make_float3(P_curve[1]) * sd->u + make_float3(P_curve[0]) * (1.0f - sd->u); return make_float3(P_curve[1]) * sd->u + make_float3(P_curve[0]) * (1.0f - sd->u);
} }

18
intern/cycles/kernel/geom/curve_intersect.h
View file

@ -817,10 +817,11 @@ ccl_device_forceinline bool curve_intersect(KernelGlobals kg,
float4 curve[4]; float4 curve[4];
if (!is_motion) { if (!is_motion) {
curve[0] = kernel_data_fetch(curve_keys, ka); const int position_offset = kernel_data_fetch(objects, object).position_offset;
curve[1] = kernel_data_fetch(curve_keys, k0); curve[0] = kernel_data_fetch(attributes_float4, position_offset + ka);
curve[2] = kernel_data_fetch(curve_keys, k1); curve[1] = kernel_data_fetch(attributes_float4, position_offset + k0);
curve[3] = kernel_data_fetch(curve_keys, kb); curve[2] = kernel_data_fetch(attributes_float4, position_offset + k1);
curve[3] = kernel_data_fetch(attributes_float4, position_offset + kb);
} }
else { else {
motion_curve_keys(kg, object, time, ka, k0, k1, kb, curve); motion_curve_keys(kg, object, time, ka, k0, k1, kb, curve);
@ -886,10 +887,11 @@ ccl_device_inline void curve_shader_setup(KernelGlobals kg,
float4 P_curve[4]; float4 P_curve[4];
if (!(sd->type & PRIMITIVE_MOTION)) { if (!(sd->type & PRIMITIVE_MOTION)) {
P_curve[0] = kernel_data_fetch(curve_keys, ka); const int position_offset = kernel_data_fetch(objects, sd->object).position_offset;
P_curve[1] = kernel_data_fetch(curve_keys, k0); P_curve[0] = kernel_data_fetch(attributes_float4, position_offset + ka);
P_curve[2] = kernel_data_fetch(curve_keys, k1); P_curve[1] = kernel_data_fetch(attributes_float4, position_offset + k0);
P_curve[3] = kernel_data_fetch(curve_keys, kb); P_curve[2] = kernel_data_fetch(attributes_float4, position_offset + k1);
P_curve[3] = kernel_data_fetch(attributes_float4, position_offset + kb);
} }
else { else {
motion_curve_keys(kg, sd->object, sd->time, ka, k0, k1, kb, P_curve); motion_curve_keys(kg, sd->object, sd->time, ka, k0, k1, kb, P_curve);

52
intern/cycles/kernel/geom/motion_curve.h
View file

@ -16,7 +16,7 @@ CCL_NAMESPACE_BEGIN
* other than the frame center. Computing the curve keys at a given ray time is * other than the frame center. Computing the curve keys at a given ray time is
* a matter of interpolation of the two steps between which the ray time lies. * a matter of interpolation of the two steps between which the ray time lies.
* *
* The extra curve keys are stored as ATTR_STD_MOTION_VERTEX_POSITION. * The extra curve keys are stored as additional motion steps in ATTR_STD_POSITION.
*/ */
#ifdef __HAIR__ #ifdef __HAIR__
@ -32,21 +32,18 @@ ccl_device_inline void motion_curve_keys_for_step_linear(KernelGlobals kg,
{ {
const int center_step = (numsteps - 1) / 2; const int center_step = (numsteps - 1) / 2;
if (step == center_step) { if (step == center_step) {
/* center step: regular key location */ /* Center step: first in the array. */
keys[0] = kernel_data_fetch(curve_keys, k0);
keys[1] = kernel_data_fetch(curve_keys, k1);
} }
else { else {
/* center step is not stored in this array */ /* Non-center step, stored after center with center index skipped. */
if (step > center_step) { if (step < center_step) {
step--; step++;
} }
offset += step * numverts; offset += step * numverts;
keys[0] = kernel_data_fetch(attributes_float4, offset + k0);
keys[1] = kernel_data_fetch(attributes_float4, offset + k1);
} }
keys[0] = kernel_data_fetch(attributes_float4, offset + k0);
keys[1] = kernel_data_fetch(attributes_float4, offset + k1);
} }
/* return 2 curve key locations */ /* return 2 curve key locations */
@ -66,11 +63,8 @@ ccl_device_inline void motion_curve_keys_linear(KernelGlobals kg,
const int step = min((int)(time * maxstep), maxstep - 1); const int step = min((int)(time * maxstep), maxstep - 1);
const float t = time * maxstep - step; const float t = time * maxstep - step;
/* find attribute */
const int offset = intersection_find_attribute(kg, object, ATTR_STD_MOTION_VERTEX_POSITION);
kernel_assert(offset != ATTR_STD_NOT_FOUND);
/* fetch key coordinates */ /* fetch key coordinates */
const int offset = kernel_data_fetch(objects, object).position_offset;
float4 next_keys[2]; float4 next_keys[2];
motion_curve_keys_for_step_linear(kg, offset, numverts, numsteps, step, k0, k1, keys); motion_curve_keys_for_step_linear(kg, offset, numverts, numsteps, step, k0, k1, keys);
@ -94,25 +88,20 @@ ccl_device_inline void motion_curve_keys_for_step(KernelGlobals kg,
{ {
const int center_step = (numsteps - 1) / 2; const int center_step = (numsteps - 1) / 2;
if (step == center_step) { if (step == center_step) {
/* center step: regular key location */ /* Center step: first in the array. */
keys[0] = kernel_data_fetch(curve_keys, k0);
keys[1] = kernel_data_fetch(curve_keys, k1);
keys[2] = kernel_data_fetch(curve_keys, k2);
keys[3] = kernel_data_fetch(curve_keys, k3);
} }
else { else {
/* center step is not stored in this array */ /* Non-center step, stored after center with center index skipped. */
if (step > center_step) { if (step < center_step) {
step--; step++;
} }
offset += step * numverts; offset += step * numverts;
keys[0] = kernel_data_fetch(attributes_float4, offset + k0);
keys[1] = kernel_data_fetch(attributes_float4, offset + k1);
keys[2] = kernel_data_fetch(attributes_float4, offset + k2);
keys[3] = kernel_data_fetch(attributes_float4, offset + k3);
} }
keys[0] = kernel_data_fetch(attributes_float4, offset + k0);
keys[1] = kernel_data_fetch(attributes_float4, offset + k1);
keys[2] = kernel_data_fetch(attributes_float4, offset + k2);
keys[3] = kernel_data_fetch(attributes_float4, offset + k3);
} }
/* return 2 curve key locations */ /* return 2 curve key locations */
@ -134,11 +123,8 @@ ccl_device_inline void motion_curve_keys(KernelGlobals kg,
const int step = min((int)(time * maxstep), maxstep - 1); const int step = min((int)(time * maxstep), maxstep - 1);
const float t = time * maxstep - step; const float t = time * maxstep - step;
/* find attribute */
const int offset = intersection_find_attribute(kg, object, ATTR_STD_MOTION_VERTEX_POSITION);
kernel_assert(offset != ATTR_STD_NOT_FOUND);
/* fetch key coordinates */ /* fetch key coordinates */
const int offset = kernel_data_fetch(objects, object).position_offset;
float4 next_keys[4]; float4 next_keys[4];
motion_curve_keys_for_step(kg, offset, numverts, numsteps, step, k0, k1, k2, k3, keys); motion_curve_keys_for_step(kg, offset, numverts, numsteps, step, k0, k1, k2, k3, keys);

21
intern/cycles/kernel/geom/motion_point.h
View file

@ -16,7 +16,7 @@ CCL_NAMESPACE_BEGIN
* other than the frame center. Computing the point at a given ray time is * 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. * 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. * The extra points are stored as additional motion steps in ATTR_STD_POSITION.
*/ */
#ifdef __POINTCLOUD__ #ifdef __POINTCLOUD__
@ -26,16 +26,16 @@ ccl_device_inline float4 motion_point_for_step(
{ {
const int center_step = (numsteps - 1) / 2; const int center_step = (numsteps - 1) / 2;
if (step == center_step) { if (step == center_step) {
/* center step: regular key location */ /* Center step: first in the array. */
return kernel_data_fetch(points, prim);
} }
/* center step is not stored in this array */ else {
if (step > center_step) { /* Non-center step, stored after center with center index skipped. */
step--; if (step < center_step) {
step++;
}
offset += step * numverts;
} }
offset += step * numverts;
return kernel_data_fetch(attributes_float4, offset + prim); return kernel_data_fetch(attributes_float4, offset + prim);
} }
@ -54,11 +54,8 @@ ccl_device_inline float4 motion_point(KernelGlobals kg,
const int step = min((int)(time * maxstep), maxstep - 1); const int step = min((int)(time * maxstep), maxstep - 1);
const float t = time * maxstep - step; const float t = time * maxstep - step;
/* find attribute */
const int offset = intersection_find_attribute(kg, object, ATTR_STD_MOTION_VERTEX_POSITION);
kernel_assert(offset != ATTR_STD_NOT_FOUND);
/* fetch key coordinates */ /* fetch key coordinates */
const int offset = kernel_data_fetch(objects, object).position_offset;
const float4 point = motion_point_for_step(kg, offset, numverts, numsteps, step, prim); const float4 point = motion_point_for_step(kg, offset, numverts, numsteps, step, prim);
const float4 next_point = motion_point_for_step(kg, offset, numverts, numsteps, step + 1, prim); const float4 next_point = motion_point_for_step(kg, offset, numverts, numsteps, step + 1, prim);

9
intern/cycles/kernel/geom/point.h
View file

@ -38,9 +38,10 @@ ccl_device float3 point_position(KernelGlobals kg, const ccl_private ShaderData
{ {
if (sd->type & PRIMITIVE_POINT) { if (sd->type & PRIMITIVE_POINT) {
/* World space center. */ /* World space center. */
const int position_offset = kernel_data_fetch(objects, sd->object).position_offset;
float3 P = (sd->type & PRIMITIVE_MOTION) ? float3 P = (sd->type & PRIMITIVE_MOTION) ?
make_float3(motion_point(kg, sd->object, sd->prim, sd->time)) : make_float3(motion_point(kg, sd->object, sd->prim, sd->time)) :
make_float3(kernel_data_fetch(points, sd->prim)); make_float3(kernel_data_fetch(attributes_float4, position_offset + sd->prim));
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) { if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
object_position_transform(kg, sd, &P); object_position_transform(kg, sd, &P);
@ -58,7 +59,8 @@ ccl_device float point_radius(KernelGlobals kg, const ccl_private ShaderData *sd
{ {
if (sd->type & PRIMITIVE_POINT) { if (sd->type & PRIMITIVE_POINT) {
/* World space radius. */ /* World space radius. */
const float r = kernel_data_fetch(points, sd->prim).w; const int position_offset = kernel_data_fetch(objects, sd->object).position_offset;
const float r = kernel_data_fetch(attributes_float4, position_offset + sd->prim).w;
if (sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED) { if (sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED) {
return r; return r;
@ -89,7 +91,8 @@ ccl_device float point_random(KernelGlobals kg, const ccl_private ShaderData *sd
ccl_device float3 point_motion_center_location(KernelGlobals kg, const ccl_private ShaderData *sd) ccl_device float3 point_motion_center_location(KernelGlobals kg, const ccl_private ShaderData *sd)
{ {
return make_float3(kernel_data_fetch(points, sd->prim)); const int position_offset = kernel_data_fetch(objects, sd->object).position_offset;
return make_float3(kernel_data_fetch(attributes_float4, position_offset + sd->prim));
} }
#endif /* __POINTCLOUD__ */ #endif /* __POINTCLOUD__ */

14
intern/cycles/kernel/geom/point_intersect.h
View file

@ -74,8 +74,10 @@ ccl_device_forceinline bool point_intersect(KernelGlobals kg,
const float time, const float time,
const int type) const int type)
{ {
const float4 point = (type & PRIMITIVE_MOTION) ? motion_point(kg, object, prim, time) : const int position_offset = kernel_data_fetch(objects, object).position_offset;
kernel_data_fetch(points, prim); const float4 point = (type & PRIMITIVE_MOTION) ?
motion_point(kg, object, prim, time) :
kernel_data_fetch(attributes_float4, position_offset + prim);
if (!point_intersect_test(point, ray_P, ray_D, ray_tmin, ray_tmax, &isect->t)) { if (!point_intersect_test(point, ray_P, ray_D, ray_tmin, ray_tmax, &isect->t)) {
return false; return false;
@ -104,9 +106,11 @@ ccl_device_inline void point_shader_setup(KernelGlobals kg,
# endif # endif
/* Compute point center for normal. */ /* Compute point center for normal. */
float3 center = make_float3((isect->type & PRIMITIVE_MOTION) ? const int position_offset = kernel_data_fetch(objects, sd->object).position_offset;
motion_point(kg, sd->object, sd->prim, sd->time) : float3 center = make_float3(
kernel_data_fetch(points, sd->prim)); (isect->type & PRIMITIVE_MOTION) ?
motion_point(kg, sd->object, sd->prim, sd->time) :
kernel_data_fetch(attributes_float4, position_offset + sd->prim));
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) { if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
object_position_transform(kg, sd, &center); object_position_transform(kg, sd, &center);
} }

29
intern/cycles/kernel/geom/primitive.h
View file

@ -222,13 +222,19 @@ ccl_device_forceinline void primitive_motion_data_without_camera(KernelGlobals k
*motion_post = *motion_center; *motion_post = *motion_center;
/* deformation motion */ /* deformation motion */
const ccl_global KernelObject &kobject = kernel_data_fetch(objects, sd->object);
const int pos_offset = kobject.position_offset;
const int numverts = kobject.numverts;
const int num_motion_steps = kobject.num_geom_steps;
#if defined(__HAIR__) || defined(__POINTCLOUD__) #if defined(__HAIR__) || defined(__POINTCLOUD__)
if (is_curve_or_point) { if (is_curve_or_point) {
AttributeDescriptor desc = find_attribute(kg, sd, ATTR_STD_MOTION_VERTEX_POSITION); if (sd->object_flag & SD_OBJECT_HAS_VERTEX_MOTION) {
if (is_attribute_found(desc)) { AttributeDescriptor desc;
const ccl_global KernelObject *kobject = &kernel_data_fetch(objects, sd->object); desc.type = NODE_ATTR_FLOAT4;
const int numverts = kobject->numverts; desc.element = (sd->type & PRIMITIVE_CURVE) ? ATTR_ELEMENT_CURVE_KEY : ATTR_ELEMENT_VERTEX;
const int num_motion_steps = kobject->num_geom_steps;
desc.offset = pos_offset + numverts;
*motion_pre = make_float3(primitive_surface_attribute<float4>(kg, sd, desc)); *motion_pre = make_float3(primitive_surface_attribute<float4>(kg, sd, desc));
if (num_motion_steps > 2) { if (num_motion_steps > 2) {
desc.offset += numverts; desc.offset += numverts;
@ -237,27 +243,18 @@ ccl_device_forceinline void primitive_motion_data_without_camera(KernelGlobals k
else { else {
object_inverse_position_transform(kg, sd, motion_post); object_inverse_position_transform(kg, sd, motion_post);
} }
/* Curve */
if ((sd->object_flag & SD_OBJECT_HAS_VERTEX_MOTION) == 0) {
object_position_transform(kg, sd, motion_pre);
object_position_transform(kg, sd, motion_post);
}
} }
} }
else else
#endif #endif
if (sd->type & PRIMITIVE_TRIANGLE) if (sd->type & PRIMITIVE_TRIANGLE)
{ {
/* Mesh: motion stored inline in ATTR_STD_POSITION, after the center step. */
if (sd->object_flag & SD_OBJECT_HAS_VERTEX_MOTION) { if (sd->object_flag & SD_OBJECT_HAS_VERTEX_MOTION) {
const KernelObject &kobject = kernel_data_fetch(objects, sd->object);
const int numverts = kobject.numverts;
const int num_motion_steps = kobject.num_geom_steps;
AttributeDescriptor desc; AttributeDescriptor desc;
desc.element = ATTR_ELEMENT_VERTEX; desc.element = ATTR_ELEMENT_VERTEX;
desc.type = NODE_ATTR_FLOAT3; desc.type = NODE_ATTR_FLOAT3;
desc.offset = kobject.position_offset + numverts;
desc.offset = pos_offset + numverts;
*motion_pre = triangle_attribute<float3>(kg, sd, desc); *motion_pre = triangle_attribute<float3>(kg, sd, desc);
if (num_motion_steps > 2) { if (num_motion_steps > 2) {
desc.offset += numverts; desc.offset += numverts;

9
intern/cycles/kernel/geom/shader_data.h
View file

@ -340,10 +340,11 @@ ccl_device void shader_setup_from_curve(KernelGlobals kg,
float4 P_curve[4]; float4 P_curve[4];
P_curve[0] = kernel_data_fetch(curve_keys, ka); const int position_offset = kernel_data_fetch(objects, object).position_offset;
P_curve[1] = kernel_data_fetch(curve_keys, k0); P_curve[0] = kernel_data_fetch(attributes_float4, position_offset + ka);
P_curve[2] = kernel_data_fetch(curve_keys, k1); P_curve[1] = kernel_data_fetch(attributes_float4, position_offset + k0);
P_curve[3] = kernel_data_fetch(curve_keys, kb); P_curve[2] = kernel_data_fetch(attributes_float4, position_offset + k1);
P_curve[3] = kernel_data_fetch(attributes_float4, position_offset + kb);
/* Interpolate position and tangent. */ /* Interpolate position and tangent. */
sd->P = (sd->type & PRIMITIVE_CURVE) == PRIMITIVE_CURVE_THICK_LINEAR ? sd->P = (sd->type & PRIMITIVE_CURVE) == PRIMITIVE_CURVE_THICK_LINEAR ?

6
intern/cycles/kernel/osl/closures_setup.h
View file

@ -1266,8 +1266,10 @@ ccl_device void osl_closure_hair_huang_setup(KernelGlobals kg,
const KernelCurve kcurve = kernel_data_fetch(curves, sd->prim); const KernelCurve kcurve = kernel_data_fetch(curves, sd->prim);
const int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type); const int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
const int k1 = k0 + 1; const int k1 = k0 + 1;
const float radius = mix( const int position_offset = kernel_data_fetch(objects, sd->object).position_offset;
kernel_data_fetch(curve_keys, k0).w, kernel_data_fetch(curve_keys, k1).w, sd->u); const float radius = mix(kernel_data_fetch(attributes_float4, position_offset + k0).w,
kernel_data_fetch(attributes_float4, position_offset + k1).w,
sd->u);
bsdf->extra->pixel_coverage = 0.5f * sd->dP / radius; bsdf->extra->pixel_coverage = 0.5f * sd->dP / radius;
} }

6
intern/cycles/kernel/svm/closure.h
View file

@ -1009,8 +1009,10 @@ ccl_device
const KernelCurve kcurve = kernel_data_fetch(curves, sd->prim); const KernelCurve kcurve = kernel_data_fetch(curves, sd->prim);
const int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type); const int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
const int k1 = k0 + 1; const int k1 = k0 + 1;
const float radius = mix( const int position_offset = kernel_data_fetch(objects, sd->object).position_offset;
kernel_data_fetch(curve_keys, k0).w, kernel_data_fetch(curve_keys, k1).w, sd->u); const float radius = mix(kernel_data_fetch(attributes_float4, position_offset + k0).w,
kernel_data_fetch(attributes_float4, position_offset + k1).w,
sd->u);
bsdf->extra->pixel_coverage = 0.5f * sd->dP / radius; bsdf->extra->pixel_coverage = 0.5f * sd->dP / radius;
} }

11
intern/cycles/kernel/types.h
View file

@ -749,20 +749,12 @@ enum AttributeElement {
/* Only these combinations are supported by the kernel and can be /* Only these combinations are supported by the kernel and can be
* created on geometry. */ * created on geometry. */
ATTR_ELEMENT_VERTEX_MOTION = ATTR_ELEMENT_VERTEX | ATTR_ELEMENT_IS_MOTION,
ATTR_ELEMENT_VERTEX_NORMAL = ATTR_ELEMENT_VERTEX | ATTR_ELEMENT_IS_NORMAL, ATTR_ELEMENT_VERTEX_NORMAL = ATTR_ELEMENT_VERTEX | ATTR_ELEMENT_IS_NORMAL,
ATTR_ELEMENT_VERTEX_NORMAL_MOTION = ATTR_ELEMENT_VERTEX | ATTR_ELEMENT_IS_NORMAL |
ATTR_ELEMENT_IS_MOTION,
ATTR_ELEMENT_CORNER_BYTE = ATTR_ELEMENT_CORNER | ATTR_ELEMENT_IS_BYTE, ATTR_ELEMENT_CORNER_BYTE = ATTR_ELEMENT_CORNER | ATTR_ELEMENT_IS_BYTE,
ATTR_ELEMENT_CORNER_NORMAL = ATTR_ELEMENT_CORNER | ATTR_ELEMENT_IS_NORMAL, ATTR_ELEMENT_CORNER_NORMAL = ATTR_ELEMENT_CORNER | ATTR_ELEMENT_IS_NORMAL,
ATTR_ELEMENT_CORNER_NORMAL_MOTION = ATTR_ELEMENT_CORNER | ATTR_ELEMENT_IS_NORMAL |
ATTR_ELEMENT_IS_MOTION,
ATTR_ELEMENT_CURVE_KEY_MOTION = ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_IS_MOTION,
ATTR_ELEMENT_CURVE_KEY_NORMAL = ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_IS_NORMAL, ATTR_ELEMENT_CURVE_KEY_NORMAL = ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_IS_NORMAL,
ATTR_ELEMENT_CURVE_KEY_NORMAL_MOTION = ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_IS_NORMAL |
ATTR_ELEMENT_IS_MOTION,
}; };
enum AttributeStandard : int { enum AttributeStandard : int {
@ -782,9 +774,6 @@ enum AttributeStandard : int {
ATTR_STD_POSITION_UNDEFORMED, ATTR_STD_POSITION_UNDEFORMED,
ATTR_STD_POSITION_UNDISPLACED, ATTR_STD_POSITION_UNDISPLACED,
ATTR_STD_NORMAL_UNDISPLACED, ATTR_STD_NORMAL_UNDISPLACED,
ATTR_STD_MOTION_VERTEX_POSITION,
ATTR_STD_MOTION_VERTEX_NORMAL,
ATTR_STD_MOTION_CORNER_NORMAL,
ATTR_STD_PARTICLE, ATTR_STD_PARTICLE,
ATTR_STD_CURVE_INTERCEPT, ATTR_STD_CURVE_INTERCEPT,
ATTR_STD_CURVE_LENGTH, ATTR_STD_CURVE_LENGTH,

155
intern/cycles/scene/attribute.cpp
View file

@ -62,6 +62,14 @@ Attribute::Attribute(ustring name,
center.sharing_info = sharing_info; center.sharing_info = sharing_info;
} }
static size_t attribute_alloc_bytes(const size_t element_size, const size_t size)
{
/* rtcSetSharedGeometryBuffer is documented as requiring 4 bytes past the
* end of a float3 for 16-byte SSE loads, so we add that for all attributes. */
static constexpr size_t ATTRIBUTE_BUFFER_PADDING = 4;
return element_size * size + ATTRIBUTE_BUFFER_PADDING;
}
static void free_step_buffer(Attribute::Buffer &buf, static void free_step_buffer(Attribute::Buffer &buf,
const AttributeElement element, const AttributeElement element,
const size_t data_sizeof, const size_t data_sizeof,
@ -79,7 +87,7 @@ static void free_step_buffer(Attribute::Buffer &buf,
} }
else if (buf.data) { else if (buf.data) {
GuardedAllocator<char>().deallocate(static_cast<char *>(const_cast<void *>(buf.data)), GuardedAllocator<char>().deallocate(static_cast<char *>(const_cast<void *>(buf.data)),
data_sizeof * size); attribute_alloc_bytes(data_sizeof, size));
} }
buf.data = nullptr; buf.data = nullptr;
buf.sharing_info = nullptr; buf.sharing_info = nullptr;
@ -97,12 +105,10 @@ Attribute::Attribute(Attribute &&other)
void Attribute::free_data() void Attribute::free_data()
{ {
const size_t sz = data_sizeof(); const size_t element_size = data_sizeof();
free_step_buffer(center, element, sz, size); free_step_buffer(center, element, element_size, size);
/* Motion steps are never voxels, strip that flag for the helper. */
const AttributeElement motion_element = AttributeElement(element & ~ATTR_ELEMENT_VOXEL);
for (Buffer &buf : motion) { for (Buffer &buf : motion) {
free_step_buffer(buf, motion_element, sz, size); free_step_buffer(buf, element, element_size, size);
} }
motion.clear(); motion.clear();
} }
@ -127,22 +133,23 @@ void Attribute::resize(const size_t num_elements)
if (num_elements == size_t(size)) { if (num_elements == size_t(size)) {
return; return;
} }
const size_t sz = data_sizeof(); const size_t element_size = data_sizeof();
const size_t copy_elems = std::min(num_elements, size_t(size)); const size_t copy_elems = std::min(num_elements, size_t(size));
const size_t alloc_bytes = attribute_alloc_bytes(element_size, num_elements);
/* Allocate and copy center step. */ /* Allocate and copy center step. */
Buffer new_center; Buffer new_center;
new_center.data = GuardedAllocator<char>().allocate(num_elements * sz); new_center.data = GuardedAllocator<char>().allocate(alloc_bytes);
if (center.data) { if (center.data) {
memcpy(const_cast<void *>(new_center.data), center.data, copy_elems * sz); memcpy(const_cast<void *>(new_center.data), center.data, copy_elems * element_size);
} }
/* Allocate and copy motion steps. */ /* Allocate and copy motion steps. */
vector<Buffer> new_motion(motion.size()); vector<Buffer> new_motion(motion.size());
for (size_t i = 0; i < motion.size(); i++) { for (size_t i = 0; i < motion.size(); i++) {
new_motion[i].data = GuardedAllocator<char>().allocate(num_elements * sz); new_motion[i].data = GuardedAllocator<char>().allocate(alloc_bytes);
if (motion[i].data) { if (motion[i].data) {
memcpy(const_cast<void *>(new_motion[i].data), motion[i].data, copy_elems * sz); memcpy(const_cast<void *>(new_motion[i].data), motion[i].data, copy_elems * element_size);
} }
} }
@ -154,30 +161,30 @@ void Attribute::resize(const size_t num_elements)
void Attribute::add_motion(const Geometry *geom) void Attribute::add_motion(const Geometry *geom)
{ {
const int new_motion_count = int(geom->get_motion_steps()) - 1; const int motion_steps = geom->get_motion_steps();
assert(new_motion_count >= 0); if (motion_steps <= 0) {
if (new_motion_count == int(motion.size())) {
return; return;
} }
const size_t sz = data_sizeof();
const AttributeElement motion_element = AttributeElement(element & ~ATTR_ELEMENT_VOXEL);
if (new_motion_count < int(motion.size())) { const int motion_size = geom->get_motion_steps() - 1;
/* Shrink: free and drop extra steps. */ if (motion_size == motion.size()) {
for (size_t i = new_motion_count; i < motion.size(); i++) { return;
free_step_buffer(motion[i], motion_element, sz, size); }
const size_t element_size = data_sizeof();
if (motion_size < motion.size()) {
for (size_t i = motion_size; i < motion.size(); i++) {
free_step_buffer(motion[i], element, element_size, size);
} }
motion.resize(new_motion_count); motion.resize(motion_size);
} }
else { else {
/* Grow: allocate fresh zero-initialized arrays for the new steps. */ motion.reserve(motion_size);
motion.reserve(new_motion_count); while (motion.size() < motion_size) {
while (int(motion.size()) < new_motion_count) {
Buffer buf; Buffer buf;
if (size > 0) { if (size > 0) {
const size_t alloc_bytes = sz * size; /* Left uninitialized, callers fill in the motion data for every step. */
buf.data = GuardedAllocator<char>().allocate(alloc_bytes); buf.data = GuardedAllocator<char>().allocate(attribute_alloc_bytes(element_size, size));
memset(const_cast<void *>(buf.data), 0, alloc_bytes);
} }
motion.push_back(buf); motion.push_back(buf);
} }
@ -191,10 +198,9 @@ void Attribute::remove_motion()
if (!has_motion()) { if (!has_motion()) {
return; return;
} }
const size_t sz = data_sizeof(); const size_t element_size = data_sizeof();
const AttributeElement motion_element = AttributeElement(element & ~ATTR_ELEMENT_VOXEL);
for (Buffer &buf : motion) { for (Buffer &buf : motion) {
free_step_buffer(buf, motion_element, sz, size); free_step_buffer(buf, element, element_size, size);
} }
motion.clear(); motion.clear();
modified = true; modified = true;
@ -210,7 +216,7 @@ void Attribute::take_motion_from(Attribute &other)
modified = true; modified = true;
} }
static char *buffer_for_write(Attribute::Buffer &buf, const size_t data_sizeof, const size_t size) static char *buffer_for_write(Attribute::Buffer &buf, const size_t element_size, const size_t size)
{ {
if (!buf.data) { if (!buf.data) {
return nullptr; return nullptr;
@ -219,8 +225,8 @@ static char *buffer_for_write(Attribute::Buffer &buf, const size_t data_sizeof,
/* Here we assume that the sharing info is not mutable. With the addition of another sharing /* Here we assume that the sharing info is not mutable. With the addition of another sharing
* info callback function pointer we could check the user count to avoid unnecessary copies. * info callback function pointer we could check the user count to avoid unnecessary copies.
* For now that isn't expected to happen in practice though. */ * For now that isn't expected to happen in practice though. */
auto *new_data = GuardedAllocator<char>().allocate(data_sizeof * size); auto *new_data = GuardedAllocator<char>().allocate(attribute_alloc_bytes(element_size, size));
memcpy(new_data, buf.data, data_sizeof * size); memcpy(new_data, buf.data, element_size * size);
g_implicit_sharing_user_remove_fn(buf.sharing_info); g_implicit_sharing_user_remove_fn(buf.sharing_info);
buf.sharing_info = nullptr; buf.sharing_info = nullptr;
buf.data = new_data; buf.data = new_data;
@ -249,10 +255,10 @@ void Attribute::set_data_from(Attribute &&other)
flags = other.flags; flags = other.flags;
const size_t sz = data_sizeof(); const size_t element_size = data_sizeof();
const AttributeElement motion_element = AttributeElement(element & ~ATTR_ELEMENT_VOXEL);
const auto take_all = [&]() { /* If topology or motion steps differ, take all data. */
if (size != other.size || motion.size() != other.motion.size()) {
free_data(); free_data();
center = other.center; center = other.center;
motion = std::move(other.motion); motion = std::move(other.motion);
@ -260,44 +266,37 @@ void Attribute::set_data_from(Attribute &&other)
other.center = Buffer(); other.center = Buffer();
other.size = 0; other.size = 0;
modified = true; modified = true;
};
/* If topology or step count differ, take the whole thing. */
if (size != other.size || motion.size() != other.motion.size()) {
take_all();
return; return;
} }
/* Same shape: compare each step independently, taking only those that differ /* Compare each step independently. */
* so that unchanged, implicitly-shared steps keep their sharing. */
const auto take_step = [&](Buffer &dst, Buffer &src, const AttributeElement step_element) { const auto take_step = [&](Buffer &dst, Buffer &src, const AttributeElement step_element) {
free_step_buffer(dst, step_element, sz, size); free_step_buffer(dst, step_element, element_size, size);
dst = src; dst = src;
src = Buffer(); src = Buffer();
modified = true; modified = true;
}; };
const auto step_differs = [&](const Buffer &a, const Buffer &b) { const auto step_equals = [&](const Buffer &a, const Buffer &b) {
if (a.data == b.data) { if (a.data == b.data) {
/* Same buffer, e.g. shared through implicit sharing. */ /* Same buffer, e.g. shared through implicit sharing. */
return false;
}
if (a.sharing_info != b.sharing_info) {
return true; return true;
} }
if (size == 0) { if (a.sharing_info != b.sharing_info) {
return false; return false;
} }
return memcmp(a.data, b.data, sz * size) != 0; if (size == 0) {
return true;
}
return memcmp(a.data, b.data, element_size * size) == 0;
}; };
if (step_differs(center, other.center)) { if (!step_equals(center, other.center)) {
take_step(center, other.center, element); take_step(center, other.center, element);
} }
for (size_t i = 0; i < motion.size(); i++) { for (size_t i = 0; i < motion.size(); i++) {
/* Motion steps are never voxels, strip that flag for the helper. */ if (!step_equals(motion[i], other.motion[i])) {
if (step_differs(motion[i], other.motion[i])) { take_step(motion[i], other.motion[i], element);
take_step(motion[i], other.motion[i], motion_element);
} }
} }
} }
@ -361,13 +360,7 @@ size_t Attribute::element_size(Geometry *geom,
size = pointcloud->num_points(); size = pointcloud->num_points();
} }
break; break;
case ATTR_ELEMENT_VERTEX_MOTION:
case ATTR_ELEMENT_VERTEX_NORMAL_MOTION:
if (geom->is_pointcloud()) {
PointCloud *pointcloud = static_cast<PointCloud *>(geom);
size = pointcloud->num_points() * (pointcloud->get_motion_steps() - 1);
}
break;
case ATTR_ELEMENT_FACE: case ATTR_ELEMENT_FACE:
if (geom->is_mesh() || geom->is_volume()) { if (geom->is_mesh() || geom->is_volume()) {
Mesh *mesh = static_cast<Mesh *>(geom); Mesh *mesh = static_cast<Mesh *>(geom);
@ -405,14 +398,6 @@ size_t Attribute::element_size(Geometry *geom,
size = hair->num_keys(); size = hair->num_keys();
} }
break; break;
case ATTR_ELEMENT_CURVE_KEY_MOTION:
case ATTR_ELEMENT_CURVE_KEY_NORMAL_MOTION:
if (geom->is_hair()) {
Hair *hair = static_cast<Hair *>(geom);
DCHECK_GT(hair->get_motion_steps(), 0);
size = hair->num_keys() * (hair->get_motion_steps() - 1);
}
break;
default: default:
break; break;
} }
@ -477,11 +462,6 @@ const char *Attribute::standard_name(AttributeStandard std)
return "undisplaced"; return "undisplaced";
case ATTR_STD_NORMAL_UNDISPLACED: case ATTR_STD_NORMAL_UNDISPLACED:
return "undisplaced_N"; return "undisplaced_N";
case ATTR_STD_MOTION_VERTEX_POSITION:
return "motion_P";
case ATTR_STD_MOTION_VERTEX_NORMAL:
case ATTR_STD_MOTION_CORNER_NORMAL:
return "motion_N";
case ATTR_STD_PARTICLE: case ATTR_STD_PARTICLE:
return "particle"; return "particle";
case ATTR_STD_CURVE_INTERCEPT: case ATTR_STD_CURVE_INTERCEPT:
@ -717,14 +697,8 @@ static TypeDesc find_type_from_geometry_std(Geometry *geometry, AttributeStandar
case ATTR_STD_POSITION_UNDEFORMED: case ATTR_STD_POSITION_UNDEFORMED:
case ATTR_STD_POSITION_UNDISPLACED: case ATTR_STD_POSITION_UNDISPLACED:
return TypePoint; return TypePoint;
case ATTR_STD_MOTION_VERTEX_POSITION:
return TypePoint;
case ATTR_STD_MOTION_VERTEX_NORMAL:
return TypeNormal;
case ATTR_STD_CORNER_NORMAL: case ATTR_STD_CORNER_NORMAL:
return TypeNormal; return TypeNormal;
case ATTR_STD_MOTION_CORNER_NORMAL:
return TypeNormal;
case ATTR_STD_PTEX_FACE_ID: case ATTR_STD_PTEX_FACE_ID:
return TypeFloat; return TypeFloat;
case ATTR_STD_PTEX_UV: case ATTR_STD_PTEX_UV:
@ -750,8 +724,6 @@ static TypeDesc find_type_from_geometry_std(Geometry *geometry, AttributeStandar
return TypeFloat2; return TypeFloat2;
case ATTR_STD_GENERATED: case ATTR_STD_GENERATED:
return TypePoint; return TypePoint;
case ATTR_STD_MOTION_VERTEX_POSITION:
return TypeFloat4;
case ATTR_STD_POINT_RANDOM: case ATTR_STD_POINT_RANDOM:
return TypeFloat; return TypeFloat;
case ATTR_STD_GENERATED_TRANSFORM: case ATTR_STD_GENERATED_TRANSFORM:
@ -796,14 +768,10 @@ static TypeDesc find_type_from_geometry_std(Geometry *geometry, AttributeStandar
return TypeFloat; return TypeFloat;
case ATTR_STD_VERTEX_NORMAL: case ATTR_STD_VERTEX_NORMAL:
return TypeNormal; return TypeNormal;
case ATTR_STD_MOTION_VERTEX_NORMAL:
return TypeNormal;
case ATTR_STD_UV: case ATTR_STD_UV:
return TypeFloat2; return TypeFloat2;
case ATTR_STD_GENERATED: case ATTR_STD_GENERATED:
return TypePoint; return TypePoint;
case ATTR_STD_MOTION_VERTEX_POSITION:
return TypeFloat4;
case ATTR_STD_CURVE_INTERCEPT: case ATTR_STD_CURVE_INTERCEPT:
return TypeFloat; return TypeFloat;
case ATTR_STD_CURVE_LENGTH: case ATTR_STD_CURVE_LENGTH:
@ -851,14 +819,8 @@ static AttributeElement find_element_from_geometry_std(Geometry *geometry, Attri
case ATTR_STD_POSITION_UNDEFORMED: case ATTR_STD_POSITION_UNDEFORMED:
case ATTR_STD_POSITION_UNDISPLACED: case ATTR_STD_POSITION_UNDISPLACED:
return ATTR_ELEMENT_VERTEX; return ATTR_ELEMENT_VERTEX;
case ATTR_STD_MOTION_VERTEX_POSITION:
return ATTR_ELEMENT_VERTEX_MOTION;
case ATTR_STD_MOTION_VERTEX_NORMAL:
return ATTR_ELEMENT_VERTEX_NORMAL_MOTION;
case ATTR_STD_CORNER_NORMAL: case ATTR_STD_CORNER_NORMAL:
return ATTR_ELEMENT_CORNER_NORMAL; return ATTR_ELEMENT_CORNER_NORMAL;
case ATTR_STD_MOTION_CORNER_NORMAL:
return ATTR_ELEMENT_CORNER_NORMAL_MOTION;
case ATTR_STD_PTEX_FACE_ID: case ATTR_STD_PTEX_FACE_ID:
return ATTR_ELEMENT_FACE; return ATTR_ELEMENT_FACE;
case ATTR_STD_PTEX_UV: case ATTR_STD_PTEX_UV:
@ -884,8 +846,6 @@ static AttributeElement find_element_from_geometry_std(Geometry *geometry, Attri
return ATTR_ELEMENT_VERTEX; return ATTR_ELEMENT_VERTEX;
case ATTR_STD_GENERATED: case ATTR_STD_GENERATED:
return ATTR_ELEMENT_VERTEX; return ATTR_ELEMENT_VERTEX;
case ATTR_STD_MOTION_VERTEX_POSITION:
return ATTR_ELEMENT_VERTEX_MOTION;
case ATTR_STD_POINT_RANDOM: case ATTR_STD_POINT_RANDOM:
return ATTR_ELEMENT_VERTEX; return ATTR_ELEMENT_VERTEX;
case ATTR_STD_GENERATED_TRANSFORM: case ATTR_STD_GENERATED_TRANSFORM:
@ -930,14 +890,10 @@ static AttributeElement find_element_from_geometry_std(Geometry *geometry, Attri
return ATTR_ELEMENT_CURVE_KEY; return ATTR_ELEMENT_CURVE_KEY;
case ATTR_STD_VERTEX_NORMAL: case ATTR_STD_VERTEX_NORMAL:
return ATTR_ELEMENT_CURVE_KEY_NORMAL; return ATTR_ELEMENT_CURVE_KEY_NORMAL;
case ATTR_STD_MOTION_VERTEX_NORMAL:
return ATTR_ELEMENT_CURVE_KEY_NORMAL_MOTION;
case ATTR_STD_UV: case ATTR_STD_UV:
return ATTR_ELEMENT_CURVE; return ATTR_ELEMENT_CURVE;
case ATTR_STD_GENERATED: case ATTR_STD_GENERATED:
return ATTR_ELEMENT_CURVE; return ATTR_ELEMENT_CURVE;
case ATTR_STD_MOTION_VERTEX_POSITION:
return ATTR_ELEMENT_CURVE_KEY_MOTION;
case ATTR_STD_CURVE_INTERCEPT: case ATTR_STD_CURVE_INTERCEPT:
return ATTR_ELEMENT_CURVE_KEY; return ATTR_ELEMENT_CURVE_KEY;
case ATTR_STD_CURVE_LENGTH: case ATTR_STD_CURVE_LENGTH:
@ -1006,6 +962,9 @@ Attribute &AttributeSet::copy(const Attribute &attr)
{ {
Attribute &copy_attr = *add(attr.name, attr.type, attr.element); Attribute &copy_attr = *add(attr.name, attr.type, attr.element);
copy_attr.std = attr.std; copy_attr.std = attr.std;
if (attr.has_motion()) {
copy_attr.add_motion(geometry);
}
return copy_attr; return copy_attr;
} }

2
intern/cycles/scene/devicescene.cpp
View file

@ -20,9 +20,7 @@ DeviceScene::DeviceScene(Device *device)
tri_shader(device, "tri_shader", MEM_GLOBAL), tri_shader(device, "tri_shader", MEM_GLOBAL),
tri_vindex(device, "tri_vindex", MEM_GLOBAL), tri_vindex(device, "tri_vindex", MEM_GLOBAL),
curves(device, "curves", MEM_GLOBAL), curves(device, "curves", MEM_GLOBAL),
curve_keys(device, "curve_keys", MEM_GLOBAL),
curve_segments(device, "curve_segments", MEM_GLOBAL), curve_segments(device, "curve_segments", MEM_GLOBAL),
points(device, "points", MEM_GLOBAL),
points_shader(device, "points_shader", MEM_GLOBAL), points_shader(device, "points_shader", MEM_GLOBAL),
objects(device, "objects", MEM_GLOBAL), objects(device, "objects", MEM_GLOBAL),
object_motion_pass(device, "object_motion_pass", MEM_GLOBAL), object_motion_pass(device, "object_motion_pass", MEM_GLOBAL),

2
intern/cycles/scene/devicescene.h
View file

@ -30,11 +30,9 @@ class DeviceScene {
device_vector<packed_uint3> tri_vindex; device_vector<packed_uint3> tri_vindex;
device_vector<KernelCurve> curves; device_vector<KernelCurve> curves;
device_vector<float4> curve_keys;
device_vector<KernelCurveSegment> curve_segments; device_vector<KernelCurveSegment> curve_segments;
/* point-cloud */ /* point-cloud */
device_vector<float4> points;
device_vector<uint> points_shader; device_vector<uint> points_shader;
/* objects */ /* objects */

51
intern/cycles/scene/geometry.cpp
View file

@ -362,7 +362,6 @@ void GeometryManager::geom_calc_offset(Scene *scene, BVHLayout bvh_layout)
size_t tri_size = 0; size_t tri_size = 0;
size_t curve_size = 0; size_t curve_size = 0;
size_t curve_key_size = 0;
size_t curve_segment_size = 0; size_t curve_segment_size = 0;
size_t point_size = 0; size_t point_size = 0;
@ -392,12 +391,10 @@ void GeometryManager::geom_calc_offset(Scene *scene, BVHLayout bvh_layout)
Hair *hair = static_cast<Hair *>(geom); Hair *hair = static_cast<Hair *>(geom);
prim_offset_changed = (hair->curve_segment_offset != curve_segment_size); prim_offset_changed = (hair->curve_segment_offset != curve_segment_size);
hair->curve_key_offset = curve_key_size;
hair->curve_segment_offset = curve_segment_size; hair->curve_segment_offset = curve_segment_size;
hair->prim_offset = curve_size; hair->prim_offset = curve_size;
curve_size += hair->num_curves(); curve_size += hair->num_curves();
curve_key_size += hair->num_keys();
curve_segment_size += hair->num_segments(); curve_segment_size += hair->num_segments();
} }
else if (geom->is_pointcloud()) { else if (geom->is_pointcloud()) {
@ -622,12 +619,10 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
if (device_update_flags & DEVICE_CURVE_DATA_NEEDS_REALLOC) { if (device_update_flags & DEVICE_CURVE_DATA_NEEDS_REALLOC) {
dscene->curves.tag_realloc(); dscene->curves.tag_realloc();
dscene->curve_keys.tag_realloc();
dscene->curve_segments.tag_realloc(); dscene->curve_segments.tag_realloc();
} }
if (device_update_flags & DEVICE_POINT_DATA_NEEDS_REALLOC) { if (device_update_flags & DEVICE_POINT_DATA_NEEDS_REALLOC) {
dscene->points.tag_realloc();
dscene->points_shader.tag_realloc(); dscene->points_shader.tag_realloc();
} }
} }
@ -691,13 +686,11 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
} }
if (device_update_flags & DEVICE_CURVE_DATA_MODIFIED) { if (device_update_flags & DEVICE_CURVE_DATA_MODIFIED) {
dscene->curve_keys.tag_modified();
dscene->curves.tag_modified(); dscene->curves.tag_modified();
dscene->curve_segments.tag_modified(); dscene->curve_segments.tag_modified();
} }
if (device_update_flags & DEVICE_POINT_DATA_MODIFIED) { if (device_update_flags & DEVICE_POINT_DATA_MODIFIED) {
dscene->points.tag_modified();
dscene->points_shader.tag_modified(); dscene->points_shader.tag_modified();
} }
@ -984,7 +977,13 @@ void GeometryManager::device_update(Device *device,
return; return;
} }
{ /* Attributes must be uploaded to the device before displacement and hair shadow
* transparency, which run shader evaluation. Otherwise the upload is deferred
* until after BVH building. */
const bool need_attributes_before_bvh = true_displacement_used ||
curve_need_update_shadow_transparency;
if (need_attributes_before_bvh) {
const scoped_callback_timer timer([scene](double time) { const scoped_callback_timer timer([scene](double time) {
if (scene->update_stats) { if (scene->update_stats) {
scene->update_stats->geometry.times.add_entry({"device_update (attributes)", time}); scene->update_stats->geometry.times.add_entry({"device_update (attributes)", time});
@ -1057,20 +1056,12 @@ void GeometryManager::device_update(Device *device,
return; return;
} }
/* Device re-update after applying transforms and displacement. */ /* Displacement and hair shadow transparency modified the geometry. Free device buffers
* that need to be reallocated, so the BVH and attributes are rebuilt from the updated
* data below. Freeing the attribute buffers uploaded earlier also keeps them from
* overlapping with the temporary BVH building buffers, lowering peak memory usage. */
if (displacement_done || curve_shadow_transparency_done) { if (displacement_done || curve_shadow_transparency_done) {
const scoped_callback_timer timer([scene](double time) {
if (scene->update_stats) {
scene->update_stats->geometry.times.add_entry(
{"device_update (displacement: attributes)", time});
}
});
device_free(device, dscene, false); device_free(device, dscene, false);
device_update_attributes(device, dscene, scene, progress);
if (progress.get_cancel()) {
return;
}
} }
/* Update the BVH even when there is no geometry so the kernel's BVH data is still valid, /* Update the BVH even when there is no geometry so the kernel's BVH data is still valid,
@ -1154,6 +1145,22 @@ void GeometryManager::device_update(Device *device,
dscene->data.bvh.bvh_layout = BVHParams::best_bvh_layout( dscene->data.bvh.bvh_layout = BVHParams::best_bvh_layout(
scene->params.bvh_layout, device->get_bvh_layout_mask(dscene->data.kernel_features)); scene->params.bvh_layout, device->get_bvh_layout_mask(dscene->data.kernel_features));
/* Upload attributes to the device.
*
* This is deferred until after BVH building so the attribute buffers do not overlap with
* the temporary buffers allocated during BVH building. */
{
const scoped_callback_timer timer([scene](double time) {
if (scene->update_stats) {
scene->update_stats->geometry.times.add_entry({"device_update (attributes)", time});
}
});
device_update_attributes(device, dscene, scene, progress);
if (progress.get_cancel()) {
return;
}
}
{ {
const scoped_callback_timer timer([scene](double time) { const scoped_callback_timer timer([scene](double time) {
if (scene->update_stats) { if (scene->update_stats) {
@ -1194,9 +1201,7 @@ void GeometryManager::device_update(Device *device,
dscene->tri_shader.clear_modified(); dscene->tri_shader.clear_modified();
dscene->tri_vindex.clear_modified(); dscene->tri_vindex.clear_modified();
dscene->curves.clear_modified(); dscene->curves.clear_modified();
dscene->curve_keys.clear_modified();
dscene->curve_segments.clear_modified(); dscene->curve_segments.clear_modified();
dscene->points.clear_modified();
dscene->points_shader.clear_modified(); dscene->points_shader.clear_modified();
dscene->attributes_map.clear_modified(); dscene->attributes_map.clear_modified();
dscene->attributes_float.clear_modified(); dscene->attributes_float.clear_modified();
@ -1220,9 +1225,7 @@ void GeometryManager::device_free(Device *device, DeviceScene *dscene, bool forc
dscene->tri_shader.free_if_need_realloc(force_free); dscene->tri_shader.free_if_need_realloc(force_free);
dscene->tri_vindex.free_if_need_realloc(force_free); dscene->tri_vindex.free_if_need_realloc(force_free);
dscene->curves.free_if_need_realloc(force_free); 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->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->points_shader.free_if_need_realloc(force_free);
dscene->attributes_map.free_if_need_realloc(force_free); dscene->attributes_map.free_if_need_realloc(force_free);
dscene->attributes_float.free_if_need_realloc(force_free); dscene->attributes_float.free_if_need_realloc(force_free);

126
intern/cycles/scene/geometry_attributes.cpp
View file

@ -327,12 +327,24 @@ class AttributeTableBuilder {
return; return;
} }
/* store element and type */ /* For hair and pointcloud, pack combined position + radius as float4. */
if (mattr->std == ATTR_STD_POSITION && (geom->is_hair() || geom->is_pointcloud())) {
add_position_radius(geom, mattr, type, desc);
return;
}
if (mattr->std == ATTR_STD_RADIUS && (geom->is_hair() || geom->is_pointcloud())) {
desc.element = ATTR_ELEMENT_NONE;
desc.offset = 0;
return;
}
/* Store element and type. */
desc.element = mattr->element; desc.element = mattr->element;
type = mattr->type; type = mattr->type;
/* store attribute data in arrays */ /* Store attribute data in arrays, including possible motion steps. */
const size_t size = Attribute::element_size(geom, mattr->element, prim); const size_t per_step = Attribute::element_size(geom, mattr->element, prim);
const int num_motion = mattr->motion.size();
const AttributeElement &element = desc.element; const AttributeElement &element = desc.element;
int &offset = desc.offset; int &offset = desc.offset;
@ -343,42 +355,53 @@ class AttributeTableBuilder {
offset = handle.kernel_id(); offset = handle.kernel_id();
} }
else if (mattr->element & ATTR_ELEMENT_IS_BYTE) { else if (mattr->element & ATTR_ELEMENT_IS_BYTE) {
offset = attr_uchar4.add(mattr->data<uchar4>(), size, mattr->modified); offset = attr_uchar4.add(mattr->data<uchar4>(), per_step, mattr->modified);
for (int step = 1; step <= num_motion; step++) {
attr_uchar4.add(mattr->data<uchar4>(step), per_step, mattr->modified);
}
} }
else if (mattr->element & ATTR_ELEMENT_IS_NORMAL) { else if (mattr->element & ATTR_ELEMENT_IS_NORMAL) {
offset = attr_normal.add(mattr->data<packed_normal>(), size, mattr->modified); offset = attr_normal.add(mattr->data<packed_normal>(), per_step, mattr->modified);
for (int step = 0; step < int(mattr->motion.size()); step++) { for (int step = 1; step <= num_motion; step++) {
attr_normal.add(mattr->data<packed_normal>(step + 1), size, mattr->modified); attr_normal.add(mattr->data<packed_normal>(step), per_step, mattr->modified);
} }
} }
else if (mattr->type == TypeFloat) { else if (mattr->type == TypeFloat) {
offset = attr_float.add(mattr->data<float>(), size, mattr->modified); offset = attr_float.add(mattr->data<float>(), per_step, mattr->modified);
for (int step = 0; step < int(mattr->motion.size()); step++) { for (int step = 1; step <= num_motion; step++) {
attr_float.add(mattr->data<float>(step + 1), size, mattr->modified); attr_float.add(mattr->data<float>(step), per_step, mattr->modified);
} }
} }
else if (mattr->type == TypeFloat2) { else if (mattr->type == TypeFloat2) {
offset = attr_float2.add(mattr->data<float2>(), size, mattr->modified); offset = attr_float2.add(mattr->data<float2>(), per_step, mattr->modified);
for (int step = 1; step <= num_motion; step++) {
attr_float2.add(mattr->data<float2>(step), per_step, mattr->modified);
}
} }
else if (mattr->type == TypeMatrix) { else if (mattr->type == TypeMatrix) {
offset = attr_float4.add((float4 *)mattr->data<Transform>(), size * 3, mattr->modified); offset = attr_float4.add(
(const float4 *)mattr->data<Transform>(), per_step * 3, mattr->modified);
for (int step = 1; step <= num_motion; step++) {
attr_float4.add(
(const float4 *)mattr->data<Transform>(step), per_step * 3, mattr->modified);
}
} }
else if (mattr->type == TypeFloat4 || mattr->type == TypeRGBA) { else if (mattr->type == TypeFloat4 || mattr->type == TypeRGBA) {
offset = attr_float4.add(mattr->data<float4>(), size, mattr->modified); offset = attr_float4.add(mattr->data<float4>(), per_step, mattr->modified);
for (int step = 0; step < int(mattr->motion.size()); step++) { for (int step = 1; step <= num_motion; step++) {
attr_float4.add(mattr->data<float4>(step + 1), size, mattr->modified); attr_float4.add(mattr->data<float4>(step), per_step, mattr->modified);
} }
} }
else { else {
offset = attr_float3.add(mattr->data<packed_float3>(), size, mattr->modified); offset = attr_float3.add(mattr->data<packed_float3>(), per_step, mattr->modified);
for (int step = 0; step < int(mattr->motion.size()); step++) { for (int step = 1; step <= num_motion; step++) {
attr_float3.add(mattr->data<packed_float3>(step + 1), size, mattr->modified); attr_float3.add(mattr->data<packed_float3>(step), per_step, mattr->modified);
} }
} }
/* mesh vertex/curve index is global, not per object, so we sneak /* Primitive index is global, not per object, so we sneak a correction
* a correction for that in here */ * for that in here. Vertex index is per object. */
if (geom->is_mesh()) { if (geom->is_mesh() || geom->is_volume()) {
Mesh *mesh = static_cast<Mesh *>(geom); Mesh *mesh = static_cast<Mesh *>(geom);
if (element & ATTR_ELEMENT_FACE) { if (element & ATTR_ELEMENT_FACE) {
offset -= mesh->prim_offset; offset -= mesh->prim_offset;
@ -392,9 +415,6 @@ class AttributeTableBuilder {
if (element & ATTR_ELEMENT_CURVE) { if (element & ATTR_ELEMENT_CURVE) {
offset -= hair->prim_offset; offset -= hair->prim_offset;
} }
else if (element & ATTR_ELEMENT_CURVE_KEY) {
offset -= hair->curve_key_offset;
}
} }
else if (geom->is_pointcloud()) { else if (geom->is_pointcloud()) {
if (element & ATTR_ELEMENT_VERTEX) { if (element & ATTR_ELEMENT_VERTEX) {
@ -409,9 +429,10 @@ class AttributeTableBuilder {
return; return;
} }
const size_t base_size = Attribute::element_size(geom, mattr->element, prim); /* Must match the number of steps written by add(), which is derived from
/* Inline motion: reserve space for center step plus each motion sub-step. */ * the attribute's own stored motion steps rather than geom->get_motion_steps(). */
const size_t size = base_size * (1 + mattr->motion.size()); const int steps = mattr->num_motion_steps();
const size_t size = Attribute::element_size(geom, mattr->element, prim) * steps;
if (mattr->element & ATTR_ELEMENT_VOXEL) { if (mattr->element & ATTR_ELEMENT_VOXEL) {
/* pass */ /* pass */
@ -439,6 +460,47 @@ class AttributeTableBuilder {
} }
} }
/* Pack combined position + radius for hair and point cloud. */
void add_position_radius(Geometry *geom,
Attribute *attr_P,
TypeDesc &type,
AttributeDescriptor &desc)
{
Attribute *attr_R = geom->attributes.find(ATTR_STD_RADIUS);
const size_t base_size = attr_P->size;
const int steps = attr_P->has_motion() ? geom->get_motion_steps() : 1;
const size_t total_size = base_size * steps;
vector<float4> combined(total_size);
for (int step = 0; step < steps; step++) {
const packed_float3 *P = attr_P->data<packed_float3>(step);
const float *R = attr_R->data<float>(step);
const size_t dst_offset = size_t(step) * base_size;
for (size_t i = 0; i < base_size; i++) {
combined[dst_offset + i] = make_float4(P[i], R[i]);
}
}
desc.element = attr_P->element;
type = TypeFloat4;
int &offset = desc.offset;
const bool modified = attr_P->modified || attr_R->modified;
offset = attr_float4.add(combined.data(), total_size, modified);
/* Pointcloud uses global primitive index. */
if (geom->is_pointcloud()) {
offset -= geom->prim_offset;
}
}
void reserve_position_radius(Geometry *geom, Attribute *attr_P)
{
const int steps = attr_P->has_motion() ? geom->get_motion_steps() : 1;
const size_t total_size = attr_P->size * steps;
attr_float4.reserve(total_size);
}
void alloc() void alloc()
{ {
attr_float.alloc(); attr_float.alloc();
@ -546,6 +608,16 @@ void GeometryManager::device_update_attributes(Device *device,
AttributeRequestSet &attributes = geom_attributes[i]; AttributeRequestSet &attributes = geom_attributes[i];
for (AttributeRequest &req : attributes.requests) { for (AttributeRequest &req : attributes.requests) {
Attribute *attr = geom->attributes.find(req); Attribute *attr = geom->attributes.find(req);
if (attr && (geom->is_hair() || geom->is_pointcloud())) {
/* Special cases for packed position + radius. */
if (attr->std == ATTR_STD_POSITION) {
builder.reserve_position_radius(geom, attr);
continue;
}
if (attr->std == ATTR_STD_RADIUS) {
continue;
}
}
builder.reserve(geom, attr, ATTR_PRIM_GEOMETRY); builder.reserve(geom, attr, ATTR_PRIM_GEOMETRY);
} }
} }

22
intern/cycles/scene/geometry_mesh.cpp
View file

@ -31,7 +31,6 @@ void GeometryManager::device_update_mesh(Device * /*unused*/,
/* Count. */ /* Count. */
size_t tri_size = 0; size_t tri_size = 0;
size_t curve_key_size = 0;
size_t curve_size = 0; size_t curve_size = 0;
size_t curve_segment_size = 0; size_t curve_segment_size = 0;
@ -46,7 +45,6 @@ void GeometryManager::device_update_mesh(Device * /*unused*/,
else if (geom->is_hair()) { else if (geom->is_hair()) {
Hair *hair = static_cast<Hair *>(geom); Hair *hair = static_cast<Hair *>(geom);
curve_key_size += hair->num_keys();
curve_size += hair->num_curves(); curve_size += hair->num_curves();
curve_segment_size += hair->num_segments(); curve_segment_size += hair->num_segments();
} }
@ -96,38 +94,31 @@ void GeometryManager::device_update_mesh(Device * /*unused*/,
if (curve_segment_size != 0) { if (curve_segment_size != 0) {
progress.set_status("Updating Mesh", "Copying Curves to device"); progress.set_status("Updating Mesh", "Copying Curves to device");
float4 *curve_keys = dscene->curve_keys.alloc(curve_key_size);
KernelCurve *curves = dscene->curves.alloc(curve_size); KernelCurve *curves = dscene->curves.alloc(curve_size);
KernelCurveSegment *curve_segments = dscene->curve_segments.alloc(curve_segment_size); KernelCurveSegment *curve_segments = dscene->curve_segments.alloc(curve_segment_size);
const bool copy_all_data = dscene->curve_keys.need_realloc() || const bool copy_all_data = dscene->curves.need_realloc() ||
dscene->curves.need_realloc() ||
dscene->curve_segments.need_realloc(); dscene->curve_segments.need_realloc();
for (Geometry *geom : scene->geometry) { for (Geometry *geom : scene->geometry) {
if (geom->is_hair()) { if (geom->is_hair()) {
Hair *hair = static_cast<Hair *>(geom); Hair *hair = static_cast<Hair *>(geom);
const bool curve_keys_co_modified = hair->radius_is_modified() ||
hair->position_is_modified();
const bool curve_data_modified = hair->curve_shader_is_modified() || const bool curve_data_modified = hair->curve_shader_is_modified() ||
hair->curve_first_key_is_modified(); hair->curve_first_key_is_modified();
if (!curve_keys_co_modified && !curve_data_modified && !copy_all_data) { if (!curve_data_modified && !copy_all_data) {
continue; continue;
} }
hair->pack_curves(scene, hair->pack_curves(
&curve_keys[hair->curve_key_offset], scene, &curves[hair->prim_offset], &curve_segments[hair->curve_segment_offset]);
&curves[hair->prim_offset],
&curve_segments[hair->curve_segment_offset]);
if (progress.get_cancel()) { if (progress.get_cancel()) {
return; return;
} }
} }
} }
dscene->curve_keys.copy_to_device_if_modified();
dscene->curves.copy_to_device_if_modified(); dscene->curves.copy_to_device_if_modified();
dscene->curve_segments.copy_to_device_if_modified(); dscene->curve_segments.copy_to_device_if_modified();
} }
@ -135,21 +126,18 @@ void GeometryManager::device_update_mesh(Device * /*unused*/,
if (point_size != 0) { if (point_size != 0) {
progress.set_status("Updating Mesh", "Copying Point clouds to device"); 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); uint *points_shader = dscene->points_shader.alloc(point_size);
for (Geometry *geom : scene->geometry) { for (Geometry *geom : scene->geometry) {
if (geom->is_pointcloud()) { if (geom->is_pointcloud()) {
PointCloud *pointcloud = static_cast<PointCloud *>(geom); PointCloud *pointcloud = static_cast<PointCloud *>(geom);
pointcloud->pack( pointcloud->pack(scene, &points_shader[pointcloud->prim_offset]);
scene, &points[pointcloud->prim_offset], &points_shader[pointcloud->prim_offset]);
if (progress.get_cancel()) { if (progress.get_cancel()) {
return; return;
} }
} }
} }
dscene->points.copy_to_device();
dscene->points_shader.copy_to_device(); dscene->points_shader.copy_to_device();
} }
} }

351
intern/cycles/scene/hair.cpp
View file

@ -42,16 +42,16 @@ void Hair::Curve::bounds_grow(const int k, const float4 *keys, BoundBox &bounds)
} }
void Hair::Curve::bounds_grow(const int k, void Hair::Curve::bounds_grow(const int k,
const float3 *positions, const packed_float3 *curve_keys,
const float *radius, const float *curve_radius,
BoundBox &bounds) const BoundBox &bounds) const
{ {
float3 P[4]; float3 P[4];
P[0] = positions[max(first_key + k - 1, first_key)]; P[0] = curve_keys[max(first_key + k - 1, first_key)];
P[1] = positions[first_key + k]; P[1] = curve_keys[first_key + k];
P[2] = positions[first_key + k + 1]; P[2] = curve_keys[first_key + k + 1];
P[3] = positions[min(first_key + k + 2, first_key + num_keys - 1)]; P[3] = curve_keys[min(first_key + k + 2, first_key + num_keys - 1)];
float3 lower; float3 lower;
float3 upper; float3 upper;
@ -60,24 +60,24 @@ void Hair::Curve::bounds_grow(const int k,
curvebounds(&lower.y, &upper.y, P, 1); curvebounds(&lower.y, &upper.y, P, 1);
curvebounds(&lower.z, &upper.z, P, 2); curvebounds(&lower.z, &upper.z, P, 2);
const float mr = max(radius[first_key + k], radius[first_key + k + 1]); const float mr = max(curve_radius[first_key + k], curve_radius[first_key + k + 1]);
bounds.grow(lower, mr); bounds.grow(lower, mr);
bounds.grow(upper, mr); bounds.grow(upper, mr);
} }
void Hair::Curve::bounds_grow(const int k, void Hair::Curve::bounds_grow(const int k,
const float3 *positions, const packed_float3 *curve_keys,
const float *radius, const float *curve_radius,
const Transform &aligned_space, const Transform &aligned_space,
BoundBox &bounds) const BoundBox &bounds) const
{ {
float3 P[4]; float3 P[4];
P[0] = positions[max(first_key + k - 1, first_key)]; P[0] = curve_keys[max(first_key + k - 1, first_key)];
P[1] = positions[first_key + k]; P[1] = curve_keys[first_key + k];
P[2] = positions[first_key + k + 1]; P[2] = curve_keys[first_key + k + 1];
P[3] = positions[min(first_key + k + 2, first_key + num_keys - 1)]; P[3] = curve_keys[min(first_key + k + 2, first_key + num_keys - 1)];
P[0] = transform_point(&aligned_space, P[0]); P[0] = transform_point(&aligned_space, P[0]);
P[1] = transform_point(&aligned_space, P[1]); P[1] = transform_point(&aligned_space, P[1]);
@ -91,63 +91,7 @@ void Hair::Curve::bounds_grow(const int k,
curvebounds(&lower.y, &upper.y, P, 1); curvebounds(&lower.y, &upper.y, P, 1);
curvebounds(&lower.z, &upper.z, P, 2); curvebounds(&lower.z, &upper.z, P, 2);
const float mr = max(radius[first_key + k], radius[first_key + k + 1]); const float mr = max(curve_radius[first_key + k], curve_radius[first_key + k + 1]);
bounds.grow(lower, mr);
bounds.grow(upper, mr);
}
void Hair::Curve::bounds_grow(const int k,
const packed_float3 *positions,
const float *radius,
BoundBox &bounds) const
{
float3 P[4];
P[0] = positions[max(first_key + k - 1, first_key)];
P[1] = positions[first_key + k];
P[2] = positions[first_key + k + 1];
P[3] = positions[min(first_key + k + 2, first_key + num_keys - 1)];
float3 lower;
float3 upper;
curvebounds(&lower.x, &upper.x, P, 0);
curvebounds(&lower.y, &upper.y, P, 1);
curvebounds(&lower.z, &upper.z, P, 2);
const float mr = max(radius[first_key + k], radius[first_key + k + 1]);
bounds.grow(lower, mr);
bounds.grow(upper, mr);
}
void Hair::Curve::bounds_grow(const int k,
const packed_float3 *positions,
const float *radius,
const Transform &aligned_space,
BoundBox &bounds) const
{
float3 P[4];
P[0] = positions[max(first_key + k - 1, first_key)];
P[1] = positions[first_key + k];
P[2] = positions[first_key + k + 1];
P[3] = positions[min(first_key + k + 2, first_key + num_keys - 1)];
P[0] = transform_point(&aligned_space, P[0]);
P[1] = transform_point(&aligned_space, P[1]);
P[2] = transform_point(&aligned_space, P[2]);
P[3] = transform_point(&aligned_space, P[3]);
float3 lower;
float3 upper;
curvebounds(&lower.x, &upper.x, P, 0);
curvebounds(&lower.y, &upper.y, P, 1);
curvebounds(&lower.z, &upper.z, P, 2);
const float mr = max(radius[first_key + k], radius[first_key + k + 1]);
bounds.grow(lower, mr); bounds.grow(lower, mr);
bounds.grow(upper, mr); bounds.grow(upper, mr);
@ -175,10 +119,22 @@ void Hair::Curve::bounds_grow(const float4 keys[4], BoundBox &bounds) const
bounds.grow(upper, mr); bounds.grow(upper, mr);
} }
void Hair::Curve::motion_keys(const packed_float3 *positions, /* Get position and radius arrays for a given time-ordered motion step. */
const float *radius, static void hair_step_buffers(const Attribute *attr_P,
const float4 *key_steps, const Attribute *attr_R,
const size_t num_positions, const size_t step,
const packed_float3 *&P,
const float *&R)
{
/* Radius motion follows the position motion steps, if it has fewer steps the
* read falls back to the center radius. */
const int num_steps = attr_P->num_motion_steps();
P = attr_P->data_at_time_step<packed_float3>(step, num_steps);
R = attr_R->data_at_time_step<float>(step, num_steps);
}
void Hair::Curve::motion_keys(const Attribute *attr_P,
const Attribute *attr_R,
const size_t num_steps, const size_t num_steps,
const float time, const float time,
size_t k0, size_t k0,
@ -192,18 +148,15 @@ void Hair::Curve::motion_keys(const packed_float3 *positions,
/* Fetch vertex coordinates. */ /* Fetch vertex coordinates. */
float4 curr_keys[2]; float4 curr_keys[2];
float4 next_keys[2]; float4 next_keys[2];
keys_for_step(positions, radius, key_steps, num_positions, num_steps, step, k0, k1, curr_keys); keys_for_step(attr_P, attr_R, step, k0, k1, curr_keys);
keys_for_step( keys_for_step(attr_P, attr_R, step + 1, k0, k1, next_keys);
positions, radius, key_steps, num_positions, num_steps, step + 1, k0, k1, next_keys);
/* Interpolate between steps. */ /* Interpolate between steps. */
r_keys[0] = (1.0f - t) * curr_keys[0] + t * next_keys[0]; r_keys[0] = (1.0f - t) * curr_keys[0] + t * next_keys[0];
r_keys[1] = (1.0f - t) * curr_keys[1] + t * next_keys[1]; r_keys[1] = (1.0f - t) * curr_keys[1] + t * next_keys[1];
} }
void Hair::Curve::cardinal_motion_keys(const packed_float3 *positions, void Hair::Curve::cardinal_motion_keys(const Attribute *attr_P,
const float *radius, const Attribute *attr_R,
const float4 *key_steps,
const size_t num_positions,
const size_t num_steps, const size_t num_steps,
const float time, const float time,
size_t k0, size_t k0,
@ -219,10 +172,8 @@ void Hair::Curve::cardinal_motion_keys(const packed_float3 *positions,
/* Fetch vertex coordinates. */ /* Fetch vertex coordinates. */
float4 curr_keys[4]; float4 curr_keys[4];
float4 next_keys[4]; float4 next_keys[4];
cardinal_keys_for_step( cardinal_keys_for_step(attr_P, attr_R, step, k0, k1, k2, k3, curr_keys);
positions, radius, key_steps, num_positions, num_steps, step, k0, k1, k2, k3, curr_keys); cardinal_keys_for_step(attr_P, attr_R, step + 1, k0, k1, k2, k3, next_keys);
cardinal_keys_for_step(
positions, radius, key_steps, num_positions, num_steps, step + 1, k0, k1, k2, k3, next_keys);
/* Interpolate between steps. */ /* Interpolate between steps. */
r_keys[0] = (1.0f - t) * curr_keys[0] + t * next_keys[0]; r_keys[0] = (1.0f - t) * curr_keys[0] + t * next_keys[0];
r_keys[1] = (1.0f - t) * curr_keys[1] + t * next_keys[1]; r_keys[1] = (1.0f - t) * curr_keys[1] + t * next_keys[1];
@ -230,47 +181,25 @@ void Hair::Curve::cardinal_motion_keys(const packed_float3 *positions,
r_keys[3] = (1.0f - t) * curr_keys[3] + t * next_keys[3]; r_keys[3] = (1.0f - t) * curr_keys[3] + t * next_keys[3];
} }
void Hair::Curve::keys_for_step(const packed_float3 *positions, void Hair::Curve::keys_for_step(const Attribute *attr_P,
const float *radius, const Attribute *attr_R,
const float4 *key_steps, const size_t step,
const size_t num_positions,
const size_t num_steps,
size_t step,
size_t k0, size_t k0,
size_t k1, size_t k1,
float4 r_keys[2]) const float4 r_keys[2]) const
{ {
k0 = max(k0, (size_t)0); k0 = max(k0, (size_t)0);
k1 = min(k1, (size_t)(num_keys - 1)); k1 = min(k1, (size_t)(num_keys - 1));
const size_t center_step = ((num_steps - 1) / 2); const packed_float3 *P;
if (step == center_step) { const float *R;
/* Center step: regular key location. */ hair_step_buffers(attr_P, attr_R, step, P, R);
r_keys[0] = make_float4(positions[first_key + k0], radius[first_key + k0]); r_keys[0] = make_float4(P[first_key + k0], R[first_key + k0]);
r_keys[1] = make_float4(positions[first_key + k1], radius[first_key + k1]); r_keys[1] = make_float4(P[first_key + k1], R[first_key + k1]);
}
else {
/* Center step is not stored in this array. */
if (step > center_step) {
step--;
}
const size_t offset = first_key + step * num_positions;
r_keys[0] = make_float4(key_steps[offset + k0].x,
key_steps[offset + k0].y,
key_steps[offset + k0].z,
radius[first_key + k0]);
r_keys[1] = make_float4(key_steps[offset + k1].x,
key_steps[offset + k1].y,
key_steps[offset + k1].z,
radius[first_key + k1]);
}
} }
void Hair::Curve::cardinal_keys_for_step(const packed_float3 *positions, void Hair::Curve::cardinal_keys_for_step(const Attribute *attr_P,
const float *radius, const Attribute *attr_R,
const float4 *key_steps, const size_t step,
const size_t num_positions,
const size_t num_steps,
size_t step,
size_t k0, size_t k0,
size_t k1, size_t k1,
size_t k2, size_t k2,
@ -279,37 +208,13 @@ void Hair::Curve::cardinal_keys_for_step(const packed_float3 *positions,
{ {
k0 = max(k0, (size_t)0); k0 = max(k0, (size_t)0);
k3 = min(k3, (size_t)(num_keys - 1)); k3 = min(k3, (size_t)(num_keys - 1));
const size_t center_step = ((num_steps - 1) / 2); const packed_float3 *P;
if (step == center_step) { const float *R;
/* Center step: regular key location. */ hair_step_buffers(attr_P, attr_R, step, P, R);
r_keys[0] = make_float4(positions[first_key + k0], radius[first_key + k0]); r_keys[0] = make_float4(P[first_key + k0], R[first_key + k0]);
r_keys[1] = make_float4(positions[first_key + k1], radius[first_key + k1]); r_keys[1] = make_float4(P[first_key + k1], R[first_key + k1]);
r_keys[2] = make_float4(positions[first_key + k2], radius[first_key + k2]); r_keys[2] = make_float4(P[first_key + k2], R[first_key + k2]);
r_keys[3] = make_float4(positions[first_key + k3], radius[first_key + k3]); r_keys[3] = make_float4(P[first_key + k3], R[first_key + k3]);
}
else {
/* Center step is not stored in this array. */
if (step > center_step) {
step--;
}
const size_t offset = first_key + step * num_positions;
r_keys[0] = make_float4(key_steps[offset + k0].x,
key_steps[offset + k0].y,
key_steps[offset + k0].z,
radius[first_key + k0]);
r_keys[1] = make_float4(key_steps[offset + k1].x,
key_steps[offset + k1].y,
key_steps[offset + k1].z,
radius[first_key + k1]);
r_keys[2] = make_float4(key_steps[offset + k2].x,
key_steps[offset + k2].y,
key_steps[offset + k2].z,
radius[first_key + k2]);
r_keys[3] = make_float4(key_steps[offset + k3].x,
key_steps[offset + k3].y,
key_steps[offset + k3].z,
radius[first_key + k3]);
}
} }
/* Hair */ /* Hair */
@ -326,7 +231,6 @@ NODE_DEFINE(Hair)
Hair::Hair() : Geometry(get_node_type(), Geometry::HAIR) Hair::Hair() : Geometry(get_node_type(), Geometry::HAIR)
{ {
curve_key_offset = 0;
curve_segment_offset = 0; curve_segment_offset = 0;
curve_shape = CURVE_RIBBON; curve_shape = CURVE_RIBBON;
@ -335,12 +239,6 @@ Hair::Hair() : Geometry(get_node_type(), Geometry::HAIR)
Hair::~Hair() = default; Hair::~Hair() = default;
size_t Hair::num_keys() const
{
const Attribute *attr = attributes.find(ATTR_STD_POSITION);
return attr ? attr->size : 0;
}
void Hair::add_builtin_attributes() void Hair::add_builtin_attributes()
{ {
attributes.add(ATTR_STD_POSITION); attributes.add(ATTR_STD_POSITION);
@ -349,8 +247,10 @@ void Hair::add_builtin_attributes()
void Hair::resize_curves(const int numcurves, const int numkeys) void Hair::resize_curves(const int numcurves, const int numkeys)
{ {
attributes.add(ATTR_STD_POSITION)->resize(numkeys); Attribute *attr_P = attributes.add(ATTR_STD_POSITION);
attributes.add(ATTR_STD_RADIUS)->resize(numkeys); attr_P->resize(numkeys);
Attribute *attr_R = attributes.add(ATTR_STD_RADIUS);
attr_R->resize(numkeys);
curve_first_key.resize(numcurves); curve_first_key.resize(numcurves);
curve_shader.resize(numcurves); curve_shader.resize(numcurves);
@ -370,19 +270,24 @@ void Hair::clear(bool preserve_shaders)
void Hair::copy_center_to_motion_step(const int motion_step) void Hair::copy_center_to_motion_step(const int motion_step)
{ {
Attribute *attr_mP = attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); const int attr_step = motion_step + 1;
if (attr_mP) { const size_t numkeys = num_keys();
const packed_float3 *keys = get_position();
const size_t numkeys = num_keys(); Attribute *attr_P = attributes.find(ATTR_STD_POSITION);
std::copy_n(keys, numkeys, attr_mP->data_for_write<packed_float3>() + motion_step * numkeys); if (attr_P->has_motion()) {
std::copy_n(get_position(), numkeys, attr_P->data_for_write<packed_float3>(attr_step));
}
Attribute *attr_R = attributes.find(ATTR_STD_RADIUS);
if (attr_R->has_motion()) {
std::copy_n(get_radius(), numkeys, attr_R->data_for_write<float>(attr_step));
} }
Attribute *attr_mvN = attributes.find(ATTR_STD_MOTION_VERTEX_NORMAL);
Attribute *attr_vN = attributes.find(ATTR_STD_VERTEX_NORMAL); Attribute *attr_vN = attributes.find(ATTR_STD_VERTEX_NORMAL);
if (attr_mvN && attr_vN) { if (attr_vN && attr_vN->has_motion()) {
const packed_normal *vN = attr_vN->data<packed_normal>(); std::copy_n(attr_vN->data<packed_normal>(),
const size_t numkeys = num_keys(); numkeys,
std::copy_n(vN, numkeys, attr_mvN->data_for_write<packed_normal>() + motion_step * numkeys); attr_vN->data_for_write<packed_normal>(attr_step));
} }
} }
@ -405,12 +310,12 @@ void Hair::get_uv_tiles(ustring map, unordered_set<int> &tiles)
void Hair::compute_bounds() void Hair::compute_bounds()
{ {
BoundBox bnds = BoundBox::empty; BoundBox bnds = BoundBox::empty;
const size_t positions_size = num_keys(); const size_t curve_keys_size = num_keys();
const packed_float3 *curve_keys_data = get_position();
const float *curve_radius_data = get_radius();
const size_t curve_num = num_curves(); const size_t curve_num = num_curves();
const packed_float3 *positions = get_position();
const float *radius = get_radius();
if (positions_size > 0) { if (curve_keys_size > 0) {
bnds.grow(parallel_reduce( bnds.grow(parallel_reduce(
blocked_range<size_t>(0, curve_num), blocked_range<size_t>(0, curve_num),
BoundBox(BoundBox::empty), BoundBox(BoundBox::empty),
@ -420,7 +325,7 @@ void Hair::compute_bounds()
const Curve curve = get_curve(i); const Curve curve = get_curve(i);
const int num_segments = curve.num_segments(); const int num_segments = curve.num_segments();
for (int k = 0; k < num_segments; k++) { for (int k = 0; k < num_segments; k++) {
curve.bounds_grow(k, positions, radius, current_bounds); curve.bounds_grow(k, curve_keys_data, curve_radius_data, current_bounds);
} }
} }
return current_bounds; return current_bounds;
@ -431,15 +336,13 @@ void Hair::compute_bounds()
return combined_bounds; return combined_bounds;
})); }));
Attribute *curve_attr = attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); const Attribute *attr_P = attributes.find(ATTR_STD_POSITION);
if (use_motion_blur && curve_attr) { if (use_motion_blur && attr_P->has_motion()) {
const size_t steps_size = positions_size * (motion_steps - 1); for (int attr_step = 1; attr_step < attr_P->num_motion_steps(); attr_step++) {
// Attribute data is stored as a float4 and is not const packed_float3 *key_step = attr_P->data<packed_float3>(attr_step);
// interchangeable with float3 for (size_t i = 0; i < curve_keys_size; i++) {
const float4 *key_steps = curve_attr->data<float4>(); bnds.grow(key_step[i]);
}
for (size_t i = 0; i < steps_size; i++) {
bnds.grow(make_float3(key_steps[i]));
} }
} }
@ -447,18 +350,16 @@ void Hair::compute_bounds()
bnds = BoundBox::empty; bnds = BoundBox::empty;
/* skip nan or inf coordinates */ /* skip nan or inf coordinates */
for (size_t i = 0; i < positions_size; i++) { for (size_t i = 0; i < curve_keys_size; i++) {
bnds.grow_safe(positions[i], radius[i]); bnds.grow_safe(curve_keys_data[i], curve_radius_data[i]);
} }
if (use_motion_blur && curve_attr) { if (use_motion_blur && attr_P->has_motion()) {
const size_t steps_size = positions_size * (motion_steps - 1); for (int attr_step = 1; attr_step < attr_P->num_motion_steps(); attr_step++) {
// Attribute data is stored as a float4 which is not const packed_float3 *key_step = attr_P->data<packed_float3>(attr_step);
// interchangeable with float4 for (size_t i = 0; i < curve_keys_size; i++) {
const float4 *key_steps = curve_attr->data<float4>(); bnds.grow_safe(key_step[i]);
}
for (size_t i = 0; i < steps_size; i++) {
bnds.grow_safe(make_float3(key_steps[i]));
} }
} }
} }
@ -481,56 +382,44 @@ void Hair::apply_transform(const Transform &tfm, const bool apply_to_motion)
const float scalar = powf(fabsf(dot(cross(c0, c1), c2)), 1.0f / 3.0f); const float scalar = powf(fabsf(dot(cross(c0, c1), c2)), 1.0f / 3.0f);
/* apply transform to curve keys */ /* apply transform to curve keys */
const size_t num_keys_local = num_keys(); packed_float3 *keys = get_position_for_write();
packed_float3 *positions_data = get_position_for_write(); float *radius = get_radius_for_write();
float *radius_data = get_radius_for_write(); const size_t numkeys = num_keys();
for (size_t i = 0; i < numkeys; i++) {
for (size_t i = 0; i < num_keys_local; i++) { const float3 co = transform_point(&tfm, keys[i]);
const float3 co = transform_point(&tfm, positions_data[i]); const float r = radius[i] * scalar;
const float radius = radius_data[i] * scalar;
/* scale for curve radius is only correct for uniform scale */ /* scale for curve radius is only correct for uniform scale */
positions_data[i] = co; keys[i] = co;
radius_data[i] = radius; radius[i] = r;
} }
tag_position_modified();
tag_radius_modified();
if (apply_to_motion) { if (apply_to_motion) {
Attribute *curve_attr = attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); Attribute *attr_P = attributes.find(ATTR_STD_POSITION);
Attribute *attr_R = attributes.find(ATTR_STD_RADIUS);
if (curve_attr) { if (attr_P->has_motion()) {
/* apply transform to motion curve keys */ const bool has_motion_radius = attr_R->has_motion();
const size_t steps_size = num_keys_local * (motion_steps - 1); const size_t nk = num_keys();
float4 *key_steps = curve_attr->data_for_write<float4>(); for (int step = 1; step <= int(attr_P->motion.size()); step++) {
packed_float3 *motion_P = attr_P->data_for_write<packed_float3>(step);
for (size_t i = 0; i < steps_size; i++) { float *motion_R = has_motion_radius ? attr_R->data_for_write<float>(step) : nullptr;
const float3 co = transform_point(&tfm, make_float3(key_steps[i])); for (size_t i = 0; i < nk; i++) {
const float radius = key_steps[i].w * scalar; motion_P[i] = transform_point(&tfm, motion_P[i]);
if (motion_R) {
/* scale for curve radius is only correct for uniform scale */ motion_R[i] *= scalar;
key_steps[i] = make_float4(co); }
key_steps[i].w = radius; }
} }
} }
} }
} }
void Hair::pack_curves(Scene *scene, void Hair::pack_curves(Scene *scene, KernelCurve *curves, KernelCurveSegment *curve_segments)
float4 *curve_key_co,
KernelCurve *curves,
KernelCurveSegment *curve_segments)
{ {
const size_t positions_size = num_keys();
/* pack curve keys */
if (positions_size) {
const packed_float3 *keys_ptr = get_position();
const float *radius_ptr = get_radius();
for (size_t i = 0; i < positions_size; i++) {
curve_key_co[i] = make_float4(float3(keys_ptr[i]), radius_ptr[i]);
}
}
/* pack curve segments */ /* pack curve segments */
const PrimitiveType type = primitive_type(); const PrimitiveType type = primitive_type();
@ -546,7 +435,7 @@ void Hair::pack_curves(Scene *scene,
shader_id = scene->shader_manager->get_shader_id(shader, false); shader_id = scene->shader_manager->get_shader_id(shader, false);
curves[i].shader_id = shader_id; curves[i].shader_id = shader_id;
curves[i].first_key = curve_key_offset + curve.first_key; curves[i].first_key = curve.first_key;
curves[i].num_keys = curve.num_keys; curves[i].num_keys = curve.num_keys;
curves[i].type = type; curves[i].type = type;

55
intern/cycles/scene/hair.h
View file

@ -25,40 +25,26 @@ class Hair : public Geometry {
} }
void bounds_grow(const int k, void bounds_grow(const int k,
const float3 *positions, const packed_float3 *curve_keys,
const float *curve_radius,
BoundBox &bounds) const;
void bounds_grow(const int k,
const packed_float3 *positions,
const float *curve_radius, const float *curve_radius,
BoundBox &bounds) const; BoundBox &bounds) const;
void bounds_grow(const int k, const float4 *keys, BoundBox &bounds) const; void bounds_grow(const int k, const float4 *keys, BoundBox &bounds) const;
void bounds_grow(const float4 keys[4], BoundBox &bounds) const; void bounds_grow(const float4 keys[4], BoundBox &bounds) const;
void bounds_grow(const float3 keys[4], BoundBox &bounds) const;
void bounds_grow(const int k, void bounds_grow(const int k,
const float3 *positions, const packed_float3 *curve_keys,
const float *curve_radius,
const Transform &aligned_space,
BoundBox &bounds) const;
void bounds_grow(const int k,
const packed_float3 *positions,
const float *curve_radius, const float *curve_radius,
const Transform &aligned_space, const Transform &aligned_space,
BoundBox &bounds) const; BoundBox &bounds) const;
void motion_keys(const packed_float3 *positions, void motion_keys(const Attribute *attr_P,
const float *curve_radius, const Attribute *attr_R,
const float4 *key_steps,
const size_t num_positions,
const size_t num_steps, const size_t num_steps,
const float time, const float time,
size_t k0, size_t k0,
size_t k1, size_t k1,
float4 r_keys[2]) const; float4 r_keys[2]) const;
void cardinal_motion_keys(const packed_float3 *positions, void cardinal_motion_keys(const Attribute *attr_P,
const float *curve_radius, const Attribute *attr_R,
const float4 *key_steps,
const size_t num_positions,
const size_t num_steps, const size_t num_steps,
const float time, const float time,
size_t k0, size_t k0,
@ -67,20 +53,14 @@ class Hair : public Geometry {
size_t k3, size_t k3,
float4 r_keys[4]) const; float4 r_keys[4]) const;
void keys_for_step(const packed_float3 *positions, void keys_for_step(const Attribute *attr_P,
const float *curve_radius, const Attribute *attr_R,
const float4 *key_steps,
const size_t num_positions,
const size_t num_steps,
const size_t step, const size_t step,
size_t k0, size_t k0,
size_t k1, size_t k1,
float4 r_keys[2]) const; float4 r_keys[2]) const;
void cardinal_keys_for_step(const packed_float3 *positions, void cardinal_keys_for_step(const Attribute *attr_P,
const float *curve_radius, const Attribute *attr_R,
const float4 *key_steps,
const size_t num_positions,
const size_t num_steps,
const size_t step, const size_t step,
size_t k0, size_t k0,
size_t k1, size_t k1,
@ -93,7 +73,6 @@ class Hair : public Geometry {
NODE_SOCKET_API_ARRAY(array<int>, curve_shader) NODE_SOCKET_API_ARRAY(array<int>, curve_shader)
/* BVH */ /* BVH */
size_t curve_key_offset;
size_t curve_segment_offset; size_t curve_segment_offset;
CurveShapeType curve_shape; CurveShapeType curve_shape;
@ -115,14 +94,17 @@ class Hair : public Geometry {
Curve get_curve(const size_t i) const Curve get_curve(const size_t i) const
{ {
const int first = curve_first_key[i]; const int first = curve_first_key[i];
const int next_first = (i + 1 < curve_first_key.size()) ? curve_first_key[i + 1] : const int next_first = (i + 1 < curve_first_key.size()) ? curve_first_key[i + 1] : num_keys();
int(num_keys());
Curve curve = {first, next_first - first}; Curve curve = {first, next_first - first};
return curve; return curve;
} }
size_t num_keys() const; size_t num_keys() const
{
const Attribute *attr = attributes.find(ATTR_STD_POSITION);
return attr ? attr->size : 0;
}
size_t num_curves() const size_t num_curves() const
{ {
@ -143,10 +125,7 @@ class Hair : public Geometry {
void get_uv_tiles(ustring map, unordered_set<int> &tiles) override; void get_uv_tiles(ustring map, unordered_set<int> &tiles) override;
/* BVH */ /* BVH */
void pack_curves(Scene *scene, void pack_curves(Scene *scene, KernelCurve *curve, KernelCurveSegment *curve_segments);
float4 *curve_key_co,
KernelCurve *curve,
KernelCurveSegment *curve_segments);
PrimitiveType primitive_type() const override; PrimitiveType primitive_type() const override;

17
intern/cycles/scene/object.cpp
View file

@ -1122,6 +1122,23 @@ void ObjectManager::device_update_geom_offsets(Device * /*unused*/,
assert(normal_offset != ATTR_STD_NOT_FOUND || assert(normal_offset != ATTR_STD_NOT_FOUND ||
static_cast<Mesh *>(geom)->num_triangles() == 0); static_cast<Mesh *>(geom)->num_triangles() == 0);
} }
else if (geom->is_hair()) {
position_offset = find_attribute(dscene->attributes_map.data(),
attr_map_offset,
PRIMITIVE_CURVE_THICK,
ATTR_STD_POSITION)
.offset;
assert(position_offset != ATTR_STD_NOT_FOUND || static_cast<Hair *>(geom)->num_keys() == 0);
}
else if (geom->is_pointcloud()) {
position_offset = find_attribute(dscene->attributes_map.data(),
attr_map_offset,
PRIMITIVE_POINT,
ATTR_STD_POSITION)
.offset;
assert(position_offset != ATTR_STD_NOT_FOUND ||
static_cast<PointCloud *>(geom)->num_points() == 0);
}
if (kobject.position_offset != position_offset) { if (kobject.position_offset != position_offset) {
kobject.position_offset = position_offset; kobject.position_offset = position_offset;
update = true; update = true;

162
intern/cycles/scene/pointcloud.cpp
View file

@ -32,10 +32,9 @@ void PointCloud::Point::bounds_grow(const float4 &point, BoundBox &bounds) const
bounds.grow(make_float3(point), point.w); bounds.grow(make_float3(point), point.w);
} }
float4 PointCloud::Point::motion_key(const packed_float3 *points, float4 PointCloud::Point::motion_key(const float *radius,
const float *radius, const Attribute *attr_P,
const float4 *point_steps, const Attribute *attr_R,
const size_t num_points,
const size_t num_steps, const size_t num_steps,
const float time, const float time,
size_t p) const size_t p) const
@ -46,33 +45,21 @@ float4 PointCloud::Point::motion_key(const packed_float3 *points,
const size_t step = min((size_t)(time * max_step), max_step - 1); const size_t step = min((size_t)(time * max_step), max_step - 1);
const float t = time * max_step - step; const float t = time * max_step - step;
/* Fetch vertex coordinates. */ /* Fetch vertex coordinates. */
const float4 curr_key = point_for_step( const float4 curr_key = point_for_step(radius, attr_P, attr_R, step, p);
points, radius, point_steps, num_points, num_steps, step, p); const float4 next_key = point_for_step(radius, attr_P, attr_R, step + 1, p);
const float4 next_key = point_for_step(
points, radius, point_steps, num_points, num_steps, step + 1, p);
/* Interpolate between steps. */ /* Interpolate between steps. */
return (1.0f - t) * curr_key + t * next_key; return (1.0f - t) * curr_key + t * next_key;
} }
float4 PointCloud::Point::point_for_step(const packed_float3 *points, float4 PointCloud::Point::point_for_step(const float *radius,
const float *radius, const Attribute *attr_P,
const float4 *point_steps, const Attribute *attr_R,
const size_t num_points, const size_t step,
const size_t num_steps, const size_t p) const
size_t step,
size_t p) const
{ {
const size_t center_step = ((num_steps - 1) / 2); const int num_steps = attr_P->num_motion_steps();
if (step == center_step) { const float r = attr_R ? attr_R->data_at_time_step<float>(step, num_steps)[p] : radius[p];
/* Center step: regular key location. */ return make_float4(float3(attr_P->data_at_time_step<packed_float3>(step, num_steps)[p]), r);
return make_float4(float3(points[p]), radius[p]);
}
/* Center step is not stored in this array. */
if (step > center_step) {
step--;
}
const size_t offset = step * num_points;
return point_steps[offset + p];
} }
/* PointCloud */ /* PointCloud */
@ -94,12 +81,6 @@ PointCloud::PointCloud() : Geometry(get_node_type(), Geometry::POINTCLOUD)
PointCloud::~PointCloud() = default; PointCloud::~PointCloud() = default;
size_t PointCloud::num_points() const
{
const Attribute *attr = attributes.find(ATTR_STD_POSITION);
return attr ? attr->size : 0;
}
void PointCloud::add_builtin_attributes() void PointCloud::add_builtin_attributes()
{ {
attributes.add(ATTR_STD_POSITION); attributes.add(ATTR_STD_POSITION);
@ -108,8 +89,10 @@ void PointCloud::add_builtin_attributes()
void PointCloud::resize(const int numpoints) void PointCloud::resize(const int numpoints)
{ {
attributes.add(ATTR_STD_POSITION)->resize(numpoints); Attribute *attr_P = attributes.add(ATTR_STD_POSITION);
attributes.add(ATTR_STD_RADIUS)->resize(numpoints); attr_P->resize(numpoints);
Attribute *attr_R = attributes.add(ATTR_STD_RADIUS);
attr_R->resize(numpoints);
shader.resize(numpoints); shader.resize(numpoints);
attributes.resize(); attributes.resize();
@ -138,18 +121,17 @@ void PointCloud::clear(const bool preserve_shaders)
void PointCloud::copy_center_to_motion_step(const int motion_step) void PointCloud::copy_center_to_motion_step(const int motion_step)
{ {
Attribute *attr_mP = attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); const int attr_step = motion_step + 1;
if (attr_mP) { const size_t numpoints = num_points();
const packed_float3 *points_data = get_position();
const size_t numpoints = num_points();
const float *radius_data = get_radius();
float4 *attrib_P = attr_mP->data_for_write<float4>() + motion_step * numpoints; Attribute *attr_P = attributes.find(ATTR_STD_POSITION);
for (size_t i = 0; i < numpoints; i++) { if (attr_P->has_motion()) {
const float3 P = float3(points_data[i]); std::copy_n(get_position(), numpoints, attr_P->data_for_write<packed_float3>(attr_step));
const float r = radius_data[i]; }
attrib_P[i] = make_float4(P, r);
} Attribute *attr_R = attributes.find(ATTR_STD_RADIUS);
if (attr_R->has_motion()) {
std::copy_n(get_radius(), numpoints, attr_R->data_for_write<float>(attr_step));
} }
} }
@ -173,21 +155,23 @@ void PointCloud::compute_bounds()
{ {
BoundBox bnds = BoundBox::empty; BoundBox bnds = BoundBox::empty;
const size_t numpoints = num_points(); const size_t numpoints = num_points();
const packed_float3 *points = get_position(); const packed_float3 *points_data = get_position();
const float *radius = get_radius(); const float *radius_data = get_radius();
if (numpoints > 0) { if (numpoints > 0) {
for (size_t i = 0; i < numpoints; i++) { for (size_t i = 0; i < numpoints; i++) {
bnds.grow(points[i], radius[i]); bnds.grow(points_data[i], radius_data[i]);
} }
Attribute *attr = attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); const Attribute *attr_P = attributes.find(ATTR_STD_POSITION);
if (use_motion_blur && attr) { const Attribute *attr_R = attributes.find(ATTR_STD_RADIUS);
const size_t steps_size = numpoints * (motion_steps - 1); if (use_motion_blur && attr_P->has_motion()) {
const float4 *point_steps = attr->data<float4>(); for (int attr_step = 1; attr_step < attr_P->num_motion_steps(); attr_step++) {
const packed_float3 *motion_P = attr_P->data<packed_float3>(attr_step);
for (size_t i = 0; i < steps_size; i++) { const float *motion_R = attr_R->data<float>(attr_step);
bnds.grow(make_float3(point_steps[i]), point_steps[i].w); for (size_t i = 0; i < numpoints; i++) {
bnds.grow(motion_P[i], motion_R[i]);
}
} }
} }
@ -196,15 +180,16 @@ void PointCloud::compute_bounds()
/* skip nan or inf coordinates */ /* skip nan or inf coordinates */
for (size_t i = 0; i < numpoints; i++) { for (size_t i = 0; i < numpoints; i++) {
bnds.grow_safe(points[i], radius[i]); bnds.grow_safe(points_data[i], radius_data[i]);
} }
if (use_motion_blur && attr) { if (use_motion_blur && attr_P->has_motion()) {
const size_t steps_size = numpoints * (motion_steps - 1); for (int attr_step = 1; attr_step < attr_P->num_motion_steps(); attr_step++) {
const float4 *point_steps = attr->data<float4>(); const packed_float3 *motion_P = attr_P->data<packed_float3>(attr_step);
const float *motion_R = attr_R->data<float>(attr_step);
for (size_t i = 0; i < steps_size; i++) { for (size_t i = 0; i < numpoints; i++) {
bnds.grow_safe(make_float3(point_steps[i]), point_steps[i].w); bnds.grow_safe(motion_P[i], motion_R[i]);
}
} }
} }
} }
@ -226,52 +211,45 @@ void PointCloud::apply_transform(const Transform &tfm, const bool apply_to_motio
const float3 c2 = transform_get_column(&tfm, 2); const float3 c2 = transform_get_column(&tfm, 2);
const float scalar = powf(fabsf(dot(cross(c0, c1), c2)), 1.0f / 3.0f); const float scalar = powf(fabsf(dot(cross(c0, c1), c2)), 1.0f / 3.0f);
/* apply transform to curve keys */ /* apply transform to points */
packed_float3 *points_data = get_position_for_write();
float *radius_data = get_radius_for_write();
const size_t numpoints = num_points(); const size_t numpoints = num_points();
packed_float3 *points = get_position_for_write();
float *radius = get_radius_for_write();
for (size_t i = 0; i < numpoints; i++) { for (size_t i = 0; i < numpoints; i++) {
const float3 co = transform_point(&tfm, points[i]); const float3 co = transform_point(&tfm, points_data[i]);
const float r = radius[i] * scalar; const float r = radius_data[i] * scalar;
/* scale for curve radius is only correct for uniform scale /* scale for radius is only correct for uniform scale */
*/ points_data[i] = co;
points[i] = co; radius_data[i] = r;
radius[i] = r;
} }
if (apply_to_motion) { if (apply_to_motion) {
Attribute *attr = attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); Attribute *attr_P = attributes.find(ATTR_STD_POSITION);
Attribute *attr_R = attributes.find(ATTR_STD_RADIUS);
if (attr) { if (attr_P->has_motion()) {
/* apply transform to motion curve keys */ const bool has_motion_radius = attr_R->has_motion();
const size_t steps_size = numpoints * (motion_steps - 1); for (int step = 1; step <= int(attr_P->motion.size()); step++) {
float4 *point_steps = attr->data_for_write<float4>(); packed_float3 *motion_P = attr_P->data_for_write<packed_float3>(step);
float *motion_R = has_motion_radius ? attr_R->data_for_write<float>(step) : nullptr;
for (size_t i = 0; i < steps_size; i++) { for (size_t i = 0; i < numpoints; i++) {
const float3 co = transform_point(&tfm, make_float3(point_steps[i])); motion_P[i] = transform_point(&tfm, motion_P[i]);
const float radius = point_steps[i].w * scalar; if (motion_R) {
/* scale for curve radius is only correct for uniform scale */
/* scale for curve radius is only correct for uniform motion_R[i] *= scalar;
* scale */ }
point_steps[i] = make_float4(co); }
point_steps[i].w = radius;
} }
} }
} }
} }
void PointCloud::pack(Scene *scene, float4 *packed_points, uint *packed_shader) void PointCloud::pack(Scene *scene, uint *packed_shader)
{ {
const size_t numpoints = num_points(); const size_t numpoints = num_points();
const packed_float3 *points_data = get_position();
const float *radius_data = get_radius();
int *shader_data = shader.data(); int *shader_data = shader.data();
for (size_t i = 0; i < numpoints; i++) {
packed_points[i] = make_float4(float3(points_data[i]), radius_data[i]);
}
uint shader_id = 0; uint shader_id = 0;
uint last_shader = -1; uint last_shader = -1;
for (size_t i = 0; i < numpoints; i++) { for (size_t i = 0; i < numpoints; i++) {

23
intern/cycles/scene/pointcloud.h
View file

@ -23,18 +23,15 @@ class PointCloud : public Geometry {
BoundBox &bounds) const; BoundBox &bounds) const;
void bounds_grow(const float4 &point, BoundBox &bounds) const; void bounds_grow(const float4 &point, BoundBox &bounds) const;
float4 motion_key(const packed_float3 *points, float4 motion_key(const float *radius,
const float *radius, const Attribute *attr_P,
const float4 *point_steps, const Attribute *attr_R,
const size_t num_points,
const size_t num_steps, const size_t num_steps,
const float time, const float time,
size_t p) const; size_t p) const;
float4 point_for_step(const packed_float3 *points, float4 point_for_step(const float *radius,
const float *radius, const Attribute *attr_P,
const float4 *point_steps, const Attribute *attr_R,
const size_t num_points,
const size_t num_steps,
const size_t step, const size_t step,
size_t p) const; size_t p) const;
}; };
@ -63,7 +60,11 @@ class PointCloud : public Geometry {
return point; return point;
} }
size_t num_points() const; size_t num_points() const
{
const Attribute *attr = attributes.find(ATTR_STD_POSITION);
return attr ? attr->size : 0;
}
size_t num_attributes() const size_t num_attributes() const
{ {
@ -76,7 +77,7 @@ class PointCloud : public Geometry {
PrimitiveType primitive_type() const override; PrimitiveType primitive_type() const override;
/* BVH */ /* BVH */
void pack(Scene *scene, float4 *packed_points, uint *packed_shader); void pack(Scene *scene, uint *packed_shader);
private: private:
void add_builtin_attributes(); void add_builtin_attributes();