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

kk: Move all resource tracking to the residency set

Browse Source 
Removes encoder's use resource utilities. All memory allocations
are now tracked in the VkDevice level residency set. This is
accomplished by tracking buffer objects at create/destroy.
Also removes all descriptor set residency tracking since it is
no longer needed.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38505>
This commit is contained in:
Aitor Camacho
2025-11-09 16:44:21 +09:00
committed by Marge Bot
parent 4779f3d0de
commit a8a3863313
20 changed files with 46 additions and 235 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
src/kosmickrisp/bridge/mtl_encoder.h
View File

@@ -50,13 +50,6 @@ void mtl_compute_set_pipeline_state(mtl_compute_encoder *encoder,
void mtl_compute_set_buffer(mtl_compute_encoder *encoder, mtl_buffer *buffer,
size_t offset, size_t index);
void mtl_compute_use_resource(mtl_compute_encoder *encoder,
mtl_resource *res_handle, uint32_t usage);
void mtl_compute_use_resources(mtl_compute_encoder *encoder,
mtl_resource **resource_handles, uint32_t count,
enum mtl_resource_usage usage);
void mtl_compute_use_heaps(mtl_compute_encoder *encoder, mtl_heap **heaps,
uint32_t count);

23
src/kosmickrisp/bridge/mtl_encoder.m
View File

@@ -191,29 +191,6 @@ mtl_compute_set_buffer(mtl_compute_encoder *encoder,
}
}
void
mtl_compute_use_resource(mtl_compute_encoder *encoder,
mtl_resource *res_handle, uint32_t usage)
{
@autoreleasepool {
id<MTLComputeCommandEncoder> enc = (id<MTLComputeCommandEncoder>)encoder;
id<MTLResource> res = (id<MTLResource>)res_handle;
[enc useResource:res usage:(MTLResourceUsage)usage];
}
}
void
mtl_compute_use_resources(mtl_compute_encoder *encoder,
mtl_resource **resource_handles, uint32_t count,
enum mtl_resource_usage usage)
{
@autoreleasepool {
id<MTLComputeCommandEncoder> enc = (id<MTLComputeCommandEncoder>)encoder;
id<MTLResource> *handles = (id<MTLResource>*)resource_handles;
[enc useResources:handles count:count usage:(MTLResourceUsage)usage];
}
}
void
mtl_compute_use_heaps(mtl_compute_encoder *encoder, mtl_heap **heaps,
uint32_t count)

13
src/kosmickrisp/bridge/stubs/mtl_encoder.c
View File

