Merge branch 'master' into tmp-widget-optitmp-widget-opti

# Conflicts:
#	source/blender/gpu/shaders/gpu_shader_2D_widget_base_vert.glsl

72 files changed, 1363 insertions, 3169 deletions

diff --git a/intern/cycles/blender/addon/__init__.py b/intern/cycles/blender/addon/__init__.py
index 3d2a52d0cf6..3ab352e52a2 100644
--- a/intern/cycles/blender/addon/__init__.py
+++ b/intern/cycles/blender/addon/__init__.py
@@ -82,8 +82,8 @@ class CyclesRender(bpy.types.RenderEngine):
     def render(self, depsgraph):
         engine.render(self, depsgraph)
 
-    def bake(self, depsgraph, obj, pass_type, pass_filter, object_id, pixel_array, num_pixels, depth, result):
-        engine.bake(self, depsgraph, obj, pass_type, pass_filter, object_id, pixel_array, num_pixels, depth, result)
+    def bake(self, depsgraph, obj, pass_type, pass_filter, width, height):
+        engine.bake(self, depsgraph, obj, pass_type, pass_filter, width, height)
 
     # viewport render
     def view_update(self, context, depsgraph):
diff --git a/intern/cycles/blender/addon/engine.py b/intern/cycles/blender/addon/engine.py
index a1b063430f5..e7ea5e7a1f6 100644
--- a/intern/cycles/blender/addon/engine.py
+++ b/intern/cycles/blender/addon/engine.py
@@ -168,11 +168,11 @@ def render(engine, depsgraph):
         _cycles.render(engine.session, depsgraph.as_pointer())
 
 
-def bake(engine, depsgraph, obj, pass_type, pass_filter, object_id, pixel_array, num_pixels, depth, result):
+def bake(engine, depsgraph, obj, pass_type, pass_filter, width, height):
     import _cycles
     session = getattr(engine, "session", None)
     if session is not None:
-        _cycles.bake(engine.session, depsgraph.as_pointer(), obj.as_pointer(), pass_type, pass_filter, object_id, pixel_array.as_pointer(), num_pixels, depth, result.as_pointer())
+        _cycles.bake(engine.session, depsgraph.as_pointer(), obj.as_pointer(), pass_type, pass_filter, width, height)
 
 
 def reset(engine, data, depsgraph):
diff --git a/intern/cycles/blender/addon/ui.py b/intern/cycles/blender/addon/ui.py
index 08641c05941..a15daee706f 100644
--- a/intern/cycles/blender/addon/ui.py
+++ b/intern/cycles/blender/addon/ui.py
@@ -1907,10 +1907,15 @@ class CYCLES_RENDER_PT_bake_selected_to_active(CyclesButtonsPanel, Panel):
 
         col.prop(cbk, "use_cage", text="Cage")
         if cbk.use_cage:
-            col.prop(cbk, "cage_extrusion", text="Extrusion")
-            col.prop(cbk, "cage_object", text="Cage Object")
+            col.prop(cbk, "cage_object")
+            col = layout.column()
+            col.prop(cbk, "cage_extrusion")
+            col.active = cbk.cage_object is None
         else:
-            col.prop(cbk, "cage_extrusion", text="Ray Distance")
+            col.prop(cbk, "cage_extrusion", text="Extrusion")
+
+        col = layout.column()
+        col.prop(cbk, "max_ray_distance")
 
 
 class CYCLES_RENDER_PT_bake_output(CyclesButtonsPanel, Panel):
diff --git a/intern/cycles/blender/blender_object.cpp b/intern/cycles/blender/blender_object.cpp
index 4b29c28913b..a461982a538 100644
--- a/intern/cycles/blender/blender_object.cpp
+++ b/intern/cycles/blender/blender_object.cpp
@@ -460,9 +460,12 @@ void BlenderSync::sync_motion(BL::RenderSettings &b_render,
     sync_objects(b_depsgraph, b_v3d, 0.0f);
   }
 
