Merge branch 'master' into simulation-access-modifier

15 files changed, 150 insertions, 74 deletions

diff --git a/intern/cycles/blender/addon/ui.py b/intern/cycles/blender/addon/ui.py
index 92f0a8fb830..08641c05941 100644
--- a/intern/cycles/blender/addon/ui.py
+++ b/intern/cycles/blender/addon/ui.py
@@ -770,8 +770,9 @@ class CYCLES_RENDER_PT_filter(CyclesButtonsPanel, Panel):
         col.prop(view_layer, "use_strand", text="Hair")
         col.prop(view_layer, "use_volumes", text="Volumes")
         if with_freestyle:
-            col.prop(view_layer, "use_freestyle", text="Freestyle")
-            col.active = rd.use_freestyle
+            sub = col.row(align=True)
+            sub.prop(view_layer, "use_freestyle", text="Freestyle")
+            sub.active = rd.use_freestyle
 
 
 class CYCLES_RENDER_PT_override(CyclesButtonsPanel, Panel):
@@ -1409,6 +1410,8 @@ class CYCLES_LIGHT_PT_nodes(CyclesButtonsPanel, Panel):
     def draw(self, context):
         layout = self.layout
 
+        layout.use_property_split = True
+
         light = context.light
         panel_node_draw(layout, light, 'OUTPUT_LIGHT', 'Surface')
 
@@ -1458,6 +1461,8 @@ class CYCLES_WORLD_PT_surface(CyclesButtonsPanel, Panel):
     def draw(self, context):
         layout = self.layout
 
+        layout.use_property_split = True
+
         world = context.world
 
         if not panel_node_draw(layout, world, 'OUTPUT_WORLD', 'Surface'):
@@ -1477,6 +1482,8 @@ class CYCLES_WORLD_PT_volume(CyclesButtonsPanel, Panel):
     def draw(self, context):
         layout = self.layout
 
+        layout.use_property_split = True
+
         world = context.world
         panel_node_draw(layout, world, 'OUTPUT_WORLD', 'Volume')
 
@@ -1664,6 +1671,8 @@ class CYCLES_MATERIAL_PT_surface(CyclesButtonsPanel, Panel):
     def draw(self, context):
         layout = self.layout
 
+        layout.use_property_split = True
+
         mat = context.material
         if not panel_node_draw(layout, mat, 'OUTPUT_MATERIAL', 'Surface'):
             layout.prop(mat, "diffuse_color")
@@ -1682,6 +1691,8 @@ class CYCLES_MATERIAL_PT_volume(CyclesButtonsPanel, Panel):
     def draw(self, context):
         layout = self.layout
 
+        layout.use_property_split = True
+
         mat = context.material
         # cmat = mat.cycles
 
@@ -1700,6 +1711,8 @@ class CYCLES_MATERIAL_PT_displacement(CyclesButtonsPanel, Panel):
     def draw(self, context):
         layout = self.layout
 
+        layout.use_property_split = True
+
         mat = context.material
         panel_node_draw(layout, mat, 'OUTPUT_MATERIAL', 'Displacement')
 
