From 0da350f8797b637ff7f073b4134cdbe75841b58a Mon Sep 17 00:00:00 2001 From: Caterina Shablia Date: Tue, 16 Dec 2025 13:36:44 +0000 Subject: [PATCH] panvk: remove AFBC header zeroing This is not actually necessary and moreover was corrupting mipmapped arrayed 2D images in cases when the transition barrier wasn't transitioning all mips, but more than one layer. Keep the layout transition infrastructure in place as we'll need it for transaction elimination CRC zeroing on v10-. Fixes: c95f8993 ("panvk: add a meta command for transitioning image layout") Reviewed-by: Boris Brezillon Part-of: --- src/panfrost/libpan/clear_afbc_metadata.cl | 21 ------ src/panfrost/libpan/meson.build | 1 - src/panfrost/vulkan/panvk_vX_cmd_meta.c | 79 ---------------------- 3 files changed, 101 deletions(-) delete mode 100644 src/panfrost/libpan/clear_afbc_metadata.cl 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}; }