← Back to branch summary

~mmach/netext73/mesa-haswell

« back to all changes in this revision

Viewing changes to src/gallium/drivers/radeonsi/si_shaderlib_nir.c

  • Committer: mmach
  • Date: 2022-09-22 19:56:13 UTC
  • Revision ID: netbit73@gmail.com-20220922195613-wtik9mmy20tmor0i
2022-09-22 21:17:09

Show diffs side-by-side

added added

removed removed

Lines of Context:
src/gallium/drivers/radeonsi/si_shaderlib_nir.c
1
 
/*
2
 
 * Copyright 2018 Advanced Micro Devices, Inc.
3
 
 * All Rights Reserved.
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
 
 * on the rights to use, copy, modify, merge, publish, distribute, sub
9
 
 * license, and/or sell copies of the Software, and to permit persons to whom
10
 
 * the 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 NON-INFRINGEMENT. IN NO EVENT SHALL
19
 
 * THE AUTHOR(S) AND/OR THEIR SUPPLIERS BE LIABLE FOR ANY CLAIM,
20
 
 * DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
21
 
 * OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
22
 
 * USE OR OTHER DEALINGS IN THE SOFTWARE.
23
 
 */
24
 
 
25
 
#define AC_SURFACE_INCLUDE_NIR
26
 
#include "ac_surface.h"
27
 
#include "si_pipe.h"
28
 
 
29
 
static void *create_nir_cs(struct si_context *sctx, nir_builder *b)
30
 
{
31
 
   nir_shader_gather_info(b->shader, nir_shader_get_entrypoint(b->shader));
32
 
 
33
 
   struct pipe_compute_state state = {0};
34
 
   state.ir_type = PIPE_SHADER_IR_NIR;
35
 
   state.prog = b->shader;
36
 
   sctx->b.screen->finalize_nir(sctx->b.screen, (void*)state.prog);
37
 
   return sctx->b.create_compute_state(&sctx->b, &state);
38
 
}
39
 
 
40
 
static nir_ssa_def *get_global_ids(nir_builder *b, unsigned num_components)
41
 
{
42
 
   unsigned mask = BITFIELD_MASK(num_components);
43
 
 
44
 
   nir_ssa_def *local_ids = nir_channels(b, nir_load_local_invocation_id(b), mask);
45
 
   nir_ssa_def *block_ids = nir_channels(b, nir_load_workgroup_id(b, 32), mask);
46
 
   nir_ssa_def *block_size = nir_channels(b, nir_load_workgroup_size(b), mask);
47
 
   return nir_iadd(b, nir_imul(b, block_ids, block_size), local_ids);
48
 
}
49
 
 
50
 
static void unpack_2x16(nir_builder *b, nir_ssa_def *src, nir_ssa_def **x, nir_ssa_def **y)
51
 
{
52
 
   *x = nir_iand(b, src, nir_imm_int(b, 0xffff));
53
 
   *y = nir_ushr(b, src, nir_imm_int(b, 16));
54
 
}
55
 
 
56
 
static nir_ssa_def *
57
 
deref_ssa(nir_builder *b, nir_variable *var)
58
 
{
59
 
   return &nir_build_deref_var(b, var)->dest.ssa;
60
 
}
61
 
 
62
 
/* Create a NIR compute shader implementing copy_image.
63
 
 *
64
 
 * This shader can handle 1D and 2D, linear and non-linear images.
65
 
 * It expects the source and destination (x,y,z) coords as user_data_amd,
66
 
 * packed into 3 SGPRs as 2x16bits per component.
67
 
 */
68
 
void *si_create_copy_image_cs(struct si_context *sctx, bool is_1D)
69
 
{
70
 
   const nir_shader_compiler_options *options =
71
 
      sctx->b.screen->get_compiler_options(sctx->b.screen, PIPE_SHADER_IR_NIR, PIPE_SHADER_COMPUTE);
72
 
 
73
 
   nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, options, "copy_image_cs");
74
 
   b.shader->info.num_images = 2;
75
 
 
76
 
   /* The workgroup size is either 8x8 for normal (non-linear) 2D images,
77
 
    * or 64x1 for 1D and linear-2D images.
78
 
    */
79
 
   b.shader->info.workgroup_size_variable = true;
80
 
 
81
 
   /* 1D uses 'x' as image coord, and 'y' as array index.
82
 
    * 2D uses 'x'&'y' as image coords, and 'z' as array index.
83
 
    */
84
 
   int n_components = is_1D ? 2 : 3;
85
 
   b.shader->info.cs.user_data_components_amd = n_components;
86
 
   nir_ssa_def *ids = get_global_ids(&b, n_components);
87
 
 
88
 
   nir_ssa_def *coord_src = NULL, *coord_dst = NULL;
89
 
   unpack_2x16(&b, nir_load_user_data_amd(&b), &coord_src, &coord_dst);
90
 
 
91
 
   coord_src = nir_iadd(&b, coord_src, ids);
92
 
   coord_dst = nir_iadd(&b, coord_dst, ids);