diff --git a/intern/cycles/blender/blender_sync.cpp b/intern/cycles/blender/blender_sync.cpp
index f4c100bcd2b..0c120b944a7 100644
--- a/intern/cycles/blender/blender_sync.cpp
+++ b/intern/cycles/blender/blender_sync.cpp
@@ -108,7 +108,7 @@ void BlenderSync::sync_recalc(BL::Depsgraph &b_depsgraph, BL::SpaceView3D &b_v3d
     }
 
     if (dicing_prop_changed) {
-      for (const pair<GeometryKey, Geometry *> &iter : geometry_map.key_to_scene_data()) {
+      for (const pair<const GeometryKey, Geometry *> &iter : geometry_map.key_to_scene_data()) {
         Geometry *geom = iter.second;
         if (geom->type == Geometry::MESH) {
           Mesh *mesh = static_cast<Mesh *>(geom);
diff --git a/intern/cycles/bvh/bvh_optix.cpp b/intern/cycles/bvh/bvh_optix.cpp
index 26b64c24db5..740994b2ebc 100644
--- a/intern/cycles/bvh/bvh_optix.cpp
+++ b/intern/cycles/bvh/bvh_optix.cpp
@@ -156,6 +156,19 @@ void BVHOptiX::pack_tlas()
     PackedBVH &bvh_pack = geom->bvh->pack;
     int geom_prim_offset = geom->prim_offset;
 
+    // Merge visibility flags of all objects and fix object indices for non-instanced geometry
+    int object_index = 0;  // Unused for instanced geometry
+    int object_visibility = 0;
+    foreach (Object *ob, objects) {
+      if (ob->geometry == geom) {
+        object_visibility |= ob->visibility_for_tracing();
+        if (!geom->is_instanced()) {
+          object_index = ob->get_device_index();
+          break;
+        }
+      }
+    }
+
     // Merge primitive, object and triangle indexes
     if (!bvh_pack.prim_index.empty()) {
       int *bvh_prim_type = &bvh_pack.prim_type[0];
@@ -174,8 +187,8 @@ void BVHOptiX::pack_tlas()
         }
 
         pack_prim_type[pack_offset] = bvh_prim_type[i];
-        pack_prim_object[pack_offset] = 0;  // Unused for instanced geometry
-        pack_prim_visibility[pack_offset] = bvh_prim_visibility[i];
+        pack_prim_object[pack_offset] = object_index;
+        pack_prim_visibility[pack_offset] = bvh_prim_visibility[i] | object_visibility;
       }
     }
 
@@ -188,27 +201,6 @@ void BVHOptiX::pack_tlas()
       pack_verts_offset += prim_tri_size;
     }
   }
-
-  // Merge visibility flags of all objects and fix object indices for non-instanced geometry
-  foreach (Object *ob, objects) {
-    Geometry *const geom = ob->geometry;
-    size_t num_primitives = 0;
-
-    if (geom->type == Geometry::MESH) {
-      num_primitives = static_cast<Mesh *const>(geom)->num_triangles();
-    }
-    else if (geom->type == Geometry::HAIR) {
-      num_primitives = static_cast<Hair *const>(geom)->num_segments();
-    }
-
-    for (size_t i = 0; i < num_primitives; ++i) {
-      if (!geom->is_instanced()) {
-        assert(pack.prim_object[geom->optix_prim_offset + i] == 0);
-        pack.prim_object[geom->optix_prim_offset + i] = ob->get_device_index();
-      }
-      pack.prim_visibility[geom->optix_prim_offset + i] |= ob->visibility_for_tracing();
-    }
-  }
 }
 
 void BVHOptiX::pack_nodes(const BVHNode *)
diff --git a/intern/cycles/kernel/kernel_compat_cuda.h b/intern/cycles/kernel/kernel_compat_cuda.h
index 3c5a10540d5..4094e173da9 100644
--- a/intern/cycles/kernel/kernel_compat_cuda.h
+++ b/intern/cycles/kernel/kernel_compat_cuda.h
@@ -71,6 +71,7 @@ __device__ half __float2half(const float f)
 #define ccl_may_alias
 #define ccl_addr_space
 #define ccl_restrict __restrict__
+#define ccl_loop_no_unroll
 /* TODO(sergey): In theory we might use references with CUDA, however
  * performance impact yet to be investigated.
  */
diff --git a/intern/cycles/kernel/kernel_compat_opencl.h b/intern/cycles/kernel/kernel_compat_opencl.h
index 4963f1cd196..35dc95ca10d 100644
--- a/intern/cycles/kernel/kernel_compat_opencl.h
+++ b/intern/cycles/kernel/kernel_compat_opencl.h
@@ -43,6 +43,7 @@
 #define ccl_local __local
 #define ccl_local_param __local
 #define ccl_private __private
