diff --git a/client/renderers/EGL/desktop.c b/client/renderers/EGL/desktop.c
index ffaf36a3..44a358b8 100644
--- a/client/renderers/EGL/desktop.c
+++ b/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;
diff --git a/client/renderers/EGL/desktop.h b/client/renderers/EGL/desktop.h
index 1d3596ec..2142ff65 100644
--- a/client/renderers/EGL/desktop.h
+++ b/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,
diff --git a/client/renderers/EGL/egl.c b/client/renderers/EGL/egl.c
index 4d25efee..167bff18 100644
--- a/client/renderers/EGL/egl.c
+++ b/client/renderers/EGL/egl.c
@@ -748,14 +748,21 @@ static bool egl_onFrame(LG_Renderer * renderer, const FrameBuffer * frame, int d
   struct Inst * this = UPCAST(struct Inst, renderer);
   egl_stateCheckShared();
 
-  const uint64_t start = nanotime();
+  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;
diff --git a/client/renderers/EGL/texture.c b/client/renderers/EGL/texture.c
index f7f9dd23..d0a093d4 100644
--- a/client/renderers/EGL/texture.c
+++ b/client/renderers/EGL/texture.c
@@ -161,27 +161,28 @@ 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 =
   {
-    .type      = EGL_TEXTYPE_FRAMEBUFFER,
-    .x         = 0,
-    .y         = 0,
-    .width     = this->format.width,
-    .height    = this->format.height,
-    .pitch     = this->format.pitch,
-    .stride    = this->format.stride,
-    .frame     = frame,
-    .rects     = damageRects,
-    .rectCount = damageRectsCount,
+    .type       = EGL_TEXTYPE_FRAMEBUFFER,
+    .x          = 0,
+    .y          = 0,
+    .width      = this->format.width,
+    .height     = this->format.height,
+    .pitch      = this->format.pitch,
+    .stride     = this->format.stride,
+    .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);
diff --git a/client/renderers/EGL/texture.h b/client/renderers/EGL/texture.h
index abe36cb9..09e7a055 100644
--- a/client/renderers/EGL/texture.h
+++ b/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);
 
diff --git a/client/renderers/EGL/texture_framebuffer.c b/client/renderers/EGL/texture_framebuffer.c
index eb434b44..c99737f5 100644
--- a/client/renderers/EGL/texture_framebuffer.c
+++ b/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
       );
     }
   }
diff --git a/client/transports/LGMP/lgmp.c b/client/transports/LGMP/lgmp.c
index 434d3dd9..579fc813 100644
--- a/client/transports/LGMP/lgmp.c
+++ b/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;
diff --git a/common/include/common/framebuffer.h b/common/include/common/framebuffer.h
index e4317d9d..9f3844f6 100644
--- a/common/include/common/framebuffer.h
+++ b/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
  */
diff --git a/common/include/common/rects.h b/common/include/common/rects.h
index e409ee54..1d0d5ae8 100644
--- a/common/include/common/rects.h
+++ b/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);
 