93
 
 
94
 
   const struct glsl_type *img_type = glsl_image_type(is_1D ? GLSL_SAMPLER_DIM_1D : GLSL_SAMPLER_DIM_2D,
95
 
                                                      /*is_array*/ true, GLSL_TYPE_FLOAT);
96
 
 
97
 
   nir_variable *img_src = nir_variable_create(b.shader, nir_var_image, img_type, "img_src");
98
 
   img_src->data.binding = 0;
99
 
 
100
 
   nir_variable *img_dst = nir_variable_create(b.shader, nir_var_image, img_type, "img_dst");
101
 
   img_dst->data.binding = 1;
102
 
 
103
 
   nir_ssa_def *undef32 = nir_ssa_undef(&b, 1, 32);
104
 
   nir_ssa_def *zero = nir_imm_int(&b, 0);
105
 
 
106
 
   nir_ssa_def *data = nir_image_deref_load(&b, /*num_components*/ 4, /*bit_size*/ 32,
107
 
      deref_ssa(&b, img_src), coord_src, undef32, zero);
108
 
 
109
 
   nir_image_deref_store(&b, deref_ssa(&b, img_dst), coord_dst, undef32, data, zero);
110
 
 
111
 
   return create_nir_cs(sctx, &b);
112
 
}
113
 
 
114
 
void *si_create_dcc_retile_cs(struct si_context *sctx, struct radeon_surf *surf)
115
 
{
116
 
   const nir_shader_compiler_options *options =
117
 
      sctx->b.screen->get_compiler_options(sctx->b.screen, PIPE_SHADER_IR_NIR, PIPE_SHADER_COMPUTE);
118
 
 
119
 
   nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, options, "dcc_retile");
120
 
   b.shader->info.workgroup_size[0] = 8;
121
 
   b.shader->info.workgroup_size[1] = 8;
122
 
   b.shader->info.workgroup_size[2] = 1;
123
 
   b.shader->info.cs.user_data_components_amd = 3;
124
 
   b.shader->info.num_ssbos = 1;
125
 
 
126
 
   /* Get user data SGPRs. */
127
 
   nir_ssa_def *user_sgprs = nir_load_user_data_amd(&b);
128
 
 
129
 
   /* Relative offset from the displayable DCC to the non-displayable DCC in the same buffer. */
130
 
   nir_ssa_def *src_dcc_offset = nir_channel(&b, user_sgprs, 0);
131
 
 
132
 
   nir_ssa_def *src_dcc_pitch, *dst_dcc_pitch, *src_dcc_height, *dst_dcc_height;
133
 
   unpack_2x16(&b, nir_channel(&b, user_sgprs, 1), &src_dcc_pitch, &src_dcc_height);
134
 
   unpack_2x16(&b, nir_channel(&b, user_sgprs, 2), &dst_dcc_pitch, &dst_dcc_height);
135
 
 
136
 
   /* Get the 2D coordinates. */
137
 
   nir_ssa_def *coord = get_global_ids(&b, 2);
138
 
   nir_ssa_def *zero = nir_imm_int(&b, 0);
139
 
 
140
 
   /* Multiply the coordinates by the DCC block size (they are DCC block coordinates). */
141
 
   coord = nir_imul(&b, coord, nir_imm_ivec2(&b, surf->u.gfx9.color.dcc_block_width,
142
 
                                             surf->u.gfx9.color.dcc_block_height));
143
 
 
144
 
   nir_ssa_def *src_offset =
145
 
      ac_nir_dcc_addr_from_coord(&b, &sctx->screen->info, surf->bpe, &surf->u.gfx9.color.dcc_equation,
146
 
                                 src_dcc_pitch, src_dcc_height, zero, /* DCC slice size */
147
 
                                 nir_channel(&b, coord, 0), nir_channel(&b, coord, 1), /* x, y */
148
 
                                 zero, zero, zero); /* z, sample, pipe_xor */
149
 
   src_offset = nir_iadd(&b, src_offset, src_dcc_offset);
150
 
   nir_ssa_def *value = nir_load_ssbo(&b, 1, 8, zero, src_offset, .align_mul=1);
151
 
 
152
 
   nir_ssa_def *dst_offset =
153
 
      ac_nir_dcc_addr_from_coord(&b, &sctx->screen->info, surf->bpe, &surf->u.gfx9.color.display_dcc_equation,
154
 
                                 dst_dcc_pitch, dst_dcc_height, zero, /* DCC slice size */
155
 
                                 nir_channel(&b, coord, 0), nir_channel(&b, coord, 1), /* x, y */
156
 
                                 zero, zero, zero); /* z, sample, pipe_xor */
157
 
   nir_store_ssbo(&b, value, zero, dst_offset, .write_mask=0x1, .align_mul=1);
158
 
 
159
 
   return create_nir_cs(sctx, &b);
160
 
}
161
 
 
162
 
void *gfx9_create_clear_dcc_msaa_cs(struct si_context *sctx, struct si_texture *tex)
163
 
{
164
 
   const nir_shader_compiler_options *options =
165
 
      sctx->b.screen->get_compiler_options(sctx->b.screen, PIPE_SHADER_IR_NIR, PIPE_SHADER_COMPUTE);
166
 
 
167
 
   nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, options, "clear_dcc_msaa");