+#define ccl_loop_no_unroll __attribute__((opencl_unroll_hint(1)))
 #define ccl_restrict restrict
 #define ccl_ref
 #define ccl_align(n) __attribute__((aligned(n)))
diff --git a/intern/cycles/kernel/kernel_compat_optix.h b/intern/cycles/kernel/kernel_compat_optix.h
index 7068acc3a32..970f5cf864c 100644
--- a/intern/cycles/kernel/kernel_compat_optix.h
+++ b/intern/cycles/kernel/kernel_compat_optix.h
@@ -70,6 +70,7 @@ __device__ half __float2half(const float f)
 #define ccl_private
 #define ccl_may_alias
 #define ccl_addr_space
+#define ccl_loop_no_unroll
 #define ccl_restrict __restrict__
 #define ccl_ref
 #define ccl_align(n) __align__(n)
diff --git a/intern/cycles/kernel/osl/osl_services.cpp b/intern/cycles/kernel/osl/osl_services.cpp
index 2857de533f3..5292b5f8055 100644
--- a/intern/cycles/kernel/osl/osl_services.cpp
+++ b/intern/cycles/kernel/osl/osl_services.cpp
@@ -1011,7 +1011,13 @@ bool OSLRenderServices::get_userdata(
   return false; /* disabled by lockgeom */
 }
 
+#if OSL_LIBRARY_VERSION_CODE >= 11100
+TextureSystem::TextureHandle *OSLRenderServices::get_texture_handle(ustring filename,
+                                                                    OSL::ShadingContext *)
+#else
+
 TextureSystem::TextureHandle *OSLRenderServices::get_texture_handle(ustring filename)