diff --git a/common/src/framebuffer.c b/common/src/framebuffer.c
index ea5c4f23..607fbc71 100644
--- a/common/src/framebuffer.c
+++ b/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,
-    size_t size)
+bool framebuffer_wait(const FrameBuffer * frame, 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,
-    size_t dstpitch, size_t height, size_t width, size_t bpp, size_t pitch)
+bool framebuffer_read_linear(const FrameBuffer * frame, void * restrict dst,
+    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)
 {
diff --git a/common/src/rects.c b/common/src/rects.c
index 780c2cf3..a10cf2e2 100644
--- a/common/src/rects.c
+++ b/common/src/rects.c
@@ -220,22 +220,37 @@ void rectsBufferToFramebuffer(FrameDamageRect * rects, int count, int bpp,
 struct FromFramebufferData
 {
   const FrameBuffer * frame;
-  int pitch;
+  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 rectsBufferCopy(rects, count, bpp, dst, dstPitch, height,
-    framebuffer_get_buffer(frame), srcPitch, &data, fbRowStart, NULL);
+  return rectsFramebufferToBufferTimed(rects, count, bpp, dst, dstPitch,
+      height, frame, srcPitch, NULL);
 }
 
 int rectsMergeOverlapping(FrameDamageRect * rects, int count)
diff --git a/host/platform/Windows/capture/D12/d12.c b/host/platform/Windows/capture/D12/d12.c
index 33887280..7d751aee 100644
--- a/host/platform/Windows/capture/D12/d12.c
+++ b/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)
diff --git a/idd/LGIdd/CD3D12CommandQueue.cpp b/idd/LGIdd/CD3D12CommandQueue.cpp
index ab18e3f8..39ca3344 100644
--- a/idd/LGIdd/CD3D12CommandQueue.cpp
+++ b/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,9 +208,19 @@ 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;
   DEBUG_INFO("Created CD3D12CommandQueue(%ls)", name);
@@ -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;
 
diff --git a/idd/LGIdd/CD3D12CommandQueue.h b/idd/LGIdd/CD3D12CommandQueue.h
index e9236e4f..81052c74 100644
--- a/idd/LGIdd/CD3D12CommandQueue.h
+++ b/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,11 +42,21 @@ class CD3D12CommandQueue
     ComPtr<ID3D12CommandList        > m_cmdList;
     ComPtr<ID3D12Fence              > m_fence;
 
-    std::atomic<bool> m_pending = false;
+    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;
-    UINT64 m_fenceValue = 0;
-    bool m_needsReset = false;
+    HANDLE                    m_waitHandle = INVALID_HANDLE_VALUE;
+    UINT64                    m_fenceValue = 0;
+    bool                      m_needsReset = false;
 
     typedef void (*CompletionFunction)(CD3D12CommandQueue * queue,
       bool result, void * param1, void * param2);
@@ -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(); }
diff --git a/idd/LGIdd/CFrameBufferResource.h b/idd/LGIdd/CFrameBufferResource.h
index 118e57f0..2126f5c4 100644
--- a/idd/LGIdd/CFrameBufferResource.h
+++ b/idd/LGIdd/CFrameBufferResource.h
@@ -33,16 +33,16 @@ using namespace Microsoft::WRL;
 class CFrameBufferResource
 {
   private:
-    bool                   m_valid           = false;
-    unsigned               m_frameIndex      = 0;
-    uint8_t              * m_base            = nullptr;
-    size_t                 m_size            = 0;
-    size_t                 m_frameSize       = 0;
-    uint64_t               m_captureTime     = 0;
-    uint64_t               m_postProcessTime = 0;
-    uint64_t               m_copyStart       = 0;
+    bool                   m_valid            = false;
+    unsigned               m_frameIndex       = 0;
+    uint8_t              * m_base             = nullptr;
+    size_t                 m_size             = 0;
+    size_t                 m_frameSize        = 0;
+    uint64_t               m_captureTime      = 0;
+    uint64_t               m_postProcessStart = 0;
+    uint64_t               m_copyStart        = 0;
     ComPtr<ID3D12Resource> m_res;
-    void                 * m_map             = nullptr;
+    void                 * m_map              = nullptr;
 
   public:
     bool Init(CSwapChainProcessor * swapChain, unsigned frameIndex, uint8_t * base, size_t size);
@@ -55,16 +55,16 @@ 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_copyStart       = copyStart;
+      m_captureTime      = captureTime;
+      m_postProcessStart = postProcessStart;
+      m_copyStart        = copyStart;
     }
-    uint64_t GetCaptureTime    () const { return m_captureTime;     }
-    uint64_t GetPostProcessTime() const { return m_postProcessTime; }
-    uint64_t GetCopyStart      () const { return m_copyStart;       }
+    uint64_t GetCaptureTime     () const { return m_captureTime;      }
+    uint64_t GetPostProcessStart() const { return m_postProcessStart; }
+    uint64_t GetCopyStart       () const { return m_copyStart;        }
 
     ComPtr<ID3D12Resource> Get() { return m_res; }
 };
diff --git a/idd/LGIdd/CSwapChainProcessor.cpp b/idd/LGIdd/CSwapChainProcessor.cpp
index d50201ec..cd237fb5 100644
--- a/idd/LGIdd/CSwapChainProcessor.cpp
+++ b/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;
 
-  auto sc    = (CSwapChainProcessor *)param1;
-  auto fbRes = (CFrameBufferResource*)param2;
+  uint64_t copyStart    = fbRes->GetCopyStart();
+  uint64_t gpuCopyStart = 0;
 
-  // fail gracefully
-  if (!result)
-  {
-    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();
-  sc->m_devContext->SetFrameTiming(fbRes->GetFrameIndex(),
-    fbRes->GetCaptureTime(), fbRes->GetPostProcessTime(), copyTime);
+  // Queue waits execute before this timestamp. Use it as the boundary so the
+  // source fence and effects are charged to Post, while Copy retains the full
+  // 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())
   {
diff --git a/idd/LGIdd/CSwapChainProcessor.h b/idd/LGIdd/CSwapChainProcessor.h
index bf76f6cc..5974bbba 100644
--- a/idd/LGIdd/CSwapChainProcessor.h
+++ b/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,