Metal: compute pass with timestampWrites but no dispatch resolves both timestamps to zero

[hughryan]

Environment

• wgpu-native v29.0.1.1 (wgpu-macos-aarch64-release prebuilt)
• macOS 26.5.2 (build 25F84), Apple M2, Metal backend
• Timestamp period reported: 1 ns/tick

Summary

A compute pass created with timestampWrites (both beginningOfPassWriteIndex and endOfPassWriteIndex set) but containing no dispatch resolves both of its timestamps to zero. The submission completes normally; no error or warning is reported.

An otherwise identical pass with one trivial workgroup dispatch, on the same device and query set, resolves to nonzero, nondecreasing timestamps — so the query set, the resolve, and the readback path are all working.

Why it matters

The zeros are indistinguishable from a legitimate timestamp value at the API level, so a caller timing a pass that happens to contain no work silently measures a garbage span rather than getting an error. If the underlying constraint is that Metal never schedules an empty encoder (so there is no moment to sample), a validation warning or a documented guarantee would make this diagnosable.

Expected vs actual

Expected: an empty pass is legal and its begin/end timestamps are still well defined (two nonzero, nondecreasing values).
Actual: both queries resolve to 0.

Possibly related to the macOS 26 counter-sampling changes discussed in gfx-rs/wgpu#9414, but I have not verified whether this reproduces on earlier macOS versions.

Reproducer

Self-contained C program below; build against a stock wgpu-native install:

cc -std=c11 -O2 main.c -o repro -I"$WGPU/include" \
   "$WGPU/lib/libwgpu_native.dylib" -rpath "$WGPU/lib"
./repro    # prints pass A (empty) = 0/0, pass B (one dispatch) = nonzero; exit 1

main.c

/*
 * wgpu-native v29.0.1.1 — a compute pass with timestampWrites but no dispatch
 * resolves both of its timestamps to zero (Metal backend).
 *
 * Environment this was captured on:
 *   macOS 26.5.2 (build 25F84), Apple M2, Metal backend,
 *   wgpu-native v29.0.1.1 (wgpu-macos-aarch64-release prebuilt).
 *
 * Sequence, twice, on the same device and query set:
 *   - begin a compute pass with timestampWrites
 *     (beginningOfPassWriteIndex, endOfPassWriteIndex both set),
 *   - pass A: end it immediately, with no dispatch,
 *     pass B: set a trivial pipeline and dispatch one workgroup, then end it,
 *   - resolveQuerySet, copy to a mappable buffer, submit, wait, map, read.
 *
 * Expected: both passes yield two nonzero, nondecreasing timestamps. An empty
 *           pass is legal and its begin/end timestamps are still well defined.
 * Actual:   pass A resolves both queries to 0 (no error, no warning, the
 *           submission completes normally). Pass B — identical apart from the
 *           one dispatch — resolves to nonzero timestamps, so the query set,
 *           the resolve and the readback are all working.
 *
 * The zeros are indistinguishable from a legitimately-zero timestamp at the
 * API level, so a caller timing a pass that happens to have no work silently
 * measures a garbage span rather than getting an error.
 *
 * Build (any stock wgpu-native install; WGPU is its root directory):
 *   cc -std=c11 -O2 main.c -o repro -I"$WGPU/include" \
 *      "$WGPU/lib/libwgpu_native.dylib" -rpath "$WGPU/lib"
 */
#include <stdint.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>

#include <webgpu/webgpu.h>
#include <webgpu/wgpu.h>

static const char* kShader =
    "@group(0) @binding(0) var<storage, read_write> data: array<u32>;\n"
    "@compute @workgroup_size(1)\n"
    "fn main() {\n"
    "    data[0] = data[0] + 1u;\n"
    "}\n";

static void die(const char* what)
{
    printf("FAIL: %s\n", what);
    exit(2);
}

static WGPUStringView view(const char* text)
{
    WGPUStringView v;
    v.data   = text;
    v.length = WGPU_STRLEN;
    return v;
}

static void print_view(const char* label, WGPUStringView v)
{
    printf("%s%.*s\n", label, (int) (v.data ? v.length : 0),
           v.data ? v.data : "");
}

static void on_uncaptured_error(WGPUDevice const* device, WGPUErrorType type,
                                WGPUStringView message, void* u1, void* u2)
{
    (void) device;
    (void) u1;
    (void) u2;
    printf("uncaptured error (type 0x%08x): %.*s\n", (unsigned) type,
           (int) (message.data ? message.length : 0),
           message.data ? message.data : "");
}