+#endif
 {
   OSLTextureHandleMap::iterator it = textures.find(filename);
 
@@ -1365,6 +1371,17 @@ bool OSLRenderServices::environment(ustring filename,
   return status;
 }
 
+#if OSL_LIBRARY_VERSION_CODE >= 11100
+bool OSLRenderServices::get_texture_info(ustring filename,
+                                         TextureHandle *texture_handle,
+                                         TexturePerthread *,
+                                         OSL::ShadingContext *,
+                                         int subimage,
+                                         ustring dataname,
+                                         TypeDesc datatype,
+                                         void *data,
+                                         ustring *)
+#else
 bool OSLRenderServices::get_texture_info(OSL::ShaderGlobals *sg,
                                          ustring filename,
                                          TextureHandle *texture_handle,
@@ -1372,6 +1389,7 @@ bool OSLRenderServices::get_texture_info(OSL::ShaderGlobals *sg,
                                          ustring dataname,
                                          TypeDesc datatype,
                                          void *data)
+#endif
 {
   OSLTextureHandle *handle = (OSLTextureHandle *)texture_handle;
 
diff --git a/intern/cycles/kernel/osl/osl_services.h b/intern/cycles/kernel/osl/osl_services.h
index d32dace23bf..894d6e471ba 100644
--- a/intern/cycles/kernel/osl/osl_services.h
+++ b/intern/cycles/kernel/osl/osl_services.h
@@ -173,7 +173,12 @@ class OSLRenderServices : public OSL::RendererServices {
                   void *val,
                   bool derivatives) override;
 
+#if OSL_LIBRARY_VERSION_CODE >= 11100
+  TextureSystem::TextureHandle *get_texture_handle(ustring filename,
+                                                   OSL::ShadingContext *context) override;
+#else
   TextureSystem::TextureHandle *get_texture_handle(ustring filename) override;
+#endif
 
   bool good(TextureSystem::TextureHandle *texture_handle) override;
 
@@ -224,6 +229,17 @@ class OSLRenderServices : public OSL::RendererServices {
                    float *dresultdt,
                    ustring *errormessage) override;
 
+#if OSL_LIBRARY_VERSION_CODE >= 11100
+  bool get_texture_info(ustring filename,
+                        TextureHandle *texture_handle,
+                        TexturePerthread *texture_thread_info,
+                        OSL::ShadingContext *shading_context,
+                        int subimage,
+                        ustring dataname,
+                        TypeDesc datatype,
+                        void *data,
+                        ustring *errormessage) override;
+#else
   bool get_texture_info(OSL::ShaderGlobals *sg,
                         ustring filename,
                         TextureHandle *texture_handle,
@@ -231,6 +247,7 @@ class OSLRenderServices : public OSL::RendererServices {
                         ustring dataname,
                         TypeDesc datatype,
                         void *data) override;
+#endif
 
   static bool get_background_attribute(
       KernelGlobals *kg, ShaderData *sd, ustring name, TypeDesc type, bool derivatives, void *val);
diff --git a/intern/cycles/kernel/svm/svm_voronoi.h b/intern/cycles/kernel/svm/svm_voronoi.h
index 2ad22592eef..f0fc0068fa2 100644
--- a/intern/cycles/kernel/svm/svm_voronoi.h
+++ b/intern/cycles/kernel/svm/svm_voronoi.h
@@ -684,7 +684,8 @@ ccl_device void voronoi_f1_4d(float4 coord,
   float4 targetPosition = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
   for (int u = -1; u <= 1; u++) {
     for (int k = -1; k <= 1; k++) {
-      for (int j = -1; j <= 1; j++) {
+      ccl_loop_no_unroll for (int j = -1; j <= 1; j++)
+      {
         for (int i = -1; i <= 1; i++) {
           float4 cellOffset = make_float4(i, j, k, u);
           float4 pointPosition = cellOffset +
@@ -722,7 +723,8 @@ ccl_device void voronoi_smooth_f1_4d(float4 coord,
   float4 smoothPosition = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
   for (int u = -2; u <= 2; u++) {
     for (int k = -2; k <= 2; k++) {
-      for (int j = -2; j <= 2; j++) {
+      ccl_loop_no_unroll for (int j = -2; j <= 2; j++)
+      {
         for (int i = -2; i <= 2; i++) {
           float4 cellOffset = make_float4(i, j, k, u);
           float4 pointPosition = cellOffset +
@@ -765,7 +767,8 @@ ccl_device void voronoi_f2_4d(float4 coord,
   float4 positionF2 = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
   for (int u = -1; u <= 1; u++) {
     for (int k = -1; k <= 1; k++) {
-      for (int j = -1; j <= 1; j++) {
+      ccl_loop_no_unroll for (int j = -1; j <= 1; j++)
+      {
         for (int i = -1; i <= 1; i++) {
           float4 cellOffset = make_float4(i, j, k, u);
           float4 pointPosition = cellOffset +
@@ -803,7 +806,8 @@ ccl_device void voronoi_distance_to_edge_4d(float4 coord, float randomness, floa
   float minDistance = 8.0f;
   for (int u = -1; u <= 1; u++) {
     for (int k = -1; k <= 1; k++) {
-      for (int j = -1; j <= 1; j++) {
+      ccl_loop_no_unroll for (int j = -1; j <= 1; j++)
+      {
         for (int i = -1; i <= 1; i++) {
           float4 cellOffset = make_float4(i, j, k, u);
           float4 vectorToPoint = cellOffset +
@@ -822,7 +826,8 @@ ccl_device void voronoi_distance_to_edge_4d(float4 coord, float randomness, floa
   minDistance = 8.0f;
   for (int u = -1; u <= 1; u++) {
     for (int k = -1; k <= 1; k++) {
-      for (int j = -1; j <= 1; j++) {
+      ccl_loop_no_unroll for (int j = -1; j <= 1; j++)
+      {
         for (int i = -1; i <= 1; i++) {
           float4 cellOffset = make_float4(i, j, k, u);
           float4 vectorToPoint = cellOffset +
@@ -851,7 +856,8 @@ ccl_device void voronoi_n_sphere_radius_4d(float4 coord, float randomness, float
   float minDistance = 8.0f;
   for (int u = -1; u <= 1; u++) {
     for (int k = -1; k <= 1; k++) {
-      for (int j = -1; j <= 1; j++) {
+      ccl_loop_no_unroll for (int j = -1; j <= 1; j++)
+      {
         for (int i = -1; i <= 1; i++) {
           float4 cellOffset = make_float4(i, j, k, u);
           float4 pointPosition = cellOffset +
@@ -871,7 +877,8 @@ ccl_device void voronoi_n_sphere_radius_4d(float4 coord, float randomness, float
   float4 closestPointToClosestPoint = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
   for (int u = -1; u <= 1; u++) {
     for (int k = -1; k <= 1; k++) {
-      for (int j = -1; j <= 1; j++) {
+      ccl_loop_no_unroll for (int j = -1; j <= 1; j++)
+      {
         for (int i = -1; i <= 1; i++) {
           if (i == 0 && j == 0 && k == 0 && u == 0) {
             continue;
diff --git a/intern/cycles/render/mesh_volume.cpp b/intern/cycles/render/mesh_volume.cpp
index 74b8fc9e5ba..607363d01c6 100644
--- a/intern/cycles/render/mesh_volume.cpp
+++ b/intern/cycles/render/mesh_volume.cpp
@@ -19,6 +19,7 @@
 #include "render/scene.h"
 
 #include "util/util_foreach.h"
+#include "util/util_hash.h"
 #include "util/util_logging.h"
 #include "util/util_progress.h"
 #include "util/util_types.h"
@@ -447,7 +448,14 @@ void GeometryManager::create_volume_mesh(Mesh *mesh, Progress &progress)
   start_point = transform_point(&itfm, start_point);
   cell_size = transform_direction(&itfm, cell_size);
 
-  volume_params.start_point = start_point;
+  /* Slightly offset vertex coordinates to avoid overlapping faces with other
+   * volumes or meshes. The proper solution would be to improve intersection in
+   * the kernel to support robust handling of multiple overlapping faces or use
+   * an all-hit intersection similar to shadows. */
+  const float3 face_overlap_avoidance = cell_size * 0.1f *
+                                        hash_uint_to_float(hash_string(mesh->name.c_str()));
+
+  volume_params.start_point = start_point + face_overlap_avoidance;
   volume_params.cell_size = cell_size;
   volume_params.pad_size = pad_size;
 
diff --git a/intern/cycles/render/session.cpp b/intern/cycles/render/session.cpp
index 58bcc7ccdfb..f7df81a0601 100644
--- a/intern/cycles/render/session.cpp
+++ b/intern/cycles/render/session.cpp
@@ -293,14 +293,12 @@ void Session::run_gpu()
        * reset and draw in between */
       thread_scoped_lock buffers_lock(buffers_mutex);
 
-      /* avoid excessive denoising in viewport after reaching a certain amount of samples */
-      bool need_denoise = tile_manager.schedule_denoising || tile_manager.state.sample < 20 ||
-                          (time_dt() - last_display_time) >= params.progressive_update_timeout;
-
       /* update status and timing */
       update_status_time();
 
       /* render */
+      bool delayed_denoise = false;
+      const bool need_denoise = render_need_denoise(delayed_denoise);
       render(need_denoise);
 
       device->task_wait();
@@ -311,7 +309,7 @@ void Session::run_gpu()
       /* update status and timing */
       update_status_time();
 
-      gpu_need_display_buffer_update = need_denoise || !params.run_denoising;
+      gpu_need_display_buffer_update = !delayed_denoise;
       gpu_draw_ready = true;
       progress.set_update();
 
@@ -477,7 +475,7 @@ void Session::update_tile_sample(RenderTile &rtile)
   update_status_time();
 }
 
-void Session::release_tile(RenderTile &rtile)
+void Session::release_tile(RenderTile &rtile, const bool need_denoise)
 {
   thread_scoped_lock tile_lock(tile_mutex);
 
@@ -485,7 +483,7 @@ void Session::release_tile(RenderTile &rtile)
 
   bool delete_tile;
 
-  if (tile_manager.finish_tile(rtile.tile_index, delete_tile)) {
+  if (tile_manager.finish_tile(rtile.tile_index, need_denoise, delete_tile)) {
     if (write_render_tile_cb && params.progressive_refine == false) {
       write_render_tile_cb(rtile);
     }
@@ -687,21 +685,19 @@ void Session::run_cpu()
        * reset and draw in between */
       thread_scoped_lock buffers_lock(buffers_mutex);
 
-      /* avoid excessive denoising in viewport after reaching a certain amount of samples */
-      bool need_denoise = tile_manager.schedule_denoising || tile_manager.state.sample < 20 ||
-                          (time_dt() - last_display_time) >= params.progressive_update_timeout;
-
       /* update status and timing */
       update_status_time();
 
       /* render */
+      bool delayed_denoise = false;
+      const bool need_denoise = render_need_denoise(delayed_denoise);
       render(need_denoise);
 
       /* update status and timing */
       update_status_time();
 
       if (!params.background)
-        need_copy_to_display_buffer = need_denoise || !params.run_denoising;
+        need_copy_to_display_buffer = !delayed_denoise;
 
       if (!device->error_message().empty())
         progress.set_error(device->error_message());
@@ -1083,7 +1079,46 @@ void Session::update_status_time(bool show_pause, bool show_done)
   progress.set_status(status, substatus);
 }
 
-void Session::render(bool with_denoising)
+bool Session::render_need_denoise(bool &delayed)
+{
+  delayed = false;
+
+  /* Denoising enabled? */
+  if (!params.run_denoising) {
+    return false;
+  }
+
+  if (params.background) {
+    /* Background render, only denoise when rendering the last sample. */
+    return tile_manager.done();
+  }
+
+  /* Viewport render. */
+
+  /* It can happen that denoising was already enabled, but the scene still needs an update. */
+  if (scene->film->need_update || !scene->film->denoising_data_offset) {
+    return false;
+  }
+
+  /* Do not denoise until the sample at which denoising should start is reached. */
+  if (tile_manager.state.sample < params.denoising_start_sample) {
+    return false;
+  }
+
+  /* Cannot denoise with resolution divider and separate denoising devices.
+   * It breaks the copy in 'MultiDevice::map_neighbor_tiles' (which operates on
+   * the full buffer dimensions and not the scaled ones). */
+  if (!params.device.denoising_devices.empty() && tile_manager.state.resolution_divider > 1) {
+    return false;
+  }
+
+  /* Avoid excessive denoising in viewport after reaching a certain amount of samples. */
+  delayed = (tile_manager.state.sample >= 20 &&
+             (time_dt() - last_display_time) < params.progressive_update_timeout);
+  return !delayed;
+}
+
+void Session::render(bool need_denoise)
 {
   if (buffers && tile_manager.state.sample == tile_manager.range_start_sample) {
     /* Clear buffers. */
@@ -1098,7 +1133,7 @@ void Session::render(bool with_denoising)
   DeviceTask task(DeviceTask::RENDER);
 
   task.acquire_tile = function_bind(&Session::acquire_tile, this, _2, _1, _3);
-  task.release_tile = function_bind(&Session::release_tile, this, _1);
+  task.release_tile = function_bind(&Session::release_tile, this, _1, need_denoise);
   task.map_neighbor_tiles = function_bind(&Session::map_neighbor_tiles, this, _1, _2);
   task.unmap_neighbor_tiles = function_bind(&Session::unmap_neighbor_tiles, this, _1, _2);
   task.get_cancel = function_bind(&Progress::get_cancel, &this->progress);
@@ -1115,27 +1150,7 @@ void Session::render(bool with_denoising)
   /* Acquire render tiles by default. */
   task.tile_types = RenderTile::PATH_TRACE;
 
-  with_denoising = params.run_denoising && with_denoising;
-  if (with_denoising) {
-    /* Do not denoise viewport until the sample at which denoising should start is reached. */
-    if (!params.background && tile_manager.state.sample < params.denoising_start_sample) {
-      with_denoising = false;
-    }
-
-    /* Cannot denoise with resolution divider and separate denoising devices.
-     * It breaks the copy in 'MultiDevice::map_neighbor_tiles' (which operates on the full buffer
-     * dimensions and not the scaled ones). */
-    if (!params.device.denoising_devices.empty() && tile_manager.state.resolution_divider > 1) {
-      with_denoising = false;
-    }
-
-    /* It can happen that denoising was already enabled, but the scene still needs an update. */
-    if (scene->film->need_update || !scene->film->denoising_data_offset) {
-      with_denoising = false;
-    }
-  }
-
-  if (with_denoising) {
+  if (need_denoise) {
     task.denoising = params.denoising;
 
     task.pass_stride = scene->film->pass_stride;
diff --git a/intern/cycles/render/session.h b/intern/cycles/render/session.h
index 61970d87e9c..f06952e8020 100644
--- a/intern/cycles/render/session.h
+++ b/intern/cycles/render/session.h
@@ -186,7 +186,7 @@ class Session {
 
   void update_status_time(bool show_pause = false, bool show_done = false);
 
-  void render(bool with_denoising);
+  void render(bool use_denoise);
   void copy_to_display_buffer(int sample);
 
   void reset_(BufferParams &params, int samples);
@@ -199,9 +199,11 @@ class Session {
   bool draw_gpu(BufferParams &params, DeviceDrawParams &draw_params);
   void reset_gpu(BufferParams &params, int samples);
 
+  bool render_need_denoise(bool &delayed);
+
   bool acquire_tile(RenderTile &tile, Device *tile_device, uint tile_types);
   void update_tile_sample(RenderTile &tile);
-  void release_tile(RenderTile &tile);
+  void release_tile(RenderTile &tile, const bool need_denoise);
 
   void map_neighbor_tiles(RenderTile *tiles, Device *tile_device);
   void unmap_neighbor_tiles(RenderTile *tiles, Device *tile_device);
diff --git a/intern/cycles/render/tile.cpp b/intern/cycles/render/tile.cpp
index 1480b6d1aab..375c9fd8e09 100644
--- a/intern/cycles/render/tile.cpp
+++ b/intern/cycles/render/tile.cpp
@@ -441,13 +441,13 @@ bool TileManager::check_neighbor_state(int index, Tile::State min_state)
 
 /* Returns whether the tile should be written (and freed if no denoising is used) instead of
  * updating. */
-bool TileManager::finish_tile(int index, bool &delete_tile)
+bool TileManager::finish_tile(const int index, const bool need_denoise, bool &delete_tile)
 {
   delete_tile = false;
 
   switch (state.tiles[index].state) {
     case Tile::RENDER: {
-      if (!schedule_denoising) {
+      if (!(schedule_denoising && need_denoise)) {
         state.tiles[index].state = Tile::DONE;
         delete_tile = !progressive;
         return true;
diff --git a/intern/cycles/render/tile.h b/intern/cycles/render/tile.h
index 9fb9c1ca782..4858a275d5c 100644
--- a/intern/cycles/render/tile.h
+++ b/intern/cycles/render/tile.h
@@ -107,7 +107,7 @@ class TileManager {
   void set_samples(int num_samples);
   bool next();
   bool next_tile(Tile *&tile, int device, uint tile_types);
-  bool finish_tile(int index, bool &delete_tile);
+  bool finish_tile(const int index, const bool need_denoise, bool &delete_tile);
   bool done();
   bool has_tiles();
 
diff --git a/intern/cycles/util/util_defines.h b/intern/cycles/util/util_defines.h
index 24a20a969ab..e8e414587fb 100644
--- a/intern/cycles/util/util_defines.h
+++ b/intern/cycles/util/util_defines.h
@@ -45,6 +45,7 @@
 #  define ccl_restrict __restrict
 #  define ccl_ref &
 #  define ccl_optional_struct_init
+#  define ccl_loop_no_unroll
 #  define __KERNEL_WITH_SSE_ALIGN__
 
 #  if defined(_WIN32) && !defined(FREE_WINDOWS)