-  /* always sample these times for camera motion */
-  motion_times.insert(-1.0f);
-  motion_times.insert(1.0f);
+  /* Insert motion times from camera. Motion times from other objects
+   * have already been added in a sync_objects call. */
+  uint camera_motion_steps = object_motion_steps(b_cam, b_cam);
+  for (size_t step = 0; step < camera_motion_steps; step++) {
+    motion_times.insert(scene->camera->motion_time(step));
+  }
 
   /* note iteration over motion_times set happens in sorted order */
   foreach (float relative_time, motion_times) {
@@ -487,10 +490,8 @@ void BlenderSync::sync_motion(BL::RenderSettings &b_render,
     b_engine.frame_set(frame, subframe);
     python_thread_state_save(python_thread_state);
 
-    /* sync camera, only supports two times at the moment */
-    if (relative_time == -1.0f || relative_time == 1.0f) {
-      sync_camera_motion(b_render, b_cam, width, height, relative_time);
-    }
+    /* Syncs camera motion if relative_time is one of the camera's motion times. */
+    sync_camera_motion(b_render, b_cam, width, height, relative_time);
 
     /* sync object */
     sync_objects(b_depsgraph, b_v3d, relative_time);
diff --git a/intern/cycles/blender/blender_python.cpp b/intern/cycles/blender/blender_python.cpp
index 8c7c0bc1daa..79c16856462 100644
--- a/intern/cycles/blender/blender_python.cpp
+++ b/intern/cycles/blender/blender_python.cpp
@@ -298,22 +298,18 @@ static PyObject *render_func(PyObject * /*self*/, PyObject *args)
 static PyObject *bake_func(PyObject * /*self*/, PyObject *args)
 {
   PyObject *pysession, *pydepsgraph, *pyobject;
-  PyObject *pypixel_array, *pyresult;
   const char *pass_type;
-  int num_pixels, depth, object_id, pass_filter;
+  int pass_filter, width, height;
 
   if (!PyArg_ParseTuple(args,
-                        "OOOsiiOiiO",
+                        "OOOsiii",
                         &pysession,
                         &pydepsgraph,
                         &pyobject,
                         &pass_type,
                         &pass_filter,
-                        &object_id,
-                        &pypixel_array,
-                        &num_pixels,
-                        &depth,
-                        &pyresult))
+                        &width,
+                        &height))
     return NULL;
 
   BlenderSession *session = (BlenderSession *)PyLong_AsVoidPtr(pysession);
@@ -326,23 +322,9 @@ static PyObject *bake_func(PyObject * /*self*/, PyObject *args)
   RNA_id_pointer_create((ID *)PyLong_AsVoidPtr(pyobject), &objectptr);
   BL::Object b_object(objectptr);
 
-  void *b_result = PyLong_AsVoidPtr(pyresult);
-
-  PointerRNA bakepixelptr;
-  RNA_pointer_create(NULL, &RNA_BakePixel, PyLong_AsVoidPtr(pypixel_array), &bakepixelptr);
-  BL::BakePixel b_bake_pixel(bakepixelptr);
-
   python_thread_state_save(&session->python_thread_state);
 
-  session->bake(b_depsgraph,
-                b_object,
-                pass_type,
-                pass_filter,
-                object_id,
-                b_bake_pixel,
-                (size_t)num_pixels,
-                depth,
-                (float *)b_result);
+  session->bake(b_depsgraph, b_object, pass_type, pass_filter, width, height);
 
   python_thread_state_restore(&session->python_thread_state);
 
diff --git a/intern/cycles/blender/blender_session.cpp b/intern/cycles/blender/blender_session.cpp
index 5ea96d6bdfd..31b09695632 100644
--- a/intern/cycles/blender/blender_session.cpp
+++ b/intern/cycles/blender/blender_session.cpp
@@ -247,9 +247,7 @@ void BlenderSession::reset_session(BL::BlendData &b_data, BL::Depsgraph &b_depsg
 
 void BlenderSession::free_session()
 {
-  if (sync)
-    delete sync;
-
+  delete sync;
   delete session;
 }
 
@@ -317,6 +315,7 @@ static void end_render_result(BL::RenderEngine &b_engine,
 
 void BlenderSession::do_write_update_render_tile(RenderTile &rtile,
                                                  bool do_update_only,
+                                                 bool do_read_only,
                                                  bool highlight)
 {
   int x = rtile.x - session->tile_manager.params.full_x;
@@ -342,7 +341,23 @@ void BlenderSession::do_write_update_render_tile(RenderTile &rtile,
 
   BL::RenderLayer b_rlay = *b_single_rlay;
 
-  if (do_update_only) {
+  if (do_read_only) {
+    /* copy each pass */
+    BL::RenderLayer::passes_iterator b_iter;
+
+    for (b_rlay.passes.begin(b_iter); b_iter != b_rlay.passes.end(); ++b_iter) {
+      BL::RenderPass b_pass(*b_iter);
+
+      /* find matching pass type */
+      PassType pass_type = BlenderSync::get_pass_type(b_pass);
+      int components = b_pass.channels();
+
+      rtile.buffers->set_pass_rect(pass_type, components, (float *)b_pass.rect());
+    }
+
+    end_render_result(b_engine, b_rr, false, false, false);
+  }
+  else if (do_update_only) {
     /* Sample would be zero at initial tile update, which is only needed
      * to tag tile form blender side as IN PROGRESS for proper highlight
      * no buffers should be sent to blender yet. For denoise we also
@@ -362,9 +377,14 @@ void BlenderSession::do_write_update_render_tile(RenderTile &rtile,
   }
 }
 
+void BlenderSession::read_render_tile(RenderTile &rtile)
+{
+  do_write_update_render_tile(rtile, false, true, false);
+}
+
 void BlenderSession::write_render_tile(RenderTile &rtile)
 {
-  do_write_update_render_tile(rtile, false, false);
+  do_write_update_render_tile(rtile, false, false, false);
 }
 
 void BlenderSession::update_render_tile(RenderTile &rtile, bool highlight)
@@ -374,9 +394,9 @@ void BlenderSession::update_render_tile(RenderTile &rtile, bool highlight)
    * would need to be investigated a bit further, but for now shall be fine
    */
   if (!b_engine.is_preview())
-    do_write_update_render_tile(rtile, true, highlight);
+    do_write_update_render_tile(rtile, true, false, highlight);
   else
-    do_write_update_render_tile(rtile, false, false);
+    do_write_update_render_tile(rtile, false, false, false);
 }
 
 static void add_cryptomatte_layer(BL::RenderResult &b_rr, string name, string manifest)
@@ -593,25 +613,6 @@ void BlenderSession::render(BL::Depsgraph &b_depsgraph_)
 #endif
 }
 
-static void populate_bake_data(BakeData *data,
-                               const int object_id,
-                               BL::BakePixel &pixel_array,
-                               const int num_pixels)
-{
-  BL::BakePixel bp = pixel_array;
-
-  int i;
-  for (i = 0; i < num_pixels; i++) {
-    if (bp.object_id() == object_id) {
-      data->set(i, bp.primitive_id(), bp.uv(), bp.du_dx(), bp.du_dy(), bp.dv_dx(), bp.dv_dy());
-    }
-    else {
-      data->set_null(i);
-    }
-    bp = bp.next();
-  }
-}
-
 static int bake_pass_filter_get(const int pass_filter)
 {
   int flag = BAKE_FILTER_NONE;
@@ -642,43 +643,26 @@ void BlenderSession::bake(BL::Depsgraph &b_depsgraph_,
                           BL::Object &b_object,
                           const string &pass_type,
                           const int pass_filter,
-                          const int object_id,
-                          BL::BakePixel &pixel_array,
-                          const size_t num_pixels,
-                          const int /*depth*/,
-                          float result[])
+                          const int bake_width,
+                          const int bake_height)
 {
   b_depsgraph = b_depsgraph_;
 
   ShaderEvalType shader_type = get_shader_type(pass_type);
-
-  /* Set baking flag in advance, so kernel loading can check if we need
-   * any baking capabilities.
-   */
-  scene->bake_manager->set_baking(true);
-
-  /* ensure kernels are loaded before we do any scene updates */
-  session->load_kernels();
-
-  if (shader_type == SHADER_EVAL_UV) {
-    /* force UV to be available */
-    Pass::add(PASS_UV, scene->film->passes);
-  }
-
   int bake_pass_filter = bake_pass_filter_get(pass_filter);
-  bake_pass_filter = BakeManager::shader_type_to_pass_filter(shader_type, bake_pass_filter);
 
-  /* force use_light_pass to be true if we bake more than just colors */
-  if (bake_pass_filter & ~BAKE_FILTER_COLOR) {
-    Pass::add(PASS_LIGHT, scene->film->passes);
-  }
+  /* Initialize bake manager, before we load the baking kernels. */
+  scene->bake_manager->set(scene, b_object.name(), shader_type, bake_pass_filter);
 
-  /* create device and update scene */
-  scene->film->tag_update(scene);
-  scene->integrator->tag_update(scene);
+  /* Passes are identified by name, so in order to return the combined pass we need to set the
+   * name. */
+  Pass::add(PASS_COMBINED, scene->film->passes, "Combined");
+
+  session->read_bake_tile_cb = function_bind(&BlenderSession::read_render_tile, this, _1);
+  session->write_render_tile_cb = function_bind(&BlenderSession::write_render_tile, this, _1);
 
   if (!session->progress.get_cancel()) {
-    /* update scene */
+    /* Sync scene. */
     BL::Object b_camera_override(b_engine.camera_override());
     sync->sync_camera(b_render, b_camera_override, width, height, "");
     sync->sync_data(
@@ -686,75 +670,43 @@ void BlenderSession::bake(BL::Depsgraph &b_depsgraph_,
     builtin_images_load();
   }
 
-  BakeData *bake_data = NULL;
+  /* Object might have been disabled for rendering or excluded in some
+   * other way, in that case Blender will report a warning afterwards. */
+  bool object_found = false;
+  foreach (Object *ob, scene->objects) {
+    if (ob->name == b_object.name()) {
+      object_found = true;
+      break;
+    }
+  }
 
-  if (!session->progress.get_cancel()) {
-    /* get buffer parameters */
+  if (object_found && !session->progress.get_cancel()) {
+    /* Get session and buffer parameters. */
     SessionParams session_params = BlenderSync::get_session_params(
         b_engine, b_userpref, b_scene, background);
-    BufferParams buffer_params = BlenderSync::get_buffer_params(
-        b_scene, b_render, b_v3d, b_rv3d, scene->camera, width, height);
+    session_params.progressive_refine = false;
 
-    scene->bake_manager->set_shader_limit((size_t)b_engine.tile_x(), (size_t)b_engine.tile_y());
+    BufferParams buffer_params;
+    buffer_params.width = bake_width;
+    buffer_params.height = bake_height;
+    buffer_params.passes = scene->film->passes;
 
-    /* set number of samples */
+    /* Update session. */
     session->tile_manager.set_samples(session_params.samples);
     session->reset(buffer_params, session_params.samples);
-    session->update_scene();
-
-    /* find object index. todo: is arbitrary - copied from mesh_displace.cpp */
-    size_t object_index = OBJECT_NONE;
-    int tri_offset = 0;
-
-    for (size_t i = 0; i < scene->objects.size(); i++) {
-      const Object *object = scene->objects[i];
-      const Geometry *geom = object->geometry;
-      if (object->name == b_object.name() && geom->type == Geometry::MESH) {
-        const Mesh *mesh = static_cast<const Mesh *>(geom);
-        object_index = i;
-        tri_offset = mesh->prim_offset;
-        break;
-      }
-    }
-
-    /* Object might have been disabled for rendering or excluded in some
-     * other way, in that case Blender will report a warning afterwards. */
-    if (object_index != OBJECT_NONE) {
-      int object = object_index;
-
-      bake_data = scene->bake_manager->init(object, tri_offset, num_pixels);
-      populate_bake_data(bake_data, object_id, pixel_array, num_pixels);
-    }
-
-    /* set number of samples */
-    session->tile_manager.set_samples(session_params.samples);
-    session->reset(buffer_params, session_params.samples);
-    session->update_scene();
 
     session->progress.set_update_callback(
         function_bind(&BlenderSession::update_bake_progress, this));
   }
 
   /* Perform bake. Check cancel to avoid crash with incomplete scene data. */
-  if (!session->progress.get_cancel() && bake_data) {
-    scene->bake_manager->bake(scene->device,
-                              &scene->dscene,
-                              scene,
-                              session->progress,
-                              shader_type,
-                              bake_pass_filter,
-                              bake_data,
-                              result);
+  if (object_found && !session->progress.get_cancel()) {
+    session->start();
+    session->wait();
   }
 
-  /* free all memory used (host and device), so we wouldn't leave render
-   * engine with extra memory allocated
-   */
-
-  session->device_free();
-
-  delete sync;
-  sync = NULL;
+  session->read_bake_tile_cb = function_null;
+  session->write_render_tile_cb = function_null;
 }
 
 void BlenderSession::do_write_update_render_result(BL::RenderLayer &b_rlay,
diff --git a/intern/cycles/blender/blender_session.h b/intern/cycles/blender/blender_session.h
index 3e6498bb655..34e952e312b 100644
--- a/intern/cycles/blender/blender_session.h
+++ b/intern/cycles/blender/blender_session.h
@@ -66,14 +66,12 @@ class BlenderSession {
             BL::Object &b_object,
             const string &pass_type,
             const int custom_flag,
-            const int object_id,
-            BL::BakePixel &pixel_array,
-            const size_t num_pixels,
-            const int depth,
-            float pixels[]);
+            const int bake_width,
+            const int bake_height);
 
   void write_render_result(BL::RenderLayer &b_rlay, RenderTile &rtile);
   void write_render_tile(RenderTile &rtile);
+  void read_render_tile(RenderTile &rtile);
 
   /* update functions are used to update display buffer only after sample was rendered
    * only needed for better visual feedback */
@@ -155,7 +153,10 @@ class BlenderSession {
   void do_write_update_render_result(BL::RenderLayer &b_rlay,
                                      RenderTile &rtile,
                                      bool do_update_only);
-  void do_write_update_render_tile(RenderTile &rtile, bool do_update_only, bool highlight);
+  void do_write_update_render_tile(RenderTile &rtile,
+                                   bool do_update_only,
+                                   bool do_read_only,
+                                   bool highlight);
 
   void builtin_images_load();
 
diff --git a/intern/cycles/blender/blender_sync.cpp b/intern/cycles/blender/blender_sync.cpp
index e8031be7dd1..f16305e737d 100644
--- a/intern/cycles/blender/blender_sync.cpp
+++ b/intern/cycles/blender/blender_sync.cpp
@@ -481,6 +481,9 @@ PassType BlenderSync::get_pass_type(BL::RenderPass &b_pass)
   MAP_PASS("AO", PASS_AO);
   MAP_PASS("Shadow", PASS_SHADOW);
 
+  MAP_PASS("BakePrimitive", PASS_BAKE_PRIMITIVE);
+  MAP_PASS("BakeDifferential", PASS_BAKE_DIFFERENTIAL);
+
 #ifdef __KERNEL_DEBUG__
   MAP_PASS("Debug BVH Traversed Nodes", PASS_BVH_TRAVERSED_NODES);
   MAP_PASS("Debug BVH Traversed Instances", PASS_BVH_TRAVERSED_INSTANCES);
diff --git a/intern/cycles/blender/blender_volume.cpp b/intern/cycles/blender/blender_volume.cpp
index 6254a1a1b24..4eed6be8c7c 100644
--- a/intern/cycles/blender/blender_volume.cpp
+++ b/intern/cycles/blender/blender_volume.cpp
@@ -34,15 +34,13 @@ CCL_NAMESPACE_BEGIN
 /* TODO: verify this is not loading unnecessary attributes. */
 class BlenderSmokeLoader : public ImageLoader {
  public:
-  BlenderSmokeLoader(const BL::Object &b_ob, AttributeStandard attribute)
-      : b_ob(b_ob), attribute(attribute)
+  BlenderSmokeLoader(BL::Object &b_ob, AttributeStandard attribute)
+      : b_domain(object_fluid_gas_domain_find(b_ob)), b_mesh(b_ob.data()), attribute(attribute)
   {
   }
 
   bool load_metadata(ImageMetaData &metadata) override
   {
-    BL::FluidDomainSettings b_domain = object_fluid_gas_domain_find(b_ob);
-
     if (!b_domain) {
       return false;
     }
@@ -79,7 +77,6 @@ class BlenderSmokeLoader : public ImageLoader {
     /* Create a matrix to transform from object space to mesh texture space.
      * This does not work with deformations but that can probably only be done
      * well with a volume grid mapping of coordinates. */
-    BL::Mesh b_mesh(b_ob.data());
     float3 loc, size;
     mesh_texture_space(b_mesh, loc, size);
     metadata.transform_3d = transform_translate(-loc) * transform_scale(size);
@@ -90,9 +87,6 @@ class BlenderSmokeLoader : public ImageLoader {
 
   bool load_pixels(const ImageMetaData &, void *pixels, const size_t, const bool) override
   {
-    /* smoke volume data */
-    BL::FluidDomainSettings b_domain = object_fluid_gas_domain_find(b_ob);
-
     if (!b_domain) {
       return false;
     }
@@ -179,10 +173,11 @@ class BlenderSmokeLoader : public ImageLoader {
   bool equals(const ImageLoader &other) const override
   {
     const BlenderSmokeLoader &other_loader = (const BlenderSmokeLoader &)other;
-    return b_ob == other_loader.b_ob && attribute == other_loader.attribute;
+    return b_domain == other_loader.b_domain && attribute == other_loader.attribute;
   }
 
-  BL::Object b_ob;
+  BL::FluidDomainSettings b_domain;
+  BL::Mesh b_mesh;
   AttributeStandard attribute;
 };
 
diff --git a/intern/cycles/device/cuda/device_cuda.h b/intern/cycles/device/cuda/device_cuda.h
index 3e397da895b..3f23f0fe4c5 100644
--- a/intern/cycles/device/cuda/device_cuda.h
+++ b/intern/cycles/device/cuda/device_cuda.h
@@ -223,7 +223,7 @@ class CUDADevice : public Device {
                               CUdeviceptr d_wtile,
                               CUstream stream = 0);
 
-  void path_trace(DeviceTask &task, RenderTile &rtile, device_vector<WorkTile> &work_tiles);
+  void render(DeviceTask &task, RenderTile &rtile, device_vector<WorkTile> &work_tiles);
 
   void film_convert(DeviceTask &task,
                     device_ptr buffer,
diff --git a/intern/cycles/device/cuda/device_cuda_impl.cpp b/intern/cycles/device/cuda/device_cuda_impl.cpp
index ba5d479e0e7..acf53c3eb1b 100644
--- a/intern/cycles/device/cuda/device_cuda_impl.cpp
+++ b/intern/cycles/device/cuda/device_cuda_impl.cpp
@@ -586,20 +586,23 @@ void CUDADevice::reserve_local_memory(const DeviceRequestedFeatures &requested_f
   cuMemGetInfo(&free_before, &total);
 
   /* Get kernel function. */
-  CUfunction cuPathTrace;
+  CUfunction cuRender;
 
-  if (requested_features.use_integrator_branched) {
-    cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_branched_path_trace"));
+  if (requested_features.use_baking) {
+    cuda_assert(cuModuleGetFunction(&cuRender, cuModule, "kernel_cuda_bake"));
+  }
+  else if (requested_features.use_integrator_branched) {
+    cuda_assert(cuModuleGetFunction(&cuRender, cuModule, "kernel_cuda_branched_path_trace"));
   }
   else {
-    cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_path_trace"));
+    cuda_assert(cuModuleGetFunction(&cuRender, cuModule, "kernel_cuda_path_trace"));
   }
 
-  cuda_assert(cuFuncSetCacheConfig(cuPathTrace, CU_FUNC_CACHE_PREFER_L1));
+  cuda_assert(cuFuncSetCacheConfig(cuRender, CU_FUNC_CACHE_PREFER_L1));
 
   int min_blocks, num_threads_per_block;
-  cuda_assert(cuOccupancyMaxPotentialBlockSize(
-      &min_blocks, &num_threads_per_block, cuPathTrace, NULL, 0, 0));
+  cuda_assert(
+      cuOccupancyMaxPotentialBlockSize(&min_blocks, &num_threads_per_block, cuRender, NULL, 0, 0));
 
   /* Launch kernel, using just 1 block appears sufficient to reserve
    * memory for all multiprocessors. It would be good to do this in
@@ -609,7 +612,7 @@ void CUDADevice::reserve_local_memory(const DeviceRequestedFeatures &requested_f
 
   void *args[] = {&d_work_tiles, &total_work_size};
 
-  cuda_assert(cuLaunchKernel(cuPathTrace, 1, 1, 1, num_threads_per_block, 1, 1, 0, 0, args, 0));
+  cuda_assert(cuLaunchKernel(cuRender, 1, 1, 1, num_threads_per_block, 1, 1, 0, 0, args, 0));
 
   cuda_assert(cuCtxSynchronize());
 
@@ -1780,9 +1783,7 @@ void CUDADevice::adaptive_sampling_post(RenderTile &rtile,
                              0));
 }
 
-void CUDADevice::path_trace(DeviceTask &task,
-                            RenderTile &rtile,
-                            device_vector<WorkTile> &work_tiles)
+void CUDADevice::render(DeviceTask &task, RenderTile &rtile, device_vector<WorkTile> &work_tiles)
 {
   scoped_timer timer(&rtile.buffers->render_time);
 
@@ -1790,21 +1791,24 @@ void CUDADevice::path_trace(DeviceTask &task,
     return;
 
   CUDAContextScope scope(this);
-  CUfunction cuPathTrace;
+  CUfunction cuRender;
 
   /* Get kernel function. */
-  if (task.integrator_branched) {
-    cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_branched_path_trace"));
+  if (rtile.task == RenderTile::BAKE) {
+    cuda_assert(cuModuleGetFunction(&cuRender, cuModule, "kernel_cuda_bake"));
+  }
+  else if (task.integrator_branched) {
+    cuda_assert(cuModuleGetFunction(&cuRender, cuModule, "kernel_cuda_branched_path_trace"));
   }
   else {
-    cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_path_trace"));
+    cuda_assert(cuModuleGetFunction(&cuRender, cuModule, "kernel_cuda_path_trace"));
   }
 
   if (have_error()) {
     return;
   }
 
-  cuda_assert(cuFuncSetCacheConfig(cuPathTrace, CU_FUNC_CACHE_PREFER_L1));
+  cuda_assert(cuFuncSetCacheConfig(cuRender, CU_FUNC_CACHE_PREFER_L1));
 
   /* Allocate work tile. */
   work_tiles.alloc(1);
@@ -1822,8 +1826,8 @@ void CUDADevice::path_trace(DeviceTask &task,
    * remain conservative for GPUs connected to a display to avoid driver
    * timeouts and display freezing. */
   int min_blocks, num_threads_per_block;
-  cuda_assert(cuOccupancyMaxPotentialBlockSize(
-      &min_blocks, &num_threads_per_block, cuPathTrace, NULL, 0, 0));
+  cuda_assert(
+      cuOccupancyMaxPotentialBlockSize(&min_blocks, &num_threads_per_block, cuRender, NULL, 0, 0));
   if (!info.display_device) {
     min_blocks *= 8;
   }
@@ -1851,7 +1855,7 @@ void CUDADevice::path_trace(DeviceTask &task,
     void *args[] = {&d_work_tiles, &total_work_size};
 
     cuda_assert(
-        cuLaunchKernel(cuPathTrace, num_blocks, 1, 1, num_threads_per_block, 1, 1, 0, 0, args, 0));
+        cuLaunchKernel(cuRender, num_blocks, 1, 1, num_threads_per_block, 1, 1, 0, 0, args, 0));
 
     /* Run the adaptive sampling kernels at selected samples aligned to step samples. */
     uint filter_sample = sample + wtile->num_samples - 1;
@@ -1957,10 +1961,7 @@ void CUDADevice::shader(DeviceTask &task)
   CUdeviceptr d_output = (CUdeviceptr)task.shader_output;
 
   /* get kernel function */
-  if (task.shader_eval_type >= SHADER_EVAL_BAKE) {
-    cuda_assert(cuModuleGetFunction(&cuShader, cuModule, "kernel_cuda_bake"));
-  }
-  else if (task.shader_eval_type == SHADER_EVAL_DISPLACE) {
+  if (task.shader_eval_type == SHADER_EVAL_DISPLACE) {
     cuda_assert(cuModuleGetFunction(&cuShader, cuModule, "kernel_cuda_displace"));
   }
   else {
@@ -2297,9 +2298,12 @@ void CUDADevice::thread_run(DeviceTask *task)
           split_kernel->path_trace(task, tile, void_buffer, void_buffer);
         }
         else {
-          path_trace(*task, tile, work_tiles);
+          render(*task, tile, work_tiles);
         }
       }
+      else if (tile.task == RenderTile::BAKE) {
+        render(*task, tile, work_tiles);
+      }
       else if (tile.task == RenderTile::DENOISE) {
         tile.sample = tile.start_sample + tile.num_samples;
 
diff --git a/intern/cycles/device/device_cpu.cpp b/intern/cycles/device/device_cpu.cpp
index c701c14318f..fc6febd8cee 100644
--- a/intern/cycles/device/device_cpu.cpp
+++ b/intern/cycles/device/device_cpu.cpp
@@ -188,6 +188,7 @@ class CPUDevice : public Device {
       convert_to_byte_kernel;
   KernelFunctions<void (*)(KernelGlobals *, uint4 *, float4 *, int, int, int, int, int)>
       shader_kernel;
+  KernelFunctions<void (*)(KernelGlobals *, float *, int, int, int, int, int)> bake_kernel;
 
   KernelFunctions<void (*)(
       int, TileInfo *, int, int, float *, float *, float *, float *, float *, int *, int, int)>
@@ -270,6 +271,7 @@ class CPUDevice : public Device {
         REGISTER_KERNEL(convert_to_half_float),
         REGISTER_KERNEL(convert_to_byte),
         REGISTER_KERNEL(shader),
+        REGISTER_KERNEL(bake),
         REGISTER_KERNEL(filter_divide_shadow),
         REGISTER_KERNEL(filter_get_feature),
         REGISTER_KERNEL(filter_write_feature),
@@ -895,7 +897,7 @@ class CPUDevice : public Device {
     }
   }
 
-  void path_trace(DeviceTask &task, RenderTile &tile, KernelGlobals *kg)
+  void render(DeviceTask &task, RenderTile &tile, KernelGlobals *kg)
   {
     const bool use_coverage = kernel_data.film.cryptomatte_passes & CRYPT_ACCURATE;
 
@@ -919,12 +921,21 @@ class CPUDevice : public Device {
           break;
       }
 
-      for (int y = tile.y; y < tile.y + tile.h; y++) {
-        for (int x = tile.x; x < tile.x + tile.w; x++) {
-          if (use_coverage) {
-            coverage.init_pixel(x, y);
+      if (tile.task == RenderTile::PATH_TRACE) {
+        for (int y = tile.y; y < tile.y + tile.h; y++) {
+          for (int x = tile.x; x < tile.x + tile.w; x++) {
+            if (use_coverage) {
+              coverage.init_pixel(x, y);
+            }
+            path_trace_kernel()(kg, render_buffer, sample, x, y, tile.offset, tile.stride);
+          }
+        }
+      }
+      else {
+        for (int y = tile.y; y < tile.y + tile.h; y++) {
+          for (int x = tile.x; x < tile.x + tile.w; x++) {
+            bake_kernel()(kg, render_buffer, sample, x, y, tile.offset, tile.stride);
           }
-          path_trace_kernel()(kg, render_buffer, sample, x, y, tile.offset, tile.stride);
         }
       }
       tile.sample = sample + 1;
@@ -1019,9 +1030,12 @@ class CPUDevice : public Device {
           split_kernel->path_trace(&task, tile, kgbuffer, void_buffer);
         }
         else {
-          path_trace(task, tile, kg);
+          render(task, tile, kg);
         }
       }
+      else if (tile.task == RenderTile::BAKE) {
+        render(task, tile, kg);
+      }
       else if (tile.task == RenderTile::DENOISE) {
         denoise(denoising, tile);
         task.update_progress(&tile, tile.w * tile.h);
diff --git a/intern/cycles/device/opencl/device_opencl.h b/intern/cycles/device/opencl/device_opencl.h
index d6f4fb43061..389268e1c2a 100644
--- a/intern/cycles/device/opencl/device_opencl.h
+++ b/intern/cycles/device/opencl/device_opencl.h
@@ -451,6 +451,7 @@ class OpenCLDevice : public Device {
                     device_ptr rgba_half);
   void shader(DeviceTask &task);
   void update_adaptive(DeviceTask &task, RenderTile &tile, int sample);
+  void bake(DeviceTask &task, RenderTile &tile);
 
   void denoise(RenderTile &tile, DenoisingTask &denoising);
 
diff --git a/intern/cycles/device/opencl/device_opencl_impl.cpp b/intern/cycles/device/opencl/device_opencl_impl.cpp
index 2766f85d17c..beb3174b111 100644
--- a/intern/cycles/device/opencl/device_opencl_impl.cpp
+++ b/intern/cycles/device/opencl/device_opencl_impl.cpp
@@ -1367,6 +1367,9 @@ void OpenCLDevice::thread_run(DeviceTask *task)
          */
         clFinish(cqCommandQueue);
       }
+      else if (tile.task == RenderTile::BAKE) {
+        bake(*task, tile);
+      }
       else if (tile.task == RenderTile::DENOISE) {
         tile.sample = tile.start_sample + tile.num_samples;
         denoise(tile, denoising);
@@ -1858,10 +1861,7 @@ void OpenCLDevice::shader(DeviceTask &task)
   cl_int d_offset = task.offset;
 
   OpenCLDevice::OpenCLProgram *program = &background_program;
-  if (task.shader_eval_type >= SHADER_EVAL_BAKE) {
-    program = &bake_program;
-  }
-  else if (task.shader_eval_type == SHADER_EVAL_DISPLACE) {
+  if (task.shader_eval_type == SHADER_EVAL_DISPLACE) {
     program = &displace_program;
   }
   program->wait_for_availability();
@@ -1892,6 +1892,51 @@ void OpenCLDevice::shader(DeviceTask &task)
   }
 }
 
+void OpenCLDevice::bake(DeviceTask &task, RenderTile &rtile)
+{
+  scoped_timer timer(&rtile.buffers->render_time);
+
+  /* Cast arguments to cl types. */
+  cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
+  cl_mem d_buffer = CL_MEM_PTR(rtile.buffer);
+  cl_int d_x = rtile.x;
+  cl_int d_y = rtile.y;
+  cl_int d_w = rtile.w;
+  cl_int d_h = rtile.h;
+  cl_int d_offset = rtile.offset;
+  cl_int d_stride = rtile.stride;
+
+  bake_program.wait_for_availability();
+  cl_kernel kernel = bake_program();
+
+  cl_uint start_arg_index = kernel_set_args(kernel, 0, d_data, d_buffer);
+
+  set_kernel_arg_buffers(kernel, &start_arg_index);
+
+  start_arg_index += kernel_set_args(
+      kernel, start_arg_index, d_x, d_y, d_w, d_h, d_offset, d_stride);
+
+  int start_sample = rtile.start_sample;
+  int end_sample = rtile.start_sample + rtile.num_samples;
+
+  for (int sample = start_sample; sample < end_sample; sample++) {
+    if (task.get_cancel()) {
+      if (task.need_finish_queue == false)
+        break;
+    }
+
+    kernel_set_args(kernel, start_arg_index, sample);
+
+    enqueue_kernel(kernel, d_w, d_h);
+
+    rtile.sample = sample + 1;
+
+    task.update_progress(&rtile, rtile.w * rtile.h);
+  }
+
+  clFinish(cqCommandQueue);
+}
+
 string OpenCLDevice::kernel_build_options(const string *debug_src)
 {
   string build_options = "-cl-no-signed-zeros -cl-mad-enable ";
diff --git a/intern/cycles/kernel/kernel_bake.h b/intern/cycles/kernel/kernel_bake.h
index f1fc697553a..2709a9da734 100644
--- a/intern/cycles/kernel/kernel_bake.h
+++ b/intern/cycles/kernel/kernel_bake.h
@@ -18,38 +18,40 @@ CCL_NAMESPACE_BEGIN
 
 #ifdef __BAKING__
 
-ccl_device_inline void compute_light_pass(
+ccl_device_noinline void compute_light_pass(
     KernelGlobals *kg, ShaderData *sd, PathRadiance *L, uint rng_hash, int pass_filter, int sample)
 {
   kernel_assert(kernel_data.film.use_light_pass);
 
-  PathRadiance L_sample;
-  PathState state;
-  Ray ray;
   float3 throughput = make_float3(1.0f, 1.0f, 1.0f);
 
-  /* emission and indirect shader data memory used by various functions */
-  ShaderData emission_sd, indirect_sd;
-
-  ray.P = sd->P + sd->Ng;
-  ray.D = -sd->Ng;
-  ray.t = FLT_MAX;
-#  ifdef __CAMERA_MOTION__
-  ray.time = 0.5f;
-#  endif
+  /* Emission and indirect shader data memory used by various functions. */
+  ShaderDataTinyStorage emission_sd_storage;
+  ShaderData *emission_sd = AS_SHADER_DATA(&emission_sd_storage);
+  ShaderData indirect_sd;
 
-  /* init radiance */
-  path_radiance_init(kg, &L_sample);
+  /* Init radiance. */
+  path_radiance_init(kg, L);
 
-  /* init path state */
-  path_state_init(kg, &emission_sd, &state, rng_hash, sample, NULL);
+  /* Init path state. */
+  PathState state;
+  path_state_init(kg, emission_sd, &state, rng_hash, sample, NULL);
 
-  /* evaluate surface shader */
+  /* Evaluate surface shader. */
   shader_eval_surface(kg, sd, &state, NULL, state.flag);
 
   /* TODO, disable more closures we don't need besides transparent */
   shader_bsdf_disable_transparency(kg, sd);
 
+  /* Init ray. */
+  Ray ray;
+  ray.P = sd->P + sd->Ng;
+  ray.D = -sd->Ng;
+  ray.t = FLT_MAX;
+#  ifdef __CAMERA_MOTION__
+  ray.time = 0.5f;
+#  endif
+
 #  ifdef __BRANCHED_PATH__
   if (!kernel_data.integrator.branched) {
     /* regular path tracer */
@@ -57,14 +59,13 @@ ccl_device_inline void compute_light_pass(
 
     /* sample ambient occlusion */
     if (pass_filter & BAKE_FILTER_AO) {
-      kernel_path_ao(
-          kg, sd, &emission_sd, &L_sample, &state, throughput, shader_bsdf_alpha(kg, sd));
+      kernel_path_ao(kg, sd, emission_sd, L, &state, throughput, shader_bsdf_alpha(kg, sd));
     }
 
     /* sample emission */
     if ((pass_filter & BAKE_FILTER_EMISSION) && (sd->flag & SD_EMISSION)) {
       float3 emission = indirect_primitive_emission(kg, sd, 0.0f, state.flag, state.ray_pdf);
-      path_radiance_accum_emission(kg, &L_sample, &state, throughput, emission);
+      path_radiance_accum_emission(kg, L, &state, throughput, emission);
     }
 
     bool is_sss_sample = false;
@@ -77,12 +78,10 @@ ccl_device_inline void compute_light_pass(
       SubsurfaceIndirectRays ss_indirect;
       kernel_path_subsurface_init_indirect(&ss_indirect);
       if (kernel_path_subsurface_scatter(
-              kg, sd, &emission_sd, &L_sample, &state, &ray, &throughput, &ss_indirect)) {
+              kg, sd, emission_sd, L, &state, &ray, &throughput, &ss_indirect)) {
         while (ss_indirect.num_rays) {
-          kernel_path_subsurface_setup_indirect(
-              kg, &ss_indirect, &state, &ray, &L_sample, &throughput);
-          kernel_path_indirect(
-              kg, &indirect_sd, &emission_sd, &ray, throughput, &state, &L_sample);
+          kernel_path_subsurface_setup_indirect(kg, &ss_indirect, &state, &ray, L, &throughput);
+          kernel_path_indirect(kg, &indirect_sd, emission_sd, &ray, throughput, &state, L);
         }
         is_sss_sample = true;
       }
@@ -91,18 +90,18 @@ ccl_device_inline void compute_light_pass(
 
     /* sample light and BSDF */
     if (!is_sss_sample && (pass_filter & (BAKE_FILTER_DIRECT | BAKE_FILTER_INDIRECT))) {
-      kernel_path_surface_connect_light(kg, sd, &emission_sd, throughput, &state, &L_sample);
+      kernel_path_surface_connect_light(kg, sd, emission_sd, throughput, &state, L);
 
-      if (kernel_path_surface_bounce(kg, sd, &throughput, &state, &L_sample.state, &ray)) {
+      if (kernel_path_surface_bounce(kg, sd, &throughput, &state, &L->state, &ray)) {
 #  ifdef __LAMP_MIS__
         state.ray_t = 0.0f;
 #  endif
         /* compute indirect light */
-        kernel_path_indirect(kg, &indirect_sd, &emission_sd, &ray, throughput, &state, &L_sample);
+        kernel_path_indirect(kg, &indirect_sd, emission_sd, &ray, throughput, &state, L);
 
         /* sum and reset indirect light pass variables for the next samples */
-        path_radiance_sum_indirect(&L_sample);
-        path_radiance_reset_indirect(&L_sample);
+        path_radiance_sum_indirect(L);
+        path_radiance_reset_indirect(L);
       }
     }
 #  ifdef __BRANCHED_PATH__
@@ -112,13 +111,13 @@ ccl_device_inline void compute_light_pass(
 
     /* sample ambient occlusion */
     if (pass_filter & BAKE_FILTER_AO) {
-      kernel_branched_path_ao(kg, sd, &emission_sd, &L_sample, &state, throughput);
+      kernel_branched_path_ao(kg, sd, emission_sd, L, &state, throughput);
     }
 
     /* sample emission */
     if ((pass_filter & BAKE_FILTER_EMISSION) && (sd->flag & SD_EMISSION)) {
       float3 emission = indirect_primitive_emission(kg, sd, 0.0f, state.flag, state.ray_pdf);
-      path_radiance_accum_emission(kg, &L_sample, &state, throughput, emission);
+      path_radiance_accum_emission(kg, L, &state, throughput, emission);
     }
 
 #    ifdef __SUBSURFACE__
@@ -127,7 +126,7 @@ ccl_device_inline void compute_light_pass(
       /* When mixing BSSRDF and BSDF closures we should skip BSDF lighting
        * if scattering was successful. */
       kernel_branched_path_subsurface_scatter(
-          kg, sd, &indirect_sd, &emission_sd, &L_sample, &state, &ray, throughput);
+          kg, sd, &indirect_sd, emission_sd, L, &state, &ray, throughput);
     }
 #    endif
 
@@ -138,19 +137,16 @@ ccl_device_inline void compute_light_pass(
       if (kernel_data.integrator.use_direct_light) {
         int all = kernel_data.integrator.sample_all_lights_direct;
         kernel_branched_path_surface_connect_light(
-            kg, sd, &emission_sd, &state, throughput, 1.0f, &L_sample, all);
+            kg, sd, emission_sd, &state, throughput, 1.0f, L, all);
       }
 #    endif
 
       /* indirect light */
       kernel_branched_path_surface_indirect_light(
-          kg, sd, &indirect_sd, &emission_sd, throughput, 1.0f, &state, &L_sample);
+          kg, sd, &indirect_sd, emission_sd, throughput, 1.0f, &state, L);
     }
   }
 #  endif
-
-  /* accumulate into master L */
-  path_radiance_accum_sample(L, &L_sample);
 }
 
 /* this helps with AA but it's not the real solution as it does not AA the geometry
@@ -225,41 +221,28 @@ ccl_device float3 kernel_bake_evaluate_direct_indirect(KernelGlobals *kg,
   return out;
 }
 
-ccl_device void kernel_bake_evaluate(KernelGlobals *kg,
-                                     ccl_global uint4 *input,
-                                     ccl_global float4 *output,
-                                     ShaderEvalType type,
-                                     int pass_filter,
-                                     int i,
-                                     int offset,
-                                     int sample)
+ccl_device void kernel_bake_evaluate(
+    KernelGlobals *kg, ccl_global float *buffer, int sample, int x, int y, int offset, int stride)
 {
-  ShaderData sd;
-  PathState state = {0};
-  uint4 in = input[i * 2];
-  uint4 diff = input[i * 2 + 1];
-
-  float3 out = make_float3(0.0f, 0.0f, 0.0f);
+  /* Setup render buffers. */
+  const int index = offset + x + y * stride;
+  const int pass_stride = kernel_data.film.pass_stride;
+  buffer += index * pass_stride;
 
-  int object = in.x;
-  int prim = in.y;
+  ccl_global float *primitive = buffer + kernel_data.film.pass_bake_primitive;
+  ccl_global float *differential = buffer + kernel_data.film.pass_bake_differential;
+  ccl_global float *output = buffer + kernel_data.film.pass_combined;
 
+  int prim = __float_as_uint(primitive[1]);
   if (prim == -1)
     return;
 
-  float u = __uint_as_float(in.z);
-  float v = __uint_as_float(in.w);
-
-  float dudx = __uint_as_float(diff.x);
-  float dudy = __uint_as_float(diff.y);
-  float dvdx = __uint_as_float(diff.z);
-  float dvdy = __uint_as_float(diff.w);
+  prim += kernel_data.bake.tri_offset;
 
+  /* Random number generator. */
+  uint rng_hash = hash_uint2(x, y) ^ kernel_data.integrator.seed;
   int num_samples = kernel_data.integrator.aa_samples;
 
-  /* random number generator */
-  uint rng_hash = cmj_hash(offset + i, kernel_data.integrator.seed);
-
   float filter_x, filter_y;
   if (sample == 0) {
     filter_x = filter_y = 0.5f;
@@ -268,23 +251,29 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg,
     path_rng_2D(kg, rng_hash, sample, num_samples, PRNG_FILTER_U, &filter_x, &filter_y);
   }
 
-  /* subpixel u/v offset */
+  /* Barycentric UV with subpixel offset. */
+  float u = primitive[2];
+  float v = primitive[3];
+
+  float dudx = differential[0];
+  float dudy = differential[1];
+  float dvdx = differential[2];
+  float dvdy = differential[3];
+
   if (sample > 0) {
     u = bake_clamp_mirror_repeat(u + dudx * (filter_x - 0.5f) + dudy * (filter_y - 0.5f), 1.0f);
     v = bake_clamp_mirror_repeat(v + dvdx * (filter_x - 0.5f) + dvdy * (filter_y - 0.5f),
                                  1.0f - u);
   }
 
-  /* triangle */
+  /* Shader data setup. */
+  int object = kernel_data.bake.object_index;
   int shader;
   float3 P, Ng;
 
   triangle_point_normal(kg, object, prim, u, v, &P, &Ng, &shader);
 
-  /* light passes */
-  PathRadiance L;
-  path_radiance_init(kg, &L);
-
+  ShaderData sd;
   shader_setup_from_sample(
       kg,
       &sd,
@@ -302,7 +291,7 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg,
       LAMP_NONE);
   sd.I = sd.N;
 
-  /* update differentials */
+  /* Setup differentials. */
   sd.dP.dx = sd.dPdu * dudx + sd.dPdv * dvdx;
   sd.dP.dy = sd.dPdu * dudy + sd.dPdv * dvdy;
   sd.du.dx = dudx;
@@ -310,17 +299,24 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg,
   sd.dv.dx = dvdx;
   sd.dv.dy = dvdy;
 
-  /* set RNG state for shaders that use sampling */
+  /* Set RNG state for shaders that use sampling. */
+  PathState state = {0};
   state.rng_hash = rng_hash;
   state.rng_offset = 0;
   state.sample = sample;
   state.num_samples = num_samples;
   state.min_ray_pdf = FLT_MAX;
 
-  /* light passes if we need more than color */
-  if (pass_filter & ~BAKE_FILTER_COLOR)
+  /* Light passes if we need more than color. */
+  PathRadiance L;
+  int pass_filter = kernel_data.bake.pass_filter;
+
+  if (kernel_data.bake.pass_filter & ~BAKE_FILTER_COLOR)
     compute_light_pass(kg, &sd, &L, rng_hash, pass_filter, sample);
 
+  float3 out = make_float3(0.0f, 0.0f, 0.0f);
+
+  ShaderEvalType type = (ShaderEvalType)kernel_data.bake.type;
   switch (type) {
     /* data passes */
     case SHADER_EVAL_NORMAL:
@@ -441,10 +437,8 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg,
   }
 
   /* write output */
-  const float output_fac = 1.0f / num_samples;
-  const float4 scaled_result = make_float4(out.x, out.y, out.z, 1.0f) * output_fac;
-
-  output[i] = (sample == 0) ? scaled_result : output[i] + scaled_result;
+  const float4 result = make_float4(out.x, out.y, out.z, 1.0f);
+  kernel_write_pass_float4(output, result);
 }
 
 #endif /* __BAKING__ */
diff --git a/intern/cycles/kernel/kernel_film.h b/intern/cycles/kernel/kernel_film.h
index 3829426f261..8344f4b4f47 100644
--- a/intern/cycles/kernel/kernel_film.h
+++ b/intern/cycles/kernel/kernel_film.h
@@ -28,13 +28,13 @@ ccl_device float4 film_get_pass_result(KernelGlobals *kg,
   int display_pass_components = kernel_data.film.display_pass_components;
 
   if (display_pass_components == 4) {
-    ccl_global float4 *in = (ccl_global float4 *)(buffer + display_pass_stride +
-                                                  index * kernel_data.film.pass_stride);
+    float4 in = *(ccl_global float4 *)(buffer + display_pass_stride +
+                                       index * kernel_data.film.pass_stride);
     float alpha = use_display_sample_scale ?
-                      (kernel_data.film.use_display_pass_alpha ? in->w : 1.0f / sample_scale) :
+                      (kernel_data.film.use_display_pass_alpha ? in.w : 1.0f / sample_scale) :
                       1.0f;
 
-    pass_result = make_float4(in->x, in->y, in->z, alpha);
+    pass_result = make_float4(in.x, in.y, in.z, alpha);
 
     int display_divide_pass_stride = kernel_data.film.display_divide_pass_stride;
     if (display_divide_pass_stride != -1) {
diff --git a/intern/cycles/kernel/kernel_light.h b/intern/cycles/kernel/kernel_light.h
index ce908ce0fe2..d918abed381 100644
--- a/intern/cycles/kernel/kernel_light.h
+++ b/intern/cycles/kernel/kernel_light.h
@@ -1041,11 +1041,19 @@ ccl_device_forceinline void triangle_light_sample(KernelGlobals *kg,
     }
   }
   else {
-    /* compute random point in triangle */
-    randu = sqrtf(randu);
+    /* compute random point in triangle. From Eric Heitz's "A Low-Distortion Map Between Triangle
+     * and Square" */
+    float u = randu;
+    float v = randv;
+    if (v > u) {
+      u *= 0.5f;
+      v -= u;
+    }
+    else {
+      v *= 0.5f;
+      u -= v;
+    }
 
-    const float u = 1.0f - randu;
-    const float v = randv * randu;
     const float t = 1.0f - u - v;
     ls->P = u * V[0] + v * V[1] + t * V[2];
     /* compute incoming direction, distance and pdf */
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index a1f8c35348d..304835a1685 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -395,6 +395,10 @@ typedef enum PassType {
   PASS_VOLUME_INDIRECT,
   /* No Scatter color since it's tricky to define what it would even mean. */
   PASS_CATEGORY_LIGHT_END = 63,
+
+  PASS_BAKE_PRIMITIVE,
+  PASS_BAKE_DIFFERENTIAL,
+  PASS_CATEGORY_BAKE_END = 95
 } PassType;
 
 #define PASS_ANY (~0)
@@ -1248,6 +1252,10 @@ typedef struct KernelFilm {
   float4 xyz_to_b;
   float4 rgb_to_y;
 
+  int pass_bake_primitive;
+  int pass_bake_differential;
+  int pad;
+
 #ifdef __KERNEL_DEBUG__
   int pass_bvh_traversed_nodes;
   int pass_bvh_traversed_instances;
@@ -1427,6 +1435,14 @@ typedef struct KernelTables {
 } KernelTables;
 static_assert_align(KernelTables, 16);
 
+typedef struct KernelBake {
+  int object_index;
+  int tri_offset;
+  int type;
+  int pass_filter;
+} KernelBake;
+static_assert_align(KernelBake, 16);
+
 typedef struct KernelData {
   KernelCamera cam;
   KernelFilm film;
@@ -1435,6 +1451,7 @@ typedef struct KernelData {
   KernelBVH bvh;
   KernelCurves curve;
   KernelTables tables;
+  KernelBake bake;
 } KernelData;
 static_assert_align(KernelData, 16);
 
diff --git a/intern/cycles/kernel/kernels/cpu/kernel_cpu.h b/intern/cycles/kernel/kernels/cpu/kernel_cpu.h
index 683f4b88d79..ea3103f12c3 100644
--- a/intern/cycles/kernel/kernels/cpu/kernel_cpu.h
+++ b/intern/cycles/kernel/kernels/cpu/kernel_cpu.h
@@ -46,6 +46,9 @@ void KERNEL_FUNCTION_FULL_NAME(shader)(KernelGlobals *kg,
                                        int offset,
                                        int sample);
 
+void KERNEL_FUNCTION_FULL_NAME(bake)(
+    KernelGlobals *kg, float *buffer, int sample, int x, int y, int offset, int stride);
+
 /* Split kernels */
 
 void KERNEL_FUNCTION_FULL_NAME(data_init)(KernelGlobals *kg,
diff --git a/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h b/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h
index 091e53cfd83..5aa3fb14318 100644
--- a/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h
+++ b/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h
@@ -132,6 +132,18 @@ void KERNEL_FUNCTION_FULL_NAME(convert_to_half_float)(KernelGlobals *kg,
 #  endif /* KERNEL_STUB */
 }
 
+/* Bake */
+
+void KERNEL_FUNCTION_FULL_NAME(bake)(
+    KernelGlobals *kg, float *buffer, int sample, int x, int y, int offset, int stride)
+{
+#  ifdef KERNEL_STUB
+  STUB_ASSERT(KERNEL_ARCH, bake);
+#  else
+  kernel_bake_evaluate(kg, buffer, sample, x, y, offset, stride);
+#  endif /* KERNEL_STUB */
+}
+
 /* Shader Evaluate */
 
 void KERNEL_FUNCTION_FULL_NAME(shader)(KernelGlobals *kg,
@@ -146,12 +158,7 @@ void KERNEL_FUNCTION_FULL_NAME(shader)(KernelGlobals *kg,
 #  ifdef KERNEL_STUB
   STUB_ASSERT(KERNEL_ARCH, shader);
 #  else
-  if (type >= SHADER_EVAL_BAKE) {
-#    ifdef __BAKING__
-    kernel_bake_evaluate(kg, input, output, (ShaderEvalType)type, filter, i, offset, sample);
-#    endif
-  }
-  else if (type == SHADER_EVAL_DISPLACE) {
+  if (type == SHADER_EVAL_DISPLACE) {
     kernel_displace_evaluate(kg, input, output, i);
   }
   else {
diff --git a/intern/cycles/kernel/kernels/cuda/kernel.cu b/intern/cycles/kernel/kernels/cuda/kernel.cu
index c4c810c6a82..d4f41132a11 100644
--- a/intern/cycles/kernel/kernels/cuda/kernel.cu
+++ b/intern/cycles/kernel/kernels/cuda/kernel.cu
@@ -214,13 +214,16 @@ kernel_cuda_background(uint4 *input,
 #ifdef __BAKING__
 extern "C" __global__ void
 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
-kernel_cuda_bake(uint4 *input, float4 *output, int type, int filter, int sx, int sw, int offset, int sample)
+kernel_cuda_bake(WorkTile *tile, uint total_work_size)
 {
-	int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
+	int work_index = ccl_global_id(0);
+
+	if(work_index < total_work_size) {
+		uint x, y, sample;
+		get_work_pixel(tile, work_index, &x, &y, &sample);
 
-	if(x < sx + sw) {
 		KernelGlobals kg;
-		kernel_bake_evaluate(&kg, input, output, (ShaderEvalType)type, filter, x, offset, sample);
+		kernel_bake_evaluate(&kg, tile->buffer, sample, x, y, tile->offset, tile->stride);
 	}
 }
 #endif
diff --git a/intern/cycles/render/bake.cpp b/intern/cycles/render/bake.cpp
index 35f942b3e9b..6044182a51a 100644
--- a/intern/cycles/render/bake.cpp
+++ b/intern/cycles/render/bake.cpp
@@ -15,6 +15,7 @@
  */
 
 #include "render/bake.h"
+#include "render/buffers.h"
 #include "render/integrator.h"
 #include "render/mesh.h"
 #include "render/object.h"
@@ -24,272 +25,130 @@
 
 CCL_NAMESPACE_BEGIN
 
-BakeData::BakeData(const int object, const size_t tri_offset, const size_t num_pixels)
-    : m_object(object), m_tri_offset(tri_offset), m_num_pixels(num_pixels)
+static int aa_samples(Scene *scene, Object *object, ShaderEvalType type)
 {
-  m_primitive.resize(num_pixels);
-  m_u.resize(num_pixels);
-  m_v.resize(num_pixels);
-  m_dudx.resize(num_pixels);
-  m_dudy.resize(num_pixels);
-  m_dvdx.resize(num_pixels);
-  m_dvdy.resize(num_pixels);
-}
-
-BakeData::~BakeData()
-{
-  m_primitive.clear();
-  m_u.clear();
-  m_v.clear();
-  m_dudx.clear();
-  m_dudy.clear();
-  m_dvdx.clear();
-  m_dvdy.clear();
-}
-
-void BakeData::set(int i, int prim, float uv[2], float dudx, float dudy, float dvdx, float dvdy)
-{
-  m_primitive[i] = (prim == -1 ? -1 : m_tri_offset + prim);
-  m_u[i] = uv[0];
-  m_v[i] = uv[1];
-  m_dudx[i] = dudx;
-  m_dudy[i] = dudy;
-  m_dvdx[i] = dvdx;
-  m_dvdy[i] = dvdy;
-}
-
-void BakeData::set_null(int i)
-{
-  m_primitive[i] = -1;
-}
-
-int BakeData::object()
-{
-  return m_object;
-}
-
-size_t BakeData::size()
-{
-  return m_num_pixels;
-}
+  if (type == SHADER_EVAL_UV || type == SHADER_EVAL_ROUGHNESS) {
+    return 1;
+  }
+  else if (type == SHADER_EVAL_NORMAL) {
+    /* Only antialias normal if mesh has bump mapping. */
+    if (object->geometry) {
+      foreach (Shader *shader, object->geometry->used_shaders) {
+        if (shader->has_bump) {
+          return scene->integrator->aa_samples;
+        }
+      }
+    }
 
-bool BakeData::is_valid(int i)
-{
-  return m_primitive[i] != -1;
+    return 1;
+  }
+  else {
+    return scene->integrator->aa_samples;
+  }
 }
 
-uint4 BakeData::data(int i)
+/* Keep it synced with kernel_bake.h logic */
+static int shader_type_to_pass_filter(ShaderEvalType type, int pass_filter)
 {
-  return make_uint4(m_object, m_primitive[i], __float_as_int(m_u[i]), __float_as_int(m_v[i]));
-}
+  const int component_flags = pass_filter &
+                              (BAKE_FILTER_DIRECT | BAKE_FILTER_INDIRECT | BAKE_FILTER_COLOR);
 
-uint4 BakeData::differentials(int i)
-{
-  return make_uint4(__float_as_int(m_dudx[i]),
-                    __float_as_int(m_dudy[i]),
-                    __float_as_int(m_dvdx[i]),
-                    __float_as_int(m_dvdy[i]));
+  switch (type) {
+    case SHADER_EVAL_AO:
+      return BAKE_FILTER_AO;
+    case SHADER_EVAL_SHADOW:
+      return BAKE_FILTER_DIRECT;
+    case SHADER_EVAL_DIFFUSE:
+      return BAKE_FILTER_DIFFUSE | component_flags;
+    case SHADER_EVAL_GLOSSY:
+      return BAKE_FILTER_GLOSSY | component_flags;
+    case SHADER_EVAL_TRANSMISSION:
+      return BAKE_FILTER_TRANSMISSION | component_flags;
+    case SHADER_EVAL_COMBINED:
+      return pass_filter;
+    default:
+      return 0;
+  }
 }
 
 BakeManager::BakeManager()
 {
-  m_bake_data = NULL;
-  m_is_baking = false;
+  type = SHADER_EVAL_BAKE;
+  pass_filter = 0;
+
   need_update = true;
-  m_shader_limit = 512 * 512;
 }
 
 BakeManager::~BakeManager()
 {
-  if (m_bake_data)
-    delete m_bake_data;
 }
 
 bool BakeManager::get_baking()
 {
-  return m_is_baking;
-}
-
-void BakeManager::set_baking(const bool value)
-{
-  m_is_baking = value;
+  return !object_name.empty();
 }
 
-BakeData *BakeManager::init(const int object, const size_t tri_offset, const size_t num_pixels)
+void BakeManager::set(Scene *scene,
+                      const std::string &object_name_,
+                      ShaderEvalType type_,
+                      int pass_filter_)
 {
-  m_bake_data = new BakeData(object, tri_offset, num_pixels);
-  return m_bake_data;
-}
-
-void BakeManager::set_shader_limit(const size_t x, const size_t y)
-{
-  m_shader_limit = x * y;
-  m_shader_limit = (size_t)pow(2, std::ceil(log(m_shader_limit) / log(2)));
-}
+  object_name = object_name_;
+  type = type_;
+  pass_filter = shader_type_to_pass_filter(type_, pass_filter_);
 
-bool BakeManager::bake(Device *device,
-                       DeviceScene *dscene,
-                       Scene *scene,
-                       Progress &progress,
-                       ShaderEvalType shader_type,
-                       const int pass_filter,
-                       BakeData *bake_data,
-                       float result[])
-{
-  size_t num_pixels = bake_data->size();
-
-  int num_samples = aa_samples(scene, bake_data, shader_type);
+  Pass::add(PASS_BAKE_PRIMITIVE, scene->film->passes);
+  Pass::add(PASS_BAKE_DIFFERENTIAL, scene->film->passes);
 
-  /* calculate the total pixel samples for the progress bar */
-  total_pixel_samples = 0;
-  for (size_t shader_offset = 0; shader_offset < num_pixels; shader_offset += m_shader_limit) {
-    size_t shader_size = (size_t)fminf(num_pixels - shader_offset, m_shader_limit);
-    total_pixel_samples += shader_size * num_samples;
+  if (type == SHADER_EVAL_UV) {
+    /* force UV to be available */
+    Pass::add(PASS_UV, scene->film->passes);
   }
-  progress.reset_sample();
-  progress.set_total_pixel_samples(total_pixel_samples);
-
-  /* needs to be up to date for baking specific AA samples */
-  dscene->data.integrator.aa_samples = num_samples;
-  device->const_copy_to("__data", &dscene->data, sizeof(dscene->data));
-
-  for (size_t shader_offset = 0; shader_offset < num_pixels; shader_offset += m_shader_limit) {
-    size_t shader_size = (size_t)fminf(num_pixels - shader_offset, m_shader_limit);
 
-    /* setup input for device task */
-    device_vector<uint4> d_input(device, "bake_input", MEM_READ_ONLY);
-    uint4 *d_input_data = d_input.alloc(shader_size * 2);
-    size_t d_input_size = 0;
-
-    for (size_t i = shader_offset; i < (shader_offset + shader_size); i++) {
-      d_input_data[d_input_size++] = bake_data->data(i);
-      d_input_data[d_input_size++] = bake_data->differentials(i);
-    }
-
-    if (d_input_size == 0) {
-      m_is_baking = false;
-      return false;
-    }
-
-    /* run device task */
-    device_vector<float4> d_output(device, "bake_output", MEM_READ_WRITE);
-    d_output.alloc(shader_size);
-    d_output.zero_to_device();
-    d_input.copy_to_device();
-
-    DeviceTask task(DeviceTask::SHADER);
-    task.shader_input = d_input.device_pointer;
-    task.shader_output = d_output.device_pointer;
-    task.shader_eval_type = shader_type;
-    task.shader_filter = pass_filter;
-    task.shader_x = 0;
-    task.offset = shader_offset;
-    task.shader_w = d_output.size();
-    task.num_samples = num_samples;
-    task.get_cancel = function_bind(&Progress::get_cancel, &progress);
-    task.update_progress_sample = function_bind(&Progress::add_samples_update, &progress, _1, _2);
-
-    device->task_add(task);
-    device->task_wait();
-
-    if (progress.get_cancel()) {
-      d_input.free();
-      d_output.free();
-      m_is_baking = false;
-      return false;
-    }
-
-    d_output.copy_from_device(0, 1, d_output.size());
-    d_input.free();
-
-    /* read result */
-    int k = 0;
-
-    float4 *offset = d_output.data();
-
-    size_t depth = 4;
-    for (size_t i = shader_offset; i < (shader_offset + shader_size); i++) {
-      size_t index = i * depth;
-      float4 out = offset[k++];
-
-      if (bake_data->is_valid(i)) {
-        for (size_t j = 0; j < 4; j++) {
-          result[index + j] = out[j];
-        }
-      }
-    }
-
-    d_output.free();
+  /* force use_light_pass to be true if we bake more than just colors */
+  if (pass_filter & ~BAKE_FILTER_COLOR) {
+    Pass::add(PASS_LIGHT, scene->film->passes);
   }
 
-  m_is_baking = false;
-  return true;
+  /* create device and update scene */
+  scene->film->tag_update(scene);
+  scene->integrator->tag_update(scene);
+
+  need_update = true;
 }
 
 void BakeManager::device_update(Device * /*device*/,
-                                DeviceScene * /*dscene*/,
-                                Scene * /*scene*/,
-                                Progress &progress)
+                                DeviceScene *dscene,
+                                Scene *scene,
+                                Progress & /* progress */)
 {
   if (!need_update)
     return;
 
-  if (progress.get_cancel())
-    return;
+  KernelIntegrator *kintegrator = &dscene->data.integrator;
+  KernelBake *kbake = &dscene->data.bake;
 
-  need_update = false;
-}
-
-void BakeManager::device_free(Device * /*device*/, DeviceScene * /*dscene*/)
-{
-}
-
-int BakeManager::aa_samples(Scene *scene, BakeData *bake_data, ShaderEvalType type)
-{
-  if (type == SHADER_EVAL_UV || type == SHADER_EVAL_ROUGHNESS) {
-    return 1;
-  }
-  else if (type == SHADER_EVAL_NORMAL) {
-    /* Only antialias normal if mesh has bump mapping. */
-    Object *object = scene->objects[bake_data->object()];
+  kbake->type = type;
+  kbake->pass_filter = pass_filter;
 
-    if (object->geometry) {
-      foreach (Shader *shader, object->geometry->used_shaders) {
-        if (shader->has_bump) {
-          return scene->integrator->aa_samples;
-        }
-      }
+  int object_index = 0;
+  foreach (Object *object, scene->objects) {
+    const Geometry *geom = object->geometry;
+    if (object->name == object_name && geom->type == Geometry::MESH) {
+      kbake->object_index = object_index;
+      kbake->tri_offset = geom->prim_offset;
+      kintegrator->aa_samples = aa_samples(scene, object, type);
+      break;
     }
 
-    return 1;
-  }
-  else {
-    return scene->integrator->aa_samples;
+    object_index++;
   }
+
+  need_update = false;
 }
 
-/* Keep it synced with kernel_bake.h logic */
-int BakeManager::shader_type_to_pass_filter(ShaderEvalType type, const int pass_filter)
+void BakeManager::device_free(Device * /*device*/, DeviceScene * /*dscene*/)
 {
-  const int component_flags = pass_filter &
-                              (BAKE_FILTER_DIRECT | BAKE_FILTER_INDIRECT | BAKE_FILTER_COLOR);
-
-  switch (type) {
-    case SHADER_EVAL_AO:
-      return BAKE_FILTER_AO;
-    case SHADER_EVAL_SHADOW:
-      return BAKE_FILTER_DIRECT;
-    case SHADER_EVAL_DIFFUSE:
-      return BAKE_FILTER_DIFFUSE | component_flags;
-    case SHADER_EVAL_GLOSSY:
-      return BAKE_FILTER_GLOSSY | component_flags;
-    case SHADER_EVAL_TRANSMISSION:
-      return BAKE_FILTER_TRANSMISSION | component_flags;
-    case SHADER_EVAL_COMBINED:
-      return pass_filter;
-    default:
-      return 0;
-  }
 }
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/render/bake.h b/intern/cycles/render/bake.h
index 88537623efb..93e664c2ab1 100644
--- a/intern/cycles/render/bake.h
+++ b/intern/cycles/render/bake.h
@@ -25,67 +25,23 @@
 
 CCL_NAMESPACE_BEGIN
 
-class BakeData {
- public:
-  BakeData(const int object, const size_t tri_offset, const size_t num_pixels);
-  ~BakeData();
-
-  void set(int i, int prim, float uv[2], float dudx, float dudy, float dvdx, float dvdy);
-  void set_null(int i);
-  int object();
-  size_t size();
-  uint4 data(int i);
-  uint4 differentials(int i);
-  bool is_valid(int i);
-
- private:
-  int m_object;
-  size_t m_tri_offset;
-  size_t m_num_pixels;
-  vector<int> m_primitive;
-  vector<float> m_u;
-  vector<float> m_v;
-  vector<float> m_dudx;
-  vector<float> m_dudy;
-  vector<float> m_dvdx;
-  vector<float> m_dvdy;
-};
-
 class BakeManager {
  public:
   BakeManager();
   ~BakeManager();
 
+  void set(Scene *scene, const std::string &object_name, ShaderEvalType type, int pass_filter);
   bool get_baking();
-  void set_baking(const bool value);
-
-  BakeData *init(const int object, const size_t tri_offset, const size_t num_pixels);
-
-  void set_shader_limit(const size_t x, const size_t y);
-
-  bool bake(Device *device,
-            DeviceScene *dscene,
-            Scene *scene,
-            Progress &progress,
-            ShaderEvalType shader_type,
-            const int pass_filter,
-            BakeData *bake_data,
-            float result[]);
 
   void device_update(Device *device, DeviceScene *dscene, Scene *scene, Progress &progress);
   void device_free(Device *device, DeviceScene *dscene);
 
-  static int shader_type_to_pass_filter(ShaderEvalType type, const int pass_filter);
-  static int aa_samples(Scene *scene, BakeData *bake_data, ShaderEvalType type);
-
   bool need_update;
 
-  size_t total_pixel_samples;
-
  private:
-  BakeData *m_bake_data;
-  bool m_is_baking;
-  size_t m_shader_limit;
+  ShaderEvalType type;
+  int pass_filter;
+  std::string object_name;
 };
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/render/buffers.cpp b/intern/cycles/render/buffers.cpp
index 2d89fb9ffba..b26366af852 100644
--- a/intern/cycles/render/buffers.cpp
+++ b/intern/cycles/render/buffers.cpp
@@ -459,6 +459,40 @@ bool RenderBuffers::get_pass_rect(
   return false;
 }
 
+bool RenderBuffers::set_pass_rect(PassType type, int components, float *pixels)
+{
+  if (buffer.data() == NULL) {
+    return false;
+  }
+
+  int pass_offset = 0;
+
+  for (size_t j = 0; j < params.passes.size(); j++) {
+    Pass &pass = params.passes[j];
+
+    if (pass.type != type) {
+      pass_offset += pass.components;
+      continue;
+    }
+
+    float *out = buffer.data() + pass_offset;
+    int pass_stride = params.get_passes_size();
+    int size = params.width * params.height;
+
+    assert(pass.components == components);
+
+    for (int i = 0; i < size; i++, out += pass_stride, pixels += components) {
+      for (int j = 0; j < components; j++) {
+        out[j] = pixels[j];
+      }
+    }
+
+    return true;
+  }
+
+  return false;
+}
+
 /* Display Buffer */
 
 DisplayBuffer::DisplayBuffer(Device *device, bool linear)
diff --git a/intern/cycles/render/buffers.h b/intern/cycles/render/buffers.h
index 42efb031843..975bae2239c 100644
--- a/intern/cycles/render/buffers.h
+++ b/intern/cycles/render/buffers.h
@@ -92,6 +92,7 @@ class RenderBuffers {
       const string &name, float exposure, int sample, int components, float *pixels);
   bool get_denoising_pass_rect(
       int offset, float exposure, int sample, int components, float *pixels);
+  bool set_pass_rect(PassType type, int components, float *pixels);
 };
 
 /* Display Buffer
@@ -130,7 +131,7 @@ class DisplayBuffer {
 
 class RenderTile {
  public:
-  typedef enum { PATH_TRACE = (1 << 0), DENOISE = (1 << 1) } Task;
+  typedef enum { PATH_TRACE = (1 << 0), BAKE = (1 << 1), DENOISE = (1 << 2) } Task;
 
   Task task;
   int x, y, w, h;
diff --git a/intern/cycles/render/film.cpp b/intern/cycles/render/film.cpp
index 26eda93fadd..d7cbf4a3581 100644
--- a/intern/cycles/render/film.cpp
+++ b/intern/cycles/render/film.cpp
@@ -196,6 +196,10 @@ void Pass::add(PassType type, vector<Pass> &passes, const char *name)
     case PASS_AOV_VALUE:
       pass.components = 1;
       break;
+    case PASS_BAKE_PRIMITIVE:
+    case PASS_BAKE_DIFFERENTIAL:
+      pass.components = 4;
+      break;
     default:
       assert(false);
       break;
@@ -386,11 +390,13 @@ void Film::device_update(Device *device, DeviceScene *dscene, Scene *scene)
     if (pass.type <= PASS_CATEGORY_MAIN_END) {
       kfilm->pass_flag |= pass_flag;
     }
-    else {
-      assert(pass.type <= PASS_CATEGORY_LIGHT_END);
+    else if (pass.type <= PASS_CATEGORY_LIGHT_END) {
       kfilm->use_light_pass = 1;
       kfilm->light_pass_flag |= pass_flag;
     }
+    else {
+      assert(pass.type <= PASS_CATEGORY_BAKE_END);
+    }
 
     switch (pass.type) {
       case PASS_COMBINED:
@@ -471,6 +477,13 @@ void Film::device_update(Device *device, DeviceScene *dscene, Scene *scene)
         kfilm->pass_volume_direct = kfilm->pass_stride;
         break;
 
+      case PASS_BAKE_PRIMITIVE:
+        kfilm->pass_bake_primitive = kfilm->pass_stride;
+        break;
+      case PASS_BAKE_DIFFERENTIAL:
+        kfilm->pass_bake_differential = kfilm->pass_stride;
+        break;
+
 #ifdef WITH_CYCLES_DEBUG
       case PASS_BVH_TRAVERSED_NODES:
         kfilm->pass_bvh_traversed_nodes = kfilm->pass_stride;
diff --git a/intern/cycles/render/session.cpp b/intern/cycles/render/session.cpp
index f7df81a0601..361a1465aac 100644
--- a/intern/cycles/render/session.cpp
+++ b/intern/cycles/render/session.cpp
@@ -410,7 +410,16 @@ bool Session::acquire_tile(RenderTile &rtile, Device *tile_device, uint tile_typ
   rtile.num_samples = tile_manager.state.num_samples;
   rtile.resolution = tile_manager.state.resolution_divider;
   rtile.tile_index = tile->index;
-  rtile.task = tile->state == Tile::DENOISE ? RenderTile::DENOISE : RenderTile::PATH_TRACE;
+
+  if (tile->state == Tile::DENOISE) {
+    rtile.task = RenderTile::DENOISE;
+  }
+  else if (read_bake_tile_cb) {
+    rtile.task = RenderTile::BAKE;
+  }
+  else {
+    rtile.task = RenderTile::PATH_TRACE;
+  }
 
   tile_lock.unlock();
 
@@ -451,11 +460,20 @@ bool Session::acquire_tile(RenderTile &rtile, Device *tile_device, uint tile_typ
   rtile.buffers = tile->buffers;
   rtile.sample = tile_manager.state.sample;
 
-  /* this will tag tile as IN PROGRESS in blender-side render pipeline,
-   * which is needed to highlight currently rendering tile before first
-   * sample was processed for it
-   */
-  update_tile_sample(rtile);
+  if (read_bake_tile_cb) {
+    /* This will read any passes needed as input for baking. */
+    {
+      thread_scoped_lock tile_lock(tile_mutex);
+      read_bake_tile_cb(rtile);
+    }
+    rtile.buffers->buffer.copy_to_device();
+  }
+  else {
+    /* This will tag tile as IN PROGRESS in blender-side render pipeline,
+     * which is needed to highlight currently rendering tile before first
+     * sample was processed for it. */
+    update_tile_sample(rtile);
+  }
 
   return true;
 }
@@ -484,6 +502,7 @@ void Session::release_tile(RenderTile &rtile, const bool need_denoise)
   bool delete_tile;
 
   if (tile_manager.finish_tile(rtile.tile_index, need_denoise, delete_tile)) {
+    /* Finished tile pixels write. */
     if (write_render_tile_cb && params.progressive_refine == false) {
       write_render_tile_cb(rtile);
     }
@@ -494,6 +513,7 @@ void Session::release_tile(RenderTile &rtile, const bool need_denoise)
     }
   }
   else {
+    /* In progress tile pixels update. */
     if (update_render_tile_cb && params.progressive_refine == false) {
       update_render_tile_cb(rtile, false);
     }
diff --git a/intern/cycles/render/session.h b/intern/cycles/render/session.h
index f06952e8020..2707eed5531 100644
--- a/intern/cycles/render/session.h
+++ b/intern/cycles/render/session.h
@@ -148,6 +148,7 @@ class Session {
 
   function<void(RenderTile &)> write_render_tile_cb;
   function<void(RenderTile &, bool)> update_render_tile_cb;
+  function<void(RenderTile &)> read_bake_tile_cb;
 
   explicit Session(const SessionParams &params);
   ~Session();
diff --git a/intern/ghost/intern/GHOST_System.cpp b/intern/ghost/intern/GHOST_System.cpp
index 587e4c28102..b0d2adff4bc 100644
--- a/intern/ghost/intern/GHOST_System.cpp
+++ b/intern/ghost/intern/GHOST_System.cpp
@@ -23,8 +23,8 @@
 
 #include "GHOST_System.h"
 
+#include <chrono>
 #include <stdio.h> /* just for printf */
-#include <time.h>
 
 #include "GHOST_DisplayManager.h"
 #include "GHOST_EventManager.h"
@@ -58,12 +58,9 @@ GHOST_System::~GHOST_System()
 
 GHOST_TUns64 GHOST_System::getMilliSeconds() const
 {
-  GHOST_TUns64 millis = ::clock();
-  if (CLOCKS_PER_SEC != 1000) {
-    millis *= 1000;
-    millis /= CLOCKS_PER_SEC;
-  }
-  return millis;
+  return std::chrono::duration_cast<std::chrono::milliseconds>(
+             std::chrono::steady_clock::now().time_since_epoch())
+      .count();
 }
 
 GHOST_ITimerTask *GHOST_System::installTimer(GHOST_TUns64 delay,
diff --git a/intern/ghost/intern/GHOST_SystemWayland.cpp b/intern/ghost/intern/GHOST_SystemWayland.cpp
index 1346e1bce20..633451feb85 100644
--- a/intern/ghost/intern/GHOST_SystemWayland.cpp
+++ b/intern/ghost/intern/GHOST_SystemWayland.cpp
@@ -25,6 +25,7 @@
 #include "GHOST_EventDragnDrop.h"
 #include "GHOST_EventKey.h"
 #include "GHOST_EventWheel.h"
+#include "GHOST_TimerManager.h"
 #include "GHOST_WindowManager.h"
 
 #include "GHOST_ContextEGL.h"
@@ -34,7 +35,7 @@
 
 #include <algorithm>
 #include <atomic>
-#include <exception>
+#include <stdexcept>
 #include <thread>
 #include <unordered_map>
 #include <unordered_set>
@@ -91,6 +92,13 @@ struct data_source_t {
   char *buffer_out;
 };
 
+struct key_repeat_payload_t {
+  GHOST_SystemWayland *system;
+  GHOST_IWindow *window;
+  GHOST_TKey key;
+  GHOST_TEventKeyData key_data;
+};
+
 struct input_t {
   GHOST_SystemWayland *system;
 
@@ -109,6 +117,17 @@ struct input_t {
 
   struct xkb_context *xkb_context;
   struct xkb_state *xkb_state;
+  struct {
+    /* Key repetition in character per second. */
+    int32_t rate;
+    /* Time (milliseconds) after which to start repeating keys. */
+    int32_t delay;
+    /* Timer for key repeats. */
+    GHOST_ITimerTask *timer = nullptr;
+  } key_repeat;
+
+  struct wl_surface *focus_pointer = nullptr;
+  struct wl_surface *focus_keyboard = nullptr;
 
   struct wl_data_device *data_device = nullptr;
   struct data_offer_t *data_offer_dnd;        /* Drag & Drop. */
@@ -174,6 +193,11 @@ static void display_destroy(display_t *d)
       }
     }
     if (input->keyboard) {
+      if (input->key_repeat.timer) {
+        delete static_cast<key_repeat_payload_t *>(input->key_repeat.timer->getUserData());
+        input->system->removeTimer(input->key_repeat.timer);
+        input->key_repeat.timer = nullptr;
+      }
       wl_keyboard_destroy(input->keyboard);
     }
     if (input->xkb_state) {
@@ -420,22 +444,26 @@ static void relative_pointer_relative_motion(
   input->x += wl_fixed_to_int(dx);
   input->y += wl_fixed_to_int(dy);
 
-  input->system->pushEvent(
-      new GHOST_EventCursor(input->system->getMilliSeconds(),
-                            GHOST_kEventCursorMove,
-                            input->system->getWindowManager()->getActiveWindow(),
-                            input->x,
-                            input->y,
-                            GHOST_TABLET_DATA_NONE));
+  GHOST_IWindow *win = static_cast<GHOST_WindowWayland *>(
+      wl_surface_get_user_data(input->focus_pointer));
+
+  input->system->pushEvent(new GHOST_EventCursor(input->system->getMilliSeconds(),
+                                                 GHOST_kEventCursorMove,
+                                                 win,
+                                                 input->x,
+                                                 input->y,
+                                                 GHOST_TABLET_DATA_NONE));
 }
 
 static const zwp_relative_pointer_v1_listener relative_pointer_listener = {
-    relative_pointer_relative_motion};
+    relative_pointer_relative_motion,
+};
 
 static void dnd_events(const input_t *const input, const GHOST_TEventType event)
 {
   const GHOST_TUns64 time = input->system->getMilliSeconds();
-  GHOST_IWindow *const window = input->system->getWindowManager()->getActiveWindow();
+  GHOST_IWindow *const window = static_cast<GHOST_WindowWayland *>(
+      wl_surface_get_user_data(input->focus_pointer));
   for (const std::string &type : mime_preference_order) {
     input->system->pushEvent(new GHOST_EventDragnDrop(time,
                                                       event,
@@ -641,7 +669,7 @@ static void data_device_drop(void *data, struct wl_data_device * /*wl_data_devic
                                                        data_offer->types.begin(),
                                                        data_offer->types.end());
 
-  auto read_uris = [](GHOST_SystemWayland *const system,
+  auto read_uris = [](input_t *const input,
                       data_offer_t *data_offer,
                       const std::string mime_receive) {
     const int x = data_offer->dnd.x;
@@ -655,6 +683,8 @@ static void data_device_drop(void *data, struct wl_data_device * /*wl_data_devic
     delete data_offer;
     data_offer = nullptr;
 
+    GHOST_SystemWayland *const system = input->system;
+
     if (mime_receive == mime_text_uri) {
       static constexpr const char *file_proto = "file://";
       static constexpr const char *crlf = "\r\n";
@@ -683,10 +713,12 @@ static void data_device_drop(void *data, struct wl_data_device * /*wl_data_devic
             malloc((uris[i].size() + 1) * sizeof(GHOST_TUns8)));
         memcpy(flist->strings[i], uris[i].data(), uris[i].size() + 1);
       }
+      GHOST_IWindow *win = static_cast<GHOST_WindowWayland *>(
+          wl_surface_get_user_data(input->focus_pointer));
       system->pushEvent(new GHOST_EventDragnDrop(system->getMilliSeconds(),
                                                  GHOST_kEventDraggingDropDone,
                                                  GHOST_kDragnDropTypeFilenames,
-                                                 system->getWindowManager()->getActiveWindow(),
+                                                 win,
                                                  x,
                                                  y,
                                                  flist));
@@ -698,7 +730,7 @@ static void data_device_drop(void *data, struct wl_data_device * /*wl_data_devic
     wl_display_roundtrip(system->display());
   };
 
-  std::thread read_thread(read_uris, input->system, data_offer, mime_receive);
+  std::thread read_thread(read_uris, input, data_offer, mime_receive);
   read_thread.detach();
 }
 
@@ -778,17 +810,24 @@ static void pointer_enter(void *data,
   input->pointer_serial = serial;
   input->x = wl_fixed_to_int(surface_x);
   input->y = wl_fixed_to_int(surface_y);
+  input->focus_pointer = surface;
 
-  static_cast<GHOST_WindowWayland *>(wl_surface_get_user_data(surface))->activate();
+  input->system->pushEvent(
+      new GHOST_EventCursor(input->system->getMilliSeconds(),
+                            GHOST_kEventCursorMove,
+                            static_cast<GHOST_WindowWayland *>(wl_surface_get_user_data(surface)),
+                            input->x,
+                            input->y,
+                            GHOST_TABLET_DATA_NONE));
 }
 
-static void pointer_leave(void * /*data*/,
+static void pointer_leave(void *data,
                           struct wl_pointer * /*wl_pointer*/,
                           uint32_t /*serial*/,
                           struct wl_surface *surface)
 {
   if (surface != nullptr) {
-    static_cast<GHOST_WindowWayland *>(wl_surface_get_user_data(surface))->deactivate();
+    static_cast<input_t *>(data)->focus_pointer = nullptr;
   }
 }
 
@@ -800,16 +839,22 @@ static void pointer_motion(void *data,
 {
   input_t *input = static_cast<input_t *>(data);
 
+  GHOST_IWindow *win = static_cast<GHOST_WindowWayland *>(
+      wl_surface_get_user_data(input->focus_pointer));
+
+  if (!win) {
+    return;
+  }
+
   input->x = wl_fixed_to_int(surface_x);
   input->y = wl_fixed_to_int(surface_y);
 
-  input->system->pushEvent(
-      new GHOST_EventCursor(input->system->getMilliSeconds(),
-                            GHOST_kEventCursorMove,
-                            input->system->getWindowManager()->getActiveWindow(),
-                            wl_fixed_to_int(surface_x),
-                            wl_fixed_to_int(surface_y),
-                            GHOST_TABLET_DATA_NONE));
+  input->system->pushEvent(new GHOST_EventCursor(input->system->getMilliSeconds(),
+                                                 GHOST_kEventCursorMove,
+                                                 win,
+                                                 wl_fixed_to_int(surface_x),
+                                                 wl_fixed_to_int(surface_y),
+                                                 GHOST_TABLET_DATA_NONE));
 }
 
 static void pointer_button(void *data,
@@ -843,14 +888,12 @@ static void pointer_button(void *data,
   }
 
   input_t *input = static_cast<input_t *>(data);
+  GHOST_IWindow *win = static_cast<GHOST_WindowWayland *>(
+      wl_surface_get_user_data(input->focus_pointer));
   input->data_source->source_serial = serial;
   input->buttons.set(ebutton, state == WL_POINTER_BUTTON_STATE_PRESSED);
-  input->system->pushEvent(
-      new GHOST_EventButton(input->system->getMilliSeconds(),
-                            etype,
-                            input->system->getWindowManager()->getActiveWindow(),
-                            ebutton,
-                            GHOST_TABLET_DATA_NONE));
+  input->system->pushEvent(new GHOST_EventButton(
+      input->system->getMilliSeconds(), etype, win, ebutton, GHOST_TABLET_DATA_NONE));
 }
 
 static void pointer_axis(void *data,
@@ -863,10 +906,10 @@ static void pointer_axis(void *data,
     return;
   }
   input_t *input = static_cast<input_t *>(data);
+  GHOST_IWindow *win = static_cast<GHOST_WindowWayland *>(
+      wl_surface_get_user_data(input->focus_pointer));
   input->system->pushEvent(
-      new GHOST_EventWheel(input->system->getMilliSeconds(),
-                           input->system->getWindowManager()->getActiveWindow(),
-                           std::signbit(value) ? +1 : -1));
+      new GHOST_EventWheel(input->system->getMilliSeconds(), win, std::signbit(value) ? +1 : -1));
 }
 
 static const struct wl_pointer_listener pointer_listener = {
@@ -913,13 +956,15 @@ static void keyboard_keymap(
  * Notification that this seat's keyboard focus is on a certain
  * surface.
  */
-static void keyboard_enter(void * /*data*/,
+static void keyboard_enter(void *data,
                            struct wl_keyboard * /*wl_keyboard*/,
                            uint32_t /*serial*/,
-                           struct wl_surface * /*surface*/,
+                           struct wl_surface *surface,
                            struct wl_array * /*keys*/)
 {
-  /* pass */
+  if (surface != nullptr) {
+    static_cast<input_t *>(data)->focus_keyboard = surface;
+  }
 }
 
 /**
@@ -928,12 +973,14 @@ static void keyboard_enter(void * /*data*/,
  * Notification that this seat's keyboard focus is no longer on a
  * certain surface.
  */
-static void keyboard_leave(void * /*data*/,
+static void keyboard_leave(void *data,
                            struct wl_keyboard * /*wl_keyboard*/,
                            uint32_t /*serial*/,
-                           struct wl_surface * /*surface*/)
+                           struct wl_surface *surface)
 {
-  /* pass */
+  if (surface != nullptr) {
+    static_cast<input_t *>(data)->focus_keyboard = nullptr;
+  }
 }
 
 /**
@@ -988,6 +1035,14 @@ static void keyboard_key(void *data,
   }
   const GHOST_TKey gkey = xkb_map_gkey(sym);
 
+  /* Delete previous timer. */
+  if (xkb_keymap_key_repeats(xkb_state_get_keymap(input->xkb_state), key + 8) &&
+      input->key_repeat.timer) {
+    delete static_cast<key_repeat_payload_t *>(input->key_repeat.timer->getUserData());
+    input->system->removeTimer(input->key_repeat.timer);
+    input->key_repeat.timer = nullptr;
+  }
+
   GHOST_TEventKeyData key_data;
 
   if (etype == GHOST_kEventKeyDown) {
@@ -999,13 +1054,38 @@ static void keyboard_key(void *data,
   }
 
   input->data_source->source_serial = serial;
-  input->system->pushEvent(new GHOST_EventKey(input->system->getMilliSeconds(),
-                                              etype,
-                                              input->system->getWindowManager()->getActiveWindow(),
-                                              gkey,
-                                              '\0',
-                                              key_data.utf8_buf,
-                                              false));
+
+  GHOST_IWindow *win = static_cast<GHOST_WindowWayland *>(
+      wl_surface_get_user_data(input->focus_keyboard));
+  input->system->pushEvent(new GHOST_EventKey(
+      input->system->getMilliSeconds(), etype, win, gkey, '\0', key_data.utf8_buf, false));
+
+  /* Start timer for repeating key, if applicable. */
+  if (input->key_repeat.rate > 0 &&
+      xkb_keymap_key_repeats(xkb_state_get_keymap(input->xkb_state), key + 8) &&
+      etype == GHOST_kEventKeyDown) {
+
+    key_repeat_payload_t *payload = new key_repeat_payload_t({
+        .system = input->system,
+        .window = win,
+        .key = gkey,
+        .key_data = key_data,
+    });
+
+    auto cb = [](GHOST_ITimerTask *task, GHOST_TUns64 /*time*/) {
+      struct key_repeat_payload_t *payload = static_cast<key_repeat_payload_t *>(
+          task->getUserData());
+      payload->system->pushEvent(new GHOST_EventKey(payload->system->getMilliSeconds(),
+                                                    GHOST_kEventKeyDown,
+                                                    payload->window,
+                                                    payload->key,
+                                                    '\0',
+                                                    payload->key_data.utf8_buf,
+                                                    true));
+    };
+    input->key_repeat.timer = input->system->installTimer(
+        input->key_repeat.delay, 1000 / input->key_repeat.rate, cb, payload);
+  }
 }
 
 static void keyboard_modifiers(void *data,
@@ -1025,12 +1105,24 @@ static void keyboard_modifiers(void *data,
                         group);
 }
 
+static void keyboard_repeat_info(void *data,
+                                 struct wl_keyboard * /*wl_keyboard*/,
+                                 int32_t rate,
+                                 int32_t delay)
+{
+  input_t *input = static_cast<input_t *>(data);
+
+  input->key_repeat.rate = rate;
+  input->key_repeat.delay = delay;
+}
+
 static const struct wl_keyboard_listener keyboard_listener = {
     keyboard_keymap,
     keyboard_enter,
     keyboard_leave,
     keyboard_key,
     keyboard_modifiers,
+    keyboard_repeat_info,
 };
 
 static void seat_capabilities(void *data, struct wl_seat *wl_seat, uint32_t capabilities)
@@ -1163,7 +1255,7 @@ static void global_add(void *data,
     input->relative_pointer = nullptr;
     input->locked_pointer = nullptr;
     input->seat = static_cast<wl_seat *>(
-        wl_registry_bind(wl_registry, name, &wl_seat_interface, 2));
+        wl_registry_bind(wl_registry, name, &wl_seat_interface, 4));
     display->inputs.push_back(input);
     wl_seat_add_listener(input->seat, &seat_listener, input);
   }
@@ -1260,10 +1352,18 @@ GHOST_SystemWayland::~GHOST_SystemWayland()
   display_destroy(d);
 }
 
-bool GHOST_SystemWayland::processEvents(bool /*waitForEvent*/)
+bool GHOST_SystemWayland::processEvents(bool waitForEvent)
 {
-  wl_display_dispatch(d->display);
-  return true;
+  const bool fired = getTimerManager()->fireTimers(getMilliSeconds());
+
+  if (waitForEvent) {
+    wl_display_dispatch(d->display);
+  }
+  else {
+    wl_display_roundtrip(d->display);
+  }
+
+  return fired || (getEventManager()->getNumEvents() > 0);
 }
 
 int GHOST_SystemWayland::toggleConsole(int /*action*/)
@@ -1352,7 +1452,7 @@ GHOST_TUns8 GHOST_SystemWayland::getNumDisplays() const
 
 GHOST_TSuccess GHOST_SystemWayland::getCursorPosition(GHOST_TInt32 &x, GHOST_TInt32 &y) const
 {
-  if (getWindowManager()->getActiveWindow() != nullptr && !d->inputs.empty()) {
+  if (!d->inputs.empty() && (d->inputs[0]->focus_pointer != nullptr)) {
     x = d->inputs[0]->x;
     y = d->inputs[0]->y;
     return GHOST_kSuccess;
diff --git a/intern/guardedalloc/CMakeLists.txt b/intern/guardedalloc/CMakeLists.txt
index 1d4f846623c..cb24df65ba0 100644
--- a/intern/guardedalloc/CMakeLists.txt
+++ b/intern/guardedalloc/CMakeLists.txt
@@ -53,6 +53,14 @@ if(WIN32 AND NOT UNIX)
 
     mmap_win.h
   )
+
+  list(APPEND INC_SYS
+    ${PTHREADS_INC}
+  )
+
+  list(APPEND LIB
+    ${PTHREADS_LIBRARIES}
+  )
 endif()
 
 # Jemalloc 5.0.0+ needs extra configuration.
diff --git a/intern/guardedalloc/MEM_guardedalloc.h b/intern/guardedalloc/MEM_guardedalloc.h
index d5b109ee59f..602297576c8 100644
--- a/intern/guardedalloc/MEM_guardedalloc.h
+++ b/intern/guardedalloc/MEM_guardedalloc.h
@@ -145,14 +145,6 @@ extern void *(*MEM_mallocN_aligned)(size_t len,
                                     const char *str) /* ATTR_MALLOC */ ATTR_WARN_UNUSED_RESULT
     ATTR_ALLOC_SIZE(1) ATTR_NONNULL(3);
 
-/**
- * Same as callocN, clears memory and uses mmap (disk cached) if supported.
- * Can be free'd with MEM_freeN as usual.
- * */
-extern void *(*MEM_mapallocN)(size_t len,
-                              const char *str) /* ATTR_MALLOC */ ATTR_WARN_UNUSED_RESULT
-    ATTR_ALLOC_SIZE(1) ATTR_NONNULL(2);
-
 /** Print a list of the names and sizes of all allocated memory
  * blocks. as a python dict for easy investigation */
 extern void (*MEM_printmemlist_pydict)(void);
@@ -176,20 +168,11 @@ extern void (*MEM_set_error_callback)(void (*func)(const char *));
  * @retval true for correct memory, false for corrupted memory. */
 extern bool (*MEM_consistency_check)(void);
 
-/** Set thread locking functions for safe memory allocation from multiple
- * threads, pass NULL pointers to disable thread locking again. */
-extern void (*MEM_set_lock_callback)(void (*lock)(void), void (*unlock)(void));
-
 /** Attempt to enforce OSX (or other OS's) to have malloc and stack nonzero */
 extern void (*MEM_set_memory_debug)(void);
 
-/**
- * Memory usage stats
- * - MEM_get_memory_in_use is all memory
- * - MEM_get_mapped_memory_in_use is a subset of all memory */
+/** Memory usage stats. */
 extern size_t (*MEM_get_memory_in_use)(void);
-/** Get mapped memory usage. */
-extern size_t (*MEM_get_mapped_memory_in_use)(void);
 /** Get amount of memory blocks in use. */
 extern unsigned int (*MEM_get_memory_blocks_in_use)(void);
 
diff --git a/intern/guardedalloc/intern/mallocn.c b/intern/guardedalloc/intern/mallocn.c
index d24437c85f2..e85f8eb03ed 100644
--- a/intern/guardedalloc/intern/mallocn.c
+++ b/intern/guardedalloc/intern/mallocn.c
@@ -48,18 +48,14 @@ void *(*MEM_malloc_arrayN)(size_t len, size_t size, const char *str) = MEM_lockf
 void *(*MEM_mallocN_aligned)(size_t len,
                              size_t alignment,
                              const char *str) = MEM_lockfree_mallocN_aligned;
-void *(*MEM_mapallocN)(size_t len, const char *str) = MEM_lockfree_mapallocN;
 void (*MEM_printmemlist_pydict)(void) = MEM_lockfree_printmemlist_pydict;
 void (*MEM_printmemlist)(void) = MEM_lockfree_printmemlist;
 void (*MEM_callbackmemlist)(void (*func)(void *)) = MEM_lockfree_callbackmemlist;
 void (*MEM_printmemlist_stats)(void) = MEM_lockfree_printmemlist_stats;
 void (*MEM_set_error_callback)(void (*func)(const char *)) = MEM_lockfree_set_error_callback;
 bool (*MEM_consistency_check)(void) = MEM_lockfree_consistency_check;
-void (*MEM_set_lock_callback)(void (*lock)(void),
-                              void (*unlock)(void)) = MEM_lockfree_set_lock_callback;
 void (*MEM_set_memory_debug)(void) = MEM_lockfree_set_memory_debug;
 size_t (*MEM_get_memory_in_use)(void) = MEM_lockfree_get_memory_in_use;
-size_t (*MEM_get_mapped_memory_in_use)(void) = MEM_lockfree_get_mapped_memory_in_use;
 unsigned int (*MEM_get_memory_blocks_in_use)(void) = MEM_lockfree_get_memory_blocks_in_use;
 void (*MEM_reset_peak_memory)(void) = MEM_lockfree_reset_peak_memory;
 size_t (*MEM_get_peak_memory)(void) = MEM_lockfree_get_peak_memory;
@@ -111,17 +107,14 @@ void MEM_use_guarded_allocator(void)
   MEM_mallocN = MEM_guarded_mallocN;
   MEM_malloc_arrayN = MEM_guarded_malloc_arrayN;
   MEM_mallocN_aligned = MEM_guarded_mallocN_aligned;
-  MEM_mapallocN = MEM_guarded_mapallocN;
   MEM_printmemlist_pydict = MEM_guarded_printmemlist_pydict;
   MEM_printmemlist = MEM_guarded_printmemlist;
   MEM_callbackmemlist = MEM_guarded_callbackmemlist;
   MEM_printmemlist_stats = MEM_guarded_printmemlist_stats;
   MEM_set_error_callback = MEM_guarded_set_error_callback;
   MEM_consistency_check = MEM_guarded_consistency_check;
-  MEM_set_lock_callback = MEM_guarded_set_lock_callback;
   MEM_set_memory_debug = MEM_guarded_set_memory_debug;
   MEM_get_memory_in_use = MEM_guarded_get_memory_in_use;
-  MEM_get_mapped_memory_in_use = MEM_guarded_get_mapped_memory_in_use;
   MEM_get_memory_blocks_in_use = MEM_guarded_get_memory_blocks_in_use;
   MEM_reset_peak_memory = MEM_guarded_reset_peak_memory;
   MEM_get_peak_memory = MEM_guarded_get_peak_memory;
diff --git a/intern/guardedalloc/intern/mallocn_guarded_impl.c b/intern/guardedalloc/intern/mallocn_guarded_impl.c
index f601609c6e0..20dcbed7235 100644
--- a/intern/guardedalloc/intern/mallocn_guarded_impl.c
+++ b/intern/guardedalloc/intern/mallocn_guarded_impl.c
@@ -28,6 +28,8 @@
 #include <string.h> /* memcpy */
 #include <sys/types.h>
 
+#include <pthread.h>
+
 #include "MEM_guardedalloc.h"
 
 /* to ensure strict conversions */
@@ -51,17 +53,6 @@
 //#define DEBUG_MEMCOUNTER
 
 /* Only for debugging:
- * defining DEBUG_THREADS will enable check whether memory manager
- * is locked with a mutex when allocation is called from non-main
- * thread.
- *
- * This helps troubleshooting memory issues caused by the fact
- * guarded allocator is not thread-safe, however this check will
- * fail to check allocations from openmp threads.
- */
-//#define DEBUG_THREADS
-
-/* Only for debugging:
  * Defining DEBUG_BACKTRACE will store a backtrace from where
  * memory block was allocated and print this trace for all
  * unfreed blocks.
@@ -104,7 +95,7 @@ typedef struct MemHead {
   const char *name;
   const char *nextname;
   int tag2;
-  short mmap;      /* if true, memory was mmapped */
+  short pad1;
   short alignment; /* if non-zero aligned alloc was used
                     * and alignment is stored here.
                     */
@@ -124,24 +115,6 @@ typedef struct MemHead {
 
 typedef MemHead MemHeadAligned;
 
-/* for openmp threading asserts, saves time troubleshooting
- * we may need to extend this if blender code starts using MEM_
- * functions inside OpenMP correctly with omp_set_lock() */
-
-#if 0 /* disable for now, only use to debug openmp code which doesn lock threads for malloc */
-#  if defined(_OPENMP) && defined(DEBUG)
-#    include <assert.h>
-#    include <omp.h>
-#    define DEBUG_OMP_MALLOC
-#  endif
-#endif
-
-#ifdef DEBUG_THREADS
-#  include <assert.h>
-#  include <pthread.h>
-static pthread_t mainid;
-#endif
-
 #ifdef DEBUG_BACKTRACE
 #  if defined(__linux__) || defined(__APPLE__)
 #    include <execinfo.h>
@@ -187,13 +160,11 @@ static const char *check_memlist(MemHead *memh);
 /* --------------------------------------------------------------------- */
 
 static unsigned int totblock = 0;
-static size_t mem_in_use = 0, mmap_in_use = 0, peak_mem = 0;
+static size_t mem_in_use = 0, peak_mem = 0;
 
 static volatile struct localListBase _membase;
 static volatile struct localListBase *membase = &_membase;
 static void (*error_callback)(const char *) = NULL;
-static void (*thread_lock_callback)(void) = NULL;
-static void (*thread_unlock_callback)(void) = NULL;
 
 static bool malloc_debug_memset = false;
 
@@ -233,40 +204,16 @@ print_error(const char *str, ...)
     fputs(buf, stderr);
 }
 
+static pthread_mutex_t thread_lock = PTHREAD_MUTEX_INITIALIZER;
+
 static void mem_lock_thread(void)
 {
-#ifdef DEBUG_THREADS
-  static int initialized = 0;
-
-  if (initialized == 0) {
-    /* assume first allocation happens from main thread */
-    mainid = pthread_self();
-    initialized = 1;
-  }
-
-  if (!pthread_equal(pthread_self(), mainid) && thread_lock_callback == NULL) {
-    assert(!"Memory function is called from non-main thread without lock");
-  }
-#endif
-
-#ifdef DEBUG_OMP_MALLOC
-  assert(omp_in_parallel() == 0);
-#endif
-
-  if (thread_lock_callback)
-    thread_lock_callback();
+  pthread_mutex_lock(&thread_lock);
 }
 
 static void mem_unlock_thread(void)
 {
-#ifdef DEBUG_THREADS
-  if (!pthread_equal(pthread_self(), mainid) && thread_lock_callback == NULL) {
-    assert(!"Thread lock was removed while allocation from thread is in progress");
-  }
-#endif
-
-  if (thread_unlock_callback)
-    thread_unlock_callback();
+  pthread_mutex_unlock(&thread_lock);
 }
 
 bool MEM_guarded_consistency_check(void)
@@ -287,12 +234,6 @@ void MEM_guarded_set_error_callback(void (*func)(const char *))
   error_callback = func;
 }
 
-void MEM_guarded_set_lock_callback(void (*lock)(void), void (*unlock)(void))
-{
-  thread_lock_callback = lock;
-  thread_unlock_callback = unlock;
-}
-
 void MEM_guarded_set_memory_debug(void)
 {
   malloc_debug_memset = true;
@@ -320,10 +261,8 @@ void *MEM_guarded_dupallocN(const void *vmemh)
     memh--;
 
 #ifndef DEBUG_MEMDUPLINAME
-    if (UNLIKELY(memh->mmap))
-      newp = MEM_guarded_mapallocN(memh->len, "dupli_mapalloc");
-    else if (LIKELY(memh->alignment == 0))
-      newp = MEM_guarded_mapallocN(memh->len, "dupli_mapalloc");
+    if (LIKELY(memh->alignment == 0))
+      newp = MEM_guarded_mallocN(memh->len, "dupli_alloc");
     else
       newp = MEM_guarded_mallocN_aligned(memh->len, (size_t)memh->alignment, "dupli_alloc");
 
@@ -334,11 +273,7 @@ void *MEM_guarded_dupallocN(const void *vmemh)
       MemHead *nmemh;
       char *name = malloc(strlen(memh->name) + 24);
 
-      if (UNLIKELY(memh->mmap)) {
-        sprintf(name, "%s %s", "dupli_mapalloc", memh->name);
-        newp = MEM_guarded_mapallocN(memh->len, name);
-      }
-      else if (LIKELY(memh->alignment == 0)) {
+      if (LIKELY(memh->alignment == 0)) {
         sprintf(name, "%s %s", "dupli_alloc", memh->name);
         newp = MEM_guarded_mallocN(memh->len, name);
       }
@@ -478,7 +413,7 @@ static void make_memhead_header(MemHead *memh, size_t len, const char *str)
   memh->name = str;
   memh->nextname = NULL;
   memh->len = len;
-  memh->mmap = 0;
+  memh->pad1 = 0;
   memh->alignment = 0;
   memh->tag2 = MEMTAG2;
 
@@ -646,58 +581,6 @@ void *MEM_guarded_calloc_arrayN(size_t len, size_t size, const char *str)
   return MEM_guarded_callocN(total_size, str);
 }
 
-/* note; mmap returns zero'd memory */
-void *MEM_guarded_mapallocN(size_t len, const char *str)
-{
-  MemHead *memh;
-
-  /* on 64 bit, simply use calloc instead, as mmap does not support
-   * allocating > 4 GB on Windows. the only reason mapalloc exists
-   * is to get around address space limitations in 32 bit OSes. */
-  if (sizeof(void *) >= 8)
-    return MEM_guarded_callocN(len, str);
-
-  len = SIZET_ALIGN_4(len);
-
-#if defined(WIN32)
-  /* our windows mmap implementation is not thread safe */
-  mem_lock_thread();
-#endif
-  memh = mmap(NULL,
-              len + sizeof(MemHead) + sizeof(MemTail),
-              PROT_READ | PROT_WRITE,
-              MAP_SHARED | MAP_ANON,
-              -1,
-              0);
-#if defined(WIN32)
-  mem_unlock_thread();
-#endif
-
-  if (memh != (MemHead *)-1) {
-    make_memhead_header(memh, len, str);
-    memh->mmap = 1;
-    atomic_add_and_fetch_z(&mmap_in_use, len);
-    mem_lock_thread();
-    peak_mem = mmap_in_use > peak_mem ? mmap_in_use : peak_mem;
-    mem_unlock_thread();
-#ifdef DEBUG_MEMCOUNTER
-    if (_mallocn_count == DEBUG_MEMCOUNTER_ERROR_VAL)
-      memcount_raise(__func__);
-    memh->_count = _mallocn_count++;
-#endif
-    return (++memh);
-  }
-  else {
-    print_error(
-        "Mapalloc returns null, fallback to regular malloc: "
-        "len=" SIZET_FORMAT " in %s, total %u\n",
-        SIZET_ARG(len),
-        str,
-        (unsigned int)mmap_in_use);
-    return MEM_guarded_callocN(len, str);
-  }
-}
-
 /* Memory statistics print */
 typedef struct MemPrintBlock {
   const char *name;
@@ -765,7 +648,7 @@ void MEM_guarded_printmemlist_stats(void)
     pb++;
 
 #ifdef USE_MALLOC_USABLE_SIZE
-    if (!membl->mmap && membl->alignment == 0) {
+    if (membl->alignment == 0) {
       mem_in_use_slop += (sizeof(MemHead) + sizeof(MemTail) + malloc_usable_size((void *)membl)) -
                          membl->len;
     }
@@ -1098,27 +981,13 @@ static void rem_memblock(MemHead *memh)
     free((char *)memh->name);
 #endif
 
-  if (memh->mmap) {
-    atomic_sub_and_fetch_z(&mmap_in_use, memh->len);
-#if defined(WIN32)
-    /* our windows mmap implementation is not thread safe */
-    mem_lock_thread();
-#endif
-    if (munmap(memh, memh->len + sizeof(MemHead) + sizeof(MemTail)))
-      printf("Couldn't unmap memory %s\n", memh->name);
-#if defined(WIN32)
-    mem_unlock_thread();
-#endif
+  if (UNLIKELY(malloc_debug_memset && memh->len))
+    memset(memh + 1, 255, memh->len);
+  if (LIKELY(memh->alignment == 0)) {
+    free(memh);
   }
   else {
-    if (UNLIKELY(malloc_debug_memset && memh->len))
-      memset(memh + 1, 255, memh->len);
-    if (LIKELY(memh->alignment == 0)) {
-      free(memh);
-    }
-    else {
-      aligned_free(MEMHEAD_REAL_PTR(memh));
-    }
+    aligned_free(MEMHEAD_REAL_PTR(memh));
   }
 }
 
@@ -1270,17 +1139,6 @@ size_t MEM_guarded_get_memory_in_use(void)
   return _mem_in_use;
 }
 
-size_t MEM_guarded_get_mapped_memory_in_use(void)
-{
-  size_t _mmap_in_use;
-
-  mem_lock_thread();
-  _mmap_in_use = mmap_in_use;
-  mem_unlock_thread();
-
-  return _mmap_in_use;
-}
-
 unsigned int MEM_guarded_get_memory_blocks_in_use(void)
 {
   unsigned int _totblock;
diff --git a/intern/guardedalloc/intern/mallocn_intern.h b/intern/guardedalloc/intern/mallocn_intern.h
index 876607fdb77..ef8845a66b3 100644
--- a/intern/guardedalloc/intern/mallocn_intern.h
+++ b/intern/guardedalloc/intern/mallocn_intern.h
@@ -24,13 +24,6 @@
 #ifndef __MALLOCN_INTERN_H__
 #define __MALLOCN_INTERN_H__
 
-/* mmap exception */
-#if defined(WIN32)
-#  include "mmap_win.h"
-#else
-#  include <sys/mman.h>
-#endif
-
 #ifdef __GNUC__
 #  define UNUSED(x) UNUSED_##x __attribute__((__unused__))
 #else
@@ -140,19 +133,14 @@ void *MEM_lockfree_mallocN_aligned(size_t len,
                                    size_t alignment,
                                    const char *UNUSED(str)) ATTR_MALLOC ATTR_WARN_UNUSED_RESULT
     ATTR_ALLOC_SIZE(1) ATTR_NONNULL(3);
-void *MEM_lockfree_mapallocN(size_t len,
-                             const char *UNUSED(str)) ATTR_MALLOC ATTR_WARN_UNUSED_RESULT
-    ATTR_ALLOC_SIZE(1) ATTR_NONNULL(2);
 void MEM_lockfree_printmemlist_pydict(void);
 void MEM_lockfree_printmemlist(void);
 void MEM_lockfree_callbackmemlist(void (*func)(void *));
 void MEM_lockfree_printmemlist_stats(void);
 void MEM_lockfree_set_error_callback(void (*func)(const char *));
 bool MEM_lockfree_consistency_check(void);
-void MEM_lockfree_set_lock_callback(void (*lock)(void), void (*unlock)(void));
 void MEM_lockfree_set_memory_debug(void);
 size_t MEM_lockfree_get_memory_in_use(void);
-size_t MEM_lockfree_get_mapped_memory_in_use(void);
 unsigned int MEM_lockfree_get_memory_blocks_in_use(void);
 void MEM_lockfree_reset_peak_memory(void);
 size_t MEM_lockfree_get_peak_memory(void) ATTR_WARN_UNUSED_RESULT;
@@ -188,19 +176,14 @@ void *MEM_guarded_mallocN_aligned(size_t len,
                                   size_t alignment,
                                   const char *UNUSED(str)) ATTR_MALLOC ATTR_WARN_UNUSED_RESULT
     ATTR_ALLOC_SIZE(1) ATTR_NONNULL(3);
-void *MEM_guarded_mapallocN(size_t len,
-                            const char *UNUSED(str)) ATTR_MALLOC ATTR_WARN_UNUSED_RESULT
-    ATTR_ALLOC_SIZE(1) ATTR_NONNULL(2);
 void MEM_guarded_printmemlist_pydict(void);
 void MEM_guarded_printmemlist(void);
 void MEM_guarded_callbackmemlist(void (*func)(void *));
 void MEM_guarded_printmemlist_stats(void);
 void MEM_guarded_set_error_callback(void (*func)(const char *));
 bool MEM_guarded_consistency_check(void);
-void MEM_guarded_set_lock_callback(void (*lock)(void), void (*unlock)(void));
 void MEM_guarded_set_memory_debug(void);
 size_t MEM_guarded_get_memory_in_use(void);
-size_t MEM_guarded_get_mapped_memory_in_use(void);
 unsigned int MEM_guarded_get_memory_blocks_in_use(void);
 void MEM_guarded_reset_peak_memory(void);
 size_t MEM_guarded_get_peak_memory(void) ATTR_WARN_UNUSED_RESULT;
diff --git a/intern/guardedalloc/intern/mallocn_lockfree_impl.c b/intern/guardedalloc/intern/mallocn_lockfree_impl.c
index ab7d9097669..205cc688d72 100644
--- a/intern/guardedalloc/intern/mallocn_lockfree_impl.c
+++ b/intern/guardedalloc/intern/mallocn_lockfree_impl.c
@@ -44,22 +44,18 @@ typedef struct MemHeadAligned {
 } MemHeadAligned;
 
 static unsigned int totblock = 0;
-static size_t mem_in_use = 0, mmap_in_use = 0, peak_mem = 0;
+static size_t mem_in_use = 0, peak_mem = 0;
 static bool malloc_debug_memset = false;
 
 static void (*error_callback)(const char *) = NULL;
-static void (*thread_lock_callback)(void) = NULL;
-static void (*thread_unlock_callback)(void) = NULL;
 
 enum {
-  MEMHEAD_MMAP_FLAG = 1,
-  MEMHEAD_ALIGN_FLAG = 2,
+  MEMHEAD_ALIGN_FLAG = 1,
 };
 
 #define MEMHEAD_FROM_PTR(ptr) (((MemHead *)ptr) - 1)
 #define PTR_FROM_MEMHEAD(memhead) (memhead + 1)
 #define MEMHEAD_ALIGNED_FROM_PTR(ptr) (((MemHeadAligned *)ptr) - 1)
-#define MEMHEAD_IS_MMAP(memhead) ((memhead)->len & (size_t)MEMHEAD_MMAP_FLAG)
 #define MEMHEAD_IS_ALIGNED(memhead) ((memhead)->len & (size_t)MEMHEAD_ALIGN_FLAG)
 
 /* Uncomment this to have proper peak counter. */
@@ -93,24 +89,10 @@ print_error(const char *str, ...)
   }
 }
 
-#if defined(WIN32)
-static void mem_lock_thread(void)
-{
-  if (thread_lock_callback)
-    thread_lock_callback();
-}
-
-static void mem_unlock_thread(void)
-{
-  if (thread_unlock_callback)
-    thread_unlock_callback();
-}
-#endif
-
 size_t MEM_lockfree_allocN_len(const void *vmemh)
 {
   if (vmemh) {
-    return MEMHEAD_FROM_PTR(vmemh)->len & ~((size_t)(MEMHEAD_MMAP_FLAG | MEMHEAD_ALIGN_FLAG));
+    return MEMHEAD_FROM_PTR(vmemh)->len & ~((size_t)(MEMHEAD_ALIGN_FLAG));
   }
   else {
     return 0;
@@ -133,29 +115,15 @@ void MEM_lockfree_freeN(void *vmemh)
   atomic_sub_and_fetch_u(&totblock, 1);
   atomic_sub_and_fetch_z(&mem_in_use, len);
 
-  if (MEMHEAD_IS_MMAP(memh)) {
-    atomic_sub_and_fetch_z(&mmap_in_use, len);
-#if defined(WIN32)
-    /* our windows mmap implementation is not thread safe */
-    mem_lock_thread();
-#endif
-    if (munmap(memh, len + sizeof(MemHead)))
-      printf("Couldn't unmap memory\n");
-#if defined(WIN32)
-    mem_unlock_thread();
-#endif
+  if (UNLIKELY(malloc_debug_memset && len)) {
+    memset(memh + 1, 255, len);
+  }
+  if (UNLIKELY(MEMHEAD_IS_ALIGNED(memh))) {
+    MemHeadAligned *memh_aligned = MEMHEAD_ALIGNED_FROM_PTR(vmemh);
+    aligned_free(MEMHEAD_REAL_PTR(memh_aligned));
   }
   else {
-    if (UNLIKELY(malloc_debug_memset && len)) {
-      memset(memh + 1, 255, len);
-    }
-    if (UNLIKELY(MEMHEAD_IS_ALIGNED(memh))) {
-      MemHeadAligned *memh_aligned = MEMHEAD_ALIGNED_FROM_PTR(vmemh);
-      aligned_free(MEMHEAD_REAL_PTR(memh_aligned));
-    }
-    else {
-      free(memh);
-    }
+    free(memh);
   }
 }
 
@@ -165,10 +133,7 @@ void *MEM_lockfree_dupallocN(const void *vmemh)
   if (vmemh) {
     MemHead *memh = MEMHEAD_FROM_PTR(vmemh);
     const size_t prev_size = MEM_lockfree_allocN_len(vmemh);
-    if (UNLIKELY(MEMHEAD_IS_MMAP(memh))) {
-      newp = MEM_lockfree_mapallocN(prev_size, "dupli_mapalloc");
-    }
-    else if (UNLIKELY(MEMHEAD_IS_ALIGNED(memh))) {
+    if (UNLIKELY(MEMHEAD_IS_ALIGNED(memh))) {
       MemHeadAligned *memh_aligned = MEMHEAD_ALIGNED_FROM_PTR(vmemh);
       newp = MEM_lockfree_mallocN_aligned(
           prev_size, (size_t)memh_aligned->alignment, "dupli_malloc");
@@ -397,47 +362,6 @@ void *MEM_lockfree_mallocN_aligned(size_t len, size_t alignment, const char *str
   return NULL;
 }
 
-void *MEM_lockfree_mapallocN(size_t len, const char *str)
-{
-  MemHead *memh;
-
-  /* on 64 bit, simply use calloc instead, as mmap does not support
-   * allocating > 4 GB on Windows. the only reason mapalloc exists
-   * is to get around address space limitations in 32 bit OSes. */
-  if (sizeof(void *) >= 8)
-    return MEM_lockfree_callocN(len, str);
-
-  len = SIZET_ALIGN_4(len);
-
-#if defined(WIN32)
-  /* our windows mmap implementation is not thread safe */
-  mem_lock_thread();
-#endif
-  memh = mmap(NULL, len + sizeof(MemHead), PROT_READ | PROT_WRITE, MAP_SHARED | MAP_ANON, -1, 0);
-#if defined(WIN32)
-  mem_unlock_thread();
-#endif
-
-  if (memh != (MemHead *)-1) {
-    memh->len = len | (size_t)MEMHEAD_MMAP_FLAG;
-    atomic_add_and_fetch_u(&totblock, 1);
-    atomic_add_and_fetch_z(&mem_in_use, len);
-    atomic_add_and_fetch_z(&mmap_in_use, len);
-
-    update_maximum(&peak_mem, mem_in_use);
-    update_maximum(&peak_mem, mmap_in_use);
-
-    return PTR_FROM_MEMHEAD(memh);
-  }
-  print_error(
-      "Mapalloc returns null, fallback to regular malloc: "
-      "len=" SIZET_FORMAT " in %s, total %u\n",
-      SIZET_ARG(len),
-      str,
-      (unsigned int)mmap_in_use);
-  return MEM_lockfree_callocN(len, str);
-}
-
 void MEM_lockfree_printmemlist_pydict(void)
 {
 }
@@ -476,12 +400,6 @@ bool MEM_lockfree_consistency_check(void)
   return true;
 }
 
-void MEM_lockfree_set_lock_callback(void (*lock)(void), void (*unlock)(void))
-{
-  thread_lock_callback = lock;
-  thread_unlock_callback = unlock;
-}
-
 void MEM_lockfree_set_memory_debug(void)
 {
   malloc_debug_memset = true;
@@ -492,11 +410,6 @@ size_t MEM_lockfree_get_memory_in_use(void)
   return mem_in_use;
 }
 
-size_t MEM_lockfree_get_mapped_memory_in_use(void)
-{
-  return mmap_in_use;
-}
-
 unsigned int MEM_lockfree_get_memory_blocks_in_use(void)
 {
   return totblock;
diff --git a/intern/libmv/libmv/autotrack/reconstruction.h b/intern/libmv/libmv/autotrack/reconstruction.h
index 7b34a0951a3..732e74063f1 100644
--- a/intern/libmv/libmv/autotrack/reconstruction.h
+++ b/intern/libmv/libmv/autotrack/reconstruction.h
@@ -23,6 +23,7 @@
 #ifndef LIBMV_AUTOTRACK_RECONSTRUCTION_H_
 #define LIBMV_AUTOTRACK_RECONSTRUCTION_H_
 
+#include "libmv/base/map.h"
 #include "libmv/base/vector.h"
 #include "libmv/numeric/numeric.h"
 #include "libmv/simple_pipeline/camera_intrinsics.h"
@@ -75,7 +76,7 @@ class Reconstruction {
   vector<CameraIntrinsics*> camera_intrinsics_;
 
   // Indexed by Marker::clip then by Marker::frame.
-  vector<vector<CameraPose> > camera_poses_;
+  vector<map<int, CameraPose>> camera_poses_;
 
   // Indexed by Marker::track.
   vector<Point> points_;
diff --git a/intern/libmv/libmv/simple_pipeline/bundle.cc b/intern/libmv/libmv/simple_pipeline/bundle.cc
index a70fdbc9888..2ecc0505e1f 100644
--- a/intern/libmv/libmv/simple_pipeline/bundle.cc
+++ b/intern/libmv/libmv/simple_pipeline/bundle.cc
@@ -24,6 +24,7 @@
 
 #include "ceres/ceres.h"
 #include "ceres/rotation.h"
+#include "libmv/base/map.h"
 #include "libmv/base/vector.h"
 #include "libmv/logging/logging.h"
 #include "libmv/multiview/fundamental.h"
@@ -407,47 +408,39 @@ void UnpackIntrinsicsFromArray(const double intrinsics_block[OFFSET_MAX],
 // Get a vector of camera's rotations denoted by angle axis
 // conjuncted with translations into single block
 //
-// Element with index i matches to a rotation+translation for
+// Element with key i matches to a rotation+translation for
 // camera at image i.
-vector<Vec6> PackCamerasRotationAndTranslation(
-    const Tracks &tracks,
+map<int, Vec6> PackCamerasRotationAndTranslation(
     const EuclideanReconstruction &reconstruction) {
-  vector<Vec6> all_cameras_R_t;
-  int max_image = tracks.MaxImage();
-
-  all_cameras_R_t.resize(max_image + 1);
-
-  for (int i = 0; i <= max_image; i++) {
-    const EuclideanCamera *camera = reconstruction.CameraForImage(i);
-
-    if (!camera) {
-      continue;
-    }
-
-    ceres::RotationMatrixToAngleAxis(&camera->R(0, 0),
-                                     &all_cameras_R_t[i](0));
-    all_cameras_R_t[i].tail<3>() = camera->t;
+  map<int, Vec6> all_cameras_R_t;
+
+  vector<EuclideanCamera> all_cameras = reconstruction.AllCameras();
+  for (const EuclideanCamera& camera : all_cameras) {
+    Vec6 camera_R_t;
+    ceres::RotationMatrixToAngleAxis(&camera.R(0, 0), &camera_R_t(0));
+    camera_R_t.tail<3>() = camera.t;
+    all_cameras_R_t.insert(make_pair(camera.image, camera_R_t));
   }
+
   return all_cameras_R_t;
 }
 
 // Convert cameras rotations fro mangle axis back to rotation matrix.
 void UnpackCamerasRotationAndTranslation(
-    const Tracks &tracks,
-    const vector<Vec6> &all_cameras_R_t,
+    const map<int, Vec6> &all_cameras_R_t,
     EuclideanReconstruction *reconstruction) {
-  int max_image = tracks.MaxImage();
 
-  for (int i = 0; i <= max_image; i++) {
-    EuclideanCamera *camera = reconstruction->CameraForImage(i);
+  for (map<int, Vec6>::value_type image_and_camera_R_T : all_cameras_R_t) {
+    const int image = image_and_camera_R_T.first;
+    const Vec6& camera_R_t = image_and_camera_R_T.second;
 
+    EuclideanCamera *camera = reconstruction->CameraForImage(image);
     if (!camera) {
       continue;
     }
 
-    ceres::AngleAxisToRotationMatrix(&all_cameras_R_t[i](0),
-                                     &camera->R(0, 0));
-    camera->t = all_cameras_R_t[i].tail<3>();
+    ceres::AngleAxisToRotationMatrix(&camera_R_t(0), &camera->R(0, 0));
+    camera->t = camera_R_t.tail<3>();
   }
 }
 
@@ -476,7 +469,7 @@ void CRSMatrixToEigenMatrix(const ceres::CRSMatrix &crs_matrix,
 
 void EuclideanBundlerPerformEvaluation(const Tracks &tracks,
                                        EuclideanReconstruction *reconstruction,
-                                       vector<Vec6> *all_cameras_R_t,
+                                       map<int, Vec6> *all_cameras_R_t,
                                        ceres::Problem *problem,
                                        BundleEvaluation *evaluation) {
   int max_track = tracks.MaxTrack();
@@ -603,7 +596,7 @@ void AddResidualBlockToProblem(const CameraIntrinsics *invariant_intrinsics,
 // are to be totally still here.
 void EuclideanBundlePointsOnly(const CameraIntrinsics *invariant_intrinsics,
                                const vector<Marker> &markers,
-                               vector<Vec6> &all_cameras_R_t,
+                               map<int, Vec6> &all_cameras_R_t,
                                double intrinsics_block[OFFSET_MAX],
                                EuclideanReconstruction *reconstruction) {
   ceres::Problem::Options problem_options;
@@ -699,8 +692,8 @@ void EuclideanBundleCommonIntrinsics(
   //
   // Block for minimization has got the following structure:
   //   <3 elements for angle-axis> <3 elements for translation>
-  vector<Vec6> all_cameras_R_t =
-    PackCamerasRotationAndTranslation(tracks, *reconstruction);
+  map<int, Vec6> all_cameras_R_t =
+    PackCamerasRotationAndTranslation(*reconstruction);
 
   // Parameterization used to restrict camera motion for modal solvers.
   ceres::SubsetParameterization *constant_translation_parameterization = NULL;
@@ -827,9 +820,7 @@ void EuclideanBundleCommonIntrinsics(
   LG << "Final report:\n" << summary.FullReport();
 
   // Copy rotations and translations back.
-  UnpackCamerasRotationAndTranslation(tracks,
-                                      all_cameras_R_t,
-                                      reconstruction);
+  UnpackCamerasRotationAndTranslation(all_cameras_R_t, reconstruction);
 
   // Copy intrinsics back.
   if (bundle_intrinsics != BUNDLE_NO_INTRINSICS)
diff --git a/intern/libmv/libmv/simple_pipeline/reconstruction.cc b/intern/libmv/libmv/simple_pipeline/reconstruction.cc
index 65e5dd27d5d..851eedb5bb1 100644
--- a/intern/libmv/libmv/simple_pipeline/reconstruction.cc
+++ b/intern/libmv/libmv/simple_pipeline/reconstruction.cc
@@ -27,14 +27,14 @@ namespace libmv {
 EuclideanReconstruction::EuclideanReconstruction() {}
 EuclideanReconstruction::EuclideanReconstruction(
     const EuclideanReconstruction &other) {
-  cameras_ = other.cameras_;
+  image_to_cameras_map_ = other.image_to_cameras_map_;
   points_ = other.points_;
 }
 
 EuclideanReconstruction &EuclideanReconstruction::operator=(
     const EuclideanReconstruction &other) {
   if (&other != this) {
-    cameras_ = other.cameras_;
+    image_to_cameras_map_ = other.image_to_cameras_map_;
     points_ = other.points_;
   }
   return *this;
@@ -44,12 +44,13 @@ void EuclideanReconstruction::InsertCamera(int image,
                                            const Mat3 &R,
                                            const Vec3 &t) {
   LG << "InsertCamera " << image << ":\nR:\n"<< R << "\nt:\n" << t;
-  if (image >= cameras_.size()) {
-    cameras_.resize(image + 1);
-  }
-  cameras_[image].image = image;
-  cameras_[image].R = R;
-  cameras_[image].t = t;
+
+  EuclideanCamera camera;
+  camera.image = image;
+  camera.R = R;
+  camera.t = t;
+
+  image_to_cameras_map_.insert(make_pair(image, camera));
 }
 
 void EuclideanReconstruction::InsertPoint(int track, const Vec3 &X) {
@@ -69,22 +70,18 @@ EuclideanCamera *EuclideanReconstruction::CameraForImage(int image) {
 
 const EuclideanCamera *EuclideanReconstruction::CameraForImage(
     int image) const {
-  if (image < 0 || image >= cameras_.size()) {
+  ImageToCameraMap::const_iterator it = image_to_cameras_map_.find(image);
+  if (it == image_to_cameras_map_.end()) {
     return NULL;
   }
-  const EuclideanCamera *camera = &cameras_[image];
-  if (camera->image == -1) {
-    return NULL;
-  }
-  return camera;
+  return &it->second;
 }
 
 vector<EuclideanCamera> EuclideanReconstruction::AllCameras() const {
   vector<EuclideanCamera> cameras;
-  for (int i = 0; i < cameras_.size(); ++i) {
-    if (cameras_[i].image != -1) {
-      cameras.push_back(cameras_[i]);
-    }
+  for (const ImageToCameraMap::value_type& image_and_camera :
+       image_to_cameras_map_) {
+    cameras.push_back(image_and_camera.second);
   }
   return cameras;
 }
@@ -115,14 +112,14 @@ vector<EuclideanPoint> EuclideanReconstruction::AllPoints() const {
   return points;
 }
 
-void ProjectiveReconstruction::InsertCamera(int image,
-                                           const Mat34 &P) {
+void ProjectiveReconstruction::InsertCamera(int image, const Mat34 &P) {
   LG << "InsertCamera " << image << ":\nP:\n"<< P;
-  if (image >= cameras_.size()) {
-    cameras_.resize(image + 1);
-  }
-  cameras_[image].image = image;
-  cameras_[image].P = P;
+
+  ProjectiveCamera camera;
+  camera.image = image;
+  camera.P = P;
+
+  image_to_cameras_map_.insert(make_pair(image, camera));
 }
 
 void ProjectiveReconstruction::InsertPoint(int track, const Vec4 &X) {
@@ -142,22 +139,18 @@ ProjectiveCamera *ProjectiveReconstruction::CameraForImage(int image) {
 
 const ProjectiveCamera *ProjectiveReconstruction::CameraForImage(
     int image) const {
-  if (image < 0 || image >= cameras_.size()) {
-    return NULL;
+  ImageToCameraMap::const_iterator it = image_to_cameras_map_.find(image);
+  if (it == image_to_cameras_map_.end()) {
+    return  NULL;
   }
-  const ProjectiveCamera *camera = &cameras_[image];
-  if (camera->image == -1) {
-    return NULL;
-  }
-  return camera;
+  return &it->second;
 }
 
 vector<ProjectiveCamera> ProjectiveReconstruction::AllCameras() const {
   vector<ProjectiveCamera> cameras;
-  for (int i = 0; i < cameras_.size(); ++i) {
-    if (cameras_[i].image != -1) {
-      cameras.push_back(cameras_[i]);
-    }
+  for (const ImageToCameraMap::value_type& image_and_camera :
+       image_to_cameras_map_) {
+    cameras.push_back(image_and_camera.second);
   }
   return cameras;
 }
diff --git a/intern/libmv/libmv/simple_pipeline/reconstruction.h b/intern/libmv/libmv/simple_pipeline/reconstruction.h
index 79c693c5e6d..544aeac042e 100644
--- a/intern/libmv/libmv/simple_pipeline/reconstruction.h
+++ b/intern/libmv/libmv/simple_pipeline/reconstruction.h
@@ -22,6 +22,7 @@
 #define LIBMV_SIMPLE_PIPELINE_RECONSTRUCTION_H_
 
 #include "libmv/base/vector.h"
+#include "libmv/base/map.h"
 #include "libmv/numeric/numeric.h"
 
 namespace libmv {
@@ -120,7 +121,11 @@ class EuclideanReconstruction {
   vector<EuclideanPoint> AllPoints() const;
 
  private:
-  vector<EuclideanCamera> cameras_;
+  // Indexed by frame number.
+  typedef map<int, EuclideanCamera> ImageToCameraMap;
+  ImageToCameraMap image_to_cameras_map_;
+
+  // Insxed by track.
   vector<EuclideanPoint> points_;
 };
 
@@ -208,7 +213,11 @@ class ProjectiveReconstruction {
   vector<ProjectivePoint> AllPoints() const;
 
  private:
-  vector<ProjectiveCamera> cameras_;
+  // Indexed by frame number.
+  typedef map<int, ProjectiveCamera> ImageToCameraMap;
+  ImageToCameraMap image_to_cameras_map_;
+
+  // Indexed by track.
   vector<ProjectivePoint> points_;
 };
 
diff --git a/intern/mantaflow/intern/MANTA_main.cpp b/intern/mantaflow/intern/MANTA_main.cpp
index 11ce32fb828..95013958561 100644
--- a/intern/mantaflow/intern/MANTA_main.cpp
+++ b/intern/mantaflow/intern/MANTA_main.cpp
@@ -81,6 +81,7 @@ MANTA::MANTA(int *res, FluidModifierData *mmd) : mCurrentID(++solverID)
   mUsingNoise = (mds->flags & FLUID_DOMAIN_USE_NOISE) && mUsingSmoke;
   mUsingFractions = (mds->flags & FLUID_DOMAIN_USE_FRACTIONS) && mUsingLiquid;
   mUsingMesh = (mds->flags & FLUID_DOMAIN_USE_MESH) && mUsingLiquid;
+  mUsingDiffusion = (mds->flags & FLUID_DOMAIN_USE_DIFFUSION) && mUsingLiquid;
   mUsingMVel = (mds->flags & FLUID_DOMAIN_USE_SPEED_VECTORS) && mUsingLiquid;
   mUsingGuiding = (mds->flags & FLUID_DOMAIN_USE_GUIDE);
   mUsingDrops = (mds->particle_type & FLUID_DOMAIN_PARTICLE_SPRAY) && mUsingLiquid;
@@ -230,6 +231,10 @@ MANTA::MANTA(int *res, FluidModifierData *mmd) : mCurrentID(++solverID)
       initLiquidMesh();
     }
 
+    if (mUsingDiffusion) {
+      initCurvature();
+    }
+
     if (mUsingGuiding) {
       mResGuiding = (mds->guide_parent) ? mds->guide_res : mds->res;
       initGuiding();
@@ -438,6 +443,16 @@ void MANTA::initLiquidMesh(FluidModifierData *mmd)
   mUsingMesh = true;
 }
 
+void MANTA::initCurvature(FluidModifierData *mmd)
+{
+  std::vector<std::string> pythonCommands;
+  std::string finalString = parseScript(liquid_alloc_curvature, mmd);
+  pythonCommands.push_back(finalString);
+
+  runPythonString(pythonCommands);
+  mUsingDiffusion = true;
+}
+
 void MANTA::initObstacle(FluidModifierData *mmd)
 {
   if (!mPhiObsIn) {
@@ -549,7 +564,7 @@ MANTA::~MANTA()
   result = runPythonString(pythonCommands);
 
   assert(result);
-  (void)result;  // not needed in release
+  UNUSED_VARS(result);
 }
 
 /**
@@ -574,8 +589,8 @@ bool MANTA::runPythonString(vector<string> commands)
     manta_main_module = PyImport_ImportModule("__main__");
   }
 
-  for (vector<std::string>::iterator it = commands.begin(); it != commands.end(); ++it) {
-    std::string command = *it;
+  for (vector<string>::iterator it = commands.begin(); it != commands.end(); ++it) {
+    string command = *it;
 
     PyObject *globals_dict = PyModule_GetDict(manta_main_module);
     PyObject *return_value = PyRun_String(
@@ -843,6 +858,23 @@ void MANTA::initializeRNAMap(FluidModifierData *mmd)
   mRNAMap["GRAVITY_Y"] = to_string(mds->gravity[1]);
   mRNAMap["GRAVITY_Z"] = to_string(mds->gravity[2]);
   mRNAMap["CACHE_DIR"] = cacheDirectory;
+  mRNAMap["NAME_DENSITY"] = FLUID_GRIDNAME_DENSITY;
+  mRNAMap["NAME_SHADOW"] = FLUID_GRIDNAME_SHADOW;
+  mRNAMap["NAME_HEAT"] = FLUID_GRIDNAME_HEAT;
+  mRNAMap["NAME_VELOCITY"] = FLUID_GRIDNAME_VELOCITY;
+  mRNAMap["NAME_COLORR"] = FLUID_GRIDNAME_COLORR;
+  mRNAMap["NAME_COLORG"] = FLUID_GRIDNAME_COLORG;
+  mRNAMap["NAME_COLORB"] = FLUID_GRIDNAME_COLORB;
+  mRNAMap["NAME_FLAME"] = FLUID_GRIDNAME_FLAME;
+  mRNAMap["NAME_FUEL"] = FLUID_GRIDNAME_FUEL;
+  mRNAMap["NAME_REACT"] = FLUID_GRIDNAME_REACT;
+  mRNAMap["NAME_DENSITYNOISE"] = FLUID_GRIDNAME_DENSITYNOISE;
+  mRNAMap["NAME_COLORRNOISE"] = FLUID_GRIDNAME_COLORRNOISE;
+  mRNAMap["NAME_COLORGNOISE"] = FLUID_GRIDNAME_COLORGNOISE;
+  mRNAMap["NAME_COLORBNOISE"] = FLUID_GRIDNAME_COLORBNOISE;
+  mRNAMap["NAME_FLAMENOISE"] = FLUID_GRIDNAME_FLAMENOISE;
+  mRNAMap["NAME_FUELNOISE"] = FLUID_GRIDNAME_FUELNOISE;
+  mRNAMap["NAME_REACTNOISE"] = FLUID_GRIDNAME_REACTNOISE;
 }
 
 string MANTA::getRealValue(const string &varName)
@@ -933,8 +965,7 @@ bool MANTA::updateFlipStructures(FluidModifierData *mmd, int framenr)
   mFlipParticleVelocity->clear();
 
   string pformat = getCacheFileEnding(mds->cache_particle_format);
-  string file = getFile(
-      mmd, FLUID_DOMAIN_DIR_DATA, FLUID_DOMAIN_FILE_PP, pformat.c_str(), framenr);
+  string file = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_FILENAME_PP, pformat, framenr);
 
   expected += 1;
   if (BLI_exists(file.c_str())) {
@@ -942,7 +973,7 @@ bool MANTA::updateFlipStructures(FluidModifierData *mmd, int framenr)
     assert(result == expected);
   }
 
-  file = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_DOMAIN_FILE_PVEL, pformat.c_str(), framenr);
+  file = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_FILENAME_PVEL, pformat, framenr);
   expected += 1;
   if (BLI_exists(file.c_str())) {
     result += updateParticlesFromFile(file, false, true);
@@ -980,7 +1011,7 @@ bool MANTA::updateMeshStructures(FluidModifierData *mmd, int framenr)
 
   string mformat = getCacheFileEnding(mds->cache_mesh_format);
   string dformat = getCacheFileEnding(mds->cache_data_format);
-  string file = getFile(mmd, FLUID_DOMAIN_DIR_MESH, FLUID_DOMAIN_FILE_MESH, mformat, framenr);
+  string file = getFile(mmd, FLUID_DOMAIN_DIR_MESH, FLUID_FILENAME_MESH, mformat, framenr);
 
   expected += 1;
   if (BLI_exists(file.c_str())) {
@@ -989,7 +1020,7 @@ bool MANTA::updateMeshStructures(FluidModifierData *mmd, int framenr)
   }
 
   if (mUsingMVel) {
-    file = getFile(mmd, FLUID_DOMAIN_DIR_MESH, FLUID_DOMAIN_FILE_MESHVEL, dformat, framenr);
+    file = getFile(mmd, FLUID_DOMAIN_DIR_MESH, FLUID_FILENAME_MESHVEL, dformat, framenr);
     expected += 1;
     if (BLI_exists(file.c_str())) {
       result += updateMeshFromFile(file);
@@ -1025,8 +1056,7 @@ bool MANTA::updateParticleStructures(FluidModifierData *mmd, int framenr)
   mSndParticleLife->clear();
 
   string pformat = getCacheFileEnding(mds->cache_particle_format);
-  string file = getFile(
-      mmd, FLUID_DOMAIN_DIR_PARTICLES, FLUID_DOMAIN_FILE_PPSND, pformat, framenr);
+  string file = getFile(mmd, FLUID_DOMAIN_DIR_PARTICLES, FLUID_FILENAME_PPSND, pformat, framenr);
 
   expected += 1;
   if (BLI_exists(file.c_str())) {
@@ -1034,14 +1064,14 @@ bool MANTA::updateParticleStructures(FluidModifierData *mmd, int framenr)
     assert(result == expected);
   }
 
-  file = getFile(mmd, FLUID_DOMAIN_DIR_PARTICLES, FLUID_DOMAIN_FILE_PVELSND, pformat, framenr);
+  file = getFile(mmd, FLUID_DOMAIN_DIR_PARTICLES, FLUID_FILENAME_PVELSND, pformat, framenr);
   expected += 1;
   if (BLI_exists(file.c_str())) {
     result += updateParticlesFromFile(file, true, true);
     assert(result == expected);
   }
 
-  file = getFile(mmd, FLUID_DOMAIN_DIR_PARTICLES, FLUID_DOMAIN_FILE_PLIFESND, pformat, framenr);
+  file = getFile(mmd, FLUID_DOMAIN_DIR_PARTICLES, FLUID_FILENAME_PLIFESND, pformat, framenr);
   expected += 1;
   if (BLI_exists(file.c_str())) {
     result += updateParticlesFromFile(file, true, false);
@@ -1051,6 +1081,26 @@ bool MANTA::updateParticleStructures(FluidModifierData *mmd, int framenr)
   return mParticlesFromFile = (result == expected);
 }
 
+static void assertGridItems(vector<MANTA::GridItem> gList)
+{
+  vector<MANTA::GridItem>::iterator gIter = gList.begin();
+  int *resPrev = (*gIter).res;
+
+  for (vector<MANTA::GridItem>::iterator it = gList.begin(); it != gList.end(); ++it) {
+    MANTA::GridItem item = *it;
+    assert(
+        ELEM(item.type, FLUID_DOMAIN_GRID_FLOAT, FLUID_DOMAIN_GRID_INT, FLUID_DOMAIN_GRID_VEC3F));
+    assert(item.pointer[0]);
+    if (item.type == FLUID_DOMAIN_GRID_VEC3F) {
+      assert(item.pointer[1] && item.pointer[2]);
+    }
+    assert(item.res[0] == resPrev[0] && item.res[1] == resPrev[1] && item.res[2] == resPrev[2]);
+    assert((item.name).compare("") != 0);
+  }
+
+  UNUSED_VARS(resPrev);
+}
+
 bool MANTA::updateSmokeStructures(FluidModifierData *mmd, int framenr)
 {
   if (MANTA::with_debug)
@@ -1065,80 +1115,106 @@ bool MANTA::updateSmokeStructures(FluidModifierData *mmd, int framenr)
     return false;
 
   int result = 0;
-  int expected = 0; /* Expected number of read successes for this frame. */
-
   string dformat = getCacheFileEnding(mds->cache_data_format);
-  string file = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_DOMAIN_FILE_DENSITY, dformat, framenr);
-
-  expected += 1;
-  if (BLI_exists(file.c_str())) {
-    result += updateGridFromFile(file, mDensity, false);
-    assert(result == expected);
-  }
-
-  file = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_DOMAIN_FILE_SHADOW, dformat, framenr);
-  expected += 1;
-  if (BLI_exists(file.c_str())) {
-    result += updateGridFromFile(file, mShadow, false);
-    assert(result == expected);
-  }
 
-  if (mUsingHeat) {
-    file = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_DOMAIN_FILE_HEAT, dformat, framenr);
-    expected += 1;
-    if (BLI_exists(file.c_str())) {
-      result += updateGridFromFile(file, mHeat, false);
-      assert(result == expected);
+  vector<FileItem> filesData;
+  vector<GridItem> gridsData;
+
+  int res[] = {mResX, mResY, mResZ};
+
+  /* Put grid pointers into pointer lists, some grids have more than 1 pointer. */
+  void *aDensity[] = {mDensity};
+  void *aShadow[] = {mShadow};
+  void *aVelocities[] = {mVelocityX, mVelocityY, mVelocityZ};
+  void *aHeat[] = {mHeat};
+  void *aColorR[] = {mColorR};
+  void *aColorG[] = {mColorG};
+  void *aColorB[] = {mColorB};
+  void *aFlame[] = {mFlame};
+  void *aFuel[] = {mFuel};
+  void *aReact[] = {mReact};
+
+  /* File names for grids. */
+  string fDensity = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_FILENAME_DENSITY, dformat, framenr);
+  string fShadow = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_FILENAME_SHADOW, dformat, framenr);
+  string fVel = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_FILENAME_VELOCITY, dformat, framenr);
+  string fHeat = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_FILENAME_HEAT, dformat, framenr);
+  string fColorR = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_FILENAME_COLORR, dformat, framenr);
+  string fColorG = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_FILENAME_COLORG, dformat, framenr);
+  string fColorB = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_FILENAME_COLORB, dformat, framenr);
+  string fFlame = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_FILENAME_FLAME, dformat, framenr);
+  string fFuel = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_FILENAME_FUEL, dformat, framenr);
+  string fReact = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_FILENAME_REACT, dformat, framenr);
+  string fFluid = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_FILENAME_DATA, dformat, framenr);
+
+  /* Prepare grid info containers. */
+  GridItem gDensity = {aDensity, FLUID_DOMAIN_GRID_FLOAT, res, FLUID_GRIDNAME_DENSITY};
+  GridItem gShadow = {aShadow, FLUID_DOMAIN_GRID_FLOAT, res, FLUID_GRIDNAME_SHADOW};
+  GridItem gVel = {aVelocities, FLUID_DOMAIN_GRID_VEC3F, res, FLUID_GRIDNAME_VELOCITY};
+  GridItem gHeat = {aHeat, FLUID_DOMAIN_GRID_FLOAT, res, FLUID_GRIDNAME_HEAT};
+  GridItem gColorR = {aColorR, FLUID_DOMAIN_GRID_FLOAT, res, FLUID_GRIDNAME_COLORR};
+  GridItem gColorG = {aColorG, FLUID_DOMAIN_GRID_FLOAT, res, FLUID_GRIDNAME_COLORG};
+  GridItem gColorB = {aColorB, FLUID_DOMAIN_GRID_FLOAT, res, FLUID_GRIDNAME_COLORB};
+  GridItem gFlame = {aFlame, FLUID_DOMAIN_GRID_FLOAT, res, FLUID_GRIDNAME_FLAME};
+  GridItem gFuel = {aFuel, FLUID_DOMAIN_GRID_FLOAT, res, FLUID_GRIDNAME_FUEL};
+  GridItem gReact = {aReact, FLUID_DOMAIN_GRID_FLOAT, res, FLUID_GRIDNAME_REACT};
+
+  /* TODO (sebbas): For now, only allow single file mode. Combined grid file export is todo. */
+  const int fileMode = FLUID_DOMAIN_CACHE_FILES_SINGLE;
+  if (fileMode == FLUID_DOMAIN_CACHE_FILES_SINGLE) {
+
+    filesData.push_back({fDensity, {gDensity}});
+    filesData.push_back({fShadow, {gShadow}});
+    filesData.push_back({fVel, {gVel}});
+    if (mUsingHeat) {
+      filesData.push_back({fHeat, {gHeat}});
     }
-  }
-
-  if (mUsingColors) {
-    file = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_DOMAIN_FILE_COLORR, dformat, framenr);
-    expected += 1;
-    if (BLI_exists(file.c_str())) {
-      result += updateGridFromFile(file, mColorR, false);
-      assert(result == expected);
-    }
-
-    file = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_DOMAIN_FILE_COLORG, dformat, framenr);
-    expected += 1;
-    if (BLI_exists(file.c_str())) {
-      result += updateGridFromFile(file, mColorG, false);
-      assert(result == expected);
+    if (mUsingColors) {
+      filesData.push_back({fColorR, {gColorR}});
+      filesData.push_back({fColorG, {gColorG}});
+      filesData.push_back({fColorB, {gColorB}});
     }
-
-    file = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_DOMAIN_FILE_COLORB, dformat, framenr);
-    expected += 1;
-    if (BLI_exists(file.c_str())) {
-      result += updateGridFromFile(file, mColorB, false);
-      assert(result == expected);
+    if (mUsingFire) {
+      filesData.push_back({fFlame, {gFlame}});
+      filesData.push_back({fFuel, {gFuel}});
+      filesData.push_back({fReact, {gReact}});
     }
   }
+  else if (fileMode == FLUID_DOMAIN_CACHE_FILES_COMBINED) {
 
-  if (mUsingFire) {
-    file = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_DOMAIN_FILE_FLAME, dformat, framenr);
-    expected += 1;
-    if (BLI_exists(file.c_str())) {
-      result += updateGridFromFile(file, mFlame, false);
-      assert(result == expected);
+    gridsData.push_back(gDensity);
+    gridsData.push_back(gShadow);
+    gridsData.push_back(gVel);
+    if (mUsingHeat) {
+      gridsData.push_back(gHeat);
+    }
+    if (mUsingColors) {
+      gridsData.push_back(gColorR);
+      gridsData.push_back(gColorG);
+      gridsData.push_back(gColorB);
+    }
+    if (mUsingFire) {
+      gridsData.push_back(gFlame);
+      gridsData.push_back(gFuel);
+      gridsData.push_back(gReact);
     }
 
-    file = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_DOMAIN_FILE_FUEL, dformat, framenr);
-    expected += 1;
-    if (BLI_exists(file.c_str())) {
-      result += updateGridFromFile(file, mFuel, false);
-      assert(result == expected);
+    if (with_debug) {
+      assertGridItems(gridsData);
     }
+    filesData.push_back({fFluid, gridsData});
+  }
 
-    file = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_DOMAIN_FILE_REACT, dformat, framenr);
-    expected += 1;
-    if (BLI_exists(file.c_str())) {
-      result += updateGridFromFile(file, mReact, false);
-      assert(result == expected);
+  /* Update files from data directory. */
+  for (vector<FileItem>::iterator it = filesData.begin(); it != filesData.end(); ++it) {
+    FileItem item = *it;
+    if (BLI_exists(item.filename.c_str())) {
+      result += updateGridsFromFile(item.filename, item.grids);
+      assert(result);
     }
   }
 
-  return mSmokeFromFile = (result == expected);
+  return mSmokeFromFile = result;
 }
 
 bool MANTA::updateNoiseStructures(FluidModifierData *mmd, int framenr)
@@ -1155,73 +1231,121 @@ bool MANTA::updateNoiseStructures(FluidModifierData *mmd, int framenr)
     return false;
 
   int result = 0;
-  int expected = 0; /* Expected number of read successes for this frame. */
-
   string dformat = getCacheFileEnding(mds->cache_data_format);
   string nformat = getCacheFileEnding(mds->cache_noise_format);
-  string file = getFile(
-      mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_DOMAIN_FILE_DENSITYNOISE, nformat, framenr);
 
-  expected += 1;
-  if (BLI_exists(file.c_str())) {
-    result += updateGridFromFile(file, mDensityHigh, true);
-    assert(result == expected);
+  vector<FileItem> filesData, filesNoise;
+  vector<GridItem> gridsData, gridsNoise;
+
+  int resData[] = {mResX, mResY, mResZ};
+  int resNoise[] = {mResXNoise, mResYNoise, mResZNoise};
+
+  /* Put grid pointers into pointer lists, some grids have more than 1 pointer. */
+  void *aShadow[] = {mShadow};
+  void *aVelocities[] = {mVelocityX, mVelocityY, mVelocityZ};
+  void *aDensity[] = {mDensityHigh};
+  void *aColorR[] = {mColorRHigh};
+  void *aColorG[] = {mColorGHigh};
+  void *aColorB[] = {mColorBHigh};
+  void *aFlame[] = {mFlameHigh};
+  void *aFuel[] = {mFuelHigh};
+  void *aReact[] = {mReactHigh};
+
+  /* File names for grids. */
+  string fShadow = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_FILENAME_SHADOW, dformat, framenr);
+  string fVel = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_FILENAME_VELOCITY, dformat, framenr);
+  string fFluid = getFile(mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_FILENAME_DATA, dformat, framenr);
+
+  string fDensity = getFile(
+      mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_FILENAME_DENSITYNOISE, nformat, framenr);
+  string fColorR = getFile(
+      mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_FILENAME_COLORRNOISE, nformat, framenr);
+  string fColorG = getFile(
+      mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_FILENAME_COLORGNOISE, nformat, framenr);
+  string fColorB = getFile(
+      mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_FILENAME_COLORBNOISE, nformat, framenr);
+  string fFlame = getFile(
+      mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_FILENAME_FLAMENOISE, nformat, framenr);
+  string fFuel = getFile(mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_FILENAME_FUELNOISE, nformat, framenr);
+  string fReact = getFile(
+      mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_FILENAME_REACTNOISE, nformat, framenr);
+  string fNoise = getFile(mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_FILENAME_NOISE, nformat, framenr);
+
+  /* Prepare grid info containers. */
+  GridItem gShadow = {aShadow, FLUID_DOMAIN_GRID_FLOAT, resData, FLUID_GRIDNAME_SHADOW};
+  GridItem gVel = {aVelocities, FLUID_DOMAIN_GRID_VEC3F, resData, FLUID_GRIDNAME_VELOCITY};
+
+  GridItem gDensity = {aDensity, FLUID_DOMAIN_GRID_FLOAT, resNoise, FLUID_GRIDNAME_DENSITYNOISE};
+  GridItem gColorR = {aColorR, FLUID_DOMAIN_GRID_FLOAT, resNoise, FLUID_GRIDNAME_COLORRNOISE};
+  GridItem gColorG = {aColorG, FLUID_DOMAIN_GRID_FLOAT, resNoise, FLUID_GRIDNAME_COLORGNOISE};
+  GridItem gColorB = {aColorB, FLUID_DOMAIN_GRID_FLOAT, resNoise, FLUID_GRIDNAME_COLORBNOISE};
+  GridItem gFlame = {aFlame, FLUID_DOMAIN_GRID_FLOAT, resNoise, FLUID_GRIDNAME_FLAMENOISE};
+  GridItem gFuel = {aFuel, FLUID_DOMAIN_GRID_FLOAT, resNoise, FLUID_GRIDNAME_FUELNOISE};
+  GridItem gReact = {aReact, FLUID_DOMAIN_GRID_FLOAT, resNoise, FLUID_GRIDNAME_REACTNOISE};
+
+  /* TODO (sebbas): For now, only allow single file mode. Combined grid file export is todo. */
+  const int fileMode = FLUID_DOMAIN_CACHE_FILES_SINGLE;
+  if (fileMode == FLUID_DOMAIN_CACHE_FILES_SINGLE) {
+
+    filesData.push_back({fShadow, {gShadow}});
+    filesData.push_back({fVel, {gVel}});
+
+    filesNoise.push_back({fDensity, {gDensity}});
+    if (mUsingColors) {
+      filesNoise.push_back({fColorR, {gColorR}});
+      filesNoise.push_back({fColorG, {gColorG}});
+      filesNoise.push_back({fColorB, {gColorB}});
+    }
+    if (mUsingFire) {
+      filesNoise.push_back({fFlame, {gFlame}});
+      filesNoise.push_back({fFuel, {gFuel}});
+      filesNoise.push_back({fReact, {gReact}});
+    }
   }
+  else if (fileMode == FLUID_DOMAIN_CACHE_FILES_COMBINED) {
 
-  file = getFile(mmd, FLUID_DOMAIN_DIR_DATA, FLUID_DOMAIN_FILE_SHADOW, dformat, framenr);
-  expected += 1;
-  if (BLI_exists(file.c_str())) {
-    result += updateGridFromFile(file, mShadow, false);
-    assert(result == expected);
-  }
+    gridsData.push_back(gShadow);
+    gridsData.push_back(gVel);
 
-  if (mUsingColors) {
-    file = getFile(mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_DOMAIN_FILE_COLORRNOISE, nformat, framenr);
-    expected += 1;
-    if (BLI_exists(file.c_str())) {
-      result += updateGridFromFile(file, mColorRHigh, true);
-      assert(result == expected);
+    gridsNoise.push_back(gDensity);
+    if (mUsingColors) {
+      gridsNoise.push_back(gColorR);
+      gridsNoise.push_back(gColorG);
+      gridsNoise.push_back(gColorB);
     }
-
-    file = getFile(mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_DOMAIN_FILE_COLORGNOISE, nformat, framenr);
-    expected += 1;
-    if (BLI_exists(file.c_str())) {
-      result += updateGridFromFile(file, mColorGHigh, true);
-      assert(result == expected);
+    if (mUsingFire) {
+      gridsNoise.push_back(gFlame);
+      gridsNoise.push_back(gFuel);
+      gridsNoise.push_back(gReact);
     }
 
-    file = getFile(mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_DOMAIN_FILE_COLORBNOISE, nformat, framenr);
-    expected += 1;
-    if (BLI_exists(file.c_str())) {
-      result += updateGridFromFile(file, mColorBHigh, true);
-      assert(result == expected);
+    if (with_debug) {
+      assertGridItems(gridsData);
+      assertGridItems(gridsNoise);
     }
+    filesData.push_back({fFluid, gridsData});
+    filesNoise.push_back({fNoise, gridsNoise});
   }
 
-  if (mUsingFire) {
-    file = getFile(mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_DOMAIN_FILE_FLAMENOISE, nformat, framenr);
-    expected += 1;
-    if (BLI_exists(file.c_str())) {
-      result += updateGridFromFile(file, mFlameHigh, true);
-      assert(result == expected);
-    }
-
-    file = getFile(mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_DOMAIN_FILE_FUELNOISE, nformat, framenr);
-    expected += 1;
-    if (BLI_exists(file.c_str())) {
-      result += updateGridFromFile(file, mFuelHigh, true);
-      assert(result == expected);
+  /* Update files from data directory. */
+  for (vector<FileItem>::iterator it = filesData.begin(); it != filesData.end(); ++it) {
+    FileItem item = *it;
+    if (BLI_exists(item.filename.c_str())) {
+      result += updateGridsFromFile(item.filename, item.grids);
+      assert(result);
     }
+  }
 
-    file = getFile(mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_DOMAIN_FILE_REACTNOISE, nformat, framenr);
-    expected += 1;
-    if (BLI_exists(file.c_str())) {
-      result += updateGridFromFile(file, mReactHigh, true);
-      assert(result == expected);
+  /* Update files from noise directory. */
+  for (vector<FileItem>::iterator it = filesNoise.begin(); it != filesNoise.end(); ++it) {
+    FileItem item = *it;
+    if (BLI_exists(item.filename.c_str())) {
+      result += updateGridsFromFile(item.filename, item.grids);
+      assert(result);
     }
   }
 
-  return mNoiseFromFile = (result == expected);
+  return mNoiseFromFile = result;
 }
 
 /* Dirty hack: Needed to format paths from python code that is run via PyRun_SimpleString */
@@ -1247,7 +1371,7 @@ bool MANTA::writeConfiguration(FluidModifierData *mmd, int framenr)
 
   string directory = getDirectory(mmd, FLUID_DOMAIN_DIR_CONFIG);
   string format = FLUID_DOMAIN_EXTENSION_UNI;
-  string file = getFile(mmd, FLUID_DOMAIN_DIR_CONFIG, FLUID_DOMAIN_FILE_CONFIG, format, framenr);
+  string file = getFile(mmd, FLUID_DOMAIN_DIR_CONFIG, FLUID_FILENAME_CONFIG, format, framenr);
 
   /* Create 'config' subdir if it does not exist already. */
   BLI_dir_create_recursive(directory.c_str());
@@ -1347,7 +1471,7 @@ bool MANTA::readConfiguration(FluidModifierData *mmd, int framenr)
 
   string directory = getDirectory(mmd, FLUID_DOMAIN_DIR_CONFIG);
   string format = FLUID_DOMAIN_EXTENSION_UNI;
-  string file = getFile(mmd, FLUID_DOMAIN_DIR_CONFIG, FLUID_DOMAIN_FILE_CONFIG, format, framenr);
+  string file = getFile(mmd, FLUID_DOMAIN_DIR_CONFIG, FLUID_FILENAME_CONFIG, format, framenr);
 
   if (!hasConfig(mmd, framenr))
     return false;
@@ -2741,14 +2865,13 @@ bool MANTA::updateParticlesFromUni(string filename, bool isSecondarySys, bool is
   return (gzclose(gzf) == Z_OK);
 }
 
-bool MANTA::updateGridFromFile(string filename, float *grid, bool isNoise)
+bool MANTA::updateGridsFromFile(string filename, vector<GridItem> grids)
 {
   if (with_debug)
-    cout << "MANTA::updateGridFromFile()" << endl;
+    cout << "MANTA::updateGridsFromFile()" << endl;
 
-  if (!grid) {
-    cerr << "Fluid Error -- updateGridFromFile(): Cannot read into uninitialized grid (grid "
-            "is null)."
+  if (grids.empty()) {
+    cerr << "Fluid Error -- updateGridsFromFile(): Cannot read into uninitialized grid vector."
          << endl;
     return false;
   }
@@ -2758,118 +2881,142 @@ bool MANTA::updateGridFromFile(string filename, float *grid, bool isNoise)
 
   idx = fname.rfind('.');
   if (idx != string::npos) {
-    string extension = fname.substr(idx + 1);
+    string extension = fname.substr(idx);
 
-    if (extension.compare("uni") == 0)
-      return updateGridFromUni(filename, grid, isNoise);
+    if (extension.compare(FLUID_DOMAIN_EXTENSION_UNI) == 0) {
+      return updateGridsFromUni(filename, grids);
+    }
 #if OPENVDB == 1
-    else if (extension.compare("vdb") == 0)
-      return updateGridFromVDB(filename, grid, isNoise);
+    else if (extension.compare(FLUID_DOMAIN_EXTENSION_OPENVDB) == 0) {
+      return updateGridsFromVDB(filename, grids);
+    }
 #endif
-    else if (extension.compare("raw") == 0)
-      return updateGridFromRaw(filename, grid, isNoise);
-    else
-      cerr << "Fluid Error -- updateGridFromFile(): Invalid file extension in file: " << filename
+    else if (extension.compare(FLUID_DOMAIN_EXTENSION_RAW) == 0) {
+      return updateGridsFromRaw(filename, grids);
+    }
+    else {
+      cerr << "Fluid Error -- updateGridsFromFile(): Invalid file extension in file: " << filename
            << endl;
+    }
     return false;
   }
   else {
-    cerr << "Fluid Error -- updateGridFromFile(): Unable to open file: " << filename << endl;
+    cerr << "Fluid Error -- updateGridsFromFile(): Unable to open file: " << filename << endl;
     return false;
   }
 }
 
-bool MANTA::updateGridFromUni(string filename, float *grid, bool isNoise)
+bool MANTA::updateGridsFromUni(string filename, vector<GridItem> grids)
 {
   if (with_debug)
-    cout << "MANTA::updateGridFromUni()" << endl;
+    cout << "MANTA::updateGridsFromUni()" << endl;
 
   gzFile gzf;
+  int expectedBytes = 0, readBytes = 0;
   int ibuffer[4];
 
   gzf = (gzFile)BLI_gzopen(filename.c_str(), "rb1");
   if (!gzf) {
-    cerr << "Fluid Error -- updateGridFromUni(): Unable to open file: " << filename << endl;
+    cerr << "Fluid Error -- updateGridsFromUni(): Unable to open file: " << filename << endl;
     return false;
   }
 
-  int readBytes = 0;
   char file_magic[5] = {0, 0, 0, 0, 0};
   readBytes = gzread(gzf, file_magic, 4);
   if (!readBytes) {
-    cerr << "Fluid Error -- updateGridFromUni(): Unable to read header in file: " << filename
-         << endl;
-    gzclose(gzf);
-    return false;
-  }
-
-  if (!strcmp(file_magic, "DDF2")) {
-    cerr << "Fluid Error -- updateGridFromUni(): Grid uni file format DDF2 not supported anymore."
-         << endl;
-    gzclose(gzf);
-    return false;
-  }
-
-  if (!strcmp(file_magic, "MNT1")) {
-    cerr << "Fluid Error -- updateGridFromUni(): Grid uni file format MNT1 not supported anymore."
-         << endl;
+    cerr << "Fluid Error -- updateGridsFromUni(): Invalid header in file: " << filename << endl;
     gzclose(gzf);
     return false;
   }
-
-  if (!strcmp(file_magic, "MNT2")) {
-    cerr << "Fluid Error -- updateGridFromUni(): Grid uni file format MNT2 not supported anymore."
+  if (!strcmp(file_magic, "DDF2") || !strcmp(file_magic, "MNT1") || !strcmp(file_magic, "MNT2")) {
+    cerr << "Fluid Error -- updateGridsFromUni(): Unsupported header in file: " << filename
          << endl;
     gzclose(gzf);
     return false;
   }
 
-  // grid uni header
-  const int STR_LEN_GRID = 252;
-  int elementType, bytesPerElement;  // data type info
-  char info[STR_LEN_GRID];           // mantaflow build information
-  int dimT;                          // optionally store forth dimension for 4d grids
-  unsigned long long timestamp;      // creation time
+  if (!strcmp(file_magic, "MNT3")) {
 
-  // read grid header
-  gzread(gzf, &ibuffer, sizeof(int) * 4);  // dimX, dimY, dimZ, gridType
-  gzread(gzf, &elementType, sizeof(int));
-  gzread(gzf, &bytesPerElement, sizeof(int));
-  gzread(gzf, &info, sizeof(info));
-  gzread(gzf, &dimT, sizeof(int));
-  gzread(gzf, &timestamp, sizeof(unsigned long long));
+    // grid uni header
+    const int STR_LEN_GRID = 252;
+    int elementType, bytesPerElement;  // data type info
+    char info[STR_LEN_GRID];           // mantaflow build information
+    int dimT;                          // optionally store forth dimension for 4d grids
+    unsigned long long timestamp;      // creation time
+
+    // read grid header
+    gzread(gzf, &ibuffer, sizeof(int) * 4);  // dimX, dimY, dimZ, gridType
+    gzread(gzf, &elementType, sizeof(int));
+    gzread(gzf, &bytesPerElement, sizeof(int));
+    gzread(gzf, &info, sizeof(info));
+    gzread(gzf, &dimT, sizeof(int));
+    gzread(gzf, &timestamp, sizeof(unsigned long long));
 
-  int resX = (isNoise) ? mResXNoise : mResX;
-  int resY = (isNoise) ? mResYNoise : mResY;
-  int resZ = (isNoise) ? mResZNoise : mResZ;
+    if (with_debug)
+      cout << "Fluid: Read " << ibuffer[3] << " grid type in file: " << filename << endl;
+
+    for (vector<GridItem>::iterator gIter = grids.begin(); gIter != grids.end(); ++gIter) {
+      GridItem gridItem = *gIter;
+      void **pointerList = gridItem.pointer;
+      int type = gridItem.type;
+      int *res = gridItem.res;
+      assert(pointerList[0]);
+      assert(res[0] == res[0] && res[1] == res[1] && res[2] == res[2]);
+      UNUSED_VARS(res);
+
+      switch (type) {
+        case FLUID_DOMAIN_GRID_VEC3F: {
+          assert(pointerList[1] && pointerList[2]);
+          float **fpointers = (float **)pointerList;
+          expectedBytes = sizeof(float) * 3 * ibuffer[0] * ibuffer[1] * ibuffer[2];
+          readBytes = 0;
+          for (int i = 0; i < ibuffer[0] * ibuffer[1] * ibuffer[2]; ++i) {
+            for (int j = 0; j < 3; ++j) {
+              readBytes += gzread(gzf, fpointers[j], sizeof(float));
+              ++fpointers[j];
+            }
+          }
+          break;
+        }
+        case FLUID_DOMAIN_GRID_FLOAT: {
+          float **fpointers = (float **)pointerList;
+          expectedBytes = sizeof(float) * ibuffer[0] * ibuffer[1] * ibuffer[2];
+          readBytes = gzread(
+              gzf, fpointers[0], sizeof(float) * ibuffer[0] * ibuffer[1] * ibuffer[2]);
+          break;
+        }
+        default: {
+          cerr << "Fluid Error -- Unknown grid type" << endl;
+        }
+      }
 
-  if (with_debug)
-    cout << "Fluid: Read " << ibuffer[3] << " grid type in file: " << filename << endl;
+      if (!readBytes) {
+        cerr << "Fluid Error -- updateGridFromRaw(): Unable to read raw file: " << filename
+             << endl;
+        gzclose(gzf);
+        return false;
+      }
+      assert(expectedBytes == readBytes);
+      UNUSED_VARS(expectedBytes);
 
-  // Sanity checks
-  if (ibuffer[0] != resX || ibuffer[1] != resY || ibuffer[2] != resZ) {
-    cout << "Fluid: Grid dim doesn't match, read: (" << ibuffer[0] << ", " << ibuffer[1] << ", "
-         << ibuffer[2] << ") vs setup: (" << resX << ", " << resY << ", " << resZ << ")" << endl;
+      if (with_debug)
+        cout << "Fluid: Read successfully: " << filename << endl;
+    }
+  }
+  else {
+    cerr << "Fluid Error -- updateGridsFromUni(): Unknown header in file: " << filename << endl;
     gzclose(gzf);
     return false;
   }
 
-  // Actual data reading
-  if (!strcmp(file_magic, "MNT3")) {
-    gzread(gzf, grid, sizeof(float) * ibuffer[0] * ibuffer[1] * ibuffer[2]);
-  }
-
-  if (with_debug)
-    cout << "Fluid: Read successfully: " << filename << endl;
-
   return (gzclose(gzf) == Z_OK);
 }
 
 #if OPENVDB == 1
-bool MANTA::updateGridFromVDB(string filename, float *grid, bool isNoise)
+bool MANTA::updateGridsFromVDB(string filename, vector<GridItem> grids)
 {
   if (with_debug)
-    cout << "MANTA::updateGridFromVDB()" << endl;
+    cout << "MANTA::updateGridsFromVDB()" << endl;
 
   openvdb::initialize();
   openvdb::io::File file(filename);
@@ -2877,66 +3024,191 @@ bool MANTA::updateGridFromVDB(string filename, float *grid, bool isNoise)
     file.open();
   }
   catch (const openvdb::IoError &) {
-    cerr << "Fluid Error -- updateGridFromVDB(): IOError, invalid OpenVDB file: " << filename
+    cerr << "Fluid Error -- updateGridsFromVDB(): IOError, invalid OpenVDB file: " << filename
          << endl;
     return false;
   }
+  if (grids.empty()) {
+    cerr << "Fluid Error -- updateGridsFromVDB(): No grids found in grid vector" << endl;
+    return false;
+  }
 
+  unordered_map<string, openvdb::FloatGrid::Accessor> floatAccessors;
+  unordered_map<string, openvdb::Vec3SGrid::Accessor> vec3fAccessors;
   openvdb::GridBase::Ptr baseGrid;
-  for (openvdb::io::File::NameIterator nameIter = file.beginName(); nameIter != file.endName();
-       ++nameIter) {
-    baseGrid = file.readGrid(nameIter.gridName());
-    break;
+
+  /* Get accessors to all grids in this OpenVDB file.*/
+  for (vector<GridItem>::iterator gIter = grids.begin(); gIter != grids.end(); ++gIter) {
+    GridItem gridItem = *gIter;
+    string itemName = gridItem.name;
+    int itemType = gridItem.type;
+
+    for (openvdb::io::File::NameIterator nameIter = file.beginName(); nameIter != file.endName();
+         ++nameIter) {
+      string vdbName = nameIter.gridName();
+      bool nameMatch = !itemName.compare(vdbName);
+
+      /* Support for <= 2.83: If file has only one grid in it, use that grid. */
+      openvdb::io::File::NameIterator peekNext = nameIter;
+      bool onlyGrid = (++peekNext == file.endName());
+      if (onlyGrid) {
+        vdbName = itemName;
+      }
+
+      if (nameMatch || onlyGrid) {
+        baseGrid = file.readGrid(nameIter.gridName());
+
+        switch (itemType) {
+          case FLUID_DOMAIN_GRID_VEC3F: {
+            openvdb::Vec3SGrid::Ptr gridVDB = openvdb::gridPtrCast<openvdb::Vec3SGrid>(baseGrid);
+            openvdb::Vec3SGrid::Accessor vdbAccessor = gridVDB->getAccessor();
+            vec3fAccessors.emplace(vdbName, vdbAccessor);
+            break;
+          }
+          case FLUID_DOMAIN_GRID_FLOAT: {
+            openvdb::FloatGrid::Ptr gridVDB = openvdb::gridPtrCast<openvdb::FloatGrid>(baseGrid);
+            openvdb::FloatGrid::Accessor vdbAccessor = gridVDB->getAccessor();
+            floatAccessors.emplace(vdbName, vdbAccessor);
+            break;
+          }
+          default: {
+            cerr << "Fluid Error -- Unknown grid type" << endl;
+          }
+        }
+      }
+      else {
+        cerr << "Fluid Error -- Could not read grid from file" << endl;
+        return false;
+      }
+    }
   }
   file.close();
-  openvdb::FloatGrid::Ptr gridVDB = openvdb::gridPtrCast<openvdb::FloatGrid>(baseGrid);
-  openvdb::FloatGrid::Accessor accessor = gridVDB->getAccessor();
-
-  int resX = (isNoise) ? mResXNoise : mResX;
-  int resY = (isNoise) ? mResYNoise : mResY;
-  int resZ = (isNoise) ? mResZNoise : mResZ;
 
   size_t index = 0;
-  for (int z = 0; z < resZ; ++z) {
-    for (int y = 0; y < resY; ++y) {
-      for (int x = 0; x < resX; ++x, ++index) {
+
+  /* Use res of first grid for grid loop. All grids must be same size anyways. */
+  vector<GridItem>::iterator gIter = grids.begin();
+  int *res = (*gIter).res;
+
+  for (int z = 0; z < res[2]; ++z) {
+    for (int y = 0; y < res[1]; ++y) {
+      for (int x = 0; x < res[0]; ++x, ++index) {
         openvdb::Coord xyz(x, y, z);
-        float v = accessor.getValue(xyz);
-        grid[index] = v;
+
+        for (vector<GridItem>::iterator gIter = grids.begin(); gIter != grids.end(); ++gIter) {
+          GridItem gridItem = *gIter;
+          void **pointerList = gridItem.pointer;
+          int type = gridItem.type;
+          int *res = gridItem.res;
+          assert(pointerList[0]);
+          assert(res[0] == res[0] && res[1] == res[1] && res[2] == res[2]);
+          UNUSED_VARS(res);
+
+          switch (type) {
+            case FLUID_DOMAIN_GRID_VEC3F: {
+              unordered_map<string, openvdb::Vec3SGrid::Accessor>::iterator it;
+              it = vec3fAccessors.find(gridItem.name);
+              if (it == vec3fAccessors.end()) {
+                cerr << "Fluid Error -- '" << gridItem.name << "' not in vdb grid map" << endl;
+                return false;
+              }
+              openvdb::Vec3f v = it->second.getValue(xyz);
+
+              assert(pointerList[1] && pointerList[2]);
+              float **fpointers = (float **)pointerList;
+              for (int j = 0; j < 3; ++j) {
+                (fpointers[j])[index] = (float)v[j];
+              }
+              break;
+            }
+            case FLUID_DOMAIN_GRID_FLOAT: {
+              unordered_map<string, openvdb::FloatGrid::Accessor>::iterator it;
+              it = floatAccessors.find(gridItem.name);
+              if (it == floatAccessors.end()) {
+                cerr << "Fluid Error -- '" << gridItem.name << "' not in vdb grid map" << endl;
+                return false;
+              }
+              float v = it->second.getValue(xyz);
+              float **fpointers = (float **)pointerList;
+              (fpointers[0])[index] = v;
+              break;
+            }
+            default: {
+              cerr << "Fluid Error -- Unknown grid type" << endl;
+            }
+          }
+        }
       }
     }
   }
+  if (with_debug)
+    cout << "Fluid: Read successfully: " << filename << endl;
+
   return true;
 }
 #endif
 
-bool MANTA::updateGridFromRaw(string filename, float *grid, bool isNoise)
+bool MANTA::updateGridsFromRaw(string filename, vector<GridItem> grids)
 {
   if (with_debug)
-    cout << "MANTA::updateGridFromRaw()" << endl;
+    cout << "MANTA::updateGridsFromRaw()" << endl;
 
   gzFile gzf;
   int expectedBytes, readBytes;
 
   gzf = (gzFile)BLI_gzopen(filename.c_str(), "rb");
   if (!gzf) {
-    cout << "MANTA::updateGridFromRaw(): unable to open file" << endl;
+    cout << "MANTA::updateGridsFromRaw(): unable to open file" << endl;
     return false;
   }
 
-  int resX = (isNoise) ? mResXNoise : mResX;
-  int resY = (isNoise) ? mResYNoise : mResY;
-  int resZ = (isNoise) ? mResZNoise : mResZ;
+  for (vector<GridItem>::iterator gIter = grids.begin(); gIter != grids.end(); ++gIter) {
+    GridItem gridItem = *gIter;
+    void **pointerList = gridItem.pointer;
+    int type = gridItem.type;
+    int *res = gridItem.res;
+    assert(pointerList[0]);
+    assert(res[0] == res[0] && res[1] == res[1] && res[2] == res[2]);
+    UNUSED_VARS(res);
+
+    switch (type) {
+      case FLUID_DOMAIN_GRID_VEC3F: {
+        assert(pointerList[1] && pointerList[2]);
+        float **fpointers = (float **)pointerList;
+        expectedBytes = sizeof(float) * 3 * res[0] * res[1] * res[2];
+        readBytes = 0;
+        for (int i = 0; i < res[0] * res[1] * res[2]; ++i) {
+          for (int j = 0; j < 3; ++j) {
+            readBytes += gzread(gzf, fpointers[j], sizeof(float));
+            ++fpointers[j];
+          }
+        }
+        break;
+      }
+      case FLUID_DOMAIN_GRID_FLOAT: {
+        float **fpointers = (float **)pointerList;
+        expectedBytes = sizeof(float) * res[0] * res[1] * res[2];
+        readBytes = gzread(gzf, fpointers[0], expectedBytes);
+        break;
+      }
+      default: {
+        cerr << "Fluid Error -- Unknown grid type" << endl;
+      }
+    }
+
+    if (!readBytes) {
+      cerr << "Fluid Error -- updateGridsFromRaw(): Unable to read raw file: " << filename << endl;
+      gzclose(gzf);
+      return false;
+    }
+    assert(expectedBytes == readBytes);
 
-  expectedBytes = sizeof(float) * resX * resY * resZ;
-  readBytes = gzread(gzf, grid, expectedBytes);
-  if (!readBytes) {
-    cerr << "Fluid Error -- updateGridFromRaw(): Unable to read raw file: " << filename << endl;
-    gzclose(gzf);
-    return false;
+    if (with_debug)
+      cout << "Fluid: Read successfully: " << filename << endl;
   }
 
-  assert(expectedBytes == readBytes);
+  if (with_debug)
+    cout << "Fluid: Read successfully: " << filename << endl;
 
   return (gzclose(gzf) == Z_OK);
 }
@@ -3084,12 +3356,12 @@ bool MANTA::hasConfig(FluidModifierData *mmd, int framenr)
 {
   string extension = FLUID_DOMAIN_EXTENSION_UNI;
   return BLI_exists(
-      getFile(mmd, FLUID_DOMAIN_DIR_CONFIG, FLUID_DOMAIN_FILE_CONFIG, extension, framenr).c_str());
+      getFile(mmd, FLUID_DOMAIN_DIR_CONFIG, FLUID_FILENAME_CONFIG, extension, framenr).c_str());
 }
 
 bool MANTA::hasData(FluidModifierData *mmd, int framenr)
 {
-  string filename = (mUsingSmoke) ? FLUID_DOMAIN_FILE_DENSITY : FLUID_DOMAIN_FILE_PP;
+  string filename = (mUsingSmoke) ? FLUID_FILENAME_DENSITY : FLUID_FILENAME_PP;
   string extension = getCacheFileEnding(mmd->domain->cache_data_format);
   return BLI_exists(getFile(mmd, FLUID_DOMAIN_DIR_DATA, filename, extension, framenr).c_str());
 }
@@ -3098,7 +3370,7 @@ bool MANTA::hasNoise(FluidModifierData *mmd, int framenr)
 {
   string extension = getCacheFileEnding(mmd->domain->cache_noise_format);
   return BLI_exists(
-      getFile(mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_DOMAIN_FILE_DENSITYNOISE, extension, framenr)
+      getFile(mmd, FLUID_DOMAIN_DIR_NOISE, FLUID_FILENAME_DENSITYNOISE, extension, framenr)
           .c_str());
 }
 
@@ -3106,21 +3378,20 @@ bool MANTA::hasMesh(FluidModifierData *mmd, int framenr)
 {
   string extension = getCacheFileEnding(mmd->domain->cache_mesh_format);
   return BLI_exists(
-      getFile(mmd, FLUID_DOMAIN_DIR_MESH, FLUID_DOMAIN_FILE_MESH, extension, framenr).c_str());
+      getFile(mmd, FLUID_DOMAIN_DIR_MESH, FLUID_FILENAME_MESH, extension, framenr).c_str());
 }
 
 bool MANTA::hasParticles(FluidModifierData *mmd, int framenr)
 {
   string extension = getCacheFileEnding(mmd->domain->cache_particle_format);
   return BLI_exists(
-      getFile(mmd, FLUID_DOMAIN_DIR_PARTICLES, FLUID_DOMAIN_FILE_PPSND, extension, framenr)
-          .c_str());
+      getFile(mmd, FLUID_DOMAIN_DIR_PARTICLES, FLUID_FILENAME_PPSND, extension, framenr).c_str());
 }
 
 bool MANTA::hasGuiding(FluidModifierData *mmd, int framenr, bool sourceDomain)
 {
   string subdirectory = (sourceDomain) ? FLUID_DOMAIN_DIR_DATA : FLUID_DOMAIN_DIR_GUIDE;
-  string filename = (sourceDomain) ? FLUID_DOMAIN_FILE_VEL : FLUID_DOMAIN_FILE_GUIDEVEL;
+  string filename = (sourceDomain) ? FLUID_FILENAME_VELOCITY : FLUID_FILENAME_GUIDEVEL;
   string extension = getCacheFileEnding(mmd->domain->cache_data_format);
   return BLI_exists(getFile(mmd, subdirectory, filename, extension, framenr).c_str());
 }
diff --git a/intern/mantaflow/intern/MANTA_main.h b/intern/mantaflow/intern/MANTA_main.h
index 9b3fd6aa141..6a8484c75d9 100644
--- a/intern/mantaflow/intern/MANTA_main.h
+++ b/intern/mantaflow/intern/MANTA_main.h
@@ -60,6 +60,19 @@ struct MANTA {
     int flags;
   } Triangle;
 
+  // Cache helper typedefs
+  typedef struct GridItem {
+    void **pointer; /* Array of pointers for this grid.*/
+    int type;
+    int *res;
+    string name;
+  } GridItem;
+
+  typedef struct FileItem {
+    string filename;
+    vector<GridItem> grids;
+  } FileItem;
+
   // Manta step, handling everything
   void step(struct FluidModifierData *mmd, int startFrame);
 
@@ -72,6 +85,7 @@ struct MANTA {
   void initLiquid(FluidModifierData *mmd = NULL);
   void initLiquidMesh(FluidModifierData *mmd = NULL);
   void initObstacle(FluidModifierData *mmd = NULL);
+  void initCurvature(FluidModifierData *mmd = NULL);
   void initGuiding(FluidModifierData *mmd = NULL);
   void initFractions(FluidModifierData *mmd = NULL);
   void initInVelocity(FluidModifierData *mmd = NULL);
@@ -761,6 +775,7 @@ struct MANTA {
   bool mUsingOutflow;
   bool mUsingNoise;
   bool mUsingMesh;
+  bool mUsingDiffusion;
   bool mUsingMVel;
   bool mUsingLiquid;
   bool mUsingSmoke;
@@ -888,12 +903,12 @@ struct MANTA {
   bool updateMeshFromObj(string filename);
   bool updateMeshFromUni(string filename);
   bool updateParticlesFromUni(string filename, bool isSecondarySys, bool isVelData);
-  bool updateGridFromUni(string filename, float *grid, bool isNoise);
-  bool updateGridFromVDB(string filename, float *grid, bool isNoise);
-  bool updateGridFromRaw(string filename, float *grid, bool isNoise);
+  bool updateGridsFromUni(string filename, vector<GridItem> grids);
+  bool updateGridsFromVDB(string filename, vector<GridItem> grids);
+  bool updateGridsFromRaw(string filename, vector<GridItem> grids);
   bool updateMeshFromFile(string filename);
   bool updateParticlesFromFile(string filename, bool isSecondarySys, bool isVelData);
-  bool updateGridFromFile(string filename, float *grid, bool isNoise);
+  bool updateGridsFromFile(string filename, vector<GridItem> grids);
   string getDirectory(struct FluidModifierData *mmd, string subdirectory);
   string getFile(struct FluidModifierData *mmd,
                  string subdirectory,
diff --git a/intern/mantaflow/intern/strings/fluid_script.h b/intern/mantaflow/intern/strings/fluid_script.h
index 922e591c001..cf99717c102 100644
--- a/intern/mantaflow/intern/strings/fluid_script.h
+++ b/intern/mantaflow/intern/strings/fluid_script.h
@@ -259,7 +259,7 @@ const std::string fluid_alloc =
     "\n\
 mantaMsg('Fluid alloc data')\n\
 flags_s$ID$       = s$ID$.create(FlagGrid)\n\
-vel_s$ID$         = s$ID$.create(MACGrid)\n\
+vel_s$ID$         = s$ID$.create(MACGrid, name='$NAME_VELOCITY$')\n\
 velTmp_s$ID$      = s$ID$.create(MACGrid)\n\
 x_vel_s$ID$       = s$ID$.create(RealGrid)\n\
 y_vel_s$ID$       = s$ID$.create(RealGrid)\n\
diff --git a/intern/mantaflow/intern/strings/liquid_script.h b/intern/mantaflow/intern/strings/liquid_script.h
index cdacea16953..e9777eb9cda 100644
--- a/intern/mantaflow/intern/strings/liquid_script.h
+++ b/intern/mantaflow/intern/strings/liquid_script.h
@@ -80,11 +80,11 @@ mantaMsg('Liquid alloc')\n\
 phiParts_s$ID$   = s$ID$.create(LevelsetGrid)\n\
 phi_s$ID$        = s$ID$.create(LevelsetGrid)\n\
 phiTmp_s$ID$     = s$ID$.create(LevelsetGrid)\n\
-curvature_s$ID$  = s$ID$.create(RealGrid)\n\
 velOld_s$ID$     = s$ID$.create(MACGrid)\n\
 velParts_s$ID$   = s$ID$.create(MACGrid)\n\
 mapWeights_s$ID$ = s$ID$.create(MACGrid)\n\
 fractions_s$ID$  = None # allocated dynamically\n\
+curvature_s$ID$  = None\n\
 \n\
 pp_s$ID$         = s$ID$.create(BasicParticleSystem)\n\
 pVel_pp$ID$      = pp_s$ID$.create(PdataVec3)\n\
@@ -124,6 +124,11 @@ liquid_mesh_dict_s$ID$ = dict(lMesh=mesh_sm$ID$)\n\
 if using_speedvectors_s$ID$:\n\
     liquid_meshvel_dict_s$ID$ = dict(lVelMesh=mVel_mesh$ID$)\n";
 
+const std::string liquid_alloc_curvature =
+    "\n\
+mantaMsg('Liquid alloc curvature')\n\
+curvature_s$ID$  = s$ID$.create(RealGrid)\n";
+
 const std::string liquid_alloc_particles =
     "\n\
 ppSnd_sp$ID$         = sp$ID$.create(BasicParticleSystem)\n\
@@ -227,7 +232,7 @@ def liquid_step_$ID$():\n\
     mantaMsg('Liquid step')\n\
     \n\
     mantaMsg('Advecting particles')\n\
-    pp_s$ID$.advectInGrid(flags=flags_s$ID$, vel=vel_s$ID$, integrationMode=IntRK4, deleteInObstacle=deleteInObstacle_s$ID$, stopInObstacle=False)\n\
+    pp_s$ID$.advectInGrid(flags=flags_s$ID$, vel=vel_s$ID$, integrationMode=IntRK4, deleteInObstacle=deleteInObstacle_s$ID$, stopInObstacle=False, skipNew=True)\n\
     \n\
     mantaMsg('Pushing particles out of obstacles')\n\
     pushOutofObs(parts=pp_s$ID$, flags=flags_s$ID$, phiObs=phiObs_s$ID$)\n\
@@ -284,12 +289,13 @@ def liquid_step_$ID$():\n\
         alphaV = viscosity_s$ID$ * s$ID$.timestep * float(res_s$ID$*res_s$ID$)\n\
         setWallBcs(flags=flags_s$ID$, vel=vel_s$ID$, obvel=None if using_fractions_s$ID$ else obvel_s$ID$, phiObs=phiObs_s$ID$, fractions=fractions_s$ID$)\n\
         cgSolveDiffusion(flags_s$ID$, vel_s$ID$, alphaV)\n\
+        \n\
+        mantaMsg('Curvature')\n\
+        getLaplacian(laplacian=curvature_s$ID$, grid=phi_s$ID$)\n\
+        curvature_s$ID$.clamp(-1.0, 1.0)\n\
     \n\
     setWallBcs(flags=flags_s$ID$, vel=vel_s$ID$, obvel=None if using_fractions_s$ID$ else obvel_s$ID$, phiObs=phiObs_s$ID$, fractions=fractions_s$ID$)\n\
     \n\
-    mantaMsg('Calculating curvature')\n\
-    getLaplacian(laplacian=curvature_s$ID$, grid=phi_s$ID$)\n\
-    \n\
     if using_guiding_s$ID$:\n\
         mantaMsg('Guiding and pressure')\n\
         PD_fluid_guiding(vel=vel_s$ID$, velT=velT_s$ID$, flags=flags_s$ID$, phi=phi_s$ID$, curv=curvature_s$ID$, surfTens=surfaceTension_s$ID$, fractions=fractions_s$ID$, weight=weightGuide_s$ID$, blurRadius=beta_sg$ID$, pressure=pressure_s$ID$, tau=tau_sg$ID$, sigma=sigma_sg$ID$, theta=theta_sg$ID$, zeroPressureFixing=not doOpen_s$ID$)\n\
@@ -343,7 +349,10 @@ def liquid_step_mesh_$ID$():\n\
         interpolateMACGrid(target=vel_sm$ID$, source=vel_s$ID$)\n\
         mVel_mesh$ID$.setSource(vel_sm$ID$, isMAC=True)\n\
     \n\
-    phi_sm$ID$.setBound(0.5,int(((upres_sm$ID$)*2)-2) )\n\
+    # Set 0.5 boundary at walls + account for extra wall thickness in fractions mode + account for grid scaling:\n\
+    # E.g. at upres=1 we expect 1 cell border (or 2 with fractions), at upres=2 we expect 2 cell border (or 4 with fractions), etc.\n\
+    # Use -1 since setBound() starts counting at 0 (and additional -1 for fractions to account for solid/fluid interface cells)\n\
+    phi_sm$ID$.setBound(value=0.5, boundaryWidth=(upres_sm$ID$*2)-2 if using_fractions_s$ID$ else upres_sm$ID$-1)\n\
     phi_sm$ID$.createMesh(mesh_sm$ID$)\n";
 
 const std::string liquid_step_particles =
diff --git a/intern/mantaflow/intern/strings/smoke_script.h b/intern/mantaflow/intern/strings/smoke_script.h
index ea2b1e9d848..72d5e20b58b 100644
--- a/intern/mantaflow/intern/strings/smoke_script.h
+++ b/intern/mantaflow/intern/strings/smoke_script.h
@@ -81,10 +81,10 @@ using_fire_s$ID$ = True\n";
 const std::string smoke_alloc =
     "\n\
 mantaMsg('Smoke alloc')\n\
-shadow_s$ID$     = s$ID$.create(RealGrid)\n\
+shadow_s$ID$     = s$ID$.create(RealGrid, name='$NAME_SHADOW$')\n\
 emission_s$ID$   = s$ID$.create(RealGrid)\n\
 emissionIn_s$ID$ = s$ID$.create(RealGrid)\n\
-density_s$ID$    = s$ID$.create(RealGrid)\n\
+density_s$ID$    = s$ID$.create(RealGrid, name='$NAME_DENSITY$')\n\
 densityIn_s$ID$  = s$ID$.create(RealGrid)\n\
 heat_s$ID$       = None # allocated dynamically\n\
 heatIn_s$ID$     = None\n\
@@ -108,7 +108,7 @@ const std::string smoke_alloc_noise =
     "\n\
 mantaMsg('Smoke alloc noise')\n\
 vel_sn$ID$        = sn$ID$.create(MACGrid)\n\
-density_sn$ID$    = sn$ID$.create(RealGrid)\n\
+density_sn$ID$    = sn$ID$.create(RealGrid, name='$NAME_DENSITYNOISE$')\n\
 phiIn_sn$ID$      = sn$ID$.create(LevelsetGrid)\n\
 phiOut_sn$ID$     = sn$ID$.create(LevelsetGrid)\n\
 phiObs_sn$ID$     = sn$ID$.create(LevelsetGrid)\n\
@@ -157,9 +157,9 @@ if 'color_g_s$ID$' in globals(): del color_g_s$ID$\n\
 if 'color_b_s$ID$' in globals(): del color_b_s$ID$\n\
 \n\
 mantaMsg('Allocating colors')\n\
-color_r_s$ID$    = s$ID$.create(RealGrid)\n\
-color_g_s$ID$    = s$ID$.create(RealGrid)\n\
-color_b_s$ID$    = s$ID$.create(RealGrid)\n\
+color_r_s$ID$    = s$ID$.create(RealGrid, name='$NAME_COLORR$')\n\
+color_g_s$ID$    = s$ID$.create(RealGrid, name='$NAME_COLORG$')\n\
+color_b_s$ID$    = s$ID$.create(RealGrid, name='$NAME_COLORB$')\n\
 color_r_in_s$ID$ = s$ID$.create(RealGrid)\n\
 color_g_in_s$ID$ = s$ID$.create(RealGrid)\n\
 color_b_in_s$ID$ = s$ID$.create(RealGrid)\n\
@@ -178,9 +178,9 @@ if 'color_g_sn$ID$' in globals(): del color_g_sn$ID$\n\
 if 'color_b_sn$ID$' in globals(): del color_b_sn$ID$\n\
 \n\
 mantaMsg('Allocating colors noise')\n\
-color_r_sn$ID$ = sn$ID$.create(RealGrid)\n\
-color_g_sn$ID$ = sn$ID$.create(RealGrid)\n\
-color_b_sn$ID$ = sn$ID$.create(RealGrid)\n\
+color_r_sn$ID$ = sn$ID$.create(RealGrid, name='$NAME_COLORRNOISE$')\n\
+color_g_sn$ID$ = sn$ID$.create(RealGrid, name='$NAME_COLORGNOISE$')\n\
+color_b_sn$ID$ = sn$ID$.create(RealGrid, name='$NAME_COLORBNOISE$')\n\
 \n\
 # Add objects to dict to load them later on\n\
 if 'smoke_noise_dict_final_s$ID$' in globals():\n\
@@ -213,7 +213,7 @@ if 'heat_s$ID$' in globals(): del heat_s$ID$\n\
 if 'heatIn_s$ID$' in globals(): del heatIn_s$ID$\n\
 \n\
 mantaMsg('Allocating heat')\n\
-heat_s$ID$   = s$ID$.create(RealGrid)\n\
+heat_s$ID$   = s$ID$.create(RealGrid, name='$NAME_HEAT$')\n\
 heatIn_s$ID$ = s$ID$.create(RealGrid)\n\
 \n\
 # Add objects to dict to load them later on\n\
@@ -232,9 +232,9 @@ if 'fuelIn_s$ID$' in globals(): del fuelIn_s$ID$\n\
 if 'reactIn_s$ID$' in globals(): del reactIn_s$ID$\n\
 \n\
 mantaMsg('Allocating fire')\n\
-flame_s$ID$   = s$ID$.create(RealGrid)\n\
-fuel_s$ID$    = s$ID$.create(RealGrid)\n\
-react_s$ID$   = s$ID$.create(RealGrid)\n\
+flame_s$ID$   = s$ID$.create(RealGrid, name='$NAME_FLAME$')\n\
+fuel_s$ID$    = s$ID$.create(RealGrid, name='$NAME_FUEL$')\n\
+react_s$ID$   = s$ID$.create(RealGrid, name='$NAME_REACT$')\n\
 fuelIn_s$ID$  = s$ID$.create(RealGrid)\n\
 reactIn_s$ID$ = s$ID$.create(RealGrid)\n\
 \n\
@@ -252,9 +252,9 @@ if 'fuel_sn$ID$' in globals(): del fuel_sn$ID$\n\
 if 'react_sn$ID$' in globals(): del react_sn$ID$\n\
 \n\
 mantaMsg('Allocating fire noise')\n\
-flame_sn$ID$ = sn$ID$.create(RealGrid)\n\
-fuel_sn$ID$  = sn$ID$.create(RealGrid)\n\
-react_sn$ID$ = sn$ID$.create(RealGrid)\n\
+flame_sn$ID$ = sn$ID$.create(RealGrid, name='$NAME_FLAMENOISE$')\n\
+fuel_sn$ID$  = sn$ID$.create(RealGrid, name='$NAME_FUELNOISE$')\n\
+react_sn$ID$ = sn$ID$.create(RealGrid, name='$NAME_REACTNOISE$')\n\
 \n\
 # Add objects to dict to load them later on\n\
 if 'smoke_noise_dict_final_s$ID$' in globals():\n\
diff --git a/intern/opensubdiv/CMakeLists.txt b/intern/opensubdiv/CMakeLists.txt
index e7292872e9c..13565a6ed28 100644
--- a/intern/opensubdiv/CMakeLists.txt
+++ b/intern/opensubdiv/CMakeLists.txt
@@ -31,7 +31,6 @@ set(SRC
   opensubdiv_capi_type.h
   opensubdiv_converter_capi.h
   opensubdiv_evaluator_capi.h
-  opensubdiv_gl_mesh_capi.h
   opensubdiv_topology_refiner_capi.h
 )
 
@@ -54,30 +53,20 @@ if(WITH_OPENSUBDIV)
     internal/opensubdiv.cc
     internal/opensubdiv_converter_factory.cc
     internal/opensubdiv_converter_internal.cc
-    internal/opensubdiv_converter_orient.cc
     internal/opensubdiv_device_context_cuda.cc
     internal/opensubdiv_device_context_opencl.cc
     internal/opensubdiv_evaluator.cc
     internal/opensubdiv_evaluator_internal.cc
-    internal/opensubdiv_gl_mesh.cc
-    internal/opensubdiv_gl_mesh_draw.cc
-    internal/opensubdiv_gl_mesh_fvar.cc
-    internal/opensubdiv_gl_mesh_internal.cc
     internal/opensubdiv_topology_refiner.cc
     internal/opensubdiv_topology_refiner_internal.cc
     internal/opensubdiv_util.cc
 
     internal/opensubdiv_converter_factory.h
     internal/opensubdiv_converter_internal.h
-    internal/opensubdiv_converter_orient.h
-    internal/opensubdiv_converter_orient_impl.h
     internal/opensubdiv_device_context_cuda.h
     internal/opensubdiv_device_context_opencl.h
     internal/opensubdiv_edge_map.h
     internal/opensubdiv_evaluator_internal.h
-    internal/opensubdiv_gl_mesh_draw.h
-    internal/opensubdiv_gl_mesh_fvar.h
-    internal/opensubdiv_gl_mesh_internal.h
     internal/opensubdiv_internal.h
     internal/opensubdiv_topology_refiner_internal.h
     internal/opensubdiv_util.h
@@ -101,11 +90,7 @@ if(WITH_OPENSUBDIV)
   OPENSUBDIV_DEFINE_COMPONENT(OPENSUBDIV_HAS_GLSL_TRANSFORM_FEEDBACK)
   OPENSUBDIV_DEFINE_COMPONENT(OPENSUBDIV_HAS_GLSL_COMPUTE)
 
-  data_to_c_simple(shader/gpu_shader_opensubdiv_vertex.glsl SRC)
-  data_to_c_simple(shader/gpu_shader_opensubdiv_geometry.glsl SRC)
-  data_to_c_simple(shader/gpu_shader_opensubdiv_fragment.glsl SRC)
-
-  add_definitions(-DGLEW_STATIC)
+  add_definitions(${GL_DEFINITIONS})
   add_definitions(-DOSD_USES_GLEW)
 
   if(WIN32)
@@ -131,7 +116,6 @@ else()
   list(APPEND SRC
     stub/opensubdiv_stub.cc
     stub/opensubdiv_evaluator_stub.cc
-    stub/opensubdiv_gl_mesh_stub.cc
     stub/opensubdiv_topology_refiner_stub.cc
   )
 endif()
diff --git a/intern/opensubdiv/internal/opensubdiv.cc b/intern/opensubdiv/internal/opensubdiv.cc
index 74b81b13351..e9f6086851b 100644
--- a/intern/opensubdiv/internal/opensubdiv.cc
+++ b/intern/opensubdiv/internal/opensubdiv.cc
@@ -24,7 +24,6 @@
 
 #include "opensubdiv_device_context_cuda.h"
 #include "opensubdiv_device_context_opencl.h"
-#include "opensubdiv_gl_mesh_capi.h"
 
 void openSubdiv_init(void)
 {
@@ -34,7 +33,6 @@ void openSubdiv_init(void)
 
 void openSubdiv_cleanup(void)
 {
-  openSubdiv_deinitGLMeshDrawingResources();
 }
 
 int openSubdiv_getAvailableEvaluators(void)
@@ -86,7 +84,7 @@ int openSubdiv_getVersionHex(void)
   }
   int major = 0, minor = 0, patch = 0;
   vector<string> tokens;
-  opensubdiv_capi::stringSplit(&tokens, version, "_", true);
+  blender::opensubdiv::stringSplit(&tokens, version, "_", true);
   if (tokens.size() == 3) {
     major = atoi(tokens[0].c_str());
     minor = atoi(tokens[1].c_str());
diff --git a/intern/opensubdiv/internal/opensubdiv_converter_factory.cc b/intern/opensubdiv/internal/opensubdiv_converter_factory.cc
index ab93b5ce952..dba2a969062 100644
--- a/intern/opensubdiv/internal/opensubdiv_converter_factory.cc
+++ b/intern/opensubdiv/internal/opensubdiv_converter_factory.cc
@@ -28,14 +28,13 @@
 #include <opensubdiv/far/topologyRefinerFactory.h>
 
 #include "internal/opensubdiv_converter_internal.h"
-#include "internal/opensubdiv_converter_orient.h"
 #include "internal/opensubdiv_internal.h"
 #include "internal/opensubdiv_util.h"
 #include "opensubdiv_converter_capi.h"
 
-using opensubdiv_capi::min;
-using opensubdiv_capi::stack;
-using opensubdiv_capi::vector;
+using blender::opensubdiv::min;
+using blender::opensubdiv::stack;
+using blender::opensubdiv::vector;
 
 struct TopologyRefinerData {
   const OpenSubdiv_Converter *converter;
@@ -246,7 +245,8 @@ inline void TopologyRefinerFactory<TopologyRefinerData>::reportInvalidTopology(
 } /* namespace OPENSUBDIV_VERSION */
 } /* namespace OpenSubdiv */
 
-namespace opensubdiv_capi {
+namespace blender {
+namespace opensubdiv {
 
 namespace {
 
@@ -292,4 +292,5 @@ OpenSubdiv::Far::TopologyRefiner *createOSDTopologyRefinerFromConverter(
   return TopologyRefinerFactory<TopologyRefinerData>::Create(cb_data, topology_options);
 }
 
-}  // namespace opensubdiv_capi
+}  // namespace opensubdiv
+}  // namespace blender
diff --git a/intern/opensubdiv/internal/opensubdiv_converter_factory.h b/intern/opensubdiv/internal/opensubdiv_converter_factory.h
index a1038474d33..3519d3059b2 100644
--- a/intern/opensubdiv/internal/opensubdiv_converter_factory.h
+++ b/intern/opensubdiv/internal/opensubdiv_converter_factory.h
@@ -27,11 +27,13 @@
 
 struct OpenSubdiv_Converter;
 
-namespace opensubdiv_capi {
+namespace blender {
+namespace opensubdiv {
 
 OpenSubdiv::Far::TopologyRefiner *createOSDTopologyRefinerFromConverter(
     struct OpenSubdiv_Converter *converter);
 
-}  // namespace opensubdiv_capi
+}  // namespace opensubdiv
+}  // namespace blender
 
 #endif  // OPENSUBDIV_CONVERTER_FACTORY_H_
diff --git a/intern/opensubdiv/internal/opensubdiv_converter_internal.cc b/intern/opensubdiv/internal/opensubdiv_converter_internal.cc
index 0335219d6b9..eedca88f77b 100644
--- a/intern/opensubdiv/internal/opensubdiv_converter_internal.cc
+++ b/intern/opensubdiv/internal/opensubdiv_converter_internal.cc
@@ -25,7 +25,8 @@
 #include <cassert>
 #include <opensubdiv/sdc/crease.h>
 
-namespace opensubdiv_capi {
+namespace blender {
+namespace opensubdiv {
 
 OpenSubdiv::Sdc::SchemeType getSchemeTypeFromCAPI(OpenSubdiv_SchemeType type)
 {
@@ -85,4 +86,5 @@ OpenSubdiv_FVarLinearInterpolation getCAPIFVarLinearInterpolationFromOSD(
   return OSD_FVAR_LINEAR_INTERPOLATION_NONE;
 }
 
-}  // namespace opensubdiv_capi
+}  // namespace opensubdiv
+}  // namespace blender
diff --git a/intern/opensubdiv/internal/opensubdiv_converter_internal.h b/intern/opensubdiv/internal/opensubdiv_converter_internal.h
index 11c6bdd7f3b..7c586b0787a 100644
--- a/intern/opensubdiv/internal/opensubdiv_converter_internal.h
+++ b/intern/opensubdiv/internal/opensubdiv_converter_internal.h
@@ -30,7 +30,8 @@
 
 struct OpenSubdiv_Converter;
 
-namespace opensubdiv_capi {
+namespace blender {
+namespace opensubdiv {
 
 // Convert scheme type from C-API enum to an OpenSubdiv native enum.
 OpenSubdiv::Sdc::SchemeType getSchemeTypeFromCAPI(OpenSubdiv_SchemeType type);
@@ -44,6 +45,7 @@ OpenSubdiv::Sdc::Options::FVarLinearInterpolation getFVarLinearInterpolationFrom
 OpenSubdiv_FVarLinearInterpolation getCAPIFVarLinearInterpolationFromOSD(
     OpenSubdiv::Sdc::Options::FVarLinearInterpolation linear_interpolation);
 
-}  // namespace opensubdiv_capi
+}  // namespace opensubdiv
+}  // namespace blender
 
 #endif  // OPENSUBDIV_CONVERTER_INTERNAL_H_
diff --git a/intern/opensubdiv/internal/opensubdiv_converter_orient.cc b/intern/opensubdiv/internal/opensubdiv_converter_orient.cc
deleted file mode 100644
index e3367fc6314..00000000000
--- a/intern/opensubdiv/internal/opensubdiv_converter_orient.cc
+++ /dev/null
@@ -1,67 +0,0 @@
-// Copyright 2018 Blender Foundation. All rights reserved.
-//
-// This program is free software; you can redistribute it and/or
-// modify it under the terms of the GNU General Public License
-// as published by the Free Software Foundation; either version 2
-// of the License, or (at your option) any later version.
-//
-// This program is distributed in the hope that it will be useful,
-// but WITHOUT ANY WARRANTY; without even the implied warranty of
-// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-// GNU General Public License for more details.
-//
-// You should have received a copy of the GNU General Public License
-// along with this program; if not, write to the Free Software Foundation,
-// Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
-//
-// Author: Sergey Sharybin
-
-#include "internal/opensubdiv_converter_orient.h"
-
-#include "internal/opensubdiv_internal.h"
-
-namespace opensubdiv_capi {
-
-void checkOrientedVertexConnectivity(const int num_vertex_edges,
-                                     const int num_vertex_faces,
-                                     const int *vertex_edges,
-                                     const int *vertex_faces,
-                                     const int *dst_vertex_edges,
-                                     const int *dst_vertex_faces)
-{
-#ifndef NDEBUG
-  for (int i = 0; i < num_vertex_faces; ++i) {
-    bool found = false;
-    for (int j = 0; j < num_vertex_faces; ++j) {
-      if (vertex_faces[i] == dst_vertex_faces[j]) {
-        found = true;
-        break;
-      }
-    }
-    if (!found) {
-      assert(!"vert-faces connectivity ruined");
-    }
-  }
-  for (int i = 0; i < num_vertex_edges; ++i) {
-    bool found = false;
-    for (int j = 0; j < num_vertex_edges; ++j) {
-      if (vertex_edges[i] == dst_vertex_edges[j]) {
-        found = true;
-        break;
-      }
-    }
-    if (!found) {
-      assert(!"vert-edges connectivity ruined");
-    }
-  }
-#else
-  (void)num_vertex_edges;
-  (void)num_vertex_faces;
-  (void)vertex_edges;
-  (void)vertex_faces;
-  (void)dst_vertex_edges;
-  (void)dst_vertex_faces;
-#endif
-}
-
-}  // namespace opensubdiv_capi
diff --git a/intern/opensubdiv/internal/opensubdiv_converter_orient.h b/intern/opensubdiv/internal/opensubdiv_converter_orient.h
deleted file mode 100644
index 967871845cb..00000000000
--- a/intern/opensubdiv/internal/opensubdiv_converter_orient.h
+++ /dev/null
@@ -1,49 +0,0 @@
-// Copyright 2018 Blender Foundation. All rights reserved.
-//
-// This program is free software; you can redistribute it and/or
-// modify it under the terms of the GNU General Public License
-// as published by the Free Software Foundation; either version 2
-// of the License, or (at your option) any later version.
-//
-// This program is distributed in the hope that it will be useful,
-// but WITHOUT ANY WARRANTY; without even the implied warranty of
-// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-// GNU General Public License for more details.
-//
-// You should have received a copy of the GNU General Public License
-// along with this program; if not, write to the Free Software Foundation,
-// Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
-//
-// Author: Sergey Sharybin
-
-#ifndef OPENSUBDIV_CONVERTER_ORIENT_H_
-#  define OPENSUBDIV_CONVERTER_ORIENT_H_
-
-#  include <opensubdiv/far/types.h>
-
-// Set of utility functions which are needed to bring topology to an orientation
-// (or, winding, if you wish) which OpenSubdiv expects.
-
-namespace opensubdiv_capi {
-
-inline void reverseFaceVertices(int *face_vertices, const int num_vertices);
-
-// TODO(sergey): Document which value corresponds to which winding.
-inline int getLoopWinding(int vert0_of_face, int vert1_of_face);
-
-inline void reverseFaceLoops(OpenSubdiv::Far::IndexArray *face_vertices,
-                             OpenSubdiv::Far::IndexArray *face_edges);
-
-// Used for debugging, checks whether orientation happened correct.
-void checkOrientedVertexConnectivity(const int num_vertex_edges,
-                                     const int num_vertex_faces,
-                                     const int *vertex_edges,
-                                     const int *vertex_faces,
-                                     const int *dst_vertex_edges,
-                                     const int *dst_vertex_faces);
-
-}  // namespace opensubdiv_capi
-
-#endif  // OPENSUBDIV_CONVERTER_ORIENT_H_
-
-#include "internal/opensubdiv_converter_orient_impl.h"
diff --git a/intern/opensubdiv/internal/opensubdiv_converter_orient_impl.h b/intern/opensubdiv/internal/opensubdiv_converter_orient_impl.h
deleted file mode 100644
index aa717f5d99d..00000000000
--- a/intern/opensubdiv/internal/opensubdiv_converter_orient_impl.h
+++ /dev/null
@@ -1,69 +0,0 @@
-// Copyright 2018 Blender Foundation. All rights reserved.
-//
-// This program is free software; you can redistribute it and/or
-// modify it under the terms of the GNU General Public License
-// as published by the Free Software Foundation; either version 2
-// of the License, or (at your option) any later version.
-//
-// This program is distributed in the hope that it will be useful,
-// but WITHOUT ANY WARRANTY; without even the implied warranty of
-// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-// GNU General Public License for more details.
-//
-// You should have received a copy of the GNU General Public License
-// along with this program; if not, write to the Free Software Foundation,
-// Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
-//
-// Author: Sergey Sharybin
-
-#ifndef OPENSUBDIV_CONVERTER_ORIENT_IMPL_H_
-#define OPENSUBDIV_CONVERTER_ORIENT_IMPL_H_
-
-#include "internal/opensubdiv_converter_orient.h"
-
-#include <cmath>
-
-#include "internal/opensubdiv_util.h"
-
-namespace opensubdiv_capi {
-
-inline void reverseFaceVertices(int *face_vertices, const int num_vertices)
-{
-  int last_vert = face_vertices[num_vertices - 1];
-  for (int i = num_vertices - 1; i > 0; --i) {
-    face_vertices[i] = face_vertices[i - 1];
-  }
-  face_vertices[0] = last_vert;
-}
-
-inline int getLoopWinding(int vert0_of_face, int vert1_of_face)
-{
-  int delta_face = vert1_of_face - vert0_of_face;
-  if (abs(delta_face) != 1) {
-    if (delta_face > 0) {
-      delta_face = -1;
-    }
-    else {
-      delta_face = 1;
-    }
-  }
-  return delta_face;
-}
-
-inline void reverseFaceLoops(OpenSubdiv::Far::IndexArray *face_vertices,
-                             OpenSubdiv::Far::IndexArray *face_edges)
-{
-  const int num_face_vertices = face_vertices->size();
-  for (int i = 0; i < num_face_vertices / 2; ++i) {
-    const int j = num_face_vertices - i - 1;
-    if (i != j) {
-      swap((*face_vertices)[i], (*face_vertices)[j]);
-      swap((*face_edges)[i], (*face_edges)[j]);
-    }
-  }
-  reverseFaceVertices(&(*face_vertices)[0], num_face_vertices);
-}
-
-}  // namespace opensubdiv_capi
-
-#endif  // OPENSUBDIV_CONVERTER_ORIENT_IMPL_H_
diff --git a/intern/opensubdiv/internal/opensubdiv_edge_map.h b/intern/opensubdiv/internal/opensubdiv_edge_map.h
index e2e6d2328fe..454068b58a4 100644
--- a/intern/opensubdiv/internal/opensubdiv_edge_map.h
+++ b/intern/opensubdiv/internal/opensubdiv_edge_map.h
@@ -21,7 +21,8 @@
 
 #include "internal/opensubdiv_util.h"
 
-namespace opensubdiv_capi {
+namespace blender {
+namespace opensubdiv {
 
 // Helper class to ease dealing with edge indexing.
 // Simply takes care of ensuring order of vertices is strictly defined.
@@ -144,12 +145,13 @@ template<typename T> typename EdgeTagMap<T>::value_type &EdgeTagMap<T>::operator
   return edge_tags_[key];
 }
 
-}  // namespace opensubdiv_capi
+}  // namespace opensubdiv
+}  // namespace blender
 
 namespace std {
 
-template<> struct hash<opensubdiv_capi::EdgeKey> {
-  std::size_t operator()(const opensubdiv_capi::EdgeKey &key) const
+template<> struct hash<blender::opensubdiv::EdgeKey> {
+  std::size_t operator()(const blender::opensubdiv::EdgeKey &key) const
   {
     return key.hash();
   }
diff --git a/intern/opensubdiv/internal/opensubdiv_evaluator_internal.cc b/intern/opensubdiv/internal/opensubdiv_evaluator_internal.cc
index c35909a045b..5279752ea4e 100644
--- a/intern/opensubdiv/internal/opensubdiv_evaluator_internal.cc
+++ b/intern/opensubdiv/internal/opensubdiv_evaluator_internal.cc
@@ -53,7 +53,8 @@ using OpenSubdiv::Osd::CpuPatchTable;
 using OpenSubdiv::Osd::CpuVertexBuffer;
 using OpenSubdiv::Osd::PatchCoord;
 
-namespace opensubdiv_capi {
+namespace blender {
+namespace opensubdiv {
 
 namespace {
 
@@ -731,7 +732,8 @@ void CpuEvalOutputAPI::evaluatePatchesLimit(const OpenSubdiv_PatchCoord *patch_c
   }
 }
 
-}  // namespace opensubdiv_capi
+}  // namespace opensubdiv
+}  // namespace blender
 
 OpenSubdiv_EvaluatorInternal::OpenSubdiv_EvaluatorInternal()
     : eval_output(NULL), patch_map(NULL), patch_table(NULL)
@@ -748,7 +750,7 @@ OpenSubdiv_EvaluatorInternal::~OpenSubdiv_EvaluatorInternal()
 OpenSubdiv_EvaluatorInternal *openSubdiv_createEvaluatorInternal(
     OpenSubdiv_TopologyRefiner *topology_refiner)
 {
-  using opensubdiv_capi::vector;
+  using blender::opensubdiv::vector;
   TopologyRefiner *refiner = topology_refiner->internal->osd_topology_refiner;
   if (refiner == NULL) {
     // Happens on bad topology.
@@ -851,13 +853,13 @@ OpenSubdiv_EvaluatorInternal *openSubdiv_createEvaluatorInternal(
   }
   // Create OpenSubdiv's CPU side evaluator.
   // TODO(sergey): Make it possible to use different evaluators.
-  opensubdiv_capi::CpuEvalOutput *eval_output = new opensubdiv_capi::CpuEvalOutput(
+  blender::opensubdiv::CpuEvalOutput *eval_output = new blender::opensubdiv::CpuEvalOutput(
       vertex_stencils, varying_stencils, all_face_varying_stencils, 2, patch_table);
   OpenSubdiv::Far::PatchMap *patch_map = new PatchMap(*patch_table);
   // Wrap everything we need into an object which we control from our side.
   OpenSubdiv_EvaluatorInternal *evaluator_descr;
   evaluator_descr = OBJECT_GUARDED_NEW(OpenSubdiv_EvaluatorInternal);
-  evaluator_descr->eval_output = new opensubdiv_capi::CpuEvalOutputAPI(eval_output, patch_map);
+  evaluator_descr->eval_output = new blender::opensubdiv::CpuEvalOutputAPI(eval_output, patch_map);
   evaluator_descr->patch_map = patch_map;
   evaluator_descr->patch_table = patch_table;
   // TOOD(sergey): Look into whether we've got duplicated stencils arrays.
diff --git a/intern/opensubdiv/internal/opensubdiv_evaluator_internal.h b/intern/opensubdiv/internal/opensubdiv_evaluator_internal.h
index 392633944c6..dbe4d88539f 100644
--- a/intern/opensubdiv/internal/opensubdiv_evaluator_internal.h
+++ b/intern/opensubdiv/internal/opensubdiv_evaluator_internal.h
@@ -29,7 +29,8 @@
 struct OpenSubdiv_PatchCoord;
 struct OpenSubdiv_TopologyRefiner;
 
-namespace opensubdiv_capi {
+namespace blender {
+namespace opensubdiv {
 
 // Anonymous forward declaration of actual evaluator implementation.
 class CpuEvalOutput;
@@ -132,14 +133,15 @@ class CpuEvalOutputAPI {
   OpenSubdiv::Far::PatchMap *patch_map_;
 };
 
-}  // namespace opensubdiv_capi
+}  // namespace opensubdiv
+}  // namespace blender
 
 struct OpenSubdiv_EvaluatorInternal {
  public:
   OpenSubdiv_EvaluatorInternal();
   ~OpenSubdiv_EvaluatorInternal();
 
-  opensubdiv_capi::CpuEvalOutputAPI *eval_output;
+  blender::opensubdiv::CpuEvalOutputAPI *eval_output;
   const OpenSubdiv::Far::PatchMap *patch_map;
   const OpenSubdiv::Far::PatchTable *patch_table;
 };
diff --git a/intern/opensubdiv/internal/opensubdiv_gl_mesh.cc b/intern/opensubdiv/internal/opensubdiv_gl_mesh.cc
deleted file mode 100644
index 6afd763a63e..00000000000
--- a/intern/opensubdiv/internal/opensubdiv_gl_mesh.cc
+++ /dev/null
@@ -1,260 +0,0 @@
-// Copyright 2013 Blender Foundation. All rights reserved.
-//
-// This program is free software; you can redistribute it and/or
-// modify it under the terms of the GNU General Public License
-// as published by the Free Software Foundation; either version 2
-// of the License, or (at your option) any later version.
-//
-// This program is distributed in the hope that it will be useful,
-// but WITHOUT ANY WARRANTY; without even the implied warranty of
-// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-// GNU General Public License for more details.
-//
-// You should have received a copy of the GNU General Public License
-// along with this program; if not, write to the Free Software Foundation,
-// Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
-
-#include "opensubdiv_gl_mesh_capi.h"
-
-#ifdef _MSC_VER
-#  include <iso646.h>
-#endif
-
-#include <opensubdiv/far/stencilTable.h>
-#include <opensubdiv/osd/glMesh.h>
-#include <opensubdiv/osd/glPatchTable.h>
-
-using OpenSubdiv::Far::StencilTable;
-using OpenSubdiv::Osd::GLMeshInterface;
-using OpenSubdiv::Osd::GLPatchTable;
-using OpenSubdiv::Osd::Mesh;
-using OpenSubdiv::Osd::MeshBitset;
-
-// CPU backend.
-#include <opensubdiv/osd/cpuEvaluator.h>
-#include <opensubdiv/osd/cpuGLVertexBuffer.h>
-using OpenSubdiv::Osd::CpuEvaluator;
-using OpenSubdiv::Osd::CpuGLVertexBuffer;
-typedef Mesh<CpuGLVertexBuffer, StencilTable, CpuEvaluator, GLPatchTable> OsdCpuMesh;
-// OpenMP backend.
-#ifdef OPENSUBDIV_HAS_OPENMP
-#  include <opensubdiv/osd/ompEvaluator.h>
-using OpenSubdiv::Osd::OmpEvaluator;
-typedef Mesh<CpuGLVertexBuffer, StencilTable, OmpEvaluator, GLPatchTable> OsdOmpMesh;
-#endif
-// OpenCL backend.
-#ifdef OPENSUBDIV_HAS_OPENCL
-#  include "opensubdiv_device_context_opencl.h"
-#  include <opensubdiv/osd/clEvaluator.h>
-#  include <opensubdiv/osd/clGLVertexBuffer.h>
-using OpenSubdiv::Osd::CLEvaluator;
-using OpenSubdiv::Osd::CLGLVertexBuffer;
-using OpenSubdiv::Osd::CLStencilTable;
-/* TODO(sergey): Use CLDeviceContext similar to OSD examples? */
-typedef Mesh<CLGLVertexBuffer, CLStencilTable, CLEvaluator, GLPatchTable, CLDeviceContext>
-    OsdCLMesh;
-static CLDeviceContext g_cl_device_context;
-#endif
-// CUDA backend.
-#ifdef OPENSUBDIV_HAS_CUDA
-#  include "opensubdiv_device_context_cuda.h"
-#  include <opensubdiv/osd/cudaEvaluator.h>
-#  include <opensubdiv/osd/cudaGLVertexBuffer.h>
-using OpenSubdiv::Osd::CudaEvaluator;
-using OpenSubdiv::Osd::CudaGLVertexBuffer;
-using OpenSubdiv::Osd::CudaStencilTable;
-typedef Mesh<CudaGLVertexBuffer, CudaStencilTable, CudaEvaluator, GLPatchTable> OsdCudaMesh;
-static CudaDeviceContext g_cuda_device_context;
-#endif
-// Transform feedback backend.
-#ifdef OPENSUBDIV_HAS_GLSL_TRANSFORM_FEEDBACK
-#  include <opensubdiv/osd/glVertexBuffer.h>
-#  include <opensubdiv/osd/glXFBEvaluator.h>
-using OpenSubdiv::Osd::GLStencilTableTBO;
-using OpenSubdiv::Osd::GLVertexBuffer;
-using OpenSubdiv::Osd::GLXFBEvaluator;
-typedef Mesh<GLVertexBuffer, GLStencilTableTBO, GLXFBEvaluator, GLPatchTable>
-    OsdGLSLTransformFeedbackMesh;
-#endif
-// GLSL compute backend.
-#ifdef OPENSUBDIV_HAS_GLSL_COMPUTE
-#  include <opensubdiv/osd/glComputeEvaluator.h>
-#  include <opensubdiv/osd/glVertexBuffer.h>
-using OpenSubdiv::Osd::GLComputeEvaluator;
-using OpenSubdiv::Osd::GLStencilTableSSBO;
-using OpenSubdiv::Osd::GLVertexBuffer;
-typedef Mesh<GLVertexBuffer, GLStencilTableSSBO, GLComputeEvaluator, GLPatchTable>
-    OsdGLSLComputeMesh;
-#endif
-
-#include "MEM_guardedalloc.h"
-
-#include "internal/opensubdiv_gl_mesh_draw.h"
-#include "internal/opensubdiv_gl_mesh_fvar.h"
-#include "internal/opensubdiv_gl_mesh_internal.h"
-#include "internal/opensubdiv_topology_refiner_internal.h"
-#include "internal/opensubdiv_util.h"
-#include "opensubdiv_topology_refiner_capi.h"
-
-using opensubdiv_capi::vector;
-
-namespace {
-
-GLMeshInterface *createGLMeshInterface(OpenSubdiv::Far::TopologyRefiner *topology_refiner,
-                                       const MeshBitset &bits,
-                                       const int num_vertex_elements,
-                                       const int num_varying_elements,
-                                       const int level,
-                                       eOpenSubdivEvaluator evaluator_type)
-{
-  GLMeshInterface *mesh = NULL;
-  switch (evaluator_type) {
-#define CHECK_EVALUATOR_TYPE(type, class) \
-  case OPENSUBDIV_EVALUATOR_##type: \
-    mesh = new class(topology_refiner, num_vertex_elements, num_varying_elements, level, bits); \
-    break;
-
-#define CHECK_EVALUATOR_TYPE_STUB(type) \
-  case OPENSUBDIV_EVALUATOR_##type: \
-    mesh = NULL; \
-    break;
-
-    CHECK_EVALUATOR_TYPE(CPU, OsdCpuMesh)
-#ifdef OPENSUBDIV_HAS_OPENMP
-    CHECK_EVALUATOR_TYPE(OPENMP, OsdOmpMesh)
-#else
-    CHECK_EVALUATOR_TYPE_STUB(OPENMP)
-#endif
-#ifdef OPENSUBDIV_HAS_OPENCL
-    CHECK_EVALUATOR_TYPE(OPENCL, OsdCLMesh)
-#else
-    CHECK_EVALUATOR_TYPE_STUB(OPENCL)
-#endif
-#ifdef OPENSUBDIV_HAS_CUDA
-    CHECK_EVALUATOR_TYPE(CUDA, OsdCudaMesh)
-#else
-    CHECK_EVALUATOR_TYPE_STUB(CUDA)
-#endif
-#ifdef OPENSUBDIV_HAS_GLSL_TRANSFORM_FEEDBACK
-    CHECK_EVALUATOR_TYPE(GLSL_TRANSFORM_FEEDBACK, OsdGLSLTransformFeedbackMesh)
-#else
-    CHECK_EVALUATOR_TYPE_STUB(GLSL_TRANSFORM_FEEDBACK)
-#endif
-#ifdef OPENSUBDIV_HAS_GLSL_COMPUTE
-    CHECK_EVALUATOR_TYPE(GLSL_COMPUTE, OsdGLSLComputeMesh)
-#else
-    CHECK_EVALUATOR_TYPE_STUB(GLSL_COMPUTE)
-#endif
-
-#undef CHECK_EVALUATOR_TYPE
-#undef CHECK_EVALUATOR_TYPE_STUB
-  }
-  return mesh;
-}
-
-////////////////////////////////////////////////////////////////////////////////
-// GLMesh structure "methods".
-
-opensubdiv_capi::GLMeshFVarData *createFVarData(OpenSubdiv::Far::TopologyRefiner *topology_refiner,
-                                                GLMeshInterface *mesh,
-                                                const float *fvar_src_buffer)
-{
-  using opensubdiv_capi::GLMeshFVarData;
-  GLMeshFVarData *fvar_data = new GLMeshFVarData();
-  fvar_data->create(topology_refiner, mesh->GetFarPatchTable(), 2, fvar_src_buffer);
-  return fvar_data;
-}
-
-unsigned int getPatchIndexBuffer(OpenSubdiv_GLMesh *gl_mesh)
-{
-  return gl_mesh->internal->mesh_interface->GetPatchTable()->GetPatchIndexBuffer();
-}
-
-void bindVertexBuffer(OpenSubdiv_GLMesh *gl_mesh)
-{
-  gl_mesh->internal->mesh_interface->BindVertexBuffer();
-}
-
-void setCoarsePositions(OpenSubdiv_GLMesh *gl_mesh,
-                        const float *positions,
-                        const int start_vertex,
-                        const int num_vertices)
-{
-  gl_mesh->internal->mesh_interface->UpdateVertexBuffer(positions, start_vertex, num_vertices);
-}
-
-void refine(OpenSubdiv_GLMesh *gl_mesh)
-{
-  gl_mesh->internal->mesh_interface->Refine();
-}
-
-void synchronize(struct OpenSubdiv_GLMesh *gl_mesh)
-{
-  gl_mesh->internal->mesh_interface->Synchronize();
-}
-
-void assignFunctionPointers(OpenSubdiv_GLMesh *gl_mesh)
-{
-  gl_mesh->getPatchIndexBuffer = getPatchIndexBuffer;
-  gl_mesh->bindVertexBuffer = bindVertexBuffer;
-  gl_mesh->setCoarsePositions = setCoarsePositions;
-  gl_mesh->refine = refine;
-  gl_mesh->synchronize = synchronize;
-
-  gl_mesh->prepareDraw = opensubdiv_capi::GLMeshDisplayPrepare;
-  gl_mesh->drawPatches = opensubdiv_capi::GLMeshDisplayDrawPatches;
-}
-
-}  // namespace
-
-struct OpenSubdiv_GLMesh *openSubdiv_createOsdGLMeshFromTopologyRefiner(
-    OpenSubdiv_TopologyRefiner *topology_refiner, eOpenSubdivEvaluator evaluator_type)
-{
-  using OpenSubdiv::Far::TopologyRefiner;
-  TopologyRefiner *osd_topology_refiner = topology_refiner->internal->osd_topology_refiner;
-  // TODO(sergey): Query this from refiner.
-  const bool is_adaptive = false;
-  MeshBitset bits;
-  bits.set(OpenSubdiv::Osd::MeshAdaptive, is_adaptive);
-  bits.set(OpenSubdiv::Osd::MeshUseSingleCreasePatch, 0);
-  bits.set(OpenSubdiv::Osd::MeshInterleaveVarying, 1);
-  bits.set(OpenSubdiv::Osd::MeshFVarData, 1);
-  bits.set(OpenSubdiv::Osd::MeshEndCapBSplineBasis, 1);
-  const int num_vertex_elements = 3;
-  const int num_varying_elements = 3;
-  GLMeshInterface *mesh = createGLMeshInterface(osd_topology_refiner,
-                                                bits,
-                                                num_vertex_elements,
-                                                num_varying_elements,
-                                                osd_topology_refiner->GetMaxLevel(),
-                                                evaluator_type);
-  if (mesh == NULL) {
-    return NULL;
-  }
-  OpenSubdiv_GLMesh *gl_mesh = OBJECT_GUARDED_NEW(OpenSubdiv_GLMesh);
-  assignFunctionPointers(gl_mesh);
-  gl_mesh->internal = new OpenSubdiv_GLMeshInternal();
-  gl_mesh->internal->evaluator_type = evaluator_type;
-  gl_mesh->internal->mesh_interface = mesh;
-  // Face-varying support.
-  // TODO(sergey): This part needs to be re-done.
-  if (osd_topology_refiner->GetNumFVarChannels() > 0) {
-    // TODO(sergey): This is a temporary stub to get things compiled. Need
-    // to store base level UVs somewhere else.
-    vector<float> uvs;
-    vector<float> fvar_data_buffer;
-    opensubdiv_capi::interpolateFVarData(*osd_topology_refiner, uvs, &fvar_data_buffer);
-    gl_mesh->internal->fvar_data = createFVarData(
-        osd_topology_refiner, mesh, &fvar_data_buffer[0]);
-  }
-  else {
-    gl_mesh->internal->fvar_data = NULL;
-  }
-  return gl_mesh;
-}
-
-void openSubdiv_deleteOsdGLMesh(OpenSubdiv_GLMesh *gl_mesh)
-{
-  delete gl_mesh->internal;
-  OBJECT_GUARDED_DELETE(gl_mesh, OpenSubdiv_GLMesh);
-}
diff --git a/intern/opensubdiv/internal/opensubdiv_gl_mesh_draw.cc b/intern/opensubdiv/internal/opensubdiv_gl_mesh_draw.cc
deleted file mode 100644
index cbccf69e02d..00000000000
--- a/intern/opensubdiv/internal/opensubdiv_gl_mesh_draw.cc
+++ /dev/null
@@ -1,577 +0,0 @@
-// Copyright 2013 Blender Foundation. All rights reserved.
-//
-// This program is free software; you can redistribute it and/or
-// modify it under the terms of the GNU General Public License
-// as published by the Free Software Foundation; either version 2
-// of the License, or (at your option) any later version.
-//
-// This program is distributed in the hope that it will be useful,
-// but WITHOUT ANY WARRANTY; without even the implied warranty of
-// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-// GNU General Public License for more details.
-//
-// You should have received a copy of the GNU General Public License
-// along with this program; if not, write to the Free Software Foundation,
-// Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
-//
-// Author: Sergey Sharybin
-
-#include "internal/opensubdiv_gl_mesh_draw.h"
-
-#ifdef _MSC_VER
-#  include <iso646.h>
-#endif
-
-#include <GL/glew.h>
-#include <cmath>
-#include <cstdio>
-
-#include <opensubdiv/osd/glMesh.h>
-
-#ifdef OPENSUBDIV_HAS_CUDA
-#  include <opensubdiv/osd/cudaGLVertexBuffer.h>
-#endif  // OPENSUBDIV_HAS_CUDA
-
-#include <opensubdiv/osd/cpuEvaluator.h>
-#include <opensubdiv/osd/cpuGLVertexBuffer.h>
-
-#include "internal/opensubdiv_gl_mesh_fvar.h"
-#include "internal/opensubdiv_gl_mesh_internal.h"
-#include "internal/opensubdiv_util.h"
-#include "opensubdiv_capi.h"
-#include "opensubdiv_gl_mesh_capi.h"
-
-using OpenSubdiv::Osd::GLMeshInterface;
-
-extern "C" char datatoc_gpu_shader_opensubdiv_vertex_glsl[];
-extern "C" char datatoc_gpu_shader_opensubdiv_geometry_glsl[];
-extern "C" char datatoc_gpu_shader_opensubdiv_fragment_glsl[];
-
-// TODO(sergey): Those are a bit of bad level calls :S
-extern "C" {
-void copy_m3_m3(float m1[3][3], float m2[3][3]);
-void copy_m3_m4(float m1[3][3], float m2[4][4]);
-void adjoint_m3_m3(float m1[3][3], float m[3][3]);
-float determinant_m3_array(float m[3][3]);
-bool invert_m3_m3(float m1[3][3], float m2[3][3]);
-bool invert_m3(float m[3][3]);
-void transpose_m3(float mat[3][3]);
-}
-
-#define MAX_LIGHTS 8
-#define SUPPORT_COLOR_MATERIAL
-
-typedef struct Light {
-  float position[4];
-  float ambient[4];
-  float diffuse[4];
-  float specular[4];
-  float spot_direction[4];
-#ifdef SUPPORT_COLOR_MATERIAL
-  float constant_attenuation;
-  float linear_attenuation;
-  float quadratic_attenuation;
-  float spot_cutoff;
-  float spot_exponent;
-  float spot_cos_cutoff;
-  float pad, pad2;
-#endif
-} Light;
-
-typedef struct Lighting {
-  Light lights[MAX_LIGHTS];
-  int num_enabled;
-} Lighting;
-
-typedef struct Transform {
-  float projection_matrix[16];
-  float model_view_matrix[16];
-  float normal_matrix[9];
-} Transform;
-
-static bool g_use_osd_glsl = false;
-static int g_active_uv_index = 0;
-
-static GLuint g_flat_fill_solid_program = 0;
-static GLuint g_flat_fill_texture2d_program = 0;
-static GLuint g_smooth_fill_solid_program = 0;
-static GLuint g_smooth_fill_texture2d_program = 0;
-
-static GLuint g_flat_fill_solid_shadeless_program = 0;
-static GLuint g_flat_fill_texture2d_shadeless_program = 0;
-static GLuint g_smooth_fill_solid_shadeless_program = 0;
-static GLuint g_smooth_fill_texture2d_shadeless_program = 0;
-
-static GLuint g_wireframe_program = 0;
-
-static GLuint g_lighting_ub = 0;
-static Lighting g_lighting_data;
-static Transform g_transform;
-
-namespace {
-
-GLuint compileShader(GLenum shaderType,
-                     const char *version,
-                     const char *define,
-                     const char *source)
-{
-  const char *sources[] = {
-      version,
-      define,
-#ifdef SUPPORT_COLOR_MATERIAL
-      "#define SUPPORT_COLOR_MATERIAL\n",
-#else
-      "",
-#endif
-      source,
-  };
-
-  GLuint shader = glCreateShader(shaderType);
-  glShaderSource(shader, 4, sources, NULL);
-  glCompileShader(shader);
-
-  GLint status;
-  glGetShaderiv(shader, GL_COMPILE_STATUS, &status);
-  if (status == GL_FALSE) {
-    GLchar emsg[1024];
-    glGetShaderInfoLog(shader, sizeof(emsg), 0, emsg);
-    fprintf(stderr, "Error compiling GLSL: %s\n", emsg);
-    fprintf(stderr, "Version: %s\n", version);
-    fprintf(stderr, "Defines: %s\n", define);
-    fprintf(stderr, "Source: %s\n", source);
-    return 0;
-  }
-
-  return shader;
-}
-
-GLuint linkProgram(const char *version, const char *define)
-{
-  GLuint vertexShader = compileShader(
-      GL_VERTEX_SHADER, version, define, datatoc_gpu_shader_opensubdiv_vertex_glsl);
-  if (vertexShader == 0) {
-    return 0;
-  }
-  GLuint geometryShader = compileShader(
-      GL_GEOMETRY_SHADER, version, define, datatoc_gpu_shader_opensubdiv_geometry_glsl);
-  if (geometryShader == 0) {
-    return 0;
-  }
-  GLuint fragmentShader = compileShader(
-      GL_FRAGMENT_SHADER, version, define, datatoc_gpu_shader_opensubdiv_fragment_glsl);
-  if (fragmentShader == 0) {
-    return 0;
-  }
-
-  GLuint program = glCreateProgram();
-
-  glAttachShader(program, vertexShader);
-  glAttachShader(program, geometryShader);
-  glAttachShader(program, fragmentShader);
-
-  glBindAttribLocation(program, 0, "position");
-  glBindAttribLocation(program, 1, "normal");
-
-  glLinkProgram(program);
-
-  glDeleteShader(vertexShader);
-  glDeleteShader(geometryShader);
-  glDeleteShader(fragmentShader);
-
-  GLint status;
-  glGetProgramiv(program, GL_LINK_STATUS, &status);
-  if (status == GL_FALSE) {
-    GLchar emsg[1024];
-    glGetProgramInfoLog(program, sizeof(emsg), 0, emsg);
-    fprintf(stderr, "Error linking GLSL program : %s\n", emsg);
-    fprintf(stderr, "Defines: %s\n", define);
-    glDeleteProgram(program);
-    return 0;
-  }
-
-  glUniformBlockBinding(program, glGetUniformBlockIndex(program, "Lighting"), 0);
-
-  if (GLEW_VERSION_4_1) {
-    glProgramUniform1i(program, glGetUniformLocation(program, "texture_buffer"), 0);
-    glProgramUniform1i(program, glGetUniformLocation(program, "FVarDataOffsetBuffer"), 30);
-    glProgramUniform1i(program, glGetUniformLocation(program, "FVarDataBuffer"), 31);
-  }
-  else {
-    glUseProgram(program);
-    glUniform1i(glGetUniformLocation(program, "texture_buffer"), 0);
-    glUniform1i(glGetUniformLocation(program, "FVarDataOffsetBuffer"), 30);
-    glUniform1i(glGetUniformLocation(program, "FVarDataBuffer"), 31);
-    glUseProgram(0);
-  }
-
-  return program;
-}
-
-void bindProgram(OpenSubdiv_GLMesh *gl_mesh, int program)
-{
-  glUseProgram(program);
-  // Matrices
-  glUniformMatrix4fv(
-      glGetUniformLocation(program, "modelViewMatrix"), 1, false, g_transform.model_view_matrix);
-  glUniformMatrix4fv(
-      glGetUniformLocation(program, "projectionMatrix"), 1, false, g_transform.projection_matrix);
-  glUniformMatrix3fv(
-      glGetUniformLocation(program, "normalMatrix"), 1, false, g_transform.normal_matrix);
-  // Lighting.
-  glBindBuffer(GL_UNIFORM_BUFFER, g_lighting_ub);
-  glBufferSubData(GL_UNIFORM_BUFFER, 0, sizeof(g_lighting_data), &g_lighting_data);
-  glBindBuffer(GL_UNIFORM_BUFFER, 0);
-  glBindBufferBase(GL_UNIFORM_BUFFER, 0, g_lighting_ub);
-  // Color.
-  {
-    // TODO(sergey): Stop using glGetMaterial.
-    float color[4];
-    glGetMaterialfv(GL_FRONT, GL_DIFFUSE, color);
-    glUniform4fv(glGetUniformLocation(program, "diffuse"), 1, color);
-    glGetMaterialfv(GL_FRONT, GL_SPECULAR, color);
-    glUniform4fv(glGetUniformLocation(program, "specular"), 1, color);
-    glGetMaterialfv(GL_FRONT, GL_SHININESS, color);
-    glUniform1f(glGetUniformLocation(program, "shininess"), color[0]);
-  }
-  // Face-vertex data.
-  opensubdiv_capi::GLMeshFVarData *fvar_data = gl_mesh->internal->fvar_data;
-  if (fvar_data != NULL) {
-    if (fvar_data->texture_buffer) {
-      glActiveTexture(GL_TEXTURE31);
-      glBindTexture(GL_TEXTURE_BUFFER, fvar_data->texture_buffer);
-      glActiveTexture(GL_TEXTURE0);
-    }
-    if (fvar_data->offset_buffer) {
-      glActiveTexture(GL_TEXTURE30);
-      glBindTexture(GL_TEXTURE_BUFFER, fvar_data->offset_buffer);
-      glActiveTexture(GL_TEXTURE0);
-    }
-    glUniform1i(glGetUniformLocation(program, "osd_fvar_count"), fvar_data->fvar_width);
-    if (fvar_data->channel_offsets.size() > 0 && g_active_uv_index >= 0) {
-      glUniform1i(glGetUniformLocation(program, "osd_active_uv_offset"),
-                  fvar_data->channel_offsets[g_active_uv_index]);
-    }
-    else {
-      glUniform1i(glGetUniformLocation(program, "osd_active_uv_offset"), 0);
-    }
-  }
-  else {
-    glUniform1i(glGetUniformLocation(program, "osd_fvar_count"), 0);
-    glUniform1i(glGetUniformLocation(program, "osd_active_uv_offset"), 0);
-  }
-}
-
-}  // namespace
-
-bool openSubdiv_initGLMeshDrawingResources(void)
-{
-  static bool need_init = true;
-  static bool init_success = false;
-  if (!need_init) {
-    return init_success;
-  }
-  // TODO(sergey): Update OSD drawing to OpenGL 3.3 core,
-  // then remove following line.
-  return false;
-  const char *version = "";
-  if (GLEW_VERSION_3_2) {
-    version = "#version 150 compatibility\n";
-  }
-  else if (GLEW_VERSION_3_1) {
-    version =
-        "#version 140\n"
-        "#extension GL_ARB_compatibility: enable\n";
-  }
-  else {
-    version = "#version 130\n";
-    // Minimum supported for OpenSubdiv.
-  }
-  g_flat_fill_solid_program = linkProgram(version,
-                                          "#define USE_COLOR_MATERIAL\n"
-                                          "#define USE_LIGHTING\n"
-                                          "#define FLAT_SHADING\n");
-  g_flat_fill_texture2d_program = linkProgram(version,
-                                              "#define USE_COLOR_MATERIAL\n"
-                                              "#define USE_LIGHTING\n"
-                                              "#define USE_TEXTURE_2D\n"
-                                              "#define FLAT_SHADING\n");
-  g_smooth_fill_solid_program = linkProgram(version,
-                                            "#define USE_COLOR_MATERIAL\n"
-                                            "#define USE_LIGHTING\n"
-                                            "#define SMOOTH_SHADING\n");
-  g_smooth_fill_texture2d_program = linkProgram(version,
-                                                "#define USE_COLOR_MATERIAL\n"
-                                                "#define USE_LIGHTING\n"
-                                                "#define USE_TEXTURE_2D\n"
-                                                "#define SMOOTH_SHADING\n");
-
-  g_flat_fill_solid_shadeless_program = linkProgram(version,
-                                                    "#define USE_COLOR_MATERIAL\n"
-                                                    "#define FLAT_SHADING\n");
-  g_flat_fill_texture2d_shadeless_program = linkProgram(version,
-                                                        "#define USE_COLOR_MATERIAL\n"
-                                                        "#define USE_TEXTURE_2D\n"
-                                                        "#define FLAT_SHADING\n");
-  g_smooth_fill_solid_shadeless_program = linkProgram(version,
-                                                      "#define USE_COLOR_MATERIAL\n"
-                                                      "#define SMOOTH_SHADING\n");
-  g_smooth_fill_texture2d_shadeless_program = linkProgram(version,
-                                                          "#define USE_COLOR_MATERIAL\n"
-                                                          "#define USE_TEXTURE_2D\n"
-                                                          "#define SMOOTH_SHADING\n");
-  g_wireframe_program = linkProgram(version, "#define WIREFRAME\n");
-
-  glGenBuffers(1, &g_lighting_ub);
-  glBindBuffer(GL_UNIFORM_BUFFER, g_lighting_ub);
-  glBufferData(GL_UNIFORM_BUFFER, sizeof(g_lighting_data), NULL, GL_STATIC_DRAW);
-  need_init = false;
-  init_success = g_flat_fill_solid_program != 0 && g_flat_fill_texture2d_program != 0 &&
-                 g_smooth_fill_solid_program != 0 && g_smooth_fill_texture2d_program != 0 &&
-                 g_wireframe_program;
-  return init_success;
-}
-
-void openSubdiv_deinitGLMeshDrawingResources(void)
-{
-  if (g_lighting_ub != 0) {
-    glDeleteBuffers(1, &g_lighting_ub);
-  }
-#define SAFE_DELETE_PROGRAM(program) \
-  do { \
-    if (program) { \
-      glDeleteProgram(program); \
-    } \
-  } while (false)
-
-  SAFE_DELETE_PROGRAM(g_flat_fill_solid_program);
-  SAFE_DELETE_PROGRAM(g_flat_fill_texture2d_program);
-  SAFE_DELETE_PROGRAM(g_smooth_fill_solid_program);
-  SAFE_DELETE_PROGRAM(g_smooth_fill_texture2d_program);
-  SAFE_DELETE_PROGRAM(g_flat_fill_solid_shadeless_program);
-  SAFE_DELETE_PROGRAM(g_flat_fill_texture2d_shadeless_program);
-  SAFE_DELETE_PROGRAM(g_smooth_fill_solid_shadeless_program);
-  SAFE_DELETE_PROGRAM(g_smooth_fill_texture2d_shadeless_program);
-  SAFE_DELETE_PROGRAM(g_wireframe_program);
-
-#undef SAFE_DELETE_PROGRAM
-}
-
-namespace opensubdiv_capi {
-
-namespace {
-
-GLuint prepare_patchDraw(OpenSubdiv_GLMesh *gl_mesh, bool fill_quads)
-{
-  GLint program = 0;
-  if (!g_use_osd_glsl) {
-    glGetIntegerv(GL_CURRENT_PROGRAM, &program);
-    if (program) {
-      GLint model;
-      glGetIntegerv(GL_SHADE_MODEL, &model);
-      GLint location = glGetUniformLocation(program, "osd_flat_shading");
-      if (location != -1) {
-        glUniform1i(location, model == GL_FLAT);
-      }
-      // Face-vertex data.
-      opensubdiv_capi::GLMeshFVarData *fvar_data = gl_mesh->internal->fvar_data;
-      if (fvar_data != NULL) {
-        if (fvar_data->texture_buffer) {
-          glActiveTexture(GL_TEXTURE31);
-          glBindTexture(GL_TEXTURE_BUFFER, fvar_data->texture_buffer);
-          glActiveTexture(GL_TEXTURE0);
-        }
-        if (fvar_data->offset_buffer) {
-          glActiveTexture(GL_TEXTURE30);
-          glBindTexture(GL_TEXTURE_BUFFER, fvar_data->offset_buffer);
-          glActiveTexture(GL_TEXTURE0);
-        }
-        GLint location = glGetUniformLocation(program, "osd_fvar_count");
-        if (location != -1) {
-          glUniform1i(location, fvar_data->fvar_width);
-        }
-        location = glGetUniformLocation(program, "osd_active_uv_offset");
-        if (location != -1) {
-          if (fvar_data->channel_offsets.size() > 0 && g_active_uv_index >= 0) {
-            glUniform1i(location, fvar_data->channel_offsets[g_active_uv_index]);
-          }
-          else {
-            glUniform1i(location, 0);
-          }
-        }
-      }
-      else {
-        glUniform1i(glGetUniformLocation(program, "osd_fvar_count"), 0);
-        glUniform1i(glGetUniformLocation(program, "osd_active_uv_offset"), 0);
-      }
-    }
-    return program;
-  }
-  if (fill_quads) {
-    int model;
-    GLboolean use_texture_2d;
-    glGetIntegerv(GL_SHADE_MODEL, &model);
-    glGetBooleanv(GL_TEXTURE_2D, &use_texture_2d);
-    if (model == GL_FLAT) {
-      if (use_texture_2d) {
-        program = g_flat_fill_texture2d_program;
-      }
-      else {
-        program = g_flat_fill_solid_program;
-      }
-    }
-    else {
-      if (use_texture_2d) {
-        program = g_smooth_fill_texture2d_program;
-      }
-      else {
-        program = g_smooth_fill_solid_program;
-      }
-    }
-  }
-  else {
-    glPolygonMode(GL_FRONT_AND_BACK, GL_LINE);
-    program = g_wireframe_program;
-  }
-  bindProgram(gl_mesh, program);
-  return program;
-}
-
-void perform_drawElements(GLuint program, int patch_index, int num_elements, int start_element)
-{
-  if (program) {
-    glUniform1i(glGetUniformLocation(program, "PrimitiveIdBase"), patch_index);
-  }
-  glDrawElements(GL_LINES_ADJACENCY,
-                 num_elements,
-                 GL_UNSIGNED_INT,
-                 reinterpret_cast<void *>(start_element * sizeof(unsigned int)));
-}
-
-void finishPatchDraw(bool fill_quads)
-{
-  // TODO(sergey): Some of the stuff could be done once after the whole
-  // mesh is displayed.
-  /// Restore state.
-  if (!fill_quads) {
-    glPolygonMode(GL_FRONT_AND_BACK, GL_FILL);
-  }
-  glBindVertexArray(0);
-  if (g_use_osd_glsl) {
-    // TODO(sergey): Store previously used program and roll back to it?
-    glUseProgram(0);
-  }
-}
-
-void drawPartitionPatchesRange(GLMeshInterface *mesh,
-                               GLuint program,
-                               int start_patch,
-                               int num_patches)
-{
-  int traversed_patches = 0, num_remained_patches = num_patches;
-  const OpenSubdiv::Osd::PatchArrayVector &patches = mesh->GetPatchTable()->GetPatchArrays();
-  for (int i = 0; i < patches.size(); ++i) {
-    const OpenSubdiv::Osd::PatchArray &patch = patches[i];
-    OpenSubdiv::Far::PatchDescriptor desc = patch.GetDescriptor();
-    OpenSubdiv::Far::PatchDescriptor::Type patchType = desc.GetType();
-    if (patchType == OpenSubdiv::Far::PatchDescriptor::QUADS) {
-      const int num_block_patches = patch.GetNumPatches();
-      if (start_patch >= traversed_patches &&
-          start_patch < traversed_patches + num_block_patches) {
-        const int num_control_verts = desc.GetNumControlVertices();
-        const int start_draw_patch = start_patch - traversed_patches;
-        const int num_draw_patches = min(num_remained_patches,
-                                         num_block_patches - start_draw_patch);
-        perform_drawElements(program,
-                             i + start_draw_patch,
-                             num_draw_patches * num_control_verts,
-                             patch.GetIndexBase() + start_draw_patch * num_control_verts);
-        num_remained_patches -= num_draw_patches;
-      }
-      if (num_remained_patches == 0) {
-        break;
-      }
-      traversed_patches += num_block_patches;
-    }
-  }
-}
-
-static void drawAllPatches(GLMeshInterface *mesh, GLuint program)
-{
-  const OpenSubdiv::Osd::PatchArrayVector &patches = mesh->GetPatchTable()->GetPatchArrays();
-  for (int i = 0; i < patches.size(); ++i) {
-    const OpenSubdiv::Osd::PatchArray &patch = patches[i];
-    OpenSubdiv::Far::PatchDescriptor desc = patch.GetDescriptor();
-    OpenSubdiv::Far::PatchDescriptor::Type patchType = desc.GetType();
-
-    if (patchType == OpenSubdiv::Far::PatchDescriptor::QUADS) {
-      perform_drawElements(
-          program, i, patch.GetNumPatches() * desc.GetNumControlVertices(), patch.GetIndexBase());
-    }
-  }
-}
-
-}  // namespace
-
-void GLMeshDisplayPrepare(struct OpenSubdiv_GLMesh * /*gl_mesh*/,
-                          const bool use_osd_glsl,
-                          const int active_uv_index)
-{
-  g_active_uv_index = active_uv_index;
-  g_use_osd_glsl = (use_osd_glsl != 0);
-  // Update transformation matrices.
-  glGetFloatv(GL_PROJECTION_MATRIX, g_transform.projection_matrix);
-  glGetFloatv(GL_MODELVIEW_MATRIX, g_transform.model_view_matrix);
-  copy_m3_m4((float(*)[3])g_transform.normal_matrix, (float(*)[4])g_transform.model_view_matrix);
-  invert_m3((float(*)[3])g_transform.normal_matrix);
-  transpose_m3((float(*)[3])g_transform.normal_matrix);
-  // Update OpenGL lights positions, colors etc.
-  g_lighting_data.num_enabled = 0;
-  for (int i = 0; i < MAX_LIGHTS; ++i) {
-    GLboolean enabled;
-    glGetBooleanv(GL_LIGHT0 + i, &enabled);
-    if (enabled) {
-      g_lighting_data.num_enabled++;
-    }
-    // TODO(sergey): Stop using glGetLight.
-    glGetLightfv(GL_LIGHT0 + i, GL_POSITION, g_lighting_data.lights[i].position);
-    glGetLightfv(GL_LIGHT0 + i, GL_AMBIENT, g_lighting_data.lights[i].ambient);
-    glGetLightfv(GL_LIGHT0 + i, GL_DIFFUSE, g_lighting_data.lights[i].diffuse);
-    glGetLightfv(GL_LIGHT0 + i, GL_SPECULAR, g_lighting_data.lights[i].specular);
-    glGetLightfv(GL_LIGHT0 + i, GL_SPOT_DIRECTION, g_lighting_data.lights[i].spot_direction);
-#ifdef SUPPORT_COLOR_MATERIAL
-    glGetLightfv(
-        GL_LIGHT0 + i, GL_CONSTANT_ATTENUATION, &g_lighting_data.lights[i].constant_attenuation);
-    glGetLightfv(
-        GL_LIGHT0 + i, GL_LINEAR_ATTENUATION, &g_lighting_data.lights[i].linear_attenuation);
-    glGetLightfv(
-        GL_LIGHT0 + i, GL_QUADRATIC_ATTENUATION, &g_lighting_data.lights[i].quadratic_attenuation);
-    glGetLightfv(GL_LIGHT0 + i, GL_SPOT_CUTOFF, &g_lighting_data.lights[i].spot_cutoff);
-    glGetLightfv(GL_LIGHT0 + i, GL_SPOT_EXPONENT, &g_lighting_data.lights[i].spot_exponent);
-    g_lighting_data.lights[i].spot_cos_cutoff = cos(g_lighting_data.lights[i].spot_cutoff);
-#endif
-  }
-}
-
-void GLMeshDisplayDrawPatches(OpenSubdiv_GLMesh *gl_mesh,
-                              const bool fill_quads,
-                              const int start_patch,
-                              const int num_patches)
-{
-  GLMeshInterface *mesh = gl_mesh->internal->mesh_interface;
-  // Make sure all global invariants are initialized.
-  if (!openSubdiv_initGLMeshDrawingResources()) {
-    return;
-  }
-  /// Setup GLSL/OpenGL to draw patches in current context.
-  GLuint program = prepare_patchDraw(gl_mesh, fill_quads != 0);
-  if (start_patch != -1) {
-    drawPartitionPatchesRange(mesh, program, start_patch, num_patches);
-  }
-  else {
-    drawAllPatches(mesh, program);
-  }
-  // Finish patch drawing by restoring all changes to the OpenGL context.
-  finishPatchDraw(fill_quads != 0);
-}
-
-}  // namespace opensubdiv_capi
diff --git a/intern/opensubdiv/internal/opensubdiv_gl_mesh_draw.h b/intern/opensubdiv/internal/opensubdiv_gl_mesh_draw.h
deleted file mode 100644
index 599ab9550e7..00000000000
--- a/intern/opensubdiv/internal/opensubdiv_gl_mesh_draw.h
+++ /dev/null
@@ -1,39 +0,0 @@
-// Copyright 2013 Blender Foundation. All rights reserved.
-//
-// This program is free software; you can redistribute it and/or
-// modify it under the terms of the GNU General Public License
-// as published by the Free Software Foundation; either version 2
-// of the License, or (at your option) any later version.
-//
-// This program is distributed in the hope that it will be useful,
-// but WITHOUT ANY WARRANTY; without even the implied warranty of
-// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-// GNU General Public License for more details.
-//
-// You should have received a copy of the GNU General Public License
-// along with this program; if not, write to the Free Software Foundation,
-// Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
-//
-// Author: Sergey Sharybin
-
-#ifndef OPENSUBDIV_GL_MESH_DRAW_H_
-#define OPENSUBDIV_GL_MESH_DRAW_H_
-
-#include <stdint.h>  // for bool
-
-struct OpenSubdiv_GLMesh;
-
-namespace opensubdiv_capi {
-
-void GLMeshDisplayPrepare(struct OpenSubdiv_GLMesh *gl_mesh,
-                          const bool use_osd_glsl,
-                          const int active_uv_index);
-
-void GLMeshDisplayDrawPatches(OpenSubdiv_GLMesh *gl_mesh,
-                              const bool fill_quads,
-                              const int start_patch,
-                              const int num_patches);
-
-}  // namespace opensubdiv_capi
-
-#endif  // OPENSUBDIV_GL_MESH_DRAW_H_
diff --git a/intern/opensubdiv/internal/opensubdiv_gl_mesh_fvar.cc b/intern/opensubdiv/internal/opensubdiv_gl_mesh_fvar.cc
deleted file mode 100644
index 6efbe93d2d8..00000000000
--- a/intern/opensubdiv/internal/opensubdiv_gl_mesh_fvar.cc
+++ /dev/null
@@ -1,171 +0,0 @@
-// Copyright 2013 Blender Foundation. All rights reserved.
-//
-// This program is free software; you can redistribute it and/or
-// modify it under the terms of the GNU General Public License
-// as published by the Free Software Foundation; either version 2
-// of the License, or (at your option) any later version.
-//
-// This program is distributed in the hope that it will be useful,
-// but WITHOUT ANY WARRANTY; without even the implied warranty of
-// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-// GNU General Public License for more details.
-//
-// You should have received a copy of the GNU General Public License
-// along with this program; if not, write to the Free Software Foundation,
-// Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
-//
-// Author: Sergey Sharybin
-
-#include "internal/opensubdiv_gl_mesh_fvar.h"
-
-#include <GL/glew.h>
-#include <opensubdiv/far/primvarRefiner.h>
-
-#include "internal/opensubdiv_util.h"
-
-namespace opensubdiv_capi {
-
-////////////////////////////////////////////////////////////////////////////////
-// GLMeshFVarData
-
-GLMeshFVarData::GLMeshFVarData() : texture_buffer(0), offset_buffer(0)
-{
-}
-
-GLMeshFVarData::~GLMeshFVarData()
-{
-  release();
-}
-
-void GLMeshFVarData::release()
-{
-  if (texture_buffer) {
-    glDeleteTextures(1, &texture_buffer);
-  }
-  if (offset_buffer) {
-    glDeleteTextures(1, &offset_buffer);
-  }
-  texture_buffer = 0;
-  offset_buffer = 0;
-  fvar_width = 0;
-  channel_offsets.clear();
-}
-
-void GLMeshFVarData::create(const OpenSubdiv::Far::TopologyRefiner *topology_refiner,
-                            const OpenSubdiv::Far::PatchTable *patch_table,
-                            int fvar_width,
-                            const float *fvar_src_data)
-{
-  release();
-  this->fvar_width = fvar_width;
-  /// Expand fvar data to per-patch array.
-  const int max_level = topology_refiner->GetMaxLevel();
-  const int num_channels = patch_table->GetNumFVarChannels();
-  vector<float> data;
-  int fvar_data_offset = 0;
-  channel_offsets.resize(num_channels);
-  for (int channel = 0; channel < num_channels; ++channel) {
-    OpenSubdiv::Far::ConstIndexArray indices = patch_table->GetFVarValues(channel);
-    channel_offsets[channel] = data.size();
-    data.reserve(data.size() + indices.size() * fvar_width);
-    for (int fvert = 0; fvert < indices.size(); ++fvert) {
-      int index = indices[fvert] * fvar_width;
-      for (int i = 0; i < fvar_width; ++i) {
-        data.push_back(fvar_src_data[fvar_data_offset + index++]);
-      }
-    }
-    if (topology_refiner->IsUniform()) {
-      const int num_values_max = topology_refiner->GetLevel(max_level).GetNumFVarValues(channel);
-      fvar_data_offset += num_values_max * fvar_width;
-    }
-    else {
-      const int num_values_total = topology_refiner->GetNumFVarValuesTotal(channel);
-      fvar_data_offset += num_values_total * fvar_width;
-    }
-  }
-  GLuint buffer;
-  glGenBuffers(1, &buffer);
-  glBindBuffer(GL_ARRAY_BUFFER, buffer);
-  glBufferData(GL_ARRAY_BUFFER, data.size() * sizeof(float), &data[0], GL_STATIC_DRAW);
-  glGenTextures(1, &texture_buffer);
-  glBindTexture(GL_TEXTURE_BUFFER, texture_buffer);
-  glTexBuffer(GL_TEXTURE_BUFFER, GL_R32F, buffer);
-  glDeleteBuffers(1, &buffer);
-  glGenBuffers(1, &buffer);
-  glBindBuffer(GL_ARRAY_BUFFER, buffer);
-  glBufferData(
-      GL_ARRAY_BUFFER, channel_offsets.size() * sizeof(int), &channel_offsets[0], GL_STATIC_DRAW);
-  glGenTextures(1, &offset_buffer);
-  glBindTexture(GL_TEXTURE_BUFFER, offset_buffer);
-  glTexBuffer(GL_TEXTURE_BUFFER, GL_R32I, buffer);
-  glBindTexture(GL_TEXTURE_BUFFER, 0);
-  glBindBuffer(GL_ARRAY_BUFFER, 0);
-}
-
-////////////////////////////////////////////////////////////////////////////////
-// Helper functions.
-
-struct FVarVertex {
-  float u, v;
-
-  void Clear()
-  {
-    u = v = 0.0f;
-  }
-
-  void AddWithWeight(FVarVertex const &src, float weight)
-  {
-    u += weight * src.u;
-    v += weight * src.v;
-  }
-};
-
-void interpolateFVarData(const OpenSubdiv::Far::TopologyRefiner &refiner,
-                         const vector<float> &uvs,
-                         vector<float> *fvar_data)
-{
-  const int fvar_width = 2;
-  const int max_level = refiner.GetMaxLevel();
-  size_t fvar_data_offset = 0, values_offset = 0;
-  for (int channel = 0; channel < refiner.GetNumFVarChannels(); ++channel) {
-    const int num_values = refiner.GetLevel(0).GetNumFVarValues(channel) * 2;
-    const int num_values_max = refiner.GetLevel(max_level).GetNumFVarValues(channel);
-    const int num_values_total = refiner.GetNumFVarValuesTotal(channel);
-    if (num_values_total <= 0) {
-      continue;
-    }
-    OpenSubdiv::Far::PrimvarRefiner primvar_refiner(refiner);
-    if (refiner.IsUniform()) {
-      // For uniform we only keep the highest level of refinement.
-      fvar_data->resize(fvar_data->size() + num_values_max * fvar_width);
-      vector<FVarVertex> buffer(num_values_total - num_values_max);
-      FVarVertex *src = &buffer[0];
-      memcpy(src, &uvs[values_offset], num_values * sizeof(float));
-      // Defer the last level to treat separately with its alternate
-      // destination.
-      for (int level = 1; level < max_level; ++level) {
-        FVarVertex *dst = src + refiner.GetLevel(level - 1).GetNumFVarValues(channel);
-        primvar_refiner.InterpolateFaceVarying(level, src, dst, channel);
-        src = dst;
-      }
-      FVarVertex *dst = reinterpret_cast<FVarVertex *>(&(*fvar_data)[fvar_data_offset]);
-      primvar_refiner.InterpolateFaceVarying(max_level, src, dst, channel);
-      fvar_data_offset += num_values_max * fvar_width;
-    }
-    else {
-      // For adaptive we keep all levels.
-      fvar_data->resize(fvar_data->size() + num_values_total * fvar_width);
-      FVarVertex *src = reinterpret_cast<FVarVertex *>(&(*fvar_data)[fvar_data_offset]);
-      memcpy(src, &uvs[values_offset], num_values * sizeof(float));
-      for (int level = 1; level <= max_level; ++level) {
-        FVarVertex *dst = src + refiner.GetLevel(level - 1).GetNumFVarValues(channel);
-        primvar_refiner.InterpolateFaceVarying(level, src, dst, channel);
-        src = dst;
-      }
-      fvar_data_offset += num_values_total * fvar_width;
-    }
-    values_offset += num_values;
-  }
-}
-
-}  // namespace opensubdiv_capi
diff --git a/intern/opensubdiv/internal/opensubdiv_gl_mesh_fvar.h b/intern/opensubdiv/internal/opensubdiv_gl_mesh_fvar.h
deleted file mode 100644
index 73a1af05605..00000000000
--- a/intern/opensubdiv/internal/opensubdiv_gl_mesh_fvar.h
+++ /dev/null
@@ -1,57 +0,0 @@
-// Copyright 2013 Blender Foundation. All rights reserved.
-//
-// This program is free software; you can redistribute it and/or
-// modify it under the terms of the GNU General Public License
-// as published by the Free Software Foundation; either version 2
-// of the License, or (at your option) any later version.
-//
-// This program is distributed in the hope that it will be useful,
-// but WITHOUT ANY WARRANTY; without even the implied warranty of
-// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-// GNU General Public License for more details.
-//
-// You should have received a copy of the GNU General Public License
-// along with this program; if not, write to the Free Software Foundation,
-// Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
-//
-// Author: Sergey Sharybin
-
-#ifndef OPENSUBDIV_GL_MESH_FVAR_H_
-#define OPENSUBDIV_GL_MESH_FVAR_H_
-
-// NOTE: This is a [sane(er)] port of previous ground work for getting UVs to
-// work. Still needs a lot of work to make it easy, correct and have proper
-// data ownership.
-
-#include <opensubdiv/far/patchTable.h>
-#include <opensubdiv/far/topologyRefiner.h>
-
-#include "internal/opensubdiv_util.h"
-
-namespace opensubdiv_capi {
-
-// The buffer which holds GPU resources for face-varying elements.
-class GLMeshFVarData {
- public:
-  GLMeshFVarData();
-  ~GLMeshFVarData();
-
-  void release();
-  void create(const OpenSubdiv::Far::TopologyRefiner *refiner,
-              const OpenSubdiv::Far::PatchTable *patch_table,
-              int fvar_width,
-              const float *fvar_src_data);
-
-  unsigned int texture_buffer;
-  unsigned int offset_buffer;
-  vector<int> channel_offsets;
-  int fvar_width;
-};
-
-void interpolateFVarData(const OpenSubdiv::Far::TopologyRefiner &refiner,
-                         const vector<float> &uvs,
-                         vector<float> *fvar_data);
-
-}  // namespace opensubdiv_capi
-
-#endif  // OPENSUBDIV_GL_MESH_FVAR_H_
diff --git a/intern/opensubdiv/internal/opensubdiv_gl_mesh_internal.cc b/intern/opensubdiv/internal/opensubdiv_gl_mesh_internal.cc
deleted file mode 100644
index 57e56bad3fb..00000000000
--- a/intern/opensubdiv/internal/opensubdiv_gl_mesh_internal.cc
+++ /dev/null
@@ -1,32 +0,0 @@
-// Copyright 2018 Blender Foundation. All rights reserved.
-//
-// This program is free software; you can redistribute it and/or
-// modify it under the terms of the GNU General Public License
-// as published by the Free Software Foundation; either version 2
-// of the License, or (at your option) any later version.
-//
-// This program is distributed in the hope that it will be useful,
-// but WITHOUT ANY WARRANTY; without even the implied warranty of
-// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-// GNU General Public License for more details.
-//
-// You should have received a copy of the GNU General Public License
-// along with this program; if not, write to the Free Software Foundation,
-// Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
-//
-// Author: Sergey Sharybin
-
-#include "internal/opensubdiv_gl_mesh_internal.h"
-
-#include "internal/opensubdiv_gl_mesh_fvar.h"
-
-OpenSubdiv_GLMeshInternal::OpenSubdiv_GLMeshInternal()
-    : evaluator_type(OPENSUBDIV_EVALUATOR_CPU), mesh_interface(NULL), fvar_data(NULL)
-{
-}
-
-OpenSubdiv_GLMeshInternal::~OpenSubdiv_GLMeshInternal()
-{
-  delete mesh_interface;
-  delete fvar_data;
-}
diff --git a/intern/opensubdiv/internal/opensubdiv_gl_mesh_internal.h b/intern/opensubdiv/internal/opensubdiv_gl_mesh_internal.h
deleted file mode 100644
index cb92fb18362..00000000000
--- a/intern/opensubdiv/internal/opensubdiv_gl_mesh_internal.h
+++ /dev/null
@@ -1,43 +0,0 @@
-// Copyright 2018 Blender Foundation. All rights reserved.
-//
-// This program is free software; you can redistribute it and/or
-// modify it under the terms of the GNU General Public License
-// as published by the Free Software Foundation; either version 2
-// of the License, or (at your option) any later version.
-//
-// This program is distributed in the hope that it will be useful,
-// but WITHOUT ANY WARRANTY; without even the implied warranty of
-// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-// GNU General Public License for more details.
-//
-// You should have received a copy of the GNU General Public License
-// along with this program; if not, write to the Free Software Foundation,
-// Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
-//
-// Author: Sergey Sharybin
-
-#ifndef OPENSUBDIV_GL_MESH_INTERNAL_H_
-#define OPENSUBDIV_GL_MESH_INTERNAL_H_
-
-#ifdef _MSC_VER
-#  include <iso646.h>
-#endif
-
-#include <opensubdiv/osd/glMesh.h>
-
-#include "opensubdiv_capi_type.h"
-
-namespace opensubdiv_capi {
-class GLMeshFVarData;
-}  // namespace opensubdiv_capi
-
-typedef struct OpenSubdiv_GLMeshInternal {
-  OpenSubdiv_GLMeshInternal();
-  ~OpenSubdiv_GLMeshInternal();
-
-  eOpenSubdivEvaluator evaluator_type;
-  OpenSubdiv::Osd::GLMeshInterface *mesh_interface;
-  opensubdiv_capi::GLMeshFVarData *fvar_data;
-} OpenSubdiv_GLMeshInternal;
-
-#endif  // OPENSUBDIV_GL_MESH_INTERNAL_H_
diff --git a/intern/opensubdiv/internal/opensubdiv_topology_refiner.cc b/intern/opensubdiv/internal/opensubdiv_topology_refiner.cc
index ac27cbdefdc..6e2dae4533a 100644
--- a/intern/opensubdiv/internal/opensubdiv_topology_refiner.cc
+++ b/intern/opensubdiv/internal/opensubdiv_topology_refiner.cc
@@ -28,7 +28,7 @@
 #include "internal/opensubdiv_topology_refiner_internal.h"
 #include "internal/opensubdiv_util.h"
 
-using opensubdiv_capi::vector;
+using blender::opensubdiv::vector;
 
 namespace {
 
@@ -182,7 +182,7 @@ int getNumFVarChannels(const struct OpenSubdiv_TopologyRefiner *topology_refiner
 OpenSubdiv_FVarLinearInterpolation getFVarLinearInterpolation(
     const struct OpenSubdiv_TopologyRefiner *topology_refiner)
 {
-  return opensubdiv_capi::getCAPIFVarLinearInterpolationFromOSD(
+  return blender::opensubdiv::getCAPIFVarLinearInterpolationFromOSD(
       getOSDTopologyRefiner(topology_refiner)->GetFVarLinearInterpolation());
 }
 
@@ -243,7 +243,7 @@ OpenSubdiv_TopologyRefiner *openSubdiv_createTopologyRefinerFromConverter(
     OpenSubdiv_Converter *converter, const OpenSubdiv_TopologyRefinerSettings *settings)
 {
   OpenSubdiv::Far::TopologyRefiner *osd_topology_refiner =
-      opensubdiv_capi::createOSDTopologyRefinerFromConverter(converter);
+      blender::opensubdiv::createOSDTopologyRefinerFromConverter(converter);
   if (osd_topology_refiner == NULL) {
     // Happens on empty or bad topology.
     return NULL;
@@ -265,7 +265,8 @@ void openSubdiv_deleteTopologyRefiner(OpenSubdiv_TopologyRefiner *topology_refin
 ////////////////////////////////////////////////////////////////////////////////
 // Comparison with converter.
 
-namespace opensubdiv_capi {
+namespace blender {
+namespace opensubdiv {
 namespace {
 
 ///////////////////////////////////////////////////////////
@@ -274,8 +275,8 @@ namespace {
 bool checkSchemeTypeMatches(const OpenSubdiv::Far::TopologyRefiner *topology_refiner,
                             const OpenSubdiv_Converter *converter)
 {
-  const OpenSubdiv::Sdc::SchemeType converter_scheme_type = opensubdiv_capi::getSchemeTypeFromCAPI(
-      converter->getSchemeType(converter));
+  const OpenSubdiv::Sdc::SchemeType converter_scheme_type =
+      blender::opensubdiv::getSchemeTypeFromCAPI(converter->getSchemeType(converter));
   return (converter_scheme_type == topology_refiner->GetSchemeType());
 }
 
@@ -286,7 +287,7 @@ bool checkOptionsMatches(const OpenSubdiv::Far::TopologyRefiner *topology_refine
   const Options options = topology_refiner->GetSchemeOptions();
   const Options::FVarLinearInterpolation fvar_interpolation = options.GetFVarLinearInterpolation();
   const Options::FVarLinearInterpolation converter_fvar_interpolation =
-      opensubdiv_capi::getFVarLinearInterpolationFromCAPI(
+      blender::opensubdiv::getFVarLinearInterpolationFromCAPI(
           converter->getFVarLinearInterpolation(converter));
   if (fvar_interpolation != converter_fvar_interpolation) {
     return false;
@@ -660,13 +661,14 @@ bool checkTopologyAttributesMatch(const OpenSubdiv::Far::TopologyRefiner *topolo
 }
 
 }  // namespace
-}  // namespace opensubdiv_capi
+}  // namespace opensubdiv
+}  // namespace blender
 
 bool openSubdiv_topologyRefinerCompareWithConverter(
     const OpenSubdiv_TopologyRefiner *topology_refiner, const OpenSubdiv_Converter *converter)
 {
   const OpenSubdiv::Far::TopologyRefiner *refiner = getOSDTopologyRefiner(topology_refiner);
-  return (opensubdiv_capi::checkPreliminaryMatches(refiner, converter) &&
-          opensubdiv_capi::checkGeometryMatches(refiner, converter) &&
-          opensubdiv_capi::checkTopologyAttributesMatch(refiner, converter));
+  return (blender::opensubdiv::checkPreliminaryMatches(refiner, converter) &&
+          blender::opensubdiv::checkGeometryMatches(refiner, converter) &&
+          blender::opensubdiv::checkTopologyAttributesMatch(refiner, converter));
 }
diff --git a/intern/opensubdiv/internal/opensubdiv_util.cc b/intern/opensubdiv/internal/opensubdiv_util.cc
index 6e6f3a0920f..ea61b21e5d0 100644
--- a/intern/opensubdiv/internal/opensubdiv_util.cc
+++ b/intern/opensubdiv/internal/opensubdiv_util.cc
@@ -23,7 +23,8 @@
 #  include <iso646.h>
 #endif
 
-namespace opensubdiv_capi {
+namespace blender {
+namespace opensubdiv {
 
 void stringSplit(vector<string> *tokens,
                  const string &str,
@@ -56,4 +57,5 @@ void stringSplit(vector<string> *tokens,
   }
 }
 
-}  // namespace opensubdiv_capi
+}  // namespace opensubdiv
+}  // namespace blender
diff --git a/intern/opensubdiv/internal/opensubdiv_util.h b/intern/opensubdiv/internal/opensubdiv_util.h
index e515859b42f..3fcfdd4c32b 100644
--- a/intern/opensubdiv/internal/opensubdiv_util.h
+++ b/intern/opensubdiv/internal/opensubdiv_util.h
@@ -27,7 +27,8 @@
 #include <utility>
 #include <vector>
 
-namespace opensubdiv_capi {
+namespace blender {
+namespace opensubdiv {
 
 using std::fill;
 using std::make_pair;
@@ -51,6 +52,7 @@ void stringSplit(vector<string> *tokens,
                  const string &separators,
                  bool skip_empty);
 
-}  // namespace opensubdiv_capi
+}  // namespace opensubdiv
+}  // namespace blender
 
 #endif  // OPENSUBDIV_UTIL_H_
diff --git a/intern/opensubdiv/opensubdiv_gl_mesh_capi.h b/intern/opensubdiv/opensubdiv_gl_mesh_capi.h
deleted file mode 100644
index f7dd6f83434..00000000000
--- a/intern/opensubdiv/opensubdiv_gl_mesh_capi.h
+++ /dev/null
@@ -1,92 +0,0 @@
-// Copyright 2013 Blender Foundation. All rights reserved.
-//
-// This program is free software; you can redistribute it and/or
-// modify it under the terms of the GNU General Public License
-// as published by the Free Software Foundation; either version 2
-// of the License, or (at your option) any later version.
-//
-// This program is distributed in the hope that it will be useful,
-// but WITHOUT ANY WARRANTY; without even the implied warranty of
-// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-// GNU General Public License for more details.
-//
-// You should have received a copy of the GNU General Public License
-// along with this program; if not, write to the Free Software Foundation,
-// Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
-//
-// Author: Sergey Sharybin
-
-#ifndef OPENSUBDIV_CAPI_GL_MESH_CAPI_H_
-#define OPENSUBDIV_CAPI_GL_MESH_CAPI_H_
-
-#include <stdint.h>  // for bool
-
-#include "opensubdiv_capi_type.h"
-
-#ifdef __cplusplus
-extern "C" {
-#endif
-
-struct OpenSubdiv_GLMeshInternal;
-
-// Mesh which is displayable in OpenGL context.
-typedef struct OpenSubdiv_GLMesh {
-  //////////////////////////////////////////////////////////////////////////////
-  // Subdivision/topology part.
-
-  // Returns the GL index buffer containing the patch control vertices.
-  unsigned int (*getPatchIndexBuffer)(struct OpenSubdiv_GLMesh *gl_mesh);
-
-  // Bind GL buffer which contains vertices (VBO).
-  // TODO(sergey): Is this a coarse vertices?
-  void (*bindVertexBuffer)(struct OpenSubdiv_GLMesh *gl_mesh);
-
-  // Set coarse positions from a continuous array of coordinates.
-  void (*setCoarsePositions)(struct OpenSubdiv_GLMesh *gl_mesh,
-                             const float *positions,
-                             const int start_vertex,
-                             const int num_vertices);
-  // TODO(sergey): setCoarsePositionsFromBuffer().
-
-  // Refine after coarse positions update.
-  void (*refine)(struct OpenSubdiv_GLMesh *gl_mesh);
-
-  // Synchronize after coarse positions update and refine.
-  void (*synchronize)(struct OpenSubdiv_GLMesh *gl_mesh);
-
-  //////////////////////////////////////////////////////////////////////////////
-  // Drawing part.
-
-  // Prepare mesh for display.
-  void (*prepareDraw)(struct OpenSubdiv_GLMesh *gl_mesh,
-                      const bool use_osd_glsl,
-                      const int active_uv_index);
-
-  // Draw given range of patches.
-  //
-  // If fill_quads is false, then patches are drawn in wireframe.
-  void (*drawPatches)(struct OpenSubdiv_GLMesh *gl_mesh,
-                      const bool fill_quads,
-                      const int start_patch,
-                      const int num_patches);
-
-  // Internal storage for the use in this module only.
-  //
-  // Tease: This contains an actual OpenSubdiv's Mesh object.
-  struct OpenSubdiv_GLMeshInternal *internal;
-} OpenSubdiv_GLMesh;
-
-OpenSubdiv_GLMesh *openSubdiv_createOsdGLMeshFromTopologyRefiner(
-    struct OpenSubdiv_TopologyRefiner *topology_refiner, eOpenSubdivEvaluator evaluator_type);
-
-void openSubdiv_deleteOsdGLMesh(OpenSubdiv_GLMesh *gl_mesh);
-
-// Global resources needed for GL mesh drawing.
-bool openSubdiv_initGLMeshDrawingResources(void);
-void openSubdiv_deinitGLMeshDrawingResources(void);
-
-#ifdef __cplusplus
-}
-#endif
-
-#endif  // OPENSUBDIV_CAPI_GL_MESH_CAPI_H_
diff --git a/intern/opensubdiv/shader/gpu_shader_opensubdiv_fragment.glsl b/intern/opensubdiv/shader/gpu_shader_opensubdiv_fragment.glsl
deleted file mode 100644
index 7f08182d78a..00000000000
--- a/intern/opensubdiv/shader/gpu_shader_opensubdiv_fragment.glsl
+++ /dev/null
@@ -1,163 +0,0 @@
-/*
- * This program is free software; you can redistribute it and/or
- * modify it under the terms of the GNU General Public License
- * as published by the Free Software Foundation; either version 2
- * of the License, or (at your option) any later version.
- *
- * This program is distributed in the hope that it will be useful,
- * but WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
- * GNU General Public License for more details.
- *
- * You should have received a copy of the GNU General Public License
- * along with this program; if not, write to the Free Software Foundation,
- * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
- *
- * The Original Code is Copyright (C) 2014 Blender Foundation.
- * All rights reserved.
- */
-
-struct VertexData {
-  vec4 position;
-  vec3 normal;
-  vec2 uv;
-};
-
-#define MAX_LIGHTS 8
-#define NUM_SOLID_LIGHTS 3
-
-struct LightSource {
-  vec4 position;
-  vec4 ambient;
-  vec4 diffuse;
-  vec4 specular;
-  vec4 spotDirection;
-#ifdef SUPPORT_COLOR_MATERIAL
-  float constantAttenuation;
-  float linearAttenuation;
-  float quadraticAttenuation;
-  float spotCutoff;
-  float spotExponent;
-  float spotCosCutoff;
-  float pad, pad2;
-#endif
-};
-
-layout(std140) uniform Lighting
-{
-  LightSource lightSource[MAX_LIGHTS];
-  int num_enabled_lights;
-};
-
-uniform vec4 diffuse;
-uniform vec4 specular;
-uniform float shininess;
-
-uniform sampler2D texture_buffer;
-
-in block
-{
-  VertexData v;
-}
-inpt;
-
-void main()
-{
-#ifdef WIREFRAME
-  gl_FragColor = diffuse;
-#else
-  vec3 N = inpt.v.normal;
-
-  if (!gl_FrontFacing)
-    N = -N;
-
-  /* Compute diffuse and specular lighting. */
-  vec3 L_diffuse = vec3(0.0);
-  vec3 L_specular = vec3(0.0);
-
-#  ifdef USE_LIGHTING
-#    ifndef USE_COLOR_MATERIAL
-  /* Assume NUM_SOLID_LIGHTS directional lights. */
-  for (int i = 0; i < NUM_SOLID_LIGHTS; i++) {
-    vec4 Plight = lightSource[i].position;
-#      ifdef USE_DIRECTIONAL_LIGHT
-    vec3 l = (Plight.w == 0.0) ? normalize(Plight.xyz) : normalize(inpt.v.position.xyz);
-#      else  /* USE_DIRECTIONAL_LIGHT */
-    /* TODO(sergey): We can normalize it outside of the shader. */
-    vec3 l = normalize(Plight.xyz);
-#      endif /* USE_DIRECTIONAL_LIGHT */
-    vec3 h = normalize(l + vec3(0, 0, 1));
-    float d = max(0.0, dot(N, l));
-    float s = pow(max(0.0, dot(N, h)), shininess);
-    L_diffuse += d * lightSource[i].diffuse.rgb;
-    L_specular += s * lightSource[i].specular.rgb;
-  }
-#    else    /* USE_COLOR_MATERIAL */
-  vec3 varying_position = inpt.v.position.xyz;
-  vec3 V = (gl_ProjectionMatrix[3][3] == 0.0) ? normalize(varying_position) : vec3(0.0, 0.0, -1.0);
-  for (int i = 0; i < num_enabled_lights; i++) {
-    /* todo: this is a slow check for disabled lights */
-    if (lightSource[i].specular.a == 0.0)
-      continue;
-
-    float intensity = 1.0;
-    vec3 light_direction;
-
-    if (lightSource[i].position.w == 0.0) {
-      /* directional light */
-      light_direction = lightSource[i].position.xyz;
-    }
-    else {
-      /* point light */
-      vec3 d = lightSource[i].position.xyz - varying_position;
-      light_direction = normalize(d);
-
-      /* spot light cone */
-      if (lightSource[i].spotCutoff < 90.0) {
-        float cosine = max(dot(light_direction, -lightSource[i].spotDirection.xyz), 0.0);
-        intensity = pow(cosine, lightSource[i].spotExponent);
-        intensity *= step(lightSource[i].spotCosCutoff, cosine);
-      }
-
-      /* falloff */
-      float distance = length(d);
-
-      intensity /= lightSource[i].constantAttenuation +
-                   lightSource[i].linearAttenuation * distance +
-                   lightSource[i].quadraticAttenuation * distance * distance;
-    }
-
-    /* diffuse light */
-    vec3 light_diffuse = lightSource[i].diffuse.rgb;
-    float diffuse_bsdf = max(dot(N, light_direction), 0.0);
-    L_diffuse += light_diffuse * diffuse_bsdf * intensity;
-
-    /* specular light */
-    vec3 light_specular = lightSource[i].specular.rgb;
-    vec3 H = normalize(light_direction - V);
-
-    float specular_bsdf = pow(max(dot(N, H), 0.0), gl_FrontMaterial.shininess);
-    L_specular += light_specular * specular_bsdf * intensity;
-  }
-#    endif   /* USE_COLOR_MATERIAL */
-#  else      /* USE_LIGHTING */
-  L_diffuse = vec3(1.0);
-#  endif
-
-  /* Compute diffuse color. */
-#  ifdef USE_TEXTURE_2D
-  L_diffuse *= texture2D(texture_buffer, inpt.v.uv).rgb;
-#  else
-  L_diffuse *= diffuse.rgb;
-#  endif
-
-  /* Sum lighting. */
-  vec3 L = L_diffuse;
-  if (shininess != 0) {
-    L += L_specular * specular.rgb;
-  }
-
-  /* Write out fragment color. */
-  gl_FragColor = vec4(L, diffuse.a);
-#endif
-}
diff --git a/intern/opensubdiv/shader/gpu_shader_opensubdiv_geometry.glsl b/intern/opensubdiv/shader/gpu_shader_opensubdiv_geometry.glsl
deleted file mode 100644
index 37bc0720113..00000000000
--- a/intern/opensubdiv/shader/gpu_shader_opensubdiv_geometry.glsl
+++ /dev/null
@@ -1,149 +0,0 @@
-/*
- * This program is free software; you can redistribute it and/or
- * modify it under the terms of the GNU General Public License
- * as published by the Free Software Foundation; either version 2
- * of the License, or (at your option) any later version.
- *
- * This program is distributed in the hope that it will be useful,
- * but WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
- * GNU General Public License for more details.
- *
- * You should have received a copy of the GNU General Public License
- * along with this program; if not, write to the Free Software Foundation,
- * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
- *
- * The Original Code is Copyright (C) 2014 Blender Foundation.
- * All rights reserved.
- */
-
-struct VertexData {
-  vec4 position;
-  vec3 normal;
-  vec2 uv;
-};
-
-layout(lines_adjacency) in;
-#ifdef WIREFRAME
-layout(line_strip, max_vertices = 8) out;
-#else
-layout(triangle_strip, max_vertices = 4) out;
-#endif
-
-uniform mat4 modelViewMatrix;
-uniform mat4 projectionMatrix;
-uniform int PrimitiveIdBase;
-uniform int osd_fvar_count;
-uniform int osd_active_uv_offset;
-
-in block
-{
-  VertexData v;
-}
-inpt[];
-
-#define INTERP_FACE_VARYING_2(result, fvarOffset, tessCoord) \
-  { \
-    vec2 v[4]; \
-    int primOffset = (gl_PrimitiveID + PrimitiveIdBase) * 4; \
-    for (int i = 0; i < 4; ++i) { \
-      int index = (primOffset + i) * osd_fvar_count + fvarOffset; \
-      v[i] = vec2(texelFetch(FVarDataBuffer, index).s, texelFetch(FVarDataBuffer, index + 1).s); \
-    } \
-    result = mix(mix(v[0], v[1], tessCoord.s), mix(v[3], v[2], tessCoord.s), tessCoord.t); \
-  }
-
-uniform samplerBuffer FVarDataBuffer;
-uniform isamplerBuffer FVarDataOffsetBuffer;
-
-out block
-{
-  VertexData v;
-}
-outpt;
-
-#ifdef FLAT_SHADING
-void emit(int index, vec3 normal)
-{
-  outpt.v.position = inpt[index].v.position;
-  outpt.v.normal = normal;
-
-  /* TODO(sergey): Only uniform subdivisions atm. */
-  vec2 quadst[4] = vec2[](vec2(0, 0), vec2(1, 0), vec2(1, 1), vec2(0, 1));
-  vec2 st = quadst[index];
-
-  INTERP_FACE_VARYING_2(outpt.v.uv, osd_active_uv_offset, st);
-
-  gl_Position = projectionMatrix * inpt[index].v.position;
-  EmitVertex();
-}
-
-#  ifdef WIREFRAME
-void emit_edge(int v0, int v1, vec3 normal)
-{
-  emit(v0, normal);
-  emit(v1, normal);
-}
-#  endif
-
-#else
-void emit(int index)
-{
-  outpt.v.position = inpt[index].v.position;
-  outpt.v.normal = inpt[index].v.normal;
-
-  /* TODO(sergey): Only uniform subdivisions atm. */
-  vec2 quadst[4] = vec2[](vec2(0, 0), vec2(1, 0), vec2(1, 1), vec2(0, 1));
-  vec2 st = quadst[index];
-
-  INTERP_FACE_VARYING_2(outpt.v.uv, osd_active_uv_offset, st);
-
-  gl_Position = projectionMatrix * inpt[index].v.position;
-  EmitVertex();
-}
-
-#  ifdef WIREFRAME
-void emit_edge(int v0, int v1)
-{
-  emit(v0);
-  emit(v1);
-}
-#  endif
-
-#endif
-
-void main()
-{
-  gl_PrimitiveID = gl_PrimitiveIDIn;
-
-#ifdef FLAT_SHADING
-  vec3 A = (inpt[0].v.position - inpt[1].v.position).xyz;
-  vec3 B = (inpt[3].v.position - inpt[1].v.position).xyz;
-  vec3 flat_normal = normalize(cross(B, A));
-#  ifndef WIREFRAME
-  emit(0, flat_normal);
-  emit(1, flat_normal);
-  emit(3, flat_normal);
-  emit(2, flat_normal);
-#  else
-  emit_edge(0, 1, flat_normal);
-  emit_edge(1, 2, flat_normal);
-  emit_edge(2, 3, flat_normal);
-  emit_edge(3, 0, flat_normal);
-#  endif
-#else
-#  ifndef WIREFRAME
-  emit(0);
-  emit(1);
-  emit(3);
-  emit(2);
-#  else
-  emit_edge(0, 1);
-  emit_edge(1, 2);
-  emit_edge(2, 3);
-  emit_edge(3, 0);
-#  endif
-#endif
-
-  EndPrimitive();
-}
diff --git a/intern/opensubdiv/shader/gpu_shader_opensubdiv_vertex.glsl b/intern/opensubdiv/shader/gpu_shader_opensubdiv_vertex.glsl
deleted file mode 100644
index 474a716e927..00000000000
--- a/intern/opensubdiv/shader/gpu_shader_opensubdiv_vertex.glsl
+++ /dev/null
@@ -1,42 +0,0 @@
-/*
- * This program is free software; you can redistribute it and/or
- * modify it under the terms of the GNU General Public License
- * as published by the Free Software Foundation; either version 2
- * of the License, or (at your option) any later version.
- *
- * This program is distributed in the hope that it will be useful,
- * but WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
- * GNU General Public License for more details.
- *
- * You should have received a copy of the GNU General Public License
- * along with this program; if not, write to the Free Software Foundation,
- * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
- *
- * The Original Code is Copyright (C) 2014 Blender Foundation.
- * All rights reserved.
- */
-
-struct VertexData {
-  vec4 position;
-  vec3 normal;
-  vec2 uv;
-};
-
-in vec3 normal;
-in vec4 position;
-
-uniform mat4 modelViewMatrix;
-uniform mat3 normalMatrix;
-
-out block
-{
-  VertexData v;
-}
-outpt;
-
-void main()
-{
-  outpt.v.position = modelViewMatrix * position;
-  outpt.v.normal = normalize(normalMatrix * normal);
-}
diff --git a/intern/opensubdiv/stub/opensubdiv_gl_mesh_stub.cc b/intern/opensubdiv/stub/opensubdiv_gl_mesh_stub.cc
deleted file mode 100644
index 91ac0676dbd..00000000000
--- a/intern/opensubdiv/stub/opensubdiv_gl_mesh_stub.cc
+++ /dev/null
@@ -1,40 +0,0 @@
-// Copyright 2018 Blender Foundation. All rights reserved.
-//
-// This program is free software; you can redistribute it and/or
-// modify it under the terms of the GNU General Public License
-// as published by the Free Software Foundation; either version 2
-// of the License, or (at your option) any later version.
-//
-// This program is distributed in the hope that it will be useful,
-// but WITHOUT ANY WARRANTY; without even the implied warranty of
-// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-// GNU General Public License for more details.
-//
-// You should have received a copy of the GNU General Public License
-// along with this program; if not, write to the Free Software Foundation,
-// Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
-//
-// Author: Sergey Sharybin
-
-#include "opensubdiv_gl_mesh_capi.h"
-
-#include <cstddef>
-
-struct OpenSubdiv_GLMesh *openSubdiv_createOsdGLMeshFromTopologyRefiner(
-    OpenSubdiv_TopologyRefiner * /*topology_refiner*/, eOpenSubdivEvaluator /*evaluator_type*/)
-{
-  return NULL;
-}
-
-void openSubdiv_deleteOsdGLMesh(OpenSubdiv_GLMesh * /*gl_mesh*/)
-{
-}
-
-bool openSubdiv_initGLMeshDrawingResources(void)
-{
-  return false;
-}
-
-void openSubdiv_deinitGLMeshDrawingResources(void)
-{
-}