mirror of
https://gitlab.freedesktop.org/mesa/mesa.git
synced 2026-04-21 16:20:39 +02:00
panfrost: Make the MTK detiling more robust
There are several problems in the current logic: - it doesn't account for the WSI stride, and just assumes things are tightly packed and the resource extent is aligned on a tile - the format re-interpretation done in panfrost_mtk_detile_compute() where the view format and the image format have a different block size is not supposed to be supported, and that's something we'd like to enforce now - we write to textures that may have no descriptors bound to. This works thanks to the robustness behind image stores/loads, but we'd probably rather discard a bunch of instructions when Y/UV planes are copied separately - the linear to tiled coordinates conversion can be simplified by using a combination of local_invocation_id+workgroup_id instead of the global_invocation_id Signed-off-by: Boris Brezillon <boris.brezillon@collabora.com> Reviewed-by: Mary Guillemard <mary.guillemard@collabora.com> Tested-by: Eric R. Smith <eric.smith@collabora.com> Reviewed-by: Eric R. Smith <eric.smith@collabora.com> Acked-by: Erik Faye-Lund <erik.faye-lund@collabora.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/35317>
This commit is contained in:
Boris Brezillon
Marge Bot
parent
a87805fdd8
commit
eb476fd2c6
5 changed files with 285 additions and 130 deletions
|
|
@ -3652,10 +3652,6 @@ panfrost_mtk_detile_compute(struct panfrost_context *ctx, struct pipe_blit_info
|
|||
|
||||
unsigned width = info->src.box.width;
|
||||
unsigned height = info->src.box.height;
|
||||
unsigned src_stride =
|
||||
pan_resource(y_src)->plane.layout.slices[0].row_stride_B;
|
||||
unsigned dst_stride =
|
||||
pan_resource(y_dst)->plane.layout.slices[0].row_stride_B;
|
||||
|
||||
/* 4 images: y_src, uv_src, y_dst, uv_dst */
|
||||
struct pipe_image_view image[4] = { 0 };
|
||||
|
|
@ -3673,6 +3669,43 @@ panfrost_mtk_detile_compute(struct panfrost_context *ctx, struct pipe_blit_info
|
|||
y_src = y_dst = NULL;
|
||||
}
|
||||
}
|
||||
|
||||
/* We're not supposed to create views with a format whose size doesn't match
|
||||
* the image format. */
|
||||
struct panfrost_resource y_src_save;
|
||||
struct panfrost_resource uv_src_save;
|
||||
struct panfrost_resource y_dst_save;
|
||||
struct panfrost_resource uv_dst_save;
|
||||
|
||||
panfrost_resource_change_format(pan_resource(y_src),
|
||||
PIPE_FORMAT_R8G8B8A8_UINT, &y_src_save);
|
||||
panfrost_resource_change_format(pan_resource(uv_src),
|
||||
PIPE_FORMAT_R8G8B8A8_UINT, &uv_src_save);
|
||||
panfrost_resource_change_format(pan_resource(y_dst),
|
||||
PIPE_FORMAT_R8G8B8A8_UINT, &y_dst_save);
|
||||
panfrost_resource_change_format(pan_resource(uv_dst),
|
||||
PIPE_FORMAT_R8G8B8A8_UINT, &uv_dst_save);
|
||||
|
||||
struct panfrost_mtk_detile_info consts = {
|
||||
.height = height,
|
||||
/* The copy width is expressed for an R8_UNORM resource, but we
|
||||
* changed the format into RGBA8_UINT, so we need to adjust the width if
|
||||
* we want the shader-side bound check to do its job. */
|
||||
.width = width / 4,
|
||||
};
|
||||
|
||||
if (y_src) {
|
||||
consts.src_y_row_stride_tl =
|
||||
pan_resource(y_src)->image.props.extent_px.height /
|
||||
DIV_ROUND_UP(y_src->height0, 32);
|
||||
}
|
||||
|
||||
if (uv_src) {
|
||||
consts.src_uv_row_stride_tl =
|
||||
pan_resource(uv_src)->image.props.extent_px.height /
|
||||
DIV_ROUND_UP(uv_src->height0, 16);
|
||||
}
|
||||
|
||||
image[0].resource = y_src;
|
||||
image[0].format = PIPE_FORMAT_R8G8B8A8_UINT;
|
||||
image[0].shader_access = image[0].access = PIPE_IMAGE_ACCESS_READ;
|
||||
|
|
@ -3701,12 +3734,6 @@ panfrost_mtk_detile_compute(struct panfrost_context *ctx, struct pipe_blit_info
|
|||
image[3].u.tex.first_layer = 0;
|
||||
image[3].u.tex.last_layer = uv_dst ? (unsigned)(uv_dst->array_size - 1) : 0;
|
||||
|
||||
struct panfrost_mtk_detile_info consts = {
|
||||
.tiles_per_stride = src_stride >> 4,
|
||||
.src_width = width,
|
||||
.src_height = height,
|
||||
.dst_stride = dst_stride,
|
||||
};
|
||||
panfrost_flush_all_batches(ctx, "mtk_detile pre-barrier");
|
||||
|
||||
struct panfrost_batch *batch = panfrost_get_batch_for_fbo(ctx);
|
||||
|
|
@ -3714,20 +3741,20 @@ panfrost_mtk_detile_compute(struct panfrost_context *ctx, struct pipe_blit_info
|
|||
|
||||
/* launch the compute shader */
|
||||
struct pan_mod_convert_shader_data *shader =
|
||||
panfrost_get_mtk_detile_shader(ctx);
|
||||
panfrost_get_mtk_detile_shader(ctx, y_src != NULL, uv_src != NULL);
|
||||
struct pipe_constant_buffer cbuf = {
|
||||
.buffer_size = sizeof(consts),
|
||||
.user_buffer = &consts};
|
||||
|
||||
struct pipe_grid_info grid_info = {
|
||||
.block[0] = 4,
|
||||
.last_block[0] = (width/4) % 4,
|
||||
.last_block[0] = (width / 4) % 4,
|
||||
.block[1] = 16,
|
||||
.last_block[1] = height % 16,
|
||||
.last_block[1] = (height / 2) % 16,
|
||||
.block[2] = 1,
|
||||
.last_block[2] = 0,
|
||||
.grid[0] = DIV_ROUND_UP(width/4, 4),
|
||||
.grid[1] = DIV_ROUND_UP(height, 16),
|
||||
.grid[0] = DIV_ROUND_UP(width / 4, 4),
|
||||
.grid[1] = DIV_ROUND_UP(height / 2, 16),
|
||||
.grid[2] = 1,
|
||||
};
|
||||
|
||||
|
|
@ -3745,6 +3772,11 @@ panfrost_mtk_detile_compute(struct panfrost_context *ctx, struct pipe_blit_info
|
|||
|
||||
pipe->bind_compute_state(pipe, saved_cso);
|
||||
pipe->set_constant_buffer(pipe, PIPE_SHADER_COMPUTE, 0, true, &saved_const);
|
||||
|
||||
panfrost_resource_restore_format(pan_resource(y_src), &y_src_save);
|
||||
panfrost_resource_restore_format(pan_resource(uv_src), &uv_src_save);
|
||||
panfrost_resource_restore_format(pan_resource(y_dst), &y_dst_save);
|
||||
panfrost_resource_restore_format(pan_resource(uv_dst), &uv_dst_save);
|
||||
}
|
||||
|
||||
static void *
|
||||
|
|
|
|||
|
|
@ -285,55 +285,21 @@ panfrost_create_afbc_pack_shader(struct panfrost_screen *screen,
|
|||
#define panfrost_mtk_detile_get_info_field(b, field) \
|
||||
panfrost_mtk_get_info_field(detile, b, field)
|
||||
|
||||
static nir_def *
|
||||
pan_mtk_tiled_from_linear(nir_builder *b, nir_def *linear, nir_def *tiles_per_stride, nir_def *width)
|
||||
static void
|
||||
copy_y_uv_texel(nir_builder *b, unsigned src_img, nir_def *src_coords,
|
||||
unsigned dst_img, nir_def *dst_coords)
|
||||
{
|
||||
nir_def *tiled;
|
||||
/* uvec2 tlc = uvec2(linear) >> uvec2(2u, 5u) */
|
||||
nir_def *tlc = nir_ushr(b, linear,
|
||||
nir_imm_ivec2(b, 2, 5));
|
||||
nir_def *sample = nir_imm_int(b, 0);
|
||||
nir_def *lod = nir_imm_int(b, 0);
|
||||
|
||||
/* uvec2 txc = uvec2(linear) & uvec2(3u, 31u) */
|
||||
nir_def *txc = nir_iand(b, linear,
|
||||
nir_imm_ivec2(b, 3, 31));
|
||||
|
||||
/* uint tlo = tlc.y * tiles_per_stride + tlc.x */
|
||||
nir_def *tlo = nir_iadd(b,
|
||||
nir_imul(b,
|
||||
nir_channel(b, tlc, 1),
|
||||
tiles_per_stride),
|
||||
nir_channel(b, tlc, 0));
|
||||
nir_def *txcx = nir_channel(b, txc, 0);
|
||||
nir_def *txcy = nir_channel(b, txc, 1);
|
||||
nir_def *txcytmp = nir_vec2(b, txcy,
|
||||
nir_ushr_imm(b, txcy, 1));
|
||||
|
||||
/* txo = (uvec2(txc.y, txc.y >> 1) << uvec2(2u)) | txc.xx */
|
||||
nir_def *txo = nir_ior(b,
|
||||
nir_ishl_imm(b, txcytmp, 2),
|
||||
nir_vec2(b, txcx, txcx));
|
||||
|
||||
/* uvec2 off = (uvec2(tlo) << uvec2(7u, 6u)) | txo */
|
||||
nir_def *off = nir_ior(b,
|
||||
nir_ishl(b,
|
||||
nir_vec2(b, tlo, tlo),
|
||||
nir_imm_ivec2(b, 7, 6)),
|
||||
txo);
|
||||
|
||||
/* convert to 2D coord
|
||||
* tiled.xy = off % (width / 4, width / 4)
|
||||
* tiled.zw = off / (width / 4, width / 4) */
|
||||
nir_def *width4 = nir_ishl_imm(b, tiles_per_stride, 2);
|
||||
width4 = nir_vec2(b, width4, width4);
|
||||
nir_def *tiled_xy = nir_umod(b, off, width4);
|
||||
nir_def *tiled_zw = nir_udiv(b, off, width4);
|
||||
tiled = nir_vec4(b,
|
||||
nir_channel(b, tiled_xy, 0),
|
||||
nir_channel(b, tiled_xy, 1),
|
||||
nir_channel(b, tiled_zw, 0),
|
||||
nir_channel(b, tiled_zw, 1));
|
||||
|
||||
return tiled;
|
||||
nir_def *val = nir_image_load(
|
||||
b, 4, 32, nir_imm_int(b, src_img), src_coords, sample, lod,
|
||||
.access = ACCESS_NON_WRITEABLE, .image_dim = GLSL_SAMPLER_DIM_2D,
|
||||
.image_array = false, .dest_type = nir_type_uint32);
|
||||
nir_image_store(b, nir_imm_int(b, dst_img), dst_coords, sample, val, lod,
|
||||
.access = ACCESS_NON_READABLE,
|
||||
.image_dim = GLSL_SAMPLER_DIM_2D, .image_array = false,
|
||||
.src_type = nir_type_uint32);
|
||||
}
|
||||
|
||||
static nir_shader *
|
||||
|
|
@ -348,6 +314,36 @@ panfrost_create_mtk_tiled_detile_shader(
|
|||
b.shader->info.workgroup_size[0] = 4;
|
||||
b.shader->info.workgroup_size[1] = 16;
|
||||
b.shader->info.workgroup_size[2] = 1;
|
||||
nir_def *intra_tile_coords = nir_trim_vector(&b, nir_load_local_invocation_id(&b), 2);
|
||||
nir_def *wg_id = nir_trim_vector(&b, nir_load_workgroup_id(&b), 2);
|
||||
nir_def *uv_linear_coords = nir_pad_vector_imm_int(
|
||||
&b,
|
||||
nir_iadd(&b, nir_imul(&b, wg_id, nir_imm_ivec2(&b, 4, 16)),
|
||||
intra_tile_coords),
|
||||
0, 4);
|
||||
nir_def *y_linear_coords =
|
||||
nir_ishl(&b, uv_linear_coords, nir_imm_ivec4(&b, 0, 1, 0, 0));
|
||||
nir_def *src_y_row_stride_tl =
|
||||
panfrost_mtk_detile_get_info_field(&b, src_y_row_stride_tl);
|
||||
nir_def *src_uv_row_stride_tl =
|
||||
panfrost_mtk_detile_get_info_field(&b, src_uv_row_stride_tl);
|
||||
nir_def *dst_extent =
|
||||
nir_vec2(&b, panfrost_mtk_detile_get_info_field(&b, width),
|
||||
panfrost_mtk_detile_get_info_field(&b, height));
|
||||
nir_def *uv_tiled_coords = nir_vec4(
|
||||
&b,
|
||||
nir_iadd(&b, nir_imul_imm(&b, nir_channel(&b, intra_tile_coords, 1), 4),
|
||||
nir_channel(&b, intra_tile_coords, 0)),
|
||||
nir_iadd(&b, nir_imul(&b, nir_channel(&b, wg_id, 1), src_uv_row_stride_tl),
|
||||
nir_channel(&b, wg_id, 0)),
|
||||
nir_imm_int(&b, 0), nir_imm_int(&b, 0));
|
||||
nir_def *y_tiled_coords = nir_vec4(
|
||||
&b,
|
||||
nir_iadd(&b, nir_imul_imm(&b, nir_channel(&b, intra_tile_coords, 1), 8),
|
||||
nir_channel(&b, intra_tile_coords, 0)),
|
||||
nir_iadd(&b, nir_imul(&b, nir_channel(&b, wg_id, 1), src_y_row_stride_tl),
|
||||
nir_channel(&b, wg_id, 0)),
|
||||
nir_imm_int(&b, 0), nir_imm_int(&b, 0));
|
||||
|
||||
const struct glsl_type *image_type =
|
||||
glsl_image_type(GLSL_SAMPLER_DIM_2D, /*is_array*/ false, GLSL_TYPE_UINT);
|
||||
|
|
@ -377,70 +373,35 @@ panfrost_create_mtk_tiled_detile_shader(
|
|||
uv_linear->data.image.format = PIPE_FORMAT_R8G8B8A8_UINT;
|
||||
BITSET_SET(b.shader->info.images_used, 3);
|
||||
|
||||
nir_def *linear = nir_load_global_invocation_id(&b, 32);
|
||||
nir_def *tiles_per_stride =
|
||||
panfrost_mtk_detile_get_info_field(&b, tiles_per_stride);
|
||||
nir_def *src_width = panfrost_mtk_detile_get_info_field(&b, src_width);
|
||||
|
||||
nir_def *zero = nir_imm_int(&b, 0);
|
||||
|
||||
nir_def *coord = nir_vec2(&b,
|
||||
nir_channel(&b, linear, 0),
|
||||
nir_channel(&b, linear, 1));
|
||||
|
||||
nir_def *tiled = pan_mtk_tiled_from_linear(&b, coord, tiles_per_stride, src_width);
|
||||
|
||||
nir_def *tiled_xz = nir_vec4(&b, nir_channel(&b, tiled, 0),
|
||||
nir_channel(&b, tiled, 2), zero, zero);
|
||||
nir_def *tiled_yw = nir_vec4(&b, nir_channel(&b, tiled, 1),
|
||||
nir_channel(&b, tiled, 3), zero, zero);
|
||||
|
||||
nir_def *yval = nir_image_load(&b, 4, 32, zero, tiled_xz,
|
||||
zero /* sample */, zero /* lod */,
|
||||
.access = ACCESS_NON_WRITEABLE,
|
||||
.image_dim = GLSL_SAMPLER_DIM_2D,
|
||||
.image_array = false,
|
||||
.dest_type = nir_type_uint32);
|
||||
nir_def *uvval;
|
||||
|
||||
nir_def *dst_y_coord = nir_vec4(&b,
|
||||
nir_channel(&b, coord, 0),
|
||||
nir_channel(&b, coord, 1),
|
||||
zero, zero);
|
||||
/* store Y data */
|
||||
nir_def *img_deref_st_y = nir_imm_int(&b, 2);
|
||||
nir_image_store(&b, img_deref_st_y, dst_y_coord, zero /* sample */,
|
||||
yval, zero /* lod */,
|
||||
.access = ACCESS_NON_READABLE,
|
||||
.image_dim = GLSL_SAMPLER_DIM_2D,
|
||||
.image_array = false, .src_type = nir_type_uint32);
|
||||
|
||||
/* store UV data */
|
||||
nir_def *odd_even_line = nir_iand_imm(&b,
|
||||
nir_channel(&b, dst_y_coord, 1),
|
||||
1);
|
||||
nir_push_if(&b, nir_ieq_imm(&b, odd_even_line, 0));
|
||||
nir_def *in_bounds = nir_ball(&b, nir_ilt(&b, y_linear_coords, dst_extent));
|
||||
nir_push_if(&b, in_bounds);
|
||||
{
|
||||
if (tint_yuv) {
|
||||
/* use just blue for chroma */
|
||||
uvval = nir_imm_ivec4(&b, 0xc0, 0x80, 0xc0, 0x80);
|
||||
} else {
|
||||
nir_def *img_deref_uv = nir_imm_int(&b, 1);
|
||||
uvval = nir_image_load(&b, 4, 32, img_deref_uv, tiled_yw,
|
||||
zero /* sample */, zero /* lod */,
|
||||
.access = ACCESS_NON_WRITEABLE,
|
||||
.image_dim = GLSL_SAMPLER_DIM_2D,
|
||||
.image_array = false,
|
||||
.dest_type = nir_type_uint32);
|
||||
if (key->mtk_tiled.has_y)
|
||||
copy_y_uv_texel(&b, 0, y_tiled_coords, 2, y_linear_coords);
|
||||
|
||||
if (key->mtk_tiled.has_uv) {
|
||||
if (!tint_yuv) {
|
||||
copy_y_uv_texel(&b, 1, uv_tiled_coords, 3, uv_linear_coords);
|
||||
} else {
|
||||
/* use just blue for chroma */
|
||||
nir_def *val = nir_imm_ivec4(&b, 0xc0, 0x80, 0xc0, 0x80);
|
||||
nir_def *sample = nir_imm_int(&b, 0);
|
||||
nir_def *lod = nir_imm_int(&b, 0);
|
||||
|
||||
nir_image_store(&b, nir_imm_int(&b, 3), uv_linear_coords, sample,
|
||||
val, lod, .access = ACCESS_NON_READABLE,
|
||||
.image_dim = GLSL_SAMPLER_DIM_2D,
|
||||
.image_array = false, .src_type = nir_type_uint32);
|
||||
}
|
||||
}
|
||||
|
||||
/* Next line of Y components, UV is vertically subsampled. */
|
||||
if (key->mtk_tiled.has_y) {
|
||||
y_linear_coords =
|
||||
nir_iadd(&b, y_linear_coords, nir_imm_ivec2(&b, 0, 1));
|
||||
y_tiled_coords = nir_iadd(&b, y_tiled_coords, nir_imm_ivec2(&b, 4, 0));
|
||||
copy_y_uv_texel(&b, 0, y_tiled_coords, 2, y_linear_coords);
|
||||
}
|
||||
nir_def *dst_uv_coord = nir_ishr(&b, dst_y_coord,
|
||||
nir_imm_ivec4(&b, 0, 1, 0, 0));
|
||||
nir_def *img_deref_st_uv = nir_imm_int(&b, 3);
|
||||
nir_image_store(&b, img_deref_st_uv, dst_uv_coord, zero /* sample */,
|
||||
uvval, zero /* lod */,
|
||||
.access = ACCESS_NON_READABLE,
|
||||
.image_dim = GLSL_SAMPLER_DIM_2D,
|
||||
.image_array = false, .src_type = nir_type_uint32);
|
||||
}
|
||||
nir_pop_if(&b, NULL);
|
||||
|
||||
|
|
@ -510,10 +471,15 @@ panfrost_get_afbc_pack_shaders(struct panfrost_context *ctx,
|
|||
}
|
||||
|
||||
struct pan_mod_convert_shader_data *
|
||||
panfrost_get_mtk_detile_shader(struct panfrost_context *ctx)
|
||||
panfrost_get_mtk_detile_shader(struct panfrost_context *ctx, bool has_y,
|
||||
bool has_uv)
|
||||
{
|
||||
struct pan_mod_convert_shader_key key = {
|
||||
.mod = DRM_FORMAT_MOD_MTK_16L_32S_TILE,
|
||||
.mtk_tiled = {
|
||||
.has_y = has_y,
|
||||
.has_uv = has_uv,
|
||||
},
|
||||
};
|
||||
|
||||
return get_mod_convert_shaders(ctx, &key);
|
||||
|
|
|
|||
|
|
@ -41,6 +41,11 @@ struct pan_mod_convert_shader_key {
|
|||
unsigned bpp;
|
||||
unsigned align;
|
||||
} afbc;
|
||||
struct {
|
||||
unsigned has_y : 1;
|
||||
unsigned has_uv : 1;
|
||||
unsigned unused : 30;
|
||||
} mtk_tiled;
|
||||
};
|
||||
};
|
||||
|
||||
|
|
@ -83,10 +88,10 @@ struct panfrost_afbc_pack_info {
|
|||
} PACKED;
|
||||
|
||||
struct panfrost_mtk_detile_info {
|
||||
uint32_t tiles_per_stride;
|
||||
uint32_t src_width;
|
||||
uint32_t src_height;
|
||||
uint32_t dst_stride;
|
||||
uint32_t src_y_row_stride_tl;
|
||||
uint32_t src_uv_row_stride_tl;
|
||||
uint32_t width;
|
||||
uint32_t height;
|
||||
} PACKED;
|
||||
|
||||
void panfrost_afbc_context_init(struct panfrost_context *ctx);
|
||||
|
|
@ -97,7 +102,8 @@ panfrost_get_afbc_pack_shaders(struct panfrost_context *ctx,
|
|||
struct panfrost_resource *rsrc, unsigned align);
|
||||
|
||||
struct pan_mod_convert_shader_data *
|
||||
panfrost_get_mtk_detile_shader(struct panfrost_context *ctx);
|
||||
panfrost_get_mtk_detile_shader(struct panfrost_context *ctx, bool has_y,
|
||||
bool has_uv);
|
||||
|
||||
#define drm_is_mtk_tiled(mod) \
|
||||
((mod >> 52) == (0 | (DRM_FORMAT_MOD_VENDOR_MTK << 4)))
|
||||
|
|
|
|||
|
|
@ -132,6 +132,54 @@ panfrost_resource_init_image(struct panfrost_resource *rsc,
|
|||
plane->image = rsc->image;
|
||||
}
|
||||
|
||||
static bool
|
||||
adjust_mtk_tiled_props(struct panfrost_resource *rsc,
|
||||
struct pan_image_props *iprops, unsigned plane_idx,
|
||||
struct pan_image_layout_constraints *explicit_layout)
|
||||
{
|
||||
bool is_uv_plane =
|
||||
iprops->format == PIPE_FORMAT_R8G8_UNORM ||
|
||||
(iprops->format == PIPE_FORMAT_R8_G8B8_420_UNORM && plane_idx > 0);
|
||||
unsigned tile_w_px, tile_h_px, blksz_B;
|
||||
|
||||
if (is_uv_plane) {
|
||||
tile_w_px = 8;
|
||||
tile_h_px = 16;
|
||||
blksz_B = 2;
|
||||
iprops->format = PIPE_FORMAT_R8G8_UNORM;
|
||||
} else {
|
||||
tile_w_px = 16;
|
||||
tile_h_px = 32;
|
||||
blksz_B = 1;
|
||||
iprops->format = PIPE_FORMAT_R8_UNORM;
|
||||
}
|
||||
|
||||
/* SW detiling on MTK_TILED resources. This forces us to treat such
|
||||
* resources as linear images with:
|
||||
* width = tile_width * tile_height
|
||||
* height = (wsi_row_stride / (tile_width * blksize)) * (height /
|
||||
* tile_height)
|
||||
*/
|
||||
iprops->extent_px.width = tile_w_px * tile_h_px;
|
||||
iprops->extent_px.height =
|
||||
(explicit_layout->wsi_row_pitch_B / (blksz_B * tile_w_px)) *
|
||||
DIV_ROUND_UP(rsc->base.height0, tile_h_px);
|
||||
|
||||
/* Reject the import if the pitch is not aligned on a tile or if it's not
|
||||
* covering the resource width. */
|
||||
unsigned min_row_pitch_B = rsc->base.width0 * blksz_B;
|
||||
unsigned row_pitch_align_req_B = blksz_B * tile_w_px;
|
||||
|
||||
if (explicit_layout->strict &&
|
||||
(explicit_layout->wsi_row_pitch_B % row_pitch_align_req_B != 0 ||
|
||||
explicit_layout->wsi_row_pitch_B < min_row_pitch_B))
|
||||
return false;
|
||||
|
||||
/* Now adjust the row pitch. */
|
||||
explicit_layout->wsi_row_pitch_B = iprops->extent_px.width * blksz_B;
|
||||
return true;
|
||||
}
|
||||
|
||||
static struct pipe_resource *
|
||||
panfrost_resource_from_handle(struct pipe_screen *pscreen,
|
||||
const struct pipe_resource *templat,
|
||||
|
|
@ -190,6 +238,14 @@ panfrost_resource_from_handle(struct pipe_screen *pscreen,
|
|||
.nr_samples = MAX2(prsc->nr_samples, 1),
|
||||
.nr_slices = 1,
|
||||
};
|
||||
|
||||
if (drm_is_mtk_tiled(mod) &&
|
||||
!adjust_mtk_tiled_props(rsc, &iprops, whandle->plane,
|
||||
&explicit_layout)) {
|
||||
FREE(rsc);
|
||||
return NULL;
|
||||
}
|
||||
|
||||
unsigned format_plane =
|
||||
util_format_get_num_planes(iprops.format) > 1 ? whandle->plane : 0;
|
||||
|
||||
|
|
@ -266,6 +322,26 @@ panfrost_resource_get_handle(struct pipe_screen *pscreen,
|
|||
&rsrc->image.planes[handle->plane]->layout, 0);
|
||||
handle->offset =
|
||||
pan_image_get_wsi_offset(&rsrc->image.planes[handle->plane]->layout, 0);
|
||||
|
||||
/* SW detiling on MTK_TILED resources. This forces us to treat such
|
||||
* resources as linear images with:
|
||||
* width = tile_width * tile_height
|
||||
* height = (wsi_row_stride / (tile_width * blksize)) * (height / tile_height)
|
||||
*
|
||||
* We need to extract the original WSI row pitch from this.
|
||||
*/
|
||||
if (drm_is_mtk_tiled(rsrc->modifier)) {
|
||||
bool subsamp = handle->plane > 0 ||
|
||||
rsrc->image.props.format == PIPE_FORMAT_R8G8_UNORM;
|
||||
unsigned blksz_B = subsamp ? 2 : 1;
|
||||
unsigned tile_w_px = 16 / (subsamp ? 2 : 1);
|
||||
unsigned tile_h_px = 32 / (subsamp ? 2 : 1);
|
||||
unsigned row_stride_tl = rsrc->image.props.extent_px.height /
|
||||
DIV_ROUND_UP(rsrc->base.height0, tile_h_px);
|
||||
|
||||
handle->stride = row_stride_tl * tile_w_px * blksz_B;
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
|
|
@ -978,6 +1054,72 @@ panfrost_resource_create_with_modifiers(struct pipe_screen *screen,
|
|||
return panfrost_resource_create(screen, template);
|
||||
}
|
||||
|
||||
void
|
||||
panfrost_resource_change_format(struct panfrost_resource *rsrc,
|
||||
enum pipe_format new_format,
|
||||
struct panfrost_resource *save)
|
||||
{
|
||||
if (!rsrc)
|
||||
return;
|
||||
|
||||
assert(rsrc->image.props.modifier == DRM_FORMAT_MOD_LINEAR);
|
||||
assert(util_format_get_num_planes(new_format) == 1);
|
||||
assert(util_format_get_blockwidth(new_format) == 1 &&
|
||||
util_format_get_blockheight(new_format) == 1);
|
||||
assert(util_format_get_blockwidth(rsrc->image.props.format) == 1 &&
|
||||
util_format_get_blockheight(rsrc->image.props.format) == 1);
|
||||
|
||||
if (new_format == rsrc->image.props.format)
|
||||
return;
|
||||
|
||||
*save = *rsrc;
|
||||
|
||||
unsigned old_res_plane_idx = pan_resource_plane_index(rsrc);
|
||||
enum pipe_format old_format =
|
||||
util_format_get_plane_format(rsrc->image.props.format, old_res_plane_idx);
|
||||
unsigned old_width =
|
||||
util_format_get_plane_width(rsrc->image.props.format, old_res_plane_idx,
|
||||
rsrc->image.props.extent_px.width);
|
||||
|
||||
unsigned old_fmt_blksize = util_format_get_blocksize(old_format);
|
||||
unsigned new_fmt_blksize = util_format_get_blocksize(new_format);
|
||||
|
||||
if (old_fmt_blksize != new_fmt_blksize) {
|
||||
assert((old_fmt_blksize * rsrc->base.width0) % new_fmt_blksize == 0);
|
||||
rsrc->base.width0 =
|
||||
(old_fmt_blksize * rsrc->base.width0) / new_fmt_blksize;
|
||||
rsrc->image.props.extent_px.width =
|
||||
(old_fmt_blksize * old_width) /
|
||||
new_fmt_blksize;
|
||||
rsrc->image.props.extent_px.height =
|
||||
util_format_get_plane_height(rsrc->image.props.format, old_res_plane_idx,
|
||||
rsrc->image.props.extent_px.height);
|
||||
}
|
||||
|
||||
rsrc->base.next = NULL;
|
||||
rsrc->base.format = new_format;
|
||||
rsrc->image.props.format = new_format;
|
||||
rsrc->image.planes[0] = &rsrc->plane;
|
||||
rsrc->image.planes[1] = NULL;
|
||||
rsrc->image.planes[2] = NULL;
|
||||
}
|
||||
|
||||
void
|
||||
panfrost_resource_restore_format(struct panfrost_resource *rsrc,
|
||||
const struct panfrost_resource *saved)
|
||||
{
|
||||
if (!rsrc)
|
||||
return;
|
||||
|
||||
rsrc->base.next = saved->base.next;
|
||||
memcpy(rsrc->image.planes, saved->image.planes, sizeof(rsrc->image.planes));
|
||||
rsrc->base.format = saved->base.format;
|
||||
rsrc->image.props.format = saved->image.props.format;
|
||||
rsrc->base.width0 = saved->base.width0;
|
||||
rsrc->image.props.extent_px.width = saved->image.props.extent_px.width;
|
||||
rsrc->image.props.extent_px.height = saved->image.props.extent_px.height;
|
||||
}
|
||||
|
||||
static void
|
||||
panfrost_resource_destroy(struct pipe_screen *screen, struct pipe_resource *pt)
|
||||
{
|
||||
|
|
|
|||
|
|
@ -197,6 +197,15 @@ panfrost_translate_texture_dimension(enum pipe_texture_target t)
|
|||
}
|
||||
}
|
||||
|
||||
void
|
||||
panfrost_resource_change_format(struct panfrost_resource *rsrc,
|
||||
enum pipe_format new_format,
|
||||
struct panfrost_resource *save);
|
||||
|
||||
void
|
||||
panfrost_resource_restore_format(struct panfrost_resource *rsrc,
|
||||
const struct panfrost_resource *saved);
|
||||
|
||||
bool panfrost_should_pack_afbc(struct panfrost_device *dev,
|
||||
const struct panfrost_resource *rsrc);
|
||||
|
||||
|
|
|
|||
Loading…
Add table
Reference in a new issue