sandey/blender
1
0
Fork 0
forked from bartvdbraak/blender
Code Pull Requests Activity

Fix Cycles OptiX crash with 3D curves after point cloud changes

Browse Source 
Includes refactoring to reduce the number of bits taken by primitive types,
so they more easily fit in the OptiX limit.
This commit is contained in:
Brecht Van Lommel 2021-12-20 02:52:56 +01:00
parent 5adc06d2d8
commit e2e7f7ea52
28 changed files with 122 additions and 121 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -656,24 +656,24 @@ bool BVHBuild::range_within_max_leaf_size(const BVHRange &range,
for (int i = 0; i < size; i++) { for (int i = 0; i < size; i++) {
const BVHReference &ref = references[range.start() + i]; const BVHReference &ref = references[range.start() + i];
if (ref.prim_type() & PRIMITIVE_ALL_CURVE) { if (ref.prim_type() & PRIMITIVE_CURVE) {
if (ref.prim_type() & PRIMITIVE_ALL_MOTION) { if (ref.prim_type() & PRIMITIVE_MOTION) {
num_motion_curves++; num_motion_curves++;
} }
else { else {
num_curves++; num_curves++;
} }
} }
else if (ref.prim_type() & PRIMITIVE_ALL_TRIANGLE) { else if (ref.prim_type() & PRIMITIVE_TRIANGLE) {
if (ref.prim_type() & PRIMITIVE_ALL_MOTION) { if (ref.prim_type() & PRIMITIVE_MOTION) {
num_motion_triangles++; num_motion_triangles++;
} }
else { else {
num_triangles++; num_triangles++;
} }
} }
else if (ref.prim_type() & PRIMITIVE_ALL_POINT) { else if (ref.prim_type() & PRIMITIVE_POINT) {
if (ref.prim_type() & PRIMITIVE_ALL_MOTION) { if (ref.prim_type() & PRIMITIVE_MOTION) {
num_motion_points++; num_motion_points++;
} }
else { else {
@ -973,7 +973,7 @@ BVHNode *BVHBuild::create_leaf_node(const BVHRange &range, const vector<BVHRefer
for (int i = 0; i < range.size(); i++) { for (int i = 0; i < range.size(); i++) {
const BVHReference &ref = references[range.start() + i]; const BVHReference &ref = references[range.start() + i];
if (ref.prim_index() != -1) { if (ref.prim_index() != -1) {
uint32_t type_index = bitscan((uint32_t)(ref.prim_type() & PRIMITIVE_ALL)); uint32_t type_index = PRIMITIVE_INDEX(ref.prim_type() & PRIMITIVE_ALL);
p_ref[type_index].push_back(ref); p_ref[type_index].push_back(ref);
p_type[type_index].push_back(ref.prim_type()); p_type[type_index].push_back(ref.prim_type());
p_index[type_index].push_back(ref.prim_index()); p_index[type_index].push_back(ref.prim_index());

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

@ -387,7 +387,7 @@ void BVH2::refit_primitives(int start, int end, BoundBox &bbox, uint &visibility
} }
else { else {
/* Primitives. */ /* Primitives. */
if (pack.prim_type[prim] & PRIMITIVE_ALL_CURVE) { if (pack.prim_type[prim] & PRIMITIVE_CURVE) {
/* Curves. */ /* Curves. */
const Hair *hair = static_cast<const Hair *>(ob->get_geometry()); const Hair *hair = static_cast<const Hair *>(ob->get_geometry());
int prim_offset = (params.top_level) ? hair->prim_offset : 0; int prim_offset = (params.top_level) ? hair->prim_offset : 0;
@ -410,7 +410,7 @@ void BVH2::refit_primitives(int start, int end, BoundBox &bbox, uint &visibility
} }
} }
} }
else if (pack.prim_type[prim] & PRIMITIVE_ALL_POINT) { else if (pack.prim_type[prim] & PRIMITIVE_POINT) {
/* Points. */ /* Points. */
const PointCloud *pointcloud = static_cast<const PointCloud *>(ob->get_geometry()); const PointCloud *pointcloud = static_cast<const PointCloud *>(ob->get_geometry());
int prim_offset = (params.top_level) ? pointcloud->prim_offset : 0; int prim_offset = (params.top_level) ? pointcloud->prim_offset : 0;
@ -590,13 +590,7 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
float2 *bvh_prim_time = bvh->pack.prim_time.size() ? &bvh->pack.prim_time[0] : NULL; float2 *bvh_prim_time = bvh->pack.prim_time.size() ? &bvh->pack.prim_time[0] : NULL;
for (size_t i = 0; i < bvh_prim_index_size; i++) { for (size_t i = 0; i < bvh_prim_index_size; i++) {
if (bvh->pack.prim_type[i] & PRIMITIVE_ALL_CURVE) { pack_prim_index[pack_prim_index_offset] = bvh_prim_index[i] + geom_prim_offset;
pack_prim_index[pack_prim_index_offset] = bvh_prim_index[i] + geom_prim_offset;
}
else {
pack_prim_index[pack_prim_index_offset] = bvh_prim_index[i] + geom_prim_offset;
}
pack_prim_type[pack_prim_index_offset] = bvh_prim_type[i]; pack_prim_type[pack_prim_index_offset] = bvh_prim_type[i];
pack_prim_visibility[pack_prim_index_offset] = bvh_prim_visibility[i]; pack_prim_visibility[pack_prim_index_offset] = bvh_prim_visibility[i];
pack_prim_object[pack_prim_index_offset] = 0; // unused for instances pack_prim_object[pack_prim_index_offset] = 0; // unused for instances

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

@ -91,7 +91,7 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
++ctx->num_hits; ++ctx->num_hits;
/* Always use baked shadow transparency for curves. */ /* Always use baked shadow transparency for curves. */
if (current_isect.type & PRIMITIVE_ALL_CURVE) { if (current_isect.type & PRIMITIVE_CURVE) {
ctx->throughput *= intersection_curve_shadow_transparency( ctx->throughput *= intersection_curve_shadow_transparency(
kg, current_isect.object, current_isect.prim, current_isect.u); kg, current_isect.object, current_isect.prim, current_isect.u);

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

@ -535,15 +535,15 @@ void BVHSpatialSplit::split_reference(const BVHBuild &builder,
/* loop over vertices/edges. */ /* loop over vertices/edges. */
const Object *ob = builder.objects[ref.prim_object()]; const Object *ob = builder.objects[ref.prim_object()];
if (ref.prim_type() & PRIMITIVE_ALL_TRIANGLE) { if (ref.prim_type() & PRIMITIVE_TRIANGLE) {
Mesh *mesh = static_cast<Mesh *>(ob->get_geometry()); Mesh *mesh = static_cast<Mesh *>(ob->get_geometry());
split_triangle_reference(ref, mesh, dim, pos, left_bounds, right_bounds); split_triangle_reference(ref, mesh, dim, pos, left_bounds, right_bounds);
} }
else if (ref.prim_type() & PRIMITIVE_ALL_CURVE) { else if (ref.prim_type() & PRIMITIVE_CURVE) {
Hair *hair = static_cast<Hair *>(ob->get_geometry()); Hair *hair = static_cast<Hair *>(ob->get_geometry());
split_curve_reference(ref, hair, dim, pos, left_bounds, right_bounds); split_curve_reference(ref, hair, dim, pos, left_bounds, right_bounds);
} }
else if (ref.prim_type() & PRIMITIVE_ALL_POINT) { else if (ref.prim_type() & PRIMITIVE_POINT) {
PointCloud *pointcloud = static_cast<PointCloud *>(ob->get_geometry()); PointCloud *pointcloud = static_cast<PointCloud *>(ob->get_geometry());
split_point_reference(ref, pointcloud, dim, pos, left_bounds, right_bounds); split_point_reference(ref, pointcloud, dim, pos, left_bounds, right_bounds);
} }

4
intern/cycles/bvh/unaligned.cpp
View File

@ -69,7 +69,7 @@ bool BVHUnaligned::compute_aligned_space(const BVHReference &ref, Transform *ali
const int packed_type = ref.prim_type(); const int packed_type = ref.prim_type();
const int type = (packed_type & PRIMITIVE_ALL); const int type = (packed_type & PRIMITIVE_ALL);
/* No motion blur curves here, we can't fit them to aligned boxes well. */ /* No motion blur curves here, we can't fit them to aligned boxes well. */
if (type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_CURVE_THICK)) { if ((type & PRIMITIVE_CURVE) && !(type & PRIMITIVE_MOTION)) {
const int curve_index = ref.prim_index(); const int curve_index = ref.prim_index();
const int segment = PRIMITIVE_UNPACK_SEGMENT(packed_type); const int segment = PRIMITIVE_UNPACK_SEGMENT(packed_type);
const Hair *hair = static_cast<const Hair *>(object->get_geometry()); const Hair *hair = static_cast<const Hair *>(object->get_geometry());
@ -95,7 +95,7 @@ BoundBox BVHUnaligned::compute_aligned_prim_boundbox(const BVHReference &prim,
const int packed_type = prim.prim_type(); const int packed_type = prim.prim_type();
const int type = (packed_type & PRIMITIVE_ALL); const int type = (packed_type & PRIMITIVE_ALL);
/* No motion blur curves here, we can't fit them to aligned boxes well. */ /* No motion blur curves here, we can't fit them to aligned boxes well. */
if (type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_CURVE_THICK)) { if ((type & PRIMITIVE_CURVE) && !(type & PRIMITIVE_MOTION)) {
const int curve_index = prim.prim_index(); const int curve_index = prim.prim_index();
const int segment = PRIMITIVE_UNPACK_SEGMENT(packed_type); const int segment = PRIMITIVE_UNPACK_SEGMENT(packed_type);
const Hair *hair = static_cast<const Hair *>(object->get_geometry()); const Hair *hair = static_cast<const Hair *>(object->get_geometry());

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

@ -174,7 +174,7 @@ ccl_device_inline
case PRIMITIVE_MOTION_CURVE_THICK: case PRIMITIVE_MOTION_CURVE_THICK:
case PRIMITIVE_CURVE_RIBBON: case PRIMITIVE_CURVE_RIBBON:
case PRIMITIVE_MOTION_CURVE_RIBBON: { case PRIMITIVE_MOTION_CURVE_RIBBON: {
if ((type & PRIMITIVE_ALL_MOTION) && kernel_data.bvh.use_bvh_steps) { if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) {
const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr); const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
if (ray->time < prim_time.x || ray->time > prim_time.y) { if (ray->time < prim_time.x || ray->time > prim_time.y) {
hit = false; hit = false;
@ -203,7 +203,7 @@ ccl_device_inline
#if BVH_FEATURE(BVH_POINTCLOUD) #if BVH_FEATURE(BVH_POINTCLOUD)
case PRIMITIVE_POINT: case PRIMITIVE_POINT:
case PRIMITIVE_MOTION_POINT: { case PRIMITIVE_MOTION_POINT: {
if ((type & PRIMITIVE_ALL_MOTION) && kernel_data.bvh.use_bvh_steps) { if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) {
const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr); const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
if (ray->time < prim_time.x || ray->time > prim_time.y) { if (ray->time < prim_time.x || ray->time > prim_time.y) {
hit = false; hit = false;
@ -255,7 +255,7 @@ ccl_device_inline
bool record_intersection = true; bool record_intersection = true;
/* Always use baked shadow transparency for curves. */ /* Always use baked shadow transparency for curves. */
if (isect.type & PRIMITIVE_ALL_CURVE) { if (isect.type & PRIMITIVE_CURVE) {
*throughput *= intersection_curve_shadow_transparency( *throughput *= intersection_curve_shadow_transparency(
kg, isect.object, isect.prim, isect.u); kg, isect.object, isect.prim, isect.u);

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

@ -166,7 +166,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
case PRIMITIVE_CURVE_RIBBON: case PRIMITIVE_CURVE_RIBBON:
case PRIMITIVE_MOTION_CURVE_RIBBON: { case PRIMITIVE_MOTION_CURVE_RIBBON: {
for (; prim_addr < prim_addr2; prim_addr++) { for (; prim_addr < prim_addr2; prim_addr++) {
if ((type & PRIMITIVE_ALL_MOTION) && kernel_data.bvh.use_bvh_steps) { if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) {
const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr); const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
if (ray->time < prim_time.x || ray->time > prim_time.y) { if (ray->time < prim_time.x || ray->time > prim_time.y) {
continue; continue;
@ -193,7 +193,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
case PRIMITIVE_POINT: case PRIMITIVE_POINT:
case PRIMITIVE_MOTION_POINT: { case PRIMITIVE_MOTION_POINT: {
for (; prim_addr < prim_addr2; prim_addr++) { for (; prim_addr < prim_addr2; prim_addr++) {
if ((type & PRIMITIVE_ALL_MOTION) && kernel_data.bvh.use_bvh_steps) { if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) {
const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr); const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
if (ray->time < prim_time.x || ray->time > prim_time.y) { if (ray->time < prim_time.x || ray->time > prim_time.y) {
continue; continue;

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

@ -118,16 +118,16 @@ ccl_device_forceinline int intersection_get_shader_flags(KernelGlobals kg,
{ {
int shader = 0; int shader = 0;
if (type & PRIMITIVE_ALL_TRIANGLE) { if (type & PRIMITIVE_TRIANGLE) {
shader = kernel_tex_fetch(__tri_shader, prim); shader = kernel_tex_fetch(__tri_shader, prim);
} }
#ifdef __POINTCLOUD__ #ifdef __POINTCLOUD__
else if (type & PRIMITIVE_ALL_POINT) { else if (type & PRIMITIVE_POINT) {
shader = kernel_tex_fetch(__points_shader, prim); shader = kernel_tex_fetch(__points_shader, prim);
} }
#endif #endif
#ifdef __HAIR__ #ifdef __HAIR__
else if (type & PRIMITIVE_ALL_CURVE) { else if (type & PRIMITIVE_CURVE) {
shader = kernel_tex_fetch(__curves, prim).shader_id; shader = kernel_tex_fetch(__curves, prim).shader_id;
} }
#endif #endif
@ -141,16 +141,16 @@ ccl_device_forceinline int intersection_get_shader_from_isect_prim(KernelGlobals
{ {
int shader = 0; int shader = 0;
if (isect_type & PRIMITIVE_ALL_TRIANGLE) { if (isect_type & PRIMITIVE_TRIANGLE) {
shader = kernel_tex_fetch(__tri_shader, prim); shader = kernel_tex_fetch(__tri_shader, prim);
} }
#ifdef __POINTCLOUD__ #ifdef __POINTCLOUD__
else if (isect_type & PRIMITIVE_ALL_POINT) { else if (isect_type & PRIMITIVE_POINT) {
shader = kernel_tex_fetch(__points_shader, prim); shader = kernel_tex_fetch(__points_shader, prim);
} }
#endif #endif
#ifdef __HAIR__ #ifdef __HAIR__
else if (isect_type & PRIMITIVE_ALL_CURVE) { else if (isect_type & PRIMITIVE_CURVE) {
shader = kernel_tex_fetch(__curves, prim).shader_id; shader = kernel_tex_fetch(__curves, prim).shader_id;
} }
#endif #endif

2
intern/cycles/kernel/closure/bsdf.h
View File

@ -124,7 +124,7 @@ ccl_device_inline int bsdf_sample(KernelGlobals kg,
/* For curves use the smooth normal, particularly for ribbons the geometric /* For curves use the smooth normal, particularly for ribbons the geometric
* normal gives too much darkening otherwise. */ * normal gives too much darkening otherwise. */
int label; int label;
const float3 Ng = (sd->type & PRIMITIVE_ALL_CURVE) ? sc->N : sd->Ng; const float3 Ng = (sd->type & PRIMITIVE_CURVE) ? sc->N : sd->Ng;
switch (sc->type) { switch (sc->type) {
case CLOSURE_BSDF_DIFFUSE_ID: case CLOSURE_BSDF_DIFFUSE_ID:

4
intern/cycles/kernel/closure/bsdf_hair_principled.h
View File

@ -213,9 +213,7 @@ ccl_device int bsdf_principled_hair_setup(ccl_private ShaderData *sd,
/* TODO: we convert this value to a cosine later and discard the sign, so /* TODO: we convert this value to a cosine later and discard the sign, so
* we could probably save some operations. */ * we could probably save some operations. */
float h = (sd->type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) ? float h = (sd->type & PRIMITIVE_CURVE_RIBBON) ? -sd->v : dot(cross(sd->Ng, X), Z);
-sd->v :
dot(cross(sd->Ng, X), Z);
kernel_assert(fabsf(h) < 1.0f + 1e-4f); kernel_assert(fabsf(h) < 1.0f + 1e-4f);
kernel_assert(isfinite3_safe(Y)); kernel_assert(isfinite3_safe(Y));

6
intern/cycles/kernel/device/metal/kernel.metal
View File

@ -211,7 +211,7 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
} }
/* Always use baked shadow transparency for curves. */ /* Always use baked shadow transparency for curves. */
if (type & PRIMITIVE_ALL_CURVE) { if (type & PRIMITIVE_CURVE) {
float throughput = payload.throughput; float throughput = payload.throughput;
throughput *= context.intersection_curve_shadow_transparency(nullptr, object, prim, u); throughput *= context.intersection_curve_shadow_transparency(nullptr, object, prim, u);
payload.throughput = throughput; payload.throughput = throughput;
@ -476,7 +476,7 @@ __intersection__curve_ribbon(constant KernelParamsMetal &launch_params_metal [[b
result.continue_search = true; result.continue_search = true;
result.distance = ray_tmax; result.distance = ray_tmax;
if (segment.type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) { if (segment.type & PRIMITIVE_CURVE_RIBBON) {
metalrt_intersection_curve(launch_params_metal, payload, object, segment.prim, segment.type, ray_origin, ray_direction, metalrt_intersection_curve(launch_params_metal, payload, object, segment.prim, segment.type, ray_origin, ray_direction,
# if defined(__METALRT_MOTION__) # if defined(__METALRT_MOTION__)
payload.time, payload.time,
@ -507,7 +507,7 @@ __intersection__curve_ribbon_shadow(constant KernelParamsMetal &launch_params_me
result.continue_search = true; result.continue_search = true;
result.distance = ray_tmax; result.distance = ray_tmax;
if (segment.type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) { if (segment.type & PRIMITIVE_CURVE_RIBBON) {
metalrt_intersection_curve_shadow(launch_params_metal, payload, object, segment.prim, segment.type, ray_origin, ray_direction, metalrt_intersection_curve_shadow(launch_params_metal, payload, object, segment.prim, segment.type, ray_origin, ray_direction,
# if defined(__METALRT_MOTION__) # if defined(__METALRT_MOTION__)
payload.time, payload.time,

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

@ -194,7 +194,7 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
type = kernel_tex_fetch(__objects, object).primitive_type; type = kernel_tex_fetch(__objects, object).primitive_type;
} }
# ifdef __HAIR__ # ifdef __HAIR__
else if (optixGetHitKind() & PRIMITIVE_ALL_CURVE) { else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) {
u = __uint_as_float(optixGetAttribute_0()); u = __uint_as_float(optixGetAttribute_0());
v = __uint_as_float(optixGetAttribute_1()); v = __uint_as_float(optixGetAttribute_1());
@ -234,7 +234,7 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
} }
/* Always use baked shadow transparency for curves. */ /* Always use baked shadow transparency for curves. */
if (type & PRIMITIVE_ALL_CURVE) { if (type & PRIMITIVE_CURVE) {
float throughput = __uint_as_float(optixGetPayload_1()); float throughput = __uint_as_float(optixGetPayload_1());
throughput *= intersection_curve_shadow_transparency(nullptr, object, prim, u); throughput *= intersection_curve_shadow_transparency(nullptr, object, prim, u);
optixSetPayload_1(__float_as_uint(throughput)); optixSetPayload_1(__float_as_uint(throughput));
@ -320,7 +320,7 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
{ {
#ifdef __HAIR__ #ifdef __HAIR__
# if OPTIX_ABI_VERSION < 55 # if OPTIX_ABI_VERSION < 55
if (optixGetHitKind() & PRIMITIVE_ALL_CURVE) { if (optixGetPrimitiveType() == OPTIX_PRIMITIVE_TYPE_ROUND_CUBIC_BSPLINE) {
/* Filter out curve endcaps. */ /* Filter out curve endcaps. */
const float u = __uint_as_float(optixGetAttribute_0()); const float u = __uint_as_float(optixGetAttribute_0());
if (u == 0.0f || u == 1.0f) { if (u == 0.0f || u == 1.0f) {
@ -359,7 +359,7 @@ extern "C" __global__ void __closesthit__kernel_optix_hit()
optixSetPayload_3(prim); optixSetPayload_3(prim);
optixSetPayload_5(kernel_tex_fetch(__objects, object).primitive_type); optixSetPayload_5(kernel_tex_fetch(__objects, object).primitive_type);
} }
else if (optixGetHitKind() & PRIMITIVE_ALL_CURVE) { else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) {
const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim); const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
optixSetPayload_1(optixGetAttribute_0()); /* Same as 'optixGetCurveParameter()' */ optixSetPayload_1(optixGetAttribute_0()); /* Same as 'optixGetCurveParameter()' */
optixSetPayload_2(optixGetAttribute_1()); optixSetPayload_2(optixGetAttribute_1());
@ -406,6 +406,7 @@ ccl_device_inline void optix_intersection_curve(const int prim, const int type)
isect.t *= len; isect.t *= len;
if (curve_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) { if (curve_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) {
static_assert(PRIMITIVE_ALL < 128, "Values >= 128 are reserved for OptiX internal use");
optixReportIntersection(isect.t / len, optixReportIntersection(isect.t / len,
type & PRIMITIVE_ALL, type & PRIMITIVE_ALL,
__float_as_int(isect.u), /* Attribute_0 */ __float_as_int(isect.u), /* Attribute_0 */
@ -418,7 +419,7 @@ extern "C" __global__ void __intersection__curve_ribbon()
const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, optixGetPrimitiveIndex()); const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, optixGetPrimitiveIndex());
const int prim = segment.prim; const int prim = segment.prim;
const int type = segment.type; const int type = segment.type;
if (type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) { if (type & PRIMITIVE_CURVE_RIBBON) {
optix_intersection_curve(prim, type); optix_intersection_curve(prim, type);
} }
} }
@ -460,6 +461,7 @@ extern "C" __global__ void __intersection__point()
} }
if (point_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) { if (point_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) {
static_assert(PRIMITIVE_ALL < 128, "Values >= 128 are reserved for OptiX internal use");
optixReportIntersection(isect.t / len, type & PRIMITIVE_ALL); optixReportIntersection(isect.t / len, type & PRIMITIVE_ALL);
} }
} }

2
intern/cycles/kernel/geom/attribute.h
View File

@ -36,7 +36,7 @@ ccl_device_inline uint subd_triangle_patch(KernelGlobals kg, ccl_private const S
ccl_device_inline uint attribute_primitive_type(KernelGlobals kg, ccl_private const ShaderData *sd) ccl_device_inline uint attribute_primitive_type(KernelGlobals kg, ccl_private const ShaderData *sd)
{ {
if ((sd->type & PRIMITIVE_ALL_TRIANGLE) && subd_triangle_patch(kg, sd) != ~0) { if ((sd->type & PRIMITIVE_TRIANGLE) && subd_triangle_patch(kg, sd) != ~0) {
return ATTR_PRIM_SUBD; return ATTR_PRIM_SUBD;
} }
else { else {

6
intern/cycles/kernel/geom/curve.h
View File

@ -205,14 +205,14 @@ ccl_device float curve_thickness(KernelGlobals kg, ccl_private const ShaderData
{ {
float r = 0.0f; float r = 0.0f;
if (sd->type & PRIMITIVE_ALL_CURVE) { if (sd->type & PRIMITIVE_CURVE) {
KernelCurve curve = kernel_tex_fetch(__curves, sd->prim); KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type); int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
int k1 = k0 + 1; int k1 = k0 + 1;
float4 P_curve[2]; float4 P_curve[2];
if (!(sd->type & PRIMITIVE_ALL_MOTION)) { if (!(sd->type & PRIMITIVE_MOTION)) {
P_curve[0] = kernel_tex_fetch(__curve_keys, k0); P_curve[0] = kernel_tex_fetch(__curve_keys, k0);
P_curve[1] = kernel_tex_fetch(__curve_keys, k1); P_curve[1] = kernel_tex_fetch(__curve_keys, k1);
} }
@ -249,7 +249,7 @@ ccl_device float3 curve_tangent_normal(KernelGlobals kg, ccl_private const Shade
{ {
float3 tgN = make_float3(0.0f, 0.0f, 0.0f); float3 tgN = make_float3(0.0f, 0.0f, 0.0f);
if (sd->type & PRIMITIVE_ALL_CURVE) { if (sd->type & PRIMITIVE_CURVE) {
tgN = -(-sd->I - sd->dPdu * (dot(sd->dPdu, -sd->I) / len_squared(sd->dPdu))); tgN = -(-sd->I - sd->dPdu * (dot(sd->dPdu, -sd->I) / len_squared(sd->dPdu)));
tgN = normalize(tgN); tgN = normalize(tgN);

8
intern/cycles/kernel/geom/curve_intersect.h
View File

@ -635,7 +635,7 @@ ccl_device_forceinline bool curve_intersect(KernelGlobals kg,
float time, float time,
int type) int type)
{ {
const bool is_motion = (type & PRIMITIVE_ALL_MOTION); const bool is_motion = (type & PRIMITIVE_MOTION);
KernelCurve kcurve = kernel_tex_fetch(__curves, prim); KernelCurve kcurve = kernel_tex_fetch(__curves, prim);
@ -655,7 +655,7 @@ ccl_device_forceinline bool curve_intersect(KernelGlobals kg,
motion_curve_keys(kg, object, prim, time, ka, k0, k1, kb, curve); motion_curve_keys(kg, object, prim, time, ka, k0, k1, kb, curve);
} }
if (type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) { if (type & PRIMITIVE_CURVE_RIBBON) {
/* todo: adaptive number of subdivisions could help performance here. */ /* todo: adaptive number of subdivisions could help performance here. */
const int subdivisions = kernel_data.bvh.curve_subdivisions; const int subdivisions = kernel_data.bvh.curve_subdivisions;
if (ribbon_intersect(P, dir, tmax, subdivisions, curve, isect)) { if (ribbon_intersect(P, dir, tmax, subdivisions, curve, isect)) {
@ -704,7 +704,7 @@ ccl_device_inline void curve_shader_setup(KernelGlobals kg,
float4 P_curve[4]; float4 P_curve[4];
if (!(sd->type & PRIMITIVE_ALL_MOTION)) { if (!(sd->type & PRIMITIVE_MOTION)) {
P_curve[0] = kernel_tex_fetch(__curve_keys, ka); P_curve[0] = kernel_tex_fetch(__curve_keys, ka);
P_curve[1] = kernel_tex_fetch(__curve_keys, k0); P_curve[1] = kernel_tex_fetch(__curve_keys, k0);
P_curve[2] = kernel_tex_fetch(__curve_keys, k1); P_curve[2] = kernel_tex_fetch(__curve_keys, k1);
@ -719,7 +719,7 @@ ccl_device_inline void curve_shader_setup(KernelGlobals kg,
const float4 dPdu4 = catmull_rom_basis_derivative(P_curve, sd->u); const float4 dPdu4 = catmull_rom_basis_derivative(P_curve, sd->u);
const float3 dPdu = float4_to_float3(dPdu4); const float3 dPdu = float4_to_float3(dPdu4);
if (sd->type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) { if (sd->type & PRIMITIVE_CURVE_RIBBON) {
/* Rounded smooth normals for ribbons, to approximate thick curve shape. */ /* Rounded smooth normals for ribbons, to approximate thick curve shape. */
const float3 tangent = normalize(dPdu); const float3 tangent = normalize(dPdu);
const float3 bitangent = normalize(cross(tangent, -D)); const float3 bitangent = normalize(cross(tangent, -D));

2
intern/cycles/kernel/geom/point.h
View File

@ -113,7 +113,7 @@ ccl_device float4 point_attribute_float4(KernelGlobals kg,
ccl_device float point_radius(KernelGlobals kg, ccl_private const ShaderData *sd) ccl_device float point_radius(KernelGlobals kg, ccl_private const ShaderData *sd)
{ {
if (sd->type & PRIMITIVE_ALL_POINT) { if (sd->type & PRIMITIVE_POINT) {
return kernel_tex_fetch(__points, sd->prim).w; return kernel_tex_fetch(__points, sd->prim).w;
} }

6
intern/cycles/kernel/geom/point_intersect.h
View File

@ -75,8 +75,8 @@ ccl_device_forceinline bool point_intersect(KernelGlobals kg,
const float time, const float time,
const int type) const int type)
{ {
const float4 point = (type & PRIMITIVE_ALL_MOTION) ? motion_point(kg, object, prim, time) : const float4 point = (type & PRIMITIVE_MOTION) ? motion_point(kg, object, prim, time) :
kernel_tex_fetch(__points, prim); kernel_tex_fetch(__points, prim);
if (!point_intersect_test(point, P, dir, tmax, &isect->t)) { if (!point_intersect_test(point, P, dir, tmax, &isect->t)) {
return false; return false;
@ -105,7 +105,7 @@ ccl_device_inline void point_shader_setup(KernelGlobals kg,
# endif # endif
/* Computer point center for normal. */ /* Computer point center for normal. */
float3 center = float4_to_float3((isect->type & PRIMITIVE_ALL_MOTION) ? float3 center = float4_to_float3((isect->type & PRIMITIVE_MOTION) ?
motion_point(kg, sd->object, sd->prim, sd->time) : motion_point(kg, sd->object, sd->prim, sd->time) :
kernel_tex_fetch(__points, sd->prim)); kernel_tex_fetch(__points, sd->prim));

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

@ -37,19 +37,19 @@ ccl_device_inline float primitive_surface_attribute_float(KernelGlobals kg,
ccl_private float *dx, ccl_private float *dx,
ccl_private float *dy) ccl_private float *dy)
{ {
if (sd->type & PRIMITIVE_ALL_TRIANGLE) { if (sd->type & PRIMITIVE_TRIANGLE) {
if (subd_triangle_patch(kg, sd) == ~0) if (subd_triangle_patch(kg, sd) == ~0)
return triangle_attribute_float(kg, sd, desc, dx, dy); return triangle_attribute_float(kg, sd, desc, dx, dy);
else else
return subd_triangle_attribute_float(kg, sd, desc, dx, dy); return subd_triangle_attribute_float(kg, sd, desc, dx, dy);
} }
#ifdef __HAIR__ #ifdef __HAIR__
else if (sd->type & PRIMITIVE_ALL_CURVE) { else if (sd->type & PRIMITIVE_CURVE) {
return curve_attribute_float(kg, sd, desc, dx, dy); return curve_attribute_float(kg, sd, desc, dx, dy);
} }
#endif #endif
#ifdef __POINTCLOUD__ #ifdef __POINTCLOUD__
else if (sd->type & PRIMITIVE_ALL_POINT) { else if (sd->type & PRIMITIVE_POINT) {
return point_attribute_float(kg, sd, desc, dx, dy); return point_attribute_float(kg, sd, desc, dx, dy);
} }
#endif #endif
@ -68,19 +68,19 @@ ccl_device_inline float2 primitive_surface_attribute_float2(KernelGlobals kg,
ccl_private float2 *dx, ccl_private float2 *dx,
ccl_private float2 *dy) ccl_private float2 *dy)
{ {
if (sd->type & PRIMITIVE_ALL_TRIANGLE) { if (sd->type & PRIMITIVE_TRIANGLE) {
if (subd_triangle_patch(kg, sd) == ~0) if (subd_triangle_patch(kg, sd) == ~0)
return triangle_attribute_float2(kg, sd, desc, dx, dy); return triangle_attribute_float2(kg, sd, desc, dx, dy);
else else
return subd_triangle_attribute_float2(kg, sd, desc, dx, dy); return subd_triangle_attribute_float2(kg, sd, desc, dx, dy);
} }
#ifdef __HAIR__ #ifdef __HAIR__
else if (sd->type & PRIMITIVE_ALL_CURVE) { else if (sd->type & PRIMITIVE_CURVE) {
return curve_attribute_float2(kg, sd, desc, dx, dy); return curve_attribute_float2(kg, sd, desc, dx, dy);
} }
#endif #endif
#ifdef __POINTCLOUD__ #ifdef __POINTCLOUD__
else if (sd->type & PRIMITIVE_ALL_POINT) { else if (sd->type & PRIMITIVE_POINT) {
return point_attribute_float2(kg, sd, desc, dx, dy); return point_attribute_float2(kg, sd, desc, dx, dy);
} }
#endif #endif
@ -99,19 +99,19 @@ ccl_device_inline float3 primitive_surface_attribute_float3(KernelGlobals kg,
ccl_private float3 *dx, ccl_private float3 *dx,
ccl_private float3 *dy) ccl_private float3 *dy)
{ {
if (sd->type & PRIMITIVE_ALL_TRIANGLE) { if (sd->type & PRIMITIVE_TRIANGLE) {
if (subd_triangle_patch(kg, sd) == ~0) if (subd_triangle_patch(kg, sd) == ~0)
return triangle_attribute_float3(kg, sd, desc, dx, dy); return triangle_attribute_float3(kg, sd, desc, dx, dy);
else else
return subd_triangle_attribute_float3(kg, sd, desc, dx, dy); return subd_triangle_attribute_float3(kg, sd, desc, dx, dy);
} }
#ifdef __HAIR__ #ifdef __HAIR__
else if (sd->type & PRIMITIVE_ALL_CURVE) { else if (sd->type & PRIMITIVE_CURVE) {
return curve_attribute_float3(kg, sd, desc, dx, dy); return curve_attribute_float3(kg, sd, desc, dx, dy);
} }
#endif #endif
#ifdef __POINTCLOUD__ #ifdef __POINTCLOUD__
else if (sd->type & PRIMITIVE_ALL_POINT) { else if (sd->type & PRIMITIVE_POINT) {
return point_attribute_float3(kg, sd, desc, dx, dy); return point_attribute_float3(kg, sd, desc, dx, dy);
} }
#endif #endif
@ -130,19 +130,19 @@ ccl_device_forceinline float4 primitive_surface_attribute_float4(KernelGlobals k
ccl_private float4 *dx, ccl_private float4 *dx,
ccl_private float4 *dy) ccl_private float4 *dy)
{ {
if (sd->type & PRIMITIVE_ALL_TRIANGLE) { if (sd->type & PRIMITIVE_TRIANGLE) {
if (subd_triangle_patch(kg, sd) == ~0) if (subd_triangle_patch(kg, sd) == ~0)
return triangle_attribute_float4(kg, sd, desc, dx, dy); return triangle_attribute_float4(kg, sd, desc, dx, dy);
else else
return subd_triangle_attribute_float4(kg, sd, desc, dx, dy); return subd_triangle_attribute_float4(kg, sd, desc, dx, dy);
} }
#ifdef __HAIR__ #ifdef __HAIR__
else if (sd->type & PRIMITIVE_ALL_CURVE) { else if (sd->type & PRIMITIVE_CURVE) {
return curve_attribute_float4(kg, sd, desc, dx, dy); return curve_attribute_float4(kg, sd, desc, dx, dy);
} }
#endif #endif
#ifdef __POINTCLOUD__ #ifdef __POINTCLOUD__
else if (sd->type & PRIMITIVE_ALL_POINT) { else if (sd->type & PRIMITIVE_POINT) {
return point_attribute_float4(kg, sd, desc, dx, dy); return point_attribute_float4(kg, sd, desc, dx, dy);
} }
#endif #endif
@ -246,7 +246,7 @@ ccl_device bool primitive_ptex(KernelGlobals kg,
ccl_device float3 primitive_tangent(KernelGlobals kg, ccl_private ShaderData *sd) ccl_device float3 primitive_tangent(KernelGlobals kg, ccl_private ShaderData *sd)
{ {
#if defined(__HAIR__) || defined(__POINTCLOUD__) #if defined(__HAIR__) || defined(__POINTCLOUD__)
if (sd->type & (PRIMITIVE_ALL_CURVE | PRIMITIVE_ALL_POINT)) if (sd->type & (PRIMITIVE_CURVE | PRIMITIVE_POINT))
# ifdef __DPDU__ # ifdef __DPDU__
return normalize(sd->dPdu); return normalize(sd->dPdu);
# else # else
@ -282,16 +282,16 @@ ccl_device_inline float4 primitive_motion_vector(KernelGlobals kg,
float3 center; float3 center;
#if defined(__HAIR__) || defined(__POINTCLOUD__) #if defined(__HAIR__) || defined(__POINTCLOUD__)
bool is_curve_or_point = sd->type & (PRIMITIVE_ALL_CURVE | PRIMITIVE_ALL_POINT); bool is_curve_or_point = sd->type & (PRIMITIVE_CURVE | PRIMITIVE_POINT);
if (is_curve_or_point) { if (is_curve_or_point) {
center = make_float3(0.0f, 0.0f, 0.0f); center = make_float3(0.0f, 0.0f, 0.0f);
if (sd->type & PRIMITIVE_ALL_CURVE) { if (sd->type & PRIMITIVE_CURVE) {
# if defined(__HAIR__) # if defined(__HAIR__)
center = curve_motion_center_location(kg, sd); center = curve_motion_center_location(kg, sd);
# endif # endif
} }
else if (sd->type & PRIMITIVE_ALL_POINT) { else if (sd->type & PRIMITIVE_POINT) {
# if defined(__POINTCLOUD__) # if defined(__POINTCLOUD__)
center = point_motion_center_location(kg, sd); center = point_motion_center_location(kg, sd);
# endif # endif
@ -331,7 +331,7 @@ ccl_device_inline float4 primitive_motion_vector(KernelGlobals kg,
} }
else else
#endif #endif
if (sd->type & PRIMITIVE_ALL_TRIANGLE) { if (sd->type & PRIMITIVE_TRIANGLE) {
/* Triangle */ /* Triangle */
if (subd_triangle_patch(kg, sd) == ~0) { if (subd_triangle_patch(kg, sd) == ~0) {
motion_pre = triangle_attribute_float3(kg, sd, desc, NULL, NULL); motion_pre = triangle_attribute_float3(kg, sd, desc, NULL, NULL);

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

@ -69,20 +69,20 @@ ccl_device_inline void shader_setup_from_ray(KernelGlobals kg,
sd->I = -ray->D; sd->I = -ray->D;
#ifdef __HAIR__ #ifdef __HAIR__
if (sd->type & PRIMITIVE_ALL_CURVE) { if (sd->type & PRIMITIVE_CURVE) {
/* curve */ /* curve */
curve_shader_setup(kg, sd, ray->P, ray->D, isect->t, isect->object, isect->prim); curve_shader_setup(kg, sd, ray->P, ray->D, isect->t, isect->object, isect->prim);
} }
else else
#endif #endif
#ifdef __POINTCLOUD__ #ifdef __POINTCLOUD__
if (sd->type & PRIMITIVE_ALL_POINT) { if (sd->type & PRIMITIVE_POINT) {
/* point */ /* point */
point_shader_setup(kg, sd, isect, ray); point_shader_setup(kg, sd, isect, ray);
} }
else else
#endif #endif
if (sd->type & PRIMITIVE_TRIANGLE) { if (sd->type == PRIMITIVE_TRIANGLE) {
/* static triangle */ /* static triangle */
float3 Ng = triangle_normal(kg, sd); float3 Ng = triangle_normal(kg, sd);
sd->shader = kernel_tex_fetch(__tri_shader, sd->prim); sd->shader = kernel_tex_fetch(__tri_shader, sd->prim);
@ -201,7 +201,7 @@ ccl_device_inline void shader_setup_from_sample(KernelGlobals kg,
object_dir_transform_auto(kg, sd, &sd->I); object_dir_transform_auto(kg, sd, &sd->I);
} }
if (sd->type & PRIMITIVE_TRIANGLE) { if (sd->type == PRIMITIVE_TRIANGLE) {
/* smooth normal */ /* smooth normal */
if (sd->shader & SHADER_SMOOTH_NORMAL) { if (sd->shader & SHADER_SMOOTH_NORMAL) {
sd->N = triangle_smooth_normal(kg, Ng, sd->prim, sd->u, sd->v); sd->N = triangle_smooth_normal(kg, Ng, sd->prim, sd->u, sd->v);

2
intern/cycles/kernel/integrator/shade_surface.h
View File

@ -82,7 +82,7 @@ ccl_device_forceinline void integrate_surface_emission(KernelGlobals kg,
# ifdef __HAIR__ # ifdef __HAIR__
if (!(path_flag & PATH_RAY_MIS_SKIP) && (sd->flag & SD_USE_MIS) && if (!(path_flag & PATH_RAY_MIS_SKIP) && (sd->flag & SD_USE_MIS) &&
(sd->type & PRIMITIVE_ALL_TRIANGLE)) (sd->type & PRIMITIVE_TRIANGLE))
# else # else
if (!(path_flag & PATH_RAY_MIS_SKIP) && (sd->flag & SD_USE_MIS)) if (!(path_flag & PATH_RAY_MIS_SKIP) && (sd->flag & SD_USE_MIS))
# endif # endif

2
intern/cycles/kernel/light/sample.h
View File

@ -191,7 +191,7 @@ ccl_device_inline float3 shadow_ray_offset(KernelGlobals kg,
float3 Ng = (transmit ? -sd->Ng : sd->Ng); float3 Ng = (transmit ? -sd->Ng : sd->Ng);
float3 P = ray_offset(sd->P, Ng); float3 P = ray_offset(sd->P, Ng);
if ((sd->type & PRIMITIVE_ALL_TRIANGLE) && (sd->shader & SHADER_SMOOTH_NORMAL)) { if ((sd->type & PRIMITIVE_TRIANGLE) && (sd->shader & SHADER_SMOOTH_NORMAL)) {
const float offset_cutoff = const float offset_cutoff =
kernel_tex_fetch(__objects, sd->object).shadow_terminator_geometry_offset; kernel_tex_fetch(__objects, sd->object).shadow_terminator_geometry_offset;
/* Do ray offset (heavy stuff) only for close to be terminated triangles: /* Do ray offset (heavy stuff) only for close to be terminated triangles:

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

@ -960,13 +960,15 @@ bool OSLRenderServices::get_object_standard_attribute(const KernelGlobalsCPU *kg
return set_attribute_int(3, type, derivatives, val); return set_attribute_int(3, type, derivatives, val);
} }
else if ((name == u_geom_trianglevertices || name == u_geom_polyvertices) && else if ((name == u_geom_trianglevertices || name == u_geom_polyvertices) &&
sd->type & PRIMITIVE_ALL_TRIANGLE) { sd->type & PRIMITIVE_TRIANGLE) {
float3 P[3]; float3 P[3];
if (sd->type & PRIMITIVE_TRIANGLE) if (sd->type & PRIMITIVE_MOTION) {
triangle_vertices(kg, sd->prim, P);
else
motion_triangle_vertices(kg, sd->object, sd->prim, sd->time, P); motion_triangle_vertices(kg, sd->object, sd->prim, sd->time, P);
}
else {
triangle_vertices(kg, sd->prim, P);
}
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) { if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
object_position_transform(kg, sd, &P[0]); object_position_transform(kg, sd, &P[0]);
@ -986,7 +988,7 @@ bool OSLRenderServices::get_object_standard_attribute(const KernelGlobalsCPU *kg
} }
/* Hair Attributes */ /* Hair Attributes */
else if (name == u_is_curve) { else if (name == u_is_curve) {
float f = (sd->type & PRIMITIVE_ALL_CURVE) != 0; float f = (sd->type & PRIMITIVE_CURVE) != 0;
return set_attribute_float(f, type, derivatives, val); return set_attribute_float(f, type, derivatives, val);
} }
else if (name == u_curve_thickness) { else if (name == u_curve_thickness) {
@ -999,7 +1001,7 @@ bool OSLRenderServices::get_object_standard_attribute(const KernelGlobalsCPU *kg
} }
/* point attributes */ /* point attributes */
else if (name == u_is_point) { else if (name == u_is_point) {
float f = (sd->type & PRIMITIVE_ALL_POINT) != 0; float f = (sd->type & PRIMITIVE_POINT) != 0;
return set_attribute_float(f, type, derivatives, val); return set_attribute_float(f, type, derivatives, val);
} }
else if (name == u_point_radius) { else if (name == u_point_radius) {
@ -1007,7 +1009,7 @@ bool OSLRenderServices::get_object_standard_attribute(const KernelGlobalsCPU *kg
return set_attribute_float(f, type, derivatives, val); return set_attribute_float(f, type, derivatives, val);
} }
else if (name == u_normal_map_normal) { else if (name == u_normal_map_normal) {
if (sd->type & PRIMITIVE_ALL_TRIANGLE) { if (sd->type & PRIMITIVE_TRIANGLE) {
float3 f = triangle_smooth_normal_unnormalized(kg, sd, sd->Ng, sd->prim, sd->u, sd->v); float3 f = triangle_smooth_normal_unnormalized(kg, sd, sd->Ng, sd->prim, sd->u, sd->v);
return set_attribute_float3(f, type, derivatives, val); return set_attribute_float3(f, type, derivatives, val);
} }

8
intern/cycles/kernel/svm/bevel.h
View File

@ -206,12 +206,12 @@ ccl_device float3 svm_bevel(
for (int hit = 0; hit < num_eval_hits; hit++) { for (int hit = 0; hit < num_eval_hits; hit++) {
/* Quickly retrieve P and Ng without setting up ShaderData. */ /* Quickly retrieve P and Ng without setting up ShaderData. */
float3 hit_P; float3 hit_P;
if (sd->type & PRIMITIVE_TRIANGLE) { if (sd->type == PRIMITIVE_TRIANGLE) {
hit_P = triangle_refine_local( hit_P = triangle_refine_local(
kg, sd, ray.P, ray.D, ray.t, isect.hits[hit].object, isect.hits[hit].prim); kg, sd, ray.P, ray.D, ray.t, isect.hits[hit].object, isect.hits[hit].prim);
} }
# ifdef __OBJECT_MOTION__ # ifdef __OBJECT_MOTION__
else if (sd->type & PRIMITIVE_MOTION_TRIANGLE) { else if (sd->type == PRIMITIVE_MOTION_TRIANGLE) {
float3 verts[3]; float3 verts[3];
motion_triangle_vertices(kg, sd->object, isect.hits[hit].prim, sd->time, verts); motion_triangle_vertices(kg, sd->object, isect.hits[hit].prim, sd->time, verts);
hit_P = motion_triangle_refine_local( hit_P = motion_triangle_refine_local(
@ -236,11 +236,11 @@ ccl_device float3 svm_bevel(
float u = isect.hits[hit].u; float u = isect.hits[hit].u;
float v = isect.hits[hit].v; float v = isect.hits[hit].v;
if (sd->type & PRIMITIVE_TRIANGLE) { if (sd->type == PRIMITIVE_TRIANGLE) {
N = triangle_smooth_normal(kg, N, prim, u, v); N = triangle_smooth_normal(kg, N, prim, u, v);
} }
# ifdef __OBJECT_MOTION__ # ifdef __OBJECT_MOTION__
else if (sd->type & PRIMITIVE_MOTION_TRIANGLE) { else if (sd->type == PRIMITIVE_MOTION_TRIANGLE) {
N = motion_triangle_smooth_normal(kg, N, sd->object, prim, u, v, sd->time); N = motion_triangle_smooth_normal(kg, N, sd->object, prim, u, v, sd->time);
} }
# endif /* __OBJECT_MOTION__ */ # endif /* __OBJECT_MOTION__ */

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

@ -107,7 +107,7 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
} }
float3 N = stack_valid(data_node.x) ? stack_load_float3(stack, data_node.x) : sd->N; float3 N = stack_valid(data_node.x) ? stack_load_float3(stack, data_node.x) : sd->N;
if (!(sd->type & PRIMITIVE_ALL_CURVE)) { if (!(sd->type & PRIMITIVE_CURVE)) {
N = ensure_valid_reflection(sd->Ng, sd->I, N); N = ensure_valid_reflection(sd->Ng, sd->I, N);
} }
@ -191,7 +191,7 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
float3 clearcoat_normal = stack_valid(data_cn_ssr.x) ? float3 clearcoat_normal = stack_valid(data_cn_ssr.x) ?
stack_load_float3(stack, data_cn_ssr.x) : stack_load_float3(stack, data_cn_ssr.x) :
sd->N; sd->N;
if (!(sd->type & PRIMITIVE_ALL_CURVE)) { if (!(sd->type & PRIMITIVE_CURVE)) {
clearcoat_normal = ensure_valid_reflection(sd->Ng, sd->I, clearcoat_normal); clearcoat_normal = ensure_valid_reflection(sd->Ng, sd->I, clearcoat_normal);
} }
float3 subsurface_radius = stack_valid(data_cn_ssr.y) ? float3 subsurface_radius = stack_valid(data_cn_ssr.y) ?
@ -902,7 +902,7 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
if (stack_valid(data_node.y)) { if (stack_valid(data_node.y)) {
bsdf->T = normalize(stack_load_float3(stack, data_node.y)); bsdf->T = normalize(stack_load_float3(stack, data_node.y));
} }
else if (!(sd->type & PRIMITIVE_ALL_CURVE)) { else if (!(sd->type & PRIMITIVE_CURVE)) {
bsdf->T = normalize(sd->dPdv); bsdf->T = normalize(sd->dPdv);
bsdf->offset = 0.0f; bsdf->offset = 0.0f;
} }

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

@ -227,7 +227,7 @@ ccl_device_noinline void svm_node_hair_info(KernelGlobals kg,
switch (type) { switch (type) {
case NODE_INFO_CURVE_IS_STRAND: { case NODE_INFO_CURVE_IS_STRAND: {
data = (sd->type & PRIMITIVE_ALL_CURVE) != 0; data = (sd->type & PRIMITIVE_CURVE) != 0;
stack_store_float(stack, out_offset, data); stack_store_float(stack, out_offset, data);
break; break;
} }

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

@ -291,7 +291,7 @@ ccl_device_noinline void svm_node_normal_map(KernelGlobals kg,
if (space == NODE_NORMAL_MAP_TANGENT) { if (space == NODE_NORMAL_MAP_TANGENT) {
/* tangent space */ /* tangent space */
if (sd->object == OBJECT_NONE || (sd->type & PRIMITIVE_ALL_TRIANGLE) == 0) { if (sd->object == OBJECT_NONE || (sd->type & PRIMITIVE_TRIANGLE) == 0) {
/* Fallback to unperturbed normal. */ /* Fallback to unperturbed normal. */
stack_store_float3(stack, normal_offset, sd->N); stack_store_float3(stack, normal_offset, sd->N);
return; return;

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

@ -43,7 +43,7 @@ ccl_device_inline float wireframe(KernelGlobals kg,
ccl_private float3 *P) ccl_private float3 *P)
{ {
#if defined(__HAIR__) || defined(__POINTCLOUD__) #if defined(__HAIR__) || defined(__POINTCLOUD__)
if (sd->prim != PRIM_NONE && sd->type & PRIMITIVE_ALL_TRIANGLE) if (sd->prim != PRIM_NONE && sd->type & PRIMITIVE_TRIANGLE)
#else #else
if (sd->prim != PRIM_NONE) if (sd->prim != PRIM_NONE)
#endif #endif
@ -54,10 +54,12 @@ ccl_device_inline float wireframe(KernelGlobals kg,
/* Triangles */ /* Triangles */
int np = 3; int np = 3;
if (sd->type & PRIMITIVE_TRIANGLE) if (sd->type & PRIMITIVE_MOTION) {
triangle_vertices(kg, sd->prim, Co);
else
motion_triangle_vertices(kg, sd->object, sd->prim, sd->time, Co); motion_triangle_vertices(kg, sd->object, sd->prim, sd->time, Co);
}
else {
triangle_vertices(kg, sd->prim, Co);
}
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) { if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
object_position_transform(kg, sd, &Co[0]); object_position_transform(kg, sd, &Co[0]);

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

@ -537,31 +537,34 @@ typedef struct Intersection {
typedef enum PrimitiveType { typedef enum PrimitiveType {
PRIMITIVE_NONE = 0, PRIMITIVE_NONE = 0,
PRIMITIVE_TRIANGLE = (1 << 0), PRIMITIVE_TRIANGLE = (1 << 0),
PRIMITIVE_MOTION_TRIANGLE = (1 << 1), PRIMITIVE_CURVE_THICK = (1 << 1),
PRIMITIVE_CURVE_THICK = (1 << 2), PRIMITIVE_CURVE_RIBBON = (1 << 2),
PRIMITIVE_MOTION_CURVE_THICK = (1 << 3), PRIMITIVE_POINT = (1 << 3),
PRIMITIVE_CURVE_RIBBON = (1 << 4), PRIMITIVE_VOLUME = (1 << 4),
PRIMITIVE_MOTION_CURVE_RIBBON = (1 << 5), PRIMITIVE_LAMP = (1 << 5),
PRIMITIVE_POINT = (1 << 6),
PRIMITIVE_MOTION_POINT = (1 << 7),
PRIMITIVE_VOLUME = (1 << 8),
PRIMITIVE_LAMP = (1 << 9),
PRIMITIVE_ALL_TRIANGLE = (PRIMITIVE_TRIANGLE | PRIMITIVE_MOTION_TRIANGLE), PRIMITIVE_MOTION = (1 << 6),
PRIMITIVE_ALL_CURVE = (PRIMITIVE_CURVE_THICK | PRIMITIVE_MOTION_CURVE_THICK | PRIMITIVE_MOTION_TRIANGLE = (PRIMITIVE_TRIANGLE | PRIMITIVE_MOTION),
PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON), PRIMITIVE_MOTION_CURVE_THICK = (PRIMITIVE_CURVE_THICK | PRIMITIVE_MOTION),
PRIMITIVE_ALL_POINT = (PRIMITIVE_POINT | PRIMITIVE_MOTION_POINT), PRIMITIVE_MOTION_CURVE_RIBBON = (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION),
PRIMITIVE_ALL_VOLUME = (PRIMITIVE_VOLUME), PRIMITIVE_MOTION_POINT = (PRIMITIVE_POINT | PRIMITIVE_MOTION),
PRIMITIVE_ALL_MOTION = (PRIMITIVE_MOTION_TRIANGLE | PRIMITIVE_MOTION_CURVE_THICK |
PRIMITIVE_MOTION_CURVE_RIBBON | PRIMITIVE_MOTION_POINT),
PRIMITIVE_ALL = (PRIMITIVE_ALL_TRIANGLE | PRIMITIVE_ALL_CURVE | PRIMITIVE_ALL_VOLUME |
PRIMITIVE_LAMP | PRIMITIVE_ALL_POINT),
PRIMITIVE_NUM = 10, PRIMITIVE_CURVE = (PRIMITIVE_CURVE_THICK | PRIMITIVE_CURVE_RIBBON),
PRIMITIVE_ALL = (PRIMITIVE_TRIANGLE | PRIMITIVE_CURVE | PRIMITIVE_POINT | PRIMITIVE_VOLUME |
PRIMITIVE_LAMP | PRIMITIVE_MOTION),
PRIMITIVE_NUM_SHAPES = 6,
PRIMITIVE_NUM_BITS = PRIMITIVE_NUM_SHAPES + 1, /* All shapes + motion bit. */
PRIMITIVE_NUM = PRIMITIVE_NUM_SHAPES * 2, /* With and without motion. */
} PrimitiveType; } PrimitiveType;
#define PRIMITIVE_PACK_SEGMENT(type, segment) ((segment << PRIMITIVE_NUM) | (type)) /* Convert type to index in range 0..PRIMITIVE_NUM-1. */
#define PRIMITIVE_UNPACK_SEGMENT(type) (type >> PRIMITIVE_NUM) #define PRIMITIVE_INDEX(type) (bitscan((uint32_t)(type)) * 2 + (((type)&PRIMITIVE_MOTION) ? 1 : 0))
/* Pack segment into type value to save space. */
#define PRIMITIVE_PACK_SEGMENT(type, segment) ((segment << PRIMITIVE_NUM_BITS) | (type))
#define PRIMITIVE_UNPACK_SEGMENT(type) (type >> PRIMITIVE_NUM_BITS)
typedef enum CurveShapeType { typedef enum CurveShapeType {
CURVE_RIBBON = 0, CURVE_RIBBON = 0,