AlexIndustrial/mesa
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

asahi: introduce agx_va data structure

Browse Source 
prep for sparse.

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/30633>
This commit is contained in:
Alyssa Rosenzweig
2024-08-01 10:43:22 -04:00
parent 5c7dd392d3
commit f4cca9d600
32 changed files with 309 additions and 259 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/asahi/lib/agx_bo.h
+16 -1
View File
@@ -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;
src/asahi/lib/agx_device.c
+24 -46
View File
@@ -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);
src/asahi/lib/agx_device.h
+6
View File
@@ -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);
src/asahi/lib/agx_device_virtio.c
+11 -16
View File
@@ -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);
}
}
src/asahi/lib/agx_linker.c
+6 -6
View File
@@ -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)
src/asahi/lib/agx_scratch.c
+10 -11
View File
@@ -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,
src/asahi/lib/agx_va.c
+64
View File
@@ -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);
}
src/asahi/lib/decode.c
+10 -10
View File
@@ -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
src/asahi/lib/meson.build
+1
View File
@@ -27,6 +27,7 @@ libasahi_lib_files = files(
'agx_nir_prolog_epilog.c',
'agx_ppp.h',
'agx_scratch.c',
'agx_va.c',
'pool.c',
)
src/asahi/lib/pool.c
+2 -2
View File
@@ -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)
src/asahi/vulkan/hk_buffer.c
+24 -38
View File
@@ -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 =
src/asahi/vulkan/hk_buffer.h
+2 -2
View File
@@ -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,
src/asahi/vulkan/hk_cmd_buffer.c
+8 -5
View File
@@ -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;
src/asahi/vulkan/hk_cmd_draw.c
+9 -9
View File
@@ -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,
};
src/asahi/vulkan/hk_cmd_pool.c
+1 -1
View File
@@ -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;
src/asahi/vulkan/hk_descriptor_set.c
+7 -7
View File
@@ -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);
}
src/asahi/vulkan/hk_descriptor_table.c
+1 -1
View File
@@ -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;
src/asahi/vulkan/hk_device.c
+9 -9
View File
@@ -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;
src/asahi/vulkan/hk_device_memory.c
+2 -2
View File
@@ -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;
}
src/asahi/vulkan/hk_event.c
+4 -4
View File
@@ -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
src/asahi/vulkan/hk_image.c
+2 -2
View File
@@ -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);
}
src/asahi/vulkan/hk_query_pool.c
+11 -11
View File
@@ -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),
src/asahi/vulkan/hk_queue.c
+10 -10
View File
@@ -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();
src/asahi/vulkan/hk_shader.c
+2 -2
View File
@@ -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) {
src/gallium/drivers/asahi/agx_batch.c
+7 -7
View File
@@ -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);
src/gallium/drivers/asahi/agx_disk_cache.c
+2 -2
View File
@@ -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));
src/gallium/drivers/asahi/agx_pipe.c
+20 -20
View File
@@ -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);
}
src/gallium/drivers/asahi/agx_query.c
+9 -6
View File
@@ -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 = {
src/gallium/drivers/asahi/agx_state.c
+23 -23
View File
@@ -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);
src/gallium/drivers/asahi/agx_state.h
+2 -2
View File
@@ -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);
}
src/gallium/drivers/asahi/agx_streamout.c
+1 -1
View File
@@ -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
src/gallium/drivers/asahi/agx_uniforms.c
+3 -3
View File
@@ -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 */