[client/host/idd] correct frame timing attribution

Use calibrated copy-queue timestamps to separate source and effect
waits from the actual framebuffer copy.

Exclude producer readiness waits from client import timing so the
same interval is not counted in both Copy and Import.

client/renderers/EGL/desktop.c

@@ -391,11 +391,13 @@ bool egl_desktopSetup(EGL_Desktop * desktop, const LG_RendererFormat format)
 }
 
 bool egl_desktopUpdate(EGL_Desktop * desktop, const FrameBuffer * frame, int dmaFd,
-    const FrameDamageRect * damageRects, int damageRectsCount)
+    const FrameDamageRect * damageRects, int damageRectsCount,
+    uint64_t * waitTimeNs)
 {
   if (likely(desktop->useDMA && dmaFd >= 0))
   {
-    if (likely(egl_textureUpdateFromDMA(desktop->texture, frame, dmaFd)))
+    if (likely(egl_textureUpdateFromDMA(
+            desktop->texture, frame, dmaFd, waitTimeNs)))
     {
       atomic_store(&desktop->processFrame, true);
       return true;
@@ -437,7 +439,7 @@ bool egl_desktopUpdate(EGL_Desktop * desktop, const FrameBuffer * frame, int dma
    * previous texture contents instead.
    */
   if (likely(egl_textureUpdateFromFrame(desktop->texture, frame,
-        damageRects, damageRectsCount)))
+        damageRects, damageRectsCount, waitTimeNs)))
     return true;
 
   return false;

client/renderers/EGL/desktop.h

@@ -49,7 +49,8 @@ void egl_desktopGetHDRMapping(EGL_Desktop * desktop, bool * enabled,
     float * gain, float * contentPeak);
 bool egl_desktopSetup (EGL_Desktop * desktop, const LG_RendererFormat format);
 bool egl_desktopUpdate(EGL_Desktop * desktop, const FrameBuffer * frame, int dmaFd,
-    const FrameDamageRect * damageRects, int damageRectsCount);
+    const FrameDamageRect * damageRects, int damageRectsCount,
+    uint64_t * waitTimeNs);
 void egl_desktopResize(EGL_Desktop * desktop, int width, int height);
 bool egl_desktopRender(EGL_Desktop * desktop, unsigned int outputWidth,
     unsigned int outputHeight, const float x, const float y,

client/renderers/EGL/egl.c

@@ -749,13 +749,20 @@ static bool egl_onFrame(LG_Renderer * renderer, const FrameBuffer * frame, int d
   egl_stateCheckShared();
 
   const uint64_t start      = nanotime();
+  uint64_t       waitTimeNs = 0;
   if (unlikely(!egl_desktopUpdate(
-          this->desktop, frame, dmaFd, damageRects, damageRectsCount)))
+          this->desktop, frame, dmaFd, damageRects, damageRectsCount,
+          &waitTimeNs)))
   {
     DEBUG_INFO("Failed to to update the desktop");
     return false;
   }
