1 /*
2 * Copyright 2021 Alyssa Rosenzweig
3 * Copyright 2020-2021 Collabora, Ltd.
4 * Copyright 2019 Sonny Jiang <[email protected]>
5 * Copyright 2019 Advanced Micro Devices, Inc.
6 * Copyright 2014 Broadcom
7 * SPDX-License-Identifier: MIT
8 */
9
10 #include <stdint.h>
11 #include "asahi/layout/layout.h"
12 #include "asahi/lib/agx_nir_passes.h"
13 #include "compiler/nir/nir_builder.h"
14 #include "compiler/nir/nir_format_convert.h"
15 #include "gallium/auxiliary/util/u_blitter.h"
16 #include "gallium/auxiliary/util/u_dump.h"
17 #include "nir/pipe_nir.h"
18 #include "pipe/p_context.h"
19 #include "pipe/p_defines.h"
20 #include "pipe/p_state.h"
21 #include "util/format/u_format.h"
22 #include "util/format/u_formats.h"
23 #include "util/hash_table.h"
24 #include "util/macros.h"
25 #include "util/ralloc.h"
26 #include "util/u_sampler.h"
27 #include "util/u_surface.h"
28 #include "agx_state.h"
29 #include "glsl_types.h"
30 #include "nir.h"
31 #include "nir_builder_opcodes.h"
32 #include "shader_enums.h"
33
34 /* For block based blit kernels, we hardcode the maximum tile size which we can
35 * always achieve. This simplifies our life.
36 */
37 #define TILE_WIDTH 32
38 #define TILE_HEIGHT 32
39
40 static enum pipe_format
effective_format(enum pipe_format format)41 effective_format(enum pipe_format format)
42 {
43 switch (format) {
44 case PIPE_FORMAT_Z32_FLOAT:
45 case PIPE_FORMAT_Z24X8_UNORM:
46 return PIPE_FORMAT_R32_FLOAT;
47 case PIPE_FORMAT_Z16_UNORM:
48 return PIPE_FORMAT_R16_UNORM;
49 case PIPE_FORMAT_S8_UINT:
50 return PIPE_FORMAT_R8_UINT;
51 default:
52 return format;
53 }
54 }
55
56 static void *
asahi_blit_compute_shader(struct pipe_context * ctx,struct asahi_blit_key * key)57 asahi_blit_compute_shader(struct pipe_context *ctx, struct asahi_blit_key *key)
58 {
59 const nir_shader_compiler_options *options =
60 ctx->screen->get_compiler_options(ctx->screen, PIPE_SHADER_IR_NIR,
61 PIPE_SHADER_COMPUTE);
62
63 nir_builder b_ =
64 nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, options, "blit_cs");
65 nir_builder *b = &b_;
66 b->shader->info.workgroup_size[0] = TILE_WIDTH;
67 b->shader->info.workgroup_size[1] = TILE_HEIGHT;
68 b->shader->info.num_ubos = 1;
69
70 BITSET_SET(b->shader->info.textures_used, 0);
71 BITSET_SET(b->shader->info.samplers_used, 0);
72 BITSET_SET(b->shader->info.images_used, 0);
73
74 nir_def *zero = nir_imm_int(b, 0);
75
76 nir_def *params[4];
77 b->shader->num_uniforms = ARRAY_SIZE(params);
78 for (unsigned i = 0; i < b->shader->num_uniforms; ++i) {
79 params[i] = nir_load_ubo(b, 2, 32, zero, nir_imm_int(b, i * 8),
80 .align_mul = 4, .range = ~0);
81 }
82
83 nir_def *trans_offs = params[0];
84 nir_def *trans_scale = params[1];
85 nir_def *dst_offs_2d = params[2];
86 nir_def *dimensions_el_2d = params[3];
87
88 nir_def *phys_id_el_nd = nir_trim_vector(
89 b, nir_load_global_invocation_id(b, 32), key->array ? 3 : 2);
90 nir_def *phys_id_el_2d = nir_trim_vector(b, phys_id_el_nd, 2);
91 nir_def *layer = key->array ? nir_channel(b, phys_id_el_nd, 2) : NULL;
92
93 /* Offset within the tile. We're dispatched for the entire tile but the
94 * beginning might be out-of-bounds, so fix up.
95 */
96 nir_def *offs_in_tile_el_2d = nir_iand_imm(b, dst_offs_2d, 31);
97 nir_def *logical_id_el_2d = nir_isub(b, phys_id_el_2d, offs_in_tile_el_2d);
98
99 nir_def *image_pos_2d = nir_iadd(b, logical_id_el_2d, dst_offs_2d);
100 nir_def *image_pos_nd = image_pos_2d;
101 if (layer) {
102 image_pos_nd =
103 nir_vector_insert_imm(b, nir_pad_vector(b, image_pos_nd, 3), layer, 2);
104 }
105
106 nir_def *in_bounds;
107 if (key->aligned) {
108 in_bounds = nir_imm_true(b);
109 } else {
110 in_bounds = nir_ige(b, logical_id_el_2d, nir_imm_ivec2(b, 0, 0));
111 in_bounds =
112 nir_iand(b, in_bounds, nir_ilt(b, logical_id_el_2d, dimensions_el_2d));
113 }
114
115 nir_def *colour0, *colour1;
116 nir_push_if(b, nir_ball(b, in_bounds));
117 {
118 /* For pixels within the copy area, texture from the source */
119 nir_def *coords_el_2d =
120 nir_ffma(b, nir_u2f32(b, logical_id_el_2d), trans_scale, trans_offs);
121
122 nir_def *coords_el_nd = coords_el_2d;
123 if (layer) {
124 coords_el_nd = nir_vector_insert_imm(
125 b, nir_pad_vector(b, coords_el_nd, 3), nir_u2f32(b, layer), 2);
126 }
127
128 nir_tex_instr *tex = nir_tex_instr_create(b->shader, 1);
129 tex->dest_type = nir_type_uint32; /* irrelevant */
130 tex->sampler_dim = GLSL_SAMPLER_DIM_2D;
131 tex->is_array = key->array;
132 tex->op = nir_texop_tex;
133 tex->src[0] = nir_tex_src_for_ssa(nir_tex_src_coord, coords_el_nd);
134 tex->backend_flags = AGX_TEXTURE_FLAG_NO_CLAMP;
135 tex->coord_components = coords_el_nd->num_components;
136 tex->texture_index = 0;
137 tex->sampler_index = 0;
138 nir_def_init(&tex->instr, &tex->def, 4, 32);
139 nir_builder_instr_insert(b, &tex->instr);
140 colour0 = &tex->def;
141 }
142 nir_push_else(b, NULL);
143 {
144 /* For out-of-bounds pixels, copy in the destination */
145 colour1 = nir_image_load(
146 b, 4, 32, nir_imm_int(b, 0), nir_pad_vec4(b, image_pos_nd), zero, zero,
147 .image_array = key->array, .image_dim = GLSL_SAMPLER_DIM_2D,
148 .access = ACCESS_IN_BOUNDS_AGX, .dest_type = nir_type_uint32);
149 }
150 nir_pop_if(b, NULL);
151 nir_def *color = nir_if_phi(b, colour0, colour1);
152
153 enum asahi_blit_clamp clamp = ASAHI_BLIT_CLAMP_NONE;
154 bool src_sint = util_format_is_pure_sint(key->src_format);
155 bool dst_sint = util_format_is_pure_sint(key->dst_format);
156 if (util_format_is_pure_integer(key->src_format) &&
157 util_format_is_pure_integer(key->dst_format)) {
158
159 if (src_sint && !dst_sint)
160 clamp = ASAHI_BLIT_CLAMP_SINT_TO_UINT;
161 else if (!src_sint && dst_sint)
162 clamp = ASAHI_BLIT_CLAMP_UINT_TO_SINT;
163 }
164
165 if (clamp == ASAHI_BLIT_CLAMP_SINT_TO_UINT)
166 color = nir_imax(b, color, nir_imm_int(b, 0));
167 else if (clamp == ASAHI_BLIT_CLAMP_UINT_TO_SINT)
168 color = nir_umin(b, color, nir_imm_int(b, INT32_MAX));
169
170 nir_def *local_offset = nir_imm_intN_t(b, 0, 16);
171 nir_def *lid = nir_trim_vector(b, nir_load_local_invocation_id(b), 2);
172 lid = nir_u2u16(b, lid);
173
174 /* Pure integer formatss need to be clamped in software, at least in some
175 * cases. We do so on store. Piglit gl-3.0-render-integer checks this, as
176 * does KHR-GL33.packed_pixels.*.
177 *
178 * TODO: Make this common code somehow.
179 */
180 const struct util_format_description *desc =
181 util_format_description(key->dst_format);
182 unsigned c = util_format_get_first_non_void_channel(key->dst_format);
183
184 if (desc->channel[c].size <= 16 &&
185 util_format_is_pure_integer(key->dst_format)) {
186
187 unsigned bits[4] = {
188 desc->channel[0].size ?: desc->channel[0].size,
189 desc->channel[1].size ?: desc->channel[0].size,
190 desc->channel[2].size ?: desc->channel[0].size,
191 desc->channel[3].size ?: desc->channel[0].size,
192 };
193
194 if (util_format_is_pure_sint(key->dst_format))
195 color = nir_format_clamp_sint(b, color, bits);
196 else
197 color = nir_format_clamp_uint(b, color, bits);
198
199 color = nir_u2u16(b, color);
200 }
201
202 /* The source texel has been converted into a 32-bit value. We need to
203 * convert it to a tilebuffer format that can then be converted to the
204 * destination format in the PBE hardware. That's the renderable format for
205 * the destination format, which must exist along this path. This mirrors the
206 * flow of fragment and end-of-tile shaders.
207 */
208 enum pipe_format tib_format =
209 ail_pixel_format[effective_format(key->dst_format)].renderable;
210
211 nir_store_local_pixel_agx(b, color, nir_imm_int(b, 1), lid, .base = 0,
212 .write_mask = 0xf, .format = tib_format,
213 .explicit_coord = true);
214
215 nir_barrier(b, .execution_scope = SCOPE_WORKGROUP);
216
217 nir_push_if(b, nir_ball(b, nir_ieq_imm(b, lid, 0)));
218 {
219 nir_def *pbe_index = nir_imm_intN_t(b, 2, 16);
220 nir_image_store_block_agx(
221 b, pbe_index, local_offset, image_pos_nd, .format = tib_format,
222 .image_dim = GLSL_SAMPLER_DIM_2D, .image_array = key->array,
223 .explicit_coord = true);
224 }
225 nir_pop_if(b, NULL);
226 b->shader->info.cs.image_block_size_per_thread_agx =
227 util_format_get_blocksize(key->dst_format);
228
229 return pipe_shader_from_nir(ctx, b->shader);
230 }
231
232 static bool
asahi_compute_blit_supported(const struct pipe_blit_info * info)233 asahi_compute_blit_supported(const struct pipe_blit_info *info)
234 {
235 return (info->src.box.depth == info->dst.box.depth) && !info->alpha_blend &&
236 !info->num_window_rectangles && !info->sample0_only &&
237 !info->scissor_enable && !info->window_rectangle_include &&
238 info->src.resource->nr_samples <= 1 &&
239 info->dst.resource->nr_samples <= 1 &&
240 !util_format_is_depth_and_stencil(info->src.format) &&
241 !util_format_is_depth_and_stencil(info->dst.format) &&
242 info->src.box.depth >= 0 &&
243 info->mask == util_format_get_mask(info->src.format) &&
244 /* XXX: texsubimage pbo failing otherwise, needs investigation */
245 info->dst.format != PIPE_FORMAT_B5G6R5_UNORM &&
246 info->dst.format != PIPE_FORMAT_B5G5R5A1_UNORM &&
247 info->dst.format != PIPE_FORMAT_B5G5R5X1_UNORM &&
248 info->dst.format != PIPE_FORMAT_R5G6B5_UNORM &&
249 info->dst.format != PIPE_FORMAT_R5G5B5A1_UNORM &&
250 info->dst.format != PIPE_FORMAT_R5G5B5X1_UNORM;
251 }
252
253 static void
asahi_compute_save(struct agx_context * ctx)254 asahi_compute_save(struct agx_context *ctx)
255 {
256 struct asahi_blitter *blitter = &ctx->compute_blitter;
257 struct agx_stage *stage = &ctx->stage[PIPE_SHADER_COMPUTE];
258
259 assert(!blitter->active && "recursion detected, driver bug");
260
261 pipe_resource_reference(&blitter->saved_cb.buffer, stage->cb[0].buffer);
262 memcpy(&blitter->saved_cb, &stage->cb[0],
263 sizeof(struct pipe_constant_buffer));
264
265 blitter->has_saved_image = stage->image_mask & BITFIELD_BIT(0);
266 if (blitter->has_saved_image) {
267 pipe_resource_reference(&blitter->saved_image.resource,
268 stage->images[0].resource);
269 memcpy(&blitter->saved_image, &stage->images[0],
270 sizeof(struct pipe_image_view));
271 }
272
273 pipe_sampler_view_reference(&blitter->saved_sampler_view,
274 &stage->textures[0]->base);
275
276 blitter->saved_num_sampler_states = stage->sampler_count;
277 memcpy(blitter->saved_sampler_states, stage->samplers,
278 stage->sampler_count * sizeof(void *));
279
280 blitter->saved_cs = stage->shader;
281 blitter->active = true;
282 }
283
284 static void
asahi_compute_restore(struct agx_context * ctx)285 asahi_compute_restore(struct agx_context *ctx)
286 {
287 struct pipe_context *pctx = &ctx->base;
288 struct asahi_blitter *blitter = &ctx->compute_blitter;
289
290 if (blitter->has_saved_image) {
291 pctx->set_shader_images(pctx, PIPE_SHADER_COMPUTE, 0, 1, 0,
292 &blitter->saved_image);
293 pipe_resource_reference(&blitter->saved_image.resource, NULL);
294 }
295
296 /* take_ownership=true so do not unreference */
297 pctx->set_constant_buffer(pctx, PIPE_SHADER_COMPUTE, 0, true,
298 &blitter->saved_cb);
299 blitter->saved_cb.buffer = NULL;
300
301 if (blitter->saved_sampler_view) {
302 pctx->set_sampler_views(pctx, PIPE_SHADER_COMPUTE, 0, 1, 0, true,
303 &blitter->saved_sampler_view);
304
305 blitter->saved_sampler_view = NULL;
306 }
307
308 if (blitter->saved_num_sampler_states) {
309 pctx->bind_sampler_states(pctx, PIPE_SHADER_COMPUTE, 0,
310 blitter->saved_num_sampler_states,
311 blitter->saved_sampler_states);
312 }
313
314 pctx->bind_compute_state(pctx, blitter->saved_cs);
315 blitter->saved_cs = NULL;
316 blitter->active = false;
317 }
318
319 static void
asahi_compute_blit(struct pipe_context * ctx,const struct pipe_blit_info * info,struct asahi_blitter * blitter)320 asahi_compute_blit(struct pipe_context *ctx, const struct pipe_blit_info *info,
321 struct asahi_blitter *blitter)
322 {
323 if (info->src.box.width == 0 || info->src.box.height == 0 ||
324 info->dst.box.width == 0 || info->dst.box.height == 0)
325 return;
326
327 assert(asahi_compute_blit_supported(info));
328 asahi_compute_save(agx_context(ctx));
329
330 unsigned depth = info->dst.box.depth;
331 bool array = depth > 1;
332
333 struct pipe_resource *src = info->src.resource;
334 struct pipe_resource *dst = info->dst.resource;
335 struct pipe_sampler_view src_templ = {0}, *src_view;
336
337 float src_width = (float)u_minify(src->width0, info->src.level);
338 float src_height = (float)u_minify(src->height0, info->src.level);
339
340 float x_scale =
341 (info->src.box.width / (float)info->dst.box.width) / src_width;
342
343 float y_scale =
344 (info->src.box.height / (float)info->dst.box.height) / src_height;
345
346 /* Expand the grid so destinations are in tiles */
347 unsigned expanded_x0 = info->dst.box.x & ~(TILE_WIDTH - 1);
348 unsigned expanded_y0 = info->dst.box.y & ~(TILE_HEIGHT - 1);
349 unsigned expanded_x1 =
350 align(info->dst.box.x + info->dst.box.width, TILE_WIDTH);
351 unsigned expanded_y1 =
352 align(info->dst.box.y + info->dst.box.height, TILE_HEIGHT);
353
354 /* But clamp to the destination size to save some redundant threads */
355 expanded_x1 =
356 MIN2(expanded_x1, u_minify(info->dst.resource->width0, info->dst.level));
357 expanded_y1 =
358 MIN2(expanded_y1, u_minify(info->dst.resource->height0, info->dst.level));
359
360 /* Calculate the width/height based on the expanded grid */
361 unsigned width = expanded_x1 - expanded_x0;
362 unsigned height = expanded_y1 - expanded_y0;
363
364 unsigned data[] = {
365 fui(0.5f * x_scale + (float)info->src.box.x / src_width),
366 fui(0.5f * y_scale + (float)info->src.box.y / src_height),
367 fui(x_scale),
368 fui(y_scale),
369 info->dst.box.x,
370 info->dst.box.y,
371 info->dst.box.width,
372 info->dst.box.height,
373 };
374
375 struct pipe_constant_buffer cb = {
376 .buffer_size = sizeof(data),
377 .user_buffer = data,
378 };
379 ctx->set_constant_buffer(ctx, PIPE_SHADER_COMPUTE, 0, false, &cb);
380
381 struct pipe_image_view image = {
382 .resource = dst,
383 .access = PIPE_IMAGE_ACCESS_WRITE | PIPE_IMAGE_ACCESS_DRIVER_INTERNAL,
384 .shader_access = PIPE_IMAGE_ACCESS_WRITE,
385 .format = info->dst.format,
386 .u.tex.level = info->dst.level,
387 .u.tex.first_layer = info->dst.box.z,
388 .u.tex.last_layer = info->dst.box.z + depth - 1,
389 .u.tex.single_layer_view = !array,
390 };
391 ctx->set_shader_images(ctx, PIPE_SHADER_COMPUTE, 0, 1, 0, &image);
392
393 if (!blitter->sampler[info->filter]) {
394 struct pipe_sampler_state sampler_state = {
395 .wrap_s = PIPE_TEX_WRAP_CLAMP_TO_EDGE,
396 .wrap_t = PIPE_TEX_WRAP_CLAMP_TO_EDGE,
397 .wrap_r = PIPE_TEX_WRAP_CLAMP_TO_EDGE,
398 .min_img_filter = info->filter,
399 .mag_img_filter = info->filter,
400 .compare_func = PIPE_FUNC_ALWAYS,
401 .seamless_cube_map = true,
402 .max_lod = 31.0f,
403 };
404
405 blitter->sampler[info->filter] =
406 ctx->create_sampler_state(ctx, &sampler_state);
407 }
408
409 ctx->bind_sampler_states(ctx, PIPE_SHADER_COMPUTE, 0, 1,
410 &blitter->sampler[info->filter]);
411
412 /* Initialize the sampler view. */
413 u_sampler_view_default_template(&src_templ, src, src->format);
414 src_templ.format = info->src.format;
415 src_templ.target = array ? PIPE_TEXTURE_2D_ARRAY : PIPE_TEXTURE_2D;
416 src_templ.swizzle_r = PIPE_SWIZZLE_X;
417 src_templ.swizzle_g = PIPE_SWIZZLE_Y;
418 src_templ.swizzle_b = PIPE_SWIZZLE_Z;
419 src_templ.swizzle_a = PIPE_SWIZZLE_W;
420 src_templ.u.tex.first_layer = info->src.box.z;
421 src_templ.u.tex.last_layer = info->src.box.z + depth - 1;
422 src_templ.u.tex.first_level = info->src.level;
423 src_templ.u.tex.last_level = info->src.level;
424 src_view = ctx->create_sampler_view(ctx, src, &src_templ);
425 ctx->set_sampler_views(ctx, PIPE_SHADER_COMPUTE, 0, 1, 0, true, &src_view);
426
427 struct asahi_blit_key key = {
428 .src_format = info->src.format,
429 .dst_format = info->dst.format,
430 .array = array,
431 .aligned = info->dst.box.width == width && info->dst.box.height == height,
432 };
433 struct hash_entry *ent = _mesa_hash_table_search(blitter->blit_cs, &key);
434 void *cs = NULL;
435
436 if (ent) {
437 cs = ent->data;
438 } else {
439 cs = asahi_blit_compute_shader(ctx, &key);
440 _mesa_hash_table_insert(
441 blitter->blit_cs, ralloc_memdup(blitter->blit_cs, &key, sizeof(key)),
442 cs);
443 }
444
445 assert(cs != NULL);
446 ctx->bind_compute_state(ctx, cs);
447
448 struct pipe_grid_info grid_info = {
449 .block = {TILE_WIDTH, TILE_HEIGHT, 1},
450 .last_block = {width % TILE_WIDTH, height % TILE_HEIGHT, 1},
451 .grid =
452 {
453 DIV_ROUND_UP(width, TILE_WIDTH),
454 DIV_ROUND_UP(height, TILE_HEIGHT),
455 depth,
456 },
457 };
458 ctx->launch_grid(ctx, &grid_info);
459 ctx->set_shader_images(ctx, PIPE_SHADER_COMPUTE, 0, 0, 1, NULL);
460 ctx->set_constant_buffer(ctx, PIPE_SHADER_COMPUTE, 0, false, NULL);
461 ctx->set_sampler_views(ctx, PIPE_SHADER_COMPUTE, 0, 0, 1, false, NULL);
462
463 asahi_compute_restore(agx_context(ctx));
464 }
465
466 void
agx_blitter_save(struct agx_context * ctx,struct blitter_context * blitter,bool render_cond)467 agx_blitter_save(struct agx_context *ctx, struct blitter_context *blitter,
468 bool render_cond)
469 {
470 util_blitter_save_vertex_buffers(blitter, ctx->vertex_buffers,
471 util_last_bit(ctx->vb_mask));
472 util_blitter_save_vertex_elements(blitter, ctx->attributes);
473 util_blitter_save_vertex_shader(blitter,
474 ctx->stage[PIPE_SHADER_VERTEX].shader);
475 util_blitter_save_tessctrl_shader(blitter,
476 ctx->stage[PIPE_SHADER_TESS_CTRL].shader);
477 util_blitter_save_tesseval_shader(blitter,
478 ctx->stage[PIPE_SHADER_TESS_EVAL].shader);
479 util_blitter_save_geometry_shader(blitter,
480 ctx->stage[PIPE_SHADER_GEOMETRY].shader);
481 util_blitter_save_rasterizer(blitter, ctx->rast);
482 util_blitter_save_viewport(blitter, &ctx->viewport[0]);
483 util_blitter_save_scissor(blitter, &ctx->scissor[0]);
484 util_blitter_save_fragment_shader(blitter,
485 ctx->stage[PIPE_SHADER_FRAGMENT].shader);
486 util_blitter_save_blend(blitter, ctx->blend);
487 util_blitter_save_depth_stencil_alpha(blitter, ctx->zs);
488 util_blitter_save_stencil_ref(blitter, &ctx->stencil_ref);
489 util_blitter_save_so_targets(blitter, ctx->streamout.num_targets,
490 ctx->streamout.targets);
491 util_blitter_save_sample_mask(blitter, ctx->sample_mask, 0);
492
493 util_blitter_save_framebuffer(blitter, &ctx->framebuffer);
494 util_blitter_save_fragment_sampler_states(
495 blitter, ctx->stage[PIPE_SHADER_FRAGMENT].sampler_count,
496 (void **)(ctx->stage[PIPE_SHADER_FRAGMENT].samplers));
497 util_blitter_save_fragment_sampler_views(
498 blitter, ctx->stage[PIPE_SHADER_FRAGMENT].texture_count,
499 (struct pipe_sampler_view **)ctx->stage[PIPE_SHADER_FRAGMENT].textures);
500 util_blitter_save_fragment_constant_buffer_slot(
501 blitter, ctx->stage[PIPE_SHADER_FRAGMENT].cb);
502
503 if (!render_cond) {
504 util_blitter_save_render_condition(blitter,
505 (struct pipe_query *)ctx->cond_query,
506 ctx->cond_cond, ctx->cond_mode);
507 }
508 }
509
510 void
agx_blit(struct pipe_context * pipe,const struct pipe_blit_info * info)511 agx_blit(struct pipe_context *pipe, const struct pipe_blit_info *info)
512 {
513 struct agx_context *ctx = agx_context(pipe);
514
515 if (info->render_condition_enable && !agx_render_condition_check(ctx))
516 return;
517
518 /* Legalize compression /before/ calling into u_blitter to avoid recursion.
519 * u_blitter bans recursive usage.
520 */
521 agx_legalize_compression(ctx, agx_resource(info->dst.resource),
522 info->dst.format);
523
524 agx_legalize_compression(ctx, agx_resource(info->src.resource),
525 info->src.format);
526
527 if (asahi_compute_blit_supported(info)) {
528 asahi_compute_blit(pipe, info, &ctx->compute_blitter);
529 return;
530 }
531
532 if (!util_blitter_is_blit_supported(ctx->blitter, info)) {
533 fprintf(stderr, "\n");
534 util_dump_blit_info(stderr, info);
535 fprintf(stderr, "\n\n");
536 unreachable("Unsupported blit");
537 }
538
539 /* Handle self-blits */
540 agx_flush_writer(ctx, agx_resource(info->dst.resource), "Blit");
541
542 agx_blitter_save(ctx, ctx->blitter, info->render_condition_enable);
543 util_blitter_blit(ctx->blitter, info, NULL);
544 }
545
546 static bool
try_copy_via_blit(struct pipe_context * pctx,struct pipe_resource * dst,unsigned dst_level,unsigned dstx,unsigned dsty,unsigned dstz,struct pipe_resource * src,unsigned src_level,const struct pipe_box * src_box)547 try_copy_via_blit(struct pipe_context *pctx, struct pipe_resource *dst,
548 unsigned dst_level, unsigned dstx, unsigned dsty,
549 unsigned dstz, struct pipe_resource *src, unsigned src_level,
550 const struct pipe_box *src_box)
551 {
552 struct agx_context *ctx = agx_context(pctx);
553
554 if (dst->target == PIPE_BUFFER)
555 return false;
556
557 /* TODO: Handle these for rusticl copies */
558 if (dst->target != src->target)
559 return false;
560
561 struct pipe_blit_info info = {
562 .dst =
563 {
564 .resource = dst,
565 .level = dst_level,
566 .box.x = dstx,
567 .box.y = dsty,
568 .box.z = dstz,
569 .box.width = src_box->width,
570 .box.height = src_box->height,
571 .box.depth = src_box->depth,
572 .format = dst->format,
573 },
574 .src =
575 {
576 .resource = src,
577 .level = src_level,
578 .box = *src_box,
579 .format = src->format,
580 },
581 .mask = util_format_get_mask(src->format),
582 .filter = PIPE_TEX_FILTER_NEAREST,
583 .scissor_enable = 0,
584 };
585
586 /* snorm formats don't round trip, so don't use them for copies */
587 if (util_format_is_snorm(info.dst.format))
588 info.dst.format = util_format_snorm_to_sint(info.dst.format);
589
590 if (util_format_is_snorm(info.src.format))
591 info.src.format = util_format_snorm_to_sint(info.src.format);
592
593 if (util_blitter_is_blit_supported(ctx->blitter, &info) &&
594 info.dst.format == info.src.format) {
595
596 agx_blit(pctx, &info);
597 return true;
598 } else {
599 return false;
600 }
601 }
602
603 void
agx_resource_copy_region(struct pipe_context * pctx,struct pipe_resource * dst,unsigned dst_level,unsigned dstx,unsigned dsty,unsigned dstz,struct pipe_resource * src,unsigned src_level,const struct pipe_box * src_box)604 agx_resource_copy_region(struct pipe_context *pctx, struct pipe_resource *dst,
605 unsigned dst_level, unsigned dstx, unsigned dsty,
606 unsigned dstz, struct pipe_resource *src,
607 unsigned src_level, const struct pipe_box *src_box)
608 {
609 if (try_copy_via_blit(pctx, dst, dst_level, dstx, dsty, dstz, src, src_level,
610 src_box))
611 return;
612
613 /* CPU fallback */
614 util_resource_copy_region(pctx, dst, dst_level, dstx, dsty, dstz, src,
615 src_level, src_box);
616 }
617