static void on_adapter(WGPURequestAdapterStatus status, WGPUAdapter adapter,
                       WGPUStringView message, void* userdata1, void* userdata2)
{
    (void) userdata2;
    if (status != WGPURequestAdapterStatus_Success)
        printf("requestAdapter failed (status 0x%08x): %.*s\n",
               (unsigned) status, (int) (message.data ? message.length : 0),
               message.data ? message.data : "");
    *(WGPUAdapter*) userdata1 = adapter;
}

static void on_device(WGPURequestDeviceStatus status, WGPUDevice device,
                      WGPUStringView message, void* userdata1, void* userdata2)
{
    (void) userdata2;
    if (status != WGPURequestDeviceStatus_Success)
        printf("requestDevice failed (status 0x%08x): %.*s\n",
               (unsigned) status, (int) (message.data ? message.length : 0),
               message.data ? message.data : "");
    *(WGPUDevice*) userdata1 = device;
}

static void on_map(WGPUMapAsyncStatus status, WGPUStringView message,
                   void* userdata1, void* userdata2)
{
    (void) userdata2;
    if (status != WGPUMapAsyncStatus_Success)
        printf("mapAsync failed (status 0x%08x): %.*s\n", (unsigned) status,
               (int) (message.data ? message.length : 0),
               message.data ? message.data : "");
    *(int*) userdata1 = 1;
}

static WGPUBuffer make_buffer(WGPUDevice device, const char* label,
                              uint64_t size, WGPUBufferUsage usage)
{
    WGPUBufferDescriptor desc;
    memset(&desc, 0, sizeof(desc));
    desc.label = view(label);
    desc.usage = usage;
    desc.size  = size;

    WGPUBuffer buffer = wgpuDeviceCreateBuffer(device, &desc);
    if (!buffer)
        die("createBuffer returned NULL");
    return buffer;
}

/*
 * Runs one compute pass whose begin/end timestamps go to querySet[first] and
 * querySet[first + 1], with a dispatch only when `dispatch` is set, and reads
 * the two resolved values back into out[0..1].
 */
static void run_pass(WGPUDevice device, WGPUInstance instance, WGPUQueue queue,
                     WGPUQuerySet querySet, uint32_t first, int dispatch,
                     WGPUComputePipeline pipeline, WGPUBindGroup bindGroup,
                     uint64_t out[2])
{
    const uint64_t kResultBytes = 2 * sizeof(uint64_t);

    WGPUBuffer resolved = make_buffer(device, "resolved queries", kResultBytes,
                                      WGPUBufferUsage_QueryResolve | WGPUBufferUsage_CopySrc);
    WGPUBuffer readback = make_buffer(device, "readback", kResultBytes,
                                      WGPUBufferUsage_CopyDst | WGPUBufferUsage_MapRead);

    WGPUCommandEncoder encoder = wgpuDeviceCreateCommandEncoder(device, NULL);

    WGPUPassTimestampWrites timestampWrites;
    memset(&timestampWrites, 0, sizeof(timestampWrites));
    timestampWrites.querySet                  = querySet;
    timestampWrites.beginningOfPassWriteIndex = first;
    timestampWrites.endOfPassWriteIndex       = first + 1;

    WGPUComputePassDescriptor passDesc;
    memset(&passDesc, 0, sizeof(passDesc));
    passDesc.label           = view("timed pass");
    passDesc.timestampWrites = &timestampWrites;

    WGPUComputePassEncoder pass = wgpuCommandEncoderBeginComputePass(encoder, &passDesc);
    if (dispatch)
    {
        wgpuComputePassEncoderSetPipeline(pass, pipeline);
        wgpuComputePassEncoderSetBindGroup(pass, 0, bindGroup, 0, NULL);
        wgpuComputePassEncoderDispatchWorkgroups(pass, 1, 1, 1);
    }
    wgpuComputePassEncoderEnd(pass);
    wgpuComputePassEncoderRelease(pass);

    wgpuCommandEncoderResolveQuerySet(encoder, querySet, first, 2, resolved, 0);
    wgpuCommandEncoderCopyBufferToBuffer(encoder, resolved, 0, readback, 0, kResultBytes);

    WGPUCommandBuffer commands = wgpuCommandEncoderFinish(encoder, NULL);
    if (!commands)
        die("commandEncoderFinish returned NULL");

    WGPUSubmissionIndex submission = wgpuQueueSubmitForIndex(queue, 1, &commands);
    wgpuDevicePoll(device, 1 /* wait */, &submission);

    int mapped = 0;
    WGPUBufferMapCallbackInfo mapCallback;
    memset(&mapCallback, 0, sizeof(mapCallback));
    mapCallback.mode      = WGPUCallbackMode_AllowProcessEvents;
    mapCallback.callback  = on_map;
    mapCallback.userdata1 = &mapped;
    wgpuBufferMapAsync(readback, WGPUMapMode_Read, 0, (size_t) kResultBytes, mapCallback);
    while (!mapped)
    {
        wgpuDevicePoll(device, 1 /* wait */, NULL);
        wgpuInstanceProcessEvents(instance);
    }

    const uint64_t* ticks =
        (const uint64_t*) wgpuBufferGetConstMappedRange(readback, 0, (size_t) kResultBytes);
    if (!ticks)
        die("getConstMappedRange returned NULL");
    out[0] = ticks[0];
    out[1] = ticks[1];

    wgpuBufferUnmap(readback);
    wgpuCommandBufferRelease(commands);
    wgpuCommandEncoderRelease(encoder);
    wgpuBufferRelease(readback);
    wgpuBufferRelease(resolved);
}

