1 files changed, 13 insertions, 8 deletions
diff --git a/intern/cycles/device/metal/queue.mm b/intern/cycles/device/metal/queue.mm
index 0e260886abb..5ac63a16c61 100644
--- a/intern/cycles/device/metal/queue.mm
+++ b/intern/cycles/device/metal/queue.mm
@@ -293,6 +293,11 @@ int MetalDeviceQueue::num_concurrent_busy_states() const
   return result;
 }
 
+int MetalDeviceQueue::num_sort_partition_elements() const
+{
+  return MetalInfo::optimal_sort_partition_elements(metal_device_->mtlDevice);
+}
+
 void MetalDeviceQueue::init_execution()
 {
   /* Synchronize all textures and memory copies before executing task. */
@@ -311,8 +316,8 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
     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);
 
@@ -358,8 +363,8 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
 
   /* Prepare any non-pointer (i.e. plain-old-data) KernelParamsMetal data */
   /* The plain-old-data is contiguous, continuing to the end of KernelParamsMetal */
-  size_t plain_old_launch_data_offset = offsetof(KernelParamsMetal, __integrator_state) +
-                                        sizeof(IntegratorStateGPU);
+  size_t plain_old_launch_data_offset = offsetof(KernelParamsMetal, integrator_state) +
+                                        offsetof(IntegratorStateGPU, sort_partition_divisor);
   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,
@@ -415,8 +420,8 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
   }
 
   /* this relies on IntegratorStateGPU layout being contiguous device_ptrs  */
-  const size_t pointer_block_end = offsetof(KernelParamsMetal, __integrator_state) +
-                                   sizeof(IntegratorStateGPU);
+  const size_t pointer_block_end = offsetof(KernelParamsMetal, integrator_state) +
+                                   offsetof(IntegratorStateGPU, sort_partition_divisor);
   for (size_t offset = 0; offset < pointer_block_end; offset += sizeof(device_ptr)) {
     int pointer_index = int(offset / sizeof(device_ptr));
     MetalDevice::MetalMem *mmem = *(
@@ -550,7 +555,7 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
     /* Enhanced command buffer errors are only available in 11.0+ */
     if (@available(macos 11.0, *)) {
       if (command_buffer.status == MTLCommandBufferStatusError && command_buffer.error != nil) {
-        printf("CommandBuffer Failed: %s\n", [kernel_name UTF8String]);
+        metal_device_->set_error(string("CommandBuffer Failed: ") + [kernel_name UTF8String]);
         NSArray<id<MTLCommandBufferEncoderInfo>> *encoderInfos = [command_buffer.error.userInfo
             valueForKey:MTLCommandBufferEncoderInfoErrorKey];
         if (encoderInfos != nil) {
@@ -564,7 +569,7 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
         }
       }
       else if (command_buffer.error) {
-        printf("CommandBuffer Failed: %s\n", [kernel_name UTF8String]);
+        metal_device_->set_error(string("CommandBuffer Failed: ") + [kernel_name UTF8String]);
       }
     }
   }];