1 files changed, 56 insertions, 19 deletions

diff --git a/intern/cycles/device/metal/queue.h b/intern/cycles/device/metal/queue.h
index 6cc84a20787..b0bd487c86d 100644
--- a/intern/cycles/device/metal/queue.h
+++ b/intern/cycles/device/metal/queue.h
@@ -12,8 +12,6 @@
 #  include "device/metal/util.h"
 #  include "kernel/device/metal/globals.h"
 
-#  define metal_printf VLOG(4) << string_printf
-
 CCL_NAMESPACE_BEGIN
 
 class MetalDevice;
@@ -40,43 +38,82 @@ 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 label_command_encoders_ = false;
+
+  /* Per-kernel profiling (see CYCLES_METAL_PROFILING). */
+
+  struct TimingData {
+    DeviceKernel kernel;
+    int work_size;
+    uint64_t 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;
+
+  /* .gputrace capture (see CYCLES_DEBUG_METAL_CAPTURE_...). */
+
+  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