-  app_setFrameImportTime(nanotime() - start);
+  const uint64_t elapsed = nanotime() - start;
+
+  /* Producer Copy already covers the interval before FrameBuffer::wp becomes
+   * ready. Exclude that overlapping wait from the client Import stage. */
+  app_setFrameImportTime(
+      elapsed > waitTimeNs ? elapsed - waitTimeNs : 0);
 
   INTERLOCKED_SECTION(this->desktopDamageLock, {
     struct DesktopDamage * damage = this->desktopDamage + this->desktopDamageIdx;

client/renderers/EGL/texture.c

@@ -161,7 +161,7 @@ bool egl_textureUpdateRect(EGL_Texture * this,
 
 bool egl_textureUpdateFromFrame(EGL_Texture * this,
     const FrameBuffer * frame, const FrameDamageRect * damageRects,
-    int damageRectsCount)
+    int damageRectsCount, uint64_t * waitTimeNs)
 {
   const struct EGL_TexUpdate update =
   {
@@ -175,13 +175,14 @@ bool egl_textureUpdateFromFrame(EGL_Texture * this,
     .frame      = frame,
     .rects      = damageRects,
     .rectCount  = damageRectsCount,
+    .waitTimeNs = waitTimeNs,
   };
 
   return this->ops.update(this, &update);
 }
 
 bool egl_textureUpdateFromDMA(EGL_Texture * this,
-    const FrameBuffer * frame, const int dmaFd)
+    const FrameBuffer * frame, const int dmaFd, uint64_t * waitTimeNs)
 {
   const struct EGL_TexUpdate update =
   {
@@ -196,7 +197,8 @@ bool egl_textureUpdateFromDMA(EGL_Texture * this,
   };
 
   /* wait for completion */
-  if (unlikely(!framebuffer_wait(frame, this->format.dataSize)))
+  if (unlikely(!framebuffer_wait_timed(
+          frame, this->format.dataSize, waitTimeNs)))
     return false;
 
   return this->ops.update(this, &update);

client/renderers/EGL/texture.h

@@ -41,6 +41,7 @@ typedef struct EGL_TexUpdate
 {
   /* the type of this update */
   EGL_TexType type;
+  uint64_t  * waitTimeNs;
 
   int x, y, width, height;
 
@@ -125,10 +126,10 @@ bool egl_textureUpdateRect(EGL_Texture * texture,
 
 bool egl_textureUpdateFromFrame(EGL_Texture * texture,
     const FrameBuffer * frame, const FrameDamageRect * damageRects,
-    int damageRectsCount);
+    int damageRectsCount, uint64_t * waitTimeNs);
 
 bool egl_textureUpdateFromDMA(EGL_Texture * texture,
-    const FrameBuffer * frame, const int dmaFd);
+    const FrameBuffer * frame, const int dmaFd, uint64_t * waitTimeNs);
 
 enum EGL_TexStatus egl_textureProcess(EGL_Texture * texture);
 

client/renderers/EGL/texture_framebuffer.c

@@ -94,14 +94,15 @@ static bool egl_texFBUpdate(EGL_Texture * texture, const EGL_TexUpdate * update)
 
   if (damageAll)
   {
-    complete = framebuffer_read(
+    complete = framebuffer_read_timed(
       update->frame,
       parent->buf[parent->bufIndex].map,
       texture->format.pitch,
       texture->format.height,
       texture->format.width,
       texture->format.bpp,
-      texture->format.pitch
+      texture->format.pitch,
+      update->waitTimeNs
     );
   }
   else
@@ -123,7 +124,7 @@ static bool egl_texFBUpdate(EGL_Texture * texture, const EGL_TexUpdate * update)
         scaledDamageRects[i] = rect;
       }
 
-      complete = rectsFramebufferToBuffer(
+      complete = rectsFramebufferToBufferTimed(
         scaledDamageRects,
         damage->count,
         texture->format.bpp,
@@ -131,12 +132,13 @@ static bool egl_texFBUpdate(EGL_Texture * texture, const EGL_TexUpdate * update)
         texture->format.pitch,
         texture->format.height,
         update->frame,
-        texture->format.pitch
+        texture->format.pitch,
+        update->waitTimeNs
       );
     }
     else
     {
-      complete = rectsFramebufferToBuffer(
+      complete = rectsFramebufferToBufferTimed(
         damage->rects,
         damage->count,
         texture->format.bpp,
@@ -144,7 +146,8 @@ static bool egl_texFBUpdate(EGL_Texture * texture, const EGL_TexUpdate * update)
         texture->format.pitch,
         texture->format.height,
         update->frame,
-        texture->format.pitch
+        texture->format.pitch,
+        update->waitTimeNs
       );
     }
   }

client/transports/LGMP/lgmp.c

@@ -622,7 +622,7 @@ static void lgmp_getFrameTiming(LG_Transport * this,
       this->pendingFrame->frameSerial != frame->serial)
     return;
 
-  /* The producer writes these immediately before completing the framebuffer.
+  /* The producer writes these immediately after completing the framebuffer.
    * nextFrame can observe the header earlier, so sample them only after the
    * renderer's onFrame call has consumed the framebuffer. */
   for (unsigned i = 0;

common/include/common/framebuffer.h

@@ -44,6 +44,13 @@ typedef bool (*FrameBufferReadFn)(void * opaque, const void * src, size_t size);
  */
 bool framebuffer_wait(const FrameBuffer * frame, size_t size);
 
+/**
+ * Wait for the framebuffer to fill to the specified size and accumulate the
+ * nanoseconds spent waiting for the producer in `waitTimeNs`.
+ */
+bool framebuffer_wait_timed(const FrameBuffer * frame, size_t size,
+    uint64_t * waitTimeNs);
+
 /**
  * Read `size` bytes from the KVMFRFrame into the dst buffer
  */
@@ -56,6 +63,14 @@ bool framebuffer_read_linear(const FrameBuffer * frame, void * restrict dst,
 bool framebuffer_read(const FrameBuffer * frame, void * dst, size_t dstpitch,
     size_t height, size_t width, size_t bpp, size_t pitch);
 
+/**
+ * Read data from the KVMFRFrame and accumulate the nanoseconds spent waiting
+ * for the producer in `waitTimeNs`.
+ */
+bool framebuffer_read_timed(const FrameBuffer * frame, void * dst,
+    size_t dstpitch, size_t height, size_t width, size_t bpp, size_t pitch,
+    uint64_t * waitTimeNs);
+
 /**
  * Read data from the KVMFRFrame using a callback
  */

common/include/common/rects.h

@@ -39,6 +39,11 @@ bool rectsFramebufferToBuffer(FrameDamageRect * rects, int count, int bpp,
   uint8_t * dst, int dstPitch, int height,
   const FrameBuffer * frame, int srcPitch);
 
+/* As above, accumulating producer wait time in nanoseconds. */
+bool rectsFramebufferToBufferTimed(FrameDamageRect * rects, int count, int bpp,
+  uint8_t * dst, int dstPitch, int height,
+  const FrameBuffer * frame, int srcPitch, uint64_t * waitTimeNs);
+
 int rectsMergeOverlapping(FrameDamageRect * rects, int count);
 int rectsRejectContained(FrameDamageRect * rects, int count);
 

common/src/framebuffer.c

@@ -21,6 +21,7 @@
 #include "common/framebuffer.h"
 #include "common/cpuinfo.h"
 #include "common/debug.h"
+#include "common/time.h"
 
 //#define FB_PROFILE
 #ifdef FB_PROFILE
@@ -33,24 +34,42 @@
 #include <immintrin.h>
 #include <unistd.h>
 
-bool framebuffer_wait(const FrameBuffer * frame, size_t size)
+bool framebuffer_wait_timed(const FrameBuffer * frame, size_t size,
+    uint64_t * waitTimeNs)
 {
+  if (atomic_load_explicit(&frame->wp, memory_order_acquire) >= size)
+    return true;
+
+  const uint64_t waitStart = waitTimeNs ? nanotime() : 0;
+
   while(atomic_load_explicit(&frame->wp, memory_order_acquire) < size)
   {
     int spinCount = 0;
     while(frame->wp < size)
     {
       if (++spinCount == FB_SPIN_LIMIT)
+      {
+        if (waitTimeNs)
+          *waitTimeNs += nanotime() - waitStart;
         return false;
+      }
       usleep(1);
     }
   }
 
+  if (waitTimeNs)
+    *waitTimeNs += nanotime() - waitStart;
+
   return true;
 }
 
-bool framebuffer_read_linear(const FrameBuffer * frame, void * restrict dst,
+bool framebuffer_wait(const FrameBuffer * frame, size_t size)
-    size_t size)
+{
+  return framebuffer_wait_timed(frame, size, NULL);
+}
+
+static bool framebuffer_read_linear_timed(const FrameBuffer * frame,
+    void * restrict dst, size_t size, uint64_t * waitTimeNs)
 {
 #ifdef FB_PROFILE
   static RunningAvg ra = NULL;
@@ -67,7 +86,7 @@ bool framebuffer_read_linear(const FrameBuffer * frame, void * restrict dst,
   while(size)
   {
     const size_t copy = size < FB_CHUNK_SIZE ? size : FB_CHUNK_SIZE;
-    if (!framebuffer_wait(frame, rp + copy))
+    if (!framebuffer_wait_timed(frame, rp + copy, waitTimeNs))
       return false;
 
     memcpy(d, frame->data + rp, copy);
@@ -85,11 +104,19 @@ bool framebuffer_read_linear(const FrameBuffer * frame, void * restrict dst,
   return true;
 }
 
-bool framebuffer_read(const FrameBuffer * frame, void * restrict dst,
+bool framebuffer_read_linear(const FrameBuffer * frame, void * restrict dst,
-    size_t dstpitch, size_t height, size_t width, size_t bpp, size_t pitch)
+    size_t size)
+{
+  return framebuffer_read_linear_timed(frame, dst, size, NULL);
+}
+
+bool framebuffer_read_timed(const FrameBuffer * frame, void * restrict dst,
+    size_t dstpitch, size_t height, size_t width, size_t bpp, size_t pitch,
+    uint64_t * waitTimeNs)
 {
   if (dstpitch == pitch)
-    return framebuffer_read_linear(frame, dst, height * pitch);
+    return framebuffer_read_linear_timed(
+        frame, dst, height * pitch, waitTimeNs);
 
 #ifdef FB_PROFILE
   static RunningAvg ra = NULL;
@@ -106,7 +133,7 @@ bool framebuffer_read(const FrameBuffer * frame, void * restrict dst,
   const size_t linewidth = width * bpp;
   for(size_t y = 0; y < height; ++y)
   {
-    if (!framebuffer_wait(frame, rp + linewidth))
+    if (!framebuffer_wait_timed(frame, rp + linewidth, waitTimeNs))
       return false;
 
     memcpy(d, frame->data + rp, dstpitch);
@@ -123,6 +150,13 @@ bool framebuffer_read(const FrameBuffer * frame, void * restrict dst,
   return true;
 }
 
+bool framebuffer_read(const FrameBuffer * frame, void * restrict dst,
+    size_t dstpitch, size_t height, size_t width, size_t bpp, size_t pitch)
+{
+  return framebuffer_read_timed(frame, dst, dstpitch, height, width, bpp,
+      pitch, NULL);
+}
+
 bool framebuffer_read_fn(const FrameBuffer * frame, size_t height, size_t width,
     size_t bpp, size_t pitch, FrameBufferReadFn fn, void * opaque)
 {

common/src/rects.c

@@ -221,21 +221,36 @@ struct FromFramebufferData
 {
   const FrameBuffer * frame;
   int                 pitch;
+  uint64_t          * waitTimeNs;
 };
 
 static bool fbRowStart(int y, void * opaque)
 {
   struct FromFramebufferData * data = opaque;
-  return framebuffer_wait(data->frame, y * data->pitch);
+  return framebuffer_wait_timed(
+      data->frame, y * data->pitch, data->waitTimeNs);
+}
+
+bool rectsFramebufferToBufferTimed(FrameDamageRect * rects, int count, int bpp,
+  uint8_t * dst, int dstPitch, int height,
+  const FrameBuffer * frame, int srcPitch, uint64_t * waitTimeNs)
+{
+  struct FromFramebufferData data =
+  {
+    .frame      = frame,
+    .pitch      = srcPitch,
+    .waitTimeNs = waitTimeNs
+  };
+  return rectsBufferCopy(rects, count, bpp, dst, dstPitch, height,
+    framebuffer_get_buffer(frame), srcPitch, &data, fbRowStart, NULL);
 }
 
 bool rectsFramebufferToBuffer(FrameDamageRect * rects, int count, int bpp,
   uint8_t * dst, int dstPitch, int height,
   const FrameBuffer * frame, int srcPitch)
 {
-  struct FromFramebufferData data = { .frame = frame, .pitch = srcPitch };
+  return rectsFramebufferToBufferTimed(rects, count, bpp, dst, dstPitch,
-  return rectsBufferCopy(rects, count, bpp, dst, dstPitch, height,
+      height, frame, srcPitch, NULL);
-    framebuffer_get_buffer(frame), srcPitch, &data, fbRowStart, NULL);
 }
 
 int rectsMergeOverlapping(FrameDamageRect * rects, int count)

host/platform/Windows/capture/D12/d12.c

@@ -58,6 +58,19 @@ struct D12Interface
   D12CommandGroup       copyCommand;
   D12CommandGroup       computeCommand;
 
+  struct
+  {
+    ID3D12QueryHeap ** heap;
+    ID3D12Resource  ** readback;
+    UINT64           * map;
+    UINT64             timestampFrequency;
+    UINT64             qpcFrequency;
+    UINT64             calibrationGPU;
+    UINT64             calibrationCPU;
+    bool               supported;
+  }
+  copyTiming;
+
   void        * ivshmemBase;
   ID3D12Heap ** ivshmemHeap;
   bool          indirectCopy;
@@ -176,6 +189,13 @@ static ID3D12Resource * d12_frameBufferToResource(
   unsigned size,
   void ** map);
 
+static bool d12_copyTimingInit(
+  ID3D12Device3 * device, ID3D12CommandQueue * queue);
+
+static void d12_copyTimingDeinit(void);
+static void d12_copyTimingUpdateCalibration(void);
+static bool d12_copyTimingGetStart(uint64_t * start);
+
 // implementation
 
 static const char * d12_getName(void)
@@ -183,6 +203,178 @@ static const char * d12_getName(void)
   return "D12";
 }
 
+static uint64_t d12_scaleTicks(
+  uint64_t ticks, uint64_t targetFrequency, uint64_t sourceFrequency)
+{
+  return ticks / sourceFrequency * targetFrequency +
+    ticks % sourceFrequency * targetFrequency / sourceFrequency;
+}
+
+static bool d12_copyTimingInit(
+  ID3D12Device3 * device, ID3D12CommandQueue * queue)
+{
+  D3D12_FEATURE_DATA_D3D12_OPTIONS3 options = {0};
+  HRESULT hr = ID3D12Device3_CheckFeatureSupport(device,
+    D3D12_FEATURE_D3D12_OPTIONS3, &options, sizeof(options));
+  if (FAILED(hr) || !options.CopyQueueTimestampQueriesSupported)
+    return false;
+
+  UINT64 timestampFrequency;
+  hr = ID3D12CommandQueue_GetTimestampFrequency(queue, &timestampFrequency);
+  if (FAILED(hr) || !timestampFrequency)
+    return false;
+
+  LARGE_INTEGER qpcFrequency;
+  if (!QueryPerformanceFrequency(&qpcFrequency) || !qpcFrequency.QuadPart)
+    return false;
+
+  UINT64 calibrationGPU;
+  UINT64 calibrationCPU;
+  hr = ID3D12CommandQueue_GetClockCalibration(
+    queue, &calibrationGPU, &calibrationCPU);
+  if (FAILED(hr))
+    return false;
+
+  bool result = false;
+  comRef_scopePush(2);
+
+  D3D12_QUERY_HEAP_DESC queryDesc =
+  {
+    .Type  = D3D12_QUERY_HEAP_TYPE_COPY_QUEUE_TIMESTAMP,
+    .Count = 1
+  };
+
+  comRef_defineLocal(ID3D12QueryHeap, heap);
+  hr = ID3D12Device3_CreateQueryHeap(device, &queryDesc,
+    &IID_ID3D12QueryHeap, (void **)heap);
+  if (FAILED(hr))
+    goto exit;
+
+  D3D12_HEAP_PROPERTIES heapProps =
+  {
+    .Type                 = D3D12_HEAP_TYPE_READBACK,
+    .CPUPageProperty      = D3D12_CPU_PAGE_PROPERTY_UNKNOWN,
+    .MemoryPoolPreference = D3D12_MEMORY_POOL_UNKNOWN,
+    .CreationNodeMask     = 1,
+    .VisibleNodeMask      = 1
+  };
+
+  D3D12_RESOURCE_DESC resourceDesc =
+  {
+    .Dimension          = D3D12_RESOURCE_DIMENSION_BUFFER,
+    .Width              = sizeof(UINT64),
+    .Height             = 1,
+    .DepthOrArraySize   = 1,
+    .MipLevels          = 1,
+    .Format             = DXGI_FORMAT_UNKNOWN,
+    .SampleDesc.Count   = 1,
+    .Layout             = D3D12_TEXTURE_LAYOUT_ROW_MAJOR,
+    .Flags              = D3D12_RESOURCE_FLAG_NONE
+  };
+
+  comRef_defineLocal(ID3D12Resource, readback);
+  hr = ID3D12Device3_CreateCommittedResource(
+    device,
+    &heapProps,
+    D3D12_HEAP_FLAG_NONE,
+    &resourceDesc,
+    D3D12_RESOURCE_STATE_COPY_DEST,
+    NULL,
+    &IID_ID3D12Resource,
+    (void **)readback);
+  if (FAILED(hr))
+    goto exit;
+
+  D3D12_RANGE readRange    = {0, sizeof(UINT64)};
+  void      * timestampMap = NULL;
+  hr = ID3D12Resource_Map(*readback, 0, &readRange, &timestampMap);
+  if (FAILED(hr))
+    goto exit;
+
+  comRef_toGlobal(this->copyTiming.heap    , heap    );
+  comRef_toGlobal(this->copyTiming.readback, readback);
+  this->copyTiming.map                = timestampMap;
+  this->copyTiming.timestampFrequency = timestampFrequency;
+  this->copyTiming.qpcFrequency       = qpcFrequency.QuadPart;
+  this->copyTiming.calibrationGPU     = calibrationGPU;
+  this->copyTiming.calibrationCPU     = calibrationCPU;
+  this->copyTiming.supported          = true;
+  result = true;
+
+exit:
+  comRef_scopePop();
+  return result;
+}
+
+static void d12_copyTimingDeinit(void)
+{
+  if (this->copyTiming.map)
+  {
+    const D3D12_RANGE writeRange = {0, 0};
+    ID3D12Resource_Unmap(
+      *this->copyTiming.readback, 0, &writeRange);
+  }
+
+  memset(&this->copyTiming, 0, sizeof(this->copyTiming));
+}
+
+static void d12_copyTimingUpdateCalibration(void)
+{
+  if (!this->copyTiming.supported)
+    return;
+
+  LARGE_INTEGER now;
+  if (QueryPerformanceCounter(&now) &&
+      (UINT64)now.QuadPart >= this->copyTiming.calibrationCPU &&
+      (UINT64)now.QuadPart - this->copyTiming.calibrationCPU <
+        this->copyTiming.qpcFrequency)
+    return;
+
+  UINT64 calibrationGPU;
+  UINT64 calibrationCPU;
+  if (SUCCEEDED(ID3D12CommandQueue_GetClockCalibration(
+        *this->copyQueue, &calibrationGPU, &calibrationCPU)))
+  {
+    this->copyTiming.calibrationGPU = calibrationGPU;
+    this->copyTiming.calibrationCPU = calibrationCPU;
+  }
+}
+
+static bool d12_copyTimingGetStart(uint64_t * start)
+{
+  if (!this->copyTiming.supported)
+    return false;
+
+  const UINT64 gpuStart = *this->copyTiming.map;
+  UINT64 cpuStart;
+  if (gpuStart < this->copyTiming.calibrationGPU)
+  {
+    const UINT64 delta = d12_scaleTicks(
+      this->copyTiming.calibrationGPU - gpuStart,
+      this->copyTiming.qpcFrequency,
+      this->copyTiming.timestampFrequency);
+    if (delta > this->copyTiming.calibrationCPU)
+      return false;
+
+    cpuStart = this->copyTiming.calibrationCPU - delta;
+  }
+  else
+  {
+    const UINT64 delta = d12_scaleTicks(
+      gpuStart - this->copyTiming.calibrationGPU,
+      this->copyTiming.qpcFrequency,
+      this->copyTiming.timestampFrequency);
+    if (UINT64_MAX - this->copyTiming.calibrationCPU < delta)
+      return false;
+
+    cpuStart = this->copyTiming.calibrationCPU + delta;
+  }
+
+  *start = d12_scaleTicks(
+    cpuStart, 1000000000ULL, this->copyTiming.qpcFrequency);
+  return true;
+}
+
 static void d12_initOptions(void)
 {
   struct Option options[] =
@@ -418,6 +610,9 @@ static bool d12_init(void * ivshmemBase, unsigned * alignSize)
     *device, D3D12_COMMAND_LIST_TYPE_COMPUTE, &this->computeCommand, L"Compute"))
     goto exit;
 
+  if (!d12_copyTimingInit(*device, *copyQueue))
+    DEBUG_WARN("Copy queue GPU timing is unavailable");
+
   comRef_defineLocal(ID3D12Heap, ivshmemHeap);
   if (!this->indirectCopy)
   {
@@ -498,6 +693,7 @@ exit:
   if (!result)
   {
     DEBUG_TRACE("Init failed");
+    d12_copyTimingDeinit();
     D12Effect * effect;
     vector_forEach(effect, &this->effects)
       d12_effectFree(&effect);
@@ -542,6 +738,8 @@ static bool d12_deinit(void)
   d12_commandGroupFree(&this->copyCommand   );
   d12_commandGroupFree(&this->computeCommand);
 
+  d12_copyTimingDeinit();
+
   DEBUG_TRACE("comRef_freeScope");
   IDXGIFactory2 * factory = *this->factory;
   IDXGIFactory2_AddRef(factory);
@@ -765,6 +963,9 @@ static CaptureResult d12_waitFrame(unsigned frameBufferIndex,
     }
   }
 
+  // Refresh before getFrame so calibration work does not perturb Post/Copy.
+  d12_copyTimingUpdateCalibration();
+
   result = CAPTURE_RESULT_OK;
 
 exit:
@@ -778,6 +979,7 @@ static CaptureResult d12_getFrame(
   const size_t   maxFrameSize,
   CaptureFrame * captureFrame)
 {
+  const uint64_t postProcessStart = nanotime();
   CaptureResult result = CAPTURE_RESULT_ERROR;
   comRef_scopePush(3);
 
@@ -812,7 +1014,6 @@ static CaptureResult d12_getFrame(
   if (result != CAPTURE_RESULT_OK)
     goto exit;
 
-  const uint64_t postProcessStart = nanotime();
   ID3D12Resource * next = *src;
   D12Effect * effect;
   vector_forEach(effect, &this->effects)
@@ -855,6 +1056,15 @@ static CaptureResult d12_getFrame(
     }
   };
 
+  // Queue waits run before this command list. This timestamp is therefore the
+  // boundary between source/effect readiness and the framebuffer copy.
+  if (this->copyTiming.supported)
+    ID3D12GraphicsCommandList_EndQuery(
+      *this->copyCommand.gfxList,
+      *this->copyTiming.heap,
+      D3D12_QUERY_TYPE_TIMESTAMP,
+      0);
+
   // if full frame damage
   if (desc.nbDirtyRects == 0)
   {
@@ -928,6 +1138,16 @@ static CaptureResult d12_getFrame(
     this->nbDirtyRects = desc.nbDirtyRects;
   }
 
+  if (this->copyTiming.supported)
+    ID3D12GraphicsCommandList_ResolveQueryData(
+      *this->copyCommand.gfxList,
+      *this->copyTiming.heap,
+      D3D12_QUERY_TYPE_TIMESTAMP,
+      0,
+      1,
+      *this->copyTiming.readback,
+      0);
+
   // execute the compute commands
   if (this->effectsActive)
   {
@@ -939,8 +1159,8 @@ static CaptureResult d12_getFrame(
     ID3D12CommandQueue_Wait(*this->copyQueue,
       *this->computeCommand.fence, this->computeCommand.fenceValue);
   }
-  captureFrame->postProcessTime = this->effectsActive ?
+
-    nanotime() - postProcessStart : 0;
+  const uint64_t fallbackCopyStart = nanotime();
 
   // execute the copy commands
   DEBUG_TRACE("Execute copy commands");
@@ -950,6 +1170,16 @@ static CaptureResult d12_getFrame(
   DEBUG_TRACE("Fence wait");
   d12_commandGroupWait(&this->copyCommand);
 
+  const uint64_t copyComplete = nanotime();
+  uint64_t copyStart;
+  if (!d12_copyTimingGetStart(&copyStart) ||
+      copyStart < postProcessStart || copyStart > copyComplete)
+    copyStart = fallbackCopyStart;
+
+  // The caller derives Copy from its getFrame wall time minus Post, preserving
+  // the exact producer total while moving the real queue wait to Post.
+  captureFrame->postProcessTime = copyStart - postProcessStart;
+
   if (this->indirectCopy)
   {
     if (rectCount == 0)

idd/LGIdd/CD3D12CommandQueue.cpp

@@ -21,7 +21,125 @@
 #include "CD3D12CommandQueue.h"
 #include "CDebug.h"
 
-bool CD3D12CommandQueue::Init(ID3D12Device3 * device, D3D12_COMMAND_LIST_TYPE type, const WCHAR* name, CallbackMode callbackMode)
+static uint64_t ScaleTicks(uint64_t ticks, uint64_t targetFrequency,
+  uint64_t sourceFrequency)
+{
+  return ticks / sourceFrequency * targetFrequency +
+    ticks % sourceFrequency * targetFrequency / sourceFrequency;
+}
+
+static uint64_t TicksToNanoseconds(uint64_t ticks, uint64_t frequency)
+{
+  return ScaleTicks(ticks, 1000000000ULL, frequency);
+}
+
+bool CD3D12CommandQueue::InitTiming(ID3D12Device3 * device,
+  D3D12_COMMAND_LIST_TYPE type)
+{
+  if (type != D3D12_COMMAND_LIST_TYPE_COPY)
+    return false;
+
+  D3D12_FEATURE_DATA_D3D12_OPTIONS3 options = {};
+  HRESULT hr = device->CheckFeatureSupport(
+    D3D12_FEATURE_D3D12_OPTIONS3, &options, sizeof(options));
+  if (FAILED(hr) || !options.CopyQueueTimestampQueriesSupported)
+    return false;
+
+  hr = m_queue->GetTimestampFrequency(&m_timestampFrequency);
+  if (FAILED(hr) || !m_timestampFrequency)
+    return false;
+
+  LARGE_INTEGER qpcFrequency;
+  if (!QueryPerformanceFrequency(&qpcFrequency))
+    return false;
+  m_qpcFrequency = (UINT64)qpcFrequency.QuadPart;
+
+  D3D12_QUERY_HEAP_DESC queryDesc = {};
+  queryDesc.Type  = D3D12_QUERY_HEAP_TYPE_COPY_QUEUE_TIMESTAMP;
+  queryDesc.Count = 1;
+
+  hr = device->CreateQueryHeap(&queryDesc, IID_PPV_ARGS(&m_timestampHeap));
+  if (FAILED(hr))
+    return false;
+
+  D3D12_HEAP_PROPERTIES heapProps = {};
+  heapProps.Type                 = D3D12_HEAP_TYPE_READBACK;
+  heapProps.CPUPageProperty      = D3D12_CPU_PAGE_PROPERTY_UNKNOWN;
+  heapProps.MemoryPoolPreference = D3D12_MEMORY_POOL_UNKNOWN;
+  heapProps.CreationNodeMask     = 1;
+  heapProps.VisibleNodeMask      = 1;
+
+  D3D12_RESOURCE_DESC resourceDesc = {};
+  resourceDesc.Dimension          = D3D12_RESOURCE_DIMENSION_BUFFER;
+  resourceDesc.Width              = sizeof(UINT64);
+  resourceDesc.Height             = 1;
+  resourceDesc.DepthOrArraySize   = 1;
+  resourceDesc.MipLevels          = 1;
+  resourceDesc.SampleDesc.Count   = 1;
+  resourceDesc.Layout             = D3D12_TEXTURE_LAYOUT_ROW_MAJOR;
+
+  hr = device->CreateCommittedResource(
+    &heapProps,
+    D3D12_HEAP_FLAG_NONE,
+    &resourceDesc,
+    D3D12_RESOURCE_STATE_COPY_DEST,
+    NULL,
+    IID_PPV_ARGS(&m_timestampReadback));
+  if (FAILED(hr))
+  {
+    m_timestampHeap.Reset();
+    return false;
+  }
+
+  D3D12_RANGE readRange    = { 0, sizeof(UINT64) };
+  void      * timestampMap = nullptr;
+  hr = m_timestampReadback->Map(0, &readRange, &timestampMap);
+  if (FAILED(hr))
+  {
+    m_timestampReadback.Reset();
+    m_timestampHeap.Reset();
+    return false;
+  }
+  m_timestampMap = static_cast<UINT64 *>(timestampMap);
+
+  hr = m_queue->GetClockCalibration(&m_calibrationGPU, &m_calibrationCPU);
+  if (FAILED(hr))
+  {
+    D3D12_RANGE writeRange = { 0, 0 };
+    m_timestampReadback->Unmap(0, &writeRange);
+    m_timestampMap = nullptr;
+    m_timestampReadback.Reset();
+    m_timestampHeap.Reset();
+    return false;
+  }
+
+  m_timingSupported = true;
+  return true;
+}
+
+void CD3D12CommandQueue::UpdateClockCalibration()
+{
+  if (!m_timingSupported)
+    return;
+
+  LARGE_INTEGER now;
+  if (QueryPerformanceCounter(&now) &&
+      (UINT64)now.QuadPart >= m_calibrationCPU &&
+      (UINT64)now.QuadPart - m_calibrationCPU < m_qpcFrequency)
+    return;
+
+  UINT64 gpu;
+  UINT64 cpu;
+  if (SUCCEEDED(m_queue->GetClockCalibration(&gpu, &cpu)))
+  {
+    m_calibrationGPU = gpu;
+    m_calibrationCPU = cpu;
+  }
+}
+
+bool CD3D12CommandQueue::Init(ID3D12Device3 * device,
+  D3D12_COMMAND_LIST_TYPE type, const WCHAR * name,
+  CallbackMode callbackMode)
 {
   HRESULT hr;
   D3D12_COMMAND_QUEUE_DESC queueDesc = {};
@@ -79,7 +197,7 @@ bool CD3D12CommandQueue::Init(ID3D12Device3 * device, D3D12_COMMAND_LIST_TYPE ty
     ULONG flags = (callbackMode == FAST) ?
       WT_EXECUTEINWAITTHREAD : WT_EXECUTEINPERSISTENTTHREAD;
 
-    RegisterWaitForSingleObject(
+    if (!RegisterWaitForSingleObject(
       &m_waitHandle,
       m_event.Get(),
       [](PVOID param, BOOLEAN timeout){
@@ -90,8 +208,18 @@ bool CD3D12CommandQueue::Init(ID3D12Device3 * device, D3D12_COMMAND_LIST_TYPE ty
       },
       this,
       INFINITE,
-      flags);
+      flags))
+    {
+      DEBUG_ERROR_HR(GetLastError(),
+        "Failed to register the completion wait (%ls)", name);
+      m_waitHandle = INVALID_HANDLE_VALUE;
+      return false;
     }
+  }
+
+  if (callbackMode != DISABLED &&
+      type == D3D12_COMMAND_LIST_TYPE_COPY && !InitTiming(device, type))
+    DEBUG_WARN("GPU timing is unavailable for CommandQueue(%ls)", name);
 
   m_name       = name;
   m_fenceValue = 0;
@@ -103,15 +231,25 @@ void CD3D12CommandQueue::DeInit()
 {
   if (m_waitHandle != INVALID_HANDLE_VALUE)
   {
+    // Queue owners drain callbacks before destruction. Keep this unregister
+    // non-blocking so a removed or hung device cannot stall teardown.
     UnregisterWait(m_waitHandle);
     m_waitHandle = INVALID_HANDLE_VALUE;
   }
+
+  if (m_timestampMap)
+  {
+    D3D12_RANGE writeRange = { 0, 0 };
+    m_timestampReadback->Unmap(0, &writeRange);
+    m_timestampMap = nullptr;
+  }
 }
 
 bool CD3D12CommandQueue::Execute()
 {
   m_needsReset       = true;
   m_completionResult = true;
+  m_pending          = m_waitHandle != INVALID_HANDLE_VALUE;
 
   HRESULT hr = m_gfxList->Close();
   if (FAILED(hr))
@@ -137,11 +275,63 @@ bool CD3D12CommandQueue::Execute()
     return false;
   }
 
-  m_pending = true;
   m_queue->Signal(m_fence.Get(), m_fenceValue);
   return true;
 }
 
+bool CD3D12CommandQueue::BeginTiming()
+{
+  m_timingActive = m_timingSupported;
+  if (!m_timingActive)
+    return false;
+
+  // Refresh after an idle period; active queues are calibrated by their
+  // completion path without adding work to frame submission.
+  UpdateClockCalibration();
+  m_gfxList->EndQuery(
+    m_timestampHeap.Get(), D3D12_QUERY_TYPE_TIMESTAMP, 0);
+  return true;
+}
+
+void CD3D12CommandQueue::EndTiming()
+{
+  if (!m_timingActive)
+    return;
+
+  m_gfxList->ResolveQueryData(
+    m_timestampHeap.Get(), D3D12_QUERY_TYPE_TIMESTAMP,
+    0, 1, m_timestampReadback.Get(), 0);
+}
+
+bool CD3D12CommandQueue::GetGPUStartTime(uint64_t& start)
+{
+  if (!m_timingActive)
+    return false;
+
+  const UINT64 gpuStart = m_timestampMap[0];
+
+  UINT64 cpuStart;
+  if (gpuStart < m_calibrationGPU)
+  {
+    const UINT64 delta = ScaleTicks(
+      m_calibrationGPU - gpuStart, m_qpcFrequency, m_timestampFrequency);
+    if (delta > m_calibrationCPU)
+      return false;
+    cpuStart = m_calibrationCPU - delta;
+  }
+  else
+  {
+    const UINT64 delta = ScaleTicks(
+      gpuStart - m_calibrationGPU, m_qpcFrequency, m_timestampFrequency);
+    if (UINT64_MAX - m_calibrationCPU < delta)
+      return false;
+    cpuStart = m_calibrationCPU + delta;
+  }
+
+  start = TicksToNanoseconds(cpuStart, m_qpcFrequency);
+  return true;
+}
+
 #if 0
 void CD3D12CommandQueue::Wait()
 {
@@ -159,6 +349,7 @@ void CD3D12CommandQueue::Wait()
 
 bool CD3D12CommandQueue::Reset()
 {
+  m_timingActive = false;
   if (!m_needsReset)
     return true;
 

idd/LGIdd/CD3D12CommandQueue.h

@@ -25,6 +25,7 @@
 #include <wrl.h>
 #include <d3d12.h>
 #include <atomic>
+#include <stdint.h>
 
 using namespace Microsoft::WRL;
 using namespace Microsoft::WRL::Wrappers;
@@ -41,6 +42,16 @@ class CD3D12CommandQueue
     ComPtr<ID3D12CommandList        > m_cmdList;
     ComPtr<ID3D12Fence              > m_fence;
 
+    ComPtr<ID3D12QueryHeap>  m_timestampHeap;
+    ComPtr<ID3D12Resource >  m_timestampReadback;
+    UINT64                 * m_timestampMap       = nullptr;
+    UINT64                   m_timestampFrequency = 0;
+    UINT64                   m_calibrationGPU     = 0;
+    UINT64                   m_calibrationCPU     = 0;
+    UINT64                   m_qpcFrequency       = 0;
+    bool                     m_timingSupported    = false;
+    bool                     m_timingActive       = false;
+
     std::atomic<bool>         m_pending    = false;
     HandleT<HANDLENullTraits> m_event;
     HANDLE                    m_waitHandle = INVALID_HANDLE_VALUE;
@@ -54,6 +65,9 @@ class CD3D12CommandQueue
     void               * m_completionParams[2];
     bool                 m_completionResult = true;
 
+    bool InitTiming(ID3D12Device3 * device, D3D12_COMMAND_LIST_TYPE type);
+    void UpdateClockCalibration();
+
     void OnCompletion()
     {
       if (m_completionCallback)
@@ -62,6 +76,7 @@ class CD3D12CommandQueue
           m_completionResult,
           m_completionParams[0],
           m_completionParams[1]);
+      UpdateClockCalibration();
       m_pending = false;
     }
 
@@ -90,6 +105,12 @@ class CD3D12CommandQueue
     bool Reset();
     bool Execute();
 
+    bool BeginTiming();
+    void EndTiming();
+
+    // Return the command-list start in QueryPerformanceCounter-domain ns.
+    bool GetGPUStartTime(uint64_t& start);
+
     //void Wait();
     bool   IsReady () const { return !m_pending   ; }
     HANDLE GetEvent() const { return m_event.Get(); }

idd/LGIdd/CFrameBufferResource.h

@@ -39,7 +39,7 @@ class CFrameBufferResource
     size_t                 m_size             = 0;
     size_t                 m_frameSize        = 0;
     uint64_t               m_captureTime      = 0;
-    uint64_t               m_postProcessTime = 0;
+    uint64_t               m_postProcessStart = 0;
     uint64_t               m_copyStart        = 0;
     ComPtr<ID3D12Resource> m_res;
     void                 * m_map              = nullptr;
@@ -55,15 +55,15 @@ class CFrameBufferResource
     size_t    GetFrameSize()  { return m_frameSize;  }
     void *    GetMap()        { return m_map;        }
 
-    void SetTiming(uint64_t captureTime, uint64_t postProcessTime,
+    void SetTiming(uint64_t captureTime, uint64_t postProcessStart,
       uint64_t copyStart)
     {
       m_captureTime      = captureTime;
-      m_postProcessTime = postProcessTime;
+      m_postProcessStart = postProcessStart;
       m_copyStart        = copyStart;
     }
     uint64_t GetCaptureTime     () const { return m_captureTime;      }
-    uint64_t GetPostProcessTime() const { return m_postProcessTime; }
+    uint64_t GetPostProcessStart() const { return m_postProcessStart; }
     uint64_t GetCopyStart       () const { return m_copyStart;        }
 
     ComPtr<ID3D12Resource> Get() { return m_res; }

idd/LGIdd/CSwapChainProcessor.cpp

@@ -289,7 +289,7 @@ bool CSwapChainProcessor::SwapChainThreadCore()
       {
         lastFrameNumber = frameNumber;
         if (!SwapChainNewFrame(surface, dirtyRectCount, moveRegionCount,
-              colorSpace, sdrWhiteLevel, Nanotime() - captureStart))
+              colorSpace, sdrWhiteLevel, captureStart))
           DEBUG_WARN("Failed to submit frame");
       }
 
@@ -315,29 +315,35 @@ bool CSwapChainProcessor::SwapChainThreadCore()
 void CSwapChainProcessor::CompletionFunction(
   CD3D12CommandQueue * queue, bool result, void * param1, void * param2)
 {
-  UNREFERENCED_PARAMETER(queue);
-
   auto sc    = (CSwapChainProcessor   *)param1;
   auto fbRes = (CFrameBufferResource *)param2;
 
-  // fail gracefully
+  uint64_t copyStart    = fbRes->GetCopyStart();
-  if (!result)
+  uint64_t gpuCopyStart = 0;
-  {
-    sc->m_devContext->SetFrameTiming(fbRes->GetFrameIndex(),
-      fbRes->GetCaptureTime(), fbRes->GetPostProcessTime(),
-      Nanotime() - fbRes->GetCopyStart());
-    sc->m_devContext->FinalizeFrameBuffer(fbRes->GetFrameIndex());
-    return;
-  }
 
-  if (sc->m_dx12Device->IsIndirectCopy())
+  if (result && sc->m_dx12Device->IsIndirectCopy())
     sc->m_devContext->WriteFrameBuffer(
       fbRes->GetFrameIndex(), fbRes->GetMap(), 0, fbRes->GetFrameSize(), false);
 
-  const uint64_t copyTime = Nanotime() - fbRes->GetCopyStart();
+  // Queue waits execute before this timestamp. Use it as the boundary so the
-  sc->m_devContext->SetFrameTiming(fbRes->GetFrameIndex(),
+  // source fence and effects are charged to Post, while Copy retains the full
-    fbRes->GetCaptureTime(), fbRes->GetPostProcessTime(), copyTime);
+  // time through buffer readiness and any indirect memcpy.
+  const bool gpuTimingValid = result && queue->GetGPUStartTime(gpuCopyStart);
+
+  // Publish readiness before sampling the endpoint. Timing has its own valid
+  // flag and is published immediately afterwards.
   sc->m_devContext->FinalizeFrameBuffer(fbRes->GetFrameIndex());
+  const uint64_t copyEnd = Nanotime();
+  if (gpuTimingValid &&
+      gpuCopyStart >= fbRes->GetPostProcessStart() &&
+      gpuCopyStart <= copyEnd)
+    copyStart = gpuCopyStart;
+
+  const uint64_t postProcessTime = copyStart -
+    fbRes->GetPostProcessStart();
+  const uint64_t copyTime        = copyEnd - copyStart;
+  sc->m_devContext->SetFrameTiming(fbRes->GetFrameIndex(),
+    fbRes->GetCaptureTime(), postProcessTime, copyTime);
 }
 
 
@@ -515,8 +521,11 @@ bool CSwapChainProcessor::GetContentHDRMetadata(D12FrameFormat& format) const
 bool CSwapChainProcessor::SwapChainNewFrame(ComPtr<IDXGIResource> acquiredBuffer,
   unsigned dirtyRectCount, unsigned moveRegionCount,
   DXGI_COLOR_SPACE_TYPE colorSpace, UINT sdrWhiteLevel,
-  uint64_t captureTime)
+  uint64_t captureStart)
 {
+  const uint64_t postProcessStart = Nanotime();
+  const uint64_t captureTime      = postProcessStart - captureStart;
+
   // Preserve the fast drop path: never hold an IddCx frame while waiting for
   // a slow or disconnected client. We have not read its rectangles, so force
   // the next published frame to invalidate the entire image.
@@ -526,8 +535,6 @@ bool CSwapChainProcessor::SwapChainNewFrame(ComPtr<IDXGIResource> acquiredBuffer
     return true;
   }
 
-  const uint64_t postProcessStart = Nanotime();
-
   ComPtr<ID3D11Texture2D> texture;
   HRESULT hr = acquiredBuffer.As(&texture);
   if (FAILED(hr))
@@ -775,7 +782,7 @@ bool CSwapChainProcessor::SwapChainNewFrame(ComPtr<IDXGIResource> acquiredBuffer
   }
 
   const uint64_t copyStart = Nanotime();
-  fbRes->SetTiming(captureTime, copyStart - postProcessStart, copyStart);
+  fbRes->SetTiming(captureTime, postProcessStart, copyStart);
 
   copyQueue->SetCompletionCallback(&CompletionFunction, this, fbRes);
 
@@ -789,6 +796,9 @@ bool CSwapChainProcessor::SwapChainNewFrame(ComPtr<IDXGIResource> acquiredBuffer
   dstLoc.Type            = D3D12_TEXTURE_COPY_TYPE_PLACED_FOOTPRINT;
   dstLoc.PlacedFootprint = layout;
 
+  // The source/compute waits were inserted directly on the queue above. The
+  // command-list timestamp therefore marks the first actual copy operation.
+  copyQueue->BeginTiming();
   if (IsFullDamage(currentDirtyRects, nbDirtyRects, dstFormat.desc) ||
       nbDirtyRects > KVMFR_MAX_DAMAGE_RECTS || m_nbDirtyRects == 0)
   {
@@ -812,6 +822,7 @@ bool CSwapChainProcessor::SwapChainNewFrame(ComPtr<IDXGIResource> acquiredBuffer
     for (const RECT * rect = currentDirtyRects; rect < currentDirtyRects + nbDirtyRects; ++rect)
       CopyDirtyRect(copyQueue->GetGfxList(), &dstLoc, &srcLoc, *rect);
   }
+  copyQueue->EndTiming();
 
   if (!copyQueue->Execute())
   {

idd/LGIdd/CSwapChainProcessor.h

@@ -104,7 +104,7 @@ private:
   bool SwapChainNewFrame(ComPtr<IDXGIResource> acquiredBuffer,
     unsigned dirtyRectCount, unsigned moveRegionCount,
     DXGI_COLOR_SPACE_TYPE colorSpace, UINT sdrWhiteLevel,
-    uint64_t captureTime);
+    uint64_t captureStart);
 
 public:
   CSwapChainProcessor(CIndirectMonitorContext * monitorContext, UINT64 assignmentGeneration,