1 /*
2 * Copyright (C) 2023 Amazon.com, Inc. or its affiliates
3 *
4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the "Software"),
6 * to deal in the Software without restriction, including without limitation
7 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
8 * and/or sell copies of the Software, and to permit persons to whom the
9 * Software is furnished to do so, subject to the following conditions:
10 *
11 * The above copyright notice and this permission notice (including the next
12 * paragraph) shall be included in all copies or substantial portions of the
13 * Software.
14 *
15 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
16 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
17 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
18 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
19 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
20 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
21 * SOFTWARE.
22 */
23
24 #include "pan_mod_conv_cso.h"
25 #include "nir/pipe_nir.h"
26 #include "nir_builder.h"
27 #include "pan_context.h"
28 #include "pan_resource.h"
29 #include "pan_screen.h"
30 #include "pan_shader.h"
31
32 #define panfrost_afbc_add_info_ubo(name, b) \
33 nir_variable *info_ubo = nir_variable_create( \
34 b.shader, nir_var_mem_ubo, \
35 glsl_array_type(glsl_uint_type(), \
36 sizeof(struct panfrost_afbc_##name##_info) / 4, 0), \
37 "info_ubo"); \
38 info_ubo->data.driver_location = 0;
39
40 #define panfrost_afbc_get_info_field(name, b, field) \
41 nir_load_ubo( \
42 (b), 1, sizeof(((struct panfrost_afbc_##name##_info *)0)->field) * 8, \
43 nir_imm_int(b, 0), \
44 nir_imm_int(b, offsetof(struct panfrost_afbc_##name##_info, field)), \
45 .align_mul = 4, .range = ~0)
46
47 #define panfrost_mtk_add_info_ubo(name, b) \
48 nir_variable *info_ubo = nir_variable_create( \
49 b.shader, nir_var_mem_ubo, \
50 glsl_array_type(glsl_uint_type(), \
51 sizeof(struct panfrost_mtk_##name##_info) / 4, 0), \
52 "info_ubo"); \
53 info_ubo->data.driver_location = 0;
54
55 #define panfrost_mtk_get_info_field(name, b, field) \
56 nir_load_ubo( \
57 (b), 1, sizeof(((struct panfrost_mtk_##name##_info *)0)->field) * 8, \
58 nir_imm_int(b, 0), \
59 nir_imm_int(b, offsetof(struct panfrost_mtk_##name##_info, field)), \
60 .align_mul = 4, .range = ~0)
61
62 static nir_def *
read_afbc_header(nir_builder * b,nir_def * buf,nir_def * idx)63 read_afbc_header(nir_builder *b, nir_def *buf, nir_def *idx)
64 {
65 nir_def *offset = nir_imul_imm(b, idx, AFBC_HEADER_BYTES_PER_TILE);
66 return nir_load_global(b, nir_iadd(b, buf, nir_u2u64(b, offset)), 16,
67 AFBC_HEADER_BYTES_PER_TILE / 4, 32);
68 }
69
70 static void
write_afbc_header(nir_builder * b,nir_def * buf,nir_def * idx,nir_def * hdr)71 write_afbc_header(nir_builder *b, nir_def *buf, nir_def *idx, nir_def *hdr)
72 {
73 nir_def *offset = nir_imul_imm(b, idx, AFBC_HEADER_BYTES_PER_TILE);
74 nir_store_global(b, nir_iadd(b, buf, nir_u2u64(b, offset)), 16, hdr, 0xF);
75 }
76
77 static nir_def *
get_morton_index(nir_builder * b,nir_def * idx,nir_def * src_stride,nir_def * dst_stride)78 get_morton_index(nir_builder *b, nir_def *idx, nir_def *src_stride,
79 nir_def *dst_stride)
80 {
81 nir_def *x = nir_umod(b, idx, dst_stride);
82 nir_def *y = nir_udiv(b, idx, dst_stride);
83
84 nir_def *offset = nir_imul(b, nir_iand_imm(b, y, ~0x7), src_stride);
85 offset = nir_iadd(b, offset, nir_ishl_imm(b, nir_ushr_imm(b, x, 3), 6));
86
87 x = nir_iand_imm(b, x, 0x7);
88 x = nir_iand_imm(b, nir_ior(b, x, nir_ishl_imm(b, x, 2)), 0x13);
89 x = nir_iand_imm(b, nir_ior(b, x, nir_ishl_imm(b, x, 1)), 0x15);
90 y = nir_iand_imm(b, y, 0x7);
91 y = nir_iand_imm(b, nir_ior(b, y, nir_ishl_imm(b, y, 2)), 0x13);
92 y = nir_iand_imm(b, nir_ior(b, y, nir_ishl_imm(b, y, 1)), 0x15);
93 nir_def *tile_idx = nir_ior(b, x, nir_ishl_imm(b, y, 1));
94
95 return nir_iadd(b, offset, tile_idx);
96 }
97
98 static nir_def *
get_superblock_size(nir_builder * b,unsigned arch,nir_def * hdr,nir_def * uncompressed_size)99 get_superblock_size(nir_builder *b, unsigned arch, nir_def *hdr,
100 nir_def *uncompressed_size)
101 {
102 nir_def *size = nir_imm_int(b, 0);
103
104 unsigned body_base_ptr_len = 32;
105 unsigned nr_subblocks = 16;
106 unsigned sz_len = 6; /* bits */
107 nir_def *words[4];
108 nir_def *mask = nir_imm_int(b, (1 << sz_len) - 1);
109 nir_def *is_solid_color = nir_imm_bool(b, false);
110
111 for (int i = 0; i < 4; i++)
112 words[i] = nir_channel(b, hdr, i);
113
114 /* Sum up all of the subblock sizes */
115 for (int i = 0; i < nr_subblocks; i++) {
116 nir_def *subblock_size;
117 unsigned bitoffset = body_base_ptr_len + (i * sz_len);
118 unsigned start = bitoffset / 32;
119 unsigned end = (bitoffset + (sz_len - 1)) / 32;
120 unsigned offset = bitoffset % 32;
121
122 /* Handle differently if the size field is split between two words
123 * of the header */
124 if (start != end) {
125 subblock_size = nir_ior(b, nir_ushr_imm(b, words[start], offset),
126 nir_ishl_imm(b, words[end], 32 - offset));
127 subblock_size = nir_iand(b, subblock_size, mask);
128 } else {
129 subblock_size =
130 nir_ubitfield_extract_imm(b, words[start], offset, sz_len);
131 }
132 subblock_size = nir_bcsel(b, nir_ieq_imm(b, subblock_size, 1),
133 uncompressed_size, subblock_size);
134 size = nir_iadd(b, size, subblock_size);
135
136 /* When the first subblock size is set to zero, the whole superblock is
137 * filled with a solid color specified in the header */
138 if (arch >= 7 && i == 0)
139 is_solid_color = nir_ieq_imm(b, size, 0);
140 }
141
142 return (arch >= 7)
143 ? nir_bcsel(b, is_solid_color, nir_imm_zero(b, 1, 32), size)
144 : size;
145 }
146
147 static nir_def *
get_packed_offset(nir_builder * b,nir_def * metadata,nir_def * idx,nir_def ** out_size)148 get_packed_offset(nir_builder *b, nir_def *metadata, nir_def *idx,
149 nir_def **out_size)
150 {
151 nir_def *metadata_offset =
152 nir_u2u64(b, nir_imul_imm(b, idx, sizeof(struct pan_afbc_block_info)));
153 nir_def *range_ptr = nir_iadd(b, metadata, metadata_offset);
154 nir_def *entry = nir_load_global(b, range_ptr, 4,
155 sizeof(struct pan_afbc_block_info) / 4, 32);
156 nir_def *offset =
157 nir_channel(b, entry, offsetof(struct pan_afbc_block_info, offset) / 4);
158
159 if (out_size)
160 *out_size =
161 nir_channel(b, entry, offsetof(struct pan_afbc_block_info, size) / 4);
162
163 return nir_u2u64(b, offset);
164 }
165
166 #define MAX_LINE_SIZE 16
167
168 static void
copy_superblock(nir_builder * b,nir_def * dst,nir_def * dst_idx,nir_def * hdr_sz,nir_def * src,nir_def * src_idx,nir_def * metadata,nir_def * meta_idx,unsigned align)169 copy_superblock(nir_builder *b, nir_def *dst, nir_def *dst_idx, nir_def *hdr_sz,
170 nir_def *src, nir_def *src_idx, nir_def *metadata,
171 nir_def *meta_idx, unsigned align)
172 {
173 nir_def *hdr = read_afbc_header(b, src, src_idx);
174 nir_def *src_body_base_ptr = nir_u2u64(b, nir_channel(b, hdr, 0));
175 nir_def *src_bodyptr = nir_iadd(b, src, src_body_base_ptr);
176
177 nir_def *size;
178 nir_def *dst_offset = get_packed_offset(b, metadata, meta_idx, &size);
179 nir_def *dst_body_base_ptr = nir_iadd(b, dst_offset, hdr_sz);
180 nir_def *dst_bodyptr = nir_iadd(b, dst, dst_body_base_ptr);
181
182 /* Replace the `base_body_ptr` field if not zero (solid color) */
183 nir_def *hdr2 =
184 nir_vector_insert_imm(b, hdr, nir_u2u32(b, dst_body_base_ptr), 0);
185 hdr = nir_bcsel(b, nir_ieq_imm(b, src_body_base_ptr, 0), hdr, hdr2);
186 write_afbc_header(b, dst, dst_idx, hdr);
187
188 nir_variable *offset_var =
189 nir_local_variable_create(b->impl, glsl_uint_type(), "offset");
190 nir_store_var(b, offset_var, nir_imm_int(b, 0), 1);
191 nir_loop *loop = nir_push_loop(b);
192 {
193 nir_def *offset = nir_load_var(b, offset_var);
194 nir_if *loop_check = nir_push_if(b, nir_uge(b, offset, size));
195 nir_jump(b, nir_jump_break);
196 nir_push_else(b, loop_check);
197 unsigned line_sz = align <= MAX_LINE_SIZE ? align : MAX_LINE_SIZE;
198 for (unsigned i = 0; i < align / line_sz; ++i) {
199 nir_def *src_line = nir_iadd(b, src_bodyptr, nir_u2u64(b, offset));
200 nir_def *dst_line = nir_iadd(b, dst_bodyptr, nir_u2u64(b, offset));
201 nir_store_global(
202 b, dst_line, line_sz,
203 nir_load_global(b, src_line, line_sz, line_sz / 4, 32), ~0);
204 offset = nir_iadd_imm(b, offset, line_sz);
205 }
206 nir_store_var(b, offset_var, offset, 0x1);
207 nir_pop_if(b, loop_check);
208 }
209 nir_pop_loop(b, loop);
210 }
211
212 #define panfrost_afbc_size_get_info_field(b, field) \
213 panfrost_afbc_get_info_field(size, b, field)
214
215 static nir_shader *
panfrost_create_afbc_size_shader(struct panfrost_screen * screen,unsigned bpp,unsigned align)216 panfrost_create_afbc_size_shader(struct panfrost_screen *screen, unsigned bpp,
217 unsigned align)
218 {
219 struct panfrost_device *dev = pan_device(&screen->base);
220
221 nir_builder b = nir_builder_init_simple_shader(
222 MESA_SHADER_COMPUTE, screen->vtbl.get_compiler_options(),
223 "panfrost_afbc_size(bpp=%d)", bpp);
224
225 panfrost_afbc_add_info_ubo(size, b);
226
227 nir_def *coord = nir_load_global_invocation_id(&b, 32);
228 nir_def *block_idx = nir_channel(&b, coord, 0);
229 nir_def *src = panfrost_afbc_size_get_info_field(&b, src);
230 nir_def *metadata = panfrost_afbc_size_get_info_field(&b, metadata);
231 nir_def *uncompressed_size = nir_imm_int(&b, 4 * 4 * bpp / 8); /* bytes */
232
233 nir_def *hdr = read_afbc_header(&b, src, block_idx);
234 nir_def *size = get_superblock_size(&b, dev->arch, hdr, uncompressed_size);
235 size = nir_iand(&b, nir_iadd(&b, size, nir_imm_int(&b, align - 1)),
236 nir_inot(&b, nir_imm_int(&b, align - 1)));
237
238 nir_def *offset = nir_u2u64(
239 &b,
240 nir_iadd(&b,
241 nir_imul_imm(&b, block_idx, sizeof(struct pan_afbc_block_info)),
242 nir_imm_int(&b, offsetof(struct pan_afbc_block_info, size))));
243 nir_store_global(&b, nir_iadd(&b, metadata, offset), 4, size, 0x1);
244
245 return b.shader;
246 }
247
248 #define panfrost_afbc_pack_get_info_field(b, field) \
249 panfrost_afbc_get_info_field(pack, b, field)
250
251 static nir_shader *
panfrost_create_afbc_pack_shader(struct panfrost_screen * screen,unsigned align,bool tiled)252 panfrost_create_afbc_pack_shader(struct panfrost_screen *screen, unsigned align,
253 bool tiled)
254 {
255 nir_builder b = nir_builder_init_simple_shader(
256 MESA_SHADER_COMPUTE, screen->vtbl.get_compiler_options(),
257 "panfrost_afbc_pack");
258
259 panfrost_afbc_add_info_ubo(pack, b);
260
261 nir_def *coord = nir_load_global_invocation_id(&b, 32);
262 nir_def *src_stride = panfrost_afbc_pack_get_info_field(&b, src_stride);
263 nir_def *dst_stride = panfrost_afbc_pack_get_info_field(&b, dst_stride);
264 nir_def *dst_idx = nir_channel(&b, coord, 0);
265 nir_def *src_idx =
266 tiled ? get_morton_index(&b, dst_idx, src_stride, dst_stride) : dst_idx;
267 nir_def *src = panfrost_afbc_pack_get_info_field(&b, src);
268 nir_def *dst = panfrost_afbc_pack_get_info_field(&b, dst);
269 nir_def *header_size =
270 nir_u2u64(&b, panfrost_afbc_pack_get_info_field(&b, header_size));
271 nir_def *metadata = panfrost_afbc_pack_get_info_field(&b, metadata);
272
273 copy_superblock(&b, dst, dst_idx, header_size, src, src_idx, metadata,
274 src_idx, align);
275
276 return b.shader;
277 }
278
279 #define panfrost_mtk_detile_get_info_field(b, field) \
280 panfrost_mtk_get_info_field(detile, b, field)
281
282 static nir_def *
pan_mtk_tiled_from_linear(nir_builder * b,nir_def * linear,nir_def * tiles_per_stride,nir_def * width)283 pan_mtk_tiled_from_linear(nir_builder *b, nir_def *linear, nir_def *tiles_per_stride, nir_def *width)
284 {
285 nir_def *tiled;
286 /* uvec2 tlc = uvec2(linear) >> uvec2(2u, 5u) */
287 nir_def *tlc = nir_ushr(b, linear,
288 nir_imm_ivec2(b, 2, 5));
289
290 /* uvec2 txc = uvec2(linear) & uvec2(3u, 31u) */
291 nir_def *txc = nir_iand(b, linear,
292 nir_imm_ivec2(b, 3, 31));
293
294 /* uint tlo = tlc.y * tiles_per_stride + tlc.x */
295 nir_def *tlo = nir_iadd(b,
296 nir_imul(b,
297 nir_channel(b, tlc, 1),
298 tiles_per_stride),
299 nir_channel(b, tlc, 0));
300 nir_def *txcx = nir_channel(b, txc, 0);
301 nir_def *txcy = nir_channel(b, txc, 1);
302 nir_def *txcytmp = nir_vec2(b, txcy,
303 nir_ushr_imm(b, txcy, 1));
304
305 /* txo = (uvec2(txc.y, txc.y >> 1) << uvec2(2u)) | txc.xx */
306 nir_def *txo = nir_ior(b,
307 nir_ishl_imm(b, txcytmp, 2),
308 nir_vec2(b, txcx, txcx));
309
310 /* uvec2 off = (uvec2(tlo) << uvec2(7u, 6u)) | txo */
311 nir_def *off = nir_ior(b,
312 nir_ishl(b,
313 nir_vec2(b, tlo, tlo),
314 nir_imm_ivec2(b, 7, 6)),
315 txo);
316
317 /* convert to 2D coord
318 * tiled.xy = off % (width / 4, width / 4)
319 * tiled.zw = off / (width / 4, width / 4) */
320 nir_def *width4 = nir_ishl_imm(b, tiles_per_stride, 2);
321 width4 = nir_vec2(b, width4, width4);
322 nir_def *tiled_xy = nir_umod(b, off, width4);
323 nir_def *tiled_zw = nir_udiv(b, off, width4);
324 tiled = nir_vec4(b,
325 nir_channel(b, tiled_xy, 0),
326 nir_channel(b, tiled_xy, 1),
327 nir_channel(b, tiled_zw, 0),
328 nir_channel(b, tiled_zw, 1));
329
330 return tiled;
331 }
332
333 static nir_shader *
panfrost_create_mtk_detile_shader(struct panfrost_screen * screen,unsigned align,bool is_tiled)334 panfrost_create_mtk_detile_shader(struct panfrost_screen *screen, unsigned align,
335 bool is_tiled)
336 {
337 const struct panfrost_device *device = &screen->dev;
338 bool tint_yuv = (device->debug & PAN_DBG_YUV) != 0;
339 nir_builder b = nir_builder_init_simple_shader(
340 MESA_SHADER_COMPUTE, screen->vtbl.get_compiler_options(),
341 "panfrost_mtk_detile");
342 b.shader->info.workgroup_size[0] = 4;
343 b.shader->info.workgroup_size[1] = 16;
344 b.shader->info.workgroup_size[2] = 1;
345
346 const struct glsl_type *image_type =
347 glsl_image_type(GLSL_SAMPLER_DIM_2D, /*is_array*/ false, GLSL_TYPE_UINT);
348
349 panfrost_mtk_add_info_ubo(detile, b);
350
351 nir_variable *y_tiled =
352 nir_variable_create(b.shader, nir_var_image, image_type, "y_tiled");
353 y_tiled->data.binding = 0;
354 y_tiled->data.image.format = PIPE_FORMAT_R8G8B8A8_UINT;
355 BITSET_SET(b.shader->info.images_used, 0);
356 nir_variable *uv_tiled =
357 nir_variable_create(b.shader, nir_var_image, image_type, "uv_tiled");
358 uv_tiled->data.binding = 1;
359 uv_tiled->data.image.format = PIPE_FORMAT_R8G8B8A8_UINT;
360 BITSET_SET(b.shader->info.images_used, 1);
361
362 nir_variable *y_linear =
363 nir_variable_create(b.shader, nir_var_image, image_type, "y_linear");
364 y_linear->data.binding = 2;
365 y_linear->data.image.format = PIPE_FORMAT_R8G8B8A8_UINT;
366 BITSET_SET(b.shader->info.images_used, 2);
367
368 nir_variable *uv_linear =
369 nir_variable_create(b.shader, nir_var_image, image_type, "uv_linear");
370 uv_linear->data.binding = 3;
371 uv_linear->data.image.format = PIPE_FORMAT_R8G8B8A8_UINT;
372 BITSET_SET(b.shader->info.images_used, 3);
373
374 nir_def *linear = nir_load_global_invocation_id(&b, 32);
375 nir_def *tiles_per_stride =
376 panfrost_mtk_detile_get_info_field(&b, tiles_per_stride);
377 nir_def *src_width = panfrost_mtk_detile_get_info_field(&b, src_width);
378
379 nir_def *zero = nir_imm_int(&b, 0);
380
381 nir_def *coord = nir_vec2(&b,
382 nir_channel(&b, linear, 0),
383 nir_channel(&b, linear, 1));
384
385 nir_def *tiled = pan_mtk_tiled_from_linear(&b, coord, tiles_per_stride, src_width);
386
387 nir_def *tiled_xz = nir_vec4(&b, nir_channel(&b, tiled, 0),
388 nir_channel(&b, tiled, 2), zero, zero);
389 nir_def *tiled_yw = nir_vec4(&b, nir_channel(&b, tiled, 1),
390 nir_channel(&b, tiled, 3), zero, zero);
391
392 nir_def *yval = nir_image_load(&b, 4, 32, zero, tiled_xz,
393 zero /* sample */, zero /* lod */,
394 .access = ACCESS_NON_WRITEABLE,
395 .image_dim = GLSL_SAMPLER_DIM_2D,
396 .image_array = false,
397 .dest_type = nir_type_uint32);
398 nir_def *uvval;
399
400 nir_def *dst_y_coord = nir_vec4(&b,
401 nir_channel(&b, coord, 0),
402 nir_channel(&b, coord, 1),
403 zero, zero);
404 /* store Y data */
405 nir_def *img_deref_st_y = nir_imm_int(&b, 2);
406 nir_image_store(&b, img_deref_st_y, dst_y_coord, zero /* sample */,
407 yval, zero /* lod */,
408 .access = ACCESS_NON_READABLE,
409 .image_dim = GLSL_SAMPLER_DIM_2D,
410 .image_array = false, .src_type = nir_type_uint32);
411
412 /* store UV data */
413 nir_def *odd_even_line = nir_iand_imm(&b,
414 nir_channel(&b, dst_y_coord, 1),
415 1);
416 nir_push_if(&b, nir_ieq_imm(&b, odd_even_line, 0));
417 {
418 if (tint_yuv) {
419 /* use just blue for chroma */
420 uvval = nir_imm_ivec4(&b, 0xc0, 0x80, 0xc0, 0x80);
421 } else {
422 nir_def *img_deref_uv = nir_imm_int(&b, 1);
423 uvval = nir_image_load(&b, 4, 32, img_deref_uv, tiled_yw,
424 zero /* sample */, zero /* lod */,
425 .access = ACCESS_NON_WRITEABLE,
426 .image_dim = GLSL_SAMPLER_DIM_2D,
427 .image_array = false,
428 .dest_type = nir_type_uint32);
429 }
430 nir_def *dst_uv_coord = nir_ishr(&b, dst_y_coord,
431 nir_imm_ivec4(&b, 0, 1, 0, 0));
432 nir_def *img_deref_st_uv = nir_imm_int(&b, 3);
433 nir_image_store(&b, img_deref_st_uv, dst_uv_coord, zero /* sample */,
434 uvval, zero /* lod */,
435 .access = ACCESS_NON_READABLE,
436 .image_dim = GLSL_SAMPLER_DIM_2D,
437 .image_array = false, .src_type = nir_type_uint32);
438 }
439 nir_pop_if(&b, NULL);
440
441 return b.shader;
442 }
443
444 struct pan_mod_convert_shader_data *
panfrost_get_mod_convert_shaders(struct panfrost_context * ctx,struct panfrost_resource * rsrc,unsigned align)445 panfrost_get_mod_convert_shaders(struct panfrost_context *ctx,
446 struct panfrost_resource *rsrc, unsigned align)
447 {
448 struct pipe_context *pctx = &ctx->base;
449 struct panfrost_screen *screen = pan_screen(ctx->base.screen);
450 bool tiled = rsrc->image.layout.modifier & AFBC_FORMAT_MOD_TILED;
451 struct pan_mod_convert_shader_key key = {
452 .bpp = util_format_get_blocksizebits(rsrc->base.format),
453 .align = align,
454 .tiled = tiled,
455 };
456
457 pthread_mutex_lock(&ctx->mod_convert_shaders.lock);
458 struct hash_entry *he =
459 _mesa_hash_table_search(ctx->mod_convert_shaders.shaders, &key);
460 struct pan_mod_convert_shader_data *shader = he ? he->data : NULL;
461 pthread_mutex_unlock(&ctx->mod_convert_shaders.lock);
462
463 if (shader)
464 return shader;
465
466 shader = rzalloc(ctx->mod_convert_shaders.shaders, struct pan_mod_convert_shader_data);
467 shader->key = key;
468 _mesa_hash_table_insert(ctx->mod_convert_shaders.shaders, &shader->key, shader);
469
470 #define COMPILE_SHADER(name, ...) \
471 { \
472 nir_shader *nir = \
473 panfrost_create_##name##_shader(screen, __VA_ARGS__); \
474 nir->info.num_ubos = 1; \
475 shader->name##_cso = pipe_shader_from_nir(pctx, nir); \
476 }
477
478 COMPILE_SHADER(afbc_size, key.bpp, key.align);
479 COMPILE_SHADER(afbc_pack, key.align, key.tiled);
480 COMPILE_SHADER(mtk_detile, key.bpp, key.align);
481
482 #undef COMPILE_SHADER
483
484 pthread_mutex_lock(&ctx->mod_convert_shaders.lock);
485 _mesa_hash_table_insert(ctx->mod_convert_shaders.shaders, &shader->key, shader);
486 pthread_mutex_unlock(&ctx->mod_convert_shaders.lock);
487
488 return shader;
489 }
490
491 DERIVE_HASH_TABLE(pan_mod_convert_shader_key);
492
493 void
panfrost_afbc_context_init(struct panfrost_context * ctx)494 panfrost_afbc_context_init(struct panfrost_context *ctx)
495 {
496 ctx->mod_convert_shaders.shaders = pan_mod_convert_shader_key_table_create(NULL);
497 pthread_mutex_init(&ctx->mod_convert_shaders.lock, NULL);
498 }
499
500 void
panfrost_afbc_context_destroy(struct panfrost_context * ctx)501 panfrost_afbc_context_destroy(struct panfrost_context *ctx)
502 {
503 _mesa_hash_table_destroy(ctx->mod_convert_shaders.shaders, NULL);
504 pthread_mutex_destroy(&ctx->mod_convert_shaders.lock);
505 }
506