From f4cca9d600eb516aa870012bb8de8bca451ef110 Mon Sep 17 00:00:00 2001
From: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Date: Thu, 1 Aug 2024 10:43:22 -0400
Subject: [PATCH] asahi: introduce agx_va data structure

prep for sparse.

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/30633>
---
 src/asahi/lib/agx_bo.h                     | 17 +++++-
 src/asahi/lib/agx_device.c                 | 70 ++++++++--------------
 src/asahi/lib/agx_device.h                 |  6 ++
 src/asahi/lib/agx_device_virtio.c          | 27 ++++-----
 src/asahi/lib/agx_linker.c                 | 12 ++--
 src/asahi/lib/agx_scratch.c                | 21 ++++---
 src/asahi/lib/agx_va.c                     | 64 ++++++++++++++++++++
 src/asahi/lib/decode.c                     | 20 +++----
 src/asahi/lib/meson.build                  |  1 +
 src/asahi/lib/pool.c                       |  4 +-
 src/asahi/vulkan/hk_buffer.c               | 62 ++++++++-----------
 src/asahi/vulkan/hk_buffer.h               |  4 +-
 src/asahi/vulkan/hk_cmd_buffer.c           | 13 ++--
 src/asahi/vulkan/hk_cmd_draw.c             | 18 +++---
 src/asahi/vulkan/hk_cmd_pool.c             |  2 +-
 src/asahi/vulkan/hk_descriptor_set.c       | 14 ++---
 src/asahi/vulkan/hk_descriptor_table.c     |  2 +-
 src/asahi/vulkan/hk_device.c               | 18 +++---
 src/asahi/vulkan/hk_device_memory.c        |  4 +-
 src/asahi/vulkan/hk_event.c                |  8 +--
 src/asahi/vulkan/hk_image.c                |  4 +-
 src/asahi/vulkan/hk_query_pool.c           | 22 +++----
 src/asahi/vulkan/hk_queue.c                | 20 +++----
 src/asahi/vulkan/hk_shader.c               |  4 +-
 src/gallium/drivers/asahi/agx_batch.c      | 14 ++---
 src/gallium/drivers/asahi/agx_disk_cache.c |  4 +-
 src/gallium/drivers/asahi/agx_pipe.c       | 40 ++++++-------
 src/gallium/drivers/asahi/agx_query.c      | 15 +++--
 src/gallium/drivers/asahi/agx_state.c      | 46 +++++++-------
 src/gallium/drivers/asahi/agx_state.h      |  4 +-
 src/gallium/drivers/asahi/agx_streamout.c  |  2 +-
 src/gallium/drivers/asahi/agx_uniforms.c   |  6 +-
 32 files changed, 309 insertions(+), 259 deletions(-)
 create mode 100644 src/asahi/lib/agx_va.c

diff --git a/src/asahi/lib/agx_bo.h b/src/asahi/lib/agx_bo.h
index 07b213dec8f..18a8390fe51 100644
--- a/src/asahi/lib/agx_bo.h
+++ b/src/asahi/lib/agx_bo.h
@@ -39,6 +39,20 @@ enum agx_bo_flags {
    AGX_BO_READONLY = 1 << 5,
 };
 
