1 /*
2 * Copyright © 2022 Intel Corporation
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
20 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
21 * IN THE SOFTWARE.
22 */
23
24 #include "anv_private.h"
25
26 #include "compiler/intel_nir.h"
27 #include "compiler/brw_compiler.h"
28 #include "compiler/brw_nir.h"
29 #include "compiler/nir/nir.h"
30 #include "compiler/nir/nir_builder.h"
31 #include "dev/intel_debug.h"
32 #include "intel/compiler/intel_nir.h"
33 #include "util/macros.h"
34
35 #include "vk_nir.h"
36
37 #include "anv_internal_kernels.h"
38
39 static bool
lower_base_workgroup_id(nir_builder * b,nir_intrinsic_instr * intrin,UNUSED void * data)40 lower_base_workgroup_id(nir_builder *b, nir_intrinsic_instr *intrin,
41 UNUSED void *data)
42 {
43 if (intrin->intrinsic != nir_intrinsic_load_base_workgroup_id)
44 return false;
45
46 b->cursor = nir_instr_remove(&intrin->instr);
47 nir_def_rewrite_uses(&intrin->def, nir_imm_zero(b, 3, 32));
48 return true;
49 }
50
51 static void
link_libanv(nir_shader * nir,const nir_shader * libanv)52 link_libanv(nir_shader *nir, const nir_shader *libanv)
53 {
54 nir_link_shader_functions(nir, libanv);
55 NIR_PASS_V(nir, nir_inline_functions);
56 NIR_PASS_V(nir, nir_remove_non_entrypoints);
57 NIR_PASS_V(nir, nir_lower_vars_to_explicit_types, nir_var_function_temp,
58 glsl_get_cl_type_size_align);
59 NIR_PASS_V(nir, nir_opt_deref);
60 NIR_PASS_V(nir, nir_lower_vars_to_ssa);
61 NIR_PASS_V(nir, nir_lower_explicit_io,
62 nir_var_shader_temp | nir_var_function_temp | nir_var_mem_shared |
63 nir_var_mem_global,
64 nir_address_format_62bit_generic);
65 }
66
67 static struct anv_shader_bin *
compile_shader(struct anv_device * device,const nir_shader * libanv,enum anv_internal_kernel_name shader_name,gl_shader_stage stage,const char * name,const void * hash_key,uint32_t hash_key_size,uint32_t sends_count_expectation)68 compile_shader(struct anv_device *device,
69 const nir_shader *libanv,
70 enum anv_internal_kernel_name shader_name,
71 gl_shader_stage stage,
72 const char *name,
73 const void *hash_key,
74 uint32_t hash_key_size,
75 uint32_t sends_count_expectation)
76 {
77 const nir_shader_compiler_options *nir_options =
78 device->physical->compiler->nir_options[stage];
79
80 nir_builder b = nir_builder_init_simple_shader(stage, nir_options,
81 "%s", name);
82
83 uint32_t uniform_size =
84 anv_genX(device->info, call_internal_shader)(&b, shader_name);
85
86 nir_shader *nir = b.shader;
87
88 link_libanv(nir, libanv);
89
90 NIR_PASS_V(nir, nir_lower_vars_to_ssa);
91 NIR_PASS_V(nir, nir_opt_cse);
92 NIR_PASS_V(nir, nir_opt_gcm, true);
93 NIR_PASS_V(nir, nir_opt_peephole_select, 1, false, false);
94
95 NIR_PASS_V(nir, nir_lower_variable_initializers, ~0);
96
97 NIR_PASS_V(nir, nir_split_var_copies);
98 NIR_PASS_V(nir, nir_split_per_member_structs);
99
100 if (stage == MESA_SHADER_COMPUTE) {
101 nir->info.workgroup_size[0] = 16;
102 nir->info.workgroup_size[1] = 1;
103 nir->info.workgroup_size[2] = 1;
104 }
105
106 struct brw_compiler *compiler = device->physical->compiler;
107 struct brw_nir_compiler_opts opts = {};
108 brw_preprocess_nir(compiler, nir, &opts);
109
110 NIR_PASS_V(nir, nir_propagate_invariant, false);
111
112 if (stage == MESA_SHADER_FRAGMENT) {
113 NIR_PASS_V(nir, nir_lower_input_attachments,
114 &(nir_input_attachment_options) {
115 .use_fragcoord_sysval = true,
116 .use_layer_id_sysval = true,
117 });
118 } else {
119 nir_lower_compute_system_values_options options = {
120 .has_base_workgroup_id = true,
121 .lower_cs_local_id_to_index = true,
122 .lower_workgroup_id_to_index = gl_shader_stage_is_mesh(stage),
123 };
124 NIR_PASS_V(nir, nir_lower_compute_system_values, &options);
125 NIR_PASS_V(nir, nir_shader_intrinsics_pass, lower_base_workgroup_id,
126 nir_metadata_control_flow, NULL);
127 }
128
129 /* Reset sizes before gathering information */
130 nir->global_mem_size = 0;
131 nir->scratch_size = 0;
132 nir->info.shared_size = 0;
133 nir_shader_gather_info(nir, nir_shader_get_entrypoint(nir));
134
135 NIR_PASS_V(nir, nir_copy_prop);
136 NIR_PASS_V(nir, nir_opt_constant_folding);
137 NIR_PASS_V(nir, nir_opt_dce);
138
139 union brw_any_prog_key key;
140 memset(&key, 0, sizeof(key));
141
142 union brw_any_prog_data prog_data;
143 memset(&prog_data, 0, sizeof(prog_data));
144
145 if (stage == MESA_SHADER_COMPUTE) {
146 NIR_PASS_V(nir, brw_nir_lower_cs_intrinsics,
147 device->info, &prog_data.cs);
148 }
149
150 /* Do vectorizing here. For some reason when trying to do it in the back
151 * this just isn't working.
152 */
153 nir_load_store_vectorize_options options = {
154 .modes = nir_var_mem_ubo | nir_var_mem_ssbo | nir_var_mem_global,
155 .callback = brw_nir_should_vectorize_mem,
156 .robust_modes = (nir_variable_mode)0,
157 };
158 NIR_PASS_V(nir, nir_opt_load_store_vectorize, &options);
159
160 nir->num_uniforms = uniform_size;
161
162 prog_data.base.nr_params = nir->num_uniforms / 4;
163
164 brw_nir_analyze_ubo_ranges(compiler, nir, prog_data.base.ubo_ranges);
165
166 void *temp_ctx = ralloc_context(NULL);
167
168 const unsigned *program;
169 if (stage == MESA_SHADER_FRAGMENT) {
170 struct brw_compile_stats stats[3];
171 struct brw_compile_fs_params params = {
172 .base = {
173 .nir = nir,
174 .log_data = device,
175 .debug_flag = DEBUG_WM,
176 .stats = stats,
177 .mem_ctx = temp_ctx,
178 },
179 .key = &key.wm,
180 .prog_data = &prog_data.wm,
181 };
182 program = brw_compile_fs(compiler, ¶ms);
183
184 if (!INTEL_DEBUG(DEBUG_SHADER_PRINT)) {
185 unsigned stat_idx = 0;
186 if (prog_data.wm.dispatch_8) {
187 assert(stats[stat_idx].spills == 0);
188 assert(stats[stat_idx].fills == 0);
189 assert(stats[stat_idx].sends == sends_count_expectation);
190 stat_idx++;
191 }
192 if (prog_data.wm.dispatch_16) {
193 assert(stats[stat_idx].spills == 0);
194 assert(stats[stat_idx].fills == 0);
195 assert(stats[stat_idx].sends == sends_count_expectation);
196 stat_idx++;
197 }
198 if (prog_data.wm.dispatch_32) {
199 assert(stats[stat_idx].spills == 0);
200 assert(stats[stat_idx].fills == 0);
201 assert(stats[stat_idx].sends ==
202 sends_count_expectation *
203 (device->info->ver < 20 ? 2 : 1));
204 stat_idx++;
205 }
206 }
207 } else {
208 struct brw_compile_stats stats;
209 struct brw_compile_cs_params params = {
210 .base = {
211 .nir = nir,
212 .stats = &stats,
213 .log_data = device,
214 .debug_flag = DEBUG_CS,
215 .mem_ctx = temp_ctx,
216 },
217 .key = &key.cs,
218 .prog_data = &prog_data.cs,
219 };
220 program = brw_compile_cs(compiler, ¶ms);
221
222 if (!INTEL_DEBUG(DEBUG_SHADER_PRINT)) {
223 assert(stats.spills == 0);
224 assert(stats.fills == 0);
225 assert(stats.sends == sends_count_expectation);
226 }
227 }
228
229 assert(prog_data.base.total_scratch == 0);
230 assert(program != NULL);
231 struct anv_shader_bin *kernel = NULL;
232 if (program == NULL)
233 goto exit;
234
235 struct anv_pipeline_bind_map empty_bind_map = {};
236 struct anv_push_descriptor_info empty_push_desc_info = {};
237 struct anv_shader_upload_params upload_params = {
238 .stage = nir->info.stage,
239 .key_data = hash_key,
240 .key_size = hash_key_size,
241 .kernel_data = program,
242 .kernel_size = prog_data.base.program_size,
243 .prog_data = &prog_data.base,
244 .prog_data_size = sizeof(prog_data),
245 .bind_map = &empty_bind_map,
246 .push_desc_info = &empty_push_desc_info,
247 };
248
249 kernel = anv_device_upload_kernel(device, device->internal_cache, &upload_params);
250
251 exit:
252 ralloc_free(temp_ctx);
253 ralloc_free(nir);
254
255 return kernel;
256 }
257
258 VkResult
anv_device_get_internal_shader(struct anv_device * device,enum anv_internal_kernel_name name,struct anv_shader_bin ** out_bin)259 anv_device_get_internal_shader(struct anv_device *device,
260 enum anv_internal_kernel_name name,
261 struct anv_shader_bin **out_bin)
262 {
263 const struct {
264 struct {
265 char name[40];
266 } key;
267
268 gl_shader_stage stage;
269
270 uint32_t send_count;
271 } internal_kernels[] = {
272 [ANV_INTERNAL_KERNEL_GENERATED_DRAWS] = {
273 .key = {
274 .name = "anv-generated-indirect-draws",
275 },
276 .stage = MESA_SHADER_FRAGMENT,
277 .send_count = (device->info->ver == 9 ?
278 /* 1 load +
279 * 4 stores +
280 * 2 * (2 loads + 2 stores) +
281 * 3 stores
282 */
283 16 :
284 /* 1 load +
285 * 2 * (2 loads + 3 stores) +
286 * 3 stores
287 */
288 14) +
289 /* 3 loads + 3 stores */
290 (intel_needs_workaround(device->info, 16011107343) ? 6 : 0) +
291 /* 3 loads + 3 stores */
292 (intel_needs_workaround(device->info, 22018402687) ? 6 : 0),
293 },
294 [ANV_INTERNAL_KERNEL_COPY_QUERY_RESULTS_COMPUTE] = {
295 .key = {
296 .name = "anv-copy-query-compute",
297 },
298 .stage = MESA_SHADER_COMPUTE,
299 .send_count = device->info->verx10 >= 125 ?
300 9 /* 4 loads + 4 stores + 1 EOT */ :
301 8 /* 3 loads + 4 stores + 1 EOT */,
302 },
303 [ANV_INTERNAL_KERNEL_COPY_QUERY_RESULTS_FRAGMENT] = {
304 .key = {
305 .name = "anv-copy-query-fragment",
306 },
307 .stage = MESA_SHADER_FRAGMENT,
308 .send_count = 8 /* 3 loads + 4 stores + 1 EOT */,
309 },
310 [ANV_INTERNAL_KERNEL_MEMCPY_COMPUTE] = {
311 .key = {
312 .name = "anv-memcpy-compute",
313 },
314 .stage = MESA_SHADER_COMPUTE,
315 .send_count = device->info->verx10 >= 125 ?
316 10 /* 5 loads (1 pull constants) + 4 stores + 1 EOT */ :
317 9 /* 4 loads + 4 stores + 1 EOT */,
318 },
319 };
320
321 struct anv_shader_bin *bin =
322 p_atomic_read(&device->internal_kernels[name]);
323 if (bin != NULL) {
324 *out_bin = bin;
325 return VK_SUCCESS;
326 }
327
328 bin =
329 anv_device_search_for_kernel(device,
330 device->internal_cache,
331 &internal_kernels[name].key,
332 sizeof(internal_kernels[name].key),
333 NULL);
334 if (bin != NULL) {
335 p_atomic_set(&device->internal_kernels[name], bin);
336 *out_bin = bin;
337 return VK_SUCCESS;
338 }
339
340 void *mem_ctx = ralloc_context(NULL);
341
342 nir_shader *libanv_shaders =
343 anv_genX(device->info, load_libanv_shader)(device, mem_ctx);
344
345 bin = compile_shader(device,
346 libanv_shaders,
347 name,
348 internal_kernels[name].stage,
349 internal_kernels[name].key.name,
350 &internal_kernels[name].key,
351 sizeof(internal_kernels[name].key),
352 internal_kernels[name].send_count);
353 if (bin == NULL)
354 return vk_errorf(device, VK_ERROR_OUT_OF_HOST_MEMORY,
355 "Unable to compiler internal kernel");
356
357 /* The cache already has a reference and it's not going anywhere so
358 * there is no need to hold a second reference.
359 */
360 anv_shader_bin_unref(device, bin);
361
362 p_atomic_set(&device->internal_kernels[name], bin);
363
364 *out_bin = bin;
365 return VK_SUCCESS;
366 }
367
368 VkResult
anv_device_init_internal_kernels(struct anv_device * device)369 anv_device_init_internal_kernels(struct anv_device *device)
370 {
371 const struct intel_l3_weights w =
372 intel_get_default_l3_weights(device->info,
373 true /* wants_dc_cache */,
374 false /* needs_slm */);
375 device->internal_kernels_l3_config = intel_get_l3_config(device->info, w);
376
377 return VK_SUCCESS;
378 }
379
380 void
anv_device_finish_internal_kernels(struct anv_device * device)381 anv_device_finish_internal_kernels(struct anv_device *device)
382 {
383 }
384