asahi: Allocate global IDs

Use the same UABI as Metal. One less hack, trying to rule out possible
differences to Metal...

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11815>

src/asahi/lib/agx_device.c

@@ -247,6 +247,39 @@ agx_bo_create(struct agx_device *dev, unsigned size, unsigned flags)
    return bo;
 }
 
+static void
+agx_get_global_ids(struct agx_device *dev)
+{
+#if __APPLE__
+   uint64_t out[2] = {};
+   size_t out_sz = sizeof(out);
+
+   ASSERTED kern_return_t ret = IOConnectCallStructMethod(dev->fd,
+                       AGX_SELECTOR_GET_GLOBAL_IDS,
+                       NULL, 0, &out, &out_sz);
+
+   assert(ret == 0);
+   assert(out_sz == sizeof(out));
+   assert(out[1] > out[0]);
+
+   dev->next_global_id = out[0];
+   dev->last_global_id = out[1];
+#else
+   dev->next_global_id = 0;
+   dev->last_global_id = 0x1000000;
+#endif
+}
+
+uint64_t
+agx_get_global_id(struct agx_device *dev)
+{
+   if (unlikely(dev->next_global_id >= dev->last_global_id)) {
+      agx_get_global_ids(dev);
+   }
+
+   return dev->next_global_id++;
+}
+
 /* Tries to open an AGX device, returns true if successful */
 
 bool
@@ -296,6 +329,8 @@ agx_open_device(void *memctx, struct agx_device *dev)
    dev->queue = agx_create_command_queue(dev);
    dev->cmdbuf = agx_shmem_alloc(dev, 0x4000, true); // length becomes kernelCommandDataSize
    dev->memmap = agx_shmem_alloc(dev, 0x4000, false);
+   agx_get_global_ids(dev);
+
    return true;
 }
 
@@ -312,29 +347,6 @@ agx_close_device(struct agx_device *dev)
 #endif
 }
 
-uint64_t
-agx_cmdbuf_global_ids(struct agx_device *dev)
-{
-#if __APPLE__
-   uint32_t out[4] = {};
-   size_t out_sz = sizeof(out);
-
-   ASSERTED kern_return_t ret = IOConnectCallStructMethod(dev->fd,
-                       0x6,
-                       NULL, 0, &out, &out_sz);
-
-   assert(ret == 0);
-   assert(out_sz == sizeof(out));
-   assert(out[2] == (out[0] + 0x1000000));
-
-   /* Returns a 32-bit but is 64-bit in Instruments, extend with the
-    * missing high bit */
-   return (out[0]) | (1ull << 32ull);
-#else
-   return 0;
-#endif
-}
-
 #if __APPLE__
 static struct agx_notification_queue
 agx_create_notification_queue(mach_port_t connection)

src/asahi/lib/agx_device.h

@@ -46,6 +46,7 @@ struct agx_device {
    /* XXX What to bind to? I don't understand the IOGPU UABI */
    struct agx_command_queue queue;
    struct agx_bo cmdbuf, memmap;
+   uint64_t next_global_id, last_global_id;
 
    /* Device handle */
 #if __APPLE__
@@ -89,7 +90,7 @@ void
 agx_shmem_free(struct agx_device *dev, unsigned handle);
 
 uint64_t
-agx_cmdbuf_global_ids(struct agx_device *dev);
+agx_get_global_id(struct agx_device *dev);
 
 struct agx_command_queue
 agx_create_command_queue(struct agx_device *dev);

src/asahi/lib/io.h

@@ -211,4 +211,7 @@ struct agx_map_entry {
 	uint32_t indices[6];
 } __attribute__((packed));
 
+uint64_t
+agx_get_global_id(struct agx_device *dev);
+
 #endif

src/gallium/drivers/asahi/agx_pipe.c

@@ -510,8 +510,8 @@ agx_flush(struct pipe_context *pctx,
    /* Size calculation should've been exact */
    assert(handle_i == handle_count);
 
-   unsigned cmdbuf_id = 0xDEADBEEF;
-   unsigned encoder_id = 0xCAFECAFE;
+   unsigned cmdbuf_id = agx_get_global_id(dev);
+   unsigned encoder_id = agx_get_global_id(dev);
 
    unsigned cmdbuf_size = demo_cmdbuf(dev->cmdbuf.ptr.cpu,
                dev->cmdbuf.size,