diff options
Diffstat (limited to 'intern')
28 files changed, 124 insertions, 123 deletions
diff --git a/intern/cycles/bvh/build.cpp b/intern/cycles/bvh/build.cpp index 91198e4e2a2..242595bee4c 100644 --- a/intern/cycles/bvh/build.cpp +++ b/intern/cycles/bvh/build.cpp @@ -656,24 +656,24 @@ bool BVHBuild::range_within_max_leaf_size(const BVHRange &range, for (int i = 0; i < size; i++) { const BVHReference &ref = references[range.start() + i]; - if (ref.prim_type() & PRIMITIVE_ALL_CURVE) { - if (ref.prim_type() & PRIMITIVE_ALL_MOTION) { + if (ref.prim_type() & PRIMITIVE_CURVE) { + if (ref.prim_type() & PRIMITIVE_MOTION) { num_motion_curves++; } else { num_curves++; } } - else if (ref.prim_type() & PRIMITIVE_ALL_TRIANGLE) { - if (ref.prim_type() & PRIMITIVE_ALL_MOTION) { + else if (ref.prim_type() & PRIMITIVE_TRIANGLE) { + if (ref.prim_type() & PRIMITIVE_MOTION) { num_motion_triangles++; } else { num_triangles++; } } - else if (ref.prim_type() & PRIMITIVE_ALL_POINT) { - if (ref.prim_type() & PRIMITIVE_ALL_MOTION) { + else if (ref.prim_type() & PRIMITIVE_POINT) { + if (ref.prim_type() & PRIMITIVE_MOTION) { num_motion_points++; } else { @@ -973,7 +973,7 @@ BVHNode *BVHBuild::create_leaf_node(const BVHRange &range, const vector<BVHRefer for (int i = 0; i < range.size(); i++) { const BVHReference &ref = references[range.start() + i]; 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_type[type_index].push_back(ref.prim_type()); p_index[type_index].push_back(ref.prim_index()); diff --git a/intern/cycles/bvh/bvh2.cpp b/intern/cycles/bvh/bvh2.cpp index 744e7fa9898..f1ea43da1d9 100644 --- a/intern/cycles/bvh/bvh2.cpp +++ b/intern/cycles/bvh/bvh2.cpp @@ -387,7 +387,7 @@ void BVH2::refit_primitives(int start, int end, BoundBox &bbox, uint &visibility } else { /* Primitives. */ - if (pack.prim_type[prim] & PRIMITIVE_ALL_CURVE) { + if (pack.prim_type[prim] & PRIMITIVE_CURVE) { /* Curves. */ const Hair *hair = static_cast<const Hair *>(ob->get_geometry()); 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. */ const PointCloud *pointcloud = static_cast<const PointCloud *>(ob->get_geometry()); 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; 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; - } - else { - 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; pack_prim_type[pack_prim_index_offset] = bvh_prim_type[i]; pack_prim_visibility[pack_prim_index_offset] = bvh_prim_visibility[i]; pack_prim_object[pack_prim_index_offset] = 0; // unused for instances diff --git a/intern/cycles/bvh/embree.cpp b/intern/cycles/bvh/embree.cpp index eab193f45cb..618dd9438d5 100644 --- a/intern/cycles/bvh/embree.cpp +++ b/intern/cycles/bvh/embree.cpp @@ -91,7 +91,7 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args) ++ctx->num_hits; /* 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( kg, current_isect.object, current_isect.prim, current_isect.u); diff --git a/intern/cycles/bvh/split.cpp b/intern/cycles/bvh/split.cpp index 34d12de97c0..e126b6f18bc 100644 --- a/intern/cycles/bvh/split.cpp +++ b/intern/cycles/bvh/split.cpp @@ -535,15 +535,15 @@ void BVHSpatialSplit::split_reference(const BVHBuild &builder, /* loop over vertices/edges. */ 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()); 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()); 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()); split_point_reference(ref, pointcloud, dim, pos, left_bounds, right_bounds); } diff --git a/intern/cycles/bvh/unaligned.cpp b/intern/cycles/bvh/unaligned.cpp index 3c4a600fe58..a8db6efb597 100644 --- a/intern/cycles/bvh/unaligned.cpp +++ b/intern/cycles/bvh/unaligned.cpp @@ -69,7 +69,7 @@ bool BVHUnaligned::compute_aligned_space(const BVHReference &ref, Transform *ali const int packed_type = ref.prim_type(); const int type = (packed_type & PRIMITIVE_ALL); /* 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 segment = PRIMITIVE_UNPACK_SEGMENT(packed_type); 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 type = (packed_type & PRIMITIVE_ALL); /* 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 segment = PRIMITIVE_UNPACK_SEGMENT(packed_type); const Hair *hair = static_cast<const Hair *>(object->get_geometry()); diff --git a/intern/cycles/kernel/bvh/shadow_all.h b/intern/cycles/kernel/bvh/shadow_all.h index caca85aac1a..b0e799675e0 100644 --- a/intern/cycles/kernel/bvh/shadow_all.h +++ b/intern/cycles/kernel/bvh/shadow_all.h @@ -174,7 +174,7 @@ ccl_device_inline case PRIMITIVE_MOTION_CURVE_THICK: case PRIMITIVE_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); if (ray->time < prim_time.x || ray->time > prim_time.y) { hit = false; @@ -203,7 +203,7 @@ ccl_device_inline #if BVH_FEATURE(BVH_POINTCLOUD) case PRIMITIVE_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); if (ray->time < prim_time.x || ray->time > prim_time.y) { hit = false; @@ -255,7 +255,7 @@ ccl_device_inline bool record_intersection = true; /* Always use baked shadow transparency for curves. */ - if (isect.type & PRIMITIVE_ALL_CURVE) { + if (isect.type & PRIMITIVE_CURVE) { *throughput *= intersection_curve_shadow_transparency( kg, isect.object, isect.prim, isect.u); diff --git a/intern/cycles/kernel/bvh/traversal.h b/intern/cycles/kernel/bvh/traversal.h index 180f19d11c5..e4bff1a8a80 100644 --- a/intern/cycles/kernel/bvh/traversal.h +++ b/intern/cycles/kernel/bvh/traversal.h @@ -166,7 +166,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg, case PRIMITIVE_CURVE_RIBBON: case PRIMITIVE_MOTION_CURVE_RIBBON: { 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); if (ray->time < prim_time.x || ray->time > prim_time.y) { continue; @@ -193,7 +193,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg, case PRIMITIVE_POINT: case PRIMITIVE_MOTION_POINT: { for (; prim_addr < prim_addr2; prim_addr++) { - if ((type & PRIMITIVE_ALL_MOTION) && kernel_data.bvh.use_bvh_steps) { + if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) { const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr); if (ray->time < prim_time.x || ray->time > prim_time.y) { continue; diff --git a/intern/cycles/kernel/bvh/util.h b/intern/cycles/kernel/bvh/util.h index 57593e42a88..bd79c6e19c6 100644 --- a/intern/cycles/kernel/bvh/util.h +++ b/intern/cycles/kernel/bvh/util.h @@ -118,16 +118,16 @@ ccl_device_forceinline int intersection_get_shader_flags(KernelGlobals kg, { int shader = 0; - if (type & PRIMITIVE_ALL_TRIANGLE) { + if (type & PRIMITIVE_TRIANGLE) { shader = kernel_tex_fetch(__tri_shader, prim); } #ifdef __POINTCLOUD__ - else if (type & PRIMITIVE_ALL_POINT) { + else if (type & PRIMITIVE_POINT) { shader = kernel_tex_fetch(__points_shader, prim); } #endif #ifdef __HAIR__ - else if (type & PRIMITIVE_ALL_CURVE) { + else if (type & PRIMITIVE_CURVE) { shader = kernel_tex_fetch(__curves, prim).shader_id; } #endif @@ -141,16 +141,16 @@ ccl_device_forceinline int intersection_get_shader_from_isect_prim(KernelGlobals { int shader = 0; - if (isect_type & PRIMITIVE_ALL_TRIANGLE) { + if (isect_type & PRIMITIVE_TRIANGLE) { shader = kernel_tex_fetch(__tri_shader, prim); } #ifdef __POINTCLOUD__ - else if (isect_type & PRIMITIVE_ALL_POINT) { + else if (isect_type & PRIMITIVE_POINT) { shader = kernel_tex_fetch(__points_shader, prim); } #endif #ifdef __HAIR__ - else if (isect_type & PRIMITIVE_ALL_CURVE) { + else if (isect_type & PRIMITIVE_CURVE) { shader = kernel_tex_fetch(__curves, prim).shader_id; } #endif diff --git a/intern/cycles/kernel/closure/bsdf.h b/intern/cycles/kernel/closure/bsdf.h index f0ce45d1c2c..2c8ef858270 100644 --- a/intern/cycles/kernel/closure/bsdf.h +++ b/intern/cycles/kernel/closure/bsdf.h @@ -124,7 +124,7 @@ ccl_device_inline int bsdf_sample(KernelGlobals kg, /* For curves use the smooth normal, particularly for ribbons the geometric * normal gives too much darkening otherwise. */ 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) { case CLOSURE_BSDF_DIFFUSE_ID: diff --git a/intern/cycles/kernel/closure/bsdf_hair_principled.h b/intern/cycles/kernel/closure/bsdf_hair_principled.h index f55ea0f6a2e..c68314889f1 100644 --- a/intern/cycles/kernel/closure/bsdf_hair_principled.h +++ b/intern/cycles/kernel/closure/bsdf_hair_principled.h @@ -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 * we could probably save some operations. */ - float h = (sd->type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) ? - -sd->v : - dot(cross(sd->Ng, X), Z); + float h = (sd->type & PRIMITIVE_CURVE_RIBBON) ? -sd->v : dot(cross(sd->Ng, X), Z); kernel_assert(fabsf(h) < 1.0f + 1e-4f); kernel_assert(isfinite3_safe(Y)); diff --git a/intern/cycles/kernel/device/metal/kernel.metal b/intern/cycles/kernel/device/metal/kernel.metal index 27dc1f44c6f..deb7dafe55e 100644 --- a/intern/cycles/kernel/device/metal/kernel.metal +++ b/intern/cycles/kernel/device/metal/kernel.metal @@ -211,7 +211,7 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal, } /* Always use baked shadow transparency for curves. */ - if (type & PRIMITIVE_ALL_CURVE) { + if (type & PRIMITIVE_CURVE) { float throughput = payload.throughput; throughput *= context.intersection_curve_shadow_transparency(nullptr, object, prim, u); payload.throughput = throughput; @@ -476,7 +476,7 @@ __intersection__curve_ribbon(constant KernelParamsMetal &launch_params_metal [[b result.continue_search = true; 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, # if defined(__METALRT_MOTION__) payload.time, @@ -507,7 +507,7 @@ __intersection__curve_ribbon_shadow(constant KernelParamsMetal &launch_params_me result.continue_search = true; 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, # if defined(__METALRT_MOTION__) payload.time, diff --git a/intern/cycles/kernel/device/optix/kernel.cu b/intern/cycles/kernel/device/optix/kernel.cu index c639dc87f35..aa210b31a95 100644 --- a/intern/cycles/kernel/device/optix/kernel.cu +++ b/intern/cycles/kernel/device/optix/kernel.cu @@ -194,7 +194,7 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() type = kernel_tex_fetch(__objects, object).primitive_type; } # ifdef __HAIR__ - else if (optixGetHitKind() & PRIMITIVE_ALL_CURVE) { + else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) { u = __uint_as_float(optixGetAttribute_0()); 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. */ - if (type & PRIMITIVE_ALL_CURVE) { + if (type & PRIMITIVE_CURVE) { float throughput = __uint_as_float(optixGetPayload_1()); throughput *= intersection_curve_shadow_transparency(nullptr, object, prim, u); optixSetPayload_1(__float_as_uint(throughput)); @@ -320,7 +320,7 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test() { #ifdef __HAIR__ # if OPTIX_ABI_VERSION < 55 - if (optixGetHitKind() & PRIMITIVE_ALL_CURVE) { + if (optixGetPrimitiveType() == OPTIX_PRIMITIVE_TYPE_ROUND_CUBIC_BSPLINE) { /* Filter out curve endcaps. */ const float u = __uint_as_float(optixGetAttribute_0()); if (u == 0.0f || u == 1.0f) { @@ -359,7 +359,7 @@ extern "C" __global__ void __closesthit__kernel_optix_hit() optixSetPayload_3(prim); 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); optixSetPayload_1(optixGetAttribute_0()); /* Same as 'optixGetCurveParameter()' */ optixSetPayload_2(optixGetAttribute_1()); @@ -406,6 +406,7 @@ ccl_device_inline void optix_intersection_curve(const int prim, const int type) isect.t *= len; 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, type & PRIMITIVE_ALL, __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 int prim = segment.prim; const int type = segment.type; - if (type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) { + if (type & PRIMITIVE_CURVE_RIBBON) { 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)) { + static_assert(PRIMITIVE_ALL < 128, "Values >= 128 are reserved for OptiX internal use"); optixReportIntersection(isect.t / len, type & PRIMITIVE_ALL); } } diff --git a/intern/cycles/kernel/geom/attribute.h b/intern/cycles/kernel/geom/attribute.h index a7ac2bd926f..8b3524beb5d 100644 --- a/intern/cycles/kernel/geom/attribute.h +++ b/intern/cycles/kernel/geom/attribute.h @@ -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) { - 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; } else { diff --git a/intern/cycles/kernel/geom/curve.h b/intern/cycles/kernel/geom/curve.h index 4b6eecf9640..8a63f01643b 100644 --- a/intern/cycles/kernel/geom/curve.h +++ b/intern/cycles/kernel/geom/curve.h @@ -205,14 +205,14 @@ ccl_device float curve_thickness(KernelGlobals kg, ccl_private const ShaderData { float r = 0.0f; - if (sd->type & PRIMITIVE_ALL_CURVE) { + if (sd->type & PRIMITIVE_CURVE) { KernelCurve curve = kernel_tex_fetch(__curves, sd->prim); int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type); int k1 = k0 + 1; 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[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); - 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 = normalize(tgN); diff --git a/intern/cycles/kernel/geom/curve_intersect.h b/intern/cycles/kernel/geom/curve_intersect.h index fb0b80b281f..99f4a452d6d 100644 --- a/intern/cycles/kernel/geom/curve_intersect.h +++ b/intern/cycles/kernel/geom/curve_intersect.h @@ -635,7 +635,7 @@ ccl_device_forceinline bool curve_intersect(KernelGlobals kg, float time, int type) { - const bool is_motion = (type & PRIMITIVE_ALL_MOTION); + const bool is_motion = (type & PRIMITIVE_MOTION); 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); } - if (type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) { + if (type & PRIMITIVE_CURVE_RIBBON) { /* todo: adaptive number of subdivisions could help performance here. */ const int subdivisions = kernel_data.bvh.curve_subdivisions; 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]; - if (!(sd->type & PRIMITIVE_ALL_MOTION)) { + if (!(sd->type & PRIMITIVE_MOTION)) { P_curve[0] = kernel_tex_fetch(__curve_keys, ka); P_curve[1] = kernel_tex_fetch(__curve_keys, k0); 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 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. */ const float3 tangent = normalize(dPdu); const float3 bitangent = normalize(cross(tangent, -D)); diff --git a/intern/cycles/kernel/geom/point.h b/intern/cycles/kernel/geom/point.h index 6d46b934f09..52a1e77d71a 100644 --- a/intern/cycles/kernel/geom/point.h +++ b/intern/cycles/kernel/geom/point.h @@ -113,7 +113,7 @@ ccl_device float4 point_attribute_float4(KernelGlobals kg, 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; } diff --git a/intern/cycles/kernel/geom/point_intersect.h b/intern/cycles/kernel/geom/point_intersect.h index 7abb0453ae5..b4d96ade7be 100644 --- a/intern/cycles/kernel/geom/point_intersect.h +++ b/intern/cycles/kernel/geom/point_intersect.h @@ -75,8 +75,8 @@ ccl_device_forceinline bool point_intersect(KernelGlobals kg, const float time, const int type) { - const float4 point = (type & PRIMITIVE_ALL_MOTION) ? motion_point(kg, object, prim, time) : - kernel_tex_fetch(__points, prim); + const float4 point = (type & PRIMITIVE_MOTION) ? motion_point(kg, object, prim, time) : + kernel_tex_fetch(__points, prim); if (!point_intersect_test(point, P, dir, tmax, &isect->t)) { return false; @@ -105,7 +105,7 @@ ccl_device_inline void point_shader_setup(KernelGlobals kg, # endif /* 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) : kernel_tex_fetch(__points, sd->prim)); diff --git a/intern/cycles/kernel/geom/primitive.h b/intern/cycles/kernel/geom/primitive.h index 9416385638c..2e8e0cda6f1 100644 --- a/intern/cycles/kernel/geom/primitive.h +++ b/intern/cycles/kernel/geom/primitive.h @@ -37,19 +37,19 @@ ccl_device_inline float primitive_surface_attribute_float(KernelGlobals kg, ccl_private float *dx, ccl_private float *dy) { - if (sd->type & PRIMITIVE_ALL_TRIANGLE) { + if (sd->type & PRIMITIVE_TRIANGLE) { if (subd_triangle_patch(kg, sd) == ~0) return triangle_attribute_float(kg, sd, desc, dx, dy); else return subd_triangle_attribute_float(kg, sd, desc, dx, dy); } #ifdef __HAIR__ - else if (sd->type & PRIMITIVE_ALL_CURVE) { + else if (sd->type & PRIMITIVE_CURVE) { return curve_attribute_float(kg, sd, desc, dx, dy); } #endif #ifdef __POINTCLOUD__ - else if (sd->type & PRIMITIVE_ALL_POINT) { + else if (sd->type & PRIMITIVE_POINT) { return point_attribute_float(kg, sd, desc, dx, dy); } #endif @@ -68,19 +68,19 @@ ccl_device_inline float2 primitive_surface_attribute_float2(KernelGlobals kg, ccl_private float2 *dx, ccl_private float2 *dy) { - if (sd->type & PRIMITIVE_ALL_TRIANGLE) { + if (sd->type & PRIMITIVE_TRIANGLE) { if (subd_triangle_patch(kg, sd) == ~0) return triangle_attribute_float2(kg, sd, desc, dx, dy); else return subd_triangle_attribute_float2(kg, sd, desc, dx, dy); } #ifdef __HAIR__ - else if (sd->type & PRIMITIVE_ALL_CURVE) { + else if (sd->type & PRIMITIVE_CURVE) { return curve_attribute_float2(kg, sd, desc, dx, dy); } #endif #ifdef __POINTCLOUD__ - else if (sd->type & PRIMITIVE_ALL_POINT) { + else if (sd->type & PRIMITIVE_POINT) { return point_attribute_float2(kg, sd, desc, dx, dy); } #endif @@ -99,19 +99,19 @@ ccl_device_inline float3 primitive_surface_attribute_float3(KernelGlobals kg, ccl_private float3 *dx, ccl_private float3 *dy) { - if (sd->type & PRIMITIVE_ALL_TRIANGLE) { + if (sd->type & PRIMITIVE_TRIANGLE) { if (subd_triangle_patch(kg, sd) == ~0) return triangle_attribute_float3(kg, sd, desc, dx, dy); else return subd_triangle_attribute_float3(kg, sd, desc, dx, dy); } #ifdef __HAIR__ - else if (sd->type & PRIMITIVE_ALL_CURVE) { + else if (sd->type & PRIMITIVE_CURVE) { return curve_attribute_float3(kg, sd, desc, dx, dy); } #endif #ifdef __POINTCLOUD__ - else if (sd->type & PRIMITIVE_ALL_POINT) { + else if (sd->type & PRIMITIVE_POINT) { return point_attribute_float3(kg, sd, desc, dx, dy); } #endif @@ -130,19 +130,19 @@ ccl_device_forceinline float4 primitive_surface_attribute_float4(KernelGlobals k ccl_private float4 *dx, ccl_private float4 *dy) { - if (sd->type & PRIMITIVE_ALL_TRIANGLE) { + if (sd->type & PRIMITIVE_TRIANGLE) { if (subd_triangle_patch(kg, sd) == ~0) return triangle_attribute_float4(kg, sd, desc, dx, dy); else return subd_triangle_attribute_float4(kg, sd, desc, dx, dy); } #ifdef __HAIR__ - else if (sd->type & PRIMITIVE_ALL_CURVE) { + else if (sd->type & PRIMITIVE_CURVE) { return curve_attribute_float4(kg, sd, desc, dx, dy); } #endif #ifdef __POINTCLOUD__ - else if (sd->type & PRIMITIVE_ALL_POINT) { + else if (sd->type & PRIMITIVE_POINT) { return point_attribute_float4(kg, sd, desc, dx, dy); } #endif @@ -246,7 +246,7 @@ ccl_device bool primitive_ptex(KernelGlobals kg, ccl_device float3 primitive_tangent(KernelGlobals kg, ccl_private ShaderData *sd) { #if defined(__HAIR__) || defined(__POINTCLOUD__) - if (sd->type & (PRIMITIVE_ALL_CURVE | PRIMITIVE_ALL_POINT)) + if (sd->type & (PRIMITIVE_CURVE | PRIMITIVE_POINT)) # ifdef __DPDU__ return normalize(sd->dPdu); # else @@ -282,16 +282,16 @@ ccl_device_inline float4 primitive_motion_vector(KernelGlobals kg, float3 center; #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) { center = make_float3(0.0f, 0.0f, 0.0f); - if (sd->type & PRIMITIVE_ALL_CURVE) { + if (sd->type & PRIMITIVE_CURVE) { # if defined(__HAIR__) center = curve_motion_center_location(kg, sd); # endif } - else if (sd->type & PRIMITIVE_ALL_POINT) { + else if (sd->type & PRIMITIVE_POINT) { # if defined(__POINTCLOUD__) center = point_motion_center_location(kg, sd); # endif @@ -331,7 +331,7 @@ ccl_device_inline float4 primitive_motion_vector(KernelGlobals kg, } else #endif - if (sd->type & PRIMITIVE_ALL_TRIANGLE) { + if (sd->type & PRIMITIVE_TRIANGLE) { /* Triangle */ if (subd_triangle_patch(kg, sd) == ~0) { motion_pre = triangle_attribute_float3(kg, sd, desc, NULL, NULL); diff --git a/intern/cycles/kernel/geom/shader_data.h b/intern/cycles/kernel/geom/shader_data.h index d3932545ef6..569393c306c 100644 --- a/intern/cycles/kernel/geom/shader_data.h +++ b/intern/cycles/kernel/geom/shader_data.h @@ -69,20 +69,20 @@ ccl_device_inline void shader_setup_from_ray(KernelGlobals kg, sd->I = -ray->D; #ifdef __HAIR__ - if (sd->type & PRIMITIVE_ALL_CURVE) { + if (sd->type & PRIMITIVE_CURVE) { /* curve */ curve_shader_setup(kg, sd, ray->P, ray->D, isect->t, isect->object, isect->prim); } else #endif #ifdef __POINTCLOUD__ - if (sd->type & PRIMITIVE_ALL_POINT) { + if (sd->type & PRIMITIVE_POINT) { /* point */ point_shader_setup(kg, sd, isect, ray); } else #endif - if (sd->type & PRIMITIVE_TRIANGLE) { + if (sd->type == PRIMITIVE_TRIANGLE) { /* static triangle */ float3 Ng = triangle_normal(kg, sd); 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); } - if (sd->type & PRIMITIVE_TRIANGLE) { + if (sd->type == PRIMITIVE_TRIANGLE) { /* smooth normal */ if (sd->shader & SHADER_SMOOTH_NORMAL) { sd->N = triangle_smooth_normal(kg, Ng, sd->prim, sd->u, sd->v); diff --git a/intern/cycles/kernel/integrator/shade_surface.h b/intern/cycles/kernel/integrator/shade_surface.h index c9c586f5ae4..3d5b65458c7 100644 --- a/intern/cycles/kernel/integrator/shade_surface.h +++ b/intern/cycles/kernel/integrator/shade_surface.h @@ -82,7 +82,7 @@ ccl_device_forceinline void integrate_surface_emission(KernelGlobals kg, # ifdef __HAIR__ if (!(path_flag & PATH_RAY_MIS_SKIP) && (sd->flag & SD_USE_MIS) && - (sd->type & PRIMITIVE_ALL_TRIANGLE)) + (sd->type & PRIMITIVE_TRIANGLE)) # else if (!(path_flag & PATH_RAY_MIS_SKIP) && (sd->flag & SD_USE_MIS)) # endif diff --git a/intern/cycles/kernel/light/sample.h b/intern/cycles/kernel/light/sample.h index 83f6d733d3a..b6662c7f6b3 100644 --- a/intern/cycles/kernel/light/sample.h +++ b/intern/cycles/kernel/light/sample.h @@ -191,7 +191,7 @@ ccl_device_inline float3 shadow_ray_offset(KernelGlobals kg, float3 Ng = (transmit ? -sd->Ng : sd->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 = kernel_tex_fetch(__objects, sd->object).shadow_terminator_geometry_offset; /* Do ray offset (heavy stuff) only for close to be terminated triangles: diff --git a/intern/cycles/kernel/osl/services.cpp b/intern/cycles/kernel/osl/services.cpp index 4007005dee7..a79fc323a13 100644 --- a/intern/cycles/kernel/osl/services.cpp +++ b/intern/cycles/kernel/osl/services.cpp @@ -960,13 +960,15 @@ bool OSLRenderServices::get_object_standard_attribute(const KernelGlobalsCPU *kg return set_attribute_int(3, type, derivatives, val); } else if ((name == u_geom_trianglevertices || name == u_geom_polyvertices) && - sd->type & PRIMITIVE_ALL_TRIANGLE) { + sd->type & PRIMITIVE_TRIANGLE) { float3 P[3]; - if (sd->type & PRIMITIVE_TRIANGLE) - triangle_vertices(kg, sd->prim, P); - else + if (sd->type & PRIMITIVE_MOTION) { 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)) { object_position_transform(kg, sd, &P[0]); @@ -986,7 +988,7 @@ bool OSLRenderServices::get_object_standard_attribute(const KernelGlobalsCPU *kg } /* Hair Attributes */ 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); } else if (name == u_curve_thickness) { @@ -999,7 +1001,7 @@ bool OSLRenderServices::get_object_standard_attribute(const KernelGlobalsCPU *kg } /* point attributes */ 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); } 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); } 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); return set_attribute_float3(f, type, derivatives, val); } diff --git a/intern/cycles/kernel/svm/bevel.h b/intern/cycles/kernel/svm/bevel.h index 6799489514f..46dfb6631da 100644 --- a/intern/cycles/kernel/svm/bevel.h +++ b/intern/cycles/kernel/svm/bevel.h @@ -206,12 +206,12 @@ ccl_device float3 svm_bevel( for (int hit = 0; hit < num_eval_hits; hit++) { /* Quickly retrieve P and Ng without setting up ShaderData. */ float3 hit_P; - if (sd->type & PRIMITIVE_TRIANGLE) { + if (sd->type == PRIMITIVE_TRIANGLE) { hit_P = triangle_refine_local( kg, sd, ray.P, ray.D, ray.t, isect.hits[hit].object, isect.hits[hit].prim); } # ifdef __OBJECT_MOTION__ - else if (sd->type & PRIMITIVE_MOTION_TRIANGLE) { + else if (sd->type == PRIMITIVE_MOTION_TRIANGLE) { float3 verts[3]; motion_triangle_vertices(kg, sd->object, isect.hits[hit].prim, sd->time, verts); hit_P = motion_triangle_refine_local( @@ -236,11 +236,11 @@ ccl_device float3 svm_bevel( float u = isect.hits[hit].u; 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); } # 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); } # endif /* __OBJECT_MOTION__ */ diff --git a/intern/cycles/kernel/svm/closure.h b/intern/cycles/kernel/svm/closure.h index 71952e9e0f8..2ca22d58191 100644 --- a/intern/cycles/kernel/svm/closure.h +++ b/intern/cycles/kernel/svm/closure.h @@ -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; - if (!(sd->type & PRIMITIVE_ALL_CURVE)) { + if (!(sd->type & PRIMITIVE_CURVE)) { 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) ? stack_load_float3(stack, data_cn_ssr.x) : sd->N; - if (!(sd->type & PRIMITIVE_ALL_CURVE)) { + if (!(sd->type & PRIMITIVE_CURVE)) { clearcoat_normal = ensure_valid_reflection(sd->Ng, sd->I, clearcoat_normal); } 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)) { 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->offset = 0.0f; } diff --git a/intern/cycles/kernel/svm/geometry.h b/intern/cycles/kernel/svm/geometry.h index 772942e0c08..2bac58b0aa2 100644 --- a/intern/cycles/kernel/svm/geometry.h +++ b/intern/cycles/kernel/svm/geometry.h @@ -227,7 +227,7 @@ ccl_device_noinline void svm_node_hair_info(KernelGlobals kg, switch (type) { 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); break; } diff --git a/intern/cycles/kernel/svm/tex_coord.h b/intern/cycles/kernel/svm/tex_coord.h index 5e0debc968a..4b12a0065a6 100644 --- a/intern/cycles/kernel/svm/tex_coord.h +++ b/intern/cycles/kernel/svm/tex_coord.h @@ -291,7 +291,7 @@ ccl_device_noinline void svm_node_normal_map(KernelGlobals kg, if (space == NODE_NORMAL_MAP_TANGENT) { /* 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. */ stack_store_float3(stack, normal_offset, sd->N); return; diff --git a/intern/cycles/kernel/svm/wireframe.h b/intern/cycles/kernel/svm/wireframe.h index 645dc59f22e..2c91ab3dedf 100644 --- a/intern/cycles/kernel/svm/wireframe.h +++ b/intern/cycles/kernel/svm/wireframe.h @@ -43,7 +43,7 @@ ccl_device_inline float wireframe(KernelGlobals kg, ccl_private float3 *P) { #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 if (sd->prim != PRIM_NONE) #endif @@ -54,10 +54,12 @@ ccl_device_inline float wireframe(KernelGlobals kg, /* Triangles */ int np = 3; - if (sd->type & PRIMITIVE_TRIANGLE) - triangle_vertices(kg, sd->prim, Co); - else + if (sd->type & PRIMITIVE_MOTION) { 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)) { object_position_transform(kg, sd, &Co[0]); diff --git a/intern/cycles/kernel/types.h b/intern/cycles/kernel/types.h index 855cc97edbf..20abea37649 100644 --- a/intern/cycles/kernel/types.h +++ b/intern/cycles/kernel/types.h @@ -537,31 +537,34 @@ typedef struct Intersection { typedef enum PrimitiveType { PRIMITIVE_NONE = 0, PRIMITIVE_TRIANGLE = (1 << 0), - PRIMITIVE_MOTION_TRIANGLE = (1 << 1), - PRIMITIVE_CURVE_THICK = (1 << 2), - PRIMITIVE_MOTION_CURVE_THICK = (1 << 3), - PRIMITIVE_CURVE_RIBBON = (1 << 4), - PRIMITIVE_MOTION_CURVE_RIBBON = (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_ALL_CURVE = (PRIMITIVE_CURVE_THICK | PRIMITIVE_MOTION_CURVE_THICK | - PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON), - PRIMITIVE_ALL_POINT = (PRIMITIVE_POINT | PRIMITIVE_MOTION_POINT), - PRIMITIVE_ALL_VOLUME = (PRIMITIVE_VOLUME), - PRIMITIVE_ALL_MOTION = (PRIMITIVE_MOTION_TRIANGLE | PRIMITIVE_MOTION_CURVE_THICK | - PRIMITIVE_MOTION_CURVE_RIBBON | PRIMITIVE_MOTION_POINT), - PRIMITIVE_ALL = (PRIMITIVE_ALL_TRIANGLE | PRIMITIVE_ALL_CURVE | PRIMITIVE_ALL_VOLUME | - PRIMITIVE_LAMP | PRIMITIVE_ALL_POINT), - - PRIMITIVE_NUM = 10, + PRIMITIVE_CURVE_THICK = (1 << 1), + PRIMITIVE_CURVE_RIBBON = (1 << 2), + PRIMITIVE_POINT = (1 << 3), + PRIMITIVE_VOLUME = (1 << 4), + PRIMITIVE_LAMP = (1 << 5), + + PRIMITIVE_MOTION = (1 << 6), + PRIMITIVE_MOTION_TRIANGLE = (PRIMITIVE_TRIANGLE | PRIMITIVE_MOTION), + PRIMITIVE_MOTION_CURVE_THICK = (PRIMITIVE_CURVE_THICK | PRIMITIVE_MOTION), + PRIMITIVE_MOTION_CURVE_RIBBON = (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION), + PRIMITIVE_MOTION_POINT = (PRIMITIVE_POINT | PRIMITIVE_MOTION), + + 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; -#define PRIMITIVE_PACK_SEGMENT(type, segment) ((segment << PRIMITIVE_NUM) | (type)) -#define PRIMITIVE_UNPACK_SEGMENT(type) (type >> PRIMITIVE_NUM) +/* Convert type to index in range 0..PRIMITIVE_NUM-1. */ +#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 { CURVE_RIBBON = 0, |