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

Includes refactoring to reduce the number of bits taken by primitive types,
so they more easily fit in the OptiX limit.

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,