Merge branch 'master' into cycles_oneapi

64 files changed, 760 insertions, 631 deletions

diff --git a/intern/cycles/blender/addon/ui.py b/intern/cycles/blender/addon/ui.py
index 8565ad7b263..5b8c3960c82 100644
--- a/intern/cycles/blender/addon/ui.py
+++ b/intern/cycles/blender/addon/ui.py
@@ -758,8 +758,6 @@ class CYCLES_RENDER_PT_filter(CyclesButtonsPanel, Panel):
         layout.use_property_split = True
         layout.use_property_decorate = False
 
-        with_freestyle = bpy.app.build_options.freestyle
-
         scene = context.scene
         rd = scene.render
         view_layer = context.view_layer
diff --git a/intern/cycles/blender/camera.cpp b/intern/cycles/blender/camera.cpp
index 402fd7c4ec6..2ab5f02a337 100644
--- a/intern/cycles/blender/camera.cpp
+++ b/intern/cycles/blender/camera.cpp
@@ -643,7 +643,7 @@ void BlenderSync::sync_camera_motion(
     /* TODO(sergey): De-duplicate calculation with camera sync. */
     float fov = 2.0f * atanf((0.5f * sensor_size) / bcam.lens / aspectratio);
     if (fov != cam->get_fov()) {
-      VLOG(3) << "Camera " << b_ob.name() << " FOV change detected.";
+      VLOG_WORK << "Camera " << b_ob.name() << " FOV change detected.";
       if (motion_time == 0.0f) {
         cam->set_fov(fov);
       }
diff --git a/intern/cycles/blender/curves.cpp b/intern/cycles/blender/curves.cpp
index 4e9f4f62087..b01cb85711a 100644
--- a/intern/cycles/blender/curves.cpp
+++ b/intern/cycles/blender/curves.cpp
@@ -341,7 +341,7 @@ static void ExportCurveSegments(Scene *scene, Hair *hair, ParticleCurveData *CDa
 
   /* check allocation */
   if ((hair->get_curve_keys().size() != num_keys) || (hair->num_curves() != num_curves)) {
-    VLOG(1) << "Hair memory allocation failed, clearing data.";
+    VLOG_WARNING << "Hair memory allocation failed, clearing data.";
     hair->clear(true);
   }
 }
@@ -397,7 +397,7 @@ static void export_hair_motion_validate_attribute(Hair *hair,
   if (num_motion_keys != num_keys || !have_motion) {
     /* No motion or hair "topology" changed, remove attributes again. */
     if (num_motion_keys != num_keys) {
-      VLOG(1) << "Hair topology changed, removing motion attribute.";
+      VLOG_WORK << "Hair topology changed, removing motion attribute.";
     }
     hair->attributes.remove(ATTR_STD_MOTION_VERTEX_POSITION);
   }
diff --git a/intern/cycles/blender/display_driver.cpp b/intern/cycles/blender/display_driver.cpp
index ee67073a9a4..a1bc064be68 100644
--- a/intern/cycles/blender/display_driver.cpp
+++ b/intern/cycles/blender/display_driver.cpp
@@ -982,10 +982,8 @@ void BlenderDisplayDriver::draw(const Params &params)
   gl_render_sync_ = glFenceSync(GL_SYNC_GPU_COMMANDS_COMPLETE, 0);
   glFlush();
 
-  if (VLOG_IS_ON(5)) {
-    VLOG(5) << "Number of textures: " << GLTexture::num_used;
-    VLOG(5) << "Number of PBOs: " << GLPixelBufferObject::num_used;
-  }
+  VLOG_DEVICE_STATS << "Display driver number of textures: " << GLTexture::num_used;
+  VLOG_DEVICE_STATS << "Display driver number of PBOs: " << GLPixelBufferObject::num_used;
 
   if (use_gl_context_) {
     gl_context_mutex_.unlock();
diff --git a/intern/cycles/blender/mesh.cpp b/intern/cycles/blender/mesh.cpp
index e2db52cc5c1..63913b7cd7f 100644
--- a/intern/cycles/blender/mesh.cpp
+++ b/intern/cycles/blender/mesh.cpp
@@ -1212,17 +1212,17 @@ void BlenderSync::sync_mesh_motion(BL::Depsgraph b_depsgraph,
           memcmp(mP, &mesh->get_verts()[0], sizeof(float3) * numverts) == 0) {
         /* no motion, remove attributes again */
         if (b_mesh.vertices.length() != numverts) {
-          VLOG(1) << "Topology differs, disabling motion blur for object " << ob_name;
+          VLOG_WARNING << "Topology differs, disabling motion blur for object " << ob_name;
         }
         else {
-          VLOG(1) << "No actual deformation motion for object " << ob_name;
+          VLOG_DEBUG << "No actual deformation motion for object " << ob_name;
         }
         mesh->attributes.remove(ATTR_STD_MOTION_VERTEX_POSITION);
         if (attr_mN)
           mesh->attributes.remove(ATTR_STD_MOTION_VERTEX_NORMAL);
       }
       else if (motion_step > 0) {
-        VLOG(1) << "Filling deformation motion for object " << ob_name;
+        VLOG_DEBUG << "Filling deformation motion for object " << ob_name;
         /* motion, fill up previous steps that we might have skipped because
          * they had no motion, but we need them anyway now */
         float3 *P = &mesh->get_verts()[0];
@@ -1236,8 +1236,8 @@ void BlenderSync::sync_mesh_motion(BL::Depsgraph b_depsgraph,
     }
     else {
       if (b_mesh.vertices.length() != numverts) {
-        VLOG(1) << "Topology differs, discarding motion blur for object " << ob_name << " at time "
-                << motion_step;
+        VLOG_WARNING << "Topology differs, discarding motion blur for object " << ob_name
+                     << " at time " << motion_step;
         memcpy(mP, &mesh->get_verts()[0], sizeof(float3) * numverts);
         if (mN != NULL) {
           memcpy(mN, attr_N->data_float3(), sizeof(float3) * numverts);
diff --git a/intern/cycles/blender/object.cpp b/intern/cycles/blender/object.cpp
index 9b08b564b25..ca1aa6329d9 100644
--- a/intern/cycles/blender/object.cpp
+++ b/intern/cycles/blender/object.cpp
@@ -762,7 +762,7 @@ void BlenderSync::sync_motion(BL::RenderSettings &b_render,
       continue;
     }
 
-    VLOG(1) << "Synchronizing motion for the relative time " << relative_time << ".";
+    VLOG_WORK << "Synchronizing motion for the relative time " << relative_time << ".";
 
     /* fixed shutter time to get previous and next frame for motion pass */
     float shuttertime = scene->motion_shutter_time();
diff --git a/intern/cycles/blender/session.cpp b/intern/cycles/blender/session.cpp
index 87f051ba50b..6d27b8e7d87 100644
--- a/intern/cycles/blender/session.cpp
+++ b/intern/cycles/blender/session.cpp
@@ -458,8 +458,8 @@ void BlenderSession::render(BL::Depsgraph &b_depsgraph_)
 
   double total_time, render_time;
   session->progress.get_time(total_time, render_time);
-  VLOG(1) << "Total render time: " << total_time;
-  VLOG(1) << "Render time (without synchronization): " << render_time;
+  VLOG_INFO << "Total render time: " << total_time;
+  VLOG_INFO << "Render time (without synchronization): " << render_time;
 }
 
 void BlenderSession::render_frame_finish()
diff --git a/intern/cycles/blender/sync.cpp b/intern/cycles/blender/sync.cpp
index 1028c940772..63e9e1e0e68 100644
--- a/intern/cycles/blender/sync.cpp
+++ b/intern/cycles/blender/sync.cpp
@@ -285,7 +285,7 @@ void BlenderSync::sync_data(BL::RenderSettings &b_render,
 
   free_data_after_sync(b_depsgraph);
 
-  VLOG(1) << "Total time spent synchronizing data: " << timer.get_time();
+  VLOG_INFO << "Total time spent synchronizing data: " << timer.get_time();
 
   has_updates_ = false;
 }
@@ -390,7 +390,7 @@ void BlenderSync::sync_integrator(BL::ViewLayer &b_view_layer, bool background)
   }
 
   if (scrambling_distance != 1.0f) {
-    VLOG(3) << "Using scrambling distance: " << scrambling_distance;
+    VLOG_INFO << "Using scrambling distance: " << scrambling_distance;
   }
   integrator->set_scrambling_distance(scrambling_distance);
 
diff --git a/intern/cycles/bvh/build.cpp b/intern/cycles/bvh/build.cpp
index 1df3517673e..cf03f05de60 100644
--- a/intern/cycles/bvh/build.cpp
+++ b/intern/cycles/bvh/build.cpp
@@ -529,7 +529,7 @@ BVHNode *BVHBuild::run()
     if (progress.get_cancel()) {
       rootnode->deleteSubtree();
       rootnode = NULL;
-      VLOG(1) << "BVH build cancelled.";
+      VLOG_WORK << "BVH build cancelled.";
     }
     else {
       /*rotate(rootnode, 4, 5);*/
@@ -537,26 +537,26 @@ BVHNode *BVHBuild::run()
       rootnode->update_time();
     }
     if (rootnode != NULL) {
-      VLOG(1) << "BVH build statistics:\n"
-              << "  Build time: " << time_dt() - build_start_time << "\n"
-              << "  Total number of nodes: "
-              << string_human_readable_number(rootnode->getSubtreeSize(BVH_STAT_NODE_COUNT))
-              << "\n"
-              << "  Number of inner nodes: "
-              << string_human_readable_number(rootnode->getSubtreeSize(BVH_STAT_INNER_COUNT))
-              << "\n"
-              << "  Number of leaf nodes: "
-              << string_human_readable_number(rootnode->getSubtreeSize(BVH_STAT_LEAF_COUNT))
-              << "\n"
-              << "  Number of unaligned nodes: "
-              << string_human_readable_number(rootnode->getSubtreeSize(BVH_STAT_UNALIGNED_COUNT))
-              << "\n"
-              << "  Allocation slop factor: "
-              << ((prim_type.capacity() != 0) ? (float)prim_type.size() / prim_type.capacity() :
-                                                1.0f)
-              << "\n"
-              << "  Maximum depth: "
-              << string_human_readable_number(rootnode->getSubtreeSize(BVH_STAT_DEPTH)) << "\n";
+      VLOG_WORK << "BVH build statistics:\n"
+                << "  Build time: " << time_dt() - build_start_time << "\n"
+                << "  Total number of nodes: "
+                << string_human_readable_number(rootnode->getSubtreeSize(BVH_STAT_NODE_COUNT))
+                << "\n"
+                << "  Number of inner nodes: "
+                << string_human_readable_number(rootnode->getSubtreeSize(BVH_STAT_INNER_COUNT))
+                << "\n"
+                << "  Number of leaf nodes: "
+                << string_human_readable_number(rootnode->getSubtreeSize(BVH_STAT_LEAF_COUNT))
+                << "\n"
+                << "  Number of unaligned nodes: "
+                << string_human_readable_number(rootnode->getSubtreeSize(BVH_STAT_UNALIGNED_COUNT))
+                << "\n"
+                << "  Allocation slop factor: "
+                << ((prim_type.capacity() != 0) ? (float)prim_type.size() / prim_type.capacity() :
+                                                  1.0f)
+                << "\n"
+                << "  Maximum depth: "
+                << string_human_readable_number(rootnode->getSubtreeSize(BVH_STAT_DEPTH)) << "\n";
     }
   }
 
diff --git a/intern/cycles/bvh/embree.cpp b/intern/cycles/bvh/embree.cpp
index 0fc71f49ce3..ea7480bd193 100644
--- a/intern/cycles/bvh/embree.cpp
+++ b/intern/cycles/bvh/embree.cpp
@@ -332,7 +332,7 @@ static bool rtc_memory_monitor_func(void *userPtr, const ssize_t bytes, const bo
 
 static void rtc_error_func(void *, enum RTCError, const char *str)
 {
-  VLOG(1) << str;
+  VLOG_WARNING << str;
 }
 
 static double progress_start_time = 0.0;
@@ -521,8 +521,8 @@ void BVHEmbree::add_triangles(const Object *ob, const Mesh *mesh, int i)
       geom_id, RTC_BUFFER_TYPE_INDEX, 0, RTC_FORMAT_UINT3, sizeof(int) * 3, num_triangles);
   assert(rtc_indices);
   if (!rtc_indices) {
-    VLOG(1) << "Embree could not create new geometry buffer for mesh " << mesh->name.c_str()
-            << ".\n";
+    VLOG_WARNING << "Embree could not create new geometry buffer for mesh " << mesh->name.c_str()
+                 << ".\n";
     return;
   }
   for (size_t j = 0; j < num_triangles; ++j) {
diff --git a/intern/cycles/device/cpu/device_impl.cpp b/intern/cycles/device/cpu/device_impl.cpp
index 612c391f7d5..0a4eb089037 100644
--- a/intern/cycles/device/cpu/device_impl.cpp
+++ b/intern/cycles/device/cpu/device_impl.cpp
@@ -55,8 +55,8 @@ CPUDevice::CPUDevice(const DeviceInfo &info_, Stats &stats_, Profiler &profiler_
 {
   /* Pick any kernel, all of them are supposed to have same level of microarchitecture
    * optimization. */
-  VLOG(1) << "Using " << get_cpu_kernels().integrator_init_from_camera.get_uarch_name()
-          << " CPU kernels.";
+  VLOG_INFO << "Using " << get_cpu_kernels().integrator_init_from_camera.get_uarch_name()
+            << " CPU kernels.";
 
   if (info.cpu_threads == 0) {
     info.cpu_threads = TaskScheduler::max_concurrency();
@@ -111,9 +111,9 @@ void CPUDevice::mem_alloc(device_memory &mem)
   }
   else {
     if (mem.name) {
-      VLOG(1) << "Buffer allocate: " << mem.name << ", "
-              << string_human_readable_number(mem.memory_size()) << " bytes. ("
-              << string_human_readable_size(mem.memory_size()) << ")";
+      VLOG_WORK << "Buffer allocate: " << mem.name << ", "
+                << string_human_readable_number(mem.memory_size()) << " bytes. ("
+                << string_human_readable_size(mem.memory_size()) << ")";
     }
 
     if (mem.type == MEM_DEVICE_ONLY || !mem.host_pointer) {
@@ -205,9 +205,9 @@ void CPUDevice::const_copy_to(const char *name, void *host, size_t size)
 
 void CPUDevice::global_alloc(device_memory &mem)
 {
-  VLOG(1) << "Global memory allocate: " << mem.name << ", "
-          << string_human_readable_number(mem.memory_size()) << " bytes. ("
-          << string_human_readable_size(mem.memory_size()) << ")";
+  VLOG_WORK << "Global memory allocate: " << mem.name << ", "
+            << string_human_readable_number(mem.memory_size()) << " bytes. ("
+            << string_human_readable_size(mem.memory_size()) << ")";
 
   kernel_global_memory_copy(&kernel_globals, mem.name, mem.host_pointer, mem.data_size);
 
@@ -227,9 +227,9 @@ void CPUDevice::global_free(device_memory &mem)
 
 void CPUDevice::tex_alloc(device_texture &mem)
 {
-  VLOG(1) << "Texture allocate: " << mem.name << ", "
-          << string_human_readable_number(mem.memory_size()) << " bytes. ("
-          << string_human_readable_size(mem.memory_size()) << ")";
+  VLOG_WORK << "Texture allocate: " << mem.name << ", "
+            << string_human_readable_number(mem.memory_size()) << " bytes. ("
+            << string_human_readable_size(mem.memory_size()) << ")";
 
   mem.device_pointer = (device_ptr)mem.host_pointer;
   mem.device_size = mem.memory_size();
diff --git a/intern/cycles/device/cuda/device.cpp b/intern/cycles/device/cuda/device.cpp
index 400490336d6..5a213c45b71 100644
--- a/intern/cycles/device/cuda/device.cpp
+++ b/intern/cycles/device/cuda/device.cpp
@@ -29,24 +29,25 @@ bool device_cuda_init()
   initialized = true;
   int cuew_result = cuewInit(CUEW_INIT_CUDA);
   if (cuew_result == CUEW_SUCCESS) {
-    VLOG(1) << "CUEW initialization succeeded";
+    VLOG_INFO << "CUEW initialization succeeded";
     if (CUDADevice::have_precompiled_kernels()) {
-      VLOG(1) << "Found precompiled kernels";
+      VLOG_INFO << "Found precompiled kernels";
       result = true;
     }
     else if (cuewCompilerPath() != NULL) {
-      VLOG(1) << "Found CUDA compiler " << cuewCompilerPath();
+      VLOG_INFO << "Found CUDA compiler " << cuewCompilerPath();
       result = true;
     }
     else {
-      VLOG(1) << "Neither precompiled kernels nor CUDA compiler was found,"
-              << " unable to use CUDA";
+      VLOG_INFO << "Neither precompiled kernels nor CUDA compiler was found,"
+                << " unable to use CUDA";
     }
   }
   else {
-    VLOG(1) << "CUEW initialization failed: "
-            << ((cuew_result == CUEW_ERROR_ATEXIT_FAILED) ? "Error setting up atexit() handler" :
-                                                            "Error opening the library");
+    VLOG_WARNING << "CUEW initialization failed: "
+                 << ((cuew_result == CUEW_ERROR_ATEXIT_FAILED) ?
+                         "Error setting up atexit() handler" :
+                         "Error opening the library");
   }
 
   return result;
@@ -121,7 +122,8 @@ void device_cuda_info(vector<DeviceInfo> &devices)
     int major;
     cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, num);
     if (major < 3) {
-      VLOG(1) << "Ignoring device \"" << name << "\", this graphics card is no longer supported.";
+      VLOG_INFO << "Ignoring device \"" << name
+                << "\", this graphics card is no longer supported.";
       continue;
     }
 
@@ -166,21 +168,21 @@ void device_cuda_info(vector<DeviceInfo> &devices)
      * Windows 10 even when it is, due to an issue in application profiles.
      * Detect case where we expect it to be available and override. */
     if (preempt_attr == 0 && (major >= 6) && system_windows_version_at_least(10, 17134)) {
-      VLOG(1) << "Assuming device has compute preemption on Windows 10.";
+      VLOG_INFO << "Assuming device has compute preemption on Windows 10.";
       preempt_attr = 1;
     }
 
     if (timeout_attr && !preempt_attr) {
-      VLOG(1) << "Device is recognized as display.";
+      VLOG_INFO << "Device is recognized as display.";
       info.description += " (Display)";
       info.display_device = true;
       display_devices.push_back(info);
     }
     else {
-      VLOG(1) << "Device has compute preemption or is not used for display.";
+      VLOG_INFO << "Device has compute preemption or is not used for display.";
       devices.push_back(info);
     }
-    VLOG(1) << "Added device \"" << name << "\" with id \"" << info.id << "\".";
+    VLOG_INFO << "Added device \"" << name << "\" with id \"" << info.id << "\".";
   }
 
   if (!display_devices.empty())
diff --git a/intern/cycles/device/cuda/device_impl.cpp b/intern/cycles/device/cuda/device_impl.cpp
index cb7e909a2d5..e75224abe90 100644
--- a/intern/cycles/device/cuda/device_impl.cpp
+++ b/intern/cycles/device/cuda/device_impl.cpp
@@ -244,9 +244,9 @@ string CUDADevice::compile_kernel(const uint kernel_features,
   if (!use_adaptive_compilation()) {
     if (!force_ptx) {
       const string cubin = path_get(string_printf("lib/%s_sm_%d%d.cubin", name, major, minor));
-      VLOG(1) << "Testing for pre-compiled kernel " << cubin << ".";
+      VLOG_INFO << "Testing for pre-compiled kernel " << cubin << ".";
       if (path_exists(cubin)) {
-        VLOG(1) << "Using precompiled kernel.";
+        VLOG_INFO << "Using precompiled kernel.";
         return cubin;
       }
     }
@@ -256,9 +256,9 @@ string CUDADevice::compile_kernel(const uint kernel_features,
     while (ptx_major >= 3) {
       const string ptx = path_get(
           string_printf("lib/%s_compute_%d%d.ptx", name, ptx_major, ptx_minor));
-      VLOG(1) << "Testing for pre-compiled kernel " << ptx << ".";
+      VLOG_INFO << "Testing for pre-compiled kernel " << ptx << ".";
       if (path_exists(ptx)) {
-        VLOG(1) << "Using precompiled kernel.";
+        VLOG_INFO << "Using precompiled kernel.";
         return ptx;
       }
 
@@ -287,9 +287,9 @@ string CUDADevice::compile_kernel(const uint kernel_features,
   const string cubin_file = string_printf(
       "cycles_%s_%s_%d%d_%s.%s", name, kernel_arch, major, minor, kernel_md5.c_str(), kernel_ext);
   const string cubin = path_cache_get(path_join("kernels", cubin_file));
-  VLOG(1) << "Testing for locally compiled kernel " << cubin << ".";
+  VLOG_INFO << "Testing for locally compiled kernel " << cubin << ".";
   if (path_exists(cubin)) {
-    VLOG(1) << "Using locally compiled kernel.";
+    VLOG_INFO << "Using locally compiled kernel.";
     return cubin;
   }
 
@@ -323,7 +323,7 @@ string CUDADevice::compile_kernel(const uint kernel_features,
   }
 
   const int nvcc_cuda_version = cuewCompilerVersion();
-  VLOG(1) << "Found nvcc " << nvcc << ", CUDA version " << nvcc_cuda_version << ".";
+  VLOG_INFO << "Found nvcc " << nvcc << ", CUDA version " << nvcc_cuda_version << ".";
   if (nvcc_cuda_version < 101) {
     printf(
         "Unsupported CUDA version %d.%d detected, "
@@ -399,7 +399,8 @@ bool CUDADevice::load_kernels(const uint kernel_features)
    */
   if (cuModule) {
     if (use_adaptive_compilation()) {
-      VLOG(1) << "Skipping CUDA kernel reload for adaptive compilation, not currently supported.";
+      VLOG_INFO
+          << "Skipping CUDA kernel reload for adaptive compilation, not currently supported.";
     }
     return true;
   }
@@ -481,8 +482,8 @@ void CUDADevice::reserve_local_memory(const uint kernel_features)
     cuMemGetInfo(&free_after, &total);
   }
 
-  VLOG(1) << "Local memory reserved " << string_human_readable_number(free_before - free_after)
-          << " bytes. (" << string_human_readable_size(free_before - free_after) << ")";
+  VLOG_INFO << "Local memory reserved " << string_human_readable_number(free_before - free_after)
+            << " bytes. (" << string_human_readable_size(free_before - free_after) << ")";
 
 #  if 0
   /* For testing mapped host memory, fill up device memory. */
@@ -513,7 +514,7 @@ void CUDADevice::init_host_memory()
     }
   }
   else {
-    VLOG(1) << "Mapped host memory disabled, failed to get system RAM";
+    VLOG_WARNING << "Mapped host memory disabled, failed to get system RAM";
     map_host_limit = 0;
   }
 
@@ -524,8 +525,8 @@ void CUDADevice::init_host_memory()
   device_working_headroom = 32 * 1024 * 1024LL;   // 32MB
   device_texture_headroom = 128 * 1024 * 1024LL;  // 128MB
 
-  VLOG(1) << "Mapped host memory limit set to " << string_human_readable_number(map_host_limit)
-          << " bytes. (" << string_human_readable_size(map_host_limit) << ")";
+  VLOG_INFO << "Mapped host memory limit set to " << string_human_readable_number(map_host_limit)
+            << " bytes. (" << string_human_readable_size(map_host_limit) << ")";
 }
 
 void CUDADevice::load_texture_info()
@@ -593,7 +594,7 @@ void CUDADevice::move_textures_to_host(size_t size, bool for_texture)
      * multiple CUDA devices could be moving the memory. The
      * first one will do it, and the rest will adopt the pointer. */
     if (max_mem) {
-      VLOG(1) << "Move memory from device to host: " << max_mem->name;
+      VLOG_WORK << "Move memory from device to host: " << max_mem->name;
 
       static thread_mutex move_mutex;
       thread_scoped_lock lock(move_mutex);
@@ -701,9 +702,9 @@ CUDADevice::CUDAMem *CUDADevice::generic_alloc(device_memory &mem, size_t pitch_
   }
 
   if (mem.name) {
-    VLOG(1) << "Buffer allocate: " << mem.name << ", "
-            << string_human_readable_number(mem.memory_size()) << " bytes. ("
-            << string_human_readable_size(mem.memory_size()) << ")" << status;
+    VLOG_WORK << "Buffer allocate: " << mem.name << ", "
+              << string_human_readable_number(mem.memory_size()) << " bytes. ("
+              << string_human_readable_size(mem.memory_size()) << ")" << status;
   }
 
   mem.device_pointer = (device_ptr)device_pointer;
@@ -1008,9 +1009,9 @@ void CUDADevice::tex_alloc(device_texture &mem)
     desc.NumChannels = mem.data_elements;
     desc.Flags = 0;
 
-    VLOG(1) << "Array 3D allocate: " << mem.name << ", "
-            << string_human_readable_number(mem.memory_size()) << " bytes. ("
-            << string_human_readable_size(mem.memory_size()) << ")";
+    VLOG_WORK << "Array 3D allocate: " << mem.name << ", "
+              << string_human_readable_number(mem.memory_size()) << " bytes. ("
+              << string_human_readable_size(mem.memory_size()) << ")";
 
     cuda_assert(cuArray3DCreate(&array_3d, &desc));
 
diff --git a/intern/cycles/device/cuda/queue.cpp b/intern/cycles/device/cuda/queue.cpp
index 38c71866ad0..5912e68a92b 100644
--- a/intern/cycles/device/cuda/queue.cpp
+++ b/intern/cycles/device/cuda/queue.cpp
@@ -39,12 +39,12 @@ int CUDADeviceQueue::num_concurrent_states(const size_t state_size) const
       num_states = max((int)(num_states * factor), 1024);
     }
     else {
-      VLOG(3) << "CYCLES_CONCURRENT_STATES_FACTOR evaluated to 0";
+      VLOG_DEVICE_STATS << "CYCLES_CONCURRENT_STATES_FACTOR evaluated to 0";
     }
   }
 
-  VLOG(3) << "GPU queue concurrent states: " << num_states << ", using up to "
-          << string_human_readable_size(num_states * state_size);
+  VLOG_DEVICE_STATS << "GPU queue concurrent states: " << num_states << ", using up to "
+                    << string_human_readable_size(num_states * state_size);
 
   return num_states;
 }
diff --git a/intern/cycles/device/device.cpp b/intern/cycles/device/device.cpp
index 4c244d5bb78..ace6ed517f5 100644
--- a/intern/cycles/device/device.cpp
+++ b/intern/cycles/device/device.cpp
@@ -364,8 +364,8 @@ DeviceInfo Device::get_multi_device(const vector<DeviceInfo> &subdevices,
         int orig_cpu_threads = (threads) ? threads : TaskScheduler::max_concurrency();
         int cpu_threads = max(orig_cpu_threads - (subdevices.size() - 1), size_t(0));
 
-        VLOG(1) << "CPU render threads reduced from " << orig_cpu_threads << " to " << cpu_threads
-                << ", to dedicate to GPU.";
+        VLOG_INFO << "CPU render threads reduced from " << orig_cpu_threads << " to "
+                  << cpu_threads << ", to dedicate to GPU.";
 
         if (cpu_threads >= 1) {
           DeviceInfo cpu_device = device;
@@ -377,7 +377,7 @@ DeviceInfo Device::get_multi_device(const vector<DeviceInfo> &subdevices,
         }
       }
       else {
-        VLOG(1) << "CPU render threads disabled for interactive render.";
+        VLOG_INFO << "CPU render threads disabled for interactive render.";
         continue;
       }
     }
diff --git a/intern/cycles/device/hip/device.cpp b/intern/cycles/device/hip/device.cpp
index d6a5ed9c419..3c9c73e7db0 100644
--- a/intern/cycles/device/hip/device.cpp
+++ b/intern/cycles/device/hip/device.cpp
@@ -29,30 +29,31 @@ bool device_hip_init()
   initialized = true;
   int hipew_result = hipewInit(HIPEW_INIT_HIP);
   if (hipew_result == HIPEW_SUCCESS) {
-    VLOG(1) << "HIPEW initialization succeeded";
+    VLOG_INFO << "HIPEW initialization succeeded";
     if (HIPDevice::have_precompiled_kernels()) {
-      VLOG(1) << "Found precompiled kernels";
+      VLOG_INFO << "Found precompiled kernels";
       result = true;
     }
     else if (hipewCompilerPath() != NULL) {
-      VLOG(1) << "Found HIPCC " << hipewCompilerPath();
+      VLOG_INFO << "Found HIPCC " << hipewCompilerPath();
       result = true;
     }
     else {
-      VLOG(1) << "Neither precompiled kernels nor HIPCC was found,"
-              << " unable to use HIP";
+      VLOG_INFO << "Neither precompiled kernels nor HIPCC was found,"
+                << " unable to use HIP";
     }
   }
   else {
     if (hipew_result == HIPEW_ERROR_ATEXIT_FAILED) {
-      VLOG(1) << "HIPEW initialization failed: Error setting up atexit() handler";
+      VLOG_WARNING << "HIPEW initialization failed: Error setting up atexit() handler";
     }
     else if (hipew_result == HIPEW_ERROR_OLD_DRIVER) {
-      VLOG(1) << "HIPEW initialization failed: Driver version too old, requires AMD Radeon Pro "
-                 "21.Q4 driver or newer";
+      VLOG_WARNING
+          << "HIPEW initialization failed: Driver version too old, requires AMD Radeon Pro "
+             "21.Q4 driver or newer";
     }
     else {
-      VLOG(1) << "HIPEW initialization failed: Error opening HIP dynamic library";
+      VLOG_WARNING << "HIPEW initialization failed: Error opening HIP dynamic library";
     }
   }
 
@@ -165,16 +166,16 @@ void device_hip_info(vector<DeviceInfo> &devices)
     hipDeviceGetAttribute(&timeout_attr, hipDeviceAttributeKernelExecTimeout, num);
 
     if (timeout_attr && !preempt_attr) {
-      VLOG(1) << "Device is recognized as display.";
+      VLOG_INFO << "Device is recognized as display.";
       info.description += " (Display)";
       info.display_device = true;
       display_devices.push_back(info);
     }
     else {
-      VLOG(1) << "Device has compute preemption or is not used for display.";
+      VLOG_INFO << "Device has compute preemption or is not used for display.";
       devices.push_back(info);
     }
-    VLOG(1) << "Added device \"" << name << "\" with id \"" << info.id << "\".";
+    VLOG_INFO << "Added device \"" << name << "\" with id \"" << info.id << "\".";
   }
 
   if (!display_devices.empty())
diff --git a/intern/cycles/device/hip/device_impl.cpp b/intern/cycles/device/hip/device_impl.cpp
index ea68c821166..652c1001f85 100644
--- a/intern/cycles/device/hip/device_impl.cpp
+++ b/intern/cycles/device/hip/device_impl.cpp
@@ -233,9 +233,9 @@ string HIPDevice::compile_kernel(const uint kernel_features, const char *name, c
   /* Attempt to use kernel provided with Blender. */
   if (!use_adaptive_compilation()) {
     const string fatbin = path_get(string_printf("lib/%s_%s.fatbin", name, arch));
-    VLOG(1) << "Testing for pre-compiled kernel " << fatbin << ".";
+    VLOG_INFO << "Testing for pre-compiled kernel " << fatbin << ".";
     if (path_exists(fatbin)) {
-      VLOG(1) << "Using precompiled kernel.";
+      VLOG_INFO << "Using precompiled kernel.";
       return fatbin;
     }
   }
@@ -265,9 +265,9 @@ string HIPDevice::compile_kernel(const uint kernel_features, const char *name, c
   const string include_path = source_path;
   const string fatbin_file = string_printf("cycles_%s_%s_%s", name, arch, kernel_md5.c_str());
   const string fatbin = path_cache_get(path_join("kernels", fatbin_file));
-  VLOG(1) << "Testing for locally compiled kernel " << fatbin << ".";
+  VLOG_INFO << "Testing for locally compiled kernel " << fatbin << ".";
   if (path_exists(fatbin)) {
-    VLOG(1) << "Using locally compiled kernel.";
+    VLOG_INFO << "Using locally compiled kernel.";
     return fatbin;
   }
 
@@ -301,7 +301,7 @@ string HIPDevice::compile_kernel(const uint kernel_features, const char *name, c
   }
 
   const int hipcc_hip_version = hipewCompilerVersion();
-  VLOG(1) << "Found hipcc " << hipcc << ", HIP version " << hipcc_hip_version << ".";
+  VLOG_INFO << "Found hipcc " << hipcc << ", HIP version " << hipcc_hip_version << ".";
   if (hipcc_hip_version < 40) {
     printf(
         "Unsupported HIP version %d.%d detected, "
@@ -361,7 +361,7 @@ bool HIPDevice::load_kernels(const uint kernel_features)
    */
   if (hipModule) {
     if (use_adaptive_compilation()) {
-      VLOG(1) << "Skipping HIP kernel reload for adaptive compilation, not currently supported.";
+      VLOG_INFO << "Skipping HIP kernel reload for adaptive compilation, not currently supported.";
     }
     return true;
   }
@@ -444,8 +444,8 @@ void HIPDevice::reserve_local_memory(const uint kernel_features)
     hipMemGetInfo(&free_after, &total);
   }
 
-  VLOG(1) << "Local memory reserved " << string_human_readable_number(free_before - free_after)
-          << " bytes. (" << string_human_readable_size(free_before - free_after) << ")";
+  VLOG_INFO << "Local memory reserved " << string_human_readable_number(free_before - free_after)
+            << " bytes. (" << string_human_readable_size(free_before - free_after) << ")";
 
 #  if 0
   /* For testing mapped host memory, fill up device memory. */
@@ -476,7 +476,7 @@ void HIPDevice::init_host_memory()
     }
   }
   else {
-    VLOG(1) << "Mapped host memory disabled, failed to get system RAM";
+    VLOG_WARNING << "Mapped host memory disabled, failed to get system RAM";
     map_host_limit = 0;
   }
 
@@ -487,8 +487,8 @@ void HIPDevice::init_host_memory()
   device_working_headroom = 32 * 1024 * 1024LL;   // 32MB
   device_texture_headroom = 128 * 1024 * 1024LL;  // 128MB
 
-  VLOG(1) << "Mapped host memory limit set to " << string_human_readable_number(map_host_limit)
-          << " bytes. (" << string_human_readable_size(map_host_limit) << ")";
+  VLOG_INFO << "Mapped host memory limit set to " << string_human_readable_number(map_host_limit)
+            << " bytes. (" << string_human_readable_size(map_host_limit) << ")";
 }
 
 void HIPDevice::load_texture_info()
@@ -556,7 +556,7 @@ void HIPDevice::move_textures_to_host(size_t size, bool for_texture)
      * multiple HIP devices could be moving the memory. The
      * first one will do it, and the rest will adopt the pointer. */
     if (max_mem) {
-      VLOG(1) << "Move memory from device to host: " << max_mem->name;
+      VLOG_WORK << "Move memory from device to host: " << max_mem->name;
 
       static thread_mutex move_mutex;
       thread_scoped_lock lock(move_mutex);
@@ -658,9 +658,9 @@ HIPDevice::HIPMem *HIPDevice::generic_alloc(device_memory &mem, size_t pitch_pad
   }
 
   if (mem.name) {
-    VLOG(1) << "Buffer allocate: " << mem.name << ", "
-            << string_human_readable_number(mem.memory_size()) << " bytes. ("
-            << string_human_readable_size(mem.memory_size()) << ")" << status;
+    VLOG_WORK << "Buffer allocate: " << mem.name << ", "
+              << string_human_readable_number(mem.memory_size()) << " bytes. ("
+              << string_human_readable_size(mem.memory_size()) << ")" << status;
   }
 
   mem.device_pointer = (device_ptr)device_pointer;
@@ -966,9 +966,9 @@ void HIPDevice::tex_alloc(device_texture &mem)
     desc.NumChannels = mem.data_elements;
     desc.Flags = 0;
 
-    VLOG(1) << "Array 3D allocate: " << mem.name << ", "
-            << string_human_readable_number(mem.memory_size()) << " bytes. ("
-            << string_human_readable_size(mem.memory_size()) << ")";
+    VLOG_WORK << "Array 3D allocate: " << mem.name << ", "
+              << string_human_readable_number(mem.memory_size()) << " bytes. ("
+              << string_human_readable_size(mem.memory_size()) << ")";
 
     hip_assert(hipArray3DCreate((hArray *)&array_3d, &desc));
 
diff --git a/intern/cycles/device/hip/queue.cpp b/intern/cycles/device/hip/queue.cpp
index 6c2c2c29624..8b3d963a32f 100644
--- a/intern/cycles/device/hip/queue.cpp
+++ b/intern/cycles/device/hip/queue.cpp
@@ -39,12 +39,12 @@ int HIPDeviceQueue::num_concurrent_states(const size_t state_size) const
       num_states = max((int)(num_states * factor), 1024);
     }
     else {
-      VLOG(3) << "CYCLES_CONCURRENT_STATES_FACTOR evaluated to 0";
+      VLOG_DEVICE_STATS << "CYCLES_CONCURRENT_STATES_FACTOR evaluated to 0";
     }
   }
 
-  VLOG(3) << "GPU queue concurrent states: " << num_states << ", using up to "
-          << string_human_readable_size(num_states * state_size);
+  VLOG_DEVICE_STATS << "GPU queue concurrent states: " << num_states << ", using up to "
+                    << string_human_readable_size(num_states * state_size);
 
   return num_states;
 }
diff --git a/intern/cycles/device/metal/device_impl.mm b/intern/cycles/device/metal/device_impl.mm
index 086bf0af979..a0ac677beda 100644
--- a/intern/cycles/device/metal/device_impl.mm
+++ b/intern/cycles/device/metal/device_impl.mm
@@ -411,9 +411,9 @@ MetalDevice::MetalMem *MetalDevice::generic_alloc(device_memory &mem)
   }
 
   if (mem.name) {
-    VLOG(2) << "Buffer allocate: " << mem.name << ", "
-            << string_human_readable_number(mem.memory_size()) << " bytes. ("
-            << string_human_readable_size(mem.memory_size()) << ")";
+    VLOG_WORK << "Buffer allocate: " << mem.name << ", "
+              << string_human_readable_number(mem.memory_size()) << " bytes. ("
+              << string_human_readable_size(mem.memory_size()) << ")";
   }
 
   mem.device_size = metal_buffer.allocatedSize;
@@ -800,9 +800,9 @@ void MetalDevice::tex_alloc(device_texture &mem)
     desc.textureType = MTLTextureType3D;
     desc.depth = mem.data_depth;
 
-    VLOG(2) << "Texture 3D allocate: " << mem.name << ", "
-            << string_human_readable_number(mem.memory_size()) << " bytes. ("
-            << string_human_readable_size(mem.memory_size()) << ")";
+    VLOG_WORK << "Texture 3D allocate: " << mem.name << ", "
+              << string_human_readable_number(mem.memory_size()) << " bytes. ("
+              << string_human_readable_size(mem.memory_size()) << ")";
 
     mtlTexture = [mtlDevice newTextureWithDescriptor:desc];
     assert(mtlTexture);
@@ -834,9 +834,9 @@ void MetalDevice::tex_alloc(device_texture &mem)
     desc.storageMode = storage_mode;
     desc.usage = MTLTextureUsageShaderRead;
 
-    VLOG(2) << "Texture 2D allocate: " << mem.name << ", "
-            << string_human_readable_number(mem.memory_size()) << " bytes. ("
-            << string_human_readable_size(mem.memory_size()) << ")";
+    VLOG_WORK << "Texture 2D allocate: " << mem.name << ", "
+              << string_human_readable_number(mem.memory_size()) << " bytes. ("
+              << string_human_readable_size(mem.memory_size()) << ")";
 
     mtlTexture = [mtlDevice newTextureWithDescriptor:desc];
     assert(mtlTexture);
diff --git a/intern/cycles/device/metal/queue.h b/intern/cycles/device/metal/queue.h
index de20514de0b..b0bd487c86d 100644
--- a/intern/cycles/device/metal/queue.h
+++ b/intern/cycles/device/metal/queue.h
@@ -38,45 +38,50 @@ class MetalDeviceQueue : public DeviceQueue {
   virtual void copy_from_device(device_memory &mem) override;
 
  protected:
+  void setup_capture();
+  void update_capture(DeviceKernel kernel);
+  void begin_capture();
+  void end_capture();
   void prepare_resources(DeviceKernel kernel);
 
   id<MTLComputeCommandEncoder> get_compute_encoder(DeviceKernel kernel);
   id<MTLBlitCommandEncoder> get_blit_encoder();
 
-  MetalDevice *metal_device;
-  MetalBufferPool temp_buffer_pool;
+  MetalDevice *metal_device_;
+  MetalBufferPool temp_buffer_pool_;
 
   API_AVAILABLE(macos(11.0), ios(14.0))
-  MTLCommandBufferDescriptor *command_buffer_desc = nullptr;
-  id<MTLDevice> mtlDevice = nil;
-  id<MTLCommandQueue> mtlCommandQueue = nil;
-  id<MTLCommandBuffer> mtlCommandBuffer = nil;
-  id<MTLComputeCommandEncoder> mtlComputeEncoder = nil;
-  id<MTLBlitCommandEncoder> mtlBlitEncoder = nil;
+  MTLCommandBufferDescriptor *command_buffer_desc_ = nullptr;
+  id<MTLDevice> mtlDevice_ = nil;
+  id<MTLCommandQueue> mtlCommandQueue_ = nil;
+  id<MTLCommandBuffer> mtlCommandBuffer_ = nil;
+  id<MTLComputeCommandEncoder> mtlComputeEncoder_ = nil;
+  id<MTLBlitCommandEncoder> mtlBlitEncoder_ = nil;
   API_AVAILABLE(macos(10.14), ios(14.0))
-  id<MTLSharedEvent> shared_event = nil;
+  id<MTLSharedEvent> shared_event_ = nil;
   API_AVAILABLE(macos(10.14), ios(14.0))
-  MTLSharedEventListener *shared_event_listener = nil;
+  MTLSharedEventListener *shared_event_listener_ = nil;
 
-  dispatch_queue_t event_queue;
-  dispatch_semaphore_t wait_semaphore;
+  dispatch_queue_t event_queue_;
+  dispatch_semaphore_t wait_semaphore_;
 
   struct CopyBack {
     void *host_pointer;
     void *gpu_mem;
     uint64_t size;
   };
-  std::vector<CopyBack> copy_back_mem;
+  std::vector<CopyBack> copy_back_mem_;
 
-  uint64_t shared_event_id;
-  uint64_t command_buffers_submitted = 0;
-  uint64_t command_buffers_completed = 0;
-  Stats &stats;
+  uint64_t shared_event_id_;
+  uint64_t command_buffers_submitted_ = 0;
+  uint64_t command_buffers_completed_ = 0;
+  Stats &stats_;
 
   void close_compute_encoder();
   void close_blit_encoder();
 
-  bool verbose_tracing = false;
+  bool verbose_tracing_ = false;
+  bool label_command_encoders_ = false;
 
   /* Per-kernel profiling (see CYCLES_METAL_PROFILING). */
 
@@ -85,28 +90,30 @@ class MetalDeviceQueue : public DeviceQueue {
     int work_size;
     uint64_t timing_id;
   };
-  std::vector<TimingData> command_encoder_labels;
-  id<MTLSharedEvent> timing_shared_event = nil;
-  uint64_t timing_shared_event_id;
-  uint64_t command_buffer_start_timing_id;
+  std::vector<TimingData> command_encoder_labels_;
+  API_AVAILABLE(macos(10.14), ios(14.0))
+  id<MTLSharedEvent> timing_shared_event_ = nil;
+  uint64_t timing_shared_event_id_;
+  uint64_t command_buffer_start_timing_id_;
 
   struct TimingStats {
     double total_time = 0.0;
     uint64_t total_work_size = 0;
     uint64_t num_dispatches = 0;
   };
-  TimingStats timing_stats[DEVICE_KERNEL_NUM];
-  double last_completion_time = 0.0;
+  TimingStats timing_stats_[DEVICE_KERNEL_NUM];
+  double last_completion_time_ = 0.0;
 
   /* .gputrace capture (see CYCLES_DEBUG_METAL_CAPTURE_...). */
 
-  id<MTLCaptureScope> mtlCaptureScope = nil;
-  DeviceKernel capture_kernel;
-  int capture_dispatch = 0;
-  int capture_dispatch_counter = 0;
-  bool is_capturing = false;
-  bool is_capturing_to_disk = false;
-  bool has_captured_to_disk = false;
+  id<MTLCaptureScope> mtlCaptureScope_ = nil;
+  DeviceKernel capture_kernel_;
+  int capture_dispatch_counter_ = 0;
+  bool capture_samples_ = false;
+  int capture_reset_counter_ = 0;
+  bool is_capturing_ = false;
+  bool is_capturing_to_disk_ = false;
+  bool has_captured_to_disk_ = false;
 };
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/device/metal/queue.mm b/intern/cycles/device/metal/queue.mm
index 8b2d5d81859..55db7c5afce 100644
--- a/intern/cycles/device/metal/queue.mm
+++ b/intern/cycles/device/metal/queue.mm
@@ -17,79 +17,180 @@ CCL_NAMESPACE_BEGIN
 /* MetalDeviceQueue */
 
 MetalDeviceQueue::MetalDeviceQueue(MetalDevice *device)
-    : DeviceQueue(device), metal_device(device), stats(device->stats)
+    : DeviceQueue(device), metal_device_(device), stats_(device->stats)
 {
   if (@available(macos 11.0, *)) {
-    command_buffer_desc = [[MTLCommandBufferDescriptor alloc] init];
-    command_buffer_desc.errorOptions = MTLCommandBufferErrorOptionEncoderExecutionStatus;
+    command_buffer_desc_ = [[MTLCommandBufferDescriptor alloc] init];
+    command_buffer_desc_.errorOptions = MTLCommandBufferErrorOptionEncoderExecutionStatus;
   }
 
-  mtlDevice = device->mtlDevice;
-  mtlCommandQueue = [mtlDevice newCommandQueue];
+  mtlDevice_ = device->mtlDevice;
+  mtlCommandQueue_ = [mtlDevice_ newCommandQueue];
 
   if (@available(macos 10.14, *)) {
-    shared_event = [mtlDevice newSharedEvent];
-    shared_event_id = 1;
+    shared_event_ = [mtlDevice_ newSharedEvent];
+    shared_event_id_ = 1;
 
     /* Shareable event listener */
-    event_queue = dispatch_queue_create("com.cycles.metal.event_queue", NULL);
-    shared_event_listener = [[MTLSharedEventListener alloc] initWithDispatchQueue:event_queue];
+    event_queue_ = dispatch_queue_create("com.cycles.metal.event_queue", NULL);
+    shared_event_listener_ = [[MTLSharedEventListener alloc] initWithDispatchQueue:event_queue_];
   }
 
-  wait_semaphore = dispatch_semaphore_create(0);
+  wait_semaphore_ = dispatch_semaphore_create(0);
 
   if (@available(macos 10.14, *)) {
     if (getenv("CYCLES_METAL_PROFILING")) {
       /* Enable per-kernel timing breakdown (shown at end of render). */
-      timing_shared_event = [mtlDevice newSharedEvent];
+      timing_shared_event_ = [mtlDevice_ newSharedEvent];
+      label_command_encoders_ = true;
     }
     if (getenv("CYCLES_METAL_DEBUG")) {
       /* Enable very verbose tracing (shows every dispatch). */
-      verbose_tracing = true;
+      verbose_tracing_ = true;
+      label_command_encoders_ = true;
     }
-    timing_shared_event_id = 1;
+    timing_shared_event_id_ = 1;
   }
 
-  capture_kernel = DeviceKernel(-1);
+  setup_capture();
+}
+
+void MetalDeviceQueue::setup_capture()
+{
+  capture_kernel_ = DeviceKernel(-1);
+
   if (auto capture_kernel_str = getenv("CYCLES_DEBUG_METAL_CAPTURE_KERNEL")) {
-    /* Enable .gputrace capture for the specified DeviceKernel. */
-    MTLCaptureManager *captureManager = [MTLCaptureManager sharedCaptureManager];
-    mtlCaptureScope = [captureManager newCaptureScopeWithDevice:mtlDevice];
-    mtlCaptureScope.label = [NSString stringWithFormat:@"Cycles kernel dispatch"];
-    [captureManager setDefaultCaptureScope:mtlCaptureScope];
+    /* CYCLES_DEBUG_METAL_CAPTURE_KERNEL captures a single dispatch of the specified kernel. */
+    capture_kernel_ = DeviceKernel(atoi(capture_kernel_str));
+    printf("Capture kernel: %d = %s\n", capture_kernel_, device_kernel_as_string(capture_kernel_));
 
-    capture_dispatch = -1;
+    capture_dispatch_counter_ = 0;
     if (auto capture_dispatch_str = getenv("CYCLES_DEBUG_METAL_CAPTURE_DISPATCH")) {
-      capture_dispatch = atoi(capture_dispatch_str);
-      capture_dispatch_counter = 0;
+      capture_dispatch_counter_ = atoi(capture_dispatch_str);
+
+      printf("Capture dispatch number %d\n", capture_dispatch_counter_);
+    }
+  }
+  else if (auto capture_samples_str = getenv("CYCLES_DEBUG_METAL_CAPTURE_SAMPLES")) {
+    /* CYCLES_DEBUG_METAL_CAPTURE_SAMPLES captures a block of dispatches from reset#(N) to
+     * reset#(N+1). */
+    capture_samples_ = true;
+    capture_reset_counter_ = atoi(capture_samples_str);
+
+    capture_dispatch_counter_ = INT_MAX;
+    if (auto capture_limit_str = getenv("CYCLES_DEBUG_METAL_CAPTURE_LIMIT")) {
+      /* CYCLES_DEBUG_METAL_CAPTURE_LIMIT sets the maximum number of dispatches to capture. */
+      capture_dispatch_counter_ = atoi(capture_limit_str);
     }
 
-    capture_kernel = DeviceKernel(atoi(capture_kernel_str));
-    printf("Capture kernel: %d = %s\n", capture_kernel, device_kernel_as_string(capture_kernel));
+    printf("Capturing sample block %d (dispatch limit: %d)\n",
+           capture_reset_counter_,
+           capture_dispatch_counter_);
+  }
+  else {
+    /* No capturing requested. */
+    return;
+  }
 
-    if (auto capture_url = getenv("CYCLES_DEBUG_METAL_CAPTURE_URL")) {
-      if (@available(macos 10.15, *)) {
-        if ([captureManager supportsDestination:MTLCaptureDestinationGPUTraceDocument]) {
+  /* Enable .gputrace capture for the specified DeviceKernel. */
+  MTLCaptureManager *captureManager = [MTLCaptureManager sharedCaptureManager];
+  mtlCaptureScope_ = [captureManager newCaptureScopeWithDevice:mtlDevice_];
+  mtlCaptureScope_.label = [NSString stringWithFormat:@"Cycles kernel dispatch"];
+  [captureManager setDefaultCaptureScope:mtlCaptureScope_];
 
-          MTLCaptureDescriptor *captureDescriptor = [[MTLCaptureDescriptor alloc] init];
-          captureDescriptor.captureObject = mtlCaptureScope;
-          captureDescriptor.destination = MTLCaptureDestinationGPUTraceDocument;
-          captureDescriptor.outputURL = [NSURL fileURLWithPath:@(capture_url)];
+  label_command_encoders_ = true;
 
-          NSError *error;
-          if (![captureManager startCaptureWithDescriptor:captureDescriptor error:&error]) {
-            NSString *err = [error localizedDescription];
-            printf("Start capture failed: %s\n", [err UTF8String]);
-          }
-          else {
-            printf("Capture started (URL: %s)\n", capture_url);
-            is_capturing_to_disk = true;
-          }
+  if (auto capture_url = getenv("CYCLES_DEBUG_METAL_CAPTURE_URL")) {
+    if (@available(macos 10.15, *)) {
+      if ([captureManager supportsDestination:MTLCaptureDestinationGPUTraceDocument]) {
+
+        MTLCaptureDescriptor *captureDescriptor = [[MTLCaptureDescriptor alloc] init];
+        captureDescriptor.captureObject = mtlCaptureScope_;
+        captureDescriptor.destination = MTLCaptureDestinationGPUTraceDocument;
+        captureDescriptor.outputURL = [NSURL fileURLWithPath:@(capture_url)];
+
+        NSError *error;
+        if (![captureManager startCaptureWithDescriptor:captureDescriptor error:&error]) {
+          NSString *err = [error localizedDescription];
+          printf("Start capture failed: %s\n", [err UTF8String]);
         }
         else {
-          printf("Capture to file is not supported\n");
+          printf("Capture started (URL: %s)\n", capture_url);
+          is_capturing_to_disk_ = true;
         }
       }
+      else {
+        printf("Capture to file is not supported\n");
+      }
+    }
+  }
+}
+
+void MetalDeviceQueue::update_capture(DeviceKernel kernel)
+{
+  /* Handle capture end triggers. */
+  if (is_capturing_) {
+    capture_dispatch_counter_ -= 1;
+    if (capture_dispatch_counter_ <= 0 || kernel == DEVICE_KERNEL_INTEGRATOR_RESET) {
+      /* End capture if we've hit the dispatch limit or we hit a "reset". */
+      end_capture();
+    }
+    return;
+  }
+
+  if (capture_dispatch_counter_ < 0) {
+    /* We finished capturing. */
+    return;
+  }
+
+  /* Handle single-capture start trigger. */
+  if (kernel == capture_kernel_) {
+    /* Start capturing when the we hit the Nth dispatch of the specified kernel. */
+    if (capture_dispatch_counter_ == 0) {
+      begin_capture();
+    }
+    capture_dispatch_counter_ -= 1;
+    return;
+  }
+
+  /* Handle multi-capture start trigger. */
+  if (capture_samples_) {
+    /* Start capturing when the reset countdown is at 0. */
+    if (capture_reset_counter_ == 0) {
+      begin_capture();
+    }
+
+    if (kernel == DEVICE_KERNEL_INTEGRATOR_RESET) {
+      capture_reset_counter_ -= 1;
+    }
+    return;
+  }
+}
+
+void MetalDeviceQueue::begin_capture()
+{
+  /* Start gputrace capture. */
+  if (mtlCommandBuffer_) {
+    synchronize();
+  }
+  [mtlCaptureScope_ beginScope];
+  printf("[mtlCaptureScope_ beginScope]\n");
+  is_capturing_ = true;
+}
+
+void MetalDeviceQueue::end_capture()
+{
+  [mtlCaptureScope_ endScope];
+  is_capturing_ = false;
+  printf("[mtlCaptureScope_ endScope]\n");
+
+  if (is_capturing_to_disk_) {
+    if (@available(macos 10.15, *)) {
+      [[MTLCaptureManager sharedCaptureManager] stopCapture];
+      has_captured_to_disk_ = true;
+      is_capturing_to_disk_ = false;
+      is_capturing_ = false;
+      printf("Capture stopped\n");
     }
   }
 }
@@ -98,34 +199,32 @@ MetalDeviceQueue::~MetalDeviceQueue()
 {
   /* Tidying up here isn't really practical - we should expect and require the work
    * queue to be empty here. */
-  assert(mtlCommandBuffer == nil);
-  assert(command_buffers_submitted == command_buffers_completed);
+  assert(mtlCommandBuffer_ == nil);
+  assert(command_buffers_submitted_ == command_buffers_completed_);
 
   if (@available(macos 10.14, *)) {
-    [shared_event_listener release];
-    [shared_event release];
+    [shared_event_listener_ release];
+    [shared_event_ release];
   }
 
   if (@available(macos 11.0, *)) {
-    [command_buffer_desc release];
+    [command_buffer_desc_ release];
   }
-  if (mtlCommandQueue) {
-    [mtlCommandQueue release];
-    mtlCommandQueue = nil;
+  if (mtlCommandQueue_) {
+    [mtlCommandQueue_ release];
+    mtlCommandQueue_ = nil;
   }
 
-  if (mtlCaptureScope) {
-    [mtlCaptureScope release];
+  if (mtlCaptureScope_) {
+    [mtlCaptureScope_ release];
   }
 
   double total_time = 0.0;
 
   /* Show per-kernel timings, if gathered (see CYCLES_METAL_PROFILING). */
-  int64_t total_work_size = 0;
   int64_t num_dispatches = 0;
-  for (auto &stat : timing_stats) {
+  for (auto &stat : timing_stats_) {
     total_time += stat.total_time;
-    total_work_size += stat.total_work_size;
     num_dispatches += stat.num_dispatches;
   }
 
@@ -142,7 +241,7 @@ MetalDeviceQueue::~MetalDeviceQueue()
     printf("%s\n%s\n%s\n", divider.c_str(), header.c_str(), divider.c_str());
 
     for (size_t i = 0; i < DEVICE_KERNEL_NUM; i++) {
-      auto &stat = timing_stats[i];
+      auto &stat = timing_stats_[i];
       if (stat.num_dispatches > 0) {
         printf("%-40s %16s %12s %12s %6.2fs %6.2f%%\n",
                device_kernel_as_string(DeviceKernel(i)),
@@ -171,10 +270,10 @@ int MetalDeviceQueue::num_concurrent_states(const size_t /*state_size*/) const
   /* TODO: compute automatically. */
   /* TODO: must have at least num_threads_per_block. */
   int result = 1048576;
-  if (metal_device->device_vendor == METAL_GPU_AMD) {
+  if (metal_device_->device_vendor == METAL_GPU_AMD) {
     result *= 2;
   }
-  else if (metal_device->device_vendor == METAL_GPU_APPLE) {
+  else if (metal_device_->device_vendor == METAL_GPU_APPLE) {
     result *= 4;
   }
   return result;
@@ -185,10 +284,10 @@ int MetalDeviceQueue::num_concurrent_busy_states() const
   /* METAL_WIP */
   /* TODO: compute automatically. */
   int result = 65536;
-  if (metal_device->device_vendor == METAL_GPU_AMD) {
+  if (metal_device_->device_vendor == METAL_GPU_AMD) {
     result *= 2;
   }
-  else if (metal_device->device_vendor == METAL_GPU_APPLE) {
+  else if (metal_device_->device_vendor == METAL_GPU_APPLE) {
     result *= 4;
   }
   return result;
@@ -197,7 +296,7 @@ int MetalDeviceQueue::num_concurrent_busy_states() const
 void MetalDeviceQueue::init_execution()
 {
   /* Synchronize all textures and memory copies before executing task. */
-  metal_device->load_texture_info();
+  metal_device_->load_texture_info();
 
   synchronize();
 }
@@ -206,30 +305,21 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
                                const int work_size,
                                DeviceKernelArguments const &args)
 {
-  if (kernel == capture_kernel) {
-    if (capture_dispatch < 0 || capture_dispatch == capture_dispatch_counter) {
-      /* Start gputrace capture. */
-      if (mtlCommandBuffer) {
-        synchronize();
-      }
-      [mtlCaptureScope beginScope];
-      printf("[mtlCaptureScope beginScope]\n");
-      is_capturing = true;
-    }
-    capture_dispatch_counter += 1;
-  }
+  update_capture(kernel);
 
-  if (metal_device->have_error()) {
+  if (metal_device_->have_error()) {
     return false;
   }
 
-  VLOG(3) << "Metal queue launch " << device_kernel_as_string(kernel) << ", work_size "
-          << work_size;
+  VLOG_DEVICE_STATS << "Metal queue launch " << device_kernel_as_string(kernel) << ", work_size "
+                    << work_size;
 
   id<MTLComputeCommandEncoder> mtlComputeCommandEncoder = get_compute_encoder(kernel);
 
-  if (timing_shared_event) {
-    command_encoder_labels.push_back({kernel, work_size, timing_shared_event_id});
+  if (@available(macos 10.14, *)) {
+    if (timing_shared_event_) {
+      command_encoder_labels_.push_back({kernel, work_size, timing_shared_event_id_});
+    }
   }
 
   /* Determine size requirement for argument buffer. */
@@ -248,8 +338,8 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
 
   /* Metal ancillary bindless pointers. */
   size_t metal_offsets = arg_buffer_length;
-  arg_buffer_length += metal_device->mtlAncillaryArgEncoder.encodedLength;
-  arg_buffer_length = round_up(arg_buffer_length, metal_device->mtlAncillaryArgEncoder.alignment);
+  arg_buffer_length += metal_device_->mtlAncillaryArgEncoder.encodedLength;
+  arg_buffer_length = round_up(arg_buffer_length, metal_device_->mtlAncillaryArgEncoder.alignment);
 
   /* Temporary buffer used to prepare arg_buffer */
   uint8_t *init_arg_buffer = (uint8_t *)alloca(arg_buffer_length);
@@ -272,19 +362,23 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
                                         sizeof(IntegratorStateGPU);
   size_t plain_old_launch_data_size = sizeof(KernelParamsMetal) - plain_old_launch_data_offset;
   memcpy(init_arg_buffer + globals_offsets + plain_old_launch_data_offset,
-         (uint8_t *)&metal_device->launch_params + plain_old_launch_data_offset,
+         (uint8_t *)&metal_device_->launch_params + plain_old_launch_data_offset,
          plain_old_launch_data_size);
 
   /* Allocate an argument buffer. */
   MTLResourceOptions arg_buffer_options = MTLResourceStorageModeManaged;
   if (@available(macOS 11.0, *)) {
-    if ([mtlDevice hasUnifiedMemory]) {
+    if ([mtlDevice_ hasUnifiedMemory]) {
       arg_buffer_options = MTLResourceStorageModeShared;
     }
   }
 
-  id<MTLBuffer> arg_buffer = temp_buffer_pool.get_buffer(
-      mtlDevice, mtlCommandBuffer, arg_buffer_length, arg_buffer_options, init_arg_buffer, stats);
+  id<MTLBuffer> arg_buffer = temp_buffer_pool_.get_buffer(mtlDevice_,
+                                                          mtlCommandBuffer_,
+                                                          arg_buffer_length,
+                                                          arg_buffer_options,
+                                                          init_arg_buffer,
+                                                          stats_);
 
   /* Encode the pointer "enqueue" arguments */
   bytes_written = 0;
@@ -292,16 +386,16 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
     size_t size_in_bytes = args.sizes[i];
     bytes_written = round_up(bytes_written, size_in_bytes);
     if (args.types[i] == DeviceKernelArguments::POINTER) {
-      [metal_device->mtlBufferKernelParamsEncoder setArgumentBuffer:arg_buffer
-                                                             offset:bytes_written];
+      [metal_device_->mtlBufferKernelParamsEncoder setArgumentBuffer:arg_buffer
+                                                              offset:bytes_written];
       if (MetalDevice::MetalMem *mmem = *(MetalDevice::MetalMem **)args.values[i]) {
         [mtlComputeCommandEncoder useResource:mmem->mtlBuffer
                                         usage:MTLResourceUsageRead | MTLResourceUsageWrite];
-        [metal_device->mtlBufferKernelParamsEncoder setBuffer:mmem->mtlBuffer offset:0 atIndex:0];
+        [metal_device_->mtlBufferKernelParamsEncoder setBuffer:mmem->mtlBuffer offset:0 atIndex:0];
       }
       else {
         if (@available(macos 12.0, *)) {
-          [metal_device->mtlBufferKernelParamsEncoder setBuffer:nil offset:0 atIndex:0];
+          [metal_device_->mtlBufferKernelParamsEncoder setBuffer:nil offset:0 atIndex:0];
         }
       }
     }
@@ -309,9 +403,10 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
   }
 
   /* Encode KernelParamsMetal buffers */
-  [metal_device->mtlBufferKernelParamsEncoder setArgumentBuffer:arg_buffer offset:globals_offsets];
+  [metal_device_->mtlBufferKernelParamsEncoder setArgumentBuffer:arg_buffer
+                                                          offset:globals_offsets];
 
-  if (verbose_tracing || timing_shared_event || is_capturing) {
+  if (label_command_encoders_) {
     /* Add human-readable labels if we're doing any form of debugging / profiling. */
     mtlComputeCommandEncoder.label = [[NSString alloc]
         initWithFormat:@"Metal queue launch %s, work_size %d",
@@ -323,43 +418,43 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
   const size_t pointer_block_end = offsetof(KernelParamsMetal, __integrator_state) +
                                    sizeof(IntegratorStateGPU);
   for (size_t offset = 0; offset < pointer_block_end; offset += sizeof(device_ptr)) {
-    int pointer_index = offset / sizeof(device_ptr);
+    int pointer_index = int(offset / sizeof(device_ptr));
     MetalDevice::MetalMem *mmem = *(
-        MetalDevice::MetalMem **)((uint8_t *)&metal_device->launch_params + offset);
+        MetalDevice::MetalMem **)((uint8_t *)&metal_device_->launch_params + offset);
     if (mmem && mmem->mem && (mmem->mtlBuffer || mmem->mtlTexture)) {
-      [metal_device->mtlBufferKernelParamsEncoder setBuffer:mmem->mtlBuffer
-                                                     offset:0
-                                                    atIndex:pointer_index];
+      [metal_device_->mtlBufferKernelParamsEncoder setBuffer:mmem->mtlBuffer
+                                                      offset:0
+                                                     atIndex:pointer_index];
     }
     else {
       if (@available(macos 12.0, *)) {
-        [metal_device->mtlBufferKernelParamsEncoder setBuffer:nil offset:0 atIndex:pointer_index];
+        [metal_device_->mtlBufferKernelParamsEncoder setBuffer:nil offset:0 atIndex:pointer_index];
       }
     }
   }
   bytes_written = globals_offsets + sizeof(KernelParamsMetal);
 
-  const MetalKernelPipeline *metal_kernel_pso = MetalDeviceKernels::get_best_pipeline(metal_device,
-                                                                                      kernel);
+  const MetalKernelPipeline *metal_kernel_pso = MetalDeviceKernels::get_best_pipeline(
+      metal_device_, kernel);
   if (!metal_kernel_pso) {
-    metal_device->set_error(
+    metal_device_->set_error(
         string_printf("No MetalKernelPipeline for %s\n", device_kernel_as_string(kernel)));
     return false;
   }
 
   /* Encode ancillaries */
-  [metal_device->mtlAncillaryArgEncoder setArgumentBuffer:arg_buffer offset:metal_offsets];
-  [metal_device->mtlAncillaryArgEncoder setBuffer:metal_device->texture_bindings_2d
-                                           offset:0
-                                          atIndex:0];
-  [metal_device->mtlAncillaryArgEncoder setBuffer:metal_device->texture_bindings_3d
-                                           offset:0
-                                          atIndex:1];
+  [metal_device_->mtlAncillaryArgEncoder setArgumentBuffer:arg_buffer offset:metal_offsets];
+  [metal_device_->mtlAncillaryArgEncoder setBuffer:metal_device_->texture_bindings_2d
+                                            offset:0
+                                           atIndex:0];
+  [metal_device_->mtlAncillaryArgEncoder setBuffer:metal_device_->texture_bindings_3d
+                                            offset:0
+                                           atIndex:1];
   if (@available(macos 12.0, *)) {
-    if (metal_device->use_metalrt) {
-      if (metal_device->bvhMetalRT) {
-        id<MTLAccelerationStructure> accel_struct = metal_device->bvhMetalRT->accel_struct;
-        [metal_device->mtlAncillaryArgEncoder setAccelerationStructure:accel_struct atIndex:2];
+    if (metal_device_->use_metalrt) {
+      if (metal_device_->bvhMetalRT) {
+        id<MTLAccelerationStructure> accel_struct = metal_device_->bvhMetalRT->accel_struct;
+        [metal_device_->mtlAncillaryArgEncoder setAccelerationStructure:accel_struct atIndex:2];
       }
 
       for (int table = 0; table < METALRT_TABLE_NUM; table++) {
@@ -367,19 +462,19 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
           [metal_kernel_pso->intersection_func_table[table] setBuffer:arg_buffer
                                                                offset:globals_offsets
                                                               atIndex:1];
-          [metal_device->mtlAncillaryArgEncoder
+          [metal_device_->mtlAncillaryArgEncoder
               setIntersectionFunctionTable:metal_kernel_pso->intersection_func_table[table]
                                    atIndex:3 + table];
           [mtlComputeCommandEncoder useResource:metal_kernel_pso->intersection_func_table[table]
                                           usage:MTLResourceUsageRead];
         }
         else {
-          [metal_device->mtlAncillaryArgEncoder setIntersectionFunctionTable:nil
-                                                                     atIndex:3 + table];
+          [metal_device_->mtlAncillaryArgEncoder setIntersectionFunctionTable:nil
+                                                                      atIndex:3 + table];
         }
       }
     }
-    bytes_written = metal_offsets + metal_device->mtlAncillaryArgEncoder.encodedLength;
+    bytes_written = metal_offsets + metal_device_->mtlAncillaryArgEncoder.encodedLength;
   }
 
   if (arg_buffer.storageMode == MTLStorageModeManaged) {
@@ -390,10 +485,10 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
   [mtlComputeCommandEncoder setBuffer:arg_buffer offset:globals_offsets atIndex:1];
   [mtlComputeCommandEncoder setBuffer:arg_buffer offset:metal_offsets atIndex:2];
 
-  if (metal_device->use_metalrt) {
+  if (metal_device_->use_metalrt) {
     if (@available(macos 12.0, *)) {
 
-      auto bvhMetalRT = metal_device->bvhMetalRT;
+      auto bvhMetalRT = metal_device_->bvhMetalRT;
       switch (kernel) {
         case DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST:
         case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW:
@@ -435,7 +530,7 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
     case DEVICE_KERNEL_INTEGRATOR_COMPACT_SHADOW_PATHS_ARRAY:
       /* See parallel_active_index.h for why this amount of shared memory is needed.
        * Rounded up to 16 bytes for Metal */
-      shared_mem_bytes = round_up((num_threads_per_block + 1) * sizeof(int), 16);
+      shared_mem_bytes = (int)round_up((num_threads_per_block + 1) * sizeof(int), 16);
       [mtlComputeCommandEncoder setThreadgroupMemoryLength:shared_mem_bytes atIndex:0];
       break;
 
@@ -449,7 +544,7 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
   [mtlComputeCommandEncoder dispatchThreadgroups:size_threadgroups_per_dispatch
                            threadsPerThreadgroup:size_threads_per_threadgroup];
 
-  [mtlCommandBuffer addCompletedHandler:^(id<MTLCommandBuffer> command_buffer) {
+  [mtlCommandBuffer_ addCompletedHandler:^(id<MTLCommandBuffer> command_buffer) {
     NSString *kernel_name = metal_kernel_pso->function.label;
 
     /* Enhanced command buffer errors are only available in 11.0+ */
@@ -474,12 +569,12 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
     }
   }];
 
-  if (verbose_tracing || is_capturing) {
+  if (verbose_tracing_ || is_capturing_) {
     /* Force a sync we've enabled step-by-step verbose tracing or if we're capturing. */
     synchronize();
 
     /* Show queue counters and dispatch timing. */
-    if (verbose_tracing) {
+    if (verbose_tracing_) {
       if (kernel == DEVICE_KERNEL_INTEGRATOR_RESET) {
         printf(
             "_____________________________________.____________________.______________.___________"
@@ -489,9 +584,9 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
       printf("%-40s| %7d threads |%5.2fms | buckets [",
              device_kernel_as_string(kernel),
              work_size,
-             last_completion_time * 1000.0);
-      std::lock_guard<std::recursive_mutex> lock(metal_device->metal_mem_map_mutex);
-      for (auto &it : metal_device->metal_mem_map) {
+             last_completion_time_ * 1000.0);
+      std::lock_guard<std::recursive_mutex> lock(metal_device_->metal_mem_map_mutex);
+      for (auto &it : metal_device_->metal_mem_map) {
         const string c_integrator_queue_counter = "integrator_queue_counter";
         if (it.first->name == c_integrator_queue_counter) {
           /* Workaround "device_copy_from" being protected. */
@@ -515,89 +610,76 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
     }
   }
 
-  return !(metal_device->have_error());
+  return !(metal_device_->have_error());
 }
 
 bool MetalDeviceQueue::synchronize()
 {
-  if (has_captured_to_disk || metal_device->have_error()) {
+  if (has_captured_to_disk_ || metal_device_->have_error()) {
     return false;
   }
 
-  if (mtlComputeEncoder) {
+  if (mtlComputeEncoder_) {
     close_compute_encoder();
   }
   close_blit_encoder();
 
-  if (mtlCommandBuffer) {
+  if (mtlCommandBuffer_) {
     scoped_timer timer;
-    if (timing_shared_event) {
-      /* For per-kernel timing, add event handlers to measure & accumulate dispatch times. */
-      __block double completion_time = 0;
-      for (uint64_t i = command_buffer_start_timing_id; i < timing_shared_event_id; i++) {
-        [timing_shared_event notifyListener:shared_event_listener
-                                    atValue:i
-                                      block:^(id<MTLSharedEvent> sharedEvent, uint64_t value) {
-                                        completion_time = timer.get_time() - completion_time;
-                                        last_completion_time = completion_time;
-                                        for (auto label : command_encoder_labels) {
-                                          if (label.timing_id == value) {
-                                            TimingStats &stat = timing_stats[label.kernel];
-                                            stat.num_dispatches++;
-                                            stat.total_time += completion_time;
-                                            stat.total_work_size += label.work_size;
-                                          }
-                                        }
-                                      }];
+
+    if (@available(macos 10.14, *)) {
+      if (timing_shared_event_) {
+        /* For per-kernel timing, add event handlers to measure & accumulate dispatch times. */
+        __block double completion_time = 0;
+        for (uint64_t i = command_buffer_start_timing_id_; i < timing_shared_event_id_; i++) {
+          [timing_shared_event_ notifyListener:shared_event_listener_
+                                       atValue:i
+                                         block:^(id<MTLSharedEvent> sharedEvent, uint64_t value) {
+                                           completion_time = timer.get_time() - completion_time;
+                                           last_completion_time_ = completion_time;
+                                           for (auto label : command_encoder_labels_) {
+                                             if (label.timing_id == value) {
+                                               TimingStats &stat = timing_stats_[label.kernel];
+                                               stat.num_dispatches++;
+                                               stat.total_time += completion_time;
+                                               stat.total_work_size += label.work_size;
+                                             }
+                                           }
+                                         }];
+        }
       }
     }
 
-    uint64_t shared_event_id = this->shared_event_id++;
+    uint64_t shared_event_id_ = this->shared_event_id_++;
 
     if (@available(macos 10.14, *)) {
-      __block dispatch_semaphore_t block_sema = wait_semaphore;
-      [shared_event notifyListener:shared_event_listener
-                           atValue:shared_event_id
-                             block:^(id<MTLSharedEvent> sharedEvent, uint64_t value) {
-                               dispatch_semaphore_signal(block_sema);
-                             }];
-
-      [mtlCommandBuffer encodeSignalEvent:shared_event value:shared_event_id];
-      [mtlCommandBuffer commit];
-      dispatch_semaphore_wait(wait_semaphore, DISPATCH_TIME_FOREVER);
-    }
-
-    if (is_capturing) {
-      [mtlCaptureScope endScope];
-      is_capturing = false;
-      printf("[mtlCaptureScope endScope]\n");
-
-      if (is_capturing_to_disk) {
-        if (@available(macos 10.15, *)) {
-          [[MTLCaptureManager sharedCaptureManager] stopCapture];
-          has_captured_to_disk = true;
-          is_capturing_to_disk = false;
-          is_capturing = false;
-          printf("Capture stopped\n");
-        }
-      }
+      __block dispatch_semaphore_t block_sema = wait_semaphore_;
+      [shared_event_ notifyListener:shared_event_listener_
+                            atValue:shared_event_id_
+                              block:^(id<MTLSharedEvent> sharedEvent, uint64_t value) {
+                                dispatch_semaphore_signal(block_sema);
+                              }];
+
+      [mtlCommandBuffer_ encodeSignalEvent:shared_event_ value:shared_event_id_];
+      [mtlCommandBuffer_ commit];
+      dispatch_semaphore_wait(wait_semaphore_, DISPATCH_TIME_FOREVER);
     }
 
-    [mtlCommandBuffer release];
+    [mtlCommandBuffer_ release];
 
-    for (const CopyBack &mmem : copy_back_mem) {
+    for (const CopyBack &mmem : copy_back_mem_) {
       memcpy((uchar *)mmem.host_pointer, (uchar *)mmem.gpu_mem, mmem.size);
     }
-    copy_back_mem.clear();
+    copy_back_mem_.clear();
 
-    temp_buffer_pool.process_command_buffer_completion(mtlCommandBuffer);
-    metal_device->flush_delayed_free_list();
+    temp_buffer_pool_.process_command_buffer_completion(mtlCommandBuffer_);
+    metal_device_->flush_delayed_free_list();
 
-    mtlCommandBuffer = nil;
-    command_encoder_labels.clear();
+    mtlCommandBuffer_ = nil;
+    command_encoder_labels_.clear();
   }
 
-  return !(metal_device->have_error());
+  return !(metal_device_->have_error());
 }
 
 void MetalDeviceQueue::zero_to_device(device_memory &mem)
@@ -610,20 +692,20 @@ void MetalDeviceQueue::zero_to_device(device_memory &mem)
 
   /* Allocate on demand. */
   if (mem.device_pointer == 0) {
-    metal_device->mem_alloc(mem);
+    metal_device_->mem_alloc(mem);
   }
 
   /* Zero memory on device. */
   assert(mem.device_pointer != 0);
 
-  std::lock_guard<std::recursive_mutex> lock(metal_device->metal_mem_map_mutex);
-  MetalDevice::MetalMem &mmem = *metal_device->metal_mem_map.at(&mem);
+  std::lock_guard<std::recursive_mutex> lock(metal_device_->metal_mem_map_mutex);
+  MetalDevice::MetalMem &mmem = *metal_device_->metal_mem_map.at(&mem);
   if (mmem.mtlBuffer) {
     id<MTLBlitCommandEncoder> blitEncoder = get_blit_encoder();
     [blitEncoder fillBuffer:mmem.mtlBuffer range:NSMakeRange(mmem.offset, mmem.size) value:0];
   }
   else {
-    metal_device->mem_zero(mem);
+    metal_device_->mem_zero(mem);
   }
 }
 
@@ -635,15 +717,15 @@ void MetalDeviceQueue::copy_to_device(device_memory &mem)
 
   /* Allocate on demand. */
   if (mem.device_pointer == 0) {
-    metal_device->mem_alloc(mem);
+    metal_device_->mem_alloc(mem);
   }
 
   assert(mem.device_pointer != 0);
   assert(mem.host_pointer != nullptr);
 
-  std::lock_guard<std::recursive_mutex> lock(metal_device->metal_mem_map_mutex);
-  auto result = metal_device->metal_mem_map.find(&mem);
-  if (result != metal_device->metal_mem_map.end()) {
+  std::lock_guard<std::recursive_mutex> lock(metal_device_->metal_mem_map_mutex);
+  auto result = metal_device_->metal_mem_map.find(&mem);
+  if (result != metal_device_->metal_mem_map.end()) {
     if (mem.host_pointer == mem.shared_pointer) {
       return;
     }
@@ -651,12 +733,12 @@ void MetalDeviceQueue::copy_to_device(device_memory &mem)
     MetalDevice::MetalMem &mmem = *result->second;
     id<MTLBlitCommandEncoder> blitEncoder = get_blit_encoder();
 
-    id<MTLBuffer> buffer = temp_buffer_pool.get_buffer(mtlDevice,
-                                                       mtlCommandBuffer,
-                                                       mmem.size,
-                                                       MTLResourceStorageModeShared,
-                                                       mem.host_pointer,
-                                                       stats);
+    id<MTLBuffer> buffer = temp_buffer_pool_.get_buffer(mtlDevice_,
+                                                        mtlCommandBuffer_,
+                                                        mmem.size,
+                                                        MTLResourceStorageModeShared,
+                                                        mem.host_pointer,
+                                                        stats_);
 
     [blitEncoder copyFromBuffer:buffer
                    sourceOffset:0
@@ -665,7 +747,7 @@ void MetalDeviceQueue::copy_to_device(device_memory &mem)
                            size:mmem.size];
   }
   else {
-    metal_device->mem_copy_to(mem);
+    metal_device_->mem_copy_to(mem);
   }
 }
 
@@ -680,8 +762,8 @@ void MetalDeviceQueue::copy_from_device(device_memory &mem)
   assert(mem.device_pointer != 0);
   assert(mem.host_pointer != nullptr);
 
-  std::lock_guard<std::recursive_mutex> lock(metal_device->metal_mem_map_mutex);
-  MetalDevice::MetalMem &mmem = *metal_device->metal_mem_map.at(&mem);
+  std::lock_guard<std::recursive_mutex> lock(metal_device_->metal_mem_map_mutex);
+  MetalDevice::MetalMem &mmem = *metal_device_->metal_mem_map.at(&mem);
   if (mmem.mtlBuffer) {
     const size_t size = mem.memory_size();
 
@@ -691,8 +773,8 @@ void MetalDeviceQueue::copy_from_device(device_memory &mem)
         [blitEncoder synchronizeResource:mmem.mtlBuffer];
       }
       if (mem.host_pointer != mmem.hostPtr) {
-        if (mtlCommandBuffer) {
-          copy_back_mem.push_back({mem.host_pointer, mmem.hostPtr, size});
+        if (mtlCommandBuffer_) {
+          copy_back_mem_.push_back({mem.host_pointer, mmem.hostPtr, size});
         }
         else {
           memcpy((uchar *)mem.host_pointer, (uchar *)mmem.hostPtr, size);
@@ -704,16 +786,16 @@ void MetalDeviceQueue::copy_from_device(device_memory &mem)
     }
   }
   else {
-    metal_device->mem_copy_from(mem);
+    metal_device_->mem_copy_from(mem);
   }
 }
 
 void MetalDeviceQueue::prepare_resources(DeviceKernel kernel)
 {
-  std::lock_guard<std::recursive_mutex> lock(metal_device->metal_mem_map_mutex);
+  std::lock_guard<std::recursive_mutex> lock(metal_device_->metal_mem_map_mutex);
 
   /* declare resource usage */
-  for (auto &it : metal_device->metal_mem_map) {
+  for (auto &it : metal_device_->metal_mem_map) {
     device_memory *mem = it.first;
 
     MTLResourceUsage usage = MTLResourceUsageRead;
@@ -723,97 +805,99 @@ void MetalDeviceQueue::prepare_resources(DeviceKernel kernel)
 
     if (it.second->mtlBuffer) {
       /* METAL_WIP - use array version (i.e. useResources) */
-      [mtlComputeEncoder useResource:it.second->mtlBuffer usage:usage];
+      [mtlComputeEncoder_ useResource:it.second->mtlBuffer usage:usage];
     }
     else if (it.second->mtlTexture) {
       /* METAL_WIP - use array version (i.e. useResources) */
-      [mtlComputeEncoder useResource:it.second->mtlTexture usage:usage | MTLResourceUsageSample];
+      [mtlComputeEncoder_ useResource:it.second->mtlTexture usage:usage | MTLResourceUsageSample];
     }
   }
 
   /* ancillaries */
-  [mtlComputeEncoder useResource:metal_device->texture_bindings_2d usage:MTLResourceUsageRead];
-  [mtlComputeEncoder useResource:metal_device->texture_bindings_3d usage:MTLResourceUsageRead];
+  [mtlComputeEncoder_ useResource:metal_device_->texture_bindings_2d usage:MTLResourceUsageRead];
+  [mtlComputeEncoder_ useResource:metal_device_->texture_bindings_3d usage:MTLResourceUsageRead];
 }
 
 id<MTLComputeCommandEncoder> MetalDeviceQueue::get_compute_encoder(DeviceKernel kernel)
 {
   bool concurrent = (kernel < DEVICE_KERNEL_INTEGRATOR_NUM);
 
-  if (timing_shared_event) {
-    /* Close the current encoder to ensure we're able to capture per-encoder timing data. */
-    if (mtlComputeEncoder) {
-      close_compute_encoder();
+  if (@available(macos 10.14, *)) {
+    if (timing_shared_event_) {
+      /* Close the current encoder to ensure we're able to capture per-encoder timing data. */
+      if (mtlComputeEncoder_) {
+        close_compute_encoder();
+      }
     }
-  }
 
-  if (@available(macos 10.14, *)) {
-    if (mtlComputeEncoder) {
-      if (mtlComputeEncoder.dispatchType == concurrent ? MTLDispatchTypeConcurrent :
-                                                         MTLDispatchTypeSerial) {
+    if (mtlComputeEncoder_) {
+      if (mtlComputeEncoder_.dispatchType == concurrent ? MTLDispatchTypeConcurrent :
+                                                          MTLDispatchTypeSerial) {
         /* declare usage of MTLBuffers etc */
         prepare_resources(kernel);
 
-        return mtlComputeEncoder;
+        return mtlComputeEncoder_;
       }
       close_compute_encoder();
     }
 
     close_blit_encoder();
 
-    if (!mtlCommandBuffer) {
-      mtlCommandBuffer = [mtlCommandQueue commandBuffer];
-      [mtlCommandBuffer retain];
+    if (!mtlCommandBuffer_) {
+      mtlCommandBuffer_ = [mtlCommandQueue_ commandBuffer];
+      [mtlCommandBuffer_ retain];
     }
 
-    mtlComputeEncoder = [mtlCommandBuffer
+    mtlComputeEncoder_ = [mtlCommandBuffer_
         computeCommandEncoderWithDispatchType:concurrent ? MTLDispatchTypeConcurrent :
                                                            MTLDispatchTypeSerial];
 
-    [mtlComputeEncoder setLabel:@(device_kernel_as_string(kernel))];
+    [mtlComputeEncoder_ setLabel:@(device_kernel_as_string(kernel))];
 
     /* declare usage of MTLBuffers etc */
     prepare_resources(kernel);
   }
 
-  return mtlComputeEncoder;
+  return mtlComputeEncoder_;
 }
 
 id<MTLBlitCommandEncoder> MetalDeviceQueue::get_blit_encoder()
 {
-  if (mtlBlitEncoder) {
-    return mtlBlitEncoder;
+  if (mtlBlitEncoder_) {
+    return mtlBlitEncoder_;
   }
 
-  if (mtlComputeEncoder) {
+  if (mtlComputeEncoder_) {
     close_compute_encoder();
   }
 
-  if (!mtlCommandBuffer) {
-    mtlCommandBuffer = [mtlCommandQueue commandBuffer];
-    [mtlCommandBuffer retain];
-    command_buffer_start_timing_id = timing_shared_event_id;
+  if (!mtlCommandBuffer_) {
+    mtlCommandBuffer_ = [mtlCommandQueue_ commandBuffer];
+    [mtlCommandBuffer_ retain];
+    command_buffer_start_timing_id_ = timing_shared_event_id_;
   }
 
-  mtlBlitEncoder = [mtlCommandBuffer blitCommandEncoder];
-  return mtlBlitEncoder;
+  mtlBlitEncoder_ = [mtlCommandBuffer_ blitCommandEncoder];
+  return mtlBlitEncoder_;
 }
 
 void MetalDeviceQueue::close_compute_encoder()
 {
-  [mtlComputeEncoder endEncoding];
-  mtlComputeEncoder = nil;
+  [mtlComputeEncoder_ endEncoding];
+  mtlComputeEncoder_ = nil;
 
-  if (timing_shared_event) {
-    [mtlCommandBuffer encodeSignalEvent:timing_shared_event value:timing_shared_event_id++];
+  if (@available(macos 10.14, *)) {
+    if (timing_shared_event_) {
+      [mtlCommandBuffer_ encodeSignalEvent:timing_shared_event_ value:timing_shared_event_id_++];
+    }
   }
 }
 
 void MetalDeviceQueue::close_blit_encoder()
 {
-  if (mtlBlitEncoder) {
-    [mtlBlitEncoder endEncoding];
-    mtlBlitEncoder = nil;
+  if (mtlBlitEncoder_) {
+    [mtlBlitEncoder_ endEncoding];
+    mtlBlitEncoder_ = nil;
   }
 }
 
diff --git a/intern/cycles/device/optix/device.cpp b/intern/cycles/device/optix/device.cpp
index 70810bae10d..68ca21374fd 100644
--- a/intern/cycles/device/optix/device.cpp
+++ b/intern/cycles/device/optix/device.cpp
@@ -31,12 +31,12 @@ bool device_optix_init()
   const OptixResult result = optixInit();
 
   if (result == OPTIX_ERROR_UNSUPPORTED_ABI_VERSION) {
-    VLOG(1) << "OptiX initialization failed because the installed NVIDIA driver is too old. "
-               "Please update to the latest driver first!";
+    VLOG_WARNING << "OptiX initialization failed because the installed NVIDIA driver is too old. "
+                    "Please update to the latest driver first!";
     return false;
   }
   else if (result != OPTIX_SUCCESS) {
-    VLOG(1) << "OptiX initialization failed with error code " << (unsigned int)result;
+    VLOG_WARNING << "OptiX initialization failed with error code " << (unsigned int)result;
     return false;
   }
 
diff --git a/intern/cycles/device/optix/device_impl.cpp b/intern/cycles/device/optix/device_impl.cpp
index 9ab9bbb59c5..53697db5c04 100644
--- a/intern/cycles/device/optix/device_impl.cpp
+++ b/intern/cycles/device/optix/device_impl.cpp
@@ -278,7 +278,7 @@ OptiXDevice::OptiXDevice(const DeviceInfo &info, Stats &stats, Profiler &profile
   };
 #  endif
   if (DebugFlags().optix.use_debug) {
-    VLOG(1) << "Using OptiX debug mode.";
+    VLOG_INFO << "Using OptiX debug mode.";
     options.validationMode = OPTIX_DEVICE_CONTEXT_VALIDATION_MODE_ALL;
   }
   optix_assert(optixDeviceContextCreate(cuContext, &options, &context));
@@ -1392,11 +1392,11 @@ bool OptiXDevice::build_optix_bvh(BVHOptiX *bvh,
       /* The build flags have to match the ones used to query the built-in curve intersection
          program (see optixBuiltinISModuleGet above) */
       build_input.type == OPTIX_BUILD_INPUT_TYPE_CURVES) {
-    VLOG(2) << "Using fast to trace OptiX BVH";
+    VLOG_INFO << "Using fast to trace OptiX BVH";
     options.buildFlags = OPTIX_BUILD_FLAG_PREFER_FAST_TRACE | OPTIX_BUILD_FLAG_ALLOW_COMPACTION;
   }
   else {
-    VLOG(2) << "Using fast to update OptiX BVH";
+    VLOG_INFO << "Using fast to update OptiX BVH";
     options.buildFlags = OPTIX_BUILD_FLAG_PREFER_FAST_BUILD | OPTIX_BUILD_FLAG_ALLOW_UPDATE;
   }
 
diff --git a/intern/cycles/device/queue.cpp b/intern/cycles/device/queue.cpp
index de65047ed6a..cc0cf0ccf84 100644
--- a/intern/cycles/device/queue.cpp
+++ b/intern/cycles/device/queue.cpp
@@ -19,7 +19,7 @@ DeviceQueue::DeviceQueue(Device *device)
 
 DeviceQueue::~DeviceQueue()
 {
-  if (VLOG_IS_ON(3)) {
+  if (VLOG_DEVICE_STATS_IS_ON) {
     /* Print kernel execution times sorted by time. */
     vector<pair<DeviceKernelMask, double>> stats_sorted;
     for (const auto &stat : stats_kernel_time_) {
@@ -32,17 +32,18 @@ DeviceQueue::~DeviceQueue()
            return a.second > b.second;
          });
 
-    VLOG(3) << "GPU queue stats:";
+    VLOG_DEVICE_STATS << "GPU queue stats:";
     for (const auto &[mask, time] : stats_sorted) {
-      VLOG(3) << "  " << std::setfill(' ') << std::setw(10) << std::fixed << std::setprecision(5)
-              << std::right << time << "s: " << device_kernel_mask_as_string(mask);
+      VLOG_DEVICE_STATS << "  " << std::setfill(' ') << std::setw(10) << std::fixed
+                        << std::setprecision(5) << std::right << time
+                        << "s: " << device_kernel_mask_as_string(mask);
     }
   }
 }
 
 void DeviceQueue::debug_init_execution()
 {
-  if (VLOG_IS_ON(3)) {
+  if (VLOG_DEVICE_STATS_IS_ON) {
     last_sync_time_ = time_dt();
   }
 
@@ -51,9 +52,9 @@ void DeviceQueue::debug_init_execution()
 
 void DeviceQueue::debug_enqueue(DeviceKernel kernel, const int work_size)
 {
-  if (VLOG_IS_ON(3)) {
-    VLOG(4) << "GPU queue launch " << device_kernel_as_string(kernel) << ", work_size "
-            << work_size;
+  if (VLOG_DEVICE_STATS_IS_ON) {
+    VLOG_DEVICE_STATS << "GPU queue launch " << device_kernel_as_string(kernel) << ", work_size "
+                      << work_size;
   }
 
   last_kernels_enqueued_ |= (uint64_t(1) << (uint64_t)kernel);
@@ -61,10 +62,10 @@ void DeviceQueue::debug_enqueue(DeviceKernel kernel, const int work_size)
 
 void DeviceQueue::debug_synchronize()
 {
-  if (VLOG_IS_ON(3)) {
+  if (VLOG_DEVICE_STATS_IS_ON) {
     const double new_time = time_dt();
     const double elapsed_time = new_time - last_sync_time_;
-    VLOG(4) << "GPU queue synchronize, elapsed " << std::setw(10) << elapsed_time << "s";
+    VLOG_DEVICE_STATS << "GPU queue synchronize, elapsed " << std::setw(10) << elapsed_time << "s";
 
     stats_kernel_time_[last_kernels_enqueued_] += elapsed_time;
 
diff --git a/intern/cycles/integrator/denoiser.cpp b/intern/cycles/integrator/denoiser.cpp
index 23ab825a4d2..94991d63e4c 100644
--- a/intern/cycles/integrator/denoiser.cpp
+++ b/intern/cycles/integrator/denoiser.cpp
@@ -58,8 +58,8 @@ bool Denoiser::load_kernels(Progress *progress)
     return false;
   }
 
-  VLOG(3) << "Will denoise on " << denoiser_device->info.description << " ("
-          << denoiser_device->info.id << ")";
+  VLOG_WORK << "Will denoise on " << denoiser_device->info.description << " ("
+            << denoiser_device->info.id << ")";
 
   return true;
 }
diff --git a/intern/cycles/integrator/denoiser_device.cpp b/intern/cycles/integrator/denoiser_device.cpp
index 595397312b3..5414f9dfb1a 100644
--- a/intern/cycles/integrator/denoiser_device.cpp
+++ b/intern/cycles/integrator/denoiser_device.cpp
@@ -48,7 +48,7 @@ bool DeviceDenoiser::denoise_buffer(const BufferParams &buffer_params,
     task.render_buffers = render_buffers;
   }
   else {
-    VLOG(3) << "Creating temporary buffer on denoiser device.";
+    VLOG_WORK << "Creating temporary buffer on denoiser device.";
 
     DeviceQueue *queue = denoiser_device->get_denoise_queue();
 
diff --git a/intern/cycles/integrator/denoiser_oidn.cpp b/intern/cycles/integrator/denoiser_oidn.cpp
index b074408e229..04e659a15e2 100644
--- a/intern/cycles/integrator/denoiser_oidn.cpp
+++ b/intern/cycles/integrator/denoiser_oidn.cpp
@@ -284,8 +284,8 @@ class OIDNDenoiseContext {
   /* Read pass pixels using PassAccessor into a temporary buffer which is owned by the pass.. */
   void read_pass_pixels_into_buffer(OIDNPass &oidn_pass)
   {
-    VLOG(3) << "Allocating temporary buffer for pass " << oidn_pass.name << " ("
-            << pass_type_as_string(oidn_pass.type) << ")";
+    VLOG_WORK << "Allocating temporary buffer for pass " << oidn_pass.name << " ("
+              << pass_type_as_string(oidn_pass.type) << ")";
 
     const int64_t width = buffer_params_.width;
     const int64_t height = buffer_params_.height;
diff --git a/intern/cycles/integrator/path_trace.cpp b/intern/cycles/integrator/path_trace.cpp
index 772c1454e02..6912bf928cd 100644
--- a/intern/cycles/integrator/path_trace.cpp
+++ b/intern/cycles/integrator/path_trace.cpp
@@ -348,8 +348,8 @@ void PathTrace::path_trace(RenderWork &render_work)
     return;
   }
 
-  VLOG(3) << "Will path trace " << render_work.path_trace.num_samples
-          << " samples at the resolution divider " << render_work.resolution_divider;
+  VLOG_WORK << "Will path trace " << render_work.path_trace.num_samples
+            << " samples at the resolution divider " << render_work.resolution_divider;
 
   const double start_time = time_dt();
 
@@ -373,9 +373,9 @@ void PathTrace::path_trace(RenderWork &render_work)
     work_balance_infos_[i].time_spent += work_time;
     work_balance_infos_[i].occupancy = statistics.occupancy;
 
-    VLOG(3) << "Rendered " << num_samples << " samples in " << work_time << " seconds ("
-            << work_time / num_samples
-            << " seconds per sample), occupancy: " << statistics.occupancy;
+    VLOG_WORK << "Rendered " << num_samples << " samples in " << work_time << " seconds ("
+              << work_time / num_samples
+              << " seconds per sample), occupancy: " << statistics.occupancy;
   });
 
   float occupancy_accum = 0.0f;
@@ -398,10 +398,10 @@ void PathTrace::adaptive_sample(RenderWork &render_work)
   bool did_reschedule_on_idle = false;
 
   while (true) {
-    VLOG(3) << "Will filter adaptive stopping buffer, threshold "
-            << render_work.adaptive_sampling.threshold;
+    VLOG_WORK << "Will filter adaptive stopping buffer, threshold "
+              << render_work.adaptive_sampling.threshold;
     if (render_work.adaptive_sampling.reset) {
-      VLOG(3) << "Will re-calculate convergency flag for currently converged pixels.";
+      VLOG_WORK << "Will re-calculate convergency flag for currently converged pixels.";
     }
 
     const double start_time = time_dt();
@@ -420,11 +420,11 @@ void PathTrace::adaptive_sample(RenderWork &render_work)
         render_work, time_dt() - start_time, is_cancel_requested());
 
     if (num_active_pixels == 0) {
-      VLOG(3) << "All pixels converged.";
+      VLOG_WORK << "All pixels converged.";
       if (!render_scheduler_.render_work_reschedule_on_converge(render_work)) {
         break;
       }
-      VLOG(3) << "Continuing with lower threshold.";
+      VLOG_WORK << "Continuing with lower threshold.";
     }
     else if (did_reschedule_on_idle) {
       break;
@@ -436,10 +436,10 @@ void PathTrace::adaptive_sample(RenderWork &render_work)
        * A better heuristic is possible here: for example, use maximum of 128^2 and percentage of
        * the final resolution. */
       if (!render_scheduler_.render_work_reschedule_on_idle(render_work)) {
-        VLOG(3) << "Rescheduling is not possible: final threshold is reached.";
+        VLOG_WORK << "Rescheduling is not possible: final threshold is reached.";
         break;
       }
-      VLOG(3) << "Rescheduling lower threshold.";
+      VLOG_WORK << "Rescheduling lower threshold.";
       did_reschedule_on_idle = true;
     }
     else {
@@ -483,7 +483,7 @@ void PathTrace::cryptomatte_postprocess(const RenderWork &render_work)
   if (!render_work.cryptomatte.postprocess) {
     return;
   }
-  VLOG(3) << "Perform cryptomatte work.";
+  VLOG_WORK << "Perform cryptomatte work.";
 
   parallel_for_each(path_trace_works_, [&](unique_ptr<PathTraceWork> &path_trace_work) {
     path_trace_work->cryptomatte_postproces();
@@ -501,7 +501,7 @@ void PathTrace::denoise(const RenderWork &render_work)
     return;
   }
 
-  VLOG(3) << "Perform denoising work.";
+  VLOG_WORK << "Perform denoising work.";
 
   const double start_time = time_dt();
 
@@ -599,26 +599,26 @@ void PathTrace::update_display(const RenderWork &render_work)
   }
 
   if (!display_ && !output_driver_) {
-    VLOG(3) << "Ignore display update.";
+    VLOG_WORK << "Ignore display update.";
     return;
   }
 
   if (full_params_.width == 0 || full_params_.height == 0) {
-    VLOG(3) << "Skipping PathTraceDisplay update due to 0 size of the render buffer.";
+    VLOG_WORK << "Skipping PathTraceDisplay update due to 0 size of the render buffer.";
     return;
   }
 
   const double start_time = time_dt();
 
   if (output_driver_) {
-    VLOG(3) << "Invoke buffer update callback.";
+    VLOG_WORK << "Invoke buffer update callback.";
 
     PathTraceTile tile(*this);
     output_driver_->update_render_tile(tile);
   }
 
   if (display_) {
-    VLOG(3) << "Perform copy to GPUDisplay work.";
+    VLOG_WORK << "Perform copy to GPUDisplay work.";
 
     const int texture_width = render_state_.effective_big_tile_params.window_width;
     const int texture_height = render_state_.effective_big_tile_params.window_height;
@@ -654,33 +654,33 @@ void PathTrace::rebalance(const RenderWork &render_work)
   const int num_works = path_trace_works_.size();
 
   if (num_works == 1) {
-    VLOG(3) << "Ignoring rebalance work due to single device render.";
+    VLOG_WORK << "Ignoring rebalance work due to single device render.";
     return;
   }
 
   const double start_time = time_dt();
 
   if (VLOG_IS_ON(3)) {
-    VLOG(3) << "Perform rebalance work.";
-    VLOG(3) << "Per-device path tracing time (seconds):";
+    VLOG_WORK << "Perform rebalance work.";
+    VLOG_WORK << "Per-device path tracing time (seconds):";
     for (int i = 0; i < num_works; ++i) {
-      VLOG(3) << path_trace_works_[i]->get_device()->info.description << ": "
-              << work_balance_infos_[i].time_spent;
+      VLOG_WORK << path_trace_works_[i]->get_device()->info.description << ": "
+                << work_balance_infos_[i].time_spent;
     }
   }
 
   const bool did_rebalance = work_balance_do_rebalance(work_balance_infos_);
 
   if (VLOG_IS_ON(3)) {
-    VLOG(3) << "Calculated per-device weights for works:";
+    VLOG_WORK << "Calculated per-device weights for works:";
     for (int i = 0; i < num_works; ++i) {
-      VLOG(3) << path_trace_works_[i]->get_device()->info.description << ": "
-              << work_balance_infos_[i].weight;
+      VLOG_WORK << path_trace_works_[i]->get_device()->info.description << ": "
+                << work_balance_infos_[i].weight;
     }
   }
 
   if (!did_rebalance) {
-    VLOG(3) << "Balance in path trace works did not change.";
+    VLOG_WORK << "Balance in path trace works did not change.";
     render_scheduler_.report_rebalance_time(render_work, time_dt() - start_time, false);
     return;
   }
@@ -704,7 +704,7 @@ void PathTrace::write_tile_buffer(const RenderWork &render_work)
     return;
   }
 
-  VLOG(3) << "Write tile result.";
+  VLOG_WORK << "Write tile result.";
 
   render_state_.tile_written = true;
 
@@ -718,14 +718,14 @@ void PathTrace::write_tile_buffer(const RenderWork &render_work)
    *
    * Important thing is: tile should be written to the software via callback only once. */
   if (!has_multiple_tiles) {
-    VLOG(3) << "Write tile result via buffer write callback.";
+    VLOG_WORK << "Write tile result via buffer write callback.";
     tile_buffer_write();
   }
 
   /* Write tile to disk, so that the render work's render buffer can be re-used for the next tile.
    */
   if (has_multiple_tiles) {
-    VLOG(3) << "Write tile result into .";
+    VLOG_WORK << "Write tile result into .";
     tile_buffer_write_to_disk();
   }
 }
@@ -736,10 +736,10 @@ void PathTrace::finalize_full_buffer_on_disk(const RenderWork &render_work)
     return;
   }
 
-  VLOG(3) << "Handle full-frame render buffer work.";
+  VLOG_WORK << "Handle full-frame render buffer work.";
 
   if (!tile_manager_.has_written_tiles()) {
-    VLOG(3) << "No tiles on disk.";
+    VLOG_WORK << "No tiles on disk.";
     return;
   }
 
@@ -935,7 +935,7 @@ static string get_layer_view_name(const RenderBuffers &buffers)
 
 void PathTrace::process_full_buffer_from_disk(string_view filename)
 {
-  VLOG(3) << "Processing full frame buffer file " << filename;
+  VLOG_WORK << "Processing full frame buffer file " << filename;
 
   progress_set_status("Reading full buffer from disk");
 
diff --git a/intern/cycles/integrator/path_trace_work_gpu.cpp b/intern/cycles/integrator/path_trace_work_gpu.cpp
index ede81705ae8..0acaeace4b0 100644
--- a/intern/cycles/integrator/path_trace_work_gpu.cpp
+++ b/intern/cycles/integrator/path_trace_work_gpu.cpp
@@ -152,7 +152,7 @@ void PathTraceWorkGPU::alloc_integrator_soa()
       total_soa_size += soa_memory->memory_size();
     }
 
-    VLOG(3) << "GPU SoA state size: " << string_human_readable_size(total_soa_size);
+    VLOG_DEVICE_STATS << "GPU SoA state size: " << string_human_readable_size(total_soa_size);
   }
 }
 
@@ -820,10 +820,10 @@ bool PathTraceWorkGPU::should_use_graphics_interop()
     interop_use_ = device->should_use_graphics_interop();
 
     if (interop_use_) {
-      VLOG(2) << "Using graphics interop GPU display update.";
+      VLOG_INFO << "Using graphics interop GPU display update.";
     }
     else {
-      VLOG(2) << "Using naive GPU display update.";
+      VLOG_INFO << "Using naive GPU display update.";
     }
 
     interop_use_checked_ = true;
diff --git a/intern/cycles/integrator/render_scheduler.cpp b/intern/cycles/integrator/render_scheduler.cpp
index ebc3170393f..e4676bd059c 100644
--- a/intern/cycles/integrator/render_scheduler.cpp
+++ b/intern/cycles/integrator/render_scheduler.cpp
@@ -225,7 +225,7 @@ bool RenderScheduler::render_work_reschedule_on_idle(RenderWork &render_work)
 
 void RenderScheduler::render_work_reschedule_on_cancel(RenderWork &render_work)
 {
-  VLOG(3) << "Schedule work for cancel.";
+  VLOG_WORK << "Schedule work for cancel.";
 
   /* Un-schedule samples: they will not be rendered and should not be counted. */
   state_.num_rendered_samples -= render_work.path_trace.num_samples;
@@ -475,14 +475,14 @@ void RenderScheduler::report_path_trace_time(const RenderWork &render_work,
 
   path_trace_time_.add_average(final_time_approx, render_work.path_trace.num_samples);
 
-  VLOG(4) << "Average path tracing time: " << path_trace_time_.get_average() << " seconds.";
+  VLOG_WORK << "Average path tracing time: " << path_trace_time_.get_average() << " seconds.";
 }
 
 void RenderScheduler::report_path_trace_occupancy(const RenderWork &render_work, float occupancy)
 {
   state_.occupancy_num_samples = render_work.path_trace.num_samples;
   state_.occupancy = occupancy;
-  VLOG(4) << "Measured path tracing occupancy: " << occupancy;
+  VLOG_WORK << "Measured path tracing occupancy: " << occupancy;
 }
 
 void RenderScheduler::report_adaptive_filter_time(const RenderWork &render_work,
@@ -503,8 +503,8 @@ void RenderScheduler::report_adaptive_filter_time(const RenderWork &render_work,
 
   adaptive_filter_time_.add_average(final_time_approx, render_work.path_trace.num_samples);
 
-  VLOG(4) << "Average adaptive sampling filter  time: " << adaptive_filter_time_.get_average()
-          << " seconds.";
+  VLOG_WORK << "Average adaptive sampling filter  time: " << adaptive_filter_time_.get_average()
+            << " seconds.";
 }
 
 void RenderScheduler::report_denoise_time(const RenderWork &render_work, double time)
@@ -523,7 +523,7 @@ void RenderScheduler::report_denoise_time(const RenderWork &render_work, double
 
   denoise_time_.add_average(final_time_approx);
 
-  VLOG(4) << "Average denoising time: " << denoise_time_.get_average() << " seconds.";
+  VLOG_WORK << "Average denoising time: " << denoise_time_.get_average() << " seconds.";
 }
 
 void RenderScheduler::report_display_update_time(const RenderWork &render_work, double time)
@@ -542,7 +542,8 @@ void RenderScheduler::report_display_update_time(const RenderWork &render_work,
 
   display_update_time_.add_average(final_time_approx);
 
-  VLOG(4) << "Average display update time: " << display_update_time_.get_average() << " seconds.";
+  VLOG_WORK << "Average display update time: " << display_update_time_.get_average()
+            << " seconds.";
 
   /* Move the display update moment further in time, so that logic which checks when last update
    * did happen have more reliable point in time (without path tracing and denoising parts of the
@@ -568,7 +569,7 @@ void RenderScheduler::report_rebalance_time(const RenderWork &render_work,
 
   state_.last_rebalance_changed = balance_changed;
 
-  VLOG(4) << "Average rebalance time: " << rebalance_time_.get_average() << " seconds.";
+  VLOG_WORK << "Average rebalance time: " << rebalance_time_.get_average() << " seconds.";
 }
 
 string RenderScheduler::full_report() const
@@ -1063,7 +1064,7 @@ void RenderScheduler::update_start_resolution_divider()
     /* Resolution divider has never been calculated before: use default resolution, so that we have
      * somewhat good initial behavior, giving a chance to collect real numbers. */
     start_resolution_divider_ = default_start_resolution_divider_;
-    VLOG(3) << "Initial resolution divider is " << start_resolution_divider_;
+    VLOG_WORK << "Initial resolution divider is " << start_resolution_divider_;
     return;
   }
 
@@ -1092,7 +1093,7 @@ void RenderScheduler::update_start_resolution_divider()
    * simple and compute device is fast). */
   start_resolution_divider_ = max(resolution_divider_for_update, pixel_size_);
 
-  VLOG(3) << "Calculated resolution divider is " << start_resolution_divider_;
+  VLOG_WORK << "Calculated resolution divider is " << start_resolution_divider_;
 }
 
 double RenderScheduler::guess_viewport_navigation_update_interval_in_seconds() const
diff --git a/intern/cycles/integrator/shader_eval.cpp b/intern/cycles/integrator/shader_eval.cpp
index 92b9d1c662d..b1450732f5c 100644
--- a/intern/cycles/integrator/shader_eval.cpp
+++ b/intern/cycles/integrator/shader_eval.cpp
@@ -31,8 +31,8 @@ bool ShaderEval::eval(const ShaderEvalType type,
 
   device_->foreach_device([&](Device *device) {
     if (!first_device) {
-      LOG(ERROR) << "Multi-devices are not yet fully implemented, will evaluate shader on a "
-                    "single device.";
+      VLOG_WORK << "Multi-devices are not yet fully implemented, will evaluate shader on a "
+                   "single device.";
       return;
     }
     first_device = false;
diff --git a/intern/cycles/integrator/work_tile_scheduler.cpp b/intern/cycles/integrator/work_tile_scheduler.cpp
index 6dc511064c9..4bc8c0c4396 100644
--- a/intern/cycles/integrator/work_tile_scheduler.cpp
+++ b/intern/cycles/integrator/work_tile_scheduler.cpp
@@ -55,7 +55,7 @@ void WorkTileScheduler::reset_scheduler_state()
   tile_size_ = tile_calculate_best_size(
       accelerated_rt_, image_size_px_, samples_num_, max_num_path_states_, scrambling_distance_);
 
-  VLOG(3) << "Will schedule tiles of size " << tile_size_;
+  VLOG_WORK << "Will schedule tiles of size " << tile_size_;
 
   if (VLOG_IS_ON(3)) {
     /* The logging is based on multiple tiles scheduled, ignoring overhead of multi-tile scheduling
@@ -63,8 +63,8 @@ void WorkTileScheduler::reset_scheduler_state()
     const int num_path_states_in_tile = tile_size_.width * tile_size_.height *
                                         tile_size_.num_samples;
     const int num_tiles = max_num_path_states_ / num_path_states_in_tile;
-    VLOG(3) << "Number of unused path states: "
-            << max_num_path_states_ - num_tiles * num_path_states_in_tile;
+    VLOG_WORK << "Number of unused path states: "
+              << max_num_path_states_ - num_tiles * num_path_states_in_tile;
   }
 
   num_tiles_x_ = divide_up(image_size_px_.x, tile_size_.width);
diff --git a/intern/cycles/kernel/bvh/util.h b/intern/cycles/kernel/bvh/util.h
index 71045157372..d53198f97a3 100644
--- a/intern/cycles/kernel/bvh/util.h
+++ b/intern/cycles/kernel/bvh/util.h
@@ -111,16 +111,16 @@ ccl_device_inline int intersection_find_attribute(KernelGlobals kg,
                                                   const uint id)
 {
   uint attr_offset = kernel_tex_fetch(__objects, object).attribute_map_offset;
-  uint4 attr_map = kernel_tex_fetch(__attributes_map, attr_offset);
+  AttributeMap attr_map = kernel_tex_fetch(__attributes_map, attr_offset);
 
-  while (attr_map.x != id) {
-    if (UNLIKELY(attr_map.x == ATTR_STD_NONE)) {
-      if (UNLIKELY(attr_map.y == 0)) {
+  while (attr_map.id != id) {
+    if (UNLIKELY(attr_map.id == ATTR_STD_NONE)) {
+      if (UNLIKELY(attr_map.element == 0)) {
         return (int)ATTR_STD_NOT_FOUND;
       }
       else {
         /* Chain jump to a different part of the table. */
-        attr_offset = attr_map.z;
+        attr_offset = attr_map.offset;
       }
     }
     else {
@@ -130,7 +130,7 @@ ccl_device_inline int intersection_find_attribute(KernelGlobals kg,
   }
 
   /* return result */
-  return (attr_map.y == ATTR_ELEMENT_NONE) ? (int)ATTR_STD_NOT_FOUND : (int)attr_map.z;
+  return (attr_map.element == ATTR_ELEMENT_NONE) ? (int)ATTR_STD_NOT_FOUND : (int)attr_map.offset;
 }
 
 /* Transparent Shadows */
diff --git a/intern/cycles/kernel/geom/attribute.h b/intern/cycles/kernel/geom/attribute.h
index da620f69e2d..774b25a76ff 100644
--- a/intern/cycles/kernel/geom/attribute.h
+++ b/intern/cycles/kernel/geom/attribute.h
@@ -56,16 +56,16 @@ ccl_device_inline AttributeDescriptor find_attribute(KernelGlobals kg,
   /* for SVM, find attribute by unique id */
   uint attr_offset = object_attribute_map_offset(kg, sd->object);
   attr_offset += attribute_primitive_type(kg, sd);
-  uint4 attr_map = kernel_tex_fetch(__attributes_map, attr_offset);
+  AttributeMap attr_map = kernel_tex_fetch(__attributes_map, attr_offset);
 
-  while (attr_map.x != id) {
-    if (UNLIKELY(attr_map.x == ATTR_STD_NONE)) {
-      if (UNLIKELY(attr_map.y == 0)) {
+  while (attr_map.id != id) {
+    if (UNLIKELY(attr_map.id == ATTR_STD_NONE)) {
+      if (UNLIKELY(attr_map.element == 0)) {
         return attribute_not_found();
       }
       else {
         /* Chain jump to a different part of the table. */
-        attr_offset = attr_map.z;
+        attr_offset = attr_map.offset;
       }
     }
     else {
@@ -75,7 +75,7 @@ ccl_device_inline AttributeDescriptor find_attribute(KernelGlobals kg,
   }
 
   AttributeDescriptor desc;
-  desc.element = (AttributeElement)attr_map.y;
+  desc.element = (AttributeElement)attr_map.element;
 
   if (sd->prim == PRIM_NONE && desc.element != ATTR_ELEMENT_MESH &&
       desc.element != ATTR_ELEMENT_VOXEL && desc.element != ATTR_ELEMENT_OBJECT) {
@@ -83,9 +83,10 @@ ccl_device_inline AttributeDescriptor find_attribute(KernelGlobals kg,
   }
 
   /* return result */
-  desc.offset = (attr_map.y == ATTR_ELEMENT_NONE) ? (int)ATTR_STD_NOT_FOUND : (int)attr_map.z;
-  desc.type = (NodeAttributeType)(attr_map.w & 0xff);
-  desc.flags = (AttributeFlag)(attr_map.w >> 8);
+  desc.offset = (attr_map.element == ATTR_ELEMENT_NONE) ? (int)ATTR_STD_NOT_FOUND :
+                                                          (int)attr_map.offset;
+  desc.type = (NodeAttributeType)attr_map.type;
+  desc.flags = (AttributeFlag)attr_map.flags;
 
   return desc;
 }
diff --git a/intern/cycles/kernel/osl/services.cpp b/intern/cycles/kernel/osl/services.cpp
index e2e10b5b83f..6e75ae54f33 100644
--- a/intern/cycles/kernel/osl/services.cpp
+++ b/intern/cycles/kernel/osl/services.cpp
@@ -132,7 +132,7 @@ OSLRenderServices::OSLRenderServices(OSL::TextureSystem *texture_system)
 OSLRenderServices::~OSLRenderServices()
 {
   if (texture_system) {
-    VLOG(2) << "OSL texture system stats:\n" << texture_system->getstats();
+    VLOG_INFO << "OSL texture system stats:\n" << texture_system->getstats();
   }
 }
 
diff --git a/intern/cycles/kernel/textures.h b/intern/cycles/kernel/textures.h
index 7deb589a0a9..d8ac9cbe51f 100644
--- a/intern/cycles/kernel/textures.h
+++ b/intern/cycles/kernel/textures.h
@@ -47,7 +47,7 @@ KERNEL_TEX(float4, __points)
 KERNEL_TEX(uint, __points_shader)
 
 /* attributes */
-KERNEL_TEX(uint4, __attributes_map)
+KERNEL_TEX(AttributeMap, __attributes_map)
 KERNEL_TEX(float, __attributes_float)
 KERNEL_TEX(float2, __attributes_float2)
 KERNEL_TEX(packed_float3, __attributes_float3)
diff --git a/intern/cycles/kernel/types.h b/intern/cycles/kernel/types.h
index a76effdd952..f2e61d25002 100644
--- a/intern/cycles/kernel/types.h
+++ b/intern/cycles/kernel/types.h
@@ -670,6 +670,16 @@ typedef struct AttributeDescriptor {
   int offset;
 } AttributeDescriptor;
 
+/* For looking up attributes on objects and geometry. */
+typedef struct AttributeMap {
+  uint id;       /* Global unique identifier. */
+  uint element;  /* AttributeElement. */
+  int offset;    /* Offset into __attributes global arrays. */
+  uint8_t type;  /* NodeAttributeType. */
+  uint8_t flags; /* AttributeFlag. */
+  uint8_t pad[2];
+} AttributeMap;
+
 /* Closure data */
 
 #ifndef __MAX_CLOSURE__
diff --git a/intern/cycles/scene/alembic.cpp b/intern/cycles/scene/alembic.cpp
index c1e2d306fcc..e6f39bf8625 100644
--- a/intern/cycles/scene/alembic.cpp
+++ b/intern/cycles/scene/alembic.cpp
@@ -1514,7 +1514,7 @@ void AlembicProcedural::build_caches(Progress &progress)
     }
   }
 
-  VLOG(1) << "AlembicProcedural memory usage : " << string_human_readable_size(memory_used);
+  VLOG_WORK << "AlembicProcedural memory usage : " << string_human_readable_size(memory_used);
 }
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/scene/camera.cpp b/intern/cycles/scene/camera.cpp
index 710f1c5ee90..eec269ab542 100644
--- a/intern/cycles/scene/camera.cpp
+++ b/intern/cycles/scene/camera.cpp
@@ -530,7 +530,7 @@ void Camera::device_update_volume(Device * /*device*/, DeviceScene *dscene, Scen
                      if (object->get_geometry()->has_volume &&
                          viewplane_boundbox.intersects(object->bounds)) {
                        /* TODO(sergey): Consider adding more grained check. */
-                       VLOG(1) << "Detected camera inside volume.";
+                       VLOG_INFO << "Detected camera inside volume.";
                        kcam->is_inside_volume = 1;
                        parallel_for_cancel();
                        break;
@@ -539,7 +539,7 @@ void Camera::device_update_volume(Device * /*device*/, DeviceScene *dscene, Scen
                  });
 
     if (!kcam->is_inside_volume) {
-      VLOG(1) << "Camera is outside of the volume.";
+      VLOG_INFO << "Camera is outside of the volume.";
     }
   }
 
diff --git a/intern/cycles/scene/colorspace.cpp b/intern/cycles/scene/colorspace.cpp
index f87b4c62ab2..189e3bc752d 100644
--- a/intern/cycles/scene/colorspace.cpp
+++ b/intern/cycles/scene/colorspace.cpp
@@ -55,8 +55,8 @@ ColorSpaceProcessor *ColorSpaceManager::get_processor(ustring colorspace)
     }
     catch (OCIO::Exception &exception) {
       cached_processors[colorspace] = OCIO::ConstProcessorRcPtr();
-      VLOG(1) << "Colorspace " << colorspace.c_str()
-              << " can't be converted to scene_linear: " << exception.what();
+      VLOG_WARNING << "Colorspace " << colorspace.c_str()
+                   << " can't be converted to scene_linear: " << exception.what();
     }
   }
 
@@ -132,12 +132,12 @@ ustring ColorSpaceManager::detect_known_colorspace(ustring colorspace,
 
     thread_scoped_lock cache_lock(cache_colorspaces_mutex);
     if (is_scene_linear) {
-      VLOG(1) << "Colorspace " << colorspace.string() << " is no-op";
+      VLOG_INFO << "Colorspace " << colorspace.string() << " is no-op";
       cached_colorspaces[colorspace] = u_colorspace_raw;
       return u_colorspace_raw;
     }
     else if (is_srgb) {
-      VLOG(1) << "Colorspace " << colorspace.string() << " is sRGB";
+      VLOG_INFO << "Colorspace " << colorspace.string() << " is sRGB";
       cached_colorspaces[colorspace] = u_colorspace_srgb;
       return u_colorspace_srgb;
     }
@@ -146,22 +146,23 @@ ustring ColorSpaceManager::detect_known_colorspace(ustring colorspace,
     if (!get_processor(colorspace)) {
       OCIO::ConstConfigRcPtr config = OCIO::GetCurrentConfig();
       if (!config || !config->getColorSpace(colorspace.c_str())) {
-        VLOG(1) << "Colorspace " << colorspace.c_str() << " not found, using raw instead";
+        VLOG_WARNING << "Colorspace " << colorspace.c_str() << " not found, using raw instead";
       }
       else {
-        VLOG(1) << "Colorspace " << colorspace.c_str()
-                << " can't be converted to scene_linear, using raw instead";
+        VLOG_WARNING << "Colorspace " << colorspace.c_str()
+                     << " can't be converted to scene_linear, using raw instead";
       }
       cached_colorspaces[colorspace] = u_colorspace_raw;
       return u_colorspace_raw;
     }
 
     /* Convert to/from colorspace with OpenColorIO. */
-    VLOG(1) << "Colorspace " << colorspace.string() << " handled through OpenColorIO";
+    VLOG_INFO << "Colorspace " << colorspace.string() << " handled through OpenColorIO";
     cached_colorspaces[colorspace] = colorspace;
     return colorspace;
 #else
-    VLOG(1) << "Colorspace " << colorspace.c_str() << " not available, built without OpenColorIO";
+    VLOG_WARNING << "Colorspace " << colorspace.c_str()
+                 << " not available, built without OpenColorIO";
     return u_colorspace_raw;
 #endif
   }
diff --git a/intern/cycles/scene/constant_fold.cpp b/intern/cycles/scene/constant_fold.cpp
index 46ffbe9043a..4bce5661f9b 100644
--- a/intern/cycles/scene/constant_fold.cpp
+++ b/intern/cycles/scene/constant_fold.cpp
@@ -30,8 +30,8 @@ bool ConstantFolder::all_inputs_constant() const
 
 void ConstantFolder::make_constant(float value) const
 {
-  VLOG(3) << "Folding " << node->name << "::" << output->name() << " to constant (" << value
-          << ").";
+  VLOG_DEBUG << "Folding " << node->name << "::" << output->name() << " to constant (" << value
+             << ").";
 
   foreach (ShaderInput *sock, output->links) {
     sock->set(value);
@@ -43,7 +43,8 @@ void ConstantFolder::make_constant(float value) const
 
 void ConstantFolder::make_constant(float3 value) const
 {
-  VLOG(3) << "Folding " << node->name << "::" << output->name() << " to constant " << value << ".";
+  VLOG_DEBUG << "Folding " << node->name << "::" << output->name() << " to constant " << value
+             << ".";
 
   foreach (ShaderInput *sock, output->links) {
     sock->set(value);
@@ -99,8 +100,8 @@ void ConstantFolder::bypass(ShaderOutput *new_output) const
 {
   assert(new_output);
 
-  VLOG(3) << "Folding " << node->name << "::" << output->name() << " to socket "
-          << new_output->parent->name << "::" << new_output->name() << ".";
+  VLOG_DEBUG << "Folding " << node->name << "::" << output->name() << " to socket "
+             << new_output->parent->name << "::" << new_output->name() << ".";
 
   /* Remove all outgoing links from socket and connect them to new_output instead.
    * The graph->relink method affects node inputs, so it's not safe to use in constant
@@ -118,7 +119,7 @@ void ConstantFolder::discard() const
 {
   assert(output->type() == SocketType::CLOSURE);
 
-  VLOG(3) << "Discarding closure " << node->name << ".";
+  VLOG_DEBUG << "Discarding closure " << node->name << ".";
 
   graph->disconnect(output);
 }
diff --git a/intern/cycles/scene/film.cpp b/intern/cycles/scene/film.cpp
index 7f69df7b321..17acb30e9e7 100644
--- a/intern/cycles/scene/film.cpp
+++ b/intern/cycles/scene/film.cpp
@@ -580,10 +580,10 @@ void Film::update_passes(Scene *scene, bool add_sample_count_pass)
   tag_modified();
 
   /* Debug logging. */
-  if (VLOG_IS_ON(2)) {
-    VLOG(2) << "Effective scene passes:";
+  if (VLOG_INFO_IS_ON) {
+    VLOG_INFO << "Effective scene passes:";
     for (const Pass *pass : scene->passes) {
-      VLOG(2) << "- " << *pass;
+      VLOG_INFO << "- " << *pass;
     }
   }
 }
diff --git a/intern/cycles/scene/geometry.cpp b/intern/cycles/scene/geometry.cpp
index 9152abacbdb..5ad0742b009 100644
--- a/intern/cycles/scene/geometry.cpp
+++ b/intern/cycles/scene/geometry.cpp
@@ -407,43 +407,47 @@ void GeometryManager::update_osl_attributes(Device *device,
 
 /* Generate a normal attribute map entry from an attribute descriptor. */
 static void emit_attribute_map_entry(
-    uint4 *attr_map, int index, uint id, TypeDesc type, const AttributeDescriptor &desc)
+    AttributeMap *attr_map, int index, uint id, TypeDesc type, const AttributeDescriptor &desc)
 {
-  attr_map[index].x = id;
-  attr_map[index].y = desc.element;
-  attr_map[index].z = as_uint(desc.offset);
+  attr_map[index].id = id;
+  attr_map[index].element = desc.element;
+  attr_map[index].offset = as_uint(desc.offset);
 
   if (type == TypeDesc::TypeFloat)
-    attr_map[index].w = NODE_ATTR_FLOAT;
+    attr_map[index].type = NODE_ATTR_FLOAT;
   else if (type == TypeDesc::TypeMatrix)
-    attr_map[index].w = NODE_ATTR_MATRIX;
+    attr_map[index].type = NODE_ATTR_MATRIX;
   else if (type == TypeFloat2)
-    attr_map[index].w = NODE_ATTR_FLOAT2;
+    attr_map[index].type = NODE_ATTR_FLOAT2;
   else if (type == TypeFloat4)
-    attr_map[index].w = NODE_ATTR_FLOAT4;
+    attr_map[index].type = NODE_ATTR_FLOAT4;
   else if (type == TypeRGBA)
-    attr_map[index].w = NODE_ATTR_RGBA;
+    attr_map[index].type = NODE_ATTR_RGBA;
   else
-    attr_map[index].w = NODE_ATTR_FLOAT3;
+    attr_map[index].type = NODE_ATTR_FLOAT3;
 
-  attr_map[index].w |= desc.flags << 8;
+  attr_map[index].flags = desc.flags;
 }
 
 /* Generate an attribute map end marker, optionally including a link to another map.
  * Links are used to connect object attribute maps to mesh attribute maps. */
-static void emit_attribute_map_terminator(uint4 *attr_map, int index, bool chain, uint chain_link)
+static void emit_attribute_map_terminator(AttributeMap *attr_map,
+                                          int index,
+                                          bool chain,
+                                          uint chain_link)
 {
   for (int j = 0; j < ATTR_PRIM_TYPES; j++) {
-    attr_map[index + j].x = ATTR_STD_NONE;
-    attr_map[index + j].y = chain;                      /* link is valid flag */
-    attr_map[index + j].z = chain ? chain_link + j : 0; /* link to the correct sub-entry */
-    attr_map[index + j].w = 0;
+    attr_map[index + j].id = ATTR_STD_NONE;
+    attr_map[index + j].element = chain;                     /* link is valid flag */
+    attr_map[index + j].offset = chain ? chain_link + j : 0; /* link to the correct sub-entry */
+    attr_map[index + j].type = 0;
+    attr_map[index + j].flags = 0;
   }
 }
 
 /* Generate all necessary attribute map entries from the attribute request. */
 static void emit_attribute_mapping(
-    uint4 *attr_map, int index, Scene *scene, AttributeRequest &req, Geometry *geom)
+    AttributeMap *attr_map, int index, Scene *scene, AttributeRequest &req, Geometry *geom)
 {
   uint id;
 
@@ -501,8 +505,8 @@ void GeometryManager::update_svm_attributes(Device *,
   }
 
   /* create attribute map */
-  uint4 *attr_map = dscene->attributes_map.alloc(attr_map_size);
-  memset(attr_map, 0, dscene->attributes_map.size() * sizeof(uint));
+  AttributeMap *attr_map = dscene->attributes_map.alloc(attr_map_size);
+  memset(attr_map, 0, dscene->attributes_map.size() * sizeof(*attr_map));
 
   for (size_t i = 0; i < scene->geometry.size(); i++) {
     Geometry *geom = scene->geometry[i];
@@ -1288,7 +1292,7 @@ void GeometryManager::device_update_bvh(Device *device,
   bparams.bvh_type = scene->params.bvh_type;
   bparams.curve_subdivisions = scene->params.curve_subdivisions();
 
-  VLOG(1) << "Using " << bvh_layout_name(bparams.bvh_layout) << " layout.";
+  VLOG_INFO << "Using " << bvh_layout_name(bparams.bvh_layout) << " layout.";
 
   const bool can_refit = scene->bvh != nullptr &&
                          (bparams.bvh_layout == BVHLayout::BVH_LAYOUT_OPTIX ||
@@ -1799,7 +1803,7 @@ void GeometryManager::device_update(Device *device,
   if (!need_update())
     return;
 
-  VLOG(1) << "Total " << scene->geometry.size() << " meshes.";
+  VLOG_INFO << "Total " << scene->geometry.size() << " meshes.";
 
   bool true_displacement_used = false;
   bool curve_shadow_transparency_used = false;
@@ -2038,7 +2042,7 @@ void GeometryManager::device_update(Device *device,
 
     TaskPool::Summary summary;
     pool.wait_work(&summary);
-    VLOG(2) << "Objects BVH build pool statistics:\n" << summary.full_report();
+    VLOG_WORK << "Objects BVH build pool statistics:\n" << summary.full_report();
   }
 
   foreach (Shader *shader, scene->shaders) {
diff --git a/intern/cycles/scene/image.cpp b/intern/cycles/scene/image.cpp
index 1b44162351a..f0d57309ffb 100644
--- a/intern/cycles/scene/image.cpp
+++ b/intern/cycles/scene/image.cpp
@@ -653,8 +653,8 @@ bool ImageManager::file_load_image(Image *img, int texture_limit)
     while (max_size * scale_factor > texture_limit) {
       scale_factor *= 0.5f;
     }
-    VLOG(1) << "Scaling image " << img->loader->name() << " by a factor of " << scale_factor
-            << ".";
+    VLOG_WORK << "Scaling image " << img->loader->name() << " by a factor of " << scale_factor
+              << ".";
     vector<StorageType> scaled_pixels;
     size_t scaled_width, scaled_height, scaled_depth;
     util_image_resize_pixels(pixels_storage,
diff --git a/intern/cycles/scene/image_oiio.cpp b/intern/cycles/scene/image_oiio.cpp
index 09676455308..500e53ed763 100644
--- a/intern/cycles/scene/image_oiio.cpp
+++ b/intern/cycles/scene/image_oiio.cpp
@@ -22,11 +22,11 @@ bool OIIOImageLoader::load_metadata(const ImageDeviceFeatures & /*features*/,
 {
   /* Perform preliminary checks, with meaningful logging. */
   if (!path_exists(filepath.string())) {
-    VLOG(1) << "File '" << filepath.string() << "' does not exist.";
+    VLOG_WARNING << "File '" << filepath.string() << "' does not exist.";
     return false;
   }
   if (path_is_directory(filepath.string())) {
-    VLOG(1) << "File '" << filepath.string() << "' is a directory, can't use as image.";
+    VLOG_WARNING << "File '" << filepath.string() << "' is a directory, can't use as image.";
     return false;
   }
 
diff --git a/intern/cycles/scene/image_vdb.cpp b/intern/cycles/scene/image_vdb.cpp
index 2209be60a97..059eb09fef4 100644
--- a/intern/cycles/scene/image_vdb.cpp
+++ b/intern/cycles/scene/image_vdb.cpp
@@ -70,7 +70,7 @@ struct ToNanoOp {
         nanogrid = nanovdb::openToNanoVDB(floatgrid);
       }
       catch (const std::exception &e) {
-        VLOG(1) << "Error converting OpenVDB to NanoVDB grid: " << e.what();
+        VLOG_WARNING << "Error converting OpenVDB to NanoVDB grid: " << e.what();
       }
       return true;
     }
diff --git a/intern/cycles/scene/integrator.cpp b/intern/cycles/scene/integrator.cpp
index fda6ecc8d14..aa11004fb48 100644
--- a/intern/cycles/scene/integrator.cpp
+++ b/intern/cycles/scene/integrator.cpp
@@ -338,7 +338,7 @@ AdaptiveSampling Integrator::get_adaptive_sampling() const
 
   if (aa_samples > 0 && adaptive_threshold == 0.0f) {
     adaptive_sampling.threshold = max(0.001f, 1.0f / (float)aa_samples);
-    VLOG(1) << "Cycles adaptive sampling: automatic threshold = " << adaptive_sampling.threshold;
+    VLOG_INFO << "Cycles adaptive sampling: automatic threshold = " << adaptive_sampling.threshold;
   }
   else {
     adaptive_sampling.threshold = adaptive_threshold;
@@ -350,8 +350,8 @@ AdaptiveSampling Integrator::get_adaptive_sampling() const
      * in various test scenes. */
     const int min_samples = (int)ceilf(16.0f / powf(adaptive_sampling.threshold, 0.3f));
     adaptive_sampling.min_samples = max(4, min_samples);
-    VLOG(1) << "Cycles adaptive sampling: automatic min samples = "
-            << adaptive_sampling.min_samples;
+    VLOG_INFO << "Cycles adaptive sampling: automatic min samples = "
+              << adaptive_sampling.min_samples;
   }
   else {
     adaptive_sampling.min_samples = max(4, adaptive_min_samples);
diff --git a/intern/cycles/scene/light.cpp b/intern/cycles/scene/light.cpp
index 5e311d3051f..2b0d0c807a7 100644
--- a/intern/cycles/scene/light.cpp
+++ b/intern/cycles/scene/light.cpp
@@ -215,7 +215,9 @@ void LightManager::test_enabled_lights(Scene *scene)
      */
     Shader *shader = scene->background->get_shader(scene);
     const bool disable_mis = !(has_portal || shader->has_surface_spatial_varying);
-    VLOG_IF(1, disable_mis) << "Background MIS has been disabled.\n";
+    if (disable_mis) {
+      VLOG_INFO << "Background MIS has been disabled.\n";
+    }
     foreach (Light *light, scene->lights) {
       if (light->light_type == LIGHT_BACKGROUND) {
         light->is_enabled = !disable_mis;
@@ -309,7 +311,7 @@ void LightManager::device_update_distribution(Device *,
   }
 
   size_t num_distribution = num_triangles + num_lights;
-  VLOG(1) << "Total " << num_distribution << " of light distribution primitives.";
+  VLOG_INFO << "Total " << num_distribution << " of light distribution primitives.";
 
   /* emission area */
   KernelLightDistribution *distribution = dscene->light_distribution.alloc(num_distribution + 1);
@@ -655,13 +657,14 @@ void LightManager::device_update_background(Device *device,
   if (res.x == 0) {
     res = environment_res;
     if (res.x > 0 && res.y > 0) {
-      VLOG(2) << "Automatically set World MIS resolution to " << res.x << " by " << res.y << "\n";
+      VLOG_INFO << "Automatically set World MIS resolution to " << res.x << " by " << res.y
+                << "\n";
     }
   }
   /* If it's still unknown, just use the default. */
   if (res.x == 0 || res.y == 0) {
     res = make_int2(1024, 512);
-    VLOG(2) << "Setting World MIS resolution to default\n";
+    VLOG_INFO << "Setting World MIS resolution to default\n";
   }
   kbackground->map_res_x = res.x;
   kbackground->map_res_y = res.y;
@@ -704,7 +707,7 @@ void LightManager::device_update_background(Device *device,
 
   marg_cdf[res.y].y = 1.0f;
 
-  VLOG(2) << "Background MIS build time " << time_dt() - time_start << "\n";
+  VLOG_WORK << "Background MIS build time " << time_dt() - time_start << "\n";
 
   /* update device */
   dscene->light_background_marginal_cdf.copy_to_device();
@@ -725,7 +728,7 @@ void LightManager::device_update_points(Device *, DeviceScene *dscene, Scene *sc
   KernelLight *klights = dscene->lights.alloc(num_lights);
 
   if (num_lights == 0) {
-    VLOG(1) << "No effective light, ignoring points update.";
+    VLOG_WORK << "No effective light, ignoring points update.";
     return;
   }
 
@@ -955,9 +958,9 @@ void LightManager::device_update_points(Device *, DeviceScene *dscene, Scene *sc
     light_index++;
   }
 
-  VLOG(1) << "Number of lights sent to the device: " << light_index;
+  VLOG_INFO << "Number of lights sent to the device: " << light_index;
 
-  VLOG(1) << "Number of lights without contribution: " << num_scene_lights - light_index;
+  VLOG_INFO << "Number of lights without contribution: " << num_scene_lights - light_index;
 
   dscene->lights.copy_to_device();
 }
@@ -976,7 +979,7 @@ void LightManager::device_update(Device *device,
     }
   });
 
-  VLOG(1) << "Total " << scene->lights.size() << " lights.";
+  VLOG_INFO << "Total " << scene->lights.size() << " lights.";
 
   /* Detect which lights are enabled, also determines if we need to update the background. */
   test_enabled_lights(scene);
diff --git a/intern/cycles/scene/object.cpp b/intern/cycles/scene/object.cpp
index ddd89a16640..3a1c177fe3b 100644
--- a/intern/cycles/scene/object.cpp
+++ b/intern/cycles/scene/object.cpp
@@ -688,7 +688,7 @@ void ObjectManager::device_update(Device *device,
     dscene->objects.tag_modified();
   }
 
-  VLOG(1) << "Total " << scene->objects.size() << " objects.";
+  VLOG_INFO << "Total " << scene->objects.size() << " objects.";
 
   device_free(device, dscene, false);
 
diff --git a/intern/cycles/scene/osl.cpp b/intern/cycles/scene/osl.cpp
index 6698e6e2cce..f5ee0c0f1d3 100644
--- a/intern/cycles/scene/osl.cpp
+++ b/intern/cycles/scene/osl.cpp
@@ -92,7 +92,7 @@ void OSLShaderManager::device_update_specific(Device *device,
     }
   });
 
-  VLOG(1) << "Total " << scene->shaders.size() << " shaders.";
+  VLOG_INFO << "Total " << scene->shaders.size() << " shaders.";
 
   device_free(device, dscene, scene);
 
@@ -240,7 +240,7 @@ void OSLShaderManager::shading_system_init()
     ss_shared->attribute("searchpath:shader", shader_path);
     ss_shared->attribute("greedyjit", 1);
 
-    VLOG(1) << "Using shader search path: " << shader_path;
+    VLOG_INFO << "Using shader search path: " << shader_path;
 
     /* our own ray types */
     static const char *raytypes[] = {
diff --git a/intern/cycles/scene/particles.cpp b/intern/cycles/scene/particles.cpp
index a05acf17ccf..b527dd0ebe8 100644
--- a/intern/cycles/scene/particles.cpp
+++ b/intern/cycles/scene/particles.cpp
@@ -105,7 +105,7 @@ void ParticleSystemManager::device_update(Device *device,
     }
   });
 
-  VLOG(1) << "Total " << scene->particle_systems.size() << " particle systems.";
+  VLOG_INFO << "Total " << scene->particle_systems.size() << " particle systems.";
 
   device_free(device, dscene);
 
diff --git a/intern/cycles/scene/scene.cpp b/intern/cycles/scene/scene.cpp
index 8b5604eba72..33ce0a45733 100644
--- a/intern/cycles/scene/scene.cpp
+++ b/intern/cycles/scene/scene.cpp
@@ -373,11 +373,11 @@ void Scene::device_update(Device *device_, Progress &progress)
     size_t mem_used = util_guarded_get_mem_used();
     size_t mem_peak = util_guarded_get_mem_peak();
 
-    VLOG(1) << "System memory statistics after full device sync:\n"
-            << "  Usage: " << string_human_readable_number(mem_used) << " ("
-            << string_human_readable_size(mem_used) << ")\n"
-            << "  Peak: " << string_human_readable_number(mem_peak) << " ("
-            << string_human_readable_size(mem_peak) << ")";
+    VLOG_INFO << "System memory statistics after full device sync:\n"
+              << "  Usage: " << string_human_readable_number(mem_used) << " ("
+              << string_human_readable_size(mem_used) << ")\n"
+              << "  Peak: " << string_human_readable_number(mem_peak) << " ("
+              << string_human_readable_size(mem_peak) << ")";
   }
 }
 
@@ -586,35 +586,38 @@ bool Scene::update(Progress &progress)
 
 static void log_kernel_features(const uint features)
 {
-  VLOG(2) << "Requested features:\n";
-  VLOG(2) << "Use BSDF " << string_from_bool(features & KERNEL_FEATURE_NODE_BSDF) << "\n";
-  VLOG(2) << "Use Principled BSDF " << string_from_bool(features & KERNEL_FEATURE_PRINCIPLED)
-          << "\n";
-  VLOG(2) << "Use Emission " << string_from_bool(features & KERNEL_FEATURE_NODE_EMISSION) << "\n";
-  VLOG(2) << "Use Volume " << string_from_bool(features & KERNEL_FEATURE_NODE_VOLUME) << "\n";
-  VLOG(2) << "Use Bump " << string_from_bool(features & KERNEL_FEATURE_NODE_BUMP) << "\n";
-  VLOG(2) << "Use Voronoi " << string_from_bool(features & KERNEL_FEATURE_NODE_VORONOI_EXTRA)
-          << "\n";
-  VLOG(2) << "Use Shader Raytrace " << string_from_bool(features & KERNEL_FEATURE_NODE_RAYTRACE)
-          << "\n";
-  VLOG(2) << "Use MNEE" << string_from_bool(features & KERNEL_FEATURE_MNEE) << "\n";
-  VLOG(2) << "Use Transparent " << string_from_bool(features & KERNEL_FEATURE_TRANSPARENT) << "\n";
-  VLOG(2) << "Use Denoising " << string_from_bool(features & KERNEL_FEATURE_DENOISING) << "\n";
-  VLOG(2) << "Use Path Tracing " << string_from_bool(features & KERNEL_FEATURE_PATH_TRACING)
-          << "\n";
-  VLOG(2) << "Use Hair " << string_from_bool(features & KERNEL_FEATURE_HAIR) << "\n";
-  VLOG(2) << "Use Pointclouds " << string_from_bool(features & KERNEL_FEATURE_POINTCLOUD) << "\n";
-  VLOG(2) << "Use Object Motion " << string_from_bool(features & KERNEL_FEATURE_OBJECT_MOTION)
-          << "\n";
-  VLOG(2) << "Use Camera Motion " << string_from_bool(features & KERNEL_FEATURE_CAMERA_MOTION)
-          << "\n";
-  VLOG(2) << "Use Baking " << string_from_bool(features & KERNEL_FEATURE_BAKING) << "\n";
-  VLOG(2) << "Use Subsurface " << string_from_bool(features & KERNEL_FEATURE_SUBSURFACE) << "\n";
-  VLOG(2) << "Use Volume " << string_from_bool(features & KERNEL_FEATURE_VOLUME) << "\n";
-  VLOG(2) << "Use Patch Evaluation "
-          << string_from_bool(features & KERNEL_FEATURE_PATCH_EVALUATION) << "\n";
-  VLOG(2) << "Use Shadow Catcher " << string_from_bool(features & KERNEL_FEATURE_SHADOW_CATCHER)
-          << "\n";
+  VLOG_INFO << "Requested features:\n";
+  VLOG_INFO << "Use BSDF " << string_from_bool(features & KERNEL_FEATURE_NODE_BSDF) << "\n";
+  VLOG_INFO << "Use Principled BSDF " << string_from_bool(features & KERNEL_FEATURE_PRINCIPLED)
+            << "\n";
+  VLOG_INFO << "Use Emission " << string_from_bool(features & KERNEL_FEATURE_NODE_EMISSION)
+            << "\n";
+  VLOG_INFO << "Use Volume " << string_from_bool(features & KERNEL_FEATURE_NODE_VOLUME) << "\n";
+  VLOG_INFO << "Use Bump " << string_from_bool(features & KERNEL_FEATURE_NODE_BUMP) << "\n";
+  VLOG_INFO << "Use Voronoi " << string_from_bool(features & KERNEL_FEATURE_NODE_VORONOI_EXTRA)
+            << "\n";
+  VLOG_INFO << "Use Shader Raytrace " << string_from_bool(features & KERNEL_FEATURE_NODE_RAYTRACE)
+            << "\n";
+  VLOG_INFO << "Use MNEE" << string_from_bool(features & KERNEL_FEATURE_MNEE) << "\n";
+  VLOG_INFO << "Use Transparent " << string_from_bool(features & KERNEL_FEATURE_TRANSPARENT)
+            << "\n";
+  VLOG_INFO << "Use Denoising " << string_from_bool(features & KERNEL_FEATURE_DENOISING) << "\n";
+  VLOG_INFO << "Use Path Tracing " << string_from_bool(features & KERNEL_FEATURE_PATH_TRACING)
+            << "\n";
+  VLOG_INFO << "Use Hair " << string_from_bool(features & KERNEL_FEATURE_HAIR) << "\n";
+  VLOG_INFO << "Use Pointclouds " << string_from_bool(features & KERNEL_FEATURE_POINTCLOUD)
+            << "\n";
+  VLOG_INFO << "Use Object Motion " << string_from_bool(features & KERNEL_FEATURE_OBJECT_MOTION)
+            << "\n";
+  VLOG_INFO << "Use Camera Motion " << string_from_bool(features & KERNEL_FEATURE_CAMERA_MOTION)
+            << "\n";
+  VLOG_INFO << "Use Baking " << string_from_bool(features & KERNEL_FEATURE_BAKING) << "\n";
+  VLOG_INFO << "Use Subsurface " << string_from_bool(features & KERNEL_FEATURE_SUBSURFACE) << "\n";
+  VLOG_INFO << "Use Volume " << string_from_bool(features & KERNEL_FEATURE_VOLUME) << "\n";
+  VLOG_INFO << "Use Patch Evaluation "
+            << string_from_bool(features & KERNEL_FEATURE_PATCH_EVALUATION) << "\n";
+  VLOG_INFO << "Use Shadow Catcher " << string_from_bool(features & KERNEL_FEATURE_SHADOW_CATCHER)
+            << "\n";
 }
 
 bool Scene::load_kernels(Progress &progress, bool lock_scene)
@@ -675,8 +678,8 @@ int Scene::get_max_closure_count()
      * closures discarded due to mixing or low weights. We need to limit
      * to MAX_CLOSURE as this is hardcoded in CPU/mega kernels, and it
      * avoids excessive memory usage for split kernels. */
-    VLOG(2) << "Maximum number of closures exceeded: " << max_closure_global << " > "
-            << MAX_CLOSURE;
+    VLOG_WARNING << "Maximum number of closures exceeded: " << max_closure_global << " > "
+                 << MAX_CLOSURE;
 
     max_closure_global = MAX_CLOSURE;
   }
@@ -723,7 +726,7 @@ int Scene::get_volume_stack_size() const
 
   volume_stack_size = min(volume_stack_size, MAX_VOLUME_STACK_SIZE);
 
-  VLOG(3) << "Detected required volume stack size " << volume_stack_size;
+  VLOG_WORK << "Detected required volume stack size " << volume_stack_size;
 
   return volume_stack_size;
 }
diff --git a/intern/cycles/scene/scene.h b/intern/cycles/scene/scene.h
index a0d2f4a6c06..d04c6a27f11 100644
--- a/intern/cycles/scene/scene.h
+++ b/intern/cycles/scene/scene.h
@@ -98,7 +98,7 @@ class DeviceScene {
   device_vector<DecomposedTransform> camera_motion;
 
   /* attributes */
-  device_vector<uint4> attributes_map;
+  device_vector<AttributeMap> attributes_map;
   device_vector<float> attributes_float;
   device_vector<float2> attributes_float2;
   device_vector<packed_float3> attributes_float3;
diff --git a/intern/cycles/scene/shader_graph.cpp b/intern/cycles/scene/shader_graph.cpp
index f25d0b7c7b9..ef3f142ed4e 100644
--- a/intern/cycles/scene/shader_graph.cpp
+++ b/intern/cycles/scene/shader_graph.cpp
@@ -659,7 +659,7 @@ void ShaderGraph::deduplicate_nodes()
   }
 
   if (num_deduplicated > 0) {
-    VLOG(1) << "Deduplicated " << num_deduplicated << " nodes.";
+    VLOG_DEBUG << "Deduplicated " << num_deduplicated << " nodes.";
   }
 }
 
@@ -700,7 +700,7 @@ void ShaderGraph::verify_volume_output()
     }
   }
   if (!has_valid_volume) {
-    VLOG(1) << "Disconnect meaningless volume output.";
+    VLOG_DEBUG << "Disconnect meaningless volume output.";
     disconnect(volume_in->link);
   }
 }
diff --git a/intern/cycles/scene/shader_nodes.cpp b/intern/cycles/scene/shader_nodes.cpp
index 3b58556f601..f93a1a5231a 100644
--- a/intern/cycles/scene/shader_nodes.cpp
+++ b/intern/cycles/scene/shader_nodes.cpp
@@ -2391,7 +2391,7 @@ void GlossyBsdfNode::simplify_settings(Scene *scene)
      * NOTE: Keep the epsilon in sync with kernel!
      */
     if (!roughness_input->link && roughness <= 1e-4f) {
-      VLOG(3) << "Using sharp glossy BSDF.";
+      VLOG_DEBUG << "Using sharp glossy BSDF.";
       distribution = CLOSURE_BSDF_REFLECTION_ID;
     }
   }
@@ -2400,7 +2400,7 @@ void GlossyBsdfNode::simplify_settings(Scene *scene)
      * benefit from closure blur to remove unwanted noise.
      */
     if (roughness_input->link == NULL && distribution == CLOSURE_BSDF_REFLECTION_ID) {
-      VLOG(3) << "Using GGX glossy with filter glossy.";
+      VLOG_DEBUG << "Using GGX glossy with filter glossy.";
       distribution = CLOSURE_BSDF_MICROFACET_GGX_ID;
       roughness = 0.0f;
     }
@@ -2484,7 +2484,7 @@ void GlassBsdfNode::simplify_settings(Scene *scene)
      * NOTE: Keep the epsilon in sync with kernel!
      */
     if (!roughness_input->link && roughness <= 1e-4f) {
-      VLOG(3) << "Using sharp glass BSDF.";
+      VLOG_DEBUG << "Using sharp glass BSDF.";
       distribution = CLOSURE_BSDF_SHARP_GLASS_ID;
     }
   }
@@ -2493,7 +2493,7 @@ void GlassBsdfNode::simplify_settings(Scene *scene)
      * benefit from closure blur to remove unwanted noise.
      */
     if (roughness_input->link == NULL && distribution == CLOSURE_BSDF_SHARP_GLASS_ID) {
-      VLOG(3) << "Using GGX glass with filter glossy.";
+      VLOG_DEBUG << "Using GGX glass with filter glossy.";
       distribution = CLOSURE_BSDF_MICROFACET_GGX_GLASS_ID;
       roughness = 0.0f;
     }
@@ -2577,7 +2577,7 @@ void RefractionBsdfNode::simplify_settings(Scene *scene)
      * NOTE: Keep the epsilon in sync with kernel!
      */
     if (!roughness_input->link && roughness <= 1e-4f) {
-      VLOG(3) << "Using sharp refraction BSDF.";
+      VLOG_DEBUG << "Using sharp refraction BSDF.";
       distribution = CLOSURE_BSDF_REFRACTION_ID;
     }
   }
@@ -2586,7 +2586,7 @@ void RefractionBsdfNode::simplify_settings(Scene *scene)
      * benefit from closure blur to remove unwanted noise.
      */
     if (roughness_input->link == NULL && distribution == CLOSURE_BSDF_REFRACTION_ID) {
-      VLOG(3) << "Using GGX refraction with filter glossy.";
+      VLOG_DEBUG << "Using GGX refraction with filter glossy.";
       distribution = CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID;
       roughness = 0.0f;
     }
diff --git a/intern/cycles/scene/svm.cpp b/intern/cycles/scene/svm.cpp
index 484a7d6de72..4bc5a1b9cc2 100644
--- a/intern/cycles/scene/svm.cpp
+++ b/intern/cycles/scene/svm.cpp
@@ -51,9 +51,9 @@ void SVMShaderManager::device_update_shader(Scene *scene,
   compiler.background = (shader == scene->background->get_shader(scene));
   compiler.compile(shader, *svm_nodes, 0, &summary);
 
-  VLOG(3) << "Compilation summary:\n"
-          << "Shader name: " << shader->name << "\n"
-          << summary.full_report();
+  VLOG_WORK << "Compilation summary:\n"
+            << "Shader name: " << shader->name << "\n"
+            << summary.full_report();
 }
 
 void SVMShaderManager::device_update_specific(Device *device,
@@ -72,7 +72,7 @@ void SVMShaderManager::device_update_specific(Device *device,
 
   const int num_shaders = scene->shaders.size();
 
-  VLOG(1) << "Total " << num_shaders << " shaders.";
+  VLOG_INFO << "Total " << num_shaders << " shaders.";
 
   double start_time = time_dt();
 
@@ -148,8 +148,8 @@ void SVMShaderManager::device_update_specific(Device *device,
 
   update_flags = UPDATE_NONE;
 
-  VLOG(1) << "Shader manager updated " << num_shaders << " shaders in " << time_dt() - start_time
-          << " seconds.";
+  VLOG_INFO << "Shader manager updated " << num_shaders << " shaders in " << time_dt() - start_time
+            << " seconds.";
 }
 
 void SVMShaderManager::device_free(Device *device, DeviceScene *dscene, Scene *scene)
diff --git a/intern/cycles/scene/tables.cpp b/intern/cycles/scene/tables.cpp
index bb35880d37e..3e57b535f35 100644
--- a/intern/cycles/scene/tables.cpp
+++ b/intern/cycles/scene/tables.cpp
@@ -34,7 +34,7 @@ void LookupTables::device_update(Device *, DeviceScene *dscene, Scene *scene)
     }
   });
 
-  VLOG(1) << "Total " << lookup_tables.size() << " lookup tables.";
+  VLOG_INFO << "Total " << lookup_tables.size() << " lookup tables.";
 
   if (lookup_tables.size() > 0)
     dscene->lookup_table.copy_to_device();
diff --git a/intern/cycles/scene/volume.cpp b/intern/cycles/scene/volume.cpp
index 39e9b0bbbf4..77955350305 100644
--- a/intern/cycles/scene/volume.cpp
+++ b/intern/cycles/scene/volume.cpp
@@ -754,11 +754,11 @@ void GeometryManager::create_volume_mesh(const Scene *scene, Volume *volume, Pro
   }
 
   /* Print stats. */
-  VLOG(1) << "Memory usage volume mesh: "
-          << ((vertices.size() + face_normals.size()) * sizeof(float3) +
-              indices.size() * sizeof(int)) /
-                 (1024.0 * 1024.0)
-          << "Mb.";
+  VLOG_WORK << "Memory usage volume mesh: "
+            << ((vertices.size() + face_normals.size()) * sizeof(float3) +
+                indices.size() * sizeof(int)) /
+                   (1024.0 * 1024.0)
+            << "Mb.";
 }
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/session/session.cpp b/intern/cycles/session/session.cpp
index ef177636046..ba30be817cf 100644
--- a/intern/cycles/session/session.cpp
+++ b/intern/cycles/session/session.cpp
@@ -146,11 +146,11 @@ void Session::run_main_render_loop()
     RenderWork render_work = run_update_for_next_iteration();
 
     if (!render_work) {
-      if (VLOG_IS_ON(2)) {
+      if (VLOG_WORK_IS_ON) {
         double total_time, render_time;
         progress.get_time(total_time, render_time);
-        VLOG(2) << "Rendering in main loop is done in " << render_time << " seconds.";
-        VLOG(2) << path_trace_->full_report();
+        VLOG_WORK << "Rendering in main loop is done in " << render_time << " seconds.";
+        VLOG_WORK << path_trace_->full_report();
       }
 
       if (params.background) {
diff --git a/intern/cycles/session/tile.cpp b/intern/cycles/session/tile.cpp
index 82272a7dbf5..f4930cbb945 100644
--- a/intern/cycles/session/tile.cpp
+++ b/intern/cycles/session/tile.cpp
@@ -335,7 +335,7 @@ int TileManager::compute_render_tile_size(const int suggested_tile_size) const
 
 void TileManager::reset_scheduling(const BufferParams &params, int2 tile_size)
 {
-  VLOG(3) << "Using tile size of " << tile_size;
+  VLOG_WORK << "Using tile size of " << tile_size;
 
   close_tile_output();
 
@@ -466,7 +466,7 @@ bool TileManager::open_tile_output()
 
   write_state_.num_tiles_written = 0;
 
-  VLOG(3) << "Opened tile file " << write_state_.filename;
+  VLOG_WORK << "Opened tile file " << write_state_.filename;
 
   return true;
 }
@@ -485,7 +485,7 @@ bool TileManager::close_tile_output()
     return false;
   }
 
-  VLOG(3) << "Tile output is closed.";
+  VLOG_WORK << "Tile output is closed.";
 
   return true;
 }
@@ -536,7 +536,7 @@ bool TileManager::write_tile(const RenderBuffers &tile_buffers)
     pixels = pixel_storage.data();
   }
 
-  VLOG(3) << "Write tile at " << tile_x << ", " << tile_y;
+  VLOG_WORK << "Write tile at " << tile_x << ", " << tile_y;
 
   /* The image tile sizes in the OpenEXR file are different from the size of our big tiles. The
    * write_tiles() method expects a contiguous image region that will be split into tiles
@@ -567,7 +567,7 @@ bool TileManager::write_tile(const RenderBuffers &tile_buffers)
 
   ++write_state_.num_tiles_written;
 
-  VLOG(3) << "Tile written in " << time_dt() - time_start << " seconds.";
+  VLOG_WORK << "Tile written in " << time_dt() - time_start << " seconds.";
 
   return true;
 }
@@ -591,7 +591,7 @@ void TileManager::finish_write_tiles()
       const int tile_x = tile.x + tile.window_x;
       const int tile_y = tile.y + tile.window_y;
 
-      VLOG(3) << "Write dummy tile at " << tile_x << ", " << tile_y;
+      VLOG_WORK << "Write dummy tile at " << tile_x << ", " << tile_y;
 
       write_state_.tile_out->write_tiles(tile_x,
                                          tile_x + tile.window_width,
@@ -610,8 +610,8 @@ void TileManager::finish_write_tiles()
     full_buffer_written_cb(write_state_.filename);
   }
 
-  VLOG(3) << "Tile file size is "
-          << string_human_readable_number(path_file_size(write_state_.filename)) << " bytes.";
+  VLOG_WORK << "Tile file size is "
+            << string_human_readable_number(path_file_size(write_state_.filename)) << " bytes.";
 
   /* Advance the counter upon explicit finish of the file.
    * Makes it possible to re-use tile manager for another scene, and avoids unnecessary increments
diff --git a/intern/cycles/test/render_graph_finalize_test.cpp b/intern/cycles/test/render_graph_finalize_test.cpp
index dac36ab0135..2dc13f0eccb 100644
--- a/intern/cycles/test/render_graph_finalize_test.cpp
+++ b/intern/cycles/test/render_graph_finalize_test.cpp
@@ -166,7 +166,7 @@ class RenderGraph : public testing::Test {
   virtual void SetUp()
   {
     util_logging_start();
-    util_logging_verbosity_set(3);
+    util_logging_verbosity_set(5);
 
     device_cpu = Device::create(device_info, stats, profiler);
     scene = new Scene(scene_params, device_cpu);
diff --git a/intern/cycles/util/debug.cpp b/intern/cycles/util/debug.cpp
index 65d108bb9d1..faa54a92c24 100644
--- a/intern/cycles/util/debug.cpp
+++ b/intern/cycles/util/debug.cpp
@@ -25,7 +25,7 @@ void DebugFlags::CPU::reset()
   do { \
     flag = (getenv(env) == NULL); \
     if (!flag) { \
-      VLOG(1) << "Disabling " << STRINGIFY(flag) << " instruction set."; \
+      VLOG_INFO << "Disabling " << STRINGIFY(flag) << " instruction set."; \
     } \
   } while (0)
 
diff --git a/intern/cycles/util/log.h b/intern/cycles/util/log.h
index b33c826d6f5..3780d03c0d1 100644
--- a/intern/cycles/util/log.h
+++ b/intern/cycles/util/log.h
@@ -69,9 +69,22 @@ class LogMessageVoidify {
 #  define LOG_ASSERT(expression) LOG_SUPPRESS()
 #endif
 
-#define VLOG_ONCE(level, flag) \
-  if (!flag) \
-  flag = true, VLOG(level)
+/* Verbose logging categories. */
+
+/* Warnings. */
+#define VLOG_WARNING VLOG(1)
+/* Info about devices, scene contents and features used. */
+#define VLOG_INFO VLOG(2)
+#define VLOG_INFO_IS_ON VLOG_IS_ON(2)
+/* Work being performed and timing/memory stats about that work. */
+#define VLOG_WORK VLOG(3)
+#define VLOG_WORK_IS_ON VLOG_IS_ON(3)
+/* Detailed device timing stats. */
+#define VLOG_DEVICE_STATS VLOG(4)
+#define VLOG_DEVICE_STATS_IS_ON VLOG_IS_ON(4)
+/* Verbose debug messages. */
+#define VLOG_DEBUG VLOG(5)
+#define VLOG_DEBUG_IS_ON VLOG_IS_ON(5)
 
 struct int2;
 struct float3;
diff --git a/intern/cycles/util/task.cpp b/intern/cycles/util/task.cpp
index 2edc82eb7c3..12f661f752d 100644
--- a/intern/cycles/util/task.cpp
+++ b/intern/cycles/util/task.cpp
@@ -70,7 +70,7 @@ void TaskScheduler::init(int num_threads)
   }
   if (num_threads > 0) {
     /* Automatic number of threads. */
-    VLOG(1) << "Overriding number of TBB threads to " << num_threads << ".";
+    VLOG_INFO << "Overriding number of TBB threads to " << num_threads << ".";
     global_control = new tbb::global_control(tbb::global_control::max_allowed_parallelism,
                                              num_threads);
     active_num_threads = num_threads;