int main(void)
{
    WGPUInstance instance = wgpuCreateInstance(NULL);
    if (!instance)
        die("createInstance returned NULL");

    WGPURequestAdapterOptions adapterOptions;
    memset(&adapterOptions, 0, sizeof(adapterOptions));
    adapterOptions.powerPreference = WGPUPowerPreference_HighPerformance;

    WGPUAdapter adapter = NULL;
    WGPURequestAdapterCallbackInfo adapterCallback;
    memset(&adapterCallback, 0, sizeof(adapterCallback));
    adapterCallback.mode      = WGPUCallbackMode_AllowProcessEvents;
    adapterCallback.callback  = on_adapter;
    adapterCallback.userdata1 = &adapter;
    wgpuInstanceRequestAdapter(instance, &adapterOptions, adapterCallback);
    for (int spins = 0; !adapter && spins < 1000; ++spins)
        wgpuInstanceProcessEvents(instance);
    if (!adapter)
        die("no adapter");

    WGPUAdapterInfo info;
    memset(&info, 0, sizeof(info));
    if (wgpuAdapterGetInfo(adapter, &info) == WGPUStatus_Success)
    {
        print_view("adapter:      ", info.device);
        printf("backend type: 0x%08x (5 = Metal)\n", (unsigned) info.backendType);
    }

    if (!wgpuAdapterHasFeature(adapter, WGPUFeatureName_TimestampQuery))
        die("adapter does not advertise timestamp-query");
    printf("adapter advertises timestamp-query: yes\n");

    WGPUFeatureName required[1];
    required[0] = WGPUFeatureName_TimestampQuery;

    WGPUDeviceDescriptor deviceDesc;
    memset(&deviceDesc, 0, sizeof(deviceDesc));
    deviceDesc.label                                = view("repro device");
    deviceDesc.requiredFeatureCount                 = 1;
    deviceDesc.requiredFeatures                     = required;
    deviceDesc.uncapturedErrorCallbackInfo.callback = on_uncaptured_error;

    WGPUDevice device = NULL;
    WGPURequestDeviceCallbackInfo deviceCallback;
    memset(&deviceCallback, 0, sizeof(deviceCallback));
    deviceCallback.mode      = WGPUCallbackMode_AllowProcessEvents;
    deviceCallback.callback  = on_device;
    deviceCallback.userdata1 = &device;
    wgpuAdapterRequestDevice(adapter, &deviceDesc, deviceCallback);
    for (int spins = 0; !device && spins < 1000; ++spins)
        wgpuInstanceProcessEvents(instance);
    if (!device)
        die("no device");

    WGPUQueue queue = wgpuDeviceGetQueue(device);
    printf("timestamp period: %g ns/tick\n",
           (double) wgpuQueueGetTimestampPeriod(queue));

    WGPUQuerySetDescriptor querySetDesc;
    memset(&querySetDesc, 0, sizeof(querySetDesc));
    querySetDesc.label = view("timestamps");
    querySetDesc.type  = WGPUQueryType_Timestamp;
    querySetDesc.count = 4;
    WGPUQuerySet querySet = wgpuDeviceCreateQuerySet(device, &querySetDesc);
    if (!querySet)
        die("createQuerySet returned NULL");

    WGPUShaderSourceWGSL wgsl;
    memset(&wgsl, 0, sizeof(wgsl));
    wgsl.chain.sType = WGPUSType_ShaderSourceWGSL;
    wgsl.code        = view(kShader);

    WGPUShaderModuleDescriptor moduleDesc;
    memset(&moduleDesc, 0, sizeof(moduleDesc));
    moduleDesc.nextInChain = &wgsl.chain;
    moduleDesc.label       = view("trivial");
    WGPUShaderModule module = wgpuDeviceCreateShaderModule(device, &moduleDesc);
    if (!module)
        die("createShaderModule returned NULL");

    WGPUComputePipelineDescriptor pipelineDesc;
    memset(&pipelineDesc, 0, sizeof(pipelineDesc));
    pipelineDesc.label              = view("trivial");
    pipelineDesc.compute.module     = module;
    pipelineDesc.compute.entryPoint = view("main");
    WGPUComputePipeline pipeline = wgpuDeviceCreateComputePipeline(device, &pipelineDesc);
    if (!pipeline)
        die("createComputePipeline returned NULL");

    WGPUBuffer storage = make_buffer(device, "storage", 16, WGPUBufferUsage_Storage);

    WGPUBindGroupEntry entry;
    memset(&entry, 0, sizeof(entry));
    entry.binding = 0;
    entry.buffer  = storage;
    entry.offset  = 0;
    entry.size    = 16;

    WGPUBindGroupDescriptor bindGroupDesc;
    memset(&bindGroupDesc, 0, sizeof(bindGroupDesc));
    bindGroupDesc.label      = view("storage");
    bindGroupDesc.layout     = wgpuComputePipelineGetBindGroupLayout(pipeline, 0);
    bindGroupDesc.entryCount = 1;
    bindGroupDesc.entries    = &entry;
    WGPUBindGroup bindGroup = wgpuDeviceCreateBindGroup(device, &bindGroupDesc);
    if (!bindGroup)
        die("createBindGroup returned NULL");

    uint64_t empty[2] = { 0, 0 };
    uint64_t work[2]  = { 0, 0 };

    run_pass(device, instance, queue, querySet, 0, 0 /* no dispatch */,
             pipeline, bindGroup, empty);
    run_pass(device, instance, queue, querySet, 2, 1 /* one dispatch */,
             pipeline, bindGroup, work);

    printf("\npass A (no dispatch):  begin = %llu ticks, end = %llu ticks\n",
           (unsigned long long) empty[0], (unsigned long long) empty[1]);
    printf("pass B (one dispatch): begin = %llu ticks, end = %llu ticks\n",
           (unsigned long long) work[0], (unsigned long long) work[1]);

    int emptyOk = empty[0] != 0 && empty[1] != 0 && empty[1] >= empty[0];
    int workOk  = work[0] != 0 && work[1] != 0 && work[1] >= work[0];

    printf("\npass A: %s\n", emptyOk ? "PASS: two nonzero, nondecreasing timestamps"
                                     : "FAIL: timestamps are zero or decreasing");
    printf("pass B: %s\n", workOk ? "PASS: two nonzero, nondecreasing timestamps"
                                  : "FAIL: timestamps are zero or decreasing");
    return (emptyOk && workOk) ? 0 : 1;
}

Observed output

# captured 2026-07-27
# macOS 26.5.2 (build 25F84), Apple M2, Metal backend
# wgpu-native v29.0.1.1 (wgpu-macos-aarch64-release prebuilt)
# cc -std=c11 -O2 -Wall -Wextra main.c -o repro (Apple clang)

$ ./repro
adapter:      Apple M2
backend type: 0x00000005 (5 = Metal)
adapter advertises timestamp-query: yes
timestamp period: 1 ns/tick

pass A (no dispatch):  begin = 0 ticks, end = 0 ticks
pass B (one dispatch): begin = 564552285134291 ticks, end = 564552285141583 ticks

pass A: FAIL: timestamps are zero or decreasing
pass B: PASS: two nonzero, nondecreasing timestamps
$ echo $?
1