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,
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();
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) {
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) {
/* Motion, fill up previous steps that we might have skipped because
* they had no motion, but we need them anyway now. */
for (int step = 0; step < motion_step; step++) {
float4 *mP = attr_mP->data_for_write<float4>() + step * num_keys;
for (int step = 1; step <= motion_step; step++) {
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++) {
mP[key] = make_float4(hair->get_position()[key]);
mP[key].w = hair->get_radius()[key];
}
float *mR = attr_R->data_for_write<float>(step);
std::copy_n(hair->get_radius(), num_keys, mR);
}
}
}
static void ExportCurveSegmentsMotion(Hair *hair, ParticleCurveData *CData, const int motion_step)
{
/* find attribute */
Attribute *attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
/* Set motion steps on position and radius attributes. */
Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
Attribute *attr_R = hair->attributes.find(ATTR_STD_RADIUS);
bool new_attribute = false;
/* add new attribute if it doesn't exist already */
if (!attr_mP) {
attr_mP = hair->attributes.add(ATTR_STD_MOTION_VERTEX_POSITION);
if (!attr_P->has_motion()) {
attr_P->add_motion(hair);
attr_R->add_motion(hair);
new_attribute = true;
}
/* export motion vectors for curve keys */
const size_t numkeys = hair->num_keys();
float4 *mP = attr_mP->data_for_write<float4>() + motion_step * numkeys;
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;
int i = 0;
int num_curves = 0;
const int num_keys = hair->num_keys();
for (int sys = 0; sys < CData->psys_firstcurve.size(); 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. */
const int curve_key_end = (num_curves + 1 < (int)hair->get_curve_first_key().size() ?
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 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++)
{
if (i < hair->num_keys()) {
mP[i] = CurveSegmentMotionCV(CData, sys, curve, curvekey);
if (i < num_keys) {
const float4 cv = CurveSegmentMotionCV(CData, sys, curve, curvekey);
mP[i] = make_float3(cv);
mR[i] = cv.w;
if (!have_motion) {
/* unlike mesh coordinates, these tend to be slightly different
* between frames due to particle transforms into/out of object
* space, so we use an epsilon to detect actual changes */
float4 curve_key = make_float4(hair->get_position()[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;
}
}
@ -547,7 +553,9 @@ static void ExportCurveSegmentsMotion(Hair *hair, ParticleCurveData *CData, cons
0.0f;
for (int step_index = 0; step_index < num_center_curve_keys; ++step_index) {
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++;
}
have_motion = true;
@ -717,19 +725,17 @@ static void attr_create_motion_from_velocity(Hair *hair,
/* Override motion steps to fixed number. */
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();
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. */
const float motion_times[2] = {-1.0f, 1.0f};
for (int step = 0; step < 2; step++) {
const float relative_time = motion_times[step] * 0.5f * motion_scale;
packed_float3 *mP = attr_mP->data_for_write<packed_float3>() + step * num_curve_keys;
for (int step = 1; step <= 2; step++) {
const float relative_time = motion_times[step - 1] * 0.5f * motion_scale;
packed_float3 *mP = attr_P->data_for_write<packed_float3>(step);
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;
@ -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);
}
@ -972,19 +981,23 @@ static void export_hair_curves_motion(Hair *hair,
const blender::bke::CurvesGeometry &b_curves,
const int motion_step)
{
/* Find or add attribute. */
Attribute *attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
/* Set motion steps on position and radius attributes. */
Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
Attribute *attr_R = hair->attributes.find(ATTR_STD_RADIUS);
bool new_attribute = false;
if (!attr_mP) {
attr_mP = hair->attributes.add(ATTR_STD_MOTION_VERTEX_POSITION);
if (!attr_P->has_motion()) {
attr_P->add_motion(hair);
attr_R->add_motion(hair);
new_attribute = true;
}
/* Export motion keys. */
const size_t num_keys = hair->num_keys();
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;
int num_motion_keys = 0;
int curve_index = 0;
@ -1009,15 +1022,17 @@ static void export_hair_curves_motion(Hair *hair,
const int point = points[i];
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++;
if (!have_motion) {
/* TODO: use epsilon for comparison? Was needed for particles due to
* transform, but ideally should not happen anymore. */
float4 curve_key = make_float4(hair->get_position()[i]);
curve_key.w = hair->get_radius()[i];
have_motion = !(mP[i] == curve_key);
float4 curve_key = make_float4(hair->get_position()[point]);
curve_key.w = hair->get_radius()[point];
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;
for (int i = 0; i < curve.num_keys; i++) {
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);
mP[num_motion_keys] = make_float3(cv);
mR[num_motion_keys] = cv.w;
num_motion_keys++;
}
have_motion = true;
@ -1100,12 +1117,13 @@ void BlenderSync::sync_hair(BObjectInfo &b_ob_info, Hair *hair)
{
new_hair.set_motion_steps(2);
Attribute *attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
Attribute *new_attr_mP = new_hair.attributes.add(ATTR_STD_MOTION_VERTEX_POSITION);
if (attr_mP) {
new_attr_mP->set_data_from(std::move(*attr_mP));
Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
Attribute *new_attr_P = new_hair.attributes.find(ATTR_STD_POSITION);
if (attr_P->has_motion()) {
new_attr_P->take_motion_from(*attr_P);
}
else {
new_attr_P->add_motion(&new_hair);
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. */
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 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. */
const float motion_times[2] = {-1.0f, 1.0f};
for (int step = 0; step < 2; step++) {
const float relative_time = motion_times[step] * 0.5f * motion_scale;
float4 *mP = attr_mP->data_for_write<float4>() + step * num_points;
for (int step = 1; step <= 2; step++) {
const float relative_time = motion_times[step - 1] * 0.5f * motion_scale;
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++) {
const float3 Pi = float3(P[i]) +
make_float3(b_attribute[i][0], b_attribute[i][1], b_attribute[i][2]) *
relative_time;
mP[i] = make_float4(Pi, radius[i]);
mP[i] = float3(P[i]) +
make_float3(b_attribute[i][0], b_attribute[i][1], b_attribute[i][2]) * relative_time;
mR[i] = radius[i];
}
}
}
@ -141,6 +140,8 @@ static void export_pointcloud(Scene *scene,
else {
std::fill(radius, radius + b_positions.size(), 0.01f);
}
pointcloud->tag_position_modified();
pointcloud->tag_radius_modified();
int *shader = pointcloud->get_shader().data();
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 int motion_step)
{
/* Find or add attribute. */
Attribute *attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
/* 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);
bool new_attribute = false;
if (!attr_mP) {
attr_mP = pointcloud->attributes.add(ATTR_STD_MOTION_VERTEX_POSITION);
if (!attr_P->has_motion()) {
attr_P->add_motion(pointcloud);
attr_R->add_motion(pointcloud);
new_attribute = true;
}
const int num_points = pointcloud->num_points();
/* Point cloud attributes are stored as float4 with the radius in the w element.
* This is explicit now as float3 is no longer interchangeable with float4 as it
* is packed now. */
float4 *mP = attr_mP->data_for_write<float4>() + motion_step * num_points;
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;
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::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++) {
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];
mP[i] = make_float4(P, radius);
have_motion = have_motion || (P != pointcloud_positions[i]);
mP[i] = P;
mR[i] = radius;
have_motion = have_motion || (P != pointcloud_points[i]);
}
/* In case of new attribute, we verify if there really was any motion. */
if (new_attribute) {
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) {
/* 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);
Attribute *attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
Attribute *new_attr_mP = new_pointcloud.attributes.add(ATTR_STD_MOTION_VERTEX_POSITION);
if (attr_mP) {
new_attr_mP->set_data_from(std::move(*attr_mP));
Attribute *attr_P = pointcloud->attributes.find(ATTR_STD_POSITION);
Attribute *new_attr_P = new_pointcloud.attributes.find(ATTR_STD_POSITION);
if (attr_P->has_motion()) {
new_attr_P->take_motion_from(*attr_P);
}
else {
new_attr_P->add_motion(&new_pointcloud);
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,
const int object_index)
{
const Attribute *curve_attr_mP = nullptr;
if (hair->has_motion_blur()) {
curve_attr_mP = 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);
const bool has_motion = attr_P->has_motion();
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 float *curve_radius = hair->get_radius();
for (int k = 0; k < curve.num_keys - 1; k++) {
if (curve_attr_mP == nullptr) {
if (!has_motion) {
/* Really simple logic for static hair. */
BoundBox bounds = BoundBox::empty;
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) {
/* 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
* rendering.
*/
/* TODO(sergey): Support motion steps for spatially split BVH. */
BoundBox bounds = BoundBox::empty;
curve.bounds_grow(k, hair->get_position(), curve_radius, bounds);
const size_t num_keys = hair->num_keys();
const size_t num_steps = hair->get_motion_steps();
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);
for (int attr_step = 0; attr_step < attr_P->num_motion_steps(); attr_step++) {
curve.bounds_grow(
k, attr_P->data<packed_float3>(attr_step), attr_R->data<float>(attr_step), bounds);
}
if (bounds.valid()) {
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 float num_bvh_steps_inv_1 = 1.0f / (num_bvh_steps - 1);
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.
* Will be reused later to avoid duplicated work on
* calculating BVH time step boundbox.
*/
float4 prev_keys[4];
curve.cardinal_motion_keys(curve_keys,
curve_radius,
key_steps,
num_keys,
num_steps,
0.0f,
k - 1,
k,
k + 1,
k + 2,
prev_keys);
curve.cardinal_motion_keys(
attr_P, attr_R, num_steps, 0.0f, k - 1, k, k + 1, k + 2, prev_keys);
BoundBox prev_bounds = BoundBox::empty;
curve.bounds_grow(prev_keys, prev_bounds);
/* Create all primitive time steps, */
for (int bvh_step = 1; bvh_step < num_bvh_steps; ++bvh_step) {
const float curr_time = (float)(bvh_step)*num_bvh_steps_inv_1;
float4 curr_keys[4];
curve.cardinal_motion_keys(curve_keys,
curve_radius,
key_steps,
num_keys,
num_steps,
curr_time,
k - 1,
k,
k + 1,
k + 2,
curr_keys);
curve.cardinal_motion_keys(
attr_P, attr_R, num_steps, curr_time, k - 1, k, k + 1, k + 2, curr_keys);
BoundBox curr_bounds = BoundBox::empty;
curve.bounds_grow(curr_keys, curr_bounds);
BoundBox bounds = prev_bounds;
@ -256,18 +231,16 @@ void BVHBuild::add_reference_points(BoundBox &root,
PointCloud *pointcloud,
const int i)
{
const Attribute *point_attr_mP = nullptr;
if (pointcloud->has_motion_blur()) {
point_attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
}
const Attribute *attr_P = pointcloud->attributes.find(ATTR_STD_POSITION);
const Attribute *attr_R = pointcloud->attributes.find(ATTR_STD_RADIUS);
const bool has_motion = attr_P->has_motion();
const packed_float3 *points_data = pointcloud->get_position();
const float *radius_data = pointcloud->get_radius();
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();
if (point_attr_mP == nullptr) {
if (!has_motion) {
/* Really simple logic for static points. */
for (uint j = 0; j < num_points; 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++) {
const PointCloud::Point point = pointcloud->get_point(j);
BoundBox bounds = BoundBox::empty;
point.bounds_grow(points_data, radius_data, bounds);
for (size_t step = 0; step < num_steps - 1; step++) {
point.bounds_grow(motion_data[step * num_points + j], bounds);
for (int attr_step = 0; attr_step < attr_P->num_motion_steps(); attr_step++) {
const packed_float3 *mP = attr_P->data<packed_float3>(attr_step);
const float *mR = attr_R->data<float>(attr_step);
point.bounds_grow(make_float4(float3(mP[j]), mR[j]), bounds);
}
if (bounds.valid()) {
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 float num_bvh_steps_inv_1 = 1.0f / (num_bvh_steps - 1);
const size_t num_steps = pointcloud->get_motion_steps();
const float4 *point_steps = point_attr_mP->data<float4>();
const Attribute *attr_R_motion = attr_R->has_motion() ? attr_R : nullptr;
for (uint j = 0; j < num_points; 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.
*/
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;
point.bounds_grow(prev_key, prev_bounds);
/* Create all primitive time steps, */
for (int bvh_step = 1; bvh_step < num_bvh_steps; ++bvh_step) {
const float curr_time = (float)(bvh_step)*num_bvh_steps_inv_1;
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;
point.bounds_grow(curr_key, curr_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. */
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) {
const size_t hair_size = hair->num_keys();
const size_t steps = hair->get_motion_steps() - 1;
const packed_float3 *key_steps = attr->data<packed_float3>();
for (size_t i = 0; i < steps; i++) {
curve.bounds_grow(k, key_steps + i * hair_size, hair->get_radius(), bbox);
if (attr_P->has_motion()) {
for (int attr_step = 1; attr_step < attr_P->num_motion_steps(); attr_step++) {
curve.bounds_grow(
k, attr_P->data<packed_float3>(attr_step), attr_R->data<float>(attr_step), bbox);
}
}
}
@ -405,17 +403,13 @@ void BVH2::refit_primitives(const int start, const int end, BoundBox &bbox, uint
/* Motion points. */
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) {
const size_t pointcloud_size = pointcloud->num_points();
const size_t steps = pointcloud->get_motion_steps() - 1;
const float4 *point_steps = attr->data<float4>();
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;
if (attr_P->has_motion()) {
const Attribute *attr_R = pointcloud->attributes.find(ATTR_STD_RADIUS);
for (int attr_step = 1; attr_step < attr_P->num_motion_steps(); attr_step++) {
const float3 P = attr_P->data<packed_float3>(attr_step)[point.index];
const float r = attr_R->data<float>(attr_step)[point.index];
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)
{
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;
if (hair->has_motion_blur()) {
attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
if (attr_mP) {
num_motion_steps = hair->get_motion_steps();
}
if (hair->has_motion_blur() && attr_P->has_motion()) {
num_motion_steps = hair->get_motion_steps();
}
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;
/* 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) {
// 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;
if (update) {
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);
if (rtc_verts) {
const size_t num_curves = hair->num_curves();
if (t == t_mid || attr_mP == nullptr) {
const packed_float3 *verts = hair->get_position();
pack_motion_verts<packed_float3>(
num_curves, hair, verts, curve_radius, rtc_verts, hair->curve_shape);
}
else {
const int t_ = (t > t_mid) ? (t - 1) : t;
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);
}
pack_motion_verts<packed_float3>(
num_curves,
hair,
attr_P->data_at_time_step<packed_float3>(t, num_motion_steps),
attr_R->data_at_time_step<float>(t, num_motion_steps),
rtc_verts,
hair->curve_shape);
}
if (update) {
@ -566,25 +554,17 @@ void BVHEmbree::set_point_vertex_buffer(RTCGeometry geom_id,
const PointCloud *pointcloud,
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;
if (pointcloud->has_motion_blur()) {
attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
if (attr_mP) {
num_motion_steps = pointcloud->get_motion_steps();
}
if (pointcloud->has_motion_blur() && attr_P->has_motion()) {
num_motion_steps = pointcloud->get_motion_steps();
}
const size_t num_points = pointcloud->num_points();
/* Copy the point data to Embree */
const int t_mid = (num_motion_steps - 1) / 2;
const float *radius = pointcloud->get_radius();
/* Copy the point data to Embree. */
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;
if (update) {
@ -612,18 +592,10 @@ void BVHEmbree::set_point_vertex_buffer(RTCGeometry geom_id,
assert(rtc_verts);
if (rtc_verts) {
if (t == t_mid || attr_mP == nullptr) {
/* Pack the motion points into a float4 as [x y z radius]. */
const packed_float3 *verts = pointcloud->get_position();
for (size_t j = 0; j < num_points; ++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);
const packed_float3 *verts = attr_P->data_at_time_step<packed_float3>(t, num_motion_steps);
const float *radius = attr_R->data_at_time_step<float>(t, num_motion_steps);
for (size_t j = 0; j < num_points; ++j) {
rtc_verts[j] = make_float4(float3(verts[j]), radius[j]);
}
}
@ -637,11 +609,10 @@ void BVHEmbree::add_points(const Object *ob, const PointCloud *pointcloud, const
{
const size_t prim_offset = pointcloud->prim_offset;
const Attribute *attr_mP = nullptr;
size_t num_motion_steps = 1;
if (pointcloud->has_motion_blur()) {
attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
if (attr_mP) {
const Attribute *attr_P = pointcloud->attributes.find(ATTR_STD_POSITION);
if (attr_P->has_motion()) {
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 Attribute *attr_mP = nullptr;
const Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
size_t num_motion_steps = 1;
if (hair->has_motion_blur()) {
attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
if (attr_mP) {
num_motion_steps = hair->get_motion_steps();
}
if (hair->has_motion_blur() && attr_P->has_motion()) {
num_motion_steps = hair->get_motion_steps();
}
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 size_t num_curves = hair->num_curves();
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()) {
curve_attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
}
if (curve_attr_mP == nullptr || bvh->params.num_motion_curve_steps == 0) {
if (!has_motion || bvh->params.num_motion_curve_steps == 0) {
bvh->custom_prim_info.resize(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();
int first_key = curve.first_key;
for (int k = 0; k < curve.num_keys - 1; k++) {
if (curve_attr_mP == nullptr) {
if (!has_motion) {
float3 current_keys[4];
current_keys[0] = curve_keys[max(first_key + k - 1, first_key)];
current_keys[1] = curve_keys[first_key + k];
@ -592,14 +590,12 @@ hiprtGeometryBuildInput HIPRTDevice::prepare_curve_blas(BVHHIPRT *bvh, Hair *hai
}
else {
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) {
BoundBox bounds = BoundBox::empty;
curve.bounds_grow(k, hair->get_position(), curve_radius, bounds);
for (size_t step = 0; step < num_steps - 1; step++) {
curve.bounds_grow(k, key_steps + step * num_keys, bounds);
for (int attr_step = 0; attr_step < attr_P->num_motion_steps(); attr_step++) {
curve.bounds_grow(
k, attr_P->data<packed_float3>(attr_step), attr_R->data<float>(attr_step), bounds);
}
if (bounds.valid()) {
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);
float4 prev_keys[4];
curve.cardinal_motion_keys(curve_keys,
curve_radius,
key_steps,
num_keys,
num_steps,
0.0f,
k - 1,
k,
k + 1,
k + 2,
prev_keys);
curve.cardinal_motion_keys(
attr_P, attr_R, num_steps, 0.0f, k - 1, k, k + 1, k + 2, prev_keys);
BoundBox prev_bounds = BoundBox::empty;
curve.bounds_grow(prev_keys, prev_bounds);
for (int bvh_step = 1; bvh_step < num_bvh_steps; ++bvh_step) {
const float curr_time = (float)(bvh_step)*num_bvh_steps_inv_1;
float4 curr_keys[4];
curve.cardinal_motion_keys(curve_keys,
curve_radius,
key_steps,
num_keys,
num_steps,
curr_time,
k - 1,
k,
k + 1,
k + 2,
curr_keys);
curve.cardinal_motion_keys(
attr_P, attr_R, num_steps, curr_time, k - 1, k, k + 1, k + 2, curr_keys);
BoundBox curr_bounds = BoundBox::empty;
curve.bounds_grow(curr_keys, curr_bounds);
BoundBox bounds = prev_bounds;
@ -688,21 +666,26 @@ hiprtGeometryBuildInput HIPRTDevice::prepare_point_blas(BVHHIPRT *bvh, PointClou
{
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()) {
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 float *radius_data = pointcloud->get_radius();
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 bool has_motion_radius = attr_R && attr_R->has_motion();
int num_bounds = 0;
float sum_area = 0.0f;
if (point_attr_mP == nullptr) {
if (attr_P == nullptr) {
bvh->custom_prim_info.resize(num_points);
bvh->custom_primitive_bound.alloc(num_points);
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++) {
const PointCloud::Point point = pointcloud->get_point(j);
BoundBox bounds = BoundBox::empty;
point.bounds_grow(points_data, radius_data, bounds);
for (size_t step = 0; step < num_steps - 1; step++) {
point.bounds_grow(motion_data[step * num_points + j], bounds);
for (int attr_step = 0; attr_step < attr_P->num_motion_steps(); attr_step++) {
point.bounds_grow(
attr_P->data<packed_float3>(attr_step), attr_R->data<float>(attr_step), bounds);
}
if (bounds.valid()) {
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->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++) {
const PointCloud::Point point = pointcloud->get_point(j);
const size_t num_steps = pointcloud->get_motion_steps();
const float4 *point_steps = point_attr_mP->data<float4>();
float4 prev_key = point.motion_key(
points_data, radius_data, point_steps, num_points, num_steps, 0.0f, j);
float4 prev_key = point.motion_key(radius_data, attr_P, attr_R_motion, num_steps, 0.0f, j);
BoundBox prev_bounds = BoundBox::empty;
point.bounds_grow(prev_key, prev_bounds);
for (int bvh_step = 1; bvh_step < num_bvh_steps; ++bvh_step) {
const float curr_time = (float)(bvh_step)*num_bvh_steps_inv_1;
float4 curr_key = point.motion_key(
points_data, radius_data, point_steps, num_points, num_steps, curr_time, j);
radius_data, attr_P, attr_R_motion, num_steps, curr_time, j);
BoundBox curr_bounds = BoundBox::empty;
point.bounds_grow(curr_key, curr_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();
size_t num_motion_steps = 1;
Attribute *motion_keys = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
if (motion_blur && hair->get_use_motion_blur() && motion_keys) {
const Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
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();
}
@ -417,10 +418,9 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
uint64_t numKeys = hair->num_keys();
uint64_t numCurves = hair->num_curves();
const float *radiuses = hair->get_radius();
/* Gather the curve geometry. */
std::vector<float3> cpData;
std::vector<packed_float3> cpData;
std::vector<int> idxData;
std::vector<float> radiusData;
cpData.reserve(numKeys);
@ -428,14 +428,9 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
std::vector<int> step_offsets;
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 = hair->get_position();
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;
}
const packed_float3 *keys = attr_P->data_at_time_step<packed_float3>(step,
num_motion_steps);
const float *radii = attr_R->data_at_time_step<float>(step, num_motion_steps);
step_offsets.push_back(cpData.size());
@ -446,20 +441,20 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
uint64_t idxBase = cpData.size();
if (hair->curve_shape != CURVE_THICK_LINEAR) {
cpData.push_back(keys[firstKey]);
radiusData.push_back(radiuses[firstKey]);
radiusData.push_back(radii[firstKey]);
}
for (int s = 0; s < segCount; ++s) {
if (step == 0) {
idxData.push_back(idxBase + 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]);
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) {
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];
cpBuffer = [mtl_device newBufferWithBytes:cpData.data()
length:cpData.size() * sizeof(float3)
length:cpData.size() * sizeof(packed_float3)
options:MTLResourceStorageModeShared];
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) {
MTLMotionKeyframeData *k = [MTLMotionKeyframeData data];
k.buffer = cpBuffer;
k.offset = step_offsets[step] * sizeof(float3);
k.offset = step_offsets[step] * sizeof(packed_float3);
cp_ptrs.push_back(k);
k = [MTLMotionKeyframeData data];
@ -501,7 +496,7 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
/* controlPointCount should specify the *per-step* control point count. */
geomDescCrv.controlPointCount = cpData.size() / num_motion_steps;
geomDescCrv.controlPointStride = sizeof(float3);
geomDescCrv.controlPointStride = sizeof(packed_float3);
geomDescCrv.controlPointFormat = MTLAttributeFormatFloat3;
geomDescCrv.radiusStride = sizeof(float);
geomDescCrv.radiusFormat = MTLAttributeFormatFloat;
@ -749,14 +744,13 @@ bool BVHMetal::build_BLAS_pointcloud(Progress &progress,
}
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();
size_t num_motion_steps = 1;
Attribute *motion_keys = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
if (motion_blur && pointcloud->get_use_motion_blur() && motion_keys) {
Attribute *attr_P = pointcloud->attributes.find(ATTR_STD_POSITION);
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();
}
@ -769,33 +763,19 @@ bool BVHMetal::build_BLAS_pointcloud(Progress &progress,
MTLAxisAlignedBoundingBox *aabb_data = (MTLAxisAlignedBoundingBox *)[aabbBuf contents];
/* 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) {
if (step == center_step) {
/* The center step for motion vertices is not stored in the attribute */
for (size_t j = 0; j < num_points; ++j) {
const PointCloud::Point point = pointcloud->get_point(j);
BoundBox bounds = BoundBox::empty;
point.bounds_grow(points, radius, bounds);
const packed_float3 *step_points = attr_P->data_at_time_step<packed_float3>(
step, num_motion_steps);
const float *step_radius = attr_R->data_at_time_step<float>(step, num_motion_steps);
const size_t index = step * num_points + j;
aabb_data[index].min = (MTLPackedFloat3 &)bounds.min;
aabb_data[index].max = (MTLPackedFloat3 &)bounds.max;
}
}
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 PointCloud::Point point = pointcloud->get_point(j);
BoundBox bounds = BoundBox::empty;
point.bounds_grow(step_points, step_radius, bounds);
for (size_t j = 0; j < num_points; ++j) {
const PointCloud::Point point = pointcloud->get_point(j);
BoundBox bounds = BoundBox::empty;
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;
}
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();
size_t num_motion_steps = 1;
Attribute *motion_keys = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
if (pipeline_options.usesMotionBlur && hair->get_use_motion_blur() && motion_keys) {
const Attribute *attr_P = hair->attributes.find(ATTR_STD_POSITION);
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();
}
@ -1360,19 +1361,9 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
/* Get AABBs for each motion step. */
for (size_t step = 0; step < num_motion_steps; ++step) {
/* The center step for motion vertices is not stored in the attribute. */
const size_t center_step = (num_motion_steps - 1) / 2;
const packed_float3 *keys_center = nullptr;
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]); };
const packed_float3 *keys = attr_P->data_at_time_step<packed_float3>(step,
num_motion_steps);
const float *curve_radius_step = attr_R->data_at_time_step<float>(step, num_motion_steps);
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;
@ -1380,7 +1371,6 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
++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) {
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) {
index_data[segment_index++] = vertex_index;
}
const float3 key = get_key(first_key_index + k);
vertex_data[vertex_index++] = make_float4(key, curve_radius[first_key_index + k]);
vertex_data[vertex_index++] = make_float4(float3(keys[first_key_index + k]),
curve_radius_step[first_key_index + k]);
}
const int last_key_index = first_key_index + curve.num_keys - 1;
{
const float3 key = get_key(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]),
curve_radius_step[last_key_index]);
}
else {
const int first_key_index = curve.first_key;
{
const float3 key = get_key(first_key_index);
vertex_data[vertex_index++] = make_float4(key, curve_radius[first_key_index]);
}
vertex_data[vertex_index++] = make_float4(float3(keys[first_key_index]),
curve_radius_step[first_key_index]);
for (int k = 0; k < curve.num_segments(); ++k) {
if (step == 0) {
index_data[segment_index++] = vertex_index - 1;
}
const float3 key = get_key(first_key_index + k);
vertex_data[vertex_index++] = make_float4(key, curve_radius[first_key_index + k]);
vertex_data[vertex_index++] = make_float4(float3(keys[first_key_index + k]),
curve_radius_step[first_key_index + k]);
}
const int last_key_index = first_key_index + curve.num_keys - 1;
{
const float3 key = get_key(last_key_index);
vertex_data[vertex_index++] = make_float4(key, curve_radius[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]),
curve_radius_step[last_key_index]);
vertex_data[vertex_index++] = make_float4(float3(keys[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) {
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;
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;
Attribute *motion_points = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
if (pipeline_options.usesMotionBlur && pointcloud->get_use_motion_blur() && motion_points) {
Attribute *attr_P = pointcloud->attributes.find(ATTR_STD_POSITION);
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();
}
@ -1584,44 +1572,22 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
/* Get AABBs for each motion step. */
for (size_t step = 0; step < num_motion_steps; ++step) {
/* The center step for motion vertices is not stored in the attribute. */
size_t center_step = (num_motion_steps - 1) / 2;
const packed_float3 *points = attr_P->data_at_time_step<packed_float3>(step,
num_motion_steps);
const float *radius = attr_R->data_at_time_step<float>(step, num_motion_steps);
if (step == center_step) {
const packed_float3 *points = pointcloud->get_position();
const float *radius = pointcloud->get_radius();
for (size_t i = 0; i < num_points; ++i) {
const PointCloud::Point point = pointcloud->get_point(i);
BoundBox bounds = BoundBox::empty;
point.bounds_grow(points, radius, bounds);
for (size_t i = 0; i < num_points; ++i) {
const PointCloud::Point point = pointcloud->get_point(i);
BoundBox bounds = BoundBox::empty;
point.bounds_grow(points, radius, bounds);
const size_t index = step * num_points + i;
aabb_data[index].minX = bounds.min.x;
aabb_data[index].minY = bounds.min.y;
aabb_data[index].minZ = bounds.min.z;
aabb_data[index].maxX = bounds.max.x;
aabb_data[index].maxY = bounds.max.y;
aabb_data[index].maxZ = bounds.max.z;
}
}
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;
}
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 */
KERNEL_DATA_ARRAY(KernelCurve, curves)
KERNEL_DATA_ARRAY(float4, curve_keys)
KERNEL_DATA_ARRAY(KernelCurveSegment, curve_segments)
/* pointclouds */
KERNEL_DATA_ARRAY(float4, points)
KERNEL_DATA_ARRAY(uint, points_shader)
/* 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
* radius. */
const int position_offset = kernel_data_fetch(objects, object).position_offset;
float radius[4];
radius[0] = kernel_data_fetch(curve_keys, ka).w;
radius[1] = kernel_data_fetch(curve_keys, k0).w;
radius[2] = kernel_data_fetch(curve_keys, k1).w;
radius[3] = kernel_data_fetch(curve_keys, kb).w;
radius[0] = kernel_data_fetch(attributes_float4, position_offset + ka).w;
radius[1] = kernel_data_fetch(attributes_float4, position_offset + k0).w;
radius[2] = kernel_data_fetch(attributes_float4, position_offset + k1).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]);
/* 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];
if (!is_motion) {
curve[0] = kernel_data_fetch(curve_keys, ka);
curve[1] = kernel_data_fetch(curve_keys, k0);
curve[2] = kernel_data_fetch(curve_keys, k1);
curve[3] = kernel_data_fetch(curve_keys, kb);
const int position_offset = kernel_data_fetch(objects, object).position_offset;
curve[0] = kernel_data_fetch(attributes_float4, position_offset + ka);
curve[1] = kernel_data_fetch(attributes_float4, position_offset + k0);
curve[2] = kernel_data_fetch(attributes_float4, position_offset + k1);
curve[3] = kernel_data_fetch(attributes_float4, position_offset + kb);
}
else {
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
# endif
{
P_curve[0] = kernel_data_fetch(curve_keys, k0);
P_curve[1] = kernel_data_fetch(curve_keys, k1);
const int position_offset = kernel_data_fetch(objects, sd->object).position_offset;
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);
@ -146,8 +147,9 @@ ccl_device float3 curve_motion_center_location(KernelGlobals kg, const ccl_priva
float4 P_curve[2];
P_curve[0] = kernel_data_fetch(curve_keys, k0);
P_curve[1] = kernel_data_fetch(curve_keys, k1);
const int position_offset = kernel_data_fetch(objects, sd->object).position_offset;
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);
}

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];
if (!is_motion) {
curve[0] = kernel_data_fetch(curve_keys, ka);
curve[1] = kernel_data_fetch(curve_keys, k0);
curve[2] = kernel_data_fetch(curve_keys, k1);
curve[3] = kernel_data_fetch(curve_keys, kb);
const int position_offset = kernel_data_fetch(objects, object).position_offset;
curve[0] = kernel_data_fetch(attributes_float4, position_offset + ka);
curve[1] = kernel_data_fetch(attributes_float4, position_offset + k0);
curve[2] = kernel_data_fetch(attributes_float4, position_offset + k1);
curve[3] = kernel_data_fetch(attributes_float4, position_offset + kb);
}
else {
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];
if (!(sd->type & PRIMITIVE_MOTION)) {
P_curve[0] = kernel_data_fetch(curve_keys, ka);
P_curve[1] = kernel_data_fetch(curve_keys, k0);
P_curve[2] = kernel_data_fetch(curve_keys, k1);
P_curve[3] = kernel_data_fetch(curve_keys, kb);
const int position_offset = kernel_data_fetch(objects, sd->object).position_offset;
P_curve[0] = kernel_data_fetch(attributes_float4, position_offset + ka);
P_curve[1] = kernel_data_fetch(attributes_float4, position_offset + k0);
P_curve[2] = kernel_data_fetch(attributes_float4, position_offset + k1);
P_curve[3] = kernel_data_fetch(attributes_float4, position_offset + kb);
}
else {
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
* 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__
@ -32,21 +32,18 @@ ccl_device_inline void motion_curve_keys_for_step_linear(KernelGlobals kg,
{
const int center_step = (numsteps - 1) / 2;
if (step == center_step) {
/* center step: regular key location */
keys[0] = kernel_data_fetch(curve_keys, k0);
keys[1] = kernel_data_fetch(curve_keys, k1);
/* Center step: first in the array. */
}
else {
/* center step is not stored in this array */
if (step > center_step) {
step--;
/* Non-center step, stored after center with center index skipped. */
if (step < center_step) {
step++;
}
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 */
@ -66,11 +63,8 @@ ccl_device_inline void motion_curve_keys_linear(KernelGlobals kg,
const int step = min((int)(time * maxstep), maxstep - 1);
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 */
const int offset = kernel_data_fetch(objects, object).position_offset;
float4 next_keys[2];
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;
if (step == center_step) {
/* center step: regular key location */
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);
/* Center step: first in the array. */
}
else {
/* center step is not stored in this array */
if (step > center_step) {
step--;
/* Non-center step, stored after center with center index skipped. */
if (step < center_step) {
step++;
}
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 */
@ -134,11 +123,8 @@ ccl_device_inline void motion_curve_keys(KernelGlobals kg,
const int step = min((int)(time * maxstep), maxstep - 1);
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 */
const int offset = kernel_data_fetch(objects, object).position_offset;
float4 next_keys[4];
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
* 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__
@ -26,16 +26,16 @@ ccl_device_inline float4 motion_point_for_step(
{
const int center_step = (numsteps - 1) / 2;
if (step == center_step) {
/* center step: regular key location */
return kernel_data_fetch(points, prim);
/* Center step: first in the array. */
}
/* center step is not stored in this array */
if (step > center_step) {
step--;
else {
/* Non-center step, stored after center with center index skipped. */
if (step < center_step) {
step++;
}
offset += step * numverts;
}
offset += step * numverts;
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 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 */
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 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) {
/* World space center. */
const int position_offset = kernel_data_fetch(objects, sd->object).position_offset;
float3 P = (sd->type & PRIMITIVE_MOTION) ?
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)) {
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) {
/* 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) {
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)
{
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__ */

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 int type)
{
const float4 point = (type & PRIMITIVE_MOTION) ? motion_point(kg, object, prim, time) :
kernel_data_fetch(points, prim);
const int position_offset = kernel_data_fetch(objects, object).position_offset;
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)) {
return false;
@ -104,9 +106,11 @@ ccl_device_inline void point_shader_setup(KernelGlobals kg,
# endif
/* Compute point center for normal. */
float3 center = make_float3((isect->type & PRIMITIVE_MOTION) ?
motion_point(kg, sd->object, sd->prim, sd->time) :
kernel_data_fetch(points, sd->prim));
const int position_offset = kernel_data_fetch(objects, sd->object).position_offset;
float3 center = make_float3(
(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)) {
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;
/* 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 (is_curve_or_point) {
AttributeDescriptor desc = find_attribute(kg, sd, ATTR_STD_MOTION_VERTEX_POSITION);
if (is_attribute_found(desc)) {
const ccl_global KernelObject *kobject = &kernel_data_fetch(objects, sd->object);
const int numverts = kobject->numverts;
const int num_motion_steps = kobject->num_geom_steps;
if (sd->object_flag & SD_OBJECT_HAS_VERTEX_MOTION) {
AttributeDescriptor desc;
desc.type = NODE_ATTR_FLOAT4;
desc.element = (sd->type & PRIMITIVE_CURVE) ? ATTR_ELEMENT_CURVE_KEY : ATTR_ELEMENT_VERTEX;
desc.offset = pos_offset + numverts;
*motion_pre = make_float3(primitive_surface_attribute<float4>(kg, sd, desc));
if (num_motion_steps > 2) {
desc.offset += numverts;
@ -237,27 +243,18 @@ ccl_device_forceinline void primitive_motion_data_without_camera(KernelGlobals k
else {
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
#endif
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) {
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;
desc.element = ATTR_ELEMENT_VERTEX;
desc.type = NODE_ATTR_FLOAT3;
desc.offset = kobject.position_offset + numverts;
desc.offset = pos_offset + numverts;
*motion_pre = triangle_attribute<float3>(kg, sd, desc);
if (num_motion_steps > 2) {
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];
P_curve[0] = kernel_data_fetch(curve_keys, ka);
P_curve[1] = kernel_data_fetch(curve_keys, k0);
P_curve[2] = kernel_data_fetch(curve_keys, k1);
P_curve[3] = kernel_data_fetch(curve_keys, kb);
const int position_offset = kernel_data_fetch(objects, object).position_offset;
P_curve[0] = kernel_data_fetch(attributes_float4, position_offset + ka);
P_curve[1] = kernel_data_fetch(attributes_float4, position_offset + k0);
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. */
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 int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
const int k1 = k0 + 1;
const float radius = mix(
kernel_data_fetch(curve_keys, k0).w, kernel_data_fetch(curve_keys, k1).w, sd->u);
const int position_offset = kernel_data_fetch(objects, sd->object).position_offset;
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;
}

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 int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
const int k1 = k0 + 1;
const float radius = mix(
kernel_data_fetch(curve_keys, k0).w, kernel_data_fetch(curve_keys, k1).w, sd->u);
const int position_offset = kernel_data_fetch(objects, sd->object).position_offset;
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;
}

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
* 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_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_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_MOTION = ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_IS_NORMAL |
ATTR_ELEMENT_IS_MOTION,
};
enum AttributeStandard : int {
@ -782,9 +774,6 @@ enum AttributeStandard : int {
ATTR_STD_POSITION_UNDEFORMED,
ATTR_STD_POSITION_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_CURVE_INTERCEPT,
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;
}
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,
const AttributeElement element,
const size_t data_sizeof,
@ -79,7 +87,7 @@ static void free_step_buffer(Attribute::Buffer &buf,
}
else if (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.sharing_info = nullptr;
@ -97,12 +105,10 @@ Attribute::Attribute(Attribute &&other)
void Attribute::free_data()
{
const size_t sz = data_sizeof();
free_step_buffer(center, element, sz, size);
/* Motion steps are never voxels, strip that flag for the helper. */
const AttributeElement motion_element = AttributeElement(element & ~ATTR_ELEMENT_VOXEL);
const size_t element_size = data_sizeof();
free_step_buffer(center, element, element_size, size);
for (Buffer &buf : motion) {
free_step_buffer(buf, motion_element, sz, size);
free_step_buffer(buf, element, element_size, size);
}
motion.clear();
}
@ -127,22 +133,23 @@ void Attribute::resize(const size_t num_elements)
if (num_elements == size_t(size)) {
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 alloc_bytes = attribute_alloc_bytes(element_size, num_elements);
/* Allocate and copy center step. */
Buffer new_center;
new_center.data = GuardedAllocator<char>().allocate(num_elements * sz);
new_center.data = GuardedAllocator<char>().allocate(alloc_bytes);
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. */
vector<Buffer> new_motion(motion.size());
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) {
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)
{
const int new_motion_count = int(geom->get_motion_steps()) - 1;
assert(new_motion_count >= 0);
if (new_motion_count == int(motion.size())) {
const int motion_steps = geom->get_motion_steps();
if (motion_steps <= 0) {
return;
}
const size_t sz = data_sizeof();
const AttributeElement motion_element = AttributeElement(element & ~ATTR_ELEMENT_VOXEL);
if (new_motion_count < int(motion.size())) {
/* Shrink: free and drop extra steps. */
for (size_t i = new_motion_count; i < motion.size(); i++) {
free_step_buffer(motion[i], motion_element, sz, size);
const int motion_size = geom->get_motion_steps() - 1;
if (motion_size == motion.size()) {
return;
}
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 {
/* Grow: allocate fresh zero-initialized arrays for the new steps. */
motion.reserve(new_motion_count);
while (int(motion.size()) < new_motion_count) {
motion.reserve(motion_size);
while (motion.size() < motion_size) {
Buffer buf;
if (size > 0) {
const size_t alloc_bytes = sz * size;
buf.data = GuardedAllocator<char>().allocate(alloc_bytes);
memset(const_cast<void *>(buf.data), 0, alloc_bytes);
/* Left uninitialized, callers fill in the motion data for every step. */
buf.data = GuardedAllocator<char>().allocate(attribute_alloc_bytes(element_size, size));
}
motion.push_back(buf);
}
@ -191,10 +198,9 @@ void Attribute::remove_motion()
if (!has_motion()) {
return;
}
const size_t sz = data_sizeof();
const AttributeElement motion_element = AttributeElement(element & ~ATTR_ELEMENT_VOXEL);
const size_t element_size = data_sizeof();
for (Buffer &buf : motion) {
free_step_buffer(buf, motion_element, sz, size);
free_step_buffer(buf, element, element_size, size);
}
motion.clear();
modified = true;
@ -210,7 +216,7 @@ void Attribute::take_motion_from(Attribute &other)
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) {
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
* 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. */
auto *new_data = GuardedAllocator<char>().allocate(data_sizeof * size);
memcpy(new_data, buf.data, data_sizeof * size);
auto *new_data = GuardedAllocator<char>().allocate(attribute_alloc_bytes(element_size, size));
memcpy(new_data, buf.data, element_size * size);
g_implicit_sharing_user_remove_fn(buf.sharing_info);
buf.sharing_info = nullptr;
buf.data = new_data;
@ -249,10 +255,10 @@ void Attribute::set_data_from(Attribute &&other)
flags = other.flags;
const size_t sz = data_sizeof();
const AttributeElement motion_element = AttributeElement(element & ~ATTR_ELEMENT_VOXEL);
const size_t element_size = data_sizeof();
const auto take_all = [&]() {
/* If topology or motion steps differ, take all data. */
if (size != other.size || motion.size() != other.motion.size()) {
free_data();
center = other.center;
motion = std::move(other.motion);
@ -260,44 +266,37 @@ void Attribute::set_data_from(Attribute &&other)
other.center = Buffer();
other.size = 0;
modified = true;
};
/* If topology or step count differ, take the whole thing. */
if (size != other.size || motion.size() != other.motion.size()) {
take_all();
return;
}
/* Same shape: compare each step independently, taking only those that differ
* so that unchanged, implicitly-shared steps keep their sharing. */
/* Compare each step independently. */
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;
src = Buffer();
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) {
/* Same buffer, e.g. shared through implicit sharing. */
return false;
}
if (a.sharing_info != b.sharing_info) {
return true;
}
if (size == 0) {
if (a.sharing_info != b.sharing_info) {
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);
}
for (size_t i = 0; i < motion.size(); i++) {
/* Motion steps are never voxels, strip that flag for the helper. */
if (step_differs(motion[i], other.motion[i])) {
take_step(motion[i], other.motion[i], motion_element);
if (!step_equals(motion[i], other.motion[i])) {
take_step(motion[i], other.motion[i], element);
}
}
}
@ -361,13 +360,7 @@ size_t Attribute::element_size(Geometry *geom,
size = pointcloud->num_points();
}
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:
if (geom->is_mesh() || geom->is_volume()) {
Mesh *mesh = static_cast<Mesh *>(geom);
@ -405,14 +398,6 @@ size_t Attribute::element_size(Geometry *geom,
size = hair->num_keys();
}
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:
break;
}
@ -477,11 +462,6 @@ const char *Attribute::standard_name(AttributeStandard std)
return "undisplaced";
case ATTR_STD_NORMAL_UNDISPLACED:
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:
return "particle";
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_UNDISPLACED:
return TypePoint;
case ATTR_STD_MOTION_VERTEX_POSITION:
return TypePoint;
case ATTR_STD_MOTION_VERTEX_NORMAL:
return TypeNormal;
case ATTR_STD_CORNER_NORMAL:
return TypeNormal;
case ATTR_STD_MOTION_CORNER_NORMAL:
return TypeNormal;
case ATTR_STD_PTEX_FACE_ID:
return TypeFloat;
case ATTR_STD_PTEX_UV:
@ -750,8 +724,6 @@ static TypeDesc find_type_from_geometry_std(Geometry *geometry, AttributeStandar
return TypeFloat2;
case ATTR_STD_GENERATED:
return TypePoint;
case ATTR_STD_MOTION_VERTEX_POSITION:
return TypeFloat4;
case ATTR_STD_POINT_RANDOM:
return TypeFloat;
case ATTR_STD_GENERATED_TRANSFORM:
@ -796,14 +768,10 @@ static TypeDesc find_type_from_geometry_std(Geometry *geometry, AttributeStandar
return TypeFloat;
case ATTR_STD_VERTEX_NORMAL:
return TypeNormal;
case ATTR_STD_MOTION_VERTEX_NORMAL:
return TypeNormal;
case ATTR_STD_UV:
return TypeFloat2;
case ATTR_STD_GENERATED:
return TypePoint;
case ATTR_STD_MOTION_VERTEX_POSITION:
return TypeFloat4;
case ATTR_STD_CURVE_INTERCEPT:
return TypeFloat;
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_UNDISPLACED:
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:
return ATTR_ELEMENT_CORNER_NORMAL;
case ATTR_STD_MOTION_CORNER_NORMAL:
return ATTR_ELEMENT_CORNER_NORMAL_MOTION;
case ATTR_STD_PTEX_FACE_ID:
return ATTR_ELEMENT_FACE;
case ATTR_STD_PTEX_UV:
@ -884,8 +846,6 @@ static AttributeElement find_element_from_geometry_std(Geometry *geometry, Attri
return ATTR_ELEMENT_VERTEX;
case ATTR_STD_GENERATED:
return ATTR_ELEMENT_VERTEX;
case ATTR_STD_MOTION_VERTEX_POSITION:
return ATTR_ELEMENT_VERTEX_MOTION;
case ATTR_STD_POINT_RANDOM:
return ATTR_ELEMENT_VERTEX;
case ATTR_STD_GENERATED_TRANSFORM:
@ -930,14 +890,10 @@ static AttributeElement find_element_from_geometry_std(Geometry *geometry, Attri
return ATTR_ELEMENT_CURVE_KEY;
case ATTR_STD_VERTEX_NORMAL:
return ATTR_ELEMENT_CURVE_KEY_NORMAL;
case ATTR_STD_MOTION_VERTEX_NORMAL:
return ATTR_ELEMENT_CURVE_KEY_NORMAL_MOTION;
case ATTR_STD_UV:
return ATTR_ELEMENT_CURVE;
case ATTR_STD_GENERATED:
return ATTR_ELEMENT_CURVE;
case ATTR_STD_MOTION_VERTEX_POSITION:
return ATTR_ELEMENT_CURVE_KEY_MOTION;
case ATTR_STD_CURVE_INTERCEPT:
return ATTR_ELEMENT_CURVE_KEY;
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);
copy_attr.std = attr.std;
if (attr.has_motion()) {
copy_attr.add_motion(geometry);
}
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_vindex(device, "tri_vindex", MEM_GLOBAL),
curves(device, "curves", MEM_GLOBAL),
curve_keys(device, "curve_keys", MEM_GLOBAL),
curve_segments(device, "curve_segments", MEM_GLOBAL),
points(device, "points", MEM_GLOBAL),
points_shader(device, "points_shader", MEM_GLOBAL),
objects(device, "objects", MEM_GLOBAL),
object_motion_pass(device, "object_motion_pass", MEM_GLOBAL),

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

@ -30,11 +30,9 @@ class DeviceScene {
device_vector<packed_uint3> tri_vindex;
device_vector<KernelCurve> curves;
device_vector<float4> curve_keys;
device_vector<KernelCurveSegment> curve_segments;
/* point-cloud */
device_vector<float4> points;
device_vector<uint> points_shader;
/* 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 curve_size = 0;
size_t curve_key_size = 0;
size_t curve_segment_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);
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->prim_offset = curve_size;
curve_size += hair->num_curves();
curve_key_size += hair->num_keys();
curve_segment_size += hair->num_segments();
}
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) {
dscene->curves.tag_realloc();
dscene->curve_keys.tag_realloc();
dscene->curve_segments.tag_realloc();
}
if (device_update_flags & DEVICE_POINT_DATA_NEEDS_REALLOC) {
dscene->points.tag_realloc();
dscene->points_shader.tag_realloc();
}
}
@ -691,13 +686,11 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
}
if (device_update_flags & DEVICE_CURVE_DATA_MODIFIED) {
dscene->curve_keys.tag_modified();
dscene->curves.tag_modified();
dscene->curve_segments.tag_modified();
}
if (device_update_flags & DEVICE_POINT_DATA_MODIFIED) {
dscene->points.tag_modified();
dscene->points_shader.tag_modified();
}
@ -984,7 +977,13 @@ void GeometryManager::device_update(Device *device,
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) {
if (scene->update_stats) {
scene->update_stats->geometry.times.add_entry({"device_update (attributes)", time});
@ -1057,20 +1056,12 @@ void GeometryManager::device_update(Device *device,
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) {
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_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,
@ -1154,6 +1145,22 @@ void GeometryManager::device_update(Device *device,
dscene->data.bvh.bvh_layout = BVHParams::best_bvh_layout(
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) {
if (scene->update_stats) {
@ -1194,9 +1201,7 @@ void GeometryManager::device_update(Device *device,
dscene->tri_shader.clear_modified();
dscene->tri_vindex.clear_modified();
dscene->curves.clear_modified();
dscene->curve_keys.clear_modified();
dscene->curve_segments.clear_modified();
dscene->points.clear_modified();
dscene->points_shader.clear_modified();
dscene->attributes_map.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_vindex.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->points.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_float.free_if_need_realloc(force_free);

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

@ -327,12 +327,24 @@ class AttributeTableBuilder {
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;
type = mattr->type;
/* store attribute data in arrays */
const size_t size = Attribute::element_size(geom, mattr->element, prim);
/* Store attribute data in arrays, including possible motion steps. */
const size_t per_step = Attribute::element_size(geom, mattr->element, prim);
const int num_motion = mattr->motion.size();
const AttributeElement &element = desc.element;
int &offset = desc.offset;
@ -343,42 +355,53 @@ class AttributeTableBuilder {
offset = handle.kernel_id();
}
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) {
offset = attr_normal.add(mattr->data<packed_normal>(), size, mattr->modified);
for (int step = 0; step < int(mattr->motion.size()); step++) {
attr_normal.add(mattr->data<packed_normal>(step + 1), size, mattr->modified);
offset = attr_normal.add(mattr->data<packed_normal>(), per_step, mattr->modified);
for (int step = 1; step <= num_motion; step++) {
attr_normal.add(mattr->data<packed_normal>(step), per_step, mattr->modified);
}
}
else if (mattr->type == TypeFloat) {
offset = attr_float.add(mattr->data<float>(), size, mattr->modified);
for (int step = 0; step < int(mattr->motion.size()); step++) {
attr_float.add(mattr->data<float>(step + 1), size, mattr->modified);
offset = attr_float.add(mattr->data<float>(), per_step, mattr->modified);
for (int step = 1; step <= num_motion; step++) {
attr_float.add(mattr->data<float>(step), per_step, mattr->modified);
}
}
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) {
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) {
offset = attr_float4.add(mattr->data<float4>(), size, mattr->modified);
for (int step = 0; step < int(mattr->motion.size()); step++) {
attr_float4.add(mattr->data<float4>(step + 1), size, mattr->modified);
offset = attr_float4.add(mattr->data<float4>(), per_step, mattr->modified);
for (int step = 1; step <= num_motion; step++) {
attr_float4.add(mattr->data<float4>(step), per_step, mattr->modified);
}
}
else {
offset = attr_float3.add(mattr->data<packed_float3>(), size, mattr->modified);
for (int step = 0; step < int(mattr->motion.size()); step++) {
attr_float3.add(mattr->data<packed_float3>(step + 1), size, mattr->modified);
offset = attr_float3.add(mattr->data<packed_float3>(), per_step, mattr->modified);
for (int step = 1; step <= num_motion; step++) {
attr_float3.add(mattr->data<packed_float3>(step), per_step, mattr->modified);
}
}
/* mesh vertex/curve index is global, not per object, so we sneak
* a correction for that in here */
if (geom->is_mesh()) {
/* Primitive index is global, not per object, so we sneak a correction
* for that in here. Vertex index is per object. */
if (geom->is_mesh() || geom->is_volume()) {
Mesh *mesh = static_cast<Mesh *>(geom);
if (element & ATTR_ELEMENT_FACE) {
offset -= mesh->prim_offset;
@ -392,9 +415,6 @@ class AttributeTableBuilder {
if (element & ATTR_ELEMENT_CURVE) {
offset -= hair->prim_offset;
}
else if (element & ATTR_ELEMENT_CURVE_KEY) {
offset -= hair->curve_key_offset;
}
}
else if (geom->is_pointcloud()) {
if (element & ATTR_ELEMENT_VERTEX) {
@ -409,9 +429,10 @@ class AttributeTableBuilder {
return;
}
const size_t base_size = Attribute::element_size(geom, mattr->element, prim);
/* Inline motion: reserve space for center step plus each motion sub-step. */
const size_t size = base_size * (1 + mattr->motion.size());
/* Must match the number of steps written by add(), which is derived from
* the attribute's own stored motion steps rather than geom->get_motion_steps(). */
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) {
/* 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()
{
attr_float.alloc();
@ -546,6 +608,16 @@ void GeometryManager::device_update_attributes(Device *device,
AttributeRequestSet &attributes = geom_attributes[i];
for (AttributeRequest &req : attributes.requests) {
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);
}
}

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

@ -31,7 +31,6 @@ void GeometryManager::device_update_mesh(Device * /*unused*/,
/* Count. */
size_t tri_size = 0;
size_t curve_key_size = 0;
size_t curve_size = 0;
size_t curve_segment_size = 0;
@ -46,7 +45,6 @@ void GeometryManager::device_update_mesh(Device * /*unused*/,
else if (geom->is_hair()) {
Hair *hair = static_cast<Hair *>(geom);
curve_key_size += hair->num_keys();
curve_size += hair->num_curves();
curve_segment_size += hair->num_segments();
}
@ -96,38 +94,31 @@ void GeometryManager::device_update_mesh(Device * /*unused*/,
if (curve_segment_size != 0) {
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);
KernelCurveSegment *curve_segments = dscene->curve_segments.alloc(curve_segment_size);
const bool copy_all_data = dscene->curve_keys.need_realloc() ||
dscene->curves.need_realloc() ||
const bool copy_all_data = dscene->curves.need_realloc() ||
dscene->curve_segments.need_realloc();
for (Geometry *geom : scene->geometry) {
if (geom->is_hair()) {
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() ||
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;
}
hair->pack_curves(scene,
&curve_keys[hair->curve_key_offset],
&curves[hair->prim_offset],
&curve_segments[hair->curve_segment_offset]);
hair->pack_curves(
scene, &curves[hair->prim_offset], &curve_segments[hair->curve_segment_offset]);
if (progress.get_cancel()) {
return;
}
}
}
dscene->curve_keys.copy_to_device_if_modified();
dscene->curves.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) {
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);
for (Geometry *geom : scene->geometry) {
if (geom->is_pointcloud()) {
PointCloud *pointcloud = static_cast<PointCloud *>(geom);
pointcloud->pack(
scene, &points[pointcloud->prim_offset], &points_shader[pointcloud->prim_offset]);
pointcloud->pack(scene, &points_shader[pointcloud->prim_offset]);
if (progress.get_cancel()) {
return;
}
}
}
dscene->points.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,
const float3 *positions,
const float *radius,
const packed_float3 *curve_keys,
const float *curve_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)];
P[0] = curve_keys[max(first_key + k - 1, first_key)];
P[1] = curve_keys[first_key + k];
P[2] = curve_keys[first_key + k + 1];
P[3] = curve_keys[min(first_key + k + 2, first_key + num_keys - 1)];
float3 lower;
float3 upper;
@ -60,24 +60,24 @@ void Hair::Curve::bounds_grow(const int k,
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]);
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 float3 *positions,
const float *radius,
const packed_float3 *curve_keys,
const float *curve_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] = curve_keys[max(first_key + k - 1, first_key)];
P[1] = curve_keys[first_key + k];
P[2] = curve_keys[first_key + k + 1];
P[3] = curve_keys[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]);
@ -91,63 +91,7 @@ void Hair::Curve::bounds_grow(const int k,
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,
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]);
const float mr = max(curve_radius[first_key + k], curve_radius[first_key + k + 1]);
bounds.grow(lower, 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);
}
void Hair::Curve::motion_keys(const packed_float3 *positions,
const float *radius,
const float4 *key_steps,
const size_t num_positions,
/* Get position and radius arrays for a given time-ordered motion step. */
static void hair_step_buffers(const Attribute *attr_P,
const Attribute *attr_R,
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 float time,
size_t k0,
@ -192,18 +148,15 @@ void Hair::Curve::motion_keys(const packed_float3 *positions,
/* Fetch vertex coordinates. */
float4 curr_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(
positions, radius, key_steps, num_positions, num_steps, step + 1, k0, k1, next_keys);
keys_for_step(attr_P, attr_R, step, k0, k1, curr_keys);
keys_for_step(attr_P, attr_R, step + 1, k0, k1, next_keys);
/* Interpolate between steps. */
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];
}
void Hair::Curve::cardinal_motion_keys(const packed_float3 *positions,
const float *radius,
const float4 *key_steps,
const size_t num_positions,
void Hair::Curve::cardinal_motion_keys(const Attribute *attr_P,
const Attribute *attr_R,
const size_t num_steps,
const float time,
size_t k0,
@ -219,10 +172,8 @@ void Hair::Curve::cardinal_motion_keys(const packed_float3 *positions,
/* Fetch vertex coordinates. */
float4 curr_keys[4];
float4 next_keys[4];
cardinal_keys_for_step(
positions, radius, key_steps, num_positions, num_steps, step, k0, k1, k2, k3, curr_keys);
cardinal_keys_for_step(
positions, radius, key_steps, num_positions, num_steps, step + 1, k0, k1, k2, k3, next_keys);
cardinal_keys_for_step(attr_P, attr_R, step, k0, k1, k2, k3, curr_keys);
cardinal_keys_for_step(attr_P, attr_R, step + 1, k0, k1, k2, k3, next_keys);
/* Interpolate between steps. */
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];
@ -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];
}
void Hair::Curve::keys_for_step(const packed_float3 *positions,
const float *radius,
const float4 *key_steps,
const size_t num_positions,
const size_t num_steps,
size_t step,
void Hair::Curve::keys_for_step(const Attribute *attr_P,
const Attribute *attr_R,
const size_t step,
size_t k0,
size_t k1,
float4 r_keys[2]) const
{
k0 = max(k0, (size_t)0);
k1 = min(k1, (size_t)(num_keys - 1));
const size_t center_step = ((num_steps - 1) / 2);
if (step == center_step) {
/* Center step: regular key location. */
r_keys[0] = make_float4(positions[first_key + k0], radius[first_key + k0]);
r_keys[1] = make_float4(positions[first_key + k1], radius[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]);
}
const packed_float3 *P;
const float *R;
hair_step_buffers(attr_P, attr_R, step, P, R);
r_keys[0] = make_float4(P[first_key + k0], R[first_key + k0]);
r_keys[1] = make_float4(P[first_key + k1], R[first_key + k1]);
}
void Hair::Curve::cardinal_keys_for_step(const packed_float3 *positions,
const float *radius,
const float4 *key_steps,
const size_t num_positions,
const size_t num_steps,
size_t step,
void Hair::Curve::cardinal_keys_for_step(const Attribute *attr_P,
const Attribute *attr_R,
const size_t step,
size_t k0,
size_t k1,
size_t k2,
@ -279,37 +208,13 @@ void Hair::Curve::cardinal_keys_for_step(const packed_float3 *positions,
{
k0 = max(k0, (size_t)0);
k3 = min(k3, (size_t)(num_keys - 1));
const size_t center_step = ((num_steps - 1) / 2);
if (step == center_step) {
/* Center step: regular key location. */
r_keys[0] = make_float4(positions[first_key + k0], radius[first_key + k0]);
r_keys[1] = make_float4(positions[first_key + k1], radius[first_key + k1]);
r_keys[2] = make_float4(positions[first_key + k2], radius[first_key + k2]);
r_keys[3] = make_float4(positions[first_key + k3], radius[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]);
}
const packed_float3 *P;
const float *R;
hair_step_buffers(attr_P, attr_R, step, P, R);
r_keys[0] = make_float4(P[first_key + k0], R[first_key + k0]);
r_keys[1] = make_float4(P[first_key + k1], R[first_key + k1]);
r_keys[2] = make_float4(P[first_key + k2], R[first_key + k2]);
r_keys[3] = make_float4(P[first_key + k3], R[first_key + k3]);
}
/* Hair */
@ -326,7 +231,6 @@ NODE_DEFINE(Hair)
Hair::Hair() : Geometry(get_node_type(), Geometry::HAIR)
{
curve_key_offset = 0;
curve_segment_offset = 0;
curve_shape = CURVE_RIBBON;
@ -335,12 +239,6 @@ Hair::Hair() : Geometry(get_node_type(), Geometry::HAIR)
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()
{
attributes.add(ATTR_STD_POSITION);
@ -349,8 +247,10 @@ void Hair::add_builtin_attributes()
void Hair::resize_curves(const int numcurves, const int numkeys)
{
attributes.add(ATTR_STD_POSITION)->resize(numkeys);
attributes.add(ATTR_STD_RADIUS)->resize(numkeys);
Attribute *attr_P = attributes.add(ATTR_STD_POSITION);
attr_P->resize(numkeys);
Attribute *attr_R = attributes.add(ATTR_STD_RADIUS);
attr_R->resize(numkeys);
curve_first_key.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)
{
Attribute *attr_mP = attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
if (attr_mP) {
const packed_float3 *keys = get_position();
const size_t numkeys = num_keys();
std::copy_n(keys, numkeys, attr_mP->data_for_write<packed_float3>() + motion_step * numkeys);
const int attr_step = motion_step + 1;
const size_t numkeys = num_keys();
Attribute *attr_P = attributes.find(ATTR_STD_POSITION);
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);
if (attr_mvN && attr_vN) {
const packed_normal *vN = attr_vN->data<packed_normal>();
const size_t numkeys = num_keys();
std::copy_n(vN, numkeys, attr_mvN->data_for_write<packed_normal>() + motion_step * numkeys);
if (attr_vN && attr_vN->has_motion()) {
std::copy_n(attr_vN->data<packed_normal>(),
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()
{
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 packed_float3 *positions = get_position();
const float *radius = get_radius();
if (positions_size > 0) {
if (curve_keys_size > 0) {
bnds.grow(parallel_reduce(
blocked_range<size_t>(0, curve_num),
BoundBox(BoundBox::empty),
@ -420,7 +325,7 @@ void Hair::compute_bounds()
const Curve curve = get_curve(i);
const int num_segments = curve.num_segments();
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;
@ -431,15 +336,13 @@ void Hair::compute_bounds()
return combined_bounds;
}));
Attribute *curve_attr = attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
if (use_motion_blur && curve_attr) {
const size_t steps_size = positions_size * (motion_steps - 1);
// Attribute data is stored as a float4 and is not
// interchangeable with float3
const float4 *key_steps = curve_attr->data<float4>();
for (size_t i = 0; i < steps_size; i++) {
bnds.grow(make_float3(key_steps[i]));
const Attribute *attr_P = attributes.find(ATTR_STD_POSITION);
if (use_motion_blur && attr_P->has_motion()) {
for (int attr_step = 1; attr_step < attr_P->num_motion_steps(); attr_step++) {
const packed_float3 *key_step = attr_P->data<packed_float3>(attr_step);
for (size_t i = 0; i < curve_keys_size; i++) {
bnds.grow(key_step[i]);
}
}
}
@ -447,18 +350,16 @@ void Hair::compute_bounds()
bnds = BoundBox::empty;
/* skip nan or inf coordinates */
for (size_t i = 0; i < positions_size; i++) {
bnds.grow_safe(positions[i], radius[i]);
for (size_t i = 0; i < curve_keys_size; i++) {
bnds.grow_safe(curve_keys_data[i], curve_radius_data[i]);
}
if (use_motion_blur && curve_attr) {
const size_t steps_size = positions_size * (motion_steps - 1);
// Attribute data is stored as a float4 which is not
// interchangeable with float4
const float4 *key_steps = curve_attr->data<float4>();
for (size_t i = 0; i < steps_size; i++) {
bnds.grow_safe(make_float3(key_steps[i]));
if (use_motion_blur && attr_P->has_motion()) {
for (int attr_step = 1; attr_step < attr_P->num_motion_steps(); attr_step++) {
const packed_float3 *key_step = attr_P->data<packed_float3>(attr_step);
for (size_t i = 0; i < curve_keys_size; i++) {
bnds.grow_safe(key_step[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);
/* apply transform to curve keys */
const size_t num_keys_local = num_keys();
packed_float3 *positions_data = get_position_for_write();
float *radius_data = get_radius_for_write();
for (size_t i = 0; i < num_keys_local; i++) {
const float3 co = transform_point(&tfm, positions_data[i]);
const float radius = radius_data[i] * scalar;
packed_float3 *keys = get_position_for_write();
float *radius = get_radius_for_write();
const size_t numkeys = num_keys();
for (size_t i = 0; i < numkeys; i++) {
const float3 co = transform_point(&tfm, keys[i]);
const float r = radius[i] * scalar;
/* scale for curve radius is only correct for uniform scale */
positions_data[i] = co;
radius_data[i] = radius;
keys[i] = co;
radius[i] = r;
}
tag_position_modified();
tag_radius_modified();
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) {
/* apply transform to motion curve keys */
const size_t steps_size = num_keys_local * (motion_steps - 1);
float4 *key_steps = curve_attr->data_for_write<float4>();
for (size_t i = 0; i < steps_size; i++) {
const float3 co = transform_point(&tfm, make_float3(key_steps[i]));
const float radius = key_steps[i].w * scalar;
/* scale for curve radius is only correct for uniform scale */
key_steps[i] = make_float4(co);
key_steps[i].w = radius;
if (attr_P->has_motion()) {
const bool has_motion_radius = attr_R->has_motion();
const size_t nk = num_keys();
for (int step = 1; step <= int(attr_P->motion.size()); step++) {
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 < nk; i++) {
motion_P[i] = transform_point(&tfm, motion_P[i]);
if (motion_R) {
motion_R[i] *= scalar;
}
}
}
}
}
}
void Hair::pack_curves(Scene *scene,
float4 *curve_key_co,
KernelCurve *curves,
KernelCurveSegment *curve_segments)
void Hair::pack_curves(Scene *scene, 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 */
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);
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].type = type;

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

@ -25,40 +25,26 @@ class Hair : public Geometry {
}
void bounds_grow(const int k,
const float3 *positions,
const float *curve_radius,
BoundBox &bounds) const;
void bounds_grow(const int k,
const packed_float3 *positions,
const packed_float3 *curve_keys,
const float *curve_radius,
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 float3 keys[4], BoundBox &bounds) const;
void bounds_grow(const int k,
const float3 *positions,
const float *curve_radius,
const Transform &aligned_space,
BoundBox &bounds) const;
void bounds_grow(const int k,
const packed_float3 *positions,
const packed_float3 *curve_keys,
const float *curve_radius,
const Transform &aligned_space,
BoundBox &bounds) const;
void motion_keys(const packed_float3 *positions,
const float *curve_radius,
const float4 *key_steps,
const size_t num_positions,
void motion_keys(const Attribute *attr_P,
const Attribute *attr_R,
const size_t num_steps,
const float time,
size_t k0,
size_t k1,
float4 r_keys[2]) const;
void cardinal_motion_keys(const packed_float3 *positions,
const float *curve_radius,
const float4 *key_steps,
const size_t num_positions,
void cardinal_motion_keys(const Attribute *attr_P,
const Attribute *attr_R,
const size_t num_steps,
const float time,
size_t k0,
@ -67,20 +53,14 @@ class Hair : public Geometry {
size_t k3,
float4 r_keys[4]) const;
void keys_for_step(const packed_float3 *positions,
const float *curve_radius,
const float4 *key_steps,
const size_t num_positions,
const size_t num_steps,
void keys_for_step(const Attribute *attr_P,
const Attribute *attr_R,
const size_t step,
size_t k0,
size_t k1,
float4 r_keys[2]) const;
void cardinal_keys_for_step(const packed_float3 *positions,
const float *curve_radius,
const float4 *key_steps,
const size_t num_positions,
const size_t num_steps,
void cardinal_keys_for_step(const Attribute *attr_P,
const Attribute *attr_R,
const size_t step,
size_t k0,
size_t k1,
@ -93,7 +73,6 @@ class Hair : public Geometry {
NODE_SOCKET_API_ARRAY(array<int>, curve_shader)
/* BVH */
size_t curve_key_offset;
size_t curve_segment_offset;
CurveShapeType curve_shape;
@ -115,14 +94,17 @@ class Hair : public Geometry {
Curve get_curve(const size_t i) const
{
const int first = curve_first_key[i];
const int next_first = (i + 1 < curve_first_key.size()) ? curve_first_key[i + 1] :
int(num_keys());
const int next_first = (i + 1 < curve_first_key.size()) ? curve_first_key[i + 1] : num_keys();
Curve curve = {first, next_first - first};
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
{
@ -143,10 +125,7 @@ class Hair : public Geometry {
void get_uv_tiles(ustring map, unordered_set<int> &tiles) override;
/* BVH */
void pack_curves(Scene *scene,
float4 *curve_key_co,
KernelCurve *curve,
KernelCurveSegment *curve_segments);
void pack_curves(Scene *scene, KernelCurve *curve, KernelCurveSegment *curve_segments);
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 ||
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) {
kobject.position_offset = position_offset;
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);
}
float4 PointCloud::Point::motion_key(const packed_float3 *points,
const float *radius,
const float4 *point_steps,
const size_t num_points,
float4 PointCloud::Point::motion_key(const float *radius,
const Attribute *attr_P,
const Attribute *attr_R,
const size_t num_steps,
const float time,
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 float t = time * max_step - step;
/* Fetch vertex coordinates. */
const float4 curr_key = point_for_step(
points, radius, point_steps, num_points, num_steps, step, p);
const float4 next_key = point_for_step(
points, radius, point_steps, num_points, num_steps, step + 1, p);
const float4 curr_key = point_for_step(radius, attr_P, attr_R, step, p);
const float4 next_key = point_for_step(radius, attr_P, attr_R, step + 1, p);
/* Interpolate between steps. */
return (1.0f - t) * curr_key + t * next_key;
}
float4 PointCloud::Point::point_for_step(const packed_float3 *points,
const float *radius,
const float4 *point_steps,
const size_t num_points,
const size_t num_steps,
size_t step,
size_t p) const
float4 PointCloud::Point::point_for_step(const float *radius,
const Attribute *attr_P,
const Attribute *attr_R,
const size_t step,
const size_t p) const
{
const size_t center_step = ((num_steps - 1) / 2);
if (step == center_step) {
/* Center step: regular key location. */
return make_float4(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];
const int num_steps = attr_P->num_motion_steps();
const float r = attr_R ? attr_R->data_at_time_step<float>(step, num_steps)[p] : radius[p];
return make_float4(float3(attr_P->data_at_time_step<packed_float3>(step, num_steps)[p]), r);
}
/* PointCloud */
@ -94,12 +81,6 @@ PointCloud::PointCloud() : Geometry(get_node_type(), Geometry::POINTCLOUD)
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()
{
attributes.add(ATTR_STD_POSITION);
@ -108,8 +89,10 @@ void PointCloud::add_builtin_attributes()
void PointCloud::resize(const int numpoints)
{
attributes.add(ATTR_STD_POSITION)->resize(numpoints);
attributes.add(ATTR_STD_RADIUS)->resize(numpoints);
Attribute *attr_P = attributes.add(ATTR_STD_POSITION);
attr_P->resize(numpoints);
Attribute *attr_R = attributes.add(ATTR_STD_RADIUS);
attr_R->resize(numpoints);
shader.resize(numpoints);
attributes.resize();
@ -138,18 +121,17 @@ void PointCloud::clear(const bool preserve_shaders)
void PointCloud::copy_center_to_motion_step(const int motion_step)
{
Attribute *attr_mP = attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
if (attr_mP) {
const packed_float3 *points_data = get_position();
const size_t numpoints = num_points();
const float *radius_data = get_radius();
const int attr_step = motion_step + 1;
const size_t numpoints = num_points();
float4 *attrib_P = attr_mP->data_for_write<float4>() + motion_step * numpoints;
for (size_t i = 0; i < numpoints; i++) {
const float3 P = float3(points_data[i]);
const float r = radius_data[i];
attrib_P[i] = make_float4(P, r);
}
Attribute *attr_P = attributes.find(ATTR_STD_POSITION);
if (attr_P->has_motion()) {
std::copy_n(get_position(), numpoints, 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(), numpoints, attr_R->data_for_write<float>(attr_step));
}
}
@ -173,21 +155,23 @@ void PointCloud::compute_bounds()
{
BoundBox bnds = BoundBox::empty;
const size_t numpoints = num_points();
const packed_float3 *points = get_position();
const float *radius = get_radius();
const packed_float3 *points_data = get_position();
const float *radius_data = get_radius();
if (numpoints > 0) {
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);
if (use_motion_blur && attr) {
const size_t steps_size = numpoints * (motion_steps - 1);
const float4 *point_steps = attr->data<float4>();
for (size_t i = 0; i < steps_size; i++) {
bnds.grow(make_float3(point_steps[i]), point_steps[i].w);
const Attribute *attr_P = attributes.find(ATTR_STD_POSITION);
const Attribute *attr_R = attributes.find(ATTR_STD_RADIUS);
if (use_motion_blur && attr_P->has_motion()) {
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);
const float *motion_R = attr_R->data<float>(attr_step);
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 */
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) {
const size_t steps_size = numpoints * (motion_steps - 1);
const float4 *point_steps = attr->data<float4>();
for (size_t i = 0; i < steps_size; i++) {
bnds.grow_safe(make_float3(point_steps[i]), point_steps[i].w);
if (use_motion_blur && attr_P->has_motion()) {
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);
const float *motion_R = attr_R->data<float>(attr_step);
for (size_t i = 0; i < numpoints; i++) {
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 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();
packed_float3 *points = get_position_for_write();
float *radius = get_radius_for_write();
for (size_t i = 0; i < numpoints; i++) {
const float3 co = transform_point(&tfm, points[i]);
const float r = radius[i] * scalar;
const float3 co = transform_point(&tfm, points_data[i]);
const float r = radius_data[i] * scalar;
/* scale for curve radius is only correct for uniform scale
*/
points[i] = co;
radius[i] = r;
/* scale for radius is only correct for uniform scale */
points_data[i] = co;
radius_data[i] = r;
}
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) {
/* apply transform to motion curve keys */
const size_t steps_size = numpoints * (motion_steps - 1);
float4 *point_steps = attr->data_for_write<float4>();
for (size_t i = 0; i < steps_size; i++) {
const float3 co = transform_point(&tfm, make_float3(point_steps[i]));
const float radius = point_steps[i].w * scalar;
/* scale for curve radius is only correct for uniform
* scale */
point_steps[i] = make_float4(co);
point_steps[i].w = radius;
if (attr_P->has_motion()) {
const bool has_motion_radius = attr_R->has_motion();
for (int step = 1; step <= int(attr_P->motion.size()); step++) {
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 < numpoints; i++) {
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;
}
}
}
}
}
}
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 packed_float3 *points_data = get_position();
const float *radius_data = get_radius();
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 last_shader = -1;
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;
void bounds_grow(const float4 &point, BoundBox &bounds) const;
float4 motion_key(const packed_float3 *points,
const float *radius,
const float4 *point_steps,
const size_t num_points,
float4 motion_key(const float *radius,
const Attribute *attr_P,
const Attribute *attr_R,
const size_t num_steps,
const float time,
size_t p) const;
float4 point_for_step(const packed_float3 *points,
const float *radius,
const float4 *point_steps,
const size_t num_points,
const size_t num_steps,
float4 point_for_step(const float *radius,
const Attribute *attr_P,
const Attribute *attr_R,
const size_t step,
size_t p) const;
};
@ -63,7 +60,11 @@ class PointCloud : public Geometry {
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
{
@ -76,7 +77,7 @@ class PointCloud : public Geometry {
PrimitiveType primitive_type() const override;
/* BVH */
void pack(Scene *scene, float4 *packed_points, uint *packed_shader);
void pack(Scene *scene, uint *packed_shader);
private:
void add_builtin_attributes();