Cycles: remove support for rendering hair as triangle and lines

Triangles were very memory intensive. The only reason they were not removed yet
is that they gave more accurate results, but there will be an accurate 3D curve
primitive added for this.

Line rendering was always poor quality since the ends do not match up. To keep CPU
and GPU compatibility we just remove them entirely. They could be brought back if
an Embree compatible implementation is added, but it's not clear to me that there
is a use case for these that we'd consider important.

Ref T73778

Reviewers: #cycles

Subscribers:

19 files changed, 171 insertions, 1157 deletions

diff --git a/intern/cycles/blender/addon/properties.py b/intern/cycles/blender/addon/properties.py
index 1635afab210..29ef1814294 100644
--- a/intern/cycles/blender/addon/properties.py
+++ b/intern/cycles/blender/addon/properties.py
@@ -78,20 +78,9 @@ enum_panorama_types = (
     ('MIRRORBALL', "Mirror Ball", "Uses the mirror ball mapping"),
 )
 
-enum_curve_primitives = (
-    ('TRIANGLES', "Triangles", "Create triangle geometry around strands"),
-    ('LINE_SEGMENTS', "Line Segments", "Use line segment primitives"),
-    ('CURVE_SEGMENTS', "Curve Segments", "Use segmented cardinal curve primitives"),
-)
-
-enum_triangle_curves = (
-    ('CAMERA_TRIANGLES', "Planes", "Create individual triangles forming planes that face camera"),
-    ('TESSELLATED_TRIANGLES', "Tessellated", "Create mesh surrounding each strand"),
-)
-
 enum_curve_shape = (
-    ('RIBBONS', "Ribbons", "Ignore thickness of each strand"),
-    ('THICK', "Thick", "Use thickness of strand when rendering"),
+    ('RIBBONS', "Ribbons", "Ignore thickness of each hair"),
+    ('THICK', "Thick", "Use thickness of hair when rendering"),
 )
 
 enum_tile_order = (
@@ -1241,12 +1230,6 @@ class CyclesObjectSettings(bpy.types.PropertyGroup):
 
 class CyclesCurveRenderSettings(bpy.types.PropertyGroup):
 
-    primitive: EnumProperty(
-        name="Primitive",
-        description="Type of primitive used for hair rendering",
-        items=enum_curve_primitives,
-        default='LINE_SEGMENTS',
-    )
     shape: EnumProperty(
         name="Shape",
         description="Form of hair",
@@ -1255,7 +1238,7 @@ class CyclesCurveRenderSettings(bpy.types.PropertyGroup):
     )
     cull_backfacing: BoolProperty(
         name="Cull Back-faces",
-        description="Do not test the back-face of each strand",
+        description="Do not test the back-face of each hair",
         default=True,
     )
     use_curves: BoolProperty(
@@ -1263,12 +1246,6 @@ class CyclesCurveRenderSettings(bpy.types.PropertyGroup):
         description="Activate Cycles hair rendering for particle system",
         default=True,
     )
-    resolution: IntProperty(
-        name="Resolution",
-        description="Resolution of generated mesh",
-        min=3, max=64,
-        default=3,
-    )
     subdivisions: IntProperty(
         name="Subdivisions",
         description="Number of subdivisions used in Cardinal curve intersection (power of 2)",
diff --git a/intern/cycles/blender/addon/ui.py b/intern/cycles/blender/addon/ui.py
index 78a44881743..b6d19b1d4a7 100644
--- a/intern/cycles/blender/addon/ui.py
+++ b/intern/cycles/blender/addon/ui.py
@@ -406,14 +406,11 @@ class CYCLES_RENDER_PT_hair(CyclesButtonsPanel, Panel):
 
         col = layout.column()
         col.prop(ccscene, "shape", text="Shape")
-        if not (ccscene.primitive in {'CURVE_SEGMENTS', 'LINE_SEGMENTS'} and ccscene.shape == 'RIBBONS'):
-            col.prop(ccscene, "cull_backfacing", text="Cull back-faces")
-        col.prop(ccscene, "primitive", text="Primitive")
-
-        if ccscene.primitive == 'TRIANGLES' and ccscene.shape == 'THICK':
-            col.prop(ccscene, "resolution", text="Resolution")
-        elif ccscene.primitive == 'CURVE_SEGMENTS':
+        if ccscene.shape == 'RIBBONS':
+            # TODO: use for embree
             col.prop(ccscene, "subdivisions", text="Curve subdivisions")
+        else:
+            col.prop(ccscene, "cull_backfacing", text="Cull back-faces")
 
 
 class CYCLES_RENDER_PT_volumes(CyclesButtonsPanel, Panel):
diff --git a/intern/cycles/blender/blender_curves.cpp b/intern/cycles/blender/blender_curves.cpp
index 847a43c5f34..43fe653fda4 100644
--- a/intern/cycles/blender/blender_curves.cpp
+++ b/intern/cycles/blender/blender_curves.cpp
@@ -18,7 +18,6 @@
 #include "render/camera.h"
 #include "render/curves.h"
 #include "render/hair.h"
-#include "render/mesh.h"
 #include "render/object.h"
 #include "render/scene.h"
 
@@ -39,27 +38,6 @@ ParticleCurveData::~ParticleCurveData()
 {
 }
 
-static void interp_weights(float t, float data[4])
-{
-  /* Cardinal curve interpolation */
-  float t2 = t * t;
-  float t3 = t2 * t;
-  float fc = 0.71f;
-
-  data[0] = -fc * t3 + 2.0f * fc * t2 - fc * t;
-  data[1] = (2.0f - fc) * t3 + (fc - 3.0f) * t2 + 1.0f;
-  data[2] = (fc - 2.0f) * t3 + (3.0f - 2.0f * fc) * t2 + fc * t;
-  data[3] = fc * t3 - fc * t2;
-}
-
-static void curveinterp_v3_v3v3v3v3(
-    float3 *p, float3 *v1, float3 *v2, float3 *v3, float3 *v4, const float w[4])
-{
-  p->x = v1->x * w[0] + v2->x * w[1] + v3->x * w[2] + v4->x * w[3];
-  p->y = v1->y * w[0] + v2->y * w[1] + v3->y * w[2] + v4->y * w[3];
-  p->z = v1->z * w[0] + v2->z * w[1] + v3->z * w[2] + v4->z * w[3];
-}
-
 static float shaperadius(float shape, float root, float tip, float time)
 {
   assert(time >= 0.0f);
@@ -77,43 +55,13 @@ static float shaperadius(float shape, float root, float tip, float time)
 
 /* curve functions */
 
-static void InterpolateKeySegments(
-    int seg, int segno, int key, int curve, float3 *keyloc, float *time, ParticleCurveData *CData)
-{
-  float3 ckey_loc1 = CData->curvekey_co[key];
-  float3 ckey_loc2 = ckey_loc1;
-  float3 ckey_loc3 = CData->curvekey_co[key + 1];
-  float3 ckey_loc4 = ckey_loc3;
-
-  if (key > CData->curve_firstkey[curve])
-    ckey_loc1 = CData->curvekey_co[key - 1];
-
-  if (key < CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 2)
-    ckey_loc4 = CData->curvekey_co[key + 2];
-
-  float time1 = CData->curvekey_time[key] / CData->curve_length[curve];
-  float time2 = CData->curvekey_time[key + 1] / CData->curve_length[curve];
-
-  float dfra = (time2 - time1) / (float)segno;
-
-  if (time)
-    *time = (dfra * seg) + time1;
-
-  float t[4];
-
-  interp_weights((float)seg / (float)segno, t);
-
-  if (keyloc)
-    curveinterp_v3_v3v3v3v3(keyloc, &ckey_loc1, &ckey_loc2, &ckey_loc3, &ckey_loc4, t);
-}
-
 static bool ObtainCacheParticleData(
-    Geometry *geom, BL::Mesh *b_mesh, BL::Object *b_ob, ParticleCurveData *CData, bool background)
+    Hair *hair, BL::Mesh *b_mesh, BL::Object *b_ob, ParticleCurveData *CData, bool background)
 {
   int curvenum = 0;
   int keyno = 0;
 
-  if (!(geom && b_mesh && b_ob && CData))
+  if (!(hair && b_mesh && b_ob && CData))
     return false;
 
   Transform tfm = get_transform(b_ob->matrix_world());
@@ -129,7 +77,7 @@ static bool ObtainCacheParticleData(
 
       if ((b_part.render_type() == BL::ParticleSettings::render_type_PATH) &&
           (b_part.type() == BL::ParticleSettings::type_HAIR)) {
-        int shader = clamp(b_part.material() - 1, 0, geom->used_shaders.size() - 1);
+        int shader = clamp(b_part.material() - 1, 0, hair->used_shaders.size() - 1);
         int display_step = background ? b_part.render_step() : b_part.display_step();
         int totparts = b_psys.particles.length();
         int totchild = background ? b_psys.child_particles.length() :
@@ -203,14 +151,14 @@ static bool ObtainCacheParticleData(
   return true;
 }
 
-static bool ObtainCacheParticleUV(Geometry *geom,
+static bool ObtainCacheParticleUV(Hair *hair,
                                   BL::Mesh *b_mesh,
                                   BL::Object *b_ob,
                                   ParticleCurveData *CData,
                                   bool background,
                                   int uv_num)
 {
-  if (!(geom && b_mesh && b_ob && CData))
+  if (!(hair && b_mesh && b_ob && CData))
     return false;
 
   CData->curve_uv.clear();
@@ -266,14 +214,14 @@ static bool ObtainCacheParticleUV(Geometry *geom,
   return true;
 }
 
-static bool ObtainCacheParticleVcol(Geometry *geom,
+static bool ObtainCacheParticleVcol(Hair *hair,
                                     BL::Mesh *b_mesh,
                                     BL::Object *b_ob,
                                     ParticleCurveData *CData,
                                     bool background,
                                     int vcol_num)
 {
-  if (!(geom && b_mesh && b_ob && CData))
+  if (!(hair && b_mesh && b_ob && CData))
     return false;
 
   CData->curve_vcol.clear();
@@ -329,272 +277,6 @@ static bool ObtainCacheParticleVcol(Geometry *geom,
   return true;
 }
 
-static void ExportCurveTrianglePlanes(Mesh *mesh,
-                                      ParticleCurveData *CData,
-                                      float3 RotCam,
-                                      bool is_ortho)
-{
-  int vertexno = mesh->verts.size();
-  int vertexindex = vertexno;
-  int numverts = 0, numtris = 0;
-
-  /* compute and reserve size of arrays */
-  for (int sys = 0; sys < CData->psys_firstcurve.size(); sys++) {
-    for (int curve = CData->psys_firstcurve[sys];
-         curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys];
-         curve++) {
-      numverts += 2 + (CData->curve_keynum[curve] - 1) * 2;
-      numtris += (CData->curve_keynum[curve] - 1) * 2;
-    }
-  }
-
-  mesh->reserve_mesh(mesh->verts.size() + numverts, mesh->num_triangles() + numtris);
-
-  /* actually export */
-  for (int sys = 0; sys < CData->psys_firstcurve.size(); sys++) {
-    for (int curve = CData->psys_firstcurve[sys];
-         curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys];
-         curve++) {
-      float3 xbasis;
-      float3 v1;
-      float time = 0.0f;
-      float3 ickey_loc = CData->curvekey_co[CData->curve_firstkey[curve]];
-      float radius = shaperadius(
-          CData->psys_shape[sys], CData->psys_rootradius[sys], CData->psys_tipradius[sys], 0.0f);
-      v1 = CData->curvekey_co[CData->curve_firstkey[curve] + 1] -
-           CData->curvekey_co[CData->curve_firstkey[curve]];
-      if (is_ortho)
-        xbasis = normalize(cross(RotCam, v1));
-      else
-        xbasis = normalize(cross(RotCam - ickey_loc, v1));
-      float3 ickey_loc_shfl = ickey_loc - radius * xbasis;
-      float3 ickey_loc_shfr = ickey_loc + radius * xbasis;
-      mesh->add_vertex(ickey_loc_shfl);
-      mesh->add_vertex(ickey_loc_shfr);
-      vertexindex += 2;
-
-      for (int curvekey = CData->curve_firstkey[curve] + 1;
-           curvekey < CData->curve_firstkey[curve] + CData->curve_keynum[curve];
-           curvekey++) {
-        ickey_loc = CData->curvekey_co[curvekey];
-
-        if (curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1)
-          v1 = CData->curvekey_co[curvekey] -
-               CData->curvekey_co[max(curvekey - 1, CData->curve_firstkey[curve])];
-        else
-          v1 = CData->curvekey_co[curvekey + 1] - CData->curvekey_co[curvekey - 1];
-
-        time = CData->curvekey_time[curvekey] / CData->curve_length[curve];
-        radius = shaperadius(
-            CData->psys_shape[sys], CData->psys_rootradius[sys], CData->psys_tipradius[sys], time);
-
-        if (curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1)
-          radius = shaperadius(CData->psys_shape[sys],
-                               CData->psys_rootradius[sys],
-                               CData->psys_tipradius[sys],
-                               0.95f);
-
-        if (CData->psys_closetip[sys] &&
-            (curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1))
-          radius = shaperadius(CData->psys_shape[sys], CData->psys_rootradius[sys], 0.0f, 0.95f);
-
-        if (is_ortho)
-          xbasis = normalize(cross(RotCam, v1));
-        else
-          xbasis = normalize(cross(RotCam - ickey_loc, v1));
-        float3 ickey_loc_shfl = ickey_loc - radius * xbasis;
-        float3 ickey_loc_shfr = ickey_loc + radius * xbasis;
-        mesh->add_vertex(ickey_loc_shfl);
-        mesh->add_vertex(ickey_loc_shfr);
-        mesh->add_triangle(
-            vertexindex - 2, vertexindex, vertexindex - 1, CData->psys_shader[sys], true);
-        mesh->add_triangle(
-            vertexindex + 1, vertexindex - 1, vertexindex, CData->psys_shader[sys], true);
-        vertexindex += 2;
-      }
-    }
-  }
-
-  mesh->resize_mesh(mesh->verts.size(), mesh->num_triangles());
-  mesh->attributes.remove(ATTR_STD_VERTEX_NORMAL);
-  mesh->attributes.remove(ATTR_STD_FACE_NORMAL);
-  mesh->add_face_normals();
-  mesh->add_vertex_normals();
-  mesh->attributes.remove(ATTR_STD_FACE_NORMAL);
-
-  /* texture coords still needed */
-}
-
-static void ExportCurveTriangleGeometry(Mesh *mesh, ParticleCurveData *CData, int resolution)
-{
-  int vertexno = mesh->verts.size();
-  int vertexindex = vertexno;
-  int numverts = 0, numtris = 0;
-
-  /* compute and reserve size of arrays */
-  for (int sys = 0; sys < CData->psys_firstcurve.size(); sys++) {
-    for (int curve = CData->psys_firstcurve[sys];
-         curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys];
-         curve++) {
-      numverts += (CData->curve_keynum[curve] - 1) * resolution + resolution;
-      numtris += (CData->curve_keynum[curve] - 1) * 2 * resolution;
-    }
-  }
-
-  mesh->reserve_mesh(mesh->verts.size() + numverts, mesh->num_triangles() + numtris);
-
-  /* actually export */
-  for (int sys = 0; sys < CData->psys_firstcurve.size(); sys++) {
-    for (int curve = CData->psys_firstcurve[sys];
-         curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys];
-         curve++) {
-      float3 firstxbasis = cross(make_float3(1.0f, 0.0f, 0.0f),
-                                 CData->curvekey_co[CData->curve_firstkey[curve] + 1] -
-                                     CData->curvekey_co[CData->curve_firstkey[curve]]);
-      if (!is_zero(firstxbasis))
-        firstxbasis = normalize(firstxbasis);
-      else
-        firstxbasis = normalize(cross(make_float3(0.0f, 1.0f, 0.0f),
-                                      CData->curvekey_co[CData->curve_firstkey[curve] + 1] -
-                                          CData->curvekey_co[CData->curve_firstkey[curve]]));
-
-      for (int curvekey = CData->curve_firstkey[curve];
-           curvekey < CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1;
-           curvekey++) {
-        float3 xbasis = firstxbasis;
-        float3 v1;
-        float3 v2;
-
-        if (curvekey == CData->curve_firstkey[curve]) {
-          v1 = CData->curvekey_co[min(
-                   curvekey + 2, CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1)] -
-               CData->curvekey_co[curvekey + 1];
-          v2 = CData->curvekey_co[curvekey + 1] - CData->curvekey_co[curvekey];
-        }
-        else if (curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1) {
-          v1 = CData->curvekey_co[curvekey] - CData->curvekey_co[curvekey - 1];
-          v2 = CData->curvekey_co[curvekey - 1] -
-               CData->curvekey_co[max(curvekey - 2, CData->curve_firstkey[curve])];
-        }
-        else {
-          v1 = CData->curvekey_co[curvekey + 1] - CData->curvekey_co[curvekey];
-          v2 = CData->curvekey_co[curvekey] - CData->curvekey_co[curvekey - 1];
-        }
-
-        xbasis = cross(v1, v2);
-
-        if (len_squared(xbasis) >= 0.05f * len_squared(v1) * len_squared(v2)) {
-          firstxbasis = normalize(xbasis);
-          break;
-        }
-      }
-
-      for (int curvekey = CData->curve_firstkey[curve];
-           curvekey < CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1;
-           curvekey++) {
-        int subv = 1;
-        float3 xbasis;
-        float3 ybasis;
-        float3 v1;
-        float3 v2;
-
-        if (curvekey == CData->curve_firstkey[curve]) {
-          subv = 0;
-          v1 = CData->curvekey_co[min(
-                   curvekey + 2, CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1)] -
-               CData->curvekey_co[curvekey + 1];
-          v2 = CData->curvekey_co[curvekey + 1] - CData->curvekey_co[curvekey];
-        }
-        else if (curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1) {
-          v1 = CData->curvekey_co[curvekey] - CData->curvekey_co[curvekey - 1];
-          v2 = CData->curvekey_co[curvekey - 1] -
-               CData->curvekey_co[max(curvekey - 2, CData->curve_firstkey[curve])];
-        }
-        else {
-          v1 = CData->curvekey_co[curvekey + 1] - CData->curvekey_co[curvekey];
-          v2 = CData->curvekey_co[curvekey] - CData->curvekey_co[curvekey - 1];
-        }
-
-        xbasis = cross(v1, v2);
-
-        if (len_squared(xbasis) >= 0.05f * len_squared(v1) * len_squared(v2)) {
-          xbasis = normalize(xbasis);
-          firstxbasis = xbasis;
-        }
-        else
-          xbasis = firstxbasis;
-
-        ybasis = normalize(cross(xbasis, v2));
-
-        for (; subv <= 1; subv++) {
-          float3 ickey_loc = make_float3(0.0f, 0.0f, 0.0f);
-          float time = 0.0f;
-
-          InterpolateKeySegments(subv, 1, curvekey, curve, &ickey_loc, &time, CData);
-
-          float radius = shaperadius(CData->psys_shape[sys],
-                                     CData->psys_rootradius[sys],
-                                     CData->psys_tipradius[sys],
-                                     time);
-
-          if ((curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 2) &&
-              (subv == 1))
-            radius = shaperadius(CData->psys_shape[sys],
-                                 CData->psys_rootradius[sys],
-                                 CData->psys_tipradius[sys],
-                                 0.95f);
-
-          if (CData->psys_closetip[sys] && (subv == 1) &&
-              (curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 2))
-            radius = shaperadius(CData->psys_shape[sys], CData->psys_rootradius[sys], 0.0f, 0.95f);
-
-          float angle = M_2PI_F / (float)resolution;
-          for (int section = 0; section < resolution; section++) {
-            float3 ickey_loc_shf = ickey_loc + radius * (cosf(angle * section) * xbasis +
-                                                         sinf(angle * section) * ybasis);
-            mesh->add_vertex(ickey_loc_shf);
-          }
-
-          if (subv != 0) {
-            for (int section = 0; section < resolution - 1; section++) {
-              mesh->add_triangle(vertexindex - resolution + section,
-                                 vertexindex + section,
-                                 vertexindex - resolution + section + 1,
-                                 CData->psys_shader[sys],
-                                 true);
-              mesh->add_triangle(vertexindex + section + 1,
-                                 vertexindex - resolution + section + 1,
-                                 vertexindex + section,
-                                 CData->psys_shader[sys],
-                                 true);
-            }
-            mesh->add_triangle(vertexindex - 1,
-                               vertexindex + resolution - 1,
-                               vertexindex - resolution,
-                               CData->psys_shader[sys],
-                               true);
-            mesh->add_triangle(vertexindex,
-                               vertexindex - resolution,
-                               vertexindex + resolution - 1,
-                               CData->psys_shader[sys],
-                               true);
-          }
-          vertexindex += resolution;
-        }
-      }
-    }
-  }
-
-  mesh->resize_mesh(mesh->verts.size(), mesh->num_triangles());
-  mesh->attributes.remove(ATTR_STD_VERTEX_NORMAL);
-  mesh->attributes.remove(ATTR_STD_FACE_NORMAL);
-  mesh->add_face_normals();
-  mesh->add_vertex_normals();
-  mesh->attributes.remove(ATTR_STD_FACE_NORMAL);
-
-  /* texture coords still needed */
-}
-
 static void ExportCurveSegments(Scene *scene, Hair *hair, ParticleCurveData *CData)
 {
   int num_keys = 0;
@@ -823,78 +505,6 @@ static void ExportCurveSegmentsMotion(Hair *hair, ParticleCurveData *CData, int
   }
 }
 
-static void ExportCurveTriangleUV(ParticleCurveData *CData, int resol, float2 *uvdata)
-{
-  if (uvdata == NULL)
-    return;
-  int vertexindex = 0;
-
-  for (int sys = 0; sys < CData->psys_firstcurve.size(); sys++) {
-    for (int curve = CData->psys_firstcurve[sys];
-         curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys];
-         curve++) {
-      for (int curvekey = CData->curve_firstkey[curve];
-           curvekey < CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1;
-           curvekey++) {
-        for (int section = 0; section < resol; section++) {
-          uvdata[vertexindex] = CData->curve_uv[curve];
-          vertexindex++;
-          uvdata[vertexindex] = CData->curve_uv[curve];
-          vertexindex++;
-          uvdata[vertexindex] = CData->curve_uv[curve];
-          vertexindex++;
-          uvdata[vertexindex] = CData->curve_uv[curve];
-          vertexindex++;
-          uvdata[vertexindex] = CData->curve_uv[curve];
-          vertexindex++;
-          uvdata[vertexindex] = CData->curve_uv[curve];
-          vertexindex++;
-        }
-      }
-    }
-  }
-}
-
-static void ExportCurveTriangleVcol(ParticleCurveData *CData, int resol, uchar4 *cdata)
-{
-  if (cdata == NULL)
-    return;
-
-  int vertexindex = 0;
-
-  for (int sys = 0; sys < CData->psys_firstcurve.size(); sys++) {
-    for (int curve = CData->psys_firstcurve[sys];
-         curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys];
-         curve++) {
-      for (int curvekey = CData->curve_firstkey[curve];
-           curvekey < CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1;
-           curvekey++) {
-        for (int section = 0; section < resol; section++) {
-          /* Encode vertex color using the sRGB curve. */
-          cdata[vertexindex] = color_float_to_byte(
-              color_srgb_to_linear_v3(CData->curve_vcol[curve]));
-          vertexindex++;
-          cdata[vertexindex] = color_float_to_byte(
-              color_srgb_to_linear_v3(CData->curve_vcol[curve]));
-          vertexindex++;
-          cdata[vertexindex] = color_float_to_byte(
-              color_srgb_to_linear_v3(CData->curve_vcol[curve]));
-          vertexindex++;
-          cdata[vertexindex] = color_float_to_byte(
-              color_srgb_to_linear_v3(CData->curve_vcol[curve]));
-          vertexindex++;
-          cdata[vertexindex] = color_float_to_byte(
-              color_srgb_to_linear_v3(CData->curve_vcol[curve]));
-          vertexindex++;
-          cdata[vertexindex] = color_float_to_byte(
-              color_srgb_to_linear_v3(CData->curve_vcol[curve]));
-          vertexindex++;
-        }
-      }
-    }
-  }
-}
-
 /* Hair Curve Sync */
 
 void BlenderSync::sync_curve_settings(BL::Depsgraph &b_depsgraph)
@@ -906,48 +516,11 @@ void BlenderSync::sync_curve_settings(BL::Depsgraph &b_depsgraph)
 
   curve_system_manager->use_curves = get_boolean(csscene, "use_curves");
 
-  curve_system_manager->primitive = (CurvePrimitiveType)get_enum(
-      csscene, "primitive", CURVE_NUM_PRIMITIVE_TYPES, CURVE_LINE_SEGMENTS);
   curve_system_manager->curve_shape = (CurveShapeType)get_enum(
       csscene, "shape", CURVE_NUM_SHAPE_TYPES, CURVE_THICK);
-  curve_system_manager->resolution = get_int(csscene, "resolution");
   curve_system_manager->subdivisions = get_int(csscene, "subdivisions");
   curve_system_manager->use_backfacing = !get_boolean(csscene, "cull_backfacing");
 
-  /* Triangles */
-  if (curve_system_manager->primitive == CURVE_TRIANGLES) {
-    /* camera facing planes */
-    if (curve_system_manager->curve_shape == CURVE_RIBBON) {
-      curve_system_manager->triangle_method = CURVE_CAMERA_TRIANGLES;
-      curve_system_manager->resolution = 1;
-    }
-    else if (curve_system_manager->curve_shape == CURVE_THICK) {
-      curve_system_manager->triangle_method = CURVE_TESSELATED_TRIANGLES;
-    }
-  }
-  /* Line Segments */
-  else if (curve_system_manager->primitive == CURVE_LINE_SEGMENTS) {
-    if (curve_system_manager->curve_shape == CURVE_RIBBON) {
-      /* tangent shading */
-      curve_system_manager->line_method = CURVE_UNCORRECTED;
-      curve_system_manager->use_encasing = true;
-      curve_system_manager->use_backfacing = false;
-      curve_system_manager->use_tangent_normal_geometry = true;
-    }
-    else if (curve_system_manager->curve_shape == CURVE_THICK) {
-      curve_system_manager->line_method = CURVE_ACCURATE;
-      curve_system_manager->use_encasing = false;
-      curve_system_manager->use_tangent_normal_geometry = false;
-    }
-  }
-  /* Curve Segments */
-  else if (curve_system_manager->primitive == CURVE_SEGMENTS) {
-    if (curve_system_manager->curve_shape == CURVE_RIBBON) {
-      curve_system_manager->primitive = CURVE_RIBBONS;
-      curve_system_manager->use_backfacing = false;
-    }
-  }
-
   if (curve_system_manager->modified_mesh(prev_curve_system_manager)) {
     BL::Depsgraph::objects_iterator b_ob;
 
@@ -994,78 +567,38 @@ bool BlenderSync::object_has_particle_hair(BL::Object b_ob)
 
 /* Old particle hair. */
 void BlenderSync::sync_particle_hair(
-    Geometry *geom, BL::Mesh &b_mesh, BL::Object &b_ob, bool motion, int motion_step)
+    Hair *hair, BL::Mesh &b_mesh, BL::Object &b_ob, bool motion, int motion_step)
 {
-  Hair *hair = (geom->type == Geometry::HAIR) ? static_cast<Hair *>(geom) : NULL;
-  Mesh *mesh = (geom->type == Geometry::MESH) ? static_cast<Mesh *>(geom) : NULL;
-
   /* obtain general settings */
   if (b_ob.mode() == b_ob.mode_PARTICLE_EDIT || b_ob.mode() == b_ob.mode_EDIT) {
     return;
   }
 
-  const int triangle_method = scene->curve_system_manager->triangle_method;
-  const int resolution = scene->curve_system_manager->resolution;
-  int used_res = 1;
-
   /* extract particle hair data - should be combined with connecting to mesh later*/
 
   ParticleCurveData CData;
 
-  ObtainCacheParticleData(geom, &b_mesh, &b_ob, &CData, !preview);
-
-  /* add hair geometry to mesh */
-  if (mesh) {
-    if (triangle_method == CURVE_CAMERA_TRIANGLES) {
-      /* obtain camera parameters */
-      float3 RotCam;
-      Camera *camera = scene->camera;
-      Transform &ctfm = camera->matrix;
-      if (camera->type == CAMERA_ORTHOGRAPHIC) {
-        RotCam = -make_float3(ctfm.x.z, ctfm.y.z, ctfm.z.z);
-      }
-      else {
-        Transform tfm = get_transform(b_ob.matrix_world());
-        Transform itfm = transform_quick_inverse(tfm);
-        RotCam = transform_point(&itfm, make_float3(ctfm.x.w, ctfm.y.w, ctfm.z.w));
-      }
-      bool is_ortho = camera->type == CAMERA_ORTHOGRAPHIC;
-      ExportCurveTrianglePlanes(mesh, &CData, RotCam, is_ortho);
-    }
-    else {
-      ExportCurveTriangleGeometry(mesh, &CData, resolution);
-      used_res = resolution;
-    }
-  }
-  else {
-    if (motion)
-      ExportCurveSegmentsMotion(hair, &CData, motion_step);
-    else
-      ExportCurveSegments(scene, hair, &CData);
-  }
+  ObtainCacheParticleData(hair, &b_mesh, &b_ob, &CData, !preview);
+
+  /* add hair geometry */
+  if (motion)
+    ExportCurveSegmentsMotion(hair, &CData, motion_step);
+  else
+    ExportCurveSegments(scene, hair, &CData);
 
   /* generated coordinates from first key. we should ideally get this from
    * blender to handle deforming objects */
   if (!motion) {
-    if (geom->need_attribute(scene, ATTR_STD_GENERATED)) {
+    if (hair->need_attribute(scene, ATTR_STD_GENERATED)) {
       float3 loc, size;
       mesh_texture_space(b_mesh, loc, size);
 
-      if (mesh) {
-        Attribute *attr_generated = mesh->attributes.add(ATTR_STD_GENERATED);
-        float3 *generated = attr_generated->data_float3();
+      Attribute *attr_generated = hair->attributes.add(ATTR_STD_GENERATED);
+      float3 *generated = attr_generated->data_float3();
 
-        for (size_t i = 0; i < mesh->verts.size(); i++)
-          generated[i] = mesh->verts[i] * size - loc;
-      }
-      else {
-        Attribute *attr_generated = hair->attributes.add(ATTR_STD_GENERATED);
-        float3 *generated = attr_generated->data_float3();
-
-        for (size_t i = 0; i < hair->num_curves(); i++) {
-          float3 co = hair->curve_keys[hair->get_curve(i).first_key];
-          generated[i] = co * size - loc;
-        }
+      for (size_t i = 0; i < hair->num_curves(); i++) {
+        float3 co = hair->curve_keys[hair->get_curve(i).first_key];
+        generated[i] = co * size - loc;
       }
     }
   }
@@ -1076,32 +609,22 @@ void BlenderSync::sync_particle_hair(
     int vcol_num = 0;
 
     for (b_mesh.vertex_colors.begin(l); l != b_mesh.vertex_colors.end(); ++l, vcol_num++) {
-      if (!geom->need_attribute(scene, ustring(l->name().c_str())))
+      if (!hair->need_attribute(scene, ustring(l->name().c_str())))
         continue;
 
-      ObtainCacheParticleVcol(geom, &b_mesh, &b_ob, &CData, !preview, vcol_num);
-
-      if (mesh) {
-        Attribute *attr_vcol = mesh->attributes.add(
-            ustring(l->name().c_str()), TypeDesc::TypeColor, ATTR_ELEMENT_CORNER_BYTE);
+      ObtainCacheParticleVcol(hair, &b_mesh, &b_ob, &CData, !preview, vcol_num);
 
-        uchar4 *cdata = attr_vcol->data_uchar4();
+      Attribute *attr_vcol = hair->attributes.add(
+          ustring(l->name().c_str()), TypeDesc::TypeColor, ATTR_ELEMENT_CURVE);
 
-        ExportCurveTriangleVcol(&CData, used_res, cdata);
-      }
-      else {
-        Attribute *attr_vcol = hair->attributes.add(
-            ustring(l->name().c_str()), TypeDesc::TypeColor, ATTR_ELEMENT_CURVE);
+      float3 *fdata = attr_vcol->data_float3();
 
-        float3 *fdata = attr_vcol->data_float3();
+      if (fdata) {
+        size_t i = 0;
 
-        if (fdata) {
-          size_t i = 0;
-
-          /* Encode vertex color using the sRGB curve. */
-          for (size_t curve = 0; curve < CData.curve_vcol.size(); curve++) {
-            fdata[i++] = color_srgb_to_linear_v3(CData.curve_vcol[curve]);
-          }
+        /* Encode vertex color using the sRGB curve. */
+        for (size_t curve = 0; curve < CData.curve_vcol.size(); curve++) {
+          fdata[i++] = color_srgb_to_linear_v3(CData.curve_vcol[curve]);
         }
       }
     }
@@ -1118,35 +641,23 @@ void BlenderSync::sync_particle_hair(
       ustring name = ustring(l->name().c_str());
 
       /* UV map */
-      if (geom->need_attribute(scene, name) || geom->need_attribute(scene, std)) {
+      if (hair->need_attribute(scene, name) || hair->need_attribute(scene, std)) {
         Attribute *attr_uv;
 
-        ObtainCacheParticleUV(geom, &b_mesh, &b_ob, &CData, !preview, uv_num);
-
-        if (mesh) {
-          if (active_render)
-            attr_uv = mesh->attributes.add(std, name);
-          else
-            attr_uv = mesh->attributes.add(name, TypeFloat2, ATTR_ELEMENT_CORNER);
+        ObtainCacheParticleUV(hair, &b_mesh, &b_ob, &CData, !preview, uv_num);
 
-          float2 *uv = attr_uv->data_float2();
-
-          ExportCurveTriangleUV(&CData, used_res, uv);
-        }
-        else {
-          if (active_render)
-            attr_uv = hair->attributes.add(std, name);
-          else
-            attr_uv = hair->attributes.add(name, TypeFloat2, ATTR_ELEMENT_CURVE);
+        if (active_render)
+          attr_uv = hair->attributes.add(std, name);
+        else
+          attr_uv = hair->attributes.add(name, TypeFloat2, ATTR_ELEMENT_CURVE);
 
-          float2 *uv = attr_uv->data_float2();
+        float2 *uv = attr_uv->data_float2();
 
-          if (uv) {
-            size_t i = 0;
+        if (uv) {
+          size_t i = 0;
 
-            for (size_t curve = 0; curve < CData.curve_uv.size(); curve++) {
-              uv[i++] = CData.curve_uv[curve];
-            }
+          for (size_t curve = 0; curve < CData.curve_uv.size(); curve++) {
+            uv[i++] = CData.curve_uv[curve];
           }
         }
       }
@@ -1344,27 +855,18 @@ void BlenderSync::sync_hair(Hair *hair, BL::Object &b_ob, bool motion, int motio
 
 void BlenderSync::sync_hair(BL::Depsgraph b_depsgraph,
                             BL::Object b_ob,
-                            Geometry *geom,
+                            Hair *hair,
                             const vector<Shader *> &used_shaders)
 {
-  Hair *hair = (geom->type == Geometry::HAIR) ? static_cast<Hair *>(geom) : NULL;
-  Mesh *mesh = (geom->type == Geometry::MESH) ? static_cast<Mesh *>(geom) : NULL;
-
   /* Compares curve_keys rather than strands in order to handle quick hair
    * adjustments in dynamic BVH - other methods could probably do this better. */
   array<float3> oldcurve_keys;
   array<float> oldcurve_radius;
-  array<int> oldtriangles;
-  if (hair) {
-    oldcurve_keys.steal_data(hair->curve_keys);
-    oldcurve_radius.steal_data(hair->curve_radius);
-  }
-  else {
-    oldtriangles.steal_data(mesh->triangles);
-  }
+  oldcurve_keys.steal_data(hair->curve_keys);
+  oldcurve_radius.steal_data(hair->curve_radius);
 
-  geom->clear();
-  geom->used_shaders = used_shaders;
+  hair->clear();
+  hair->used_shaders = used_shaders;
 
   if (view_layer.use_hair && scene->curve_system_manager->use_curves) {
 #ifdef WITH_NEW_OBJECT_TYPES
@@ -1377,35 +879,31 @@ void BlenderSync::sync_hair(BL::Depsgraph b_depsgraph,
 #endif
     {
       /* Particle hair. */
-      bool need_undeformed = geom->need_attribute(scene, ATTR_STD_GENERATED);
+      bool need_undeformed = hair->need_attribute(scene, ATTR_STD_GENERATED);
       BL::Mesh b_mesh = object_to_mesh(
           b_data, b_ob, b_depsgraph, need_undeformed, Mesh::SUBDIVISION_NONE);
 
       if (b_mesh) {
-        sync_particle_hair(geom, b_mesh, b_ob, false);
+        sync_particle_hair(hair, b_mesh, b_ob, false);
         free_object_to_mesh(b_data, b_ob, b_mesh);
       }
     }
   }
 
   /* tag update */
-  const bool rebuild = (hair && ((oldcurve_keys != hair->curve_keys) ||
-                                 (oldcurve_radius != hair->curve_radius))) ||
-                       (mesh && (oldtriangles != mesh->triangles));
+  const bool rebuild = ((oldcurve_keys != hair->curve_keys) ||
+                        (oldcurve_radius != hair->curve_radius));
 
-  geom->tag_update(scene, rebuild);
+  hair->tag_update(scene, rebuild);
 }
 
 void BlenderSync::sync_hair_motion(BL::Depsgraph b_depsgraph,
                                    BL::Object b_ob,
-                                   Geometry *geom,
+                                   Hair *hair,
                                    int motion_step)
 {
-  Hair *hair = (geom->type == Geometry::HAIR) ? static_cast<Hair *>(geom) : NULL;
-  Mesh *mesh = (geom->type == Geometry::MESH) ? static_cast<Mesh *>(geom) : NULL;
-
   /* Skip if nothing exported. */
-  if ((hair && hair->num_keys() == 0) || (mesh && mesh->verts.size() == 0)) {
+  if (hair->num_keys() == 0) {
     return;
   }
 
@@ -1424,7 +922,7 @@ void BlenderSync::sync_hair_motion(BL::Depsgraph b_depsgraph,
       /* Particle hair. */
       BL::Mesh b_mesh = object_to_mesh(b_data, b_ob, b_depsgraph, false, Mesh::SUBDIVISION_NONE);
       if (b_mesh) {
-        sync_particle_hair(geom, b_mesh, b_ob, true, motion_step);
+        sync_particle_hair(hair, b_mesh, b_ob, true, motion_step);
         free_object_to_mesh(b_data, b_ob, b_mesh);
         return;
       }
@@ -1432,12 +930,7 @@ void BlenderSync::sync_hair_motion(BL::Depsgraph b_depsgraph,
   }
 
   /* No deformation on this frame, copy coordinates if other frames did have it. */
-  if (hair) {
-    hair->copy_center_to_motion_step(motion_step);
-  }
-  else {
-    mesh->copy_center_to_motion_step(motion_step);
-  }
+  hair->copy_center_to_motion_step(motion_step);
 }
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/blender/blender_geometry.cpp b/intern/cycles/blender/blender_geometry.cpp
index 7ca35cff961..7d63899cb99 100644
--- a/intern/cycles/blender/blender_geometry.cpp
+++ b/intern/cycles/blender/blender_geometry.cpp
@@ -41,15 +41,11 @@ Geometry *BlenderSync::sync_geometry(BL::Depsgraph &b_depsgraph,
   Shader *default_shader = (b_ob.type() == BL::Object::type_VOLUME) ? scene->default_volume :
                                                                       scene->default_surface;
 #ifdef WITH_NEW_OBJECT_TYPES
-  Geometry::Type geom_type = ((b_ob.type() == BL::Object::type_HAIR || use_particle_hair) &&
-                              (scene->curve_system_manager->primitive != CURVE_TRIANGLES)) ?
+  Geometry::Type geom_type = (b_ob.type() == BL::Object::type_HAIR || use_particle_hair) ?
                                  Geometry::HAIR :
                                  Geometry::MESH;
 #else
-  Geometry::Type geom_type = ((use_particle_hair) &&
-                              (scene->curve_system_manager->primitive != CURVE_TRIANGLES)) ?
-                                 Geometry::HAIR :
-                                 Geometry::MESH;
+  Geometry::Type geom_type = (use_particle_hair) ? Geometry::HAIR : Geometry::MESH;
 #endif
 
   /* Find shader indices. */
@@ -134,7 +130,8 @@ Geometry *BlenderSync::sync_geometry(BL::Depsgraph &b_depsgraph,
 #else
   if (use_particle_hair) {
 #endif
-    sync_hair(b_depsgraph, b_ob, geom, used_shaders);
+    Hair *hair = static_cast<Hair *>(geom);
+    sync_hair(b_depsgraph, b_ob, hair, used_shaders);
   }
   else if (b_ob.type() == BL::Object::type_VOLUME || object_fluid_gas_domain_find(b_ob)) {
     Mesh *mesh = static_cast<Mesh *>(geom);
@@ -178,7 +175,8 @@ void BlenderSync::sync_geometry_motion(BL::Depsgraph &b_depsgraph,
 #else
   if (use_particle_hair) {
 #endif
-    sync_hair_motion(b_depsgraph, b_ob, geom, motion_step);
+    Hair *hair = static_cast<Hair *>(geom);
+    sync_hair_motion(b_depsgraph, b_ob, hair, motion_step);
   }
   else if (b_ob.type() == BL::Object::type_VOLUME || object_fluid_gas_domain_find(b_ob)) {
     /* No volume motion blur support yet. */
diff --git a/intern/cycles/blender/blender_sync.h b/intern/cycles/blender/blender_sync.h
index 341281b18ee..3c39cfd9446 100644
--- a/intern/cycles/blender/blender_sync.h
+++ b/intern/cycles/blender/blender_sync.h
@@ -153,15 +153,12 @@ class BlenderSync {
   /* Hair */
   void sync_hair(BL::Depsgraph b_depsgraph,
                  BL::Object b_ob,
-                 Geometry *geom,
+                 Hair *hair,
                  const vector<Shader *> &used_shaders);
-  void sync_hair_motion(BL::Depsgraph b_depsgraph,
-                        BL::Object b_ob,
-                        Geometry *geom,
-                        int motion_step);
+  void sync_hair_motion(BL::Depsgraph b_depsgraph, BL::Object b_ob, Hair *hair, int motion_step);
   void sync_hair(Hair *hair, BL::Object &b_ob, bool motion, int motion_step = 0);
   void sync_particle_hair(
-      Geometry *geom, BL::Mesh &b_mesh, BL::Object &b_ob, bool motion, int motion_step = 0);
+      Hair *hair, BL::Mesh &b_mesh, BL::Object &b_ob, bool motion, int motion_step = 0);
   void sync_curve_settings(BL::Depsgraph &b_depsgraph);
   bool object_has_particle_hair(BL::Object b_ob);
 
diff --git a/intern/cycles/bvh/bvh_embree.cpp b/intern/cycles/bvh/bvh_embree.cpp
index ebefc54be88..4e064f54742 100644
--- a/intern/cycles/bvh/bvh_embree.cpp
+++ b/intern/cycles/bvh/bvh_embree.cpp
@@ -78,8 +78,7 @@ static void rtc_filter_func(const RTCFilterFunctionNArguments *args)
   KernelGlobals *kg = ctx->kg;
 
   /* Check if there is backfacing hair to ignore. */
-  if (IS_HAIR(hit->geomID) && (kernel_data.curve.curveflags & CURVE_KN_INTERPOLATE) &&
-      !(kernel_data.curve.curveflags & CURVE_KN_BACKFACING) &&
+  if (IS_HAIR(hit->geomID) && !(kernel_data.curve.curveflags & CURVE_KN_BACKFACING) &&
       !(kernel_data.curve.curveflags & CURVE_KN_RIBBONS)) {
     if (dot(make_float3(ray->dir_x, ray->dir_y, ray->dir_z),
             make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z)) > 0.0f) {
@@ -99,8 +98,7 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
   KernelGlobals *kg = ctx->kg;
 
   /* For all ray types: Check if there is backfacing hair to ignore */
-  if (IS_HAIR(hit->geomID) && (kernel_data.curve.curveflags & CURVE_KN_INTERPOLATE) &&
-      !(kernel_data.curve.curveflags & CURVE_KN_BACKFACING) &&
+  if (IS_HAIR(hit->geomID) && !(kernel_data.curve.curveflags & CURVE_KN_BACKFACING) &&
       !(kernel_data.curve.curveflags & CURVE_KN_RIBBONS)) {
     if (dot(make_float3(ray->dir_x, ray->dir_y, ray->dir_z),
             make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z)) > 0.0f) {
@@ -328,7 +326,6 @@ BVHEmbree::BVHEmbree(const BVHParams &params_,
       stats(NULL),
       curve_subdivisions(params.curve_subdivisions),
       build_quality(RTC_BUILD_QUALITY_REFIT),
-      use_curves(params_.curve_flags & CURVE_KN_INTERPOLATE),
       use_ribbons(params.curve_flags & CURVE_KN_RIBBONS),
       dynamic_scene(true)
 {
@@ -726,9 +723,7 @@ void BVHEmbree::update_curve_vertex_buffer(RTCGeometry geom_id, const Hair *hair
 
   /* Catmull-Rom splines need extra CVs at the beginning and end of each curve. */
   size_t num_keys_embree = num_keys;
-  if (use_curves) {
-    num_keys_embree += num_curves * 2;
-  }
+  num_keys_embree += num_curves * 2;
 
   /* Copy the CV data to Embree */
   const int t_mid = (num_motion_steps - 1) / 2;
@@ -748,45 +743,22 @@ void BVHEmbree::update_curve_vertex_buffer(RTCGeometry geom_id, const Hair *hair
 
     assert(rtc_verts);
     if (rtc_verts) {
-      if (use_curves) {
-        const size_t num_curves = hair->num_curves();
-        for (size_t j = 0; j < num_curves; ++j) {
-          Hair::Curve c = hair->get_curve(j);
-          int fk = c.first_key;
-          int k = 1;
-          for (; k < c.num_keys + 1; ++k, ++fk) {
-            rtc_verts[k] = float3_to_float4(verts[fk]);
-            rtc_verts[k].w = curve_radius[fk];
-          }
-          /* Duplicate Embree's Catmull-Rom spline CVs at the start and end of each curve. */
-          rtc_verts[0] = rtc_verts[1];
-          rtc_verts[k] = rtc_verts[k - 1];
-          rtc_verts += c.num_keys + 2;
-        }
-      }
-      else {
-        for (size_t j = 0; j < num_keys_embree; ++j) {
-          rtc_verts[j] = float3_to_float4(verts[j]);
-          rtc_verts[j].w = curve_radius[j];
+      const size_t num_curves = hair->num_curves();
+      for (size_t j = 0; j < num_curves; ++j) {
+        Hair::Curve c = hair->get_curve(j);
+        int fk = c.first_key;
+        int k = 1;
+        for (; k < c.num_keys + 1; ++k, ++fk) {
+          rtc_verts[k] = float3_to_float4(verts[fk]);
+          rtc_verts[k].w = curve_radius[fk];
         }
+        /* Duplicate Embree's Catmull-Rom spline CVs at the start and end of each curve. */
+        rtc_verts[0] = rtc_verts[1];
+        rtc_verts[k] = rtc_verts[k - 1];
+        rtc_verts += c.num_keys + 2;
       }
     }
   }
-#  if RTC_VERSION >= 30900
-  if (!use_curves) {
-    unsigned char *flags = (unsigned char *)rtcSetNewGeometryBuffer(geom_id,
-                                                                    RTC_BUFFER_TYPE_FLAGS,
-                                                                    0,
-                                                                    RTC_FORMAT_UCHAR,
-                                                                    sizeof(unsigned char),
-                                                                    num_keys_embree);
-    flags[0] = RTC_CURVE_FLAG_NEIGHBOR_RIGHT;
-    ::memset(flags + 1,
-             RTC_CURVE_FLAG_NEIGHBOR_RIGHT | RTC_CURVE_FLAG_NEIGHBOR_RIGHT,
-             num_keys_embree - 2);
-    flags[num_keys_embree - 1] = RTC_CURVE_FLAG_NEIGHBOR_LEFT;
-  }
-#  endif
 }
 
 void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i)
@@ -822,18 +794,8 @@ void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i)
   size_t prim_tri_index_size = pack.prim_index.size();
   pack.prim_tri_index.resize(prim_tri_index_size + num_segments);
 
-#  if RTC_VERSION >= 30900
-  enum RTCGeometryType type = (!use_curves) ?
-                                  (use_ribbons ? RTC_GEOMETRY_TYPE_FLAT_LINEAR_CURVE :
-                                                 RTC_GEOMETRY_TYPE_ROUND_LINEAR_CURVE) :
-                                  (use_ribbons ? RTC_GEOMETRY_TYPE_FLAT_CATMULL_ROM_CURVE :
-                                                 RTC_GEOMETRY_TYPE_ROUND_CATMULL_ROM_CURVE);
-#  else
-  enum RTCGeometryType type = (!use_curves) ?
-                                  RTC_GEOMETRY_TYPE_FLAT_LINEAR_CURVE :
-                                  (use_ribbons ? RTC_GEOMETRY_TYPE_FLAT_CATMULL_ROM_CURVE :
-                                                 RTC_GEOMETRY_TYPE_ROUND_CATMULL_ROM_CURVE);
-#  endif
+  enum RTCGeometryType type = (use_ribbons ? RTC_GEOMETRY_TYPE_FLAT_CATMULL_ROM_CURVE :
+                                             RTC_GEOMETRY_TYPE_ROUND_CATMULL_ROM_CURVE);
 
   RTCGeometry geom_id = rtcNewGeometry(rtc_shared_device, type);
   rtcSetGeometryTessellationRate(geom_id, curve_subdivisions);
@@ -844,10 +806,8 @@ void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i)
     Hair::Curve c = hair->get_curve(j);
     for (size_t k = 0; k < c.num_segments(); ++k) {
       rtc_indices[rtc_index] = c.first_key + k;
-      if (use_curves) {
-        /* Room for extra CVs at Catmull-Rom splines. */
-        rtc_indices[rtc_index] += j * 2;
-      }
+      /* Room for extra CVs at Catmull-Rom splines. */
+      rtc_indices[rtc_index] += j * 2;
       /* Cycles specific data. */
       pack.prim_object[prim_object_size + rtc_index] = i;
       pack.prim_type[prim_type_size + rtc_index] = (PRIMITIVE_PACK_SEGMENT(
diff --git a/intern/cycles/bvh/bvh_embree.h b/intern/cycles/bvh/bvh_embree.h
index eb121d060b7..b10a741fd59 100644
--- a/intern/cycles/bvh/bvh_embree.h
+++ b/intern/cycles/bvh/bvh_embree.h
@@ -81,7 +81,7 @@ class BVHEmbree : public BVH {
   vector<RTCScene> delayed_delete_scenes;
   int curve_subdivisions;
   enum RTCBuildQuality build_quality;
-  bool use_curves, use_ribbons, dynamic_scene;
+  bool use_ribbons, dynamic_scene;
 };
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/bvh/bvh_shadow_all.h b/intern/cycles/kernel/bvh/bvh_shadow_all.h
index 268bb149970..385481983e4 100644
--- a/intern/cycles/kernel/bvh/bvh_shadow_all.h
+++ b/intern/cycles/kernel/bvh/bvh_shadow_all.h
@@ -210,28 +210,8 @@ ccl_device_inline
               case PRIMITIVE_CURVE:
               case PRIMITIVE_MOTION_CURVE: {
                 const uint curve_type = kernel_tex_fetch(__prim_type, prim_addr);
-                if (kernel_data.curve.curveflags & CURVE_KN_INTERPOLATE) {
-                  hit = cardinal_curve_intersect(kg,
-                                                 isect_array,
-                                                 P,
-                                                 dir,
-                                                 visibility,
-                                                 object,
-                                                 prim_addr,
-                                                 ray->time,
-                                                 curve_type);
-                }
-                else {
-                  hit = curve_intersect(kg,
-                                        isect_array,
-                                        P,
-                                        dir,
-                                        visibility,
-                                        object,
-                                        prim_addr,
-                                        ray->time,
-                                        curve_type);
-                }
+                hit = curve_intersect(
+                    kg, isect_array, P, dir, visibility, object, prim_addr, ray->time, curve_type);
                 break;
               }
 #endif
diff --git a/intern/cycles/kernel/bvh/bvh_traversal.h b/intern/cycles/kernel/bvh/bvh_traversal.h
index 18afc6ae4eb..ba83690dfef 100644
--- a/intern/cycles/kernel/bvh/bvh_traversal.h
+++ b/intern/cycles/kernel/bvh/bvh_traversal.h
@@ -237,15 +237,8 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
                 BVH_DEBUG_NEXT_INTERSECTION();
                 const uint curve_type = kernel_tex_fetch(__prim_type, prim_addr);
                 kernel_assert((curve_type & PRIMITIVE_ALL) == (type & PRIMITIVE_ALL));
-                bool hit;
-                if (kernel_data.curve.curveflags & CURVE_KN_INTERPOLATE) {
-                  hit = cardinal_curve_intersect(
-                      kg, isect, P, dir, visibility, object, prim_addr, ray->time, curve_type);
-                }
-                else {
-                  hit = curve_intersect(
-                      kg, isect, P, dir, visibility, object, prim_addr, ray->time, curve_type);
-                }
+                bool hit = curve_intersect(
+                    kg, isect, P, dir, visibility, object, prim_addr, ray->time, curve_type);
                 if (hit) {
                   /* shadow ray early termination */
 #  if defined(__KERNEL_SSE2__)
diff --git a/intern/cycles/kernel/bvh/obvh_shadow_all.h b/intern/cycles/kernel/bvh/obvh_shadow_all.h
index b7ab75b723c..ff82d5105a4 100644
--- a/intern/cycles/kernel/bvh/obvh_shadow_all.h
+++ b/intern/cycles/kernel/bvh/obvh_shadow_all.h
@@ -494,28 +494,15 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(OBVH)(KernelGlobals *kg,
                 case PRIMITIVE_CURVE:
                 case PRIMITIVE_MOTION_CURVE: {
                   const uint curve_type = kernel_tex_fetch(__prim_type, prim_addr);
-                  if (kernel_data.curve.curveflags & CURVE_KN_INTERPOLATE) {
-                    hit = cardinal_curve_intersect(kg,
-                                                   isect_array,
-                                                   P,
-                                                   dir,
-                                                   PATH_RAY_SHADOW,
-                                                   object,
-                                                   prim_addr,
-                                                   ray->time,
-                                                   curve_type);
-                  }
-                  else {
-                    hit = curve_intersect(kg,
-                                          isect_array,
-                                          P,
-                                          dir,
-                                          PATH_RAY_SHADOW,
-                                          object,
-                                          prim_addr,
-                                          ray->time,
-                                          curve_type);
-                  }
+                  hit = curve_intersect(kg,
+                                                 isect_array,
+                                                 P,
+                                                 dir,
+                                                 PATH_RAY_SHADOW,
+                                                 object,
+                                                 prim_addr,
+                                                 ray->time,
+                                                 curve_type);
                   break;
                 }
 #endif
diff --git a/intern/cycles/kernel/bvh/obvh_traversal.h b/intern/cycles/kernel/bvh/obvh_traversal.h
index 9095233f8b6..3f1e03693c3 100644
--- a/intern/cycles/kernel/bvh/obvh_traversal.h
+++ b/intern/cycles/kernel/bvh/obvh_traversal.h
@@ -457,15 +457,8 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(OBVH)(KernelGlobals *kg,
                 BVH_DEBUG_NEXT_INTERSECTION();
                 const uint curve_type = kernel_tex_fetch(__prim_type, prim_addr);
                 kernel_assert((curve_type & PRIMITIVE_ALL) == (type & PRIMITIVE_ALL));
-                bool hit;
-                if (kernel_data.curve.curveflags & CURVE_KN_INTERPOLATE) {
-                  hit = cardinal_curve_intersect(
-                      kg, isect, P, dir, visibility, object, prim_addr, ray->time, curve_type);
-                }
-                else {
-                  hit = curve_intersect(
-                      kg, isect, P, dir, visibility, object, prim_addr, ray->time, curve_type);
-                }
+                bool hit = curve_intersect(
+                    kg, isect, P, dir, visibility, object, prim_addr, ray->time, curve_type);
                 if (hit) {
                   tfar = avxf(isect->t);
                   /* Shadow ray early termination. */
diff --git a/intern/cycles/kernel/bvh/qbvh_shadow_all.h b/intern/cycles/kernel/bvh/qbvh_shadow_all.h
index 682251bf25b..9d428c3e1a7 100644
--- a/intern/cycles/kernel/bvh/qbvh_shadow_all.h
+++ b/intern/cycles/kernel/bvh/qbvh_shadow_all.h
@@ -284,28 +284,8 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
               case PRIMITIVE_CURVE:
               case PRIMITIVE_MOTION_CURVE: {
                 const uint curve_type = kernel_tex_fetch(__prim_type, prim_addr);
-                if (kernel_data.curve.curveflags & CURVE_KN_INTERPOLATE) {
-                  hit = cardinal_curve_intersect(kg,
-                                                 isect_array,
-                                                 P,
-                                                 dir,
-                                                 visibility,
-                                                 object,
-                                                 prim_addr,
-                                                 ray->time,
-                                                 curve_type);
-                }
-                else {
-                  hit = curve_intersect(kg,
-                                        isect_array,
-                                        P,
-                                        dir,
-                                        visibility,
-                                        object,
-                                        prim_addr,
-                                        ray->time,
-                                        curve_type);
-                }
+                hit = curve_intersect(
+                    kg, isect_array, P, dir, visibility, object, prim_addr, ray->time, curve_type);
                 break;
               }
 #endif
diff --git a/intern/cycles/kernel/bvh/qbvh_traversal.h b/intern/cycles/kernel/bvh/qbvh_traversal.h
index f43e84bf368..f68579b4d69 100644
--- a/intern/cycles/kernel/bvh/qbvh_traversal.h
+++ b/intern/cycles/kernel/bvh/qbvh_traversal.h
@@ -320,15 +320,8 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
                 BVH_DEBUG_NEXT_INTERSECTION();
                 const uint curve_type = kernel_tex_fetch(__prim_type, prim_addr);
                 kernel_assert((curve_type & PRIMITIVE_ALL) == (type & PRIMITIVE_ALL));
-                bool hit;
-                if (kernel_data.curve.curveflags & CURVE_KN_INTERPOLATE) {
-                  hit = cardinal_curve_intersect(
-                      kg, isect, P, dir, visibility, object, prim_addr, ray->time, curve_type);
-                }
-                else {
-                  hit = curve_intersect(
-                      kg, isect, P, dir, visibility, object, prim_addr, ray->time, curve_type);
-                }
+                bool hit = curve_intersect(
+                    kg, isect, P, dir, visibility, object, prim_addr, ray->time, curve_type);
                 if (hit) {
                   tfar = ssef(isect->t);
                   /* Shadow ray early termination. */
diff --git a/intern/cycles/kernel/geom/geom_curve_intersect.h b/intern/cycles/kernel/geom/geom_curve_intersect.h
index 7a770470150..971f194a7dd 100644
--- a/intern/cycles/kernel/geom/geom_curve_intersect.h
+++ b/intern/cycles/kernel/geom/geom_curve_intersect.h
@@ -26,7 +26,7 @@ ccl_device_inline ssef transform_point_T3(const ssef t[3], const ssef &a)
 #  endif
 
 /* On CPU pass P and dir by reference to aligned vector. */
-ccl_device_forceinline bool cardinal_curve_intersect(KernelGlobals *kg,
+ccl_device_forceinline bool curve_intersect(KernelGlobals *kg,
                                                      Intersection *isect,
                                                      const float3 ccl_ref P,
                                                      const float3 ccl_ref dir,
@@ -82,7 +82,7 @@ ccl_device_forceinline bool cardinal_curve_intersect(KernelGlobals *kg,
     }
     else {
       int fobject = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, curveAddr) : object;
-      motion_cardinal_curve_keys_avx(
+      motion_curve_keys_avx(
           kg, fobject, prim, time, ka, k0, k1, kb, &P_curve_0_1, &P_curve_2_3);
     }
 #    else  /* __KERNEL_AVX2__ */
@@ -96,7 +96,7 @@ ccl_device_forceinline bool cardinal_curve_intersect(KernelGlobals *kg,
     }
     else {
       int fobject = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, curveAddr) : object;
-      motion_cardinal_curve_keys(kg, fobject, prim, time, ka, k0, k1, kb, (float4 *)&P_curve);
+      motion_curve_keys(kg, fobject, prim, time, ka, k0, k1, kb, (float4 *)&P_curve);
     }
 #    endif /* __KERNEL_AVX2__ */
 
@@ -197,7 +197,7 @@ ccl_device_forceinline bool cardinal_curve_intersect(KernelGlobals *kg,
     }
     else {
       int fobject = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, curveAddr) : object;
-      motion_cardinal_curve_keys(kg, fobject, prim, time, ka, k0, k1, kb, P_curve);
+      motion_curve_keys(kg, fobject, prim, time, ka, k0, k1, kb, P_curve);
     }
 
     float3 p0 = transform_point(&htfm, float4_to_float3(P_curve[0]) - P);
@@ -488,239 +488,11 @@ ccl_device_forceinline bool cardinal_curve_intersect(KernelGlobals *kg,
   return hit;
 }
 
-ccl_device_forceinline bool curve_intersect(KernelGlobals *kg,
-                                            Intersection *isect,
-                                            float3 P,
-                                            float3 direction,
-                                            uint visibility,
-                                            int object,
-                                            int curveAddr,
-                                            float time,
-                                            int type)
-{
-  /* define few macros to minimize code duplication for SSE */
-#  ifndef __KERNEL_SSE2__
-#    define len3_squared(x) len_squared(x)
-#    define len3(x) len(x)
-#    define dot3(x, y) dot(x, y)
-#  endif
-
-  const bool is_curve_primitive = (type & PRIMITIVE_CURVE);
-
-#  ifndef __KERNEL_OPTIX__ /* see OptiX motion flag OPTIX_MOTION_FLAG_[START|END]_VANISH */
-  if (!is_curve_primitive && kernel_data.bvh.use_bvh_steps) {
-    const float2 prim_time = kernel_tex_fetch(__prim_time, curveAddr);
-    if (time < prim_time.x || time > prim_time.y) {
-      return false;
-    }
-  }
-#  endif
-
-  int segment = PRIMITIVE_UNPACK_SEGMENT(type);
-  /* curve Intersection check */
-  int flags = kernel_data.curve.curveflags;
-
-  int prim = kernel_tex_fetch(__prim_index, curveAddr);
-  float4 v00 = kernel_tex_fetch(__curves, prim);
-
-  int cnum = __float_as_int(v00.x);
-  int k0 = cnum + segment;
-  int k1 = k0 + 1;
-
-#  ifndef __KERNEL_SSE2__
-  float4 P_curve[2];
-
-  if (is_curve_primitive) {
-    P_curve[0] = kernel_tex_fetch(__curve_keys, k0);
-    P_curve[1] = kernel_tex_fetch(__curve_keys, k1);
-  }
-  else {
-    int fobject = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, curveAddr) : object;
-    motion_curve_keys(kg, fobject, prim, time, k0, k1, P_curve);
-  }
-
-  float r1 = P_curve[0].w;
-  float r2 = P_curve[1].w;
-  float3 p1 = float4_to_float3(P_curve[0]);
-  float3 p2 = float4_to_float3(P_curve[1]);
-
-  /* minimum width extension */
-  float3 dif = P - p1;
-  float3 dif_second = P - p2;
-
-  float3 p21_diff = p2 - p1;
-  float3 sphere_dif1 = (dif + dif_second) * 0.5f;
-  float3 dir = direction;
-  float sphere_b_tmp = dot3(dir, sphere_dif1);
-  float3 sphere_dif2 = sphere_dif1 - sphere_b_tmp * dir;
-#  else
-  ssef P_curve[2];
-
-  if (is_curve_primitive) {
-    P_curve[0] = load4f(&kg->__curve_keys.data[k0].x);
-    P_curve[1] = load4f(&kg->__curve_keys.data[k1].x);
-  }
-  else {
-    int fobject = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, curveAddr) : object;
-    motion_curve_keys(kg, fobject, prim, time, k0, k1, (float4 *)&P_curve);
-  }
-
-  ssef r12 = shuffle<3, 3, 3, 3>(P_curve[0], P_curve[1]);
-  const ssef vP = load4f(P);
-  const ssef dif = vP - P_curve[0];
-  const ssef dif_second = vP - P_curve[1];
-  float r1 = extract<0>(r12), r2 = extract<0>(shuffle<2>(r12));
-
-  const ssef p21_diff = P_curve[1] - P_curve[0];
-  const ssef sphere_dif1 = (dif + dif_second) * 0.5f;
-  const ssef dir = load4f(direction);
-  const ssef sphere_b_tmp = dot3_splat(dir, sphere_dif1);
-  const ssef sphere_dif2 = nmadd(sphere_b_tmp, dir, sphere_dif1);
-#  endif
-
-  float mr = max(r1, r2);
-  float l = len3(p21_diff);
-  float invl = 1.0f / l;
-  float sp_r = mr + 0.5f * l;
-
-  float sphere_b = dot3(dir, sphere_dif2);
-  float sdisc = sphere_b * sphere_b - len3_squared(sphere_dif2) + sp_r * sp_r;
-
-  if (sdisc < 0.0f)
-    return false;
-
-    /* obtain parameters and test midpoint distance for suitable modes */
-#  ifndef __KERNEL_SSE2__
-  float3 tg = p21_diff * invl;
-#  else
-  const ssef tg = p21_diff * invl;
-#  endif
-  float gd = (r2 - r1) * invl;
-
-  float dirz = dot3(dir, tg);
-  float difz = dot3(dif, tg);
-
-  float a = 1.0f - (dirz * dirz * (1 + gd * gd));
-
-  float halfb = dot3(dir, dif) - dirz * (difz + gd * (difz * gd + r1));
-
-  float tcentre = -halfb / a;
-  float zcentre = difz + (dirz * tcentre);
-
-  if ((tcentre > isect->t) && !(flags & CURVE_KN_ACCURATE))
-    return false;
-  if ((zcentre < 0 || zcentre > l) && !(flags & CURVE_KN_ACCURATE) &&
-      !(flags & CURVE_KN_INTERSECTCORRECTION))
-    return false;
-
-    /* test minimum separation */
-#  ifndef __KERNEL_SSE2__
-  float3 cprod = cross(tg, dir);
-  float cprod2sq = len3_squared(cross(tg, dif));
-#  else
-  const ssef cprod = cross(tg, dir);
-  float cprod2sq = len3_squared(cross_zxy(tg, dif));
-#  endif
-  float cprodsq = len3_squared(cprod);
-  float distscaled = dot3(cprod, dif);
-
-  if (cprodsq == 0)
-    distscaled = cprod2sq;
-  else
-    distscaled = (distscaled * distscaled) / cprodsq;
-
-  if (distscaled > mr * mr)
-    return false;
-
-    /* calculate true intersection */
-#  ifndef __KERNEL_SSE2__
-  float3 tdif = dif + tcentre * dir;
-#  else
-  const ssef tdif = madd(ssef(tcentre), dir, dif);
-#  endif
-  float tdifz = dot3(tdif, tg);
-  float tdifma = tdifz * gd + r1;
-  float tb = 2 * (dot3(dir, tdif) - dirz * (tdifz + gd * tdifma));
-  float tc = dot3(tdif, tdif) - tdifz * tdifz - tdifma * tdifma;
-  float td = tb * tb - 4 * a * tc;
-
-  if (td < 0.0f)
-    return false;
-
-  float rootd = 0.0f;
-  float correction = 0.0f;
-  if (flags & CURVE_KN_ACCURATE) {
-    rootd = sqrtf(td);
-    correction = ((-tb - rootd) / (2 * a));
-  }
-
-  float t = tcentre + correction;
-
-  if (t < isect->t) {
-
-    if (flags & CURVE_KN_INTERSECTCORRECTION) {
-      rootd = sqrtf(td);
-      correction = ((-tb - rootd) / (2 * a));
-      t = tcentre + correction;
-    }
-
-    float z = zcentre + (dirz * correction);
-    // bool backface = false;
-
-    if (flags & CURVE_KN_BACKFACING && (t < 0.0f || z < 0 || z > l)) {
-      // backface = true;
-      correction = ((-tb + rootd) / (2 * a));
-      t = tcentre + correction;
-      z = zcentre + (dirz * correction);
-    }
-
-    if (t > 0.0f && t < isect->t && z >= 0 && z <= l) {
-
-      if (flags & CURVE_KN_ENCLOSEFILTER) {
-        float enc_ratio = 1.01f;
-        if ((difz > -r1 * enc_ratio) && (dot3(dif_second, tg) < r2 * enc_ratio)) {
-          float a2 = 1.0f - (dirz * dirz * (1 + gd * gd * enc_ratio * enc_ratio));
-          float c2 = dot3(dif, dif) - difz * difz * (1 + gd * gd * enc_ratio * enc_ratio) -
-                     r1 * r1 * enc_ratio * enc_ratio - 2 * r1 * difz * gd * enc_ratio;
-          if (a2 * c2 < 0.0f)
-            return false;
-        }
-      }
-
-#  ifdef __VISIBILITY_FLAG__
-      /* visibility flag test. we do it here under the assumption
-       * that most triangles are culled by node flags */
-      if (kernel_tex_fetch(__prim_visibility, curveAddr) & visibility)
-#  endif
-      {
-        /* record intersection */
-        isect->t = t;
-        isect->u = z * invl;
-        isect->v = gd;
-        isect->prim = curveAddr;
-        isect->object = object;
-        isect->type = type;
-
-        return true;
-      }
-    }
-  }
-
-  return false;
-
-#  ifndef __KERNEL_SSE2__
-#    undef len3_squared
-#    undef len3
-#    undef dot3
-#  endif
-}
-
 ccl_device_inline float3 curve_refine(KernelGlobals *kg,
                                       ShaderData *sd,
                                       const Intersection *isect,
                                       const Ray *ray)
 {
-  int flag = kernel_data.curve.curveflags;
   float t = isect->t;
   float3 P = ray->P;
   float3 D = ray->D;
@@ -743,114 +515,66 @@ ccl_device_inline float3 curve_refine(KernelGlobals *kg,
   int k0 = __float_as_int(v00.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
   int k1 = k0 + 1;
 
-  float3 tg;
-
-  if (flag & CURVE_KN_INTERPOLATE) {
-    int ka = max(k0 - 1, __float_as_int(v00.x));
-    int kb = min(k1 + 1, __float_as_int(v00.x) + __float_as_int(v00.y) - 1);
+  int ka = max(k0 - 1, __float_as_int(v00.x));
+  int kb = min(k1 + 1, __float_as_int(v00.x) + __float_as_int(v00.y) - 1);
 
-    float4 P_curve[4];
+  float4 P_curve[4];
 
-    if (sd->type & PRIMITIVE_CURVE) {
-      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);
-      P_curve[3] = kernel_tex_fetch(__curve_keys, kb);
-    }
-    else {
-      motion_cardinal_curve_keys(kg, sd->object, sd->prim, sd->time, ka, k0, k1, kb, P_curve);
-    }
+  if (sd->type & PRIMITIVE_CURVE) {
+    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);
+    P_curve[3] = kernel_tex_fetch(__curve_keys, kb);
+  }
+  else {
+    motion_curve_keys(kg, sd->object, sd->prim, sd->time, ka, k0, k1, kb, P_curve);
+  }
 
-    float3 p[4];
-    p[0] = float4_to_float3(P_curve[0]);
-    p[1] = float4_to_float3(P_curve[1]);
-    p[2] = float4_to_float3(P_curve[2]);
-    p[3] = float4_to_float3(P_curve[3]);
+  float3 p[4];
+  p[0] = float4_to_float3(P_curve[0]);
+  p[1] = float4_to_float3(P_curve[1]);
+  p[2] = float4_to_float3(P_curve[2]);
+  p[3] = float4_to_float3(P_curve[3]);
 
-    P = P + D * t;
+  P = P + D * t;
 
 #  ifdef __UV__
-    sd->u = isect->u;
-    sd->v = 0.0f;
-#  endif
-
-    tg = normalize(curvetangent(isect->u, p[0], p[1], p[2], p[3]));
-
-    if (kernel_data.curve.curveflags & CURVE_KN_RIBBONS) {
-      sd->Ng = normalize(-(D - tg * (dot(tg, D))));
-    }
-    else {
-#  ifdef __EMBREE__
-      if (kernel_data.bvh.scene) {
-        sd->Ng = normalize(isect->Ng);
-      }
-      else
+  sd->u = isect->u;
+  sd->v = 0.0f;
 #  endif
-      {
-        /* direction from inside to surface of curve */
-        float3 p_curr = curvepoint(isect->u, p[0], p[1], p[2], p[3]);
-        sd->Ng = normalize(P - p_curr);
 
-        /* adjustment for changing radius */
-        float gd = isect->v;
+  float3 tg = normalize(curvetangent(isect->u, p[0], p[1], p[2], p[3]));
 
-        if (gd != 0.0f) {
-          sd->Ng = sd->Ng - gd * tg;
-          sd->Ng = normalize(sd->Ng);
-        }
-      }
-    }
-
-    /* todo: sometimes the normal is still so that this is detected as
-     * backfacing even if cull backfaces is enabled */
-
-    sd->N = sd->Ng;
+  if (kernel_data.curve.curveflags & CURVE_KN_RIBBONS) {
+    sd->Ng = normalize(-(D - tg * (dot(tg, D))));
   }
   else {
-    float4 P_curve[2];
-
-    if (sd->type & PRIMITIVE_CURVE) {
-      P_curve[0] = kernel_tex_fetch(__curve_keys, k0);
-      P_curve[1] = kernel_tex_fetch(__curve_keys, k1);
-    }
-    else {
-      motion_curve_keys(kg, sd->object, sd->prim, sd->time, k0, k1, P_curve);
+#  ifdef __EMBREE__
+    if (kernel_data.bvh.scene) {
+      sd->Ng = normalize(isect->Ng);
     }
-
-    float l = 1.0f;
-    tg = normalize_len(float4_to_float3(P_curve[1] - P_curve[0]), &l);
-
-    P = P + D * t;
-
-    float3 dif = P - float4_to_float3(P_curve[0]);
-
-#  ifdef __UV__
-    sd->u = dot(dif, tg) / l;
-    sd->v = 0.0f;
+    else
 #  endif
-
-    if (flag & CURVE_KN_TRUETANGENTGNORMAL) {
-      sd->Ng = -(D - tg * dot(tg, D));
-      sd->Ng = normalize(sd->Ng);
-    }
-    else {
-      float gd = isect->v;
-
+    {
       /* direction from inside to surface of curve */
-      float denom = fmaxf(P_curve[0].w + sd->u * l * gd, 1e-8f);
-      sd->Ng = (dif - tg * sd->u * l) / denom;
+      float3 p_curr = curvepoint(isect->u, p[0], p[1], p[2], p[3]);
+      sd->Ng = normalize(P - p_curr);
 
       /* adjustment for changing radius */
+      float gd = isect->v;
+
       if (gd != 0.0f) {
         sd->Ng = sd->Ng - gd * tg;
+        sd->Ng = normalize(sd->Ng);
       }
-
-      sd->Ng = normalize(sd->Ng);
     }
-
-    sd->N = sd->Ng;
   }
 
+  /* todo: sometimes the normal is still so that this is detected as
+   * backfacing even if cull backfaces is enabled */
+
+  sd->N = sd->Ng;
+
 #  ifdef __DPDU__
   /* dPdu/dPdv */
   sd->dPdu = tg;
diff --git a/intern/cycles/kernel/geom/geom_motion_curve.h b/intern/cycles/kernel/geom/geom_motion_curve.h
index 0e2a00e9d2e..dd7429c02bd 100644
--- a/intern/cycles/kernel/geom/geom_motion_curve.h
+++ b/intern/cycles/kernel/geom/geom_motion_curve.h
@@ -105,7 +105,7 @@ ccl_device_inline void motion_curve_keys(
   keys[1] = (1.0f - t) * keys[1] + t * next_keys[1];
 }
 
-ccl_device_inline void motion_cardinal_curve_keys_for_step(KernelGlobals *kg,
+ccl_device_inline void motion_curve_keys_for_step(KernelGlobals *kg,
                                                            int offset,
                                                            int numkeys,
                                                            int numsteps,
@@ -138,7 +138,7 @@ ccl_device_inline void motion_cardinal_curve_keys_for_step(KernelGlobals *kg,
 }
 
 /* return 2 curve key locations */
-ccl_device_inline void motion_cardinal_curve_keys(KernelGlobals *kg,
+ccl_device_inline void motion_curve_keys(KernelGlobals *kg,
                                                   int object,
                                                   int prim,
                                                   float time,
@@ -165,8 +165,8 @@ ccl_device_inline void motion_cardinal_curve_keys(KernelGlobals *kg,
   /* fetch key coordinates */
   float4 next_keys[4];
 
-  motion_cardinal_curve_keys_for_step(kg, offset, numkeys, numsteps, step, k0, k1, k2, k3, keys);
-  motion_cardinal_curve_keys_for_step(
+  motion_curve_keys_for_step(kg, offset, numkeys, numsteps, step, k0, k1, k2, k3, keys);
+  motion_curve_keys_for_step(
       kg, offset, numkeys, numsteps, step + 1, k0, k1, k2, k3, next_keys);
 
   /* interpolate between steps */
@@ -180,7 +180,7 @@ ccl_device_inline void motion_cardinal_curve_keys(KernelGlobals *kg,
 /* Similar to above, but returns keys as pair of two AVX registers with each
  * holding two float4.
  */
-ccl_device_inline void motion_cardinal_curve_keys_avx(KernelGlobals *kg,
+ccl_device_inline void motion_curve_keys_avx(KernelGlobals *kg,
                                                       int object,
                                                       int prim,
                                                       float time,
@@ -208,8 +208,8 @@ ccl_device_inline void motion_cardinal_curve_keys_avx(KernelGlobals *kg,
   /* Fetch key coordinates. */
   float4 next_keys[4];
   float4 keys[4];
-  motion_cardinal_curve_keys_for_step(kg, offset, numkeys, numsteps, step, k0, k1, k2, k3, keys);
-  motion_cardinal_curve_keys_for_step(
+  motion_curve_keys_for_step(kg, offset, numkeys, numsteps, step, k0, k1, k2, k3, keys);
+  motion_curve_keys_for_step(
       kg, offset, numkeys, numsteps, step + 1, k0, k1, k2, k3, next_keys);
 
   const avxf keys_0_1 = avxf(keys[0].m128, keys[1].m128);
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index a692d7a844f..75d988365a2 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -1428,13 +1428,8 @@ static_assert_align(KernelBVH, 16);
 
 typedef enum CurveFlag {
   /* runtime flags */
-  CURVE_KN_BACKFACING = 1,           /* backside of cylinder? */
-  CURVE_KN_ENCLOSEFILTER = 2,        /* don't consider strands surrounding start point? */
-  CURVE_KN_INTERPOLATE = 4,          /* render as a curve? */
-  CURVE_KN_ACCURATE = 8,             /* use accurate intersections test? */
-  CURVE_KN_INTERSECTCORRECTION = 16, /* correct for width after determing closest midpoint? */
-  CURVE_KN_TRUETANGENTGNORMAL = 32,  /* use tangent normal for geometry? */
-  CURVE_KN_RIBBONS = 64,             /* use flat curve ribbons */
+  CURVE_KN_BACKFACING = 1, /* backside of cylinder? */
+  CURVE_KN_RIBBONS = 2,    /* use flat curve ribbons */
 } CurveFlag;
 
 typedef struct KernelCurves {
diff --git a/intern/cycles/kernel/kernels/optix/kernel_optix.cu b/intern/cycles/kernel/kernels/optix/kernel_optix.cu
index e03504316ad..59fece73b3c 100644
--- a/intern/cycles/kernel/kernels/optix/kernel_optix.cu
+++ b/intern/cycles/kernel/kernels/optix/kernel_optix.cu
@@ -282,9 +282,7 @@ extern "C" __global__ void __intersection__curve()
   if (isect.t != FLT_MAX)
     isect.t *= len;
 
-  if (!(kernel_data.curve.curveflags & CURVE_KN_INTERPOLATE) ?
-          curve_intersect(NULL, &isect, P, dir, visibility, object, prim, time, type) :
-          cardinal_curve_intersect(NULL, &isect, P, dir, visibility, object, prim, time, type)) {
+  if (curve_intersect(NULL, &isect, P, dir, visibility, object, prim, time, type)) {
     optixReportIntersection(isect.t / len,
                             type & PRIMITIVE_ALL,
                             __float_as_int(isect.u),   // Attribute_0
diff --git a/intern/cycles/render/curves.cpp b/intern/cycles/render/curves.cpp
index 1907bb33d06..a223702cfab 100644
--- a/intern/cycles/render/curves.cpp
+++ b/intern/cycles/render/curves.cpp
@@ -81,17 +81,11 @@ void curvebounds(float *lower, float *upper, float3 *p, int dim)
 
 CurveSystemManager::CurveSystemManager()
 {
-  primitive = CURVE_LINE_SEGMENTS;
   curve_shape = CURVE_THICK;
-  line_method = CURVE_CORRECTED;
-  triangle_method = CURVE_CAMERA_TRIANGLES;
-  resolution = 3;
   subdivisions = 3;
 
   use_curves = true;
-  use_encasing = true;
   use_backfacing = false;
-  use_tangent_normal_geometry = false;
 
   need_update = true;
   need_mesh_update = false;
@@ -118,22 +112,12 @@ void CurveSystemManager::device_update(Device *device,
   kcurve->curveflags = 0;
 
   if (use_curves) {
-    if (primitive == CURVE_SEGMENTS || primitive == CURVE_RIBBONS)
-      kcurve->curveflags |= CURVE_KN_INTERPOLATE;
-    if (primitive == CURVE_RIBBONS)
+    if (curve_shape == CURVE_RIBBON) {
       kcurve->curveflags |= CURVE_KN_RIBBONS;
-
-    if (line_method == CURVE_ACCURATE)
-      kcurve->curveflags |= CURVE_KN_ACCURATE;
-    else if (line_method == CURVE_CORRECTED)
-      kcurve->curveflags |= CURVE_KN_INTERSECTCORRECTION;
-
-    if (use_tangent_normal_geometry)
-      kcurve->curveflags |= CURVE_KN_TRUETANGENTGNORMAL;
-    if (use_backfacing)
+    }
+    else if (use_backfacing) {
       kcurve->curveflags |= CURVE_KN_BACKFACING;
-    if (use_encasing)
-      kcurve->curveflags |= CURVE_KN_ENCLOSEFILTER;
+    }
 
     kcurve->subdivisions = subdivisions;
   }
@@ -150,23 +134,14 @@ void CurveSystemManager::device_free(Device * /*device*/, DeviceScene * /*dscene
 
 bool CurveSystemManager::modified(const CurveSystemManager &CurveSystemManager)
 {
-  return !(
-      curve_shape == CurveSystemManager.curve_shape &&
-      line_method == CurveSystemManager.line_method && primitive == CurveSystemManager.primitive &&
-      use_encasing == CurveSystemManager.use_encasing &&
-      use_tangent_normal_geometry == CurveSystemManager.use_tangent_normal_geometry &&
-      use_backfacing == CurveSystemManager.use_backfacing &&
-      triangle_method == CurveSystemManager.triangle_method &&
-      resolution == CurveSystemManager.resolution && use_curves == CurveSystemManager.use_curves &&
-      subdivisions == CurveSystemManager.subdivisions);
+  return !(use_backfacing == CurveSystemManager.use_backfacing &&
+           use_curves == CurveSystemManager.use_curves &&
+           subdivisions == CurveSystemManager.subdivisions);
 }
 
 bool CurveSystemManager::modified_mesh(const CurveSystemManager &CurveSystemManager)
 {
-  return !(
-      primitive == CurveSystemManager.primitive && curve_shape == CurveSystemManager.curve_shape &&
-      triangle_method == CurveSystemManager.triangle_method &&
-      resolution == CurveSystemManager.resolution && use_curves == CurveSystemManager.use_curves);
+  return !(use_curves == CurveSystemManager.use_curves);
 }
 
 void CurveSystemManager::tag_update(Scene * /*scene*/)
diff --git a/intern/cycles/render/curves.h b/intern/cycles/render/curves.h
index ade289a402e..047db32b44b 100644
--- a/intern/cycles/render/curves.h
+++ b/intern/cycles/render/curves.h
@@ -29,15 +29,6 @@ class Scene;
 
 void curvebounds(float *lower, float *upper, float3 *p, int dim);
 
-typedef enum CurvePrimitiveType {
-  CURVE_TRIANGLES = 0,
-  CURVE_LINE_SEGMENTS = 1,
-  CURVE_SEGMENTS = 2,
-  CURVE_RIBBONS = 3,
-
-  CURVE_NUM_PRIMITIVE_TYPES,
-} CurvePrimitiveType;
-
 typedef enum CurveShapeType {
   CURVE_RIBBON = 0,
   CURVE_THICK = 1,
@@ -45,17 +36,6 @@ typedef enum CurveShapeType {
   CURVE_NUM_SHAPE_TYPES,
 } CurveShapeType;
 
-typedef enum CurveTriangleMethod {
-  CURVE_CAMERA_TRIANGLES,
-  CURVE_TESSELATED_TRIANGLES
-} CurveTriangleMethod;
-
-typedef enum CurveLineMethod {
-  CURVE_ACCURATE,
-  CURVE_CORRECTED,
-  CURVE_UNCORRECTED
-} CurveLineMethod;
-
 class ParticleCurveData {
 
  public:
@@ -85,17 +65,11 @@ class ParticleCurveData {
 
 class CurveSystemManager {
  public:
-  CurvePrimitiveType primitive;
   CurveShapeType curve_shape;
-  CurveLineMethod line_method;
-  CurveTriangleMethod triangle_method;
-  int resolution;
   int subdivisions;
 
   bool use_curves;
-  bool use_encasing;
   bool use_backfacing;
-  bool use_tangent_normal_geometry;
 
   bool need_update;
   bool need_mesh_update;