1 /*
2 * Copyright 2021 Alyssa Rosenzweig
3 * Copyright 2020-2021 Collabora, Ltd.
4 * Copyright 2019 Sonny Jiang <sonnyj608@gmail.com>
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/macros.h"
24 #include "util/u_sampler.h"
25 #include "util/u_surface.h"
26 #include "agx_formats.h"
27 #include "agx_state.h"
28 #include "shader_enums.h"
29
30 #define BLIT_WG_SIZE 32
31
32 static void *
asahi_blit_compute_shader(struct pipe_context * ctx,enum asahi_blit_clamp clamp,bool array)33 asahi_blit_compute_shader(struct pipe_context *ctx, enum asahi_blit_clamp clamp,
34 bool array)
35 {
36 const nir_shader_compiler_options *options =
37 ctx->screen->get_compiler_options(ctx->screen, PIPE_SHADER_IR_NIR,
38 PIPE_SHADER_COMPUTE);
39
40 nir_builder b_ =
41 nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, options, "blit_cs");
42 nir_builder *b = &b_;
43 b->shader->info.workgroup_size[0] = BLIT_WG_SIZE;
44 b->shader->info.workgroup_size[1] = BLIT_WG_SIZE;
45 b->shader->info.num_ubos = 1;
46
47 BITSET_SET(b->shader->info.textures_used, 0);
48 BITSET_SET(b->shader->info.samplers_used, 0);
49 BITSET_SET(b->shader->info.images_used, 0);
50
51 nir_def *zero = nir_imm_int(b, 0);
52
53 nir_def *params[3];
54 b->shader->num_uniforms = ARRAY_SIZE(params);
55 for (unsigned i = 0; i < b->shader->num_uniforms; ++i) {
56 params[i] = nir_load_ubo(b, 2, 32, zero, nir_imm_int(b, i * 8),
57 .align_mul = 4, .range = ~0);
58 }
59
60 nir_def *ids =
61 nir_trim_vector(b, nir_load_global_invocation_id(b, 32), array ? 3 : 2);
62
63 nir_def *tex_pos = nir_u2f32(b, ids);
64 nir_def *pos2 =
65 nir_ffma(b, nir_trim_vector(b, tex_pos, 2), params[1], params[0]);
66 if (array) {
67 tex_pos = nir_vector_insert_imm(b, nir_pad_vector(b, pos2, 3),
68 nir_channel(b, tex_pos, 2), 2);
69 } else {
70 tex_pos = pos2;
71 }
72
73 nir_tex_instr *tex = nir_tex_instr_create(b->shader, 1);
74 tex->dest_type = nir_type_uint32; /* irrelevant */
75 tex->sampler_dim = GLSL_SAMPLER_DIM_2D;
76 tex->is_array = array;
77 tex->op = nir_texop_tex;
78 tex->src[0] = nir_tex_src_for_ssa(nir_tex_src_coord, tex_pos);
79 tex->backend_flags = AGX_TEXTURE_FLAG_NO_CLAMP;
80 tex->coord_components = array ? 3 : 2;
81 tex->texture_index = 0;
82 tex->sampler_index = 0;
83 nir_def_init(&tex->instr, &tex->def, 4, 32);
84 nir_builder_instr_insert(b, &tex->instr);
85 nir_def *color = &tex->def;
86
87 if (clamp == ASAHI_BLIT_CLAMP_SINT_TO_UINT)
88 color = nir_imax(b, color, nir_imm_int(b, 0));
89 else if (clamp == ASAHI_BLIT_CLAMP_UINT_TO_SINT)
90 color = nir_umin(b, color, nir_imm_int(b, INT32_MAX));
91
92 nir_def *image_pos =
93 nir_iadd(b, ids, nir_pad_vector_imm_int(b, params[2], 0, array ? 3 : 2));
94
95 nir_image_store(b, nir_imm_int(b, 0), nir_pad_vec4(b, image_pos), zero,
96 color, zero, .image_dim = GLSL_SAMPLER_DIM_2D,
97 .access = ACCESS_NON_READABLE, .image_array = array);
98
99 return pipe_shader_from_nir(ctx, b->shader);
100 }
101
102 static bool
asahi_compute_blit_supported(const struct pipe_blit_info * info)103 asahi_compute_blit_supported(const struct pipe_blit_info *info)
104 {
105 return (info->src.box.depth == info->dst.box.depth) && !info->alpha_blend &&
106 !info->num_window_rectangles && !info->sample0_only &&
107 !info->scissor_enable && !info->window_rectangle_include &&
108 info->src.resource->nr_samples <= 1 &&
109 info->dst.resource->nr_samples <= 1 &&
110 !util_format_is_depth_and_stencil(info->src.format) &&
111 !util_format_is_depth_and_stencil(info->dst.format) &&
112 info->src.box.depth >= 0 &&
113 info->mask == util_format_get_mask(info->src.format) &&
114 /* XXX: texsubimage pbo failing otherwise, needs investigation */
115 info->dst.format != PIPE_FORMAT_B5G6R5_UNORM &&
116 info->dst.format != PIPE_FORMAT_B5G5R5A1_UNORM &&
117 info->dst.format != PIPE_FORMAT_B5G5R5X1_UNORM &&
118 info->dst.format != PIPE_FORMAT_R5G6B5_UNORM &&
119 info->dst.format != PIPE_FORMAT_R5G5B5A1_UNORM &&
120 info->dst.format != PIPE_FORMAT_R5G5B5X1_UNORM;
121 }
122
123 static void
asahi_compute_save(struct agx_context * ctx)124 asahi_compute_save(struct agx_context *ctx)
125 {
126 struct asahi_blitter *blitter = &ctx->compute_blitter;
127 struct agx_stage *stage = &ctx->stage[PIPE_SHADER_COMPUTE];
128
129 assert(!blitter->active && "recursion detected, driver bug");
130
131 pipe_resource_reference(&blitter->saved_cb.buffer, stage->cb[0].buffer);
132 memcpy(&blitter->saved_cb, &stage->cb[0],
133 sizeof(struct pipe_constant_buffer));
134
135 blitter->has_saved_image = stage->image_mask & BITFIELD_BIT(0);
136 if (blitter->has_saved_image) {
137 pipe_resource_reference(&blitter->saved_image.resource,
138 stage->images[0].resource);
139 memcpy(&blitter->saved_image, &stage->images[0],
140 sizeof(struct pipe_image_view));
141 }
142
143 pipe_sampler_view_reference(&blitter->saved_sampler_view,
144 &stage->textures[0]->base);
145
146 blitter->saved_num_sampler_states = stage->sampler_count;
147 memcpy(blitter->saved_sampler_states, stage->samplers,
148 stage->sampler_count * sizeof(void *));
149
150 blitter->saved_cs = stage->shader;
151 blitter->active = true;
152 }
153
154 static void
asahi_compute_restore(struct agx_context * ctx)155 asahi_compute_restore(struct agx_context *ctx)
156 {
157 struct pipe_context *pctx = &ctx->base;
158 struct asahi_blitter *blitter = &ctx->compute_blitter;
159
160 if (blitter->has_saved_image) {
161 pctx->set_shader_images(pctx, PIPE_SHADER_COMPUTE, 0, 1, 0,
162 &blitter->saved_image);
163 pipe_resource_reference(&blitter->saved_image.resource, NULL);
164 }
165
166 /* take_ownership=true so do not unreference */
167 pctx->set_constant_buffer(pctx, PIPE_SHADER_COMPUTE, 0, true,
168 &blitter->saved_cb);
169 blitter->saved_cb.buffer = NULL;
170
171 if (blitter->saved_sampler_view) {
172 pctx->set_sampler_views(pctx, PIPE_SHADER_COMPUTE, 0, 1, 0, true,
173 &blitter->saved_sampler_view);
174
175 blitter->saved_sampler_view = NULL;
176 }
177
178 if (blitter->saved_num_sampler_states) {
179 pctx->bind_sampler_states(pctx, PIPE_SHADER_COMPUTE, 0,
180 blitter->saved_num_sampler_states,
181 blitter->saved_sampler_states);
182 }
183
184 pctx->bind_compute_state(pctx, blitter->saved_cs);
185 blitter->saved_cs = NULL;
186 blitter->active = false;
187 }
188
189 static void
asahi_compute_blit(struct pipe_context * ctx,const struct pipe_blit_info * info,struct asahi_blitter * blitter)190 asahi_compute_blit(struct pipe_context *ctx, const struct pipe_blit_info *info,
191 struct asahi_blitter *blitter)
192 {
193 if (info->src.box.width == 0 || info->src.box.height == 0 ||
194 info->dst.box.width == 0 || info->dst.box.height == 0)
195 return;
196
197 assert(asahi_compute_blit_supported(info));
198 asahi_compute_save(agx_context(ctx));
199
200 unsigned depth = info->dst.box.depth;
201 bool array = depth > 1;
202
203 struct pipe_resource *src = info->src.resource;
204 struct pipe_resource *dst = info->dst.resource;
205 struct pipe_sampler_view src_templ = {0}, *src_view;
206 unsigned width = info->dst.box.width;
207 unsigned height = info->dst.box.height;
208
209 float src_width = (float)u_minify(src->width0, info->src.level);
210 float src_height = (float)u_minify(src->height0, info->src.level);
211
212 float x_scale = (info->src.box.width / (float)width) / src_width;
213 float y_scale = (info->src.box.height / (float)height) / src_height;
214
215 unsigned data[] = {
216 fui(0.5f * x_scale + (float)info->src.box.x / src_width),
217 fui(0.5f * y_scale + (float)info->src.box.y / src_height),
218 fui(x_scale),
219 fui(y_scale),
220 info->dst.box.x,
221 info->dst.box.y,
222 };
223
224 struct pipe_constant_buffer cb = {
225 .buffer_size = sizeof(data),
226 .user_buffer = data,
227 };
228 ctx->set_constant_buffer(ctx, PIPE_SHADER_COMPUTE, 0, false, &cb);
229
230 struct pipe_image_view image = {
231 .resource = dst,
232 .access = PIPE_IMAGE_ACCESS_WRITE | PIPE_IMAGE_ACCESS_DRIVER_INTERNAL,
233 .shader_access = PIPE_IMAGE_ACCESS_WRITE,
234 .format = info->dst.format,
235 .u.tex.level = info->dst.level,
236 .u.tex.first_layer = info->dst.box.z,
237 .u.tex.last_layer = info->dst.box.z + depth - 1,
238 .u.tex.single_layer_view = !array,
239 };
240 ctx->set_shader_images(ctx, PIPE_SHADER_COMPUTE, 0, 1, 0, &image);
241
242 if (!blitter->sampler[info->filter]) {
243 struct pipe_sampler_state sampler_state = {
244 .wrap_s = PIPE_TEX_WRAP_CLAMP_TO_EDGE,
245 .wrap_t = PIPE_TEX_WRAP_CLAMP_TO_EDGE,
246 .wrap_r = PIPE_TEX_WRAP_CLAMP_TO_EDGE,
247 .min_img_filter = info->filter,
248 .mag_img_filter = info->filter,
249 .compare_func = PIPE_FUNC_ALWAYS,
250 .seamless_cube_map = true,
251 .max_lod = 31.0f,
252 };
253
254 blitter->sampler[info->filter] =
255 ctx->create_sampler_state(ctx, &sampler_state);
256 }
257
258 ctx->bind_sampler_states(ctx, PIPE_SHADER_COMPUTE, 0, 1,
259 &blitter->sampler[info->filter]);
260
261 /* Initialize the sampler view. */
262 u_sampler_view_default_template(&src_templ, src, src->format);
263 src_templ.format = info->src.format;
264 src_templ.target = array ? PIPE_TEXTURE_2D_ARRAY : PIPE_TEXTURE_2D;
265 src_templ.swizzle_r = PIPE_SWIZZLE_X;
266 src_templ.swizzle_g = PIPE_SWIZZLE_Y;
267 src_templ.swizzle_b = PIPE_SWIZZLE_Z;
268 src_templ.swizzle_a = PIPE_SWIZZLE_W;
269 src_templ.u.tex.first_layer = info->src.box.z;
270 src_templ.u.tex.last_layer = info->src.box.z + depth - 1;
271 src_templ.u.tex.first_level = info->src.level;
272 src_templ.u.tex.last_level = info->src.level;
273 src_view = ctx->create_sampler_view(ctx, src, &src_templ);
274 ctx->set_sampler_views(ctx, PIPE_SHADER_COMPUTE, 0, 1, 0, true, &src_view);
275
276 enum asahi_blit_clamp clamp = ASAHI_BLIT_CLAMP_NONE;
277 bool src_sint = util_format_is_pure_sint(info->src.format);
278 bool dst_sint = util_format_is_pure_sint(info->dst.format);
279 if (util_format_is_pure_integer(info->src.format) &&
280 util_format_is_pure_integer(info->dst.format)) {
281
282 if (src_sint && !dst_sint)
283 clamp = ASAHI_BLIT_CLAMP_SINT_TO_UINT;
284 else if (!src_sint && dst_sint)
285 clamp = ASAHI_BLIT_CLAMP_UINT_TO_SINT;
286 }
287
288 if (!blitter->blit_cs[clamp][array]) {
289 blitter->blit_cs[clamp][array] =
290 asahi_blit_compute_shader(ctx, clamp, array);
291 }
292
293 ctx->bind_compute_state(ctx, blitter->blit_cs[clamp][array]);
294
295 struct pipe_grid_info grid_info = {
296 .block = {BLIT_WG_SIZE, BLIT_WG_SIZE, 1},
297 .last_block = {width % BLIT_WG_SIZE, height % BLIT_WG_SIZE, 1},
298 .grid =
299 {
300 DIV_ROUND_UP(width, BLIT_WG_SIZE),
301 DIV_ROUND_UP(height, BLIT_WG_SIZE),
302 depth,
303 },
304 };
305 ctx->launch_grid(ctx, &grid_info);
306 ctx->set_shader_images(ctx, PIPE_SHADER_COMPUTE, 0, 0, 1, NULL);
307 ctx->set_constant_buffer(ctx, PIPE_SHADER_COMPUTE, 0, false, NULL);
308 ctx->set_sampler_views(ctx, PIPE_SHADER_COMPUTE, 0, 0, 1, false, NULL);
309
310 asahi_compute_restore(agx_context(ctx));
311 }
312
313 void
agx_blitter_save(struct agx_context * ctx,struct blitter_context * blitter,bool render_cond)314 agx_blitter_save(struct agx_context *ctx, struct blitter_context *blitter,
315 bool render_cond)
316 {
317 util_blitter_save_vertex_buffers(blitter, ctx->vertex_buffers,
318 util_last_bit(ctx->vb_mask));
319 util_blitter_save_vertex_elements(blitter, ctx->attributes);
320 util_blitter_save_vertex_shader(blitter,
321 ctx->stage[PIPE_SHADER_VERTEX].shader);
322 util_blitter_save_tessctrl_shader(blitter,
323 ctx->stage[PIPE_SHADER_TESS_CTRL].shader);
324 util_blitter_save_tesseval_shader(blitter,
325 ctx->stage[PIPE_SHADER_TESS_EVAL].shader);
326 util_blitter_save_geometry_shader(blitter,
327 ctx->stage[PIPE_SHADER_GEOMETRY].shader);
328 util_blitter_save_rasterizer(blitter, ctx->rast);
329 util_blitter_save_viewport(blitter, &ctx->viewport[0]);
330 util_blitter_save_scissor(blitter, &ctx->scissor[0]);
331 util_blitter_save_fragment_shader(blitter,
332 ctx->stage[PIPE_SHADER_FRAGMENT].shader);
333 util_blitter_save_blend(blitter, ctx->blend);
334 util_blitter_save_depth_stencil_alpha(blitter, ctx->zs);
335 util_blitter_save_stencil_ref(blitter, &ctx->stencil_ref);
336 util_blitter_save_so_targets(blitter, ctx->streamout.num_targets,
337 ctx->streamout.targets);
338 util_blitter_save_sample_mask(blitter, ctx->sample_mask, 0);
339
340 util_blitter_save_framebuffer(blitter, &ctx->framebuffer);
341 util_blitter_save_fragment_sampler_states(
342 blitter, ctx->stage[PIPE_SHADER_FRAGMENT].sampler_count,
343 (void **)(ctx->stage[PIPE_SHADER_FRAGMENT].samplers));
344 util_blitter_save_fragment_sampler_views(
345 blitter, ctx->stage[PIPE_SHADER_FRAGMENT].texture_count,
346 (struct pipe_sampler_view **)ctx->stage[PIPE_SHADER_FRAGMENT].textures);
347 util_blitter_save_fragment_constant_buffer_slot(
348 blitter, ctx->stage[PIPE_SHADER_FRAGMENT].cb);
349
350 if (!render_cond) {
351 util_blitter_save_render_condition(blitter,
352 (struct pipe_query *)ctx->cond_query,
353 ctx->cond_cond, ctx->cond_mode);
354 }
355 }
356
357 void
agx_blit(struct pipe_context * pipe,const struct pipe_blit_info * info)358 agx_blit(struct pipe_context *pipe, const struct pipe_blit_info *info)
359 {
360 struct agx_context *ctx = agx_context(pipe);
361
362 if (info->render_condition_enable && !agx_render_condition_check(ctx))
363 return;
364
365 if (!util_blitter_is_blit_supported(ctx->blitter, info)) {
366 fprintf(stderr, "\n");
367 util_dump_blit_info(stderr, info);
368 fprintf(stderr, "\n\n");
369 unreachable("Unsupported blit");
370 }
371
372 /* Legalize compression /before/ calling into u_blitter to avoid recursion.
373 * u_blitter bans recursive usage.
374 */
375 agx_legalize_compression(ctx, agx_resource(info->dst.resource),
376 info->dst.format);
377
378 agx_legalize_compression(ctx, agx_resource(info->src.resource),
379 info->src.format);
380
381 if (asahi_compute_blit_supported(info) &&
382 (agx_device(pipe->screen)->debug & AGX_DBG_COMPBLIT) &&
383 !(ail_is_compressed(&agx_resource(info->dst.resource)->layout) &&
384 util_format_get_blocksize(info->dst.format) == 16)) {
385
386 asahi_compute_blit(pipe, info, &ctx->compute_blitter);
387 return;
388 }
389
390 /* Handle self-blits */
391 agx_flush_writer(ctx, agx_resource(info->dst.resource), "Blit");
392
393 agx_blitter_save(ctx, ctx->blitter, info->render_condition_enable);
394 util_blitter_blit(ctx->blitter, info);
395 }
396
397 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)398 try_copy_via_blit(struct pipe_context *pctx, struct pipe_resource *dst,
399 unsigned dst_level, unsigned dstx, unsigned dsty,
400 unsigned dstz, struct pipe_resource *src, unsigned src_level,
401 const struct pipe_box *src_box)
402 {
403 struct agx_context *ctx = agx_context(pctx);
404
405 if (dst->target == PIPE_BUFFER)
406 return false;
407
408 /* TODO: Handle these for rusticl copies */
409 if (dst->target != src->target)
410 return false;
411
412 struct pipe_blit_info info = {
413 .dst =
414 {
415 .resource = dst,
416 .level = dst_level,
417 .box.x = dstx,
418 .box.y = dsty,
419 .box.z = dstz,
420 .box.width = src_box->width,
421 .box.height = src_box->height,
422 .box.depth = src_box->depth,
423 .format = dst->format,
424 },
425 .src =
426 {
427 .resource = src,
428 .level = src_level,
429 .box = *src_box,
430 .format = src->format,
431 },
432 .mask = util_format_get_mask(src->format),
433 .filter = PIPE_TEX_FILTER_NEAREST,
434 .scissor_enable = 0,
435 };
436
437 /* snorm formats don't round trip, so don't use them for copies */
438 if (util_format_is_snorm(info.dst.format))
439 info.dst.format = util_format_snorm_to_sint(info.dst.format);
440
441 if (util_format_is_snorm(info.src.format))
442 info.src.format = util_format_snorm_to_sint(info.src.format);
443
444 if (util_blitter_is_blit_supported(ctx->blitter, &info) &&
445 info.dst.format == info.src.format) {
446
447 agx_blit(pctx, &info);
448 return true;
449 } else {
450 return false;
451 }
452 }
453
454 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)455 agx_resource_copy_region(struct pipe_context *pctx, struct pipe_resource *dst,
456 unsigned dst_level, unsigned dstx, unsigned dsty,
457 unsigned dstz, struct pipe_resource *src,
458 unsigned src_level, const struct pipe_box *src_box)
459 {
460 if (try_copy_via_blit(pctx, dst, dst_level, dstx, dsty, dstz, src, src_level,
461 src_box))
462 return;
463
464 /* CPU fallback */
465 util_resource_copy_region(pctx, dst, dst_level, dstx, dsty, dstz, src,
466 src_level, src_box);
467 }
468