1 /**************************************************************************
2 *
3 * Copyright © 2022 Intel Corporation
4 *
5 * Permission is hereby granted, free of charge, to any person obtaining a
6 * copy of this software and associated documentation files (the "Software"),
7 * to deal in the Software without restriction, including without limitation
8 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
9 * and/or sell copies of the Software, and to permit persons to whom the
10 * Software is furnished to do so, subject to the following conditions:
11 *
12 * The above copyright notice and this permission notice (including the next
13 * paragraph) shall be included in all copies or substantial portions of the
14 * Software.
15 *
16 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
19 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
20 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
21 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
22 * DEALINGS IN THE SOFTWARE.
23 *
24 **************************************************************************/
25
26 #include "compiler/glsl/astc_glsl.h"
27 #include "compiler/glsl/bc1_glsl.h"
28 #include "compiler/glsl/bc4_glsl.h"
29 #include "compiler/glsl/cross_platform_settings_piece_all.h"
30 #include "compiler/glsl/etc2_rgba_stitch_glsl.h"
31
32 #include "main/context.h"
33 #include "main/shaderapi.h"
34 #include "main/shaderobj.h"
35 #include "main/texcompress_astc.h"
36 #include "util/texcompress_astc_luts_wrap.h"
37 #include "main/uniforms.h"
38
39 #include "state_tracker/st_atom_constbuf.h"
40 #include "state_tracker/st_bc1_tables.h"
41 #include "state_tracker/st_context.h"
42 #include "state_tracker/st_program.h"
43 #include "state_tracker/st_texcompress_compute.h"
44 #include "state_tracker/st_texture.h"
45
46 #include "util/u_hash_table.h"
47 #include "util/u_string.h"
48
49 enum compute_program_id {
50 COMPUTE_PROGRAM_BC1,
51 COMPUTE_PROGRAM_BC4,
52 COMPUTE_PROGRAM_STITCH,
53 COMPUTE_PROGRAM_ASTC_4x4,
54 COMPUTE_PROGRAM_ASTC_5x4,
55 COMPUTE_PROGRAM_ASTC_5x5,
56 COMPUTE_PROGRAM_ASTC_6x5,
57 COMPUTE_PROGRAM_ASTC_6x6,
58 COMPUTE_PROGRAM_ASTC_8x5,
59 COMPUTE_PROGRAM_ASTC_8x6,
60 COMPUTE_PROGRAM_ASTC_8x8,
61 COMPUTE_PROGRAM_ASTC_10x5,
62 COMPUTE_PROGRAM_ASTC_10x6,
63 COMPUTE_PROGRAM_ASTC_10x8,
64 COMPUTE_PROGRAM_ASTC_10x10,
65 COMPUTE_PROGRAM_ASTC_12x10,
66 COMPUTE_PROGRAM_ASTC_12x12,
67 COMPUTE_PROGRAM_COUNT
68 };
69
70 static struct gl_program * PRINTFLIKE(3, 4)
get_compute_program(struct st_context * st,enum compute_program_id prog_id,const char * source_fmt,...)71 get_compute_program(struct st_context *st,
72 enum compute_program_id prog_id,
73 const char *source_fmt, ...)
74 {
75 /* Try to get the program from the cache. */
76 assert(prog_id < COMPUTE_PROGRAM_COUNT);
77 if (st->texcompress_compute.progs[prog_id])
78 return st->texcompress_compute.progs[prog_id];
79
80 /* Cache miss. Create the final source string. */
81 char *source_str;
82 va_list ap;
83 va_start(ap, source_fmt);
84 int num_printed_bytes = vasprintf(&source_str, source_fmt, ap);
85 va_end(ap);
86 if (num_printed_bytes == -1)
87 return NULL;
88
89 /* Compile and link the shader. Then, destroy the shader string. */
90 const char *strings[] = { source_str };
91 GLuint program =
92 _mesa_CreateShaderProgramv_impl(st->ctx, GL_COMPUTE_SHADER, 1, strings);
93 free(source_str);
94
95 struct gl_shader_program *shProg =
96 _mesa_lookup_shader_program(st->ctx, program);
97 if (!shProg)
98 return NULL;
99
100 if (shProg->data->LinkStatus == LINKING_FAILURE) {
101 fprintf(stderr, "Linking failed:\n%s\n", shProg->data->InfoLog);
102 _mesa_reference_shader_program(st->ctx, &shProg, NULL);
103 return NULL;
104 }
105
106 /* Cache the program and return it. */
107 return st->texcompress_compute.progs[prog_id] =
108 shProg->_LinkedShaders[MESA_SHADER_COMPUTE]->Program;
109 }
110
111 static struct pipe_resource *
create_bc1_endpoint_ssbo(struct pipe_context * pipe)112 create_bc1_endpoint_ssbo(struct pipe_context *pipe)
113 {
114 struct pipe_resource *buffer =
115 pipe_buffer_create(pipe->screen, PIPE_BIND_SHADER_BUFFER,
116 PIPE_USAGE_IMMUTABLE, sizeof(float) *
117 (sizeof(stb__OMatch5) + sizeof(stb__OMatch6)));
118
119 if (!buffer)
120 return NULL;
121
122 struct pipe_transfer *transfer;
123 float (*buffer_map)[2] = pipe_buffer_map(pipe, buffer,
124 PIPE_MAP_WRITE |
125 PIPE_MAP_DISCARD_WHOLE_RESOURCE,
126 &transfer);
127 if (!buffer_map) {
128 pipe_resource_reference(&buffer, NULL);
129 return NULL;
130 }
131
132 for (int i = 0; i < 256; i++) {
133 for (int j = 0; j < 2; j++) {
134 buffer_map[i][j] = (float) stb__OMatch5[i][j];
135 buffer_map[i + 256][j] = (float) stb__OMatch6[i][j];
136 }
137 }
138
139 pipe_buffer_unmap(pipe, transfer);
140
141 return buffer;
142 }
143
144 static void
bind_compute_state(struct st_context * st,struct gl_program * prog,struct pipe_sampler_view ** sampler_views,const struct pipe_shader_buffer * shader_buffers,const struct pipe_image_view * image_views,bool cs_handle_from_prog,bool constbuf0_from_prog)145 bind_compute_state(struct st_context *st,
146 struct gl_program *prog,
147 struct pipe_sampler_view **sampler_views,
148 const struct pipe_shader_buffer *shader_buffers,
149 const struct pipe_image_view *image_views,
150 bool cs_handle_from_prog,
151 bool constbuf0_from_prog)
152 {
153 assert(prog->info.stage == PIPE_SHADER_COMPUTE);
154
155 /* Set compute states in the same order as defined in st_atom_list.h */
156
157 assert(prog->affected_states & ST_NEW_CS_STATE);
158 assert(st->shader_has_one_variant[PIPE_SHADER_COMPUTE]);
159 cso_set_compute_shader_handle(st->cso_context,
160 cs_handle_from_prog ?
161 prog->variants->driver_shader : NULL);
162
163 if (prog->affected_states & ST_NEW_CS_SAMPLER_VIEWS) {
164 st->pipe->set_sampler_views(st->pipe, prog->info.stage, 0,
165 prog->info.num_textures, 0, false,
166 sampler_views);
167 }
168
169 if (prog->affected_states & ST_NEW_CS_SAMPLERS) {
170 /* Programs seem to set this bit more often than needed. For example, if
171 * a program only uses texelFetch, this shouldn't be needed. Section
172 * "11.1.3.2 Texel Fetches", of the GL 4.6 spec says:
173 *
174 * Texel fetch proceeds similarly to the steps described for texture
175 * access in section 11.1.3.5, with the exception that none of the
176 * operations controlled by sampler object state are performed,
177 *
178 * We assume that the program is using texelFetch or doesn't care about
179 * this state for a similar reason.
180 *
181 * See https://gitlab.freedesktop.org/mesa/mesa/-/issues/8014.
182 */
183 }
184
185 if (prog->affected_states & ST_NEW_CS_CONSTANTS) {
186 st_upload_constants(st, constbuf0_from_prog ? prog : NULL,
187 prog->info.stage);
188 }
189
190 if (prog->affected_states & ST_NEW_CS_UBOS) {
191 unreachable("Uniform buffer objects not handled");
192 }
193
194 if (prog->affected_states & ST_NEW_CS_ATOMICS) {
195 unreachable("Atomic buffer objects not handled");
196 }
197
198 if (prog->affected_states & ST_NEW_CS_SSBOS) {
199 st->pipe->set_shader_buffers(st->pipe, prog->info.stage, 0,
200 prog->info.num_ssbos, shader_buffers,
201 prog->sh.ShaderStorageBlocksWriteAccess);
202 }
203
204 if (prog->affected_states & ST_NEW_CS_IMAGES) {
205 st->pipe->set_shader_images(st->pipe, prog->info.stage, 0,
206 prog->info.num_images, 0, image_views);
207 }
208 }
209
210 static void
dispatch_compute_state(struct st_context * st,struct gl_program * prog,struct pipe_sampler_view ** sampler_views,const struct pipe_shader_buffer * shader_buffers,const struct pipe_image_view * image_views,unsigned num_workgroups_x,unsigned num_workgroups_y,unsigned num_workgroups_z)211 dispatch_compute_state(struct st_context *st,
212 struct gl_program *prog,
213 struct pipe_sampler_view **sampler_views,
214 const struct pipe_shader_buffer *shader_buffers,
215 const struct pipe_image_view *image_views,
216 unsigned num_workgroups_x,
217 unsigned num_workgroups_y,
218 unsigned num_workgroups_z)
219 {
220 assert(prog->info.stage == PIPE_SHADER_COMPUTE);
221
222 /* Bind the state */
223 bind_compute_state(st, prog, sampler_views, shader_buffers, image_views,
224 true, true);
225
226 /* Launch the grid */
227 const struct pipe_grid_info info = {
228 .block[0] = prog->info.workgroup_size[0],
229 .block[1] = prog->info.workgroup_size[1],
230 .block[2] = prog->info.workgroup_size[2],
231 .grid[0] = num_workgroups_x,
232 .grid[1] = num_workgroups_y,
233 .grid[2] = num_workgroups_z,
234 };
235
236 st->pipe->launch_grid(st->pipe, &info);
237
238 /* Unbind the state */
239 bind_compute_state(st, prog, NULL, NULL, NULL, false, false);
240
241 /* If the previously used compute program was relying on any state that was
242 * trampled on by these state changes, dirty the relevant flags.
243 */
244 if (st->cp) {
245 st->ctx->NewDriverState |=
246 st->cp->affected_states & prog->affected_states;
247 }
248 }
249
250 static struct pipe_resource *
cs_encode_bc1(struct st_context * st,struct pipe_resource * rgba8_tex)251 cs_encode_bc1(struct st_context *st,
252 struct pipe_resource *rgba8_tex)
253 {
254 /* Create the required compute state */
255 struct gl_program *prog =
256 get_compute_program(st, COMPUTE_PROGRAM_BC1, bc1_source,
257 cross_platform_settings_piece_all_header);
258 if (!prog)
259 return NULL;
260
261 /* ... complete the program setup by defining the number of refinements to
262 * do on the created blocks. The program will attempt to create a more
263 * accurate encoding on each iteration. Doing at least one refinement
264 * provides a significant improvement in quality and is needed to give a
265 * result comparable to the CPU encoder (according to piglit tests).
266 * Additional refinements don't help as much.
267 */
268 const unsigned num_refinements = 1;
269 _mesa_uniform(0, 1, &num_refinements, st->ctx, prog->shader_program,
270 GLSL_TYPE_UINT, 1);
271
272 const struct pipe_sampler_view templ = {
273 .target = PIPE_TEXTURE_2D,
274 .format = PIPE_FORMAT_R8G8B8A8_UNORM,
275 .swizzle_r = PIPE_SWIZZLE_X,
276 .swizzle_g = PIPE_SWIZZLE_Y,
277 .swizzle_b = PIPE_SWIZZLE_Z,
278 .swizzle_a = PIPE_SWIZZLE_W,
279 };
280 struct pipe_sampler_view *rgba8_view =
281 st->pipe->create_sampler_view(st->pipe, rgba8_tex, &templ);
282 if (!rgba8_view)
283 return NULL;
284
285 const struct pipe_shader_buffer ssbo = {
286 .buffer = st->texcompress_compute.bc1_endpoint_buf,
287 .buffer_size = st->texcompress_compute.bc1_endpoint_buf->width0,
288 };
289
290 struct pipe_resource *bc1_tex =
291 st_texture_create(st, PIPE_TEXTURE_2D, PIPE_FORMAT_R32G32_UINT, 0,
292 DIV_ROUND_UP(rgba8_tex->width0, 4),
293 DIV_ROUND_UP(rgba8_tex->height0, 4), 1, 1, 0,
294 PIPE_BIND_SHADER_IMAGE |
295 PIPE_BIND_SAMPLER_VIEW, false,
296 PIPE_COMPRESSION_FIXED_RATE_NONE);
297 if (!bc1_tex)
298 goto release_sampler_views;
299
300 const struct pipe_image_view image = {
301 .resource = bc1_tex,
302 .format = PIPE_FORMAT_R16G16B16A16_UINT,
303 .access = PIPE_IMAGE_ACCESS_WRITE,
304 .shader_access = PIPE_IMAGE_ACCESS_WRITE,
305 };
306
307 /* Dispatch the compute state */
308 dispatch_compute_state(st, prog, &rgba8_view, &ssbo, &image,
309 DIV_ROUND_UP(rgba8_tex->width0, 32),
310 DIV_ROUND_UP(rgba8_tex->height0, 32), 1);
311
312 release_sampler_views:
313 pipe_sampler_view_reference(&rgba8_view, NULL);
314
315 return bc1_tex;
316 }
317
318 static struct pipe_resource *
cs_encode_bc4(struct st_context * st,struct pipe_resource * rgba8_tex,enum pipe_swizzle component,bool use_snorm)319 cs_encode_bc4(struct st_context *st,
320 struct pipe_resource *rgba8_tex,
321 enum pipe_swizzle component, bool use_snorm)
322 {
323 /* Create the required compute state */
324 struct gl_program *prog =
325 get_compute_program(st, COMPUTE_PROGRAM_BC4, bc4_source,
326 cross_platform_settings_piece_all_header);
327 if (!prog)
328 return NULL;
329
330 /* ... complete the program setup by picking the channel to encode and
331 * whether to encode it as snorm. The shader doesn't actually support
332 * channel index 2. So, pick index 0 and rely on swizzling instead.
333 */
334 const unsigned params[] = { 0, use_snorm };
335 _mesa_uniform(0, 1, params, st->ctx, prog->shader_program,
336 GLSL_TYPE_UINT, 2);
337
338 const struct pipe_sampler_view templ = {
339 .target = PIPE_TEXTURE_2D,
340 .format = PIPE_FORMAT_R8G8B8A8_UNORM,
341 .swizzle_r = component,
342 .swizzle_g = PIPE_SWIZZLE_0,
343 .swizzle_b = PIPE_SWIZZLE_0,
344 .swizzle_a = PIPE_SWIZZLE_1,
345 };
346 struct pipe_sampler_view *rgba8_view =
347 st->pipe->create_sampler_view(st->pipe, rgba8_tex, &templ);
348 if (!rgba8_view)
349 return NULL;
350
351 struct pipe_resource *bc4_tex =
352 st_texture_create(st, PIPE_TEXTURE_2D, PIPE_FORMAT_R32G32_UINT, 0,
353 DIV_ROUND_UP(rgba8_tex->width0, 4),
354 DIV_ROUND_UP(rgba8_tex->height0, 4), 1, 1, 0,
355 PIPE_BIND_SHADER_IMAGE |
356 PIPE_BIND_SAMPLER_VIEW, false,
357 PIPE_COMPRESSION_FIXED_RATE_NONE);
358 if (!bc4_tex)
359 goto release_sampler_views;
360
361 const struct pipe_image_view image = {
362 .resource = bc4_tex,
363 .format = PIPE_FORMAT_R16G16B16A16_UINT,
364 .access = PIPE_IMAGE_ACCESS_WRITE,
365 .shader_access = PIPE_IMAGE_ACCESS_WRITE,
366 };
367
368 /* Dispatch the compute state */
369 dispatch_compute_state(st, prog, &rgba8_view, NULL, &image, 1,
370 DIV_ROUND_UP(rgba8_tex->width0, 16),
371 DIV_ROUND_UP(rgba8_tex->height0, 16));
372
373 release_sampler_views:
374 pipe_sampler_view_reference(&rgba8_view, NULL);
375
376 return bc4_tex;
377 }
378
379 static struct pipe_resource *
cs_stitch_64bpb_textures(struct st_context * st,struct pipe_resource * tex_hi,struct pipe_resource * tex_lo)380 cs_stitch_64bpb_textures(struct st_context *st,
381 struct pipe_resource *tex_hi,
382 struct pipe_resource *tex_lo)
383 {
384 assert(util_format_get_blocksizebits(tex_hi->format) == 64);
385 assert(util_format_get_blocksizebits(tex_lo->format) == 64);
386 assert(tex_hi->width0 == tex_lo->width0);
387 assert(tex_hi->height0 == tex_lo->height0);
388
389 struct pipe_resource *stitched_tex = NULL;
390
391 /* Create the required compute state */
392 struct gl_program *prog =
393 get_compute_program(st, COMPUTE_PROGRAM_STITCH, etc2_rgba_stitch_source,
394 cross_platform_settings_piece_all_header);
395 if (!prog)
396 return NULL;
397
398 const struct pipe_sampler_view templ = {
399 .target = PIPE_TEXTURE_2D,
400 .format = PIPE_FORMAT_R32G32_UINT,
401 .swizzle_r = PIPE_SWIZZLE_X,
402 .swizzle_g = PIPE_SWIZZLE_Y,
403 .swizzle_b = PIPE_SWIZZLE_0,
404 .swizzle_a = PIPE_SWIZZLE_1,
405 };
406 struct pipe_sampler_view *rg32_views[2] = {
407 [0] = st->pipe->create_sampler_view(st->pipe, tex_hi, &templ),
408 [1] = st->pipe->create_sampler_view(st->pipe, tex_lo, &templ),
409 };
410 if (!rg32_views[0] || !rg32_views[1])
411 goto release_sampler_views;
412
413 stitched_tex =
414 st_texture_create(st, PIPE_TEXTURE_2D, PIPE_FORMAT_R32G32B32A32_UINT, 0,
415 tex_hi->width0,
416 tex_hi->height0, 1, 1, 0,
417 PIPE_BIND_SHADER_IMAGE |
418 PIPE_BIND_SAMPLER_VIEW, false,
419 PIPE_COMPRESSION_FIXED_RATE_NONE);
420 if (!stitched_tex)
421 goto release_sampler_views;
422
423 const struct pipe_image_view image = {
424 .resource = stitched_tex,
425 .format = PIPE_FORMAT_R32G32B32A32_UINT,
426 .access = PIPE_IMAGE_ACCESS_WRITE,
427 .shader_access = PIPE_IMAGE_ACCESS_WRITE,
428 };
429
430 /* Dispatch the compute state */
431 dispatch_compute_state(st, prog, rg32_views, NULL, &image,
432 DIV_ROUND_UP(tex_hi->width0, 8),
433 DIV_ROUND_UP(tex_hi->height0, 8), 1);
434
435 release_sampler_views:
436 pipe_sampler_view_reference(&rg32_views[0], NULL);
437 pipe_sampler_view_reference(&rg32_views[1], NULL);
438
439 return stitched_tex;
440 }
441
442 static struct pipe_resource *
cs_encode_bc3(struct st_context * st,struct pipe_resource * rgba8_tex)443 cs_encode_bc3(struct st_context *st,
444 struct pipe_resource *rgba8_tex)
445 {
446 struct pipe_resource *bc3_tex = NULL;
447
448 /* Encode RGB channels as BC1. */
449 struct pipe_resource *bc1_tex = cs_encode_bc1(st, rgba8_tex);
450 if (!bc1_tex)
451 return NULL;
452
453 /* Encode alpha channels as BC4. */
454 struct pipe_resource *bc4_tex =
455 cs_encode_bc4(st, rgba8_tex, PIPE_SWIZZLE_W, false);
456 if (!bc4_tex)
457 goto release_textures;
458
459 st->pipe->memory_barrier(st->pipe, PIPE_BARRIER_TEXTURE);
460
461 /* Combine BC1 and BC4 to create BC3. */
462 bc3_tex = cs_stitch_64bpb_textures(st, bc1_tex, bc4_tex);
463 if (!bc3_tex)
464 goto release_textures;
465
466 release_textures:
467 pipe_resource_reference(&bc1_tex, NULL);
468 pipe_resource_reference(&bc4_tex, NULL);
469
470 return bc3_tex;
471 }
472
473 static struct pipe_resource *
sw_decode_astc(struct st_context * st,uint8_t * astc_data,unsigned astc_stride,mesa_format astc_format,unsigned width_px,unsigned height_px)474 sw_decode_astc(struct st_context *st,
475 uint8_t *astc_data,
476 unsigned astc_stride,
477 mesa_format astc_format,
478 unsigned width_px, unsigned height_px)
479 {
480 /* Create the destination */
481 struct pipe_resource *rgba8_tex =
482 st_texture_create(st, PIPE_TEXTURE_2D, PIPE_FORMAT_R8G8B8A8_UNORM, 0,
483 width_px, height_px, 1, 1, 0,
484 PIPE_BIND_SAMPLER_VIEW, false,
485 PIPE_COMPRESSION_FIXED_RATE_NONE);
486 if (!rgba8_tex)
487 return NULL;
488
489 /* Temporarily map the destination and decode into the returned pointer */
490 struct pipe_transfer *rgba8_xfer;
491 void *rgba8_map = pipe_texture_map(st->pipe, rgba8_tex, 0, 0,
492 PIPE_MAP_WRITE, 0, 0,
493 width_px, height_px, &rgba8_xfer);
494 if (!rgba8_map) {
495 pipe_resource_reference(&rgba8_tex, NULL);
496 return NULL;
497 }
498
499 _mesa_unpack_astc_2d_ldr(rgba8_map, rgba8_xfer->stride,
500 astc_data, astc_stride,
501 width_px, height_px, astc_format);
502
503 pipe_texture_unmap(st->pipe, rgba8_xfer);
504
505 return rgba8_tex;
506 }
507
508 static struct pipe_sampler_view *
create_astc_cs_payload_view(struct st_context * st,uint8_t * data,unsigned stride,uint32_t width_el,uint32_t height_el)509 create_astc_cs_payload_view(struct st_context *st,
510 uint8_t *data, unsigned stride,
511 uint32_t width_el, uint32_t height_el)
512 {
513 const struct pipe_resource src_templ = {
514 .target = PIPE_TEXTURE_2D,
515 .format = PIPE_FORMAT_R32G32B32A32_UINT,
516 .bind = PIPE_BIND_SAMPLER_VIEW,
517 .usage = PIPE_USAGE_STAGING,
518 .width0 = width_el,
519 .height0 = height_el,
520 .depth0 = 1,
521 .array_size = 1,
522 };
523
524 struct pipe_resource *payload_res =
525 st->screen->resource_create(st->screen, &src_templ);
526
527 if (!payload_res)
528 return NULL;
529
530 struct pipe_box box;
531 u_box_origin_2d(width_el, height_el, &box);
532
533 st->pipe->texture_subdata(st->pipe, payload_res, 0, 0,
534 &box,
535 data,
536 stride,
537 0 /* unused */);
538
539 const struct pipe_sampler_view view_templ = {
540 .target = PIPE_TEXTURE_2D,
541 .format = payload_res->format,
542 .swizzle_r = PIPE_SWIZZLE_X,
543 .swizzle_g = PIPE_SWIZZLE_Y,
544 .swizzle_b = PIPE_SWIZZLE_Z,
545 .swizzle_a = PIPE_SWIZZLE_W,
546 };
547
548 struct pipe_sampler_view *view =
549 st->pipe->create_sampler_view(st->pipe, payload_res, &view_templ);
550
551 pipe_resource_reference(&payload_res, NULL);
552
553 return view;
554 }
555
556 static struct pipe_sampler_view *
get_astc_partition_table_view(struct st_context * st,unsigned block_w,unsigned block_h)557 get_astc_partition_table_view(struct st_context *st,
558 unsigned block_w,
559 unsigned block_h)
560 {
561 unsigned lut_width;
562 unsigned lut_height;
563 struct pipe_box ptable_box;
564 void *ptable_data =
565 _mesa_get_astc_decoder_partition_table(block_w, block_h, &lut_width, &lut_height);
566 u_box_origin_2d(lut_width, lut_height, &ptable_box);
567
568 struct pipe_sampler_view *view =
569 util_hash_table_get(st->texcompress_compute.astc_partition_tables,
570 ptable_data);
571
572 if (view)
573 return view;
574
575 struct pipe_resource *res =
576 st_texture_create(st, PIPE_TEXTURE_2D, PIPE_FORMAT_R8_UINT, 0,
577 ptable_box.width, ptable_box.height,
578 1, 1, 0,
579 PIPE_BIND_SAMPLER_VIEW, false,
580 PIPE_COMPRESSION_FIXED_RATE_NONE);
581 if (!res)
582 return NULL;
583
584 st->pipe->texture_subdata(st->pipe, res, 0, 0,
585 &ptable_box,
586 ptable_data,
587 ptable_box.width,
588 0 /* unused */);
589
590 const struct pipe_sampler_view templ = {
591 .target = PIPE_TEXTURE_2D,
592 .format = res->format,
593 .swizzle_r = PIPE_SWIZZLE_X,
594 .swizzle_g = PIPE_SWIZZLE_Y,
595 .swizzle_b = PIPE_SWIZZLE_Z,
596 .swizzle_a = PIPE_SWIZZLE_W,
597 };
598
599 view = st->pipe->create_sampler_view(st->pipe, res, &templ);
600
601 pipe_resource_reference(&res, NULL);
602
603 if (view) {
604 _mesa_hash_table_insert(st->texcompress_compute.astc_partition_tables,
605 ptable_data, view);
606 ASSERTED const unsigned max_entries =
607 COMPUTE_PROGRAM_ASTC_12x12 - COMPUTE_PROGRAM_ASTC_4x4 + 1;
608 assert(_mesa_hash_table_num_entries(
609 st->texcompress_compute.astc_partition_tables) < max_entries);
610 }
611
612 return view;
613 }
614
615 static struct pipe_resource *
cs_decode_astc(struct st_context * st,uint8_t * astc_data,unsigned astc_stride,mesa_format astc_format,unsigned width_px,unsigned height_px)616 cs_decode_astc(struct st_context *st,
617 uint8_t *astc_data,
618 unsigned astc_stride,
619 mesa_format astc_format,
620 unsigned width_px, unsigned height_px)
621 {
622 const enum compute_program_id astc_id = COMPUTE_PROGRAM_ASTC_4x4 +
623 util_format_linear(astc_format) - PIPE_FORMAT_ASTC_4x4;
624
625 unsigned block_w, block_h;
626 _mesa_get_format_block_size(astc_format, &block_w, &block_h);
627
628 struct gl_program *prog =
629 get_compute_program(st, astc_id, astc_source, block_w, block_h);
630
631 if (!prog)
632 return NULL;
633
634 struct pipe_sampler_view *ptable_view =
635 get_astc_partition_table_view(st, block_w, block_h);
636
637 if (!ptable_view)
638 return NULL;
639
640 struct pipe_sampler_view *payload_view =
641 create_astc_cs_payload_view(st, astc_data, astc_stride,
642 DIV_ROUND_UP(width_px, block_w),
643 DIV_ROUND_UP(height_px, block_h));
644
645 if (!payload_view)
646 return NULL;
647
648 /* Create the destination */
649 struct pipe_resource *rgba8_tex =
650 st_texture_create(st, PIPE_TEXTURE_2D, PIPE_FORMAT_R8G8B8A8_UNORM, 0,
651 width_px, height_px, 1, 1, 0,
652 PIPE_BIND_SAMPLER_VIEW, false,
653 PIPE_COMPRESSION_FIXED_RATE_NONE);
654
655 if (!rgba8_tex)
656 goto release_payload_view;
657
658 const struct pipe_image_view image = {
659 .resource = rgba8_tex,
660 .format = PIPE_FORMAT_R8G8B8A8_UINT,
661 .access = PIPE_IMAGE_ACCESS_WRITE,
662 .shader_access = PIPE_IMAGE_ACCESS_WRITE,
663 };
664
665 struct pipe_sampler_view *sampler_views[] = {
666 st->texcompress_compute.astc_luts[0],
667 st->texcompress_compute.astc_luts[1],
668 st->texcompress_compute.astc_luts[2],
669 st->texcompress_compute.astc_luts[3],
670 st->texcompress_compute.astc_luts[4],
671 ptable_view,
672 payload_view,
673 };
674
675 dispatch_compute_state(st, prog, sampler_views, NULL, &image,
676 DIV_ROUND_UP(payload_view->texture->width0, 2),
677 DIV_ROUND_UP(payload_view->texture->height0, 2),
678 1);
679
680 release_payload_view:
681 pipe_sampler_view_reference(&payload_view, NULL);
682
683 return rgba8_tex;
684 }
685
686 static struct pipe_sampler_view *
get_sampler_view_for_lut(struct pipe_context * pipe,const astc_decoder_lut * lut)687 get_sampler_view_for_lut(struct pipe_context *pipe,
688 const astc_decoder_lut *lut)
689 {
690 struct pipe_resource *res =
691 pipe_buffer_create_with_data(pipe,
692 PIPE_BIND_SAMPLER_VIEW,
693 PIPE_USAGE_DEFAULT,
694 lut->size_B,
695 lut->data);
696 if (!res)
697 return NULL;
698
699 const struct pipe_sampler_view templ = {
700 .format = lut->format,
701 .target = PIPE_BUFFER,
702 .swizzle_r = PIPE_SWIZZLE_X,
703 .swizzle_g = PIPE_SWIZZLE_Y,
704 .swizzle_b = PIPE_SWIZZLE_Z,
705 .swizzle_a = PIPE_SWIZZLE_W,
706 .u.buf.offset = 0,
707 .u.buf.size = lut->size_B,
708 };
709
710 struct pipe_sampler_view *view =
711 pipe->create_sampler_view(pipe, res, &templ);
712
713 pipe_resource_reference(&res, NULL);
714
715 return view;
716 }
717
718 /* Initializes required resources for Granite ASTC GPU decode.
719 *
720 * There are 5 texture buffer objects and one additional texture required.
721 * We initialize 5 tbo's here and a single texture later during runtime.
722 */
723 static bool
initialize_astc_decoder(struct st_context * st)724 initialize_astc_decoder(struct st_context *st)
725 {
726 astc_decoder_lut_holder astc_lut_holder;
727 _mesa_init_astc_decoder_luts(&astc_lut_holder);
728
729 const astc_decoder_lut *luts[] = {
730 &astc_lut_holder.color_endpoint,
731 &astc_lut_holder.color_endpoint_unquant,
732 &astc_lut_holder.weights,
733 &astc_lut_holder.weights_unquant,
734 &astc_lut_holder.trits_quints,
735 };
736
737 for (unsigned i = 0; i < ARRAY_SIZE(luts); i++) {
738 st->texcompress_compute.astc_luts[i] =
739 get_sampler_view_for_lut(st->pipe, luts[i]);
740 if (!st->texcompress_compute.astc_luts[i])
741 return false;
742 }
743
744 st->texcompress_compute.astc_partition_tables =
745 _mesa_pointer_hash_table_create(NULL);
746
747 if (!st->texcompress_compute.astc_partition_tables)
748 return false;
749
750 return true;
751 }
752
753 bool
st_init_texcompress_compute(struct st_context * st)754 st_init_texcompress_compute(struct st_context *st)
755 {
756 st->texcompress_compute.progs =
757 calloc(COMPUTE_PROGRAM_COUNT, sizeof(struct gl_program *));
758 if (!st->texcompress_compute.progs)
759 return false;
760
761 st->texcompress_compute.bc1_endpoint_buf =
762 create_bc1_endpoint_ssbo(st->pipe);
763 if (!st->texcompress_compute.bc1_endpoint_buf)
764 return false;
765
766 if (!initialize_astc_decoder(st))
767 return false;
768
769 return true;
770 }
771
772 static void
destroy_astc_decoder(struct st_context * st)773 destroy_astc_decoder(struct st_context *st)
774 {
775 for (unsigned i = 0; i < ARRAY_SIZE(st->texcompress_compute.astc_luts); i++)
776 pipe_sampler_view_reference(&st->texcompress_compute.astc_luts[i], NULL);
777
778 if (st->texcompress_compute.astc_partition_tables) {
779 hash_table_foreach(st->texcompress_compute.astc_partition_tables,
780 entry) {
781 pipe_sampler_view_reference(
782 (struct pipe_sampler_view **)&entry->data, NULL);
783 }
784 }
785
786 _mesa_hash_table_destroy(st->texcompress_compute.astc_partition_tables,
787 NULL);
788 }
789
790 void
st_destroy_texcompress_compute(struct st_context * st)791 st_destroy_texcompress_compute(struct st_context *st)
792 {
793 /* The programs in the array are part of the gl_context (in st->ctx).They
794 * are automatically destroyed when the context is destroyed (via
795 * _mesa_free_context_data -> ... -> free_shader_program_data_cb).
796 */
797 free(st->texcompress_compute.progs);
798
799 /* Destroy the SSBO used by the BC1 shader program. */
800 pipe_resource_reference(&st->texcompress_compute.bc1_endpoint_buf, NULL);
801
802 destroy_astc_decoder(st);
803 }
804
805 /* See st_texcompress_compute.h for more information. */
806 bool
st_compute_transcode_astc_to_dxt5(struct st_context * st,uint8_t * astc_data,unsigned astc_stride,mesa_format astc_format,struct pipe_resource * dxt5_tex,unsigned dxt5_level,unsigned dxt5_layer)807 st_compute_transcode_astc_to_dxt5(struct st_context *st,
808 uint8_t *astc_data,
809 unsigned astc_stride,
810 mesa_format astc_format,
811 struct pipe_resource *dxt5_tex,
812 unsigned dxt5_level,
813 unsigned dxt5_layer)
814 {
815 assert(_mesa_has_compute_shaders(st->ctx));
816 assert(_mesa_is_format_astc_2d(astc_format));
817 assert(dxt5_tex->format == PIPE_FORMAT_DXT5_RGBA ||
818 dxt5_tex->format == PIPE_FORMAT_DXT5_SRGBA);
819 assert(dxt5_level <= dxt5_tex->last_level);
820 assert(dxt5_layer <= util_max_layer(dxt5_tex, dxt5_level));
821
822 bool success = false;
823
824 /* Decode ASTC to RGBA8. */
825 struct pipe_resource *rgba8_tex =
826 cs_decode_astc(st, astc_data, astc_stride, astc_format,
827 u_minify(dxt5_tex->width0, dxt5_level),
828 u_minify(dxt5_tex->height0, dxt5_level));
829 if (!rgba8_tex)
830 return false;
831
832 st->pipe->memory_barrier(st->pipe, PIPE_BARRIER_TEXTURE);
833
834 /* Encode RGBA8 to BC3. */
835 struct pipe_resource *bc3_tex = cs_encode_bc3(st, rgba8_tex);
836 if (!bc3_tex)
837 goto release_textures;
838
839 /* Upload the result. */
840 struct pipe_box src_box;
841 u_box_origin_2d(bc3_tex->width0, bc3_tex->height0, &src_box);
842 st->pipe->resource_copy_region(st->pipe, dxt5_tex, dxt5_level,
843 0, 0, dxt5_layer, bc3_tex, 0, &src_box);
844
845 success = true;
846
847 release_textures:
848 pipe_resource_reference(&rgba8_tex, NULL);
849 pipe_resource_reference(&bc3_tex, NULL);
850
851 return success;
852 }
853