168
 
   b.shader->info.workgroup_size[0] = 8;
169
 
   b.shader->info.workgroup_size[1] = 8;
170
 
   b.shader->info.workgroup_size[2] = 1;
171
 
   b.shader->info.cs.user_data_components_amd = 2;
172
 
   b.shader->info.num_ssbos = 1;
173
 
 
174
 
   /* Get user data SGPRs. */
175
 
   nir_ssa_def *user_sgprs = nir_load_user_data_amd(&b);
176
 
   nir_ssa_def *dcc_pitch, *dcc_height, *clear_value, *pipe_xor;
177
 
   unpack_2x16(&b, nir_channel(&b, user_sgprs, 0), &dcc_pitch, &dcc_height);
178
 
   unpack_2x16(&b, nir_channel(&b, user_sgprs, 1), &clear_value, &pipe_xor);
179
 
   clear_value = nir_u2u16(&b, clear_value);
180
 
 
181
 
   /* Get the 2D coordinates. */
182
 
   nir_ssa_def *coord = get_global_ids(&b, 3);
183
 
   nir_ssa_def *zero = nir_imm_int(&b, 0);
184
 
 
185
 
   /* Multiply the coordinates by the DCC block size (they are DCC block coordinates). */
186
 
   coord = nir_imul(&b, coord,
187
 
                    nir_channels(&b, nir_imm_ivec4(&b, tex->surface.u.gfx9.color.dcc_block_width,
188
 
                                                   tex->surface.u.gfx9.color.dcc_block_height,
189
 
                                                   tex->surface.u.gfx9.color.dcc_block_depth, 0), 0x7));
190
 
 
191
 
   nir_ssa_def *offset =
192
 
      ac_nir_dcc_addr_from_coord(&b, &sctx->screen->info, tex->surface.bpe,
193
 
                                 &tex->surface.u.gfx9.color.dcc_equation,
194
 
                                 dcc_pitch, dcc_height, zero, /* DCC slice size */
195
 
                                 nir_channel(&b, coord, 0), nir_channel(&b, coord, 1), /* x, y */
196
 
                                 tex->buffer.b.b.array_size > 1 ? nir_channel(&b, coord, 2) : zero, /* z */
197
 
                                 zero, pipe_xor); /* sample, pipe_xor */
198
 
 
199
 
   /* The trick here is that DCC elements for an even and the next odd sample are next to each other
200
 
    * in memory, so we only need to compute the address for sample 0 and the next DCC byte is always
201
 
    * sample 1. That's why the clear value has 2 bytes - we're clearing 2 samples at the same time.
202
 
    */
203
 
   nir_store_ssbo(&b, clear_value, zero, offset, .write_mask=0x1, .align_mul=2);
204
 
 
205
 
   return create_nir_cs(sctx, &b);
206
 
}
207
 
 
208
 
/* Create a compute shader implementing clear_buffer or copy_buffer. */
209
 
void *si_create_clear_buffer_rmw_cs(struct si_context *sctx)
210
 
{
211
 
   const nir_shader_compiler_options *options =
212
 
      sctx->b.screen->get_compiler_options(sctx->b.screen, PIPE_SHADER_IR_NIR, PIPE_SHADER_COMPUTE);
213
 
 
214
 
   nir_builder b =
215
 
      nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, options, "clear_buffer_rmw_cs");
216
 
   b.shader->info.workgroup_size[0] = 64;
217
 
   b.shader->info.workgroup_size[1] = 1;
218
 
   b.shader->info.workgroup_size[2] = 1;
219
 
   b.shader->info.cs.user_data_components_amd = 2;
220
 
   b.shader->info.num_ssbos = 1;
221
 
 
222
 
   /* address = blockID * 64 + threadID; */
223
 
   nir_ssa_def *address = get_global_ids(&b, 1);
224
 
 
225
 
   /* address = address * 16; (byte offset, loading one vec4 per thread) */
226
 
   address = nir_ishl(&b, address, nir_imm_int(&b, 4));
227
 
   
228
 
   nir_ssa_def *zero = nir_imm_int(&b, 0);
229
 
   nir_ssa_def *data = nir_load_ssbo(&b, 4, 32, zero, address, .align_mul = 4);
230
 
 
231
 
   /* Get user data SGPRs. */
232
 
   nir_ssa_def *user_sgprs = nir_load_user_data_amd(&b);
233
 
 
234
 
   /* data &= inverted_writemask; */
235
 
   data = nir_iand(&b, data, nir_channel(&b, user_sgprs, 1));
236
 
   /* data |= clear_value_masked; */
237
 
   data = nir_ior(&b, data, nir_channel(&b, user_sgprs, 0));
238
 
 
239
 
   nir_store_ssbo(&b, data, zero, address,
240
 
      .access = SI_COMPUTE_DST_CACHE_POLICY != L2_LRU ? ACCESS_STREAM_CACHE_POLICY : 0,
241
 
      .align_mul = 4);
242
 
 
243
 
   return create_nir_cs(sctx, &b);
244
 
}
245