diff --git a/src/panfrost/libpan/clear_afbc_metadata.cl b/src/panfrost/libpan/clear_afbc_metadata.cl deleted file mode 100644 index 195e985fe0c..00000000000 --- a/src/panfrost/libpan/clear_afbc_metadata.cl +++ /dev/null @@ -1,21 +0,0 @@ -/* - * Copyright 2024 Collabora Ltd. - * SPDX-License-Identifier: MIT - */ -#include "compiler/libcl/libcl.h" -#include "compiler/libcl/libcl_vk.h" -#include "genxml/gen_macros.h" -#include "lib/pan_encoder.h" - -#if PAN_ARCH >= 6 -KERNEL(1) -panlib_clear_afbc_metadata(global uint8_t *p, - uint32_t layer_or_slice_stride) -{ - uint32_t item = get_global_id(0); - uint32_t layer_or_slice = get_global_id(1); - - uint4 *q = p + layer_or_slice * layer_or_slice_stride; - q[item] = 0; -} -#endif diff --git a/src/panfrost/libpan/meson.build b/src/panfrost/libpan/meson.build index 5eb21bf4bb9..9bf43632d28 100644 --- a/src/panfrost/libpan/meson.build +++ b/src/panfrost/libpan/meson.build @@ -3,7 +3,6 @@ libpan_shader_files = files( - 'clear_afbc_metadata.cl', 'query_pool.cl', 'draw_helper.cl', 'indirect_dispatch.cl', diff --git a/src/panfrost/vulkan/panvk_vX_cmd_meta.c b/src/panfrost/vulkan/panvk_vX_cmd_meta.c index 970b2b78fd6..1022274a6a8 100644 --- a/src/panfrost/vulkan/panvk_vX_cmd_meta.c +++ b/src/panfrost/vulkan/panvk_vX_cmd_meta.c @@ -592,74 +592,6 @@ panvk_image_has_afbc(struct panvk_image *img, VkImageSubresourceRange range) return false; } -static void -cmd_clear_afbc_metadata(VkCommandBuffer _cmdbuf, - const VkImageMemoryBarrier2 *barrier) -{ - VK_FROM_HANDLE(panvk_cmd_buffer, cmdbuf, _cmdbuf); - struct panvk_precomp_ctx precomp_ctx = panvk_per_arch(precomp_cs)(cmdbuf); - struct panvk_device *dev = to_panvk_device(cmdbuf->vk.base.device); - VK_FROM_HANDLE(panvk_image, img, barrier->image); - VkImageSubresourceRange range = barrier->subresourceRange; - VkImageAspectFlags aspect_mask = - vk_image_expand_aspect_mask(&img->vk, range.aspectMask); - uint32_t level_count = vk_image_subresource_level_count(&img->vk, &range); - struct panvk_cmd_meta_compute_save_ctx save = {0}; - - panvk_per_arch(cmd_meta_compute_start)(cmdbuf, &save); - - u_foreach_bit(aspect, aspect_mask) { - unsigned plane_index = panvk_plane_index(img, 1u << aspect); - struct panvk_image_plane *plane = &img->planes[plane_index]; - - if (!drm_is_afbc(plane->image.props.modifier)) - continue; - - for (uint32_t level = range.baseMipLevel; - level < range.baseMipLevel + level_count; - level++) { - const struct pan_image_slice_layout *slayout = - &plane->plane.layout.slices[level]; - - uint32_t layers_or_slices; - if (img->vk.image_type == VK_IMAGE_TYPE_2D) { - layers_or_slices = - vk_image_subresource_layer_count(&img->vk, &range); - } else if (img->vk.image_type == VK_IMAGE_TYPE_3D) { - layers_or_slices = - vk_image_subresource_slice_count(&dev->vk, - &img->vk, - &(VkImageSubresourceLayers) { - .mipLevel = level, - .baseArrayLayer = range.baseArrayLayer, - .layerCount = range.layerCount, - }); - } else { - UNREACHABLE("Unsupported image type"); - } - - uint32_t layer_or_slice_stride = slayout->afbc.surface_stride_B; - - uint32_t ptr = plane->plane.base + slayout->offset_B + - range.baseArrayLayer * layer_or_slice_stride; - - struct panlib_clear_afbc_metadata_args args = { - .p = ptr, - .layer_or_slice_stride = layer_or_slice_stride, - }; - panlib_clear_afbc_metadata_struct( - &precomp_ctx, - panlib_3d(slayout->afbc.header.surface_size_B / 16, - layers_or_slices, 1), - PAN_ARCH >= 10 ? PANLIB_BARRIER_CSF_SYNC : - PANLIB_BARRIER_JM_BARRIER, - args); - } - } - - panvk_per_arch(cmd_meta_compute_end)(cmdbuf, &save); -} - static bool panvk_acquire_unmodified(const VkImageMemoryBarrier2 *barrier) { @@ -685,21 +617,10 @@ struct panvk_image_layout_transition_handler { static struct panvk_image_layout_transition_handler panvk_get_image_layout_transition_handler(const VkImageMemoryBarrier2 *barrier) { - VK_FROM_HANDLE(panvk_image, img, barrier->image); - if (barrier->oldLayout == barrier->newLayout || panvk_acquire_unmodified(barrier)) return (struct panvk_image_layout_transition_handler){0}; - if (barrier->oldLayout == VK_IMAGE_LAYOUT_UNDEFINED && - panvk_image_has_afbc(img, barrier->subresourceRange)) { - return (struct panvk_image_layout_transition_handler){ - .cmd = cmd_clear_afbc_metadata, - .stages = VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT, - .access = VK_ACCESS_2_MEMORY_WRITE_BIT, - }; - } - return (struct panvk_image_layout_transition_handler){0}; }