@@ -88,19 +88,6 @@ mtl_compute_set_buffer(mtl_compute_encoder *encoder, mtl_buffer *buffer,
{
}
void
mtl_compute_use_resource(mtl_compute_encoder *encoder, mtl_resource *res_handle,
uint32_t usage)
{
}
void
mtl_compute_use_resources(mtl_compute_encoder *encoder,
mtl_resource **resource_handles, uint32_t count,
enum mtl_resource_usage usage)
{
}
void
mtl_compute_use_heaps(mtl_compute_encoder *encoder, mtl_heap **heaps,
uint32_t count)

3
src/kosmickrisp/vulkan/kk_bo.c
View File

@@ -50,6 +50,8 @@ kk_alloc_bo(struct kk_device *dev, struct vk_object_base *log_obj,
bo->gpu = mtl_buffer_get_gpu_address(map);
bo->cpu = mtl_get_contents(map);
kk_device_add_heap_to_residency_set(dev, handle);
*bo_out = bo;
return result;
@@ -64,6 +66,7 @@ fail_heap:
void
kk_destroy_bo(struct kk_device *dev, struct kk_bo *bo)
{
kk_device_remove_heap_from_residency_set(dev, bo->mtl_handle);
mtl_release(bo->map);
mtl_release(bo->mtl_handle);
FREE(bo);

76
src/kosmickrisp/vulkan/kk_cmd_buffer.c
View File

@@ -29,7 +29,6 @@ kk_descriptor_state_fini(struct kk_cmd_buffer *cmd,
desc->push[i] = NULL;
desc->sets[i] = NULL; /* We also need to set sets to NULL so state doesn't
propagate if we reset it */
desc->sets_not_resident = 0u;
}
}
@@ -204,8 +203,6 @@ kk_bind_descriptor_sets(struct kk_descriptor_state *desc,
}
desc->sets[s] = set;
desc->sets_not_resident |= BITFIELD_BIT(s);
/* Binding descriptors invalidates push descriptors */
desc->push_dirty &= ~BITFIELD_BIT(s);
}
@@ -262,11 +259,9 @@ kk_cmd_push_descriptors(struct kk_cmd_buffer *cmd,
struct kk_descriptor_set_layout *set_layout,
uint32_t set)
{
struct kk_device *dev = kk_cmd_buffer_device(cmd);
assert(set < KK_MAX_SETS);
if (unlikely(desc->push[set] == NULL)) {
size_t size = sizeof(*desc->push[set]) +
(sizeof(mtl_resource *) * set_layout->descriptor_count);
size_t size = sizeof(*desc->push[set]);
desc->push[set] = vk_zalloc(&cmd->vk.pool->alloc, size, 8,
VK_SYSTEM_ALLOCATION_SCOPE_OBJECT);
if (unlikely(desc->push[set] == NULL)) {
@@ -274,14 +269,11 @@ kk_cmd_push_descriptors(struct kk_cmd_buffer *cmd,
return NULL;
}
desc->push[set]->layout = set_layout;
for (uint32_t i = 0u; i < set_layout->descriptor_count; ++i)
desc->push[set]->mtl_resources[i] = dev->null_descriptor->map;
}
/* Pushing descriptors replaces whatever sets are bound */
desc->sets[set] = NULL;
desc->push_dirty |= BITFIELD_BIT(set);
desc->sets_not_resident |= BITFIELD_BIT(set);
return desc->push[set];
}
@@ -425,7 +417,6 @@ kk_cmd_buffer_flush_push_descriptors(struct kk_cmd_buffer *cmd,
return;
memcpy(bo->cpu, push_set->data, sizeof(push_set->data));
push_set->mtl_descriptor_buffer = bo->map;
desc->root.sets[set_idx] = bo->gpu;
desc->set_sizes[set_idx] = sizeof(push_set->data);
}
@@ -434,77 +425,12 @@ kk_cmd_buffer_flush_push_descriptors(struct kk_cmd_buffer *cmd,
desc->push_dirty = 0;
}
static void
kk_make_graphics_descriptor_resources_resident(struct kk_cmd_buffer *cmd)
{
struct kk_descriptor_state *desc = &cmd->state.gfx.descriptors;
mtl_render_encoder *encoder = kk_render_encoder(cmd);
/* Make resources resident as required by Metal */
u_foreach_bit(set_index, desc->sets_not_resident) {
mtl_resource *descriptor_buffer = NULL;
/* If we have no set, it means it was a push set */
if (desc->sets[set_index]) {
struct kk_descriptor_set *set = desc->sets[set_index];
descriptor_buffer = set->mtl_descriptor_buffer;
} else {
struct kk_push_descriptor_set *push_set = desc->push[set_index];
descriptor_buffer = push_set->mtl_descriptor_buffer;
}
/* We could have empty descriptor sets for some reason... */
if (descriptor_buffer) {
mtl_render_use_resource(encoder, descriptor_buffer,
MTL_RESOURCE_USAGE_READ);
}
}
desc->sets_not_resident = 0u;
}
static void
kk_make_compute_descriptor_resources_resident(struct kk_cmd_buffer *cmd)
{
struct kk_descriptor_state *desc = &cmd->state.cs.descriptors;
mtl_compute_encoder *encoder = kk_compute_encoder(cmd);
u_foreach_bit(set_index, desc->sets_not_resident) {
/* Make resources resident as required by Metal */
mtl_resource *descriptor_buffer = NULL;
if (desc->sets[set_index]) {
struct kk_descriptor_set *set = desc->sets[set_index];
descriptor_buffer = set->mtl_descriptor_buffer;
} else {
struct kk_push_descriptor_set *push_set = desc->push[set_index];
descriptor_buffer = push_set->mtl_descriptor_buffer;
}
/* We could have empty descriptor sets for some reason... */
if (descriptor_buffer) {
mtl_compute_use_resource(encoder, descriptor_buffer,
MTL_RESOURCE_USAGE_READ);
}
}
desc->sets_not_resident = 0u;
}
void
kk_make_descriptor_resources_resident(struct kk_cmd_buffer *cmd,
VkPipelineBindPoint bind_point)
{
if (bind_point == VK_PIPELINE_BIND_POINT_GRAPHICS)
kk_make_graphics_descriptor_resources_resident(cmd);
else if (bind_point == VK_PIPELINE_BIND_POINT_COMPUTE)
kk_make_compute_descriptor_resources_resident(cmd);
}
void
kk_cmd_write(struct kk_cmd_buffer *cmd, mtl_buffer *buffer, uint64_t addr,
uint64_t value)
{
util_dynarray_append(&cmd->encoder->imm_writes, addr);
util_dynarray_append(&cmd->encoder->imm_writes, value);
util_dynarray_append(&cmd->encoder->resident_buffers, buffer);
}
VKAPI_ATTR void VKAPI_CALL

7
src/kosmickrisp/vulkan/kk_cmd_buffer.h
View File

@@ -61,10 +61,6 @@ struct kk_descriptor_state {
uint32_t set_sizes[KK_MAX_SETS];
struct kk_descriptor_set *sets[KK_MAX_SETS];
mtl_resource **resources[KK_MAX_SETS];
/* Non resident sets can either be sets or push. If sets[index] == NULL, then
* push[index] != NULL */
uint32_t sets_not_resident;
uint32_t push_dirty;
struct kk_push_descriptor_set *push[KK_MAX_SETS];
@@ -255,9 +251,6 @@ uint64_t kk_upload_descriptor_root(struct kk_cmd_buffer *cmd,
void kk_cmd_buffer_flush_push_descriptors(struct kk_cmd_buffer *cmd,
struct kk_descriptor_state *desc);
void kk_make_descriptor_resources_resident(struct kk_cmd_buffer *cmd,
VkPipelineBindPoint bind_point);
void kk_cmd_write(struct kk_cmd_buffer *cmd, mtl_buffer *buffer, uint64_t addr,
uint64_t value);

6
src/kosmickrisp/vulkan/kk_cmd_dispatch.c
View File

@@ -72,18 +72,12 @@ static void
kk_flush_compute_state(struct kk_cmd_buffer *cmd)
{
mtl_compute_encoder *enc = kk_compute_encoder(cmd);
struct kk_device *dev = kk_cmd_buffer_device(cmd);
// Fill Metal argument buffer with descriptor set addresses
struct kk_descriptor_state *desc = &cmd->state.cs.descriptors;
if (desc->push_dirty)
kk_cmd_buffer_flush_push_descriptors(cmd, desc);
/* After push descriptors' buffers are created. Otherwise, the buffer where
* they live will not be created and cannot make it resident */
if (desc->sets_not_resident)
kk_make_descriptor_resources_resident(cmd,
VK_PIPELINE_BIND_POINT_COMPUTE);
if (desc->root_dirty)
kk_upload_descriptor_root(cmd, VK_PIPELINE_BIND_POINT_COMPUTE);

5
src/kosmickrisp/vulkan/kk_cmd_draw.c
View File

@@ -816,11 +816,6 @@ kk_flush_draw_state(struct kk_cmd_buffer *cmd)
if (desc->push_dirty)
kk_cmd_buffer_flush_push_descriptors(cmd, desc);
/* After push descriptors' buffers are created. Otherwise, the buffer where
* they live will not be created and cannot make it resident */
if (desc->sets_not_resident)
kk_make_descriptor_resources_resident(cmd,
VK_PIPELINE_BIND_POINT_GRAPHICS);
if (desc->root_dirty)
kk_upload_descriptor_root(cmd, VK_PIPELINE_BIND_POINT_GRAPHICS);

11
src/kosmickrisp/vulkan/kk_cmd_meta.c
View File

@@ -35,8 +35,7 @@ kk_cmd_bind_map_buffer(struct vk_command_buffer *vk_cmd,
buffer->mtl_handle = bo->map;
buffer->vk.device_address = bo->gpu;
*map_out = bo->cpu;
mtl_compute_use_resource(cmd->encoder->main.encoder, buffer->mtl_handle,
MTL_RESOURCE_USAGE_WRITE | MTL_RESOURCE_USAGE_READ);
return VK_SUCCESS;
}
@@ -147,11 +146,9 @@ kk_meta_end(struct kk_cmd_buffer *cmd, struct kk_meta_save *save,
desc->sets[0] = save->desc0;
desc->root.sets[0] = save->desc0->addr;
desc->set_sizes[0] = save->desc0->size;
desc->sets_not_resident |= BITFIELD_BIT(0);
desc->push_dirty &= ~BITFIELD_BIT(0);
} else if (save->has_push_desc0) {
desc->push[0] = save->push_desc0;
desc->sets_not_resident |= BITFIELD_BIT(0);
desc->push_dirty |= BITFIELD_BIT(0);
}
@@ -197,13 +194,10 @@ kk_CmdFillBuffer(VkCommandBuffer commandBuffer, VkBuffer dstBuffer,
VkDeviceSize dstOffset, VkDeviceSize dstRange, uint32_t data)
{
VK_FROM_HANDLE(kk_cmd_buffer, cmd, commandBuffer);
VK_FROM_HANDLE(kk_buffer, buf, dstBuffer);
struct kk_device *dev = kk_cmd_buffer_device(cmd);
struct kk_meta_save save;
kk_meta_begin(cmd, &save, VK_PIPELINE_BIND_POINT_COMPUTE);
mtl_compute_use_resource(kk_compute_encoder(cmd), buf->mtl_handle,
MTL_RESOURCE_USAGE_WRITE);
vk_meta_fill_buffer(&cmd->vk, &dev->meta, dstBuffer, dstOffset, dstRange,
data);
kk_meta_end(cmd, &save, VK_PIPELINE_BIND_POINT_COMPUTE);
@@ -215,13 +209,10 @@ kk_CmdUpdateBuffer(VkCommandBuffer commandBuffer, VkBuffer dstBuffer,
const void *pData)
{
VK_FROM_HANDLE(kk_cmd_buffer, cmd, commandBuffer);
VK_FROM_HANDLE(kk_buffer, buf, dstBuffer);
struct kk_device *dev = kk_cmd_buffer_device(cmd);
struct kk_meta_save save;
kk_meta_begin(cmd, &save, VK_PIPELINE_BIND_POINT_COMPUTE);
mtl_compute_use_resource(kk_compute_encoder(cmd), buf->mtl_handle,
MTL_RESOURCE_USAGE_WRITE);
vk_meta_update_buffer(&cmd->vk, &dev->meta, dstBuffer, dstOffset, dstRange,
pData);
kk_meta_end(cmd, &save, VK_PIPELINE_BIND_POINT_COMPUTE);

1
src/kosmickrisp/vulkan/kk_descriptor_set.c
View File

@@ -579,7 +579,6 @@ kk_descriptor_set_create(struct kk_device *dev, struct kk_descriptor_pool *pool,
vk_object_free(&dev->vk, NULL, set);
return result;
}
set->mtl_descriptor_buffer = pool->bo->map;
}
vk_descriptor_set_layout_ref(&layout->vk);

4
src/kosmickrisp/vulkan/kk_descriptor_set.h
View File

@@ -41,7 +41,6 @@ struct kk_descriptor_set {
struct list_head link;
struct kk_descriptor_set_layout *layout;
mtl_resource *mtl_descriptor_buffer;
void *mapped_ptr;
uint64_t addr;
uint32_t size;
@@ -64,9 +63,6 @@ kk_descriptor_set_addr(const struct kk_descriptor_set *set)
struct kk_push_descriptor_set {
uint8_t data[KK_PUSH_DESCRIPTOR_SET_SIZE];
struct kk_descriptor_set_layout *layout;
mtl_resource *mtl_descriptor_buffer;
uint32_t resource_count;
mtl_resource *mtl_resources[];
};
void kk_push_descriptor_set_update(struct kk_push_descriptor_set *push_set,

3
src/kosmickrisp/vulkan/kk_descriptor_set_layout.c
View File

@@ -176,7 +176,6 @@ kk_CreateDescriptorSetLayout(VkDevice device,
uint32_t buffer_size = 0;
uint32_t max_variable_descriptor_size = 0;
uint8_t dynamic_buffer_count = 0;
uint32_t total_descriptor_count = 0u;
for (uint32_t b = 0; b < num_bindings; b++) {
/* We stashed the pCreateInfo->pBindings[] index (plus one) in the
* immutable_samplers pointer. Check for NULL (empty binding) and then
@@ -195,7 +194,6 @@ kk_CreateDescriptorSetLayout(VkDevice device,
continue;
layout->binding[b].type = binding->descriptorType;
layout->binding[b].mtl_resources_index = total_descriptor_count;
layout->descriptor_count += binding->descriptorCount;
if (binding_flags_info && binding_flags_info->bindingCount > 0) {
@@ -250,7 +248,6 @@ kk_CreateDescriptorSetLayout(VkDevice device,
stride *= max_plane_count;
layout->binding[b].count_per_element = max_plane_count;
total_descriptor_count += max_plane_count * binding->descriptorCount;
if (stride > 0) {
assert(stride <= UINT8_MAX);

3
src/kosmickrisp/vulkan/kk_descriptor_set_layout.h
View File

@@ -38,9 +38,6 @@ struct kk_descriptor_set_binding_layout {
/* Offset into the descriptor buffer where this descriptor lives */
uint32_t offset;
/* Offset to the mtl_resource_ids array where this descriptor stores them */
uint32_t mtl_resources_index;
/* Stride between array elements in the descriptor buffer */
uint8_t stride;

64
src/kosmickrisp/vulkan/kk_device.c
View File

@@ -205,10 +205,10 @@ kk_CreateDevice(VkPhysicalDevice physicalDevice,
dev->vk.command_buffer_ops = &kk_cmd_buffer_ops;
dev->vk.command_dispatch_table = &dev->vk.dispatch_table;
/* Buffer to use as null descriptor */
result = kk_alloc_bo(dev, &dev->vk.base, sizeof(uint64_t) * 8, 8u,
&dev->null_descriptor);
if (result != VK_SUCCESS)
/* We need to initialize the device residency set before any bo is created. */
simple_mtx_init(&dev->residency_set.mutex, mtx_plain);
dev->residency_set.handle = mtl_new_residency_set(dev->mtl_handle);
if (dev->residency_set.handle == NULL)
goto fail_init;
result =
@@ -233,10 +233,6 @@ kk_CreateDevice(VkPhysicalDevice physicalDevice,
if (result != VK_SUCCESS)
goto fail_sampler_heap;
simple_mtx_init(&dev->user_residency_set.mutex, mtx_plain);
dev->user_residency_set.residency_set =
mtl_new_residency_set(dev->mtl_handle);
kk_parse_device_environment_options(dev);
*pDevice = kk_device_to_handle(dev);
@@ -252,7 +248,8 @@ fail_meta:
fail_mem_cache:
kk_queue_finish(dev, &dev->queue);
fail_vab_memory:
kk_destroy_bo(dev, dev->null_descriptor);
mtl_release(dev->residency_set.handle);
simple_mtx_destroy(&dev->residency_set.mutex);
fail_init:
vk_device_finish(&dev->vk);
fail_alloc:
@@ -268,25 +265,24 @@ kk_DestroyDevice(VkDevice _device, const VkAllocationCallbacks *pAllocator)
if (!dev)
return;
/* End capture before we start releasing resources. Otherwise, Metal capture
* may run into issues. */
if (dev->gpu_capture_enabled) {
mtl_stop_gpu_capture();
}
/* Meta first since it may destroy Vulkan objects */
kk_device_finish_meta(dev);
/* Need to end the residency otherwise stopping a capture crashes the
* program... */
mtl_residency_set_end_residency(dev->user_residency_set.residency_set);
mtl_release(dev->user_residency_set.residency_set);
simple_mtx_destroy(&dev->user_residency_set.mutex);
kk_device_finish_lib(dev);
kk_query_table_finish(dev, &dev->occlusion_queries);
kk_destroy_sampler_heap(dev, &dev->samplers);
kk_queue_finish(dev, &dev->queue);
kk_destroy_bo(dev, dev->null_descriptor);
vk_device_finish(&dev->vk);
/* Release the residency set last once all BOs are released. */
mtl_release(dev->residency_set.handle);
simple_mtx_destroy(&dev->residency_set.mutex);
if (dev->gpu_capture_enabled) {
mtl_stop_gpu_capture();
}
kk_queue_finish(dev, &dev->queue);
vk_device_finish(&dev->vk);
vk_free(&dev->vk.alloc, dev);
}
@@ -358,28 +354,26 @@ kk_GetDeviceProcAddr(VkDevice _device, const char *pName)
}
void
kk_device_add_user_heap(struct kk_device *dev, mtl_heap *heap)
kk_device_add_heap_to_residency_set(struct kk_device *dev, mtl_heap *heap)
{
simple_mtx_lock(&dev->user_residency_set.mutex);
mtl_residency_set_add_allocation(dev->user_residency_set.residency_set,
heap);
simple_mtx_unlock(&dev->user_residency_set.mutex);
simple_mtx_lock(&dev->residency_set.mutex);
mtl_residency_set_add_allocation(dev->residency_set.handle, heap);
simple_mtx_unlock(&dev->residency_set.mutex);
}
void
kk_device_remove_user_heap(struct kk_device *dev, mtl_heap *heap)
kk_device_remove_heap_from_residency_set(struct kk_device *dev, mtl_heap *heap)
{
simple_mtx_lock(&dev->user_residency_set.mutex);
mtl_residency_set_remove_allocation(dev->user_residency_set.residency_set,
heap);
simple_mtx_unlock(&dev->user_residency_set.mutex);
simple_mtx_lock(&dev->residency_set.mutex);
mtl_residency_set_remove_allocation(dev->residency_set.handle, heap);
simple_mtx_unlock(&dev->residency_set.mutex);
}
void
kk_device_make_resources_resident(struct kk_device *dev)
{
simple_mtx_lock(&dev->user_residency_set.mutex);
mtl_residency_set_commit(dev->user_residency_set.residency_set);
mtl_residency_set_request_residency(dev->user_residency_set.residency_set);
simple_mtx_unlock(&dev->user_residency_set.mutex);
simple_mtx_lock(&dev->residency_set.mutex);
mtl_residency_set_commit(dev->residency_set.handle);
mtl_residency_set_request_residency(dev->residency_set.handle);
simple_mtx_unlock(&dev->residency_set.mutex);
}

13
src/kosmickrisp/vulkan/kk_device.h
View File

@@ -32,9 +32,9 @@ enum kk_device_lib_pipeline {
KK_LIB_COUNT,
};
struct kk_user_residency_set {
struct kk_residency_set {
simple_mtx_t mutex;
mtl_residency_set *residency_set;
mtl_residency_set *handle;
};
struct mtl_sampler_packed {
@@ -84,14 +84,12 @@ struct kk_device {
* commands due to Metal limitations */
struct vk_device_dispatch_table exposed_dispatch_table;
struct kk_bo *null_descriptor;
struct kk_sampler_heap samplers;
struct kk_query_table occlusion_queries;
/* Track all heaps the user allocated so we can set them all as resident when
* recording as required by Metal. */
struct kk_user_residency_set user_residency_set;
struct kk_residency_set residency_set;
mtl_compute_pipeline_state *lib_pipelines[KK_LIB_COUNT];
@@ -123,8 +121,9 @@ VkResult kk_device_init_meta(struct kk_device *dev);
void kk_device_finish_meta(struct kk_device *dev);
VkResult kk_device_init_lib(struct kk_device *dev);
void kk_device_finish_lib(struct kk_device *dev);
void kk_device_add_user_heap(struct kk_device *dev, mtl_heap *heap);
void kk_device_remove_user_heap(struct kk_device *dev, mtl_heap *heap);
void kk_device_add_heap_to_residency_set(struct kk_device *dev, mtl_heap *heap);
void kk_device_remove_heap_from_residency_set(struct kk_device *dev,
mtl_heap *heap);
void kk_device_make_resources_resident(struct kk_device *dev);
/* Required to create a sampler */

5
src/kosmickrisp/vulkan/kk_device_memory.c
View File

@@ -94,6 +94,7 @@ kk_AllocateMemory(VkDevice device, const VkMemoryAllocateInfo *pAllocateInfo,
mem->bo->gpu = mtl_buffer_get_gpu_address(mem->bo->map);
mem->bo->cpu = mtl_get_contents(mem->bo->map);
mem->bo->size_B = mtl_heap_get_size(mem->bo->mtl_handle);
kk_device_add_heap_to_residency_set(dev, mem->bo->mtl_handle);
} else {
result =
kk_alloc_bo(dev, &dev->vk.base, aligned_size, alignment, &mem->bo);
@@ -104,8 +105,6 @@ kk_AllocateMemory(VkDevice device, const VkMemoryAllocateInfo *pAllocateInfo,
struct kk_memory_heap *heap = &pdev->mem_heaps[type->heapIndex];
p_atomic_add(&heap->used, mem->bo->size_B);
kk_device_add_user_heap(dev, mem->bo->mtl_handle);
*pMem = kk_device_memory_to_handle(mem);
return VK_SUCCESS;
@@ -126,8 +125,6 @@ kk_FreeMemory(VkDevice device, VkDeviceMemory _mem,
if (!mem)
return;
kk_device_remove_user_heap(dev, mem->bo->mtl_handle);
const VkMemoryType *type = &pdev->mem_types[mem->vk.memory_type_index];
struct kk_memory_heap *heap = &pdev->mem_heaps[type->heapIndex];
p_atomic_add(&heap->used, -((int64_t)mem->bo->size_B));

20
src/kosmickrisp/vulkan/kk_encoder.c
View File

@@ -39,7 +39,6 @@ kk_encoder_init(mtl_device *device, struct kk_queue *queue,
kk_encoder_start_internal(&enc->pre_gfx, device, queue->pre_gfx.mtl_handle);
enc->event = mtl_new_event(device);
enc->imm_writes = UTIL_DYNARRAY_INIT;
enc->resident_buffers = UTIL_DYNARRAY_INIT;
enc->copy_query_pool_result_infos = UTIL_DYNARRAY_INIT;
*encoder = enc;
@@ -180,12 +179,6 @@ upload_queue_writes(struct kk_cmd_buffer *cmd)
if (!bo)
return;
memcpy(bo->cpu, enc->imm_writes.data, enc->imm_writes.size);
uint32_t buffer_count =
util_dynarray_num_elements(&enc->resident_buffers, mtl_buffer *);
mtl_compute_use_resource(compute, bo->map, MTL_RESOURCE_USAGE_READ);
mtl_compute_use_resources(
compute, enc->resident_buffers.data, buffer_count,
MTL_RESOURCE_USAGE_READ | MTL_RESOURCE_USAGE_WRITE);
struct kk_imm_write_push push_data = {
.buffer_address = bo->gpu,
.count = count,
@@ -193,7 +186,6 @@ upload_queue_writes(struct kk_cmd_buffer *cmd)
kk_cmd_dispatch_pipeline(cmd, compute,
kk_device_lib_pipeline(dev, KK_LIB_IMM_WRITE),
&push_data, sizeof(push_data), count, 1, 1);
enc->resident_buffers.size = 0u;
enc->imm_writes.size = 0u;
}
@@ -201,11 +193,6 @@ upload_queue_writes(struct kk_cmd_buffer *cmd)
struct kk_copy_query_pool_results_info);
if (count != 0u) {
mtl_compute_encoder *compute = kk_compute_encoder(cmd);
uint32_t buffer_count =
util_dynarray_num_elements(&enc->resident_buffers, mtl_buffer *);
mtl_compute_use_resources(
compute, enc->resident_buffers.data, buffer_count,
MTL_RESOURCE_USAGE_READ | MTL_RESOURCE_USAGE_WRITE);
for (uint32_t i = 0u; i < count; ++i) {
struct kk_copy_query_pool_results_info *push_data =
@@ -216,7 +203,6 @@ upload_queue_writes(struct kk_cmd_buffer *cmd)
cmd, compute, kk_device_lib_pipeline(dev, KK_LIB_COPY_QUERY),
push_data, sizeof(*push_data), push_data->query_count, 1, 1);
}
enc->resident_buffers.size = 0u;
enc->copy_query_pool_result_infos.size = 0u;
}
@@ -292,7 +278,6 @@ kk_post_execution_release(void *data)
kk_post_execution_release_internal(&encoder->pre_gfx);
mtl_release(encoder->event);
util_dynarray_fini(&encoder->imm_writes);
util_dynarray_fini(&encoder->resident_buffers);
util_dynarray_fini(&encoder->copy_query_pool_result_infos);
free(encoder);
}
@@ -416,11 +401,6 @@ kk_encoder_render_triangle_fan_common(struct kk_cmd_buffer *cmd,
info->out_el_size_B = out_el_size_B;
info->flatshade_first = true;
mtl_compute_encoder *encoder = kk_encoder_pre_gfx_encoder(cmd);
if (index)
mtl_compute_use_resource(encoder, index, MTL_RESOURCE_USAGE_READ);
mtl_compute_use_resource(encoder, indirect, MTL_RESOURCE_USAGE_READ);
mtl_compute_use_resource(encoder, index_buffer->map,
MTL_RESOURCE_USAGE_WRITE);
struct kk_device *dev = kk_cmd_buffer_device(cmd);
kk_cmd_dispatch_pipeline(cmd, encoder,

3
src/kosmickrisp/vulkan/kk_encoder.h
View File

@@ -70,9 +70,6 @@ struct kk_encoder {
/* uint64_t pairs with first being the address, second being the value to
* write */
struct util_dynarray imm_writes;
/* mtl_buffers (destination buffers) so we can make them resident before the
* dispatch */
struct util_dynarray resident_buffers;
/* Array of kk_copy_quer_pool_results_info structs */
struct util_dynarray copy_query_pool_result_infos;
};

4
src/kosmickrisp/vulkan/kk_query_pool.c
View File

@@ -418,10 +418,6 @@ kk_CmdCopyQueryPoolResults(VkCommandBuffer commandBuffer, VkQueryPool queryPool,
};
util_dynarray_append(&cmd->encoder->copy_query_pool_result_infos, info);
util_dynarray_append(&cmd->encoder->resident_buffers, dst_buf->mtl_handle);
util_dynarray_append(&cmd->encoder->resident_buffers, pool->bo->map);
util_dynarray_append(&cmd->encoder->resident_buffers,
dev->occlusion_queries.bo->map);
/* If we are not mid encoder, just upload the writes */
if (cmd->encoder->main.last_used == KK_ENC_NONE)
upload_queue_writes(cmd);

10
src/kosmickrisp/vulkan/kk_queue.c
View File

@@ -31,11 +31,6 @@ kk_queue_submit(struct vk_queue *vk_queue, struct vk_queue_submit *submit)
if (result != VK_SUCCESS)
return result;
/* Ensure any changes to residency are propagated before we submit any work.
* All resources should have been allocated before submission. Otherwise,
* users are playing with fire. */
kk_device_make_resources_resident(dev);
/* Chain with previous sumbission */
if (queue->wait_fence) {
util_dynarray_append(&encoder->main.fences, queue->wait_fence);
@@ -78,6 +73,11 @@ kk_queue_submit(struct vk_queue *vk_queue, struct vk_queue_submit *submit)
signal->signal_value);
}
/* Ensure any changes to residency are propagated before we submit any work.
* All resources should have been allocated before submission. Otherwise,
* users are playing with fire. */
kk_device_make_resources_resident(dev);
/* Steal the last fence to chain with the next submission */
if (util_dynarray_num_elements(&encoder->main.fences, mtl_fence *) > 0)
queue->wait_fence = util_dynarray_pop(&encoder->main.fences, mtl_fence *);