+enum agx_va_flags {
+   /* VA must be inside the USC region, otherwise unrestricted. */
+   AGX_VA_USC = (1 << 0),
+
+   /* VA must be fixed, otherwise allocated by the driver. */
+   AGX_VA_FIXED = (1 << 1),
+};
+
+struct agx_va {
+   enum agx_va_flags flags;
+   uint64_t addr;
+   uint64_t size_B;
+};
+
 struct agx_ptr {
    /* If CPU mapped, CPU address. NULL if not mapped */
    void *cpu;
@@ -63,7 +77,8 @@ struct agx_bo {
    size_t align;
 
    /* Mapping */
-   struct agx_ptr ptr;
+   struct agx_va *va;
+   void *map;
 
    /* Process-local index */
    uint32_t handle;
diff --git a/src/asahi/lib/agx_device.c b/src/asahi/lib/agx_device.c
index f6a6b5b4de6..c533caecb95 100644
--- a/src/asahi/lib/agx_device.c
+++ b/src/asahi/lib/agx_device.c
@@ -75,26 +75,13 @@ agx_bo_free(struct agx_device *dev, struct agx_bo *bo)
 {
    const uint64_t handle = bo->handle;
 
-   if (bo->ptr.cpu)
-      munmap(bo->ptr.cpu, bo->size);
+   if (bo->map)
+      munmap(bo->map, bo->size);
 
-   if (bo->ptr.gpu) {
-      struct util_vma_heap *heap;
-      uint64_t bo_addr = bo->ptr.gpu;
-
-      if (bo->flags & AGX_BO_LOW_VA) {
-         heap = &dev->usc_heap;
-      } else {
-         heap = &dev->main_heap;
-      }
-
-      simple_mtx_lock(&dev->vma_lock);
-      util_vma_heap_free(heap, bo_addr, bo->size + dev->guard_size);
-      simple_mtx_unlock(&dev->vma_lock);
-
-      /* No need to unmap the BO, as the kernel will take care of that when we
-       * close it. */
-   }
+   /* Free the VA. No need to unmap the BO, as the kernel will take care of that
+    * when we close it.
+    */
+   agx_va_free(dev, bo->va);
 
    if (bo->prime_fd != -1)
       close(bo->prime_fd);
@@ -176,13 +163,9 @@ agx_bo_alloc(struct agx_device *dev, size_t size, size_t align,
    bo->handle = handle;
    bo->prime_fd = -1;
 
-   struct util_vma_heap *heap =
-      (flags & AGX_BO_LOW_VA) ? &dev->usc_heap : &dev->main_heap;
-
-   simple_mtx_lock(&dev->vma_lock);
-   bo->ptr.gpu = util_vma_heap_alloc(heap, size + dev->guard_size, bo->align);
-   simple_mtx_unlock(&dev->vma_lock);
-   if (!bo->ptr.gpu) {
+   enum agx_va_flags va_flags = flags & AGX_BO_LOW_VA ? AGX_VA_USC : 0;
+   bo->va = agx_va_alloc(dev, size, bo->align, va_flags, 0);
+   if (!bo->va) {
       fprintf(stderr, "Failed to allocate BO VMA\n");
       agx_bo_free(dev, bo);
       return NULL;
@@ -193,7 +176,7 @@ agx_bo_alloc(struct agx_device *dev, size_t size, size_t align,
       bind |= ASAHI_BIND_WRITE;
    }
 
-   ret = dev->ops.bo_bind(dev, bo, bo->ptr.gpu, bind);
+   ret = dev->ops.bo_bind(dev, bo, bo->va->addr, bind);
    if (ret) {
       agx_bo_free(dev, bo);
       return NULL;
@@ -209,7 +192,7 @@ agx_bo_mmap(struct agx_device *dev, struct agx_bo *bo)
    struct drm_asahi_gem_mmap_offset gem_mmap_offset = {.handle = bo->handle};
    int ret;
 
-   if (bo->ptr.cpu)
+   if (bo->map)
       return;
 
    ret = drmIoctl(dev->fd, DRM_IOCTL_ASAHI_GEM_MMAP_OFFSET, &gem_mmap_offset);
@@ -218,14 +201,13 @@ agx_bo_mmap(struct agx_device *dev, struct agx_bo *bo)
       assert(0);
    }
 
-   bo->ptr.cpu = os_mmap(NULL, bo->size, PROT_READ | PROT_WRITE, MAP_SHARED,
-                         dev->fd, gem_mmap_offset.offset);
-   if (bo->ptr.cpu == MAP_FAILED) {
-      bo->ptr.cpu = NULL;
-
+   bo->map = os_mmap(NULL, bo->size, PROT_READ | PROT_WRITE, MAP_SHARED,
+                     dev->fd, gem_mmap_offset.offset);
+   if (bo->map == MAP_FAILED) {
+      bo->map = NULL;
       fprintf(stderr,
               "mmap failed: result=%p size=0x%llx fd=%i offset=0x%llx %m\n",
-              bo->ptr.cpu, (long long)bo->size, dev->fd,
+              bo->map, (long long)bo->size, dev->fd,
               (long long)gem_mmap_offset.offset);
    }
 }
@@ -276,13 +258,9 @@ agx_bo_import(struct agx_device *dev, int fd)
       assert(bo->prime_fd >= 0);
 
       p_atomic_set(&bo->refcnt, 1);
+      bo->va = agx_va_alloc(dev, bo->size, bo->align, 0, 0);
 
-      simple_mtx_lock(&dev->vma_lock);
-      bo->ptr.gpu = util_vma_heap_alloc(
-         &dev->main_heap, bo->size + dev->guard_size, dev->params.vm_page_size);
-      simple_mtx_unlock(&dev->vma_lock);
-
-      if (!bo->ptr.gpu) {
+      if (!bo->va) {
          fprintf(
             stderr,
             "import failed: Could not allocate from VMA heap (0x%llx bytes)\n",
@@ -294,11 +272,11 @@ agx_bo_import(struct agx_device *dev, int fd)
          bo->vbo_res_id = vdrm_handle_to_res_id(dev->vdrm, bo->handle);
       }
 
-      ret = dev->ops.bo_bind(dev, bo, bo->ptr.gpu,
+      ret = dev->ops.bo_bind(dev, bo, bo->va->addr,
                              ASAHI_BIND_READ | ASAHI_BIND_WRITE);
       if (ret) {
          fprintf(stderr, "import failed: Could not bind BO at 0x%llx\n",
-                 (long long)bo->ptr.gpu);
+                 (long long)bo->va->addr);
          abort();
       }
    } else {
@@ -691,22 +669,22 @@ agx_debug_fault(struct agx_device *dev, uint64_t addr)
 
    for (uint32_t handle = 0; handle < dev->max_handle; handle++) {
       struct agx_bo *bo = agx_lookup_bo(dev, handle);
-      uint64_t bo_addr = bo->ptr.gpu;
+      uint64_t bo_addr = bo->va->addr;
       if (bo->flags & AGX_BO_LOW_VA)
          bo_addr += dev->shader_base;
 
       if (!bo->size || bo_addr > addr)
          continue;
 
-      if (!best || bo_addr > best->ptr.gpu)
+      if (!best || bo_addr > best->va->addr)
          best = bo;
    }
 
    if (!best) {
       mesa_logw("Address 0x%" PRIx64 " is unknown\n", addr);
    } else {
-      uint64_t start = best->ptr.gpu;
-      uint64_t end = best->ptr.gpu + best->size;
+      uint64_t start = best->va->addr;
+      uint64_t end = best->va->addr + best->size;
       if (addr > (end + 1024 * 1024 * 1024)) {
          /* 1GiB max as a sanity check */
          mesa_logw("Address 0x%" PRIx64 " is unknown\n", addr);
diff --git a/src/asahi/lib/agx_device.h b/src/asahi/lib/agx_device.h
index 067f238e0d5..c9ab7eb0829 100644
--- a/src/asahi/lib/agx_device.h
+++ b/src/asahi/lib/agx_device.h
@@ -7,6 +7,7 @@
 
 #include <stdint.h>
 #include <xf86drm.h>
+#include "util/ralloc.h"
 #include "util/simple_mtx.h"
 #include "util/sparse_array.h"
 #include "util/timespec.h"
@@ -180,3 +181,8 @@ agx_gpu_time_to_ns(struct agx_device *dev, uint64_t gpu_time)
 
 void agx_get_device_uuid(const struct agx_device *dev, void *uuid);
 void agx_get_driver_uuid(void *uuid);
+
+struct agx_va *agx_va_alloc(struct agx_device *dev, uint32_t size_B,
+                            uint32_t align_B, enum agx_va_flags flags,
+                            uint64_t fixed_va);
+void agx_va_free(struct agx_device *dev, struct agx_va *va);
diff --git a/src/asahi/lib/agx_device_virtio.c b/src/asahi/lib/agx_device_virtio.c
index d5384e30796..1435aa90a4c 100644
--- a/src/asahi/lib/agx_device_virtio.c
+++ b/src/asahi/lib/agx_device_virtio.c
@@ -58,7 +58,6 @@ agx_virtio_bo_alloc(struct agx_device *dev, size_t size, size_t align,
 {
    struct agx_bo *bo;
    unsigned handle = 0;
-   uint64_t ptr_gpu;
 
    size = ALIGN_POT(size, dev->params.vm_page_size);
 
@@ -83,19 +82,15 @@ agx_virtio_bo_alloc(struct agx_device *dev, size_t size, size_t align,
 
    uint32_t blob_id = p_atomic_inc_return(&dev->next_blob_id);
 
-   struct util_vma_heap *heap =
-      (flags & AGX_BO_LOW_VA) ? &dev->usc_heap : &dev->main_heap;
-
-   simple_mtx_lock(&dev->vma_lock);
-   ptr_gpu = util_vma_heap_alloc(heap, size + dev->guard_size,
-                                 dev->params.vm_page_size);
-   simple_mtx_unlock(&dev->vma_lock);
-   if (!ptr_gpu) {
+   enum agx_va_flags va_flags = flags & AGX_BO_LOW_VA ? AGX_VA_USC : 0;
+   struct agx_va *va =
+      agx_va_alloc(dev, size, dev->params.vm_page_size, va_flags, 0);
+   if (!va) {
       fprintf(stderr, "Failed to allocate BO VMA\n");
       return NULL;
    }
 
-   req.addr = ptr_gpu;
+   req.addr = va->addr;
    req.blob_id = blob_id;
    req.vm_id = dev->vm_id;
 
@@ -119,7 +114,7 @@ agx_virtio_bo_alloc(struct agx_device *dev, size_t size, size_t align,
    bo->handle = handle;
    bo->prime_fd = -1;
    bo->blob_id = blob_id;
-   bo->ptr.gpu = ptr_gpu;
+   bo->va = va;
    bo->vbo_res_id = vdrm_handle_to_res_id(dev->vdrm, handle);
 
    dev->ops.bo_mmap(dev, bo);
@@ -153,14 +148,14 @@ agx_virtio_bo_bind(struct agx_device *dev, struct agx_bo *bo, uint64_t addr,
 static void
 agx_virtio_bo_mmap(struct agx_device *dev, struct agx_bo *bo)
 {
-   if (bo->ptr.cpu) {
+   if (bo->map) {
       return;
    }
 
-   bo->ptr.cpu = vdrm_bo_map(dev->vdrm, bo->handle, bo->size, NULL);
-   if (bo->ptr.cpu == MAP_FAILED) {
-      bo->ptr.cpu = NULL;
-      fprintf(stderr, "mmap failed: result=%p size=0x%llx fd=%i\n", bo->ptr.cpu,
+   bo->map = vdrm_bo_map(dev->vdrm, bo->handle, bo->size, NULL);
+   if (bo->map == MAP_FAILED) {
+      bo->map = NULL;
+      fprintf(stderr, "mmap failed: result=%p size=0x%llx fd=%i\n", bo->map,
               (long long)bo->size, dev->fd);
    }
 }
diff --git a/src/asahi/lib/agx_linker.c b/src/asahi/lib/agx_linker.c
index 4933a3aec79..a88b2d6691c 100644
--- a/src/asahi/lib/agx_linker.c
+++ b/src/asahi/lib/agx_linker.c
@@ -123,12 +123,12 @@ agx_fast_link(struct agx_linked_shader *linked, struct agx_device *dev,
    /* FS prolog happens per-pixel, outside the sample loop */
    if (prolog) {
       size_t sz = prolog->info.main_size;
-      memcpy((uint8_t *)linked->bo->ptr.cpu + offset, prolog->binary, sz);
+      memcpy((uint8_t *)linked->bo->map + offset, prolog->binary, sz);
       offset += sz;
    }
 
    if (nr_samples_shaded) {
-      memcpy((uint8_t *)linked->bo->ptr.cpu + offset, sample_loop_header,
+      memcpy((uint8_t *)linked->bo->map + offset, sample_loop_header,
              sizeof(sample_loop_header));
       offset += sizeof(sample_loop_header);
    }
@@ -142,7 +142,7 @@ agx_fast_link(struct agx_linked_shader *linked, struct agx_device *dev,
          continue;
 
       size_t sz = part->info.main_size;
-      memcpy((uint8_t *)linked->bo->ptr.cpu + offset, part->binary, sz);
+      memcpy((uint8_t *)linked->bo->map + offset, part->binary, sz);
       offset += sz;
    }
 
@@ -165,18 +165,18 @@ agx_fast_link(struct agx_linked_shader *linked, struct agx_device *dev,
       *target = branch_offs;
 
       /* Copy in the patched footer */
-      memcpy((uint8_t *)linked->bo->ptr.cpu + offset, footer, sizeof(footer));
+      memcpy((uint8_t *)linked->bo->map + offset, footer, sizeof(footer));
       offset += sizeof(footer);
    } else if (nr_samples_shaded) {
       /* Just end after the first sample, no need to loop for a single sample */
-      memcpy((uint8_t *)linked->bo->ptr.cpu + offset, stop, sizeof(stop));
+      memcpy((uint8_t *)linked->bo->map + offset, stop, sizeof(stop));
       offset += sizeof(stop);
    }
 
    assert(offset == size);
 
    agx_pack(&linked->shader, USC_SHADER, cfg) {
-      cfg.code = agx_usc_addr(dev, linked->bo->ptr.gpu);
+      cfg.code = agx_usc_addr(dev, linked->bo->va->addr);
       cfg.unk_2 = fragment ? 2 : 3;
 
       if (fragment)
diff --git a/src/asahi/lib/agx_scratch.c b/src/asahi/lib/agx_scratch.c
index d652977cedc..d6d3ec59a90 100644
--- a/src/asahi/lib/agx_scratch.c
+++ b/src/asahi/lib/agx_scratch.c
@@ -36,10 +36,10 @@ agx_build_helper(struct agx_device *dev)
       dev, sizeof(libagx_g13_helper), 0,
       AGX_BO_READONLY | AGX_BO_EXEC | AGX_BO_LOW_VA, "Helper shader");
    assert(bo);
-   memcpy(bo->ptr.cpu, libagx_g13_helper, sizeof(libagx_g13_helper));
+   memcpy(bo->map, libagx_g13_helper, sizeof(libagx_g13_helper));
 
    if (dev->debug & AGX_DBG_SCRATCH)
-      fprintf(stderr, "Helper: 0x%" PRIx64 "\n", bo->ptr.gpu);
+      fprintf(stderr, "Helper: 0x%" PRIx64 "\n", bo->va->addr);
 
    return bo;
 }
@@ -132,22 +132,21 @@ agx_scratch_realloc(struct agx_scratch *scratch)
 #endif
    scratch->buf = agx_bo_create(scratch->dev, total_alloc, block_size_bytes,
                                 flags, "Scratch");
-   memset(scratch->buf->ptr.cpu, 0, blocks_off);
+   memset(scratch->buf->map, 0, blocks_off);
 
-   struct agx_helper_header *hdr = scratch->buf->ptr.cpu;
+   struct agx_helper_header *hdr = scratch->buf->map;
    scratch->header = hdr;
 
-   uint64_t blocklist_gpu = scratch->buf->ptr.gpu + blocklist_off;
-   struct agx_helper_block *blocklist_cpu =
-      scratch->buf->ptr.cpu + blocklist_off;
+   uint64_t blocklist_gpu = scratch->buf->va->addr + blocklist_off;
+   struct agx_helper_block *blocklist_cpu = scratch->buf->map + blocklist_off;
 
 #ifdef SCRATCH_DEBUG
    scratch->blocklist = blocklist_cpu;
-   scratch->data = scratch->buf->ptr.cpu + blocks_off;
+   scratch->data = scratch->buf->map + blocks_off;
    scratch->core_size = block_size_bytes * block_count * scratch->subgroups;
 #endif
 
-   uint64_t blocks_gpu = scratch->buf->ptr.gpu + blocks_off;
+   uint64_t blocks_gpu = scratch->buf->va->addr + blocks_off;
 
    hdr->subgroups = scratch->subgroups;
 
@@ -197,7 +196,7 @@ agx_scratch_realloc(struct agx_scratch *scratch)
 
    if (scratch->dev->debug & AGX_DBG_SCRATCH)
       fprintf(stderr, "New Scratch @ 0x%" PRIx64 " (size: 0x%zx)\n",
-              scratch->buf->ptr.gpu, scratch->buf->size);
+              scratch->buf->va->addr, scratch->buf->size);
 }
 
 void
@@ -252,7 +251,7 @@ agx_scratch_debug_post(struct agx_scratch *scratch)
    if (!scratch->buf)
       return;
 
-   fprintf(stderr, "Scratch @ 0x%" PRIx64 "\n", scratch->buf->ptr.gpu);
+   fprintf(stderr, "Scratch @ 0x%" PRIx64 "\n", scratch->buf->va->addr);
 
    for (int core = 0; core < scratch->max_core_id; core++) {
       fprintf(stderr, "Core %3d: max %d, failed %d, counts:", core,
diff --git a/src/asahi/lib/agx_va.c b/src/asahi/lib/agx_va.c
new file mode 100644
index 00000000000..6cc344b215d
--- /dev/null
+++ b/src/asahi/lib/agx_va.c
@@ -0,0 +1,64 @@
+/*
+ * Copyright 2024 Valve Corporation
+ * SPDX-License-Identifier: MIT
+ */
+
+#include "agx_bo.h"
+#include "agx_device.h"
+
+static struct util_vma_heap *
+agx_vma_heap(struct agx_device *dev, enum agx_va_flags flags)
+{
+   return (flags & AGX_VA_USC) ? &dev->usc_heap : &dev->main_heap;
+}
+
+struct agx_va *
+agx_va_alloc(struct agx_device *dev, uint32_t size_B, uint32_t align_B,
+             enum agx_va_flags flags, uint64_t fixed_va)
+{
+   assert((fixed_va != 0) == !!(flags & AGX_VA_FIXED));
+   assert((fixed_va % align_B) == 0);
+
+   /* All allocations need a guard at the end to prevent overreads.
+    *
+    * TODO: Even with soft fault?
+    */
+   size_B += dev->guard_size;
+
+   struct util_vma_heap *heap = agx_vma_heap(dev, flags);
+   uint64_t addr = 0;
+
+   simple_mtx_lock(&dev->vma_lock);
+   if (flags & AGX_VA_FIXED) {
+      if (util_vma_heap_alloc_addr(heap, fixed_va, size_B))
+         addr = fixed_va;
+   } else {
+      addr = util_vma_heap_alloc(heap, size_B, align_B);
+   }
+   simple_mtx_unlock(&dev->vma_lock);
+
+   if (addr == 0)
+      return NULL;
+
+   struct agx_va *va = malloc(sizeof(struct agx_va));
+   *va = (struct agx_va){
+      .flags = flags,
+      .size_B = size_B,
+      .addr = addr,
+   };
+   return va;
+}
+
+void
+agx_va_free(struct agx_device *dev, struct agx_va *va)
+{
+   if (!va)
+      return;
+
+   struct util_vma_heap *heap = agx_vma_heap(dev, va->flags);
+
+   simple_mtx_lock(&dev->vma_lock);
+   util_vma_heap_free(heap, va->addr, va->size_B);
+   simple_mtx_unlock(&dev->vma_lock);
+   free(va);
+}
diff --git a/src/asahi/lib/decode.c b/src/asahi/lib/decode.c
index ee493157d2e..417f6d865c6 100644
--- a/src/asahi/lib/decode.c
+++ b/src/asahi/lib/decode.c
@@ -65,7 +65,7 @@ agxdecode_find_mapped_gpu_mem_containing(struct agxdecode_ctx *ctx,
                                          uint64_t addr)
 {
    util_dynarray_foreach(&ctx->mmap_array, struct agx_bo, it) {
-      if (addr >= it->ptr.gpu && (addr - it->ptr.gpu) < it->size)
+      if (it->va && addr >= it->va->addr && (addr - it->va->addr) < it->size)
          return it;
    }
 
@@ -103,17 +103,17 @@ __agxdecode_fetch_gpu_mem(struct agxdecode_ctx *ctx, const struct agx_bo *mem,
 
    assert(mem);
 
-   if (size + (gpu_va - mem->ptr.gpu) > mem->size) {
+   if (size + (gpu_va - mem->va->addr) > mem->size) {
       fprintf(stderr,
               "Overflowing to unknown memory %" PRIx64
               " of size %zu (max size %zu) in %s:%d\n",
-              gpu_va, size, (size_t)(mem->size - (gpu_va - mem->ptr.gpu)),
+              gpu_va, size, (size_t)(mem->size - (gpu_va - mem->va->addr)),
               filename, line);
       fflush(agxdecode_dump_stream);
       assert(0);
    }
 
-   memcpy(buf, mem->ptr.cpu + gpu_va - mem->ptr.gpu, size);
+   memcpy(buf, mem->map + gpu_va - mem->va->addr, size);
 
    return size;
 }
@@ -173,7 +173,7 @@ agxdecode_stateful(struct agxdecode_ctx *ctx, uint64_t va, const char *label,
       assert(alloc != NULL && "nonexistent object");
       fprintf(agxdecode_dump_stream, "%s (%" PRIx64 ", handle %u)\n", label, va,
               alloc->handle);
-      size = MIN2(size, alloc->size - (va - alloc->ptr.gpu));
+      size = MIN2(size, alloc->size - (va - alloc->va->addr));
    } else {
       fprintf(agxdecode_dump_stream, "%s (%" PRIx64 ")\n", label, va);
    }
@@ -1007,15 +1007,15 @@ agxdecode_cmdstream(struct agxdecode_ctx *ctx, unsigned cmdbuf_handle,
    assert(map != NULL && "nonexistent mapping");
 
    /* Print the IOGPU stuff */
-   agx_unpack(agxdecode_dump_stream, cmdbuf->ptr.cpu, IOGPU_HEADER, cmd);
+   agx_unpack(agxdecode_dump_stream, cmdbuf->map, IOGPU_HEADER, cmd);
    DUMP_UNPACKED(IOGPU_HEADER, cmd, "IOGPU Header\n");
 
    DUMP_CL(IOGPU_ATTACHMENT_COUNT,
-           ((uint8_t *)cmdbuf->ptr.cpu + cmd.attachment_offset),
+           ((uint8_t *)cmdbuf->map + cmd.attachment_offset),
            "Attachment count");
 
    uint32_t *attachments =
-      (uint32_t *)((uint8_t *)cmdbuf->ptr.cpu + cmd.attachment_offset);
+      (uint32_t *)((uint8_t *)cmdbuf->map + cmd.attachment_offset);
    unsigned attachment_count = attachments[3];
    for (unsigned i = 0; i < attachment_count; ++i) {
       uint32_t *ptr = attachments + 4 + (i * AGX_IOGPU_ATTACHMENT_LENGTH / 4);
@@ -1027,9 +1027,9 @@ agxdecode_cmdstream(struct agxdecode_ctx *ctx, unsigned cmdbuf_handle,
    chip_id_to_params(&params, 0x8103);
 
    if (cmd.unk_5 == 3)
-      agxdecode_cs((uint32_t *)cmdbuf->ptr.cpu, cmd.encoder, verbose, &params);
+      agxdecode_cs((uint32_t *)cmdbuf->map, cmd.encoder, verbose, &params);
    else
-      agxdecode_gfx((uint32_t *)cmdbuf->ptr.cpu, cmd.encoder, verbose, &params);
+      agxdecode_gfx((uint32_t *)cmdbuf->map, cmd.encoder, verbose, &params);
 }
 
 #endif
diff --git a/src/asahi/lib/meson.build b/src/asahi/lib/meson.build
index b29d71f115d..d9ab2c3dc41 100644
--- a/src/asahi/lib/meson.build
+++ b/src/asahi/lib/meson.build
@@ -27,6 +27,7 @@ libasahi_lib_files = files(
   'agx_nir_prolog_epilog.c',
   'agx_ppp.h',
   'agx_scratch.c',
+  'agx_va.c',
   'pool.c',
 )
 
diff --git a/src/asahi/lib/pool.c b/src/asahi/lib/pool.c
index cff079f99d2..3280fe2d055 100644
--- a/src/asahi/lib/pool.c
+++ b/src/asahi/lib/pool.c
@@ -80,8 +80,8 @@ agx_pool_alloc_aligned_with_bo(struct agx_pool *pool, size_t sz,
    pool->transient_offset = offset + sz;
 
    struct agx_ptr ret = {
-      .cpu = bo->ptr.cpu + offset,
-      .gpu = bo->ptr.gpu + offset,
+      .cpu = bo->map + offset,
+      .gpu = bo->va->addr + offset,
    };
 
    if (out_bo)
diff --git a/src/asahi/vulkan/hk_buffer.c b/src/asahi/vulkan/hk_buffer.c
index 63bec5a0f70..6a56b2c3edb 100644
--- a/src/asahi/vulkan/hk_buffer.c
+++ b/src/asahi/vulkan/hk_buffer.c
@@ -6,6 +6,8 @@
  */
 #include "hk_buffer.h"
 
+#include "agx_bo.h"
+#include "agx_device.h"
 #include "hk_device.h"
 #include "hk_device_memory.h"
 #include "hk_entrypoints.h"
@@ -95,34 +97,33 @@ hk_CreateBuffer(VkDevice device, const VkBufferCreateInfo *pCreateInfo,
         (VK_BUFFER_CREATE_SPARSE_BINDING_BIT |
          VK_BUFFER_CREATE_DEVICE_ADDRESS_CAPTURE_REPLAY_BIT))) {
 
-      unreachable("todo");
-#if 0
       const uint32_t alignment =
          hk_get_buffer_alignment(hk_device_physical(dev),
                                   buffer->vk.usage,
                                   buffer->vk.create_flags);
       assert(alignment >= 4096);
-      buffer->vma_size_B = align64(buffer->vk.size, alignment);
+      uint64_t vma_size_B = align64(buffer->vk.size, alignment);
 
-      const bool sparse_residency =
-         buffer->vk.create_flags & VK_BUFFER_CREATE_SPARSE_RESIDENCY_BIT;
       const bool bda_capture_replay =
          buffer->vk.create_flags & VK_BUFFER_CREATE_DEVICE_ADDRESS_CAPTURE_REPLAY_BIT;
 
-      uint64_t bda_replay_addr = 0;
-      if (bda_capture_replay)
-         bda_replay_addr = hk_get_bda_replay_addr(pCreateInfo);
+      enum agx_va_flags flags = 0;
+      uint64_t bda_fixed_addr = 0;
+      if (bda_capture_replay) {
+         bda_fixed_addr = hk_get_bda_replay_addr(pCreateInfo);
+         if (bda_fixed_addr != 0)
+            flags |= AGX_VA_FIXED;
+      }
 
-      buffer->addr = nouveau_ws_alloc_vma(dev->ws_dev, bda_replay_addr,
-                                          buffer->vma_size_B,
-                                          alignment, bda_capture_replay,
-                                          sparse_residency);
-#endif
-      if (buffer->addr == 0) {
+      buffer->va =
+         agx_va_alloc(&dev->dev, vma_size_B, alignment, flags, bda_fixed_addr);
+
+      if (!buffer->va) {
          vk_buffer_destroy(&dev->vk, pAllocator, &buffer->vk);
          return vk_errorf(dev, VK_ERROR_OUT_OF_DEVICE_MEMORY,
                           "Sparse VMA allocation failed");
       }
+      buffer->addr = buffer->va->addr;
    }
 
    *pBuffer = hk_buffer_to_handle(buffer);
@@ -140,19 +141,10 @@ hk_DestroyBuffer(VkDevice device, VkBuffer _buffer,
    if (!buffer)
       return;
 
-   if (buffer->vma_size_B > 0) {
-      unreachable("todo");
-#if 0
-      const bool sparse_residency =
-         buffer->vk.create_flags & VK_BUFFER_CREATE_SPARSE_RESIDENCY_BIT;
-      const bool bda_capture_replay =
-         buffer->vk.create_flags &
-         VK_BUFFER_CREATE_DEVICE_ADDRESS_CAPTURE_REPLAY_BIT;
-
-      agx_bo_unbind_vma(dev->ws_dev, buffer->addr, buffer->vma_size_B);
-      nouveau_ws_free_vma(dev->ws_dev, buffer->addr, buffer->vma_size_B,
-                          bda_capture_replay, sparse_residency);
-#endif
+   if (buffer->va) {
+      // TODO
+      // agx_bo_unbind_vma(dev->ws_dev, buffer->addr, buffer->vma_size_B);
+      agx_va_free(&dev->dev, buffer->va);
    }
 
    vk_buffer_destroy(&dev->vk, pAllocator, &buffer->vk);
@@ -244,19 +236,13 @@ hk_BindBufferMemory2(VkDevice device, uint32_t bindInfoCount,
       VK_FROM_HANDLE(hk_device_memory, mem, pBindInfos[i].memory);
       VK_FROM_HANDLE(hk_buffer, buffer, pBindInfos[i].buffer);
 
-      if (buffer->vma_size_B) {
-         unreachable("todo");
-#if 0
+      if (buffer->va) {
          VK_FROM_HANDLE(hk_device, dev, device);
-         agx_bo_bind_vma(dev->ws_dev,
-                                mem->bo,
-                                buffer->addr,
-                                buffer->vma_size_B,
-                                pBindInfos[i].memoryOffset,
-                                0 /* pte_kind */);
-#endif
+         /* XXX: offset, range */
+         dev->dev.ops.bo_bind(&dev->dev, mem->bo, buffer->addr,
+                              ASAHI_BIND_READ | ASAHI_BIND_WRITE);
       } else {
-         buffer->addr = mem->bo->ptr.gpu + pBindInfos[i].memoryOffset;
+         buffer->addr = mem->bo->va->addr + pBindInfos[i].memoryOffset;
       }
 
       const VkBindMemoryStatusKHR *status =
diff --git a/src/asahi/vulkan/hk_buffer.h b/src/asahi/vulkan/hk_buffer.h
index f349a3df0e2..ed75eab0651 100644
--- a/src/asahi/vulkan/hk_buffer.h
+++ b/src/asahi/vulkan/hk_buffer.h
@@ -18,8 +18,8 @@ struct hk_buffer {
    struct vk_buffer vk;
    uint64_t addr;
 
-   /** Size of the reserved VMA range for sparse buffers, zero otherwise. */
-   uint64_t vma_size_B;
+   /** Reserved VA for sparse buffers, NULL otherwise. */
+   struct agx_va *va;
 };
 
 VK_DEFINE_NONDISP_HANDLE_CASTS(hk_buffer, vk.base, VkBuffer,
diff --git a/src/asahi/vulkan/hk_cmd_buffer.c b/src/asahi/vulkan/hk_cmd_buffer.c
index 525307fd6e9..b776f6061b2 100644
--- a/src/asahi/vulkan/hk_cmd_buffer.c
+++ b/src/asahi/vulkan/hk_cmd_buffer.c
@@ -184,7 +184,10 @@ hk_pool_alloc_internal(struct hk_cmd_buffer *cmd, uint32_t size,
          agx_bo_create(&dev->dev, size, flags, 0, "Large pool allocation");
 
       util_dynarray_append(&cmd->large_bos, struct agx_bo *, bo);
-      return bo->ptr;
+      return (struct agx_ptr){
+         .gpu = bo->va->addr,
+         .cpu = bo->map,
+      };
    }
 
    assert(size <= HK_CMD_BO_SIZE);
@@ -216,13 +219,13 @@ hk_pool_alloc_internal(struct hk_cmd_buffer *cmd, uint32_t size,
     * BO.
     */
    if (uploader->map == NULL || size < uploader->offset) {
-      uploader->map = bo->bo->ptr.cpu;
-      uploader->base = bo->bo->ptr.gpu;
+      uploader->map = bo->bo->map;
+      uploader->base = bo->bo->va->addr;
       uploader->offset = size;
    }
 
    return (struct agx_ptr){
-      .gpu = bo->bo->ptr.gpu,
+      .gpu = bo->bo->va->addr,
       .cpu = bo->map,
    };
 }
@@ -527,7 +530,7 @@ hk_cmd_buffer_upload_root(struct hk_cmd_buffer *cmd,
    struct hk_root_descriptor_table *root = &desc->root;
 
    struct agx_ptr root_ptr = hk_pool_alloc(cmd, sizeof(*root), 8);
-   if (!root_ptr.cpu)
+   if (!root_ptr.gpu)
       return 0;
 
    root->root_desc_addr = root_ptr.gpu;
diff --git a/src/asahi/vulkan/hk_cmd_draw.c b/src/asahi/vulkan/hk_cmd_draw.c
index 23b6635d35c..6fdebe73c42 100644
--- a/src/asahi/vulkan/hk_cmd_draw.c
+++ b/src/asahi/vulkan/hk_cmd_draw.c
@@ -347,7 +347,7 @@ hk_build_bg_eot(struct hk_cmd_buffer *cmd, const VkRenderingInfo *info,
             /* Shifted to match eMRT indexing, could be optimized */
             cfg.start = rt * 2;
             cfg.count = 1;
-            cfg.buffer = dev->images.bo->ptr.gpu + index * AGX_TEXTURE_LENGTH;
+            cfg.buffer = dev->images.bo->va->addr + index * AGX_TEXTURE_LENGTH;
          }
 
          nr_tex = (rt * 2) + 1;
@@ -366,7 +366,7 @@ hk_build_bg_eot(struct hk_cmd_buffer *cmd, const VkRenderingInfo *info,
          agx_usc_pack(&b, TEXTURE, cfg) {
             cfg.start = rt;
             cfg.count = 1;
-            cfg.buffer = dev->images.bo->ptr.gpu + index * AGX_TEXTURE_LENGTH;
+            cfg.buffer = dev->images.bo->va->addr + index * AGX_TEXTURE_LENGTH;
          }
 
          nr_tex = rt + 1;
@@ -876,11 +876,11 @@ hk_geometry_state(struct hk_cmd_buffer *cmd)
       /* The geometry state buffer is initialized here and then is treated by
        * the CPU as rodata, even though the GPU uses it for scratch internally.
        */
-      off_t off = dev->rodata.geometry_state - dev->rodata.bo->ptr.gpu;
-      struct agx_geometry_state *map = dev->rodata.bo->ptr.cpu + off;
+      off_t off = dev->rodata.geometry_state - dev->rodata.bo->va->addr;
+      struct agx_geometry_state *map = dev->rodata.bo->map + off;
 
       *map = (struct agx_geometry_state){
-         .heap = dev->heap->ptr.gpu,
+         .heap = dev->heap->va->addr,
          .heap_size = size,
       };
    }
@@ -1109,7 +1109,7 @@ hk_upload_tess_params(struct hk_cmd_buffer *cmd, struct hk_draw draw)
    size_t draw_stride_B = draw_stride_el * sizeof(uint32_t);
 
    /* heap is allocated by hk_geometry_state */
-   args.patch_coord_buffer = dev->heap->ptr.gpu;
+   args.patch_coord_buffer = dev->heap->va->addr;
 
    if (!draw.b.indirect) {
       unsigned in_patches = draw.b.count[0] / args.input_patch_size;
@@ -1323,7 +1323,7 @@ hk_draw_without_restart(struct hk_cmd_buffer *cmd, struct hk_cs *cs,
                         hk_grid(1024, 1, 1));
 
    struct hk_addr_range out_index = {
-      .addr = dev->heap->ptr.gpu,
+      .addr = dev->heap->va->addr,
       .range = dev->heap->size,
    };
 
@@ -1423,7 +1423,7 @@ hk_launch_gs_prerast(struct hk_cmd_buffer *cmd, struct hk_cs *cs,
    hk_dispatch_with_local_size(cmd, cs, main, grid_gs, hk_grid(1, 1, 1));
 
    struct hk_addr_range range = (struct hk_addr_range){
-      .addr = dev->heap->ptr.gpu,
+      .addr = dev->heap->va->addr,
       .range = dev->heap->size,
    };
 
@@ -1586,7 +1586,7 @@ hk_launch_tess(struct hk_cmd_buffer *cmd, struct hk_cs *cs, struct hk_draw draw)
    }
 
    struct hk_addr_range range = (struct hk_addr_range){
-      .addr = dev->heap->ptr.gpu,
+      .addr = dev->heap->va->addr,
       .range = dev->heap->size,
    };
 
diff --git a/src/asahi/vulkan/hk_cmd_pool.c b/src/asahi/vulkan/hk_cmd_pool.c
index 4b44703b3bb..8834b5543d2 100644
--- a/src/asahi/vulkan/hk_cmd_pool.c
+++ b/src/asahi/vulkan/hk_cmd_pool.c
@@ -29,7 +29,7 @@ hk_cmd_bo_create(struct hk_cmd_pool *pool, bool usc, struct hk_cmd_bo **bo_out)
       return vk_error(pool, VK_ERROR_OUT_OF_DEVICE_MEMORY);
    }
 
-   bo->map = bo->bo->ptr.cpu;
+   bo->map = bo->bo->map;
 
    *bo_out = bo;
    return VK_SUCCESS;
diff --git a/src/asahi/vulkan/hk_descriptor_set.c b/src/asahi/vulkan/hk_descriptor_set.c
index 7852a84a069..f22945feb5b 100644
--- a/src/asahi/vulkan/hk_descriptor_set.c
+++ b/src/asahi/vulkan/hk_descriptor_set.c
@@ -489,13 +489,13 @@ hk_CreateDescriptorPool(VkDevice _device,
          return vk_error(dev, VK_ERROR_OUT_OF_DEVICE_MEMORY);
       }
 
-      pool->mapped_ptr = pool->bo->ptr.cpu;
+      pool->mapped_ptr = pool->bo->map;
 
       /* The BO may be larger thanks to GPU page alignment.  We may as well
        * make that extra space available to the client.
        */
       assert(pool->bo->size >= bo_size);
-      util_vma_heap_init(&pool->heap, pool->bo->ptr.gpu, pool->bo->size);
+      util_vma_heap_init(&pool->heap, pool->bo->va->addr, pool->bo->size);
    } else {
       util_vma_heap_init(&pool->heap, 0, 0);
    }
@@ -513,9 +513,9 @@ hk_descriptor_pool_alloc(struct hk_descriptor_pool *pool, uint64_t size,
    if (addr == 0)
       return VK_ERROR_OUT_OF_POOL_MEMORY;
 
-   assert(addr >= pool->bo->ptr.gpu);
-   assert(addr + size <= pool->bo->ptr.gpu + pool->bo->size);
-   uint64_t offset = addr - pool->bo->ptr.gpu;
+   assert(addr >= pool->bo->va->addr);
+   assert(addr + size <= pool->bo->va->addr + pool->bo->size);
+   uint64_t offset = addr - pool->bo->va->addr;
 
    *addr_out = addr;
    *map_out = pool->mapped_ptr + offset;
@@ -528,8 +528,8 @@ hk_descriptor_pool_free(struct hk_descriptor_pool *pool, uint64_t addr,
                         uint64_t size)
 {
    assert(size > 0);
-   assert(addr >= pool->bo->ptr.gpu);
-   assert(addr + size <= pool->bo->ptr.gpu + pool->bo->size);
+   assert(addr >= pool->bo->va->addr);
+   assert(addr + size <= pool->bo->va->addr + pool->bo->size);
    util_vma_heap_free(&pool->heap, addr, size);
 }
 
diff --git a/src/asahi/vulkan/hk_descriptor_table.c b/src/asahi/vulkan/hk_descriptor_table.c
index cf50a97cfaa..28e26a19cef 100644
--- a/src/asahi/vulkan/hk_descriptor_table.c
+++ b/src/asahi/vulkan/hk_descriptor_table.c
@@ -30,7 +30,7 @@ hk_descriptor_table_grow_locked(struct hk_device *dev,
                        "Failed to allocate the descriptor table");
    }
 
-   void *new_map = new_bo->ptr.cpu;
+   void *new_map = new_bo->map;
 
    assert(table->bo == NULL && "not yet implemented sparse binding");
    table->bo = new_bo;
diff --git a/src/asahi/vulkan/hk_device.c b/src/asahi/vulkan/hk_device.c
index fb324f26b55..6a002c2d135 100644
--- a/src/asahi/vulkan/hk_device.c
+++ b/src/asahi/vulkan/hk_device.c
@@ -48,14 +48,14 @@ hk_upload_rodata(struct hk_device *dev)
    if (!dev->rodata.bo)
       return VK_ERROR_OUT_OF_HOST_MEMORY;
 
-   uint8_t *map = dev->rodata.bo->ptr.cpu;
+   uint8_t *map = dev->rodata.bo->map;
    uint32_t offs = 0;
 
    offs = align(offs, 8);
    agx_pack(&dev->rodata.txf_sampler, USC_SAMPLER, cfg) {
       cfg.start = 0;
       cfg.count = 1;
-      cfg.buffer = dev->rodata.bo->ptr.gpu + offs;
+      cfg.buffer = dev->rodata.bo->va->addr + offs;
    }
 
    agx_pack(map + offs, SAMPLER, cfg) {
@@ -81,11 +81,11 @@ hk_upload_rodata(struct hk_device *dev)
    agx_pack(&dev->rodata.image_heap, USC_UNIFORM, cfg) {
       cfg.start_halfs = HK_IMAGE_HEAP_UNIFORM;
       cfg.size_halfs = 4;
-      cfg.buffer = dev->rodata.bo->ptr.gpu + offs;
+      cfg.buffer = dev->rodata.bo->va->addr + offs;
    }
 
-   uint64_t *image_heap_ptr = dev->rodata.bo->ptr.cpu + offs;
-   *image_heap_ptr = dev->images.bo->ptr.gpu;
+   uint64_t *image_heap_ptr = dev->rodata.bo->map + offs;
+   *image_heap_ptr = dev->images.bo->va->addr;
    offs += sizeof(uint64_t);
 
    /* The geometry state buffer isn't strictly readonly data, but we only have a
@@ -97,15 +97,15 @@ hk_upload_rodata(struct hk_device *dev)
     * So, we allocate it here for convenience.
     */
    offs = align(offs, sizeof(uint64_t));
-   dev->rodata.geometry_state = dev->rodata.bo->ptr.gpu + offs;
+   dev->rodata.geometry_state = dev->rodata.bo->va->addr + offs;
    offs += sizeof(struct agx_geometry_state);
 
    /* For null readonly buffers, we need to allocate 16 bytes of zeroes for
     * robustness2 semantics on read.
     */
    offs = align(offs, 16);
-   dev->rodata.zero_sink = dev->rodata.bo->ptr.gpu + offs;
-   memset(dev->rodata.bo->ptr.cpu + offs, 0, 16);
+   dev->rodata.zero_sink = dev->rodata.bo->va->addr + offs;
+   memset(dev->rodata.bo->map + offs, 0, 16);
    offs += 16;
 
    /* For null storage descriptors, we need to reserve 16 bytes to catch writes.
@@ -113,7 +113,7 @@ hk_upload_rodata(struct hk_device *dev)
     * without more work.
     */
    offs = align(offs, 16);
-   dev->rodata.null_sink = dev->rodata.bo->ptr.gpu + offs;
+   dev->rodata.null_sink = dev->rodata.bo->va->addr + offs;
    offs += 16;
 
    return VK_SUCCESS;
diff --git a/src/asahi/vulkan/hk_device_memory.c b/src/asahi/vulkan/hk_device_memory.c
index 672ac03c4fe..a0922e80176 100644
--- a/src/asahi/vulkan/hk_device_memory.c
+++ b/src/asahi/vulkan/hk_device_memory.c
@@ -243,7 +243,7 @@ hk_MapMemory2KHR(VkDevice device, const VkMemoryMapInfoKHR *pMemoryMapInfo,
                        "Memory object already mapped.");
    }
 
-   mem->map = mem->bo->ptr.cpu;
+   mem->map = mem->bo->map;
    *ppData = mem->map + offset;
 
    return VK_SUCCESS;
@@ -326,5 +326,5 @@ hk_GetDeviceMemoryOpaqueCaptureAddress(
 {
    VK_FROM_HANDLE(hk_device_memory, mem, pInfo->memory);
 
-   return mem->bo->ptr.gpu;
+   return mem->bo->va->addr;
 }
diff --git a/src/asahi/vulkan/hk_event.c b/src/asahi/vulkan/hk_event.c
index 8626977d123..6057b81a2f3 100644
--- a/src/asahi/vulkan/hk_event.c
+++ b/src/asahi/vulkan/hk_event.c
@@ -32,8 +32,8 @@ hk_CreateEvent(VkDevice device, const VkEventCreateInfo *pCreateInfo,
     */
    event->bo =
       agx_bo_create(&dev->dev, HK_EVENT_MEM_SIZE, 0, AGX_BO_WRITEBACK, "Event");
-   event->status = event->bo->ptr.cpu;
-   event->addr = event->bo->ptr.gpu;
+   event->status = event->bo->map;
+   event->addr = event->bo->va->addr;
 
    *event->status = VK_EVENT_RESET;
 
@@ -91,7 +91,7 @@ hk_CmdSetEvent2(VkCommandBuffer commandBuffer, VkEvent _event,
    VK_FROM_HANDLE(hk_cmd_buffer, cmd, commandBuffer);
    VK_FROM_HANDLE(hk_event, event, _event);
 
-   hk_queue_write(cmd, event->bo->ptr.gpu, VK_EVENT_SET, false);
+   hk_queue_write(cmd, event->bo->va->addr, VK_EVENT_SET, false);
 }
 
 VKAPI_ATTR void VKAPI_CALL
@@ -101,7 +101,7 @@ hk_CmdResetEvent2(VkCommandBuffer commandBuffer, VkEvent _event,
    VK_FROM_HANDLE(hk_cmd_buffer, cmd, commandBuffer);
    VK_FROM_HANDLE(hk_event, event, _event);
 
-   hk_queue_write(cmd, event->bo->ptr.gpu, VK_EVENT_RESET, false);
+   hk_queue_write(cmd, event->bo->va->addr, VK_EVENT_RESET, false);
 }
 
 VKAPI_ATTR void VKAPI_CALL
diff --git a/src/asahi/vulkan/hk_image.c b/src/asahi/vulkan/hk_image.c
index 6187eff40a8..1f203996faf 100644
--- a/src/asahi/vulkan/hk_image.c
+++ b/src/asahi/vulkan/hk_image.c
@@ -1137,8 +1137,8 @@ hk_image_plane_bind(struct hk_device *dev, struct hk_image_plane *plane,
 #endif
       unreachable("todo");
    } else {
-      plane->addr = mem->bo->ptr.gpu + *offset_B;
-      plane->map = mem->bo->ptr.cpu + *offset_B;
+      plane->addr = mem->bo->va->addr + *offset_B;
+      plane->map = mem->bo->map + *offset_B;
       plane->rem = mem->bo->size - (*offset_B);
    }
 
diff --git a/src/asahi/vulkan/hk_query_pool.c b/src/asahi/vulkan/hk_query_pool.c
index b58dcd58500..8a97bb4f30b 100644
--- a/src/asahi/vulkan/hk_query_pool.c
+++ b/src/asahi/vulkan/hk_query_pool.c
@@ -39,7 +39,7 @@ struct hk_query_report {
 static uint16_t *
 hk_pool_oq_index_ptr(const struct hk_query_pool *pool)
 {
-   return (uint16_t *)(pool->bo->ptr.cpu + pool->query_start);
+   return (uint16_t *)(pool->bo->map + pool->query_start);
 }
 
 static uint32_t
@@ -148,14 +148,14 @@ static uint64_t
 hk_query_available_addr(struct hk_query_pool *pool, uint32_t query)
 {
    assert(query < pool->vk.query_count);
-   return pool->bo->ptr.gpu + query * sizeof(uint32_t);
+   return pool->bo->va->addr + query * sizeof(uint32_t);
 }
 
 static uint32_t *
 hk_query_available_map(struct hk_query_pool *pool, uint32_t query)
 {
    assert(query < pool->vk.query_count);
-   return (uint32_t *)pool->bo->ptr.cpu + query;
+   return (uint32_t *)pool->bo->map + query;
 }
 
 static uint64_t
@@ -171,10 +171,10 @@ hk_query_report_addr(struct hk_device *dev, struct hk_query_pool *pool,
 {
    if (pool->oq_queries) {
       uint16_t *oq_index = hk_pool_oq_index_ptr(pool);
-      return dev->occlusion_queries.bo->ptr.gpu +
+      return dev->occlusion_queries.bo->va->addr +
              (oq_index[query] * sizeof(uint64_t));
    } else {
-      return pool->bo->ptr.gpu + hk_query_offset(pool, query);
+      return pool->bo->va->addr + hk_query_offset(pool, query);
    }
 }
 
@@ -183,12 +183,12 @@ hk_query_report_map(struct hk_device *dev, struct hk_query_pool *pool,
                     uint32_t query)
 {
    if (pool->oq_queries) {
-      uint64_t *queries = (uint64_t *)dev->occlusion_queries.bo->ptr.cpu;
+      uint64_t *queries = (uint64_t *)dev->occlusion_queries.bo->map;
       uint16_t *oq_index = hk_pool_oq_index_ptr(pool);
 
       return (struct hk_query_report *)&queries[oq_index[query]];
    } else {
-      return (void *)((char *)pool->bo->ptr.cpu + hk_query_offset(pool, query));
+      return (void *)((char *)pool->bo->map + hk_query_offset(pool, query));
    }
 }
 
@@ -556,10 +556,10 @@ hk_CmdCopyQueryPoolResults(VkCommandBuffer commandBuffer, VkQueryPool queryPool,
    hk_ensure_cs_has_space(cmd, cs, 0x2000 /* TODO */);
 
    const struct libagx_copy_query_push info = {
-      .availability = pool->bo->ptr.gpu,
-      .results = pool->oq_queries ? dev->occlusion_queries.bo->ptr.gpu
-                                  : pool->bo->ptr.gpu + pool->query_start,
-      .oq_index = pool->oq_queries ? pool->bo->ptr.gpu + pool->query_start : 0,
+      .availability = pool->bo->va->addr,
+      .results = pool->oq_queries ? dev->occlusion_queries.bo->va->addr
+                                  : pool->bo->va->addr + pool->query_start,
+      .oq_index = pool->oq_queries ? pool->bo->va->addr + pool->query_start : 0,
 
       .first_query = firstQuery,
       .dst_addr = hk_buffer_address(dst_buffer, dstOffset),
diff --git a/src/asahi/vulkan/hk_queue.c b/src/asahi/vulkan/hk_queue.c
index 49226517554..78733e54700 100644
--- a/src/asahi/vulkan/hk_queue.c
+++ b/src/asahi/vulkan/hk_queue.c
@@ -75,7 +75,7 @@ asahi_fill_cdm_command(struct hk_device *dev, struct hk_cs *cs,
       .encoder_ptr = cs->addr,
       .encoder_end = cs->addr + len,
 
-      .sampler_array = dev->samplers.table.bo->ptr.gpu,
+      .sampler_array = dev->samplers.table.bo->va->addr,
       .sampler_count = dev->samplers.table.alloc,
       .sampler_max = dev->samplers.table.alloc + 1,
 
@@ -87,9 +87,9 @@ asahi_fill_cdm_command(struct hk_device *dev, struct hk_cs *cs,
    };
 
    if (cs->scratch.cs.main || cs->scratch.cs.preamble) {
-      cmd->helper_arg = dev->scratch.cs.buf->ptr.gpu;
+      cmd->helper_arg = dev->scratch.cs.buf->va->addr;
       cmd->helper_cfg = cs->scratch.cs.preamble << 16;
-      cmd->helper_program = dev->dev.helper->ptr.gpu | 1;
+      cmd->helper_program = dev->dev.helper->va->addr | 1;
    }
 }
 
@@ -228,7 +228,7 @@ asahi_fill_vdm_command(struct hk_device *dev, struct hk_cs *cs,
    c->scissor_array = cs->uploaded_scissor;
    c->depth_bias_array = cs->uploaded_zbias;
 
-   c->vertex_sampler_array = dev->samplers.table.bo->ptr.gpu;
+   c->vertex_sampler_array = dev->samplers.table.bo->va->addr;
    c->vertex_sampler_count = dev->samplers.table.alloc;
    c->vertex_sampler_max = dev->samplers.table.alloc + 1;
 
@@ -236,7 +236,7 @@ asahi_fill_vdm_command(struct hk_device *dev, struct hk_cs *cs,
    c->fragment_sampler_count = c->vertex_sampler_count;
    c->fragment_sampler_max = c->vertex_sampler_max;
 
-   c->visibility_result_buffer = dev->occlusion_queries.bo->ptr.gpu;
+   c->visibility_result_buffer = dev->occlusion_queries.bo->va->addr;
 
    /* If a tile is empty, we do not want to process it, as the redundant
     * roundtrip of memory-->tilebuffer-->memory wastes a tremendous amount of
@@ -252,15 +252,15 @@ asahi_fill_vdm_command(struct hk_device *dev, struct hk_cs *cs,
 
    if (cs->scratch.vs.main || cs->scratch.vs.preamble) {
       c->flags |= ASAHI_RENDER_VERTEX_SPILLS;
-      c->vertex_helper_arg = dev->scratch.vs.buf->ptr.gpu;
+      c->vertex_helper_arg = dev->scratch.vs.buf->va->addr;
       c->vertex_helper_cfg = cs->scratch.vs.preamble << 16;
-      c->vertex_helper_program = dev->dev.helper->ptr.gpu | 1;
+      c->vertex_helper_program = dev->dev.helper->va->addr | 1;
    }
 
    if (cs->scratch.fs.main || cs->scratch.fs.preamble) {
-      c->fragment_helper_arg = dev->scratch.fs.buf->ptr.gpu;
+      c->fragment_helper_arg = dev->scratch.fs.buf->va->addr;
       c->fragment_helper_cfg = cs->scratch.fs.preamble << 16;
-      c->fragment_helper_program = dev->dev.helper->ptr.gpu | 1;
+      c->fragment_helper_program = dev->dev.helper->va->addr | 1;
    }
 }
 
@@ -502,7 +502,7 @@ queue_submit(struct hk_device *dev, struct hk_queue *queue,
          }
       }
 
-      agxdecode_image_heap(dev->dev.agxdecode, dev->images.bo->ptr.gpu,
+      agxdecode_image_heap(dev->dev.agxdecode, dev->images.bo->va->addr,
                            dev->images.alloc);
 
       agxdecode_next_frame();
diff --git a/src/asahi/vulkan/hk_shader.c b/src/asahi/vulkan/hk_shader.c
index 044e532470c..99038ddeeb8 100644
--- a/src/asahi/vulkan/hk_shader.c
+++ b/src/asahi/vulkan/hk_shader.c
@@ -673,8 +673,8 @@ hk_upload_shader(struct hk_device *dev, struct hk_shader *shader)
 
       shader->bo = agx_bo_create(&dev->dev, size, 0,
                                  AGX_BO_EXEC | AGX_BO_LOW_VA, "Preamble");
-      memcpy(shader->bo->ptr.cpu, shader->b.binary + offs, size);
-      shader->preamble_addr = shader->bo->ptr.gpu;
+      memcpy(shader->bo->map, shader->b.binary + offs, size);
+      shader->preamble_addr = shader->bo->va->addr;
    }
 
    if (!shader->linked.ht) {
diff --git a/src/gallium/drivers/asahi/agx_batch.c b/src/gallium/drivers/asahi/agx_batch.c
index 5f739a6f7cd..455b0f1c16a 100644
--- a/src/gallium/drivers/asahi/agx_batch.c
+++ b/src/gallium/drivers/asahi/agx_batch.c
@@ -83,8 +83,8 @@ agx_encoder_allocate(struct agx_batch *batch, struct agx_device *dev)
 
    return (struct agx_encoder){
       .bo = bo,
-      .current = bo->ptr.cpu,
-      .end = (uint8_t *)bo->ptr.cpu + bo->size,
+      .current = bo->map,
+      .end = (uint8_t *)bo->map + bo->size,
    };
 }
 
@@ -162,7 +162,7 @@ agx_batch_init(struct agx_context *ctx,
    batch->result_off =
       (2 * sizeof(union agx_batch_result)) * agx_batch_idx(batch);
    batch->result =
-      (void *)(((uint8_t *)ctx->result_buf->ptr.cpu) + batch->result_off);
+      (void *)(((uint8_t *)ctx->result_buf->map) + batch->result_off);
    memset(batch->result, 0, sizeof(union agx_batch_result) * 2);
 
    agx_batch_mark_active(batch);
@@ -837,7 +837,7 @@ agx_batch_submit(struct agx_context *ctx, struct agx_batch *batch,
       struct agx_bo *bo = agx_lookup_bo(dev, handle);
 
       if (bo->flags & AGX_BO_SHARED) {
-         batch_debug(batch, "Waits on shared BO @ 0x%" PRIx64, bo->ptr.gpu);
+         batch_debug(batch, "Waits on shared BO @ 0x%" PRIx64, bo->va->addr);
 
          /* Get a sync file fd from the buffer */
          int in_sync_fd = agx_export_sync_file(dev, bo);
@@ -869,7 +869,7 @@ agx_batch_submit(struct agx_context *ctx, struct agx_batch *batch,
          if (writer && queue_id != ctx->queue_id) {
             batch_debug(
                batch, "Waits on inter-context BO @ 0x%" PRIx64 " from queue %u",
-               bo->ptr.gpu, queue_id);
+               bo->va->addr, queue_id);
 
             agx_add_sync(in_syncs, &in_sync_count,
                          agx_bo_writer_syncobj(writer));
@@ -981,7 +981,7 @@ agx_batch_submit(struct agx_context *ctx, struct agx_batch *batch,
             continue;
 
          batch_debug(batch, "Signals shared BO @ 0x%" PRIx64,
-                     shared_bos[i]->ptr.gpu);
+                     shared_bos[i]->va->addr);
 
          /* Free the in_sync handle we just acquired */
          ret = drmSyncobjDestroy(dev->fd, in_syncs[i].handle);
@@ -1011,7 +1011,7 @@ agx_batch_submit(struct agx_context *ctx, struct agx_batch *batch,
       /* But any BOs written by active batches are ours */
       assert(writer == batch && "exclusive writer");
       p_atomic_set(&bo->writer, agx_bo_writer(ctx->queue_id, batch->syncobj));
-      batch_debug(batch, "Writes to BO @ 0x%" PRIx64, bo->ptr.gpu);
+      batch_debug(batch, "Writes to BO @ 0x%" PRIx64, bo->va->addr);
    }
 
    free(in_syncs);
diff --git a/src/gallium/drivers/asahi/agx_disk_cache.c b/src/gallium/drivers/asahi/agx_disk_cache.c
index e95a4414727..906daa01ce6 100644
--- a/src/gallium/drivers/asahi/agx_disk_cache.c
+++ b/src/gallium/drivers/asahi/agx_disk_cache.c
@@ -109,12 +109,12 @@ read_shader(struct agx_screen *screen, struct blob_reader *blob,
       if (size) {
          binary->bo = agx_bo_create(&screen->dev, size, 0,
                                     AGX_BO_EXEC | AGX_BO_LOW_VA, "Executable");
-         memcpy(binary->bo->ptr.cpu, binary->b.binary, size);
+         memcpy(binary->bo->map, binary->b.binary, size);
       }
    } else if (size) {
       binary->bo = agx_bo_create(&screen->dev, size, 0,
                                  AGX_BO_EXEC | AGX_BO_LOW_VA, "Executable");
-      blob_copy_bytes(blob, binary->bo->ptr.cpu, size);
+      blob_copy_bytes(blob, binary->bo->map, size);
    }
 
    blob_copy_bytes(blob, &binary->b.info, sizeof(binary->b.info));
diff --git a/src/gallium/drivers/asahi/agx_pipe.c b/src/gallium/drivers/asahi/agx_pipe.c
index e66f48fc973..a6ce72b0f81 100644
--- a/src/gallium/drivers/asahi/agx_pipe.c
+++ b/src/gallium/drivers/asahi/agx_pipe.c
@@ -139,9 +139,9 @@ agx_resource_debug(struct agx_resource *res, const char *msg)
       (long long)res->layout.linear_stride_B,
       (long long)res->layout.layer_stride_B,
       (long long)res->layout.compression_layer_stride_B,
-      (long long)res->bo->ptr.gpu, (long long)res->layout.size_B,
+      (long long)res->bo->va->addr, (long long)res->layout.size_B,
       res->layout.metadata_offset_B
-         ? ((long long)res->bo->ptr.gpu + res->layout.metadata_offset_B)
+         ? ((long long)res->bo->va->addr + res->layout.metadata_offset_B)
          : 0,
       (long long)res->layout.metadata_offset_B, res->bo->label,
       res->bo->flags & AGX_BO_SHARED ? "SH " : "",
@@ -713,7 +713,7 @@ agx_shadow(struct agx_context *ctx, struct agx_resource *rsrc, bool needs_copy)
                      (old->flags & AGX_BO_WRITEBACK) ? "cached" : "uncached");
       agx_resource_debug(rsrc, "Shadowed: ");
 
-      memcpy(new_->ptr.cpu, old->ptr.cpu, size);
+      memcpy(new_->map, old->map, size);
    }
 
    /* Swap the pointers, dropping a reference */
@@ -1001,7 +1001,7 @@ agx_transfer_map(struct pipe_context *pctx, struct pipe_resource *resource,
       }
 
       dev->ops.bo_mmap(dev, staging->bo);
-      return staging->bo->ptr.cpu;
+      return staging->bo->map;
    }
 
    dev->ops.bo_mmap(dev, rsrc->bo);
@@ -1046,7 +1046,7 @@ agx_transfer_map(struct pipe_context *pctx, struct pipe_resource *resource,
       uint32_t offset =
          ail_get_linear_pixel_B(&rsrc->layout, level, box->x, box->y, box->z);
 
-      return ((uint8_t *)rsrc->bo->ptr.cpu) + offset;
+      return ((uint8_t *)rsrc->bo->map) + offset;
    }
 }
 
@@ -1245,7 +1245,7 @@ asahi_add_attachment(struct attachments *att, struct agx_resource *rsrc,
    int idx = att->count++;
 
    att->list[idx].size = rsrc->layout.size_B;
-   att->list[idx].pointer = rsrc->bo->ptr.gpu;
+   att->list[idx].pointer = rsrc->bo->va->addr;
    att->list[idx].order = 1; // TODO: What does this do?
    att->list[idx].flags = 0;
 }
@@ -1495,7 +1495,7 @@ agx_cmdbuf(struct agx_device *dev, struct drm_asahi_cmd_render *c,
    c->visibility_result_buffer = visibility_result_ptr;
 
    c->vertex_sampler_array =
-      batch->sampler_heap.bo ? batch->sampler_heap.bo->ptr.gpu : 0;
+      batch->sampler_heap.bo ? batch->sampler_heap.bo->va->addr : 0;
    c->vertex_sampler_count = batch->sampler_heap.count;
    c->vertex_sampler_max = batch->sampler_heap.count + 1;
 
@@ -1539,14 +1539,14 @@ agx_cmdbuf(struct agx_device *dev, struct drm_asahi_cmd_render *c,
 
    if (batch->vs_scratch) {
       c->flags |= ASAHI_RENDER_VERTEX_SPILLS;
-      c->vertex_helper_arg = batch->ctx->scratch_vs.buf->ptr.gpu;
+      c->vertex_helper_arg = batch->ctx->scratch_vs.buf->va->addr;
       c->vertex_helper_cfg = batch->vs_preamble_scratch << 16;
-      c->vertex_helper_program = dev->helper->ptr.gpu | 1;
+      c->vertex_helper_program = dev->helper->va->addr | 1;
    }
    if (batch->fs_scratch) {
-      c->fragment_helper_arg = batch->ctx->scratch_fs.buf->ptr.gpu;
+      c->fragment_helper_arg = batch->ctx->scratch_fs.buf->va->addr;
       c->fragment_helper_cfg = batch->fs_preamble_scratch << 16;
-      c->fragment_helper_program = dev->helper->ptr.gpu | 1;
+      c->fragment_helper_program = dev->helper->va->addr | 1;
    }
 }
 
@@ -1642,16 +1642,16 @@ agx_flush_compute(struct agx_context *ctx, struct agx_batch *batch,
 
    *cmdbuf = (struct drm_asahi_cmd_compute){
       .flags = 0,
-      .encoder_ptr = batch->cdm.bo->ptr.gpu,
-      .encoder_end = batch->cdm.bo->ptr.gpu +
-                     (batch->cdm.current - (uint8_t *)batch->cdm.bo->ptr.cpu),
+      .encoder_ptr = batch->cdm.bo->va->addr,
+      .encoder_end = batch->cdm.bo->va->addr +
+                     (batch->cdm.current - (uint8_t *)batch->cdm.bo->map),
       .usc_base = dev->shader_base,
       .helper_arg = 0,
       .helper_cfg = 0,
       .helper_program = 0,
       .iogpu_unk_40 = 0,
       .sampler_array =
-         batch->sampler_heap.bo ? batch->sampler_heap.bo->ptr.gpu : 0,
+         batch->sampler_heap.bo ? batch->sampler_heap.bo->va->addr : 0,
       .sampler_count = batch->sampler_heap.count,
       .sampler_max = batch->sampler_heap.count + 1,
       .encoder_id = encoder_id,
@@ -1665,10 +1665,10 @@ agx_flush_compute(struct agx_context *ctx, struct agx_batch *batch,
       // helper. Disable them for now.
 
       // cmdbuf->iogpu_unk_40 = 0x1c;
-      cmdbuf->helper_arg = ctx->scratch_cs.buf->ptr.gpu;
+      cmdbuf->helper_arg = ctx->scratch_cs.buf->va->addr;
       cmdbuf->helper_cfg = batch->cs_preamble_scratch << 16;
       // cmdbuf->helper_cfg |= 0x40;
-      cmdbuf->helper_program = dev->helper->ptr.gpu | 1;
+      cmdbuf->helper_program = dev->helper->va->addr | 1;
    }
 }
 
@@ -1729,9 +1729,9 @@ agx_flush_render(struct agx_context *ctx, struct agx_batch *batch,
    unsigned encoder_id = agx_get_global_id(dev);
 
    agx_cmdbuf(dev, cmdbuf, att, &batch->pool, batch, &batch->key,
-              batch->vdm.bo->ptr.gpu, encoder_id, cmd_ta_id, cmd_3d_id, scissor,
-              zbias, agx_get_occlusion_heap(batch), pipeline_background,
-              pipeline_background_partial, pipeline_store,
+              batch->vdm.bo->va->addr, encoder_id, cmd_ta_id, cmd_3d_id,
+              scissor, zbias, agx_get_occlusion_heap(batch),
+              pipeline_background, pipeline_background_partial, pipeline_store,
               clear_pipeline_textures, batch->clear_depth, batch->clear_stencil,
               &batch->tilebuffer_layout);
 }
diff --git a/src/gallium/drivers/asahi/agx_query.c b/src/gallium/drivers/asahi/agx_query.c
index 12fcd7c6339..200aef54659 100644
--- a/src/gallium/drivers/asahi/agx_query.c
+++ b/src/gallium/drivers/asahi/agx_query.c
@@ -105,8 +105,8 @@ agx_alloc_oq(struct agx_context *ctx)
    unsigned offset = index * sizeof(uint64_t);
 
    return (struct agx_ptr){
-      (uint8_t *)heap->bo->ptr.cpu + offset,
-      heap->bo->ptr.gpu + offset,
+      (uint8_t *)heap->bo->map + offset,
+      heap->bo->va->addr + offset,
    };
 }
 
@@ -115,7 +115,7 @@ agx_oq_index(struct agx_context *ctx, struct agx_query *q)
 {
    assert(is_occlusion(q));
 
-   return (q->ptr.gpu - ctx->oq->bo->ptr.gpu) / sizeof(uint64_t);
+   return (q->ptr.gpu - ctx->oq->bo->va->addr) / sizeof(uint64_t);
 }
 
 static void
@@ -139,7 +139,7 @@ agx_get_occlusion_heap(struct agx_batch *batch)
    struct agx_bo *bo = batch->ctx->oq->bo;
 
    if (agx_batch_uses_bo(batch, bo))
-      return bo->ptr.gpu;
+      return bo->va->addr;
    else
       return 0;
 }
@@ -167,7 +167,10 @@ agx_create_query(struct pipe_context *ctx, unsigned query_type, unsigned index)
        */
       query->bo = agx_bo_create(agx_device(ctx->screen), sizeof(uint64_t) * 2,
                                 0, AGX_BO_WRITEBACK, "Query");
-      query->ptr = query->bo->ptr;
+      query->ptr = (struct agx_ptr){
+         .gpu = query->bo->va->addr,
+         .cpu = query->bo->map,
+      };
    }
 
    if (!query->ptr.gpu) {
@@ -546,7 +549,7 @@ agx_get_query_result_resource_gpu(struct agx_context *ctx,
    memcpy(&saved_cb, &stage->cb[0], sizeof(struct pipe_constant_buffer));
 
    /* Set params */
-   uint64_t params[2] = {query->ptr.gpu, rsrc->bo->ptr.gpu + offset};
+   uint64_t params[2] = {query->ptr.gpu, rsrc->bo->va->addr + offset};
    agx_batch_writes_range(batch, rsrc, offset, result_type_size(result_type));
 
    struct pipe_constant_buffer cb = {
diff --git a/src/gallium/drivers/asahi/agx_state.c b/src/gallium/drivers/asahi/agx_state.c
index 16fe388024d..f207397f690 100644
--- a/src/gallium/drivers/asahi/agx_state.c
+++ b/src/gallium/drivers/asahi/agx_state.c
@@ -1593,8 +1593,7 @@ agx_compile_nir(struct agx_device *dev, nir_shader *nir,
       compiled->bo = agx_bo_create(dev, compiled->b.binary_size, 0,
                                    AGX_BO_EXEC | AGX_BO_LOW_VA, "Executable");
 
-      memcpy(compiled->bo->ptr.cpu, compiled->b.binary,
-             compiled->b.binary_size);
+      memcpy(compiled->bo->map, compiled->b.binary, compiled->b.binary_size);
    }
 
    return compiled;
@@ -2808,7 +2807,7 @@ agx_sampler_heap_add(struct agx_device *dev, struct agx_sampler_heap *heap,
 
    /* Precondition: there is room in the heap */
    assert(heap->count < AGX_SAMPLER_HEAP_SIZE);
-   struct agx_sampler_packed *samplers = heap->bo->ptr.cpu;
+   struct agx_sampler_packed *samplers = heap->bo->map;
    memcpy(samplers + heap->count, sampler, sizeof(*sampler));
 
    return heap->count++;
@@ -3031,7 +3030,8 @@ agx_build_pipeline(struct agx_batch *batch, struct agx_compiled_shader *cs,
          agx_usc_push_packed(&b, FRAGMENT_PROPERTIES, linked->fragment_props);
    } else {
       agx_usc_pack(&b, SHADER, cfg) {
-         cfg.code = agx_usc_addr(dev, cs->bo->ptr.gpu + cs->b.info.main_offset);
+         cfg.code =
+            agx_usc_addr(dev, cs->bo->va->addr + cs->b.info.main_offset);
          cfg.unk_2 = 3;
       }
 
@@ -3046,7 +3046,7 @@ agx_build_pipeline(struct agx_batch *batch, struct agx_compiled_shader *cs,
    if (cs->b.info.has_preamble) {
       agx_usc_pack(&b, PRESHADER, cfg) {
          cfg.code =
-            agx_usc_addr(dev, cs->bo->ptr.gpu + cs->b.info.preamble_offset);
+            agx_usc_addr(dev, cs->bo->va->addr + cs->b.info.preamble_offset);
       }
    } else {
       agx_usc_pack(&b, NO_PRESHADER, cfg)
@@ -3083,7 +3083,7 @@ agx_build_internal_usc(struct agx_batch *batch, struct agx_compiled_shader *cs,
    }
 
    agx_usc_pack(&b, SHADER, cfg) {
-      cfg.code = agx_usc_addr(dev, cs->bo->ptr.gpu + cs->b.info.main_offset);
+      cfg.code = agx_usc_addr(dev, cs->bo->va->addr + cs->b.info.main_offset);
       cfg.unk_2 = 3;
    }
 
@@ -3095,7 +3095,7 @@ agx_build_internal_usc(struct agx_batch *batch, struct agx_compiled_shader *cs,
    if (cs->b.info.has_preamble) {
       agx_usc_pack(&b, PRESHADER, cfg) {
          cfg.code =
-            agx_usc_addr(dev, cs->bo->ptr.gpu + cs->b.info.preamble_offset);
+            agx_usc_addr(dev, cs->bo->va->addr + cs->b.info.preamble_offset);
       }
    } else {
       agx_usc_pack(&b, NO_PRESHADER, cfg)
@@ -3804,7 +3804,7 @@ agx_index_buffer_rsrc_ptr(struct agx_batch *batch,
    agx_batch_reads(batch, rsrc);
 
    *extent = ALIGN_POT(rsrc->layout.size_B, 4);
-   return rsrc->bo->ptr.gpu;
+   return rsrc->bo->va->addr;
 }
 
 static uint64_t
@@ -3938,7 +3938,7 @@ agx_batch_geometry_state(struct agx_batch *batch)
       }
 
       struct agx_geometry_state state = {
-         .heap = agx_resource(ctx->heap)->bo->ptr.gpu,
+         .heap = agx_resource(ctx->heap)->bo->va->addr,
          .heap_size = size,
       };
 
@@ -3985,7 +3985,7 @@ agx_batch_geometry_params(struct agx_batch *batch, uint64_t input_index_buffer,
       params.xfb_size[i] = size;
 
       if (rsrc) {
-         params.xfb_offs_ptrs[i] = rsrc->bo->ptr.gpu;
+         params.xfb_offs_ptrs[i] = rsrc->bo->va->addr;
          agx_batch_writes(batch, rsrc, 0);
          batch->incoherent_writes = true;
       } else {
@@ -4065,7 +4065,7 @@ agx_indirect_buffer_ptr(struct agx_batch *batch,
 
    struct agx_resource *rsrc = agx_resource(indirect->buffer);
    agx_batch_reads(batch, rsrc);
-   return rsrc->bo->ptr.gpu + indirect->offset;
+   return rsrc->bo->va->addr + indirect->offset;
 }
 
 static void
@@ -4236,7 +4236,7 @@ agx_draw_without_restart(struct agx_batch *batch,
       .restart_index = info->restart_index,
       .index_buffer_size_el = ib_extent / info->index_size,
       .flatshade_first = batch->ctx->rast->base.flatshade_first,
-      .draws = indirect_rsrc->bo->ptr.gpu + indirect->offset,
+      .draws = indirect_rsrc->bo->va->addr + indirect->offset,
    };
 
    /* Unroll the index buffer for each draw */
@@ -4258,7 +4258,7 @@ agx_draw_without_restart(struct agx_batch *batch,
 
    struct pipe_draw_indirect_info new_indirect = *indirect;
    new_indirect.buffer = &out_draws_rsrc.base;
-   new_indirect.offset = out_draws.gpu - out_draws_rsrc.bo->ptr.gpu;
+   new_indirect.offset = out_draws.gpu - out_draws_rsrc.bo->va->addr;
    new_indirect.stride = 5 * sizeof(uint32_t);
 
    ctx->active_draw_without_restart = true;
@@ -4455,7 +4455,7 @@ agx_upload_draw_params(struct agx_batch *batch,
 {
    if (indirect) {
       struct agx_resource *indirect_rsrc = agx_resource(indirect->buffer);
-      uint64_t address = indirect_rsrc->bo->ptr.gpu + indirect->offset;
+      uint64_t address = indirect_rsrc->bo->va->addr + indirect->offset;
       agx_batch_reads(batch, indirect_rsrc);
 
       /* To implement draw parameters, we use the last 2 words of the
@@ -4601,7 +4601,7 @@ agx_draw_patches(struct agx_context *ctx, const struct pipe_draw_info *info,
       .output_patch_size = tcs->tess.output_patch_size,
       .tcs_patch_constants = tcs->tess.nr_patch_outputs,
       .tcs_per_vertex_outputs = tcs->tess.per_vertex_outputs,
-      .patch_coord_buffer = agx_resource(ctx->heap)->bo->ptr.gpu,
+      .patch_coord_buffer = agx_resource(ctx->heap)->bo->va->addr,
    };
 
    memcpy(&args.tess_level_outer_default, ctx->default_outer_level,
@@ -4791,7 +4791,7 @@ agx_draw_patches(struct agx_context *ctx, const struct pipe_draw_info *info,
 
    struct pipe_draw_indirect_info copy_indirect = {
       .buffer = &indirect_rsrc.base,
-      .offset = args.out_draws - draw_bo->ptr.gpu,
+      .offset = args.out_draws - draw_bo->va->addr,
       .stride = draw_stride,
       .draw_count = 1,
    };
@@ -5123,7 +5123,7 @@ agx_draw_vbo(struct pipe_context *pctx, const struct pipe_draw_info *info,
       indirect_gs = (struct pipe_draw_indirect_info){
          .draw_count = 1,
          .buffer = &indirect_rsrc.base,
-         .offset = batch->geom_indirect - indirect_rsrc.bo->ptr.gpu,
+         .offset = batch->geom_indirect - indirect_rsrc.bo->va->addr,
       };
 
       info = &info_gs;
@@ -5134,7 +5134,7 @@ agx_draw_vbo(struct pipe_context *pctx, const struct pipe_draw_info *info,
       ctx->dirty |= AGX_DIRTY_PRIM;
 
       if (info_gs.index_size) {
-         ib = agx_resource(ctx->heap)->bo->ptr.gpu;
+         ib = agx_resource(ctx->heap)->bo->va->addr;
          ib_extent = agx_resource(ctx->heap)->bo->size;
       } else {
          ib = 0;
@@ -5182,7 +5182,7 @@ agx_draw_vbo(struct pipe_context *pctx, const struct pipe_draw_info *info,
 
    if (ctx->in_generated_vdm) {
       struct agx_resource *indirect_rsrc = agx_resource(indirect->buffer);
-      uint64_t address = indirect_rsrc->bo->ptr.gpu + indirect->offset;
+      uint64_t address = indirect_rsrc->bo->va->addr + indirect->offset;
 
       agx_push(out, VDM_STREAM_LINK, cfg) {
          cfg.target_lo = address & BITFIELD_MASK(32);
@@ -5227,7 +5227,7 @@ agx_draw_vbo(struct pipe_context *pctx, const struct pipe_draw_info *info,
 
       if (indirect) {
          struct agx_resource *indirect_rsrc = agx_resource(indirect->buffer);
-         uint64_t address = indirect_rsrc->bo->ptr.gpu + indirect->offset;
+         uint64_t address = indirect_rsrc->bo->va->addr + indirect->offset;
 
          agx_push(out, INDEX_LIST_INDIRECT_BUFFER, cfg) {
             cfg.address_hi = address >> 32;
@@ -5460,7 +5460,7 @@ agx_launch_grid(struct pipe_context *pipe, const struct pipe_grid_info *info)
    if (info->indirect) {
       struct agx_resource *rsrc = agx_resource(info->indirect);
       agx_batch_reads(batch, rsrc);
-      indirect = rsrc->bo->ptr.gpu + info->indirect_offset;
+      indirect = rsrc->bo->va->addr + info->indirect_offset;
    }
 
    /* Increment the pipeline stats query.
@@ -5511,7 +5511,7 @@ agx_launch_grid(struct pipe_context *pipe, const struct pipe_grid_info *info)
       agx_batch_reads(batch, indirect);
 
       grid.mode = AGX_CDM_MODE_INDIRECT_GLOBAL;
-      grid.indirect = indirect->bo->ptr.gpu + info->indirect_offset;
+      grid.indirect = indirect->bo->va->addr + info->indirect_offset;
    } else {
       grid.mode = AGX_CDM_MODE_DIRECT;
 
@@ -5576,7 +5576,7 @@ agx_set_global_binding(struct pipe_context *pipe, unsigned first,
          struct agx_resource *rsrc = agx_resource(resources[i]);
 
          memcpy(&addr, handles[i], sizeof(addr));
-         addr += rsrc->bo->ptr.gpu;
+         addr += rsrc->bo->va->addr;
          memcpy(handles[i], &addr, sizeof(addr));
       } else {
          pipe_resource_reference(res, NULL);
diff --git a/src/gallium/drivers/asahi/agx_state.h b/src/gallium/drivers/asahi/agx_state.h
index 8240f794b14..496a50b1b2b 100644
--- a/src/gallium/drivers/asahi/agx_state.h
+++ b/src/gallium/drivers/asahi/agx_state.h
@@ -1019,14 +1019,14 @@ agx_resource_valid(struct agx_resource *rsrc, int level)
 static inline void *
 agx_map_texture_cpu(struct agx_resource *rsrc, unsigned level, unsigned z)
 {
-   return ((uint8_t *)rsrc->bo->ptr.cpu) +
+   return ((uint8_t *)rsrc->bo->map) +
           ail_get_layer_level_B(&rsrc->layout, z, level);
 }
 
 static inline uint64_t
 agx_map_texture_gpu(struct agx_resource *rsrc, unsigned z)
 {
-   return rsrc->bo->ptr.gpu +
+   return rsrc->bo->va->addr +
           (uint64_t)ail_get_layer_offset_B(&rsrc->layout, z);
 }
 
diff --git a/src/gallium/drivers/asahi/agx_streamout.c b/src/gallium/drivers/asahi/agx_streamout.c
index 0204430fe4d..f5278c775bb 100644
--- a/src/gallium/drivers/asahi/agx_streamout.c
+++ b/src/gallium/drivers/asahi/agx_streamout.c
@@ -115,7 +115,7 @@ agx_batch_get_so_address(struct agx_batch *batch, unsigned buffer,
                           target->buffer_size);
 
    *size = target->buffer_size;
-   return rsrc->bo->ptr.gpu + target->buffer_offset;
+   return rsrc->bo->va->addr + target->buffer_offset;
 }
 
 void
diff --git a/src/gallium/drivers/asahi/agx_uniforms.c b/src/gallium/drivers/asahi/agx_uniforms.c
index da164c629e0..7a056734ad0 100644
--- a/src/gallium/drivers/asahi/agx_uniforms.c
+++ b/src/gallium/drivers/asahi/agx_uniforms.c
@@ -18,7 +18,7 @@ agx_const_buffer_ptr(struct agx_batch *batch, struct pipe_constant_buffer *cb)
       struct agx_resource *rsrc = agx_resource(cb->buffer);
       agx_batch_reads(batch, rsrc);
 
-      return rsrc->bo->ptr.gpu + cb->buffer_offset;
+      return rsrc->bo->va->addr + cb->buffer_offset;
    } else {
       return 0;
    }
@@ -43,7 +43,7 @@ agx_upload_vbos(struct agx_batch *batch)
          struct agx_resource *rsrc = agx_resource(vb.buffer.resource);
          agx_batch_reads(batch, rsrc);
 
-         buffers[vbo] = rsrc->bo->ptr.gpu + vb.buffer_offset;
+         buffers[vbo] = rsrc->bo->va->addr + vb.buffer_offset;
          buf_sizes[vbo] = rsrc->layout.size_B - vb.buffer_offset;
       }
    }
@@ -148,7 +148,7 @@ agx_set_ssbo_uniforms(struct agx_batch *batch, enum pipe_shader_type stage)
             agx_batch_reads(batch, rsrc);
          }
 
-         unif->ssbo_base[cb] = rsrc->bo->ptr.gpu + sb->buffer_offset;
+         unif->ssbo_base[cb] = rsrc->bo->va->addr + sb->buffer_offset;
          unif->ssbo_size[cb] = st->ssbo[cb].buffer_size;
       } else {
          /* Invalid, so use the sink */