1 /**************************************************************************
2 *
3 * Copyright 2019 Red Hat.
4 * All Rights Reserved.
5 *
6 * Permission is hereby granted, free of charge, to any person obtaining a
7 * copy of this software and associated documentation files (the "Software"),
8 * to deal in the Software without restriction, including without limitation
9 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
10 * and/or sell copies of the Software, and to permit persons to whom the
11 * Software is furnished to do so, subject to the following conditions:
12 *
13 * The above copyright notice and this permission notice shall be included
14 * in all copies or substantial portions of the 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 "util/u_memory.h"
27 #include "util/os_time.h"
28 #include "util/u_dump.h"
29 #include "util/u_string.h"
30 #include "gallivm/lp_bld_const.h"
31 #include "gallivm/lp_bld_debug.h"
32 #include "gallivm/lp_bld_intr.h"
33 #include "gallivm/lp_bld_flow.h"
34 #include "gallivm/lp_bld_pack.h"
35 #include "gallivm/lp_bld_gather.h"
36 #include "gallivm/lp_bld_coro.h"
37 #include "gallivm/lp_bld_nir.h"
38 #include "gallivm/lp_bld_jit_sample.h"
39 #include "lp_state_cs.h"
40 #include "lp_context.h"
41 #include "lp_setup_context.h"
42 #include "lp_debug.h"
43 #include "lp_state.h"
44 #include "lp_perf.h"
45 #include "lp_screen.h"
46 #include "lp_memory.h"
47 #include "lp_query.h"
48 #include "lp_cs_tpool.h"
49 #include "frontend/sw_winsys.h"
50 #include "nir/nir_to_tgsi_info.h"
51 #include "nir/tgsi_to_nir.h"
52 #include "util/mesa-sha1.h"
53 #include "nir_serialize.h"
54
55 #include "draw/draw_context.h"
56 #include "draw/draw_llvm.h"
57 #include "draw/draw_mesh_prim.h"
58
59 /** Fragment shader number (for debugging) */
60 static unsigned cs_no = 0;
61 static unsigned task_no = 0;
62 static unsigned mesh_no = 0;
63
64 struct lp_cs_job_info {
65 unsigned grid_size[3];
66 unsigned iter_size[3];
67 unsigned grid_base[3];
68 unsigned block_size[3];
69 unsigned req_local_mem;
70 unsigned work_dim;
71 unsigned draw_id;
72 bool zero_initialize_shared_memory;
73 bool use_iters;
74 struct lp_cs_exec *current;
75 struct vertex_header *io;
76 size_t io_stride;
77 void *payload;
78 size_t payload_stride;
79 };
80
81 enum {
82 CS_ARG_CONTEXT,
83 CS_ARG_RESOURCES,
84 CS_ARG_BLOCK_X_SIZE,
85 CS_ARG_BLOCK_Y_SIZE,
86 CS_ARG_BLOCK_Z_SIZE,
87 CS_ARG_GRID_X,
88 CS_ARG_GRID_Y,
89 CS_ARG_GRID_Z,
90 CS_ARG_GRID_SIZE_X,
91 CS_ARG_GRID_SIZE_Y,
92 CS_ARG_GRID_SIZE_Z,
93 CS_ARG_WORK_DIM,
94 CS_ARG_DRAW_ID,
95 CS_ARG_VERTEX_DATA,
96 CS_ARG_PER_THREAD_DATA,
97 CS_ARG_OUTER_COUNT,
98 CS_ARG_CORO_SUBGROUP_COUNT = CS_ARG_OUTER_COUNT,
99 CS_ARG_CORO_PARTIALS,
100 CS_ARG_CORO_BLOCK_X_SIZE,
101 CS_ARG_CORO_BLOCK_Y_SIZE,
102 CS_ARG_CORO_BLOCK_Z_SIZE,
103 CS_ARG_CORO_IDX,
104 CS_ARG_CORO_MEM,
105 CS_ARG_CORO_OUTPUTS,
106 CS_ARG_MAX,
107 };
108
109 struct lp_mesh_llvm_iface {
110 struct lp_build_mesh_iface base;
111
112 LLVMValueRef vertex_count;
113 LLVMValueRef prim_count;
114 LLVMValueRef outputs;
115 };
116
117 static inline const struct lp_mesh_llvm_iface *
lp_mesh_llvm_iface(const struct lp_build_mesh_iface * iface)118 lp_mesh_llvm_iface(const struct lp_build_mesh_iface *iface)
119 {
120 return (const struct lp_mesh_llvm_iface *)iface;
121 }
122
123
124 static LLVMTypeRef
create_mesh_jit_output_type_deref(struct gallivm_state * gallivm)125 create_mesh_jit_output_type_deref(struct gallivm_state *gallivm)
126 {
127 LLVMTypeRef float_type = LLVMFloatTypeInContext(gallivm->context);
128 LLVMTypeRef output_array;
129
130 output_array = LLVMArrayType(float_type, TGSI_NUM_CHANNELS); /* num channels */
131 output_array = LLVMArrayType(output_array, PIPE_MAX_SHADER_OUTPUTS); /* num attrs per vertex */
132 return output_array;
133 }
134
135 static void
lp_mesh_llvm_emit_store_output(const struct lp_build_mesh_iface * mesh_iface,struct lp_build_context * bld,unsigned name,bool is_vindex_indirect,LLVMValueRef vertex_index,bool is_aindex_indirect,LLVMValueRef attrib_index,bool is_sindex_indirect,LLVMValueRef swizzle_index,LLVMValueRef value,LLVMValueRef mask_vec)136 lp_mesh_llvm_emit_store_output(const struct lp_build_mesh_iface *mesh_iface,
137 struct lp_build_context *bld,
138 unsigned name,
139 bool is_vindex_indirect,
140 LLVMValueRef vertex_index,
141 bool is_aindex_indirect,
142 LLVMValueRef attrib_index,
143 bool is_sindex_indirect,
144 LLVMValueRef swizzle_index,
145 LLVMValueRef value,
146 LLVMValueRef mask_vec)
147 {
148 const struct lp_mesh_llvm_iface *mesh = lp_mesh_llvm_iface(mesh_iface);
149 struct gallivm_state *gallivm = bld->gallivm;
150 LLVMBuilderRef builder = gallivm->builder;
151 LLVMValueRef indices[3];
152 LLVMValueRef res;
153 struct lp_type type = bld->type;
154 LLVMTypeRef output_type = create_mesh_jit_output_type_deref(gallivm);
155
156 if (is_vindex_indirect || is_aindex_indirect || is_sindex_indirect) {
157 for (int i = 0; i < type.length; ++i) {
158 LLVMValueRef idx = lp_build_const_int32(gallivm, i);
159 LLVMValueRef vert_chan_index = vertex_index ? vertex_index : lp_build_const_int32(gallivm, 0);
160 LLVMValueRef attr_chan_index = attrib_index;
161 LLVMValueRef swiz_chan_index = swizzle_index;
162 LLVMValueRef channel_vec;
163
164 if (is_vindex_indirect) {
165 vert_chan_index = LLVMBuildExtractElement(builder,
166 vertex_index, idx, "");
167 }
168 if (is_aindex_indirect) {
169 attr_chan_index = LLVMBuildExtractElement(builder,
170 attrib_index, idx, "");
171 }
172
173 if (is_sindex_indirect) {
174 swiz_chan_index = LLVMBuildExtractElement(builder,
175 swizzle_index, idx, "");
176 }
177
178 indices[0] = vert_chan_index;
179 indices[1] = attr_chan_index;
180 indices[2] = swiz_chan_index;
181
182 channel_vec = LLVMBuildGEP2(builder, output_type, mesh->outputs, indices, 3, "");
183
184 res = LLVMBuildExtractElement(builder, value, idx, "");
185
186 struct lp_build_if_state ifthen;
187 LLVMValueRef cond = LLVMBuildICmp(gallivm->builder, LLVMIntNE, mask_vec, lp_build_const_int_vec(gallivm, bld->type, 0), "");
188 cond = LLVMBuildExtractElement(gallivm->builder, cond, idx, "");
189 lp_build_if(&ifthen, gallivm, cond);
190 LLVMBuildStore(builder, res, channel_vec);
191 lp_build_endif(&ifthen);
192 }
193 } else {
194 indices[0] = vertex_index ? vertex_index : lp_build_const_int32(gallivm, 0);
195 indices[1] = attrib_index;
196 indices[2] = swizzle_index;
197
198 res = LLVMBuildGEP2(builder, output_type, mesh->outputs, indices, 3, "");
199 for (unsigned i = 0; i < type.length; ++i) {
200 LLVMValueRef idx = lp_build_const_int32(gallivm, i);
201 LLVMValueRef val = LLVMBuildExtractElement(builder, value, idx, "");
202
203 struct lp_build_if_state ifthen;
204 LLVMValueRef cond = LLVMBuildICmp(gallivm->builder, LLVMIntNE, mask_vec, lp_build_const_int_vec(gallivm, bld->type, 0), "");
205 cond = LLVMBuildExtractElement(gallivm->builder, cond, idx, "");
206 lp_build_if(&ifthen, gallivm, cond);
207 LLVMBuildStore(builder, val, res);
208 lp_build_endif(&ifthen);
209 }
210 }
211 }
212
213 static void
lp_mesh_emit_vertex_and_primitive_count(const struct lp_build_mesh_iface * mesh_iface,struct lp_build_context * bld,LLVMValueRef vertices_count,LLVMValueRef primitives_count)214 lp_mesh_emit_vertex_and_primitive_count(const struct lp_build_mesh_iface *mesh_iface,
215 struct lp_build_context *bld,
216 LLVMValueRef vertices_count,
217 LLVMValueRef primitives_count)
218 {
219 const struct lp_mesh_llvm_iface *mesh = lp_mesh_llvm_iface(mesh_iface);
220 struct gallivm_state *gallivm = bld->gallivm;
221
222 LLVMBuildStore(gallivm->builder, vertices_count, mesh->vertex_count);
223 LLVMBuildStore(gallivm->builder, primitives_count, mesh->prim_count);
224 }
225
226 static void
mesh_convert_to_aos(struct gallivm_state * gallivm,nir_shader * nir,bool vert_only,LLVMTypeRef io_type,LLVMValueRef io,LLVMValueRef outputs,LLVMValueRef clipmask,LLVMValueRef vertex_index,struct lp_type soa_type,int primid_slot,bool need_edgeflag)227 mesh_convert_to_aos(struct gallivm_state *gallivm,
228 nir_shader *nir,
229 bool vert_only,
230 LLVMTypeRef io_type,
231 LLVMValueRef io,
232 LLVMValueRef outputs,
233 LLVMValueRef clipmask,
234 LLVMValueRef vertex_index,
235 struct lp_type soa_type,
236 int primid_slot,
237 bool need_edgeflag)
238 {
239 LLVMBuilderRef builder = gallivm->builder;
240 LLVMValueRef inds[3];
241 LLVMTypeRef output_type = create_mesh_jit_output_type_deref(gallivm);
242 #if DEBUG_STORE
243 lp_build_printf(gallivm, " # storing begin\n");
244 #endif
245 int first_per_prim_attrib = -1;
246 nir_foreach_shader_out_variable(var, nir) {
247 if (var->data.per_primitive) {
248 first_per_prim_attrib = var->data.driver_location;
249 break;
250 }
251 }
252 nir_foreach_shader_out_variable(var, nir) {
253
254 if (vert_only && var->data.per_primitive)
255 continue;
256 if (!vert_only && !var->data.per_primitive)
257 continue;
258 int attrib = var->data.driver_location;
259 int slots = glsl_count_attribute_slots(glsl_get_array_element(var->type), false);
260
261 for (unsigned s = 0; s < slots; s++) {
262 LLVMValueRef soa[TGSI_NUM_CHANNELS];
263 LLVMValueRef aos[LP_MAX_VECTOR_WIDTH / 32];
264 for (unsigned chan = 0; chan < TGSI_NUM_CHANNELS; ++chan) {
265 inds[0] = vertex_index;
266 inds[1] = lp_build_const_int32(gallivm, attrib);
267 inds[2] = lp_build_const_int32(gallivm, chan);
268
269 LLVMValueRef res = LLVMBuildGEP2(builder, output_type, outputs, inds, 3, "");
270 LLVMTypeRef single_type = (attrib == primid_slot) ? lp_build_int_elem_type(gallivm, soa_type) : lp_build_elem_type(gallivm, soa_type);
271 LLVMValueRef out = LLVMBuildLoad2(builder, single_type, res, "");
272 lp_build_name(out, "output%u.%c", attrib, "xyzw"[chan]);
273 #if DEBUG_STORE
274 lp_build_printf(gallivm, "output %d : %d ",
275 LLVMConstInt(LLVMInt32TypeInContext(gallivm->context),
276 attrib, 0),
277 LLVMConstInt(LLVMInt32TypeInContext(gallivm->context),
278 chan, 0));
279 lp_build_print_value(gallivm, "val = ", out);
280 {
281 LLVMValueRef iv =
282 LLVMBuildBitCast(builder, out, lp_build_int_elem_type(gallivm, soa_type), "");
283
284 lp_build_print_value(gallivm, " ival = ", iv);
285 }
286 #endif
287 soa[chan] = out;
288 }
289 LLVMTypeRef float_type = LLVMFloatTypeInContext(gallivm->context);
290 aos[0] = LLVMGetUndef(LLVMVectorType(float_type, 4));
291 for (unsigned i = 0; i < 4; i++)
292 aos[0] = LLVMBuildInsertElement(builder, aos[0], soa[i], lp_build_const_int32(gallivm, i), "");
293 int aos_attrib = attrib;
294 if (var->data.per_primitive)
295 aos_attrib -= first_per_prim_attrib;
296 draw_store_aos_array(gallivm,
297 soa_type,
298 io_type,
299 io,
300 NULL,
301 aos,
302 aos_attrib,
303 clipmask,
304 need_edgeflag, var->data.per_primitive);
305 attrib++;
306 }
307 }
308 #if DEBUG_STORE
309 lp_build_printf(gallivm, " # storing end\n");
310 #endif
311 }
312
313 static void
generate_compute(struct llvmpipe_context * lp,struct lp_compute_shader * shader,struct lp_compute_shader_variant * variant)314 generate_compute(struct llvmpipe_context *lp,
315 struct lp_compute_shader *shader,
316 struct lp_compute_shader_variant *variant)
317 {
318 struct gallivm_state *gallivm = variant->gallivm;
319 struct nir_shader *nir = shader->base.ir.nir;
320 const struct lp_compute_shader_variant_key *key = &variant->key;
321 char func_name[64], func_name_coro[64];
322 LLVMTypeRef arg_types[CS_ARG_MAX];
323 LLVMTypeRef func_type, coro_func_type;
324 LLVMTypeRef int32_type = LLVMInt32TypeInContext(gallivm->context);
325 LLVMValueRef context_ptr, resources_ptr;
326 LLVMValueRef block_x_size_arg, block_y_size_arg, block_z_size_arg;
327 LLVMValueRef grid_x_arg, grid_y_arg, grid_z_arg;
328 LLVMValueRef grid_size_x_arg, grid_size_y_arg, grid_size_z_arg;
329 LLVMValueRef work_dim_arg, draw_id_arg, thread_data_ptr, io_ptr;
330 LLVMBasicBlockRef block;
331 LLVMBuilderRef builder;
332 struct lp_build_sampler_soa *sampler;
333 struct lp_build_image_soa *image;
334 LLVMValueRef function, coro;
335 struct lp_type cs_type;
336 struct lp_mesh_llvm_iface mesh_iface;
337 bool is_mesh = nir->info.stage == MESA_SHADER_MESH;
338 unsigned i;
339
340 LLVMValueRef output_array = NULL;
341
342 /*
343 * This function has two parts
344 * a) setup the coroutine execution environment loop.
345 * b) build the compute shader llvm for use inside the coroutine.
346 */
347 assert(lp_native_vector_width / 32 >= 4);
348
349 memset(&cs_type, 0, sizeof cs_type);
350 cs_type.floating = true; /* floating point values */
351 cs_type.sign = true; /* values are signed */
352 cs_type.norm = false; /* values are not limited to [0,1] or [-1,1] */
353 cs_type.width = 32; /* 32-bit float */
354 cs_type.length = MIN2(lp_native_vector_width / 32, 16); /* n*4 elements per vector */
355 snprintf(func_name, sizeof(func_name), "cs_variant");
356
357 snprintf(func_name_coro, sizeof(func_name), "cs_co_variant");
358
359 arg_types[CS_ARG_CONTEXT] = variant->jit_cs_context_ptr_type; /* context */
360 arg_types[CS_ARG_RESOURCES]= variant->jit_resources_ptr_type;
361 arg_types[CS_ARG_BLOCK_X_SIZE] = int32_type; /* block_x_size */
362 arg_types[CS_ARG_BLOCK_Y_SIZE] = int32_type; /* block_y_size */
363 arg_types[CS_ARG_BLOCK_Z_SIZE] = int32_type; /* block_z_size */
364 arg_types[CS_ARG_GRID_X] = int32_type; /* grid_x */
365 arg_types[CS_ARG_GRID_Y] = int32_type; /* grid_y */
366 arg_types[CS_ARG_GRID_Z] = int32_type; /* grid_z */
367 arg_types[CS_ARG_GRID_SIZE_X] = int32_type; /* grid_size_x */
368 arg_types[CS_ARG_GRID_SIZE_Y] = int32_type; /* grid_size_y */
369 arg_types[CS_ARG_GRID_SIZE_Z] = int32_type; /* grid_size_z */
370 arg_types[CS_ARG_WORK_DIM] = int32_type; /* work dim */
371 arg_types[CS_ARG_DRAW_ID] = int32_type; /* draw id */
372 if (variant->jit_vertex_header_ptr_type)
373 arg_types[CS_ARG_VERTEX_DATA] = variant->jit_vertex_header_ptr_type; /* mesh shaders only */
374 else
375 arg_types[CS_ARG_VERTEX_DATA] = LLVMPointerType(LLVMInt8TypeInContext(gallivm->context), 0); /* mesh shaders only */
376 arg_types[CS_ARG_PER_THREAD_DATA] = variant->jit_cs_thread_data_ptr_type; /* per thread data */
377 arg_types[CS_ARG_CORO_SUBGROUP_COUNT] = int32_type; /* coro only - subgroup count */
378 arg_types[CS_ARG_CORO_PARTIALS] = int32_type; /* coro only - partials */
379 arg_types[CS_ARG_CORO_BLOCK_X_SIZE] = int32_type; /* coro block_x_size */
380 arg_types[CS_ARG_CORO_BLOCK_Y_SIZE] = int32_type; /* coro block_y_size */
381 arg_types[CS_ARG_CORO_BLOCK_Z_SIZE] = int32_type; /* coro block_z_size */
382 arg_types[CS_ARG_CORO_IDX] = int32_type; /* coro idx */
383 arg_types[CS_ARG_CORO_MEM] = LLVMPointerType(LLVMPointerType(LLVMInt8TypeInContext(gallivm->context), 0), 0);
384 arg_types[CS_ARG_CORO_OUTPUTS] = LLVMPointerType(LLVMInt8TypeInContext(gallivm->context), 0); /* mesh shaders only */
385
386 func_type = LLVMFunctionType(LLVMVoidTypeInContext(gallivm->context),
387 arg_types, CS_ARG_OUTER_COUNT, 0);
388
389 coro_func_type = LLVMFunctionType(LLVMPointerType(LLVMInt8TypeInContext(gallivm->context), 0),
390 arg_types, CS_ARG_MAX - (!is_mesh), 0);
391
392 function = LLVMAddFunction(gallivm->module, func_name, func_type);
393 LLVMSetFunctionCallConv(function, LLVMCCallConv);
394
395 coro = LLVMAddFunction(gallivm->module, func_name_coro, coro_func_type);
396 LLVMSetFunctionCallConv(coro, LLVMCCallConv);
397 lp_build_coro_add_presplit(coro);
398
399 variant->function = function;
400
401 for (i = 0; i < CS_ARG_MAX - !is_mesh; ++i) {
402 if (LLVMGetTypeKind(arg_types[i]) == LLVMPointerTypeKind) {
403 lp_add_function_attr(coro, i + 1, LP_FUNC_ATTR_NOALIAS);
404 if (i < CS_ARG_OUTER_COUNT)
405 lp_add_function_attr(function, i + 1, LP_FUNC_ATTR_NOALIAS);
406 }
407 }
408
409 if (variant->gallivm->cache->data_size)
410 return;
411
412 context_ptr = LLVMGetParam(function, CS_ARG_CONTEXT);
413 resources_ptr = LLVMGetParam(function, CS_ARG_RESOURCES);
414 block_x_size_arg = LLVMGetParam(function, CS_ARG_BLOCK_X_SIZE);
415 block_y_size_arg = LLVMGetParam(function, CS_ARG_BLOCK_Y_SIZE);
416 block_z_size_arg = LLVMGetParam(function, CS_ARG_BLOCK_Z_SIZE);
417 grid_x_arg = LLVMGetParam(function, CS_ARG_GRID_X);
418 grid_y_arg = LLVMGetParam(function, CS_ARG_GRID_Y);
419 grid_z_arg = LLVMGetParam(function, CS_ARG_GRID_Z);
420 grid_size_x_arg = LLVMGetParam(function, CS_ARG_GRID_SIZE_X);
421 grid_size_y_arg = LLVMGetParam(function, CS_ARG_GRID_SIZE_Y);
422 grid_size_z_arg = LLVMGetParam(function, CS_ARG_GRID_SIZE_Z);
423 work_dim_arg = LLVMGetParam(function, CS_ARG_WORK_DIM);
424 draw_id_arg = LLVMGetParam(function, CS_ARG_DRAW_ID);
425 io_ptr = LLVMGetParam(function, CS_ARG_VERTEX_DATA);
426 thread_data_ptr = LLVMGetParam(function, CS_ARG_PER_THREAD_DATA);
427
428 lp_build_name(context_ptr, "context");
429 lp_build_name(resources_ptr, "resources");
430 lp_build_name(block_x_size_arg, "x_size");
431 lp_build_name(block_y_size_arg, "y_size");
432 lp_build_name(block_z_size_arg, "z_size");
433 lp_build_name(grid_x_arg, "grid_x");
434 lp_build_name(grid_y_arg, "grid_y");
435 lp_build_name(grid_z_arg, "grid_z");
436 lp_build_name(grid_size_x_arg, "grid_size_x");
437 lp_build_name(grid_size_y_arg, "grid_size_y");
438 lp_build_name(grid_size_z_arg, "grid_size_z");
439 lp_build_name(work_dim_arg, "work_dim");
440 lp_build_name(draw_id_arg, "draw_id");
441 lp_build_name(thread_data_ptr, "thread_data");
442 lp_build_name(io_ptr, "vertex_io");
443
444 lp_build_nir_prepasses(nir);
445 struct hash_table *fns = _mesa_pointer_hash_table_create(NULL);
446
447 if (exec_list_length(&nir->functions) > 1) {
448 LLVMTypeRef call_context_type = lp_build_cs_func_call_context(gallivm, cs_type.length,
449 variant->jit_cs_context_type,
450 variant->jit_resources_type);
451 nir_foreach_function(func, nir) {
452 if (func->is_entrypoint)
453 continue;
454
455 LLVMTypeRef args[32];
456 int num_args;
457
458 num_args = func->num_params + LP_RESV_FUNC_ARGS;
459
460 args[0] = LLVMVectorType(LLVMInt32TypeInContext(gallivm->context), cs_type.length); /* mask */
461 args[1] = LLVMPointerType(call_context_type, 0);
462 for (int i = 0; i < func->num_params; i++) {
463 args[i + LP_RESV_FUNC_ARGS] = LLVMVectorType(LLVMIntTypeInContext(gallivm->context, func->params[i].bit_size), cs_type.length);
464 if (func->params[i].num_components > 1)
465 args[i + LP_RESV_FUNC_ARGS] = LLVMArrayType(args[i + LP_RESV_FUNC_ARGS], func->params[i].num_components);
466 }
467
468 LLVMTypeRef func_type = LLVMFunctionType(LLVMVoidTypeInContext(gallivm->context),
469 args, num_args, 0);
470 LLVMValueRef lfunc = LLVMAddFunction(gallivm->module, func->name, func_type);
471 LLVMSetFunctionCallConv(lfunc, LLVMCCallConv);
472
473 struct lp_build_fn *new_fn = ralloc(fns, struct lp_build_fn);
474 new_fn->fn_type = func_type;
475 new_fn->fn = lfunc;
476 _mesa_hash_table_insert(fns, func, new_fn);
477 }
478
479 nir_foreach_function(func, nir) {
480 if (func->is_entrypoint)
481 continue;
482
483 struct hash_entry *entry = _mesa_hash_table_search(fns, func);
484 assert(entry);
485 struct lp_build_fn *new_fn = entry->data;
486 LLVMValueRef lfunc = new_fn->fn;
487 block = LLVMAppendBasicBlockInContext(gallivm->context, lfunc, "entry");
488
489 builder = gallivm->builder;
490 LLVMPositionBuilderAtEnd(builder, block);
491 LLVMValueRef mask_param = LLVMGetParam(lfunc, 0);
492 LLVMValueRef call_context_ptr = LLVMGetParam(lfunc, 1);
493 LLVMValueRef call_context = LLVMBuildLoad2(builder, call_context_type, call_context_ptr, "");
494 struct lp_build_mask_context mask;
495 struct lp_bld_tgsi_system_values system_values;
496
497 memset(&system_values, 0, sizeof(system_values));
498
499 lp_build_mask_begin(&mask, gallivm, cs_type, mask_param);
500 lp_build_mask_check(&mask);
501
502 struct lp_build_tgsi_params params;
503 memset(¶ms, 0, sizeof(params));
504 params.type = cs_type;
505 params.mask = &mask;
506 params.fns = fns;
507 params.current_func = lfunc;
508 params.context_type = variant->jit_cs_context_type;
509 params.resources_type = variant->jit_resources_type;
510 params.call_context_ptr = call_context_ptr;
511 params.context_ptr = LLVMBuildExtractValue(builder, call_context, LP_NIR_CALL_CONTEXT_CONTEXT, "");
512 params.resources_ptr = LLVMBuildExtractValue(builder, call_context, LP_NIR_CALL_CONTEXT_RESOURCES, "");
513 params.shared_ptr = LLVMBuildExtractValue(builder, call_context, LP_NIR_CALL_CONTEXT_SHARED, "");
514 params.scratch_ptr = LLVMBuildExtractValue(builder, call_context, LP_NIR_CALL_CONTEXT_SCRATCH, "");
515 system_values.work_dim = LLVMBuildExtractValue(builder, call_context, LP_NIR_CALL_CONTEXT_WORK_DIM, "");
516 system_values.thread_id[0] = LLVMBuildExtractValue(builder, call_context, LP_NIR_CALL_CONTEXT_THREAD_ID_0, "");
517 system_values.thread_id[1] = LLVMBuildExtractValue(builder, call_context, LP_NIR_CALL_CONTEXT_THREAD_ID_1, "");
518 system_values.thread_id[2] = LLVMBuildExtractValue(builder, call_context, LP_NIR_CALL_CONTEXT_THREAD_ID_2, "");
519 system_values.block_id[0] = LLVMBuildExtractValue(builder, call_context, LP_NIR_CALL_CONTEXT_BLOCK_ID_0, "");
520 system_values.block_id[1] = LLVMBuildExtractValue(builder, call_context, LP_NIR_CALL_CONTEXT_BLOCK_ID_1, "");
521 system_values.block_id[2] = LLVMBuildExtractValue(builder, call_context, LP_NIR_CALL_CONTEXT_BLOCK_ID_2, "");
522 system_values.grid_size[0] = LLVMBuildExtractValue(builder, call_context, LP_NIR_CALL_CONTEXT_GRID_SIZE_0, "");
523 system_values.grid_size[1] = LLVMBuildExtractValue(builder, call_context, LP_NIR_CALL_CONTEXT_GRID_SIZE_1, "");
524 system_values.grid_size[2] = LLVMBuildExtractValue(builder, call_context, LP_NIR_CALL_CONTEXT_GRID_SIZE_2, "");
525 system_values.block_size[0] = LLVMBuildExtractValue(builder, call_context, LP_NIR_CALL_CONTEXT_BLOCK_SIZE_0, "");
526 system_values.block_size[1] = LLVMBuildExtractValue(builder, call_context, LP_NIR_CALL_CONTEXT_BLOCK_SIZE_1, "");
527 system_values.block_size[2] = LLVMBuildExtractValue(builder, call_context, LP_NIR_CALL_CONTEXT_BLOCK_SIZE_2, "");
528
529 params.system_values = &system_values;
530
531 params.consts_ptr = lp_jit_resources_constants(gallivm,
532 variant->jit_resources_type,
533 params.resources_ptr);
534 params.ssbo_ptr = lp_jit_resources_ssbos(gallivm,
535 variant->jit_resources_type,
536 params.resources_ptr);
537 lp_build_nir_soa_func(gallivm, shader->base.ir.nir,
538 func->impl,
539 ¶ms,
540 NULL);
541
542 lp_build_mask_end(&mask);
543
544 LLVMBuildRetVoid(builder);
545 gallivm_verify_function(gallivm, lfunc);
546 }
547 }
548
549 block = LLVMAppendBasicBlockInContext(gallivm->context, function, "entry");
550 builder = gallivm->builder;
551 assert(builder);
552 LLVMPositionBuilderAtEnd(builder, block);
553 sampler = lp_llvm_sampler_soa_create(lp_cs_variant_key_samplers(key),
554 MAX2(key->nr_samplers,
555 key->nr_sampler_views));
556 image = lp_bld_llvm_image_soa_create(lp_cs_variant_key_images(key), key->nr_images);
557
558 if (is_mesh) {
559 LLVMTypeRef output_type = create_mesh_jit_output_type_deref(gallivm);
560 output_array = lp_build_array_alloca(gallivm, output_type, lp_build_const_int32(gallivm, align(MAX2(nir->info.mesh.max_primitives_out, nir->info.mesh.max_vertices_out), 8)), "outputs");
561 }
562
563 struct lp_build_loop_state loop_state[2];
564
565 LLVMValueRef vec_length = lp_build_const_int32(gallivm, cs_type.length);
566
567 LLVMValueRef invocation_count = LLVMBuildMul(gallivm->builder, block_x_size_arg, block_y_size_arg, "");
568 invocation_count = LLVMBuildMul(gallivm->builder, invocation_count, block_z_size_arg, "");
569
570 LLVMValueRef partials = LLVMBuildURem(gallivm->builder, invocation_count, vec_length, "");
571
572 LLVMValueRef num_subgroup_loop = LLVMBuildAdd(gallivm->builder, invocation_count, lp_build_const_int32(gallivm, cs_type.length - 1), "");
573 num_subgroup_loop = LLVMBuildUDiv(gallivm->builder, num_subgroup_loop, vec_length, "");
574
575 /* build a ptr in memory to store all the frames in later. */
576 LLVMTypeRef hdl_ptr_type = LLVMPointerType(LLVMInt8TypeInContext(gallivm->context), 0);
577 LLVMValueRef coro_mem = LLVMBuildAlloca(gallivm->builder, hdl_ptr_type, "coro_mem");
578 LLVMBuildStore(builder, LLVMConstNull(hdl_ptr_type), coro_mem);
579
580 LLVMValueRef coro_hdls = LLVMBuildArrayAlloca(gallivm->builder, hdl_ptr_type, num_subgroup_loop, "coro_hdls");
581
582 unsigned end_coroutine = INT_MAX;
583
584 /*
585 * This is the main coroutine execution loop. It iterates over the dimensions
586 * and calls the coroutine main entrypoint on the first pass, but in subsequent
587 * passes it checks if the coroutine has completed and resumes it if not.
588 */
589 lp_build_loop_begin(&loop_state[1], gallivm,
590 lp_build_const_int32(gallivm, 0)); /* coroutine reentry loop */
591 lp_build_loop_begin(&loop_state[0], gallivm,
592 lp_build_const_int32(gallivm, 0)); /* subgroup loop */
593 {
594 LLVMValueRef args[CS_ARG_MAX];
595 args[CS_ARG_CONTEXT] = context_ptr;
596 args[CS_ARG_RESOURCES] = resources_ptr;
597 args[CS_ARG_BLOCK_X_SIZE] = LLVMGetUndef(int32_type);
598 args[CS_ARG_BLOCK_Y_SIZE] = LLVMGetUndef(int32_type);
599 args[CS_ARG_BLOCK_Z_SIZE] = LLVMGetUndef(int32_type);
600 args[CS_ARG_GRID_X] = grid_x_arg;
601 args[CS_ARG_GRID_Y] = grid_y_arg;
602 args[CS_ARG_GRID_Z] = grid_z_arg;
603 args[CS_ARG_GRID_SIZE_X] = grid_size_x_arg;
604 args[CS_ARG_GRID_SIZE_Y] = grid_size_y_arg;
605 args[CS_ARG_GRID_SIZE_Z] = grid_size_z_arg;
606 args[CS_ARG_WORK_DIM] = work_dim_arg;
607 args[CS_ARG_DRAW_ID] = draw_id_arg;
608 args[CS_ARG_VERTEX_DATA] = io_ptr;
609 args[CS_ARG_PER_THREAD_DATA] = thread_data_ptr;
610 args[CS_ARG_CORO_SUBGROUP_COUNT] = num_subgroup_loop;
611 args[CS_ARG_CORO_PARTIALS] = partials;
612 args[CS_ARG_CORO_BLOCK_X_SIZE] = block_x_size_arg;
613 args[CS_ARG_CORO_BLOCK_Y_SIZE] = block_y_size_arg;
614 args[CS_ARG_CORO_BLOCK_Z_SIZE] = block_z_size_arg;
615
616 args[CS_ARG_CORO_IDX] = loop_state[0].counter;
617
618 args[CS_ARG_CORO_MEM] = coro_mem;
619
620 if (is_mesh)
621 args[CS_ARG_CORO_OUTPUTS] = output_array;
622
623 LLVMValueRef coro_entry = LLVMBuildGEP2(gallivm->builder, hdl_ptr_type, coro_hdls, &loop_state[0].counter, 1, "");
624
625 LLVMValueRef coro_hdl = LLVMBuildLoad2(gallivm->builder, hdl_ptr_type, coro_entry, "coro_hdl");
626
627 struct lp_build_if_state ifstate;
628 LLVMValueRef cmp = LLVMBuildICmp(gallivm->builder, LLVMIntEQ, loop_state[1].counter,
629 lp_build_const_int32(gallivm, 0), "");
630 /* first time here - call the coroutine function entry point */
631 lp_build_if(&ifstate, gallivm, cmp);
632 LLVMValueRef coro_ret = LLVMBuildCall2(gallivm->builder, coro_func_type, coro, args, CS_ARG_MAX - !is_mesh, "");
633 LLVMBuildStore(gallivm->builder, coro_ret, coro_entry);
634 lp_build_else(&ifstate);
635 /* subsequent calls for this invocation - check if done. */
636 LLVMValueRef coro_done = lp_build_coro_done(gallivm, coro_hdl);
637 struct lp_build_if_state ifstate2;
638 lp_build_if(&ifstate2, gallivm, coro_done);
639 /* if done destroy and force loop exit */
640 lp_build_coro_destroy(gallivm, coro_hdl);
641 lp_build_loop_force_set_counter(&loop_state[1], lp_build_const_int32(gallivm, end_coroutine - 1));
642 lp_build_else(&ifstate2);
643 /* otherwise resume the coroutine */
644 lp_build_coro_resume(gallivm, coro_hdl);
645 lp_build_endif(&ifstate2);
646 lp_build_endif(&ifstate);
647 lp_build_loop_force_reload_counter(&loop_state[1]);
648 }
649 lp_build_loop_end_cond(&loop_state[0],
650 num_subgroup_loop,
651 NULL, LLVMIntUGE);
652 lp_build_loop_end_cond(&loop_state[1],
653 lp_build_const_int32(gallivm, end_coroutine),
654 NULL, LLVMIntEQ);
655
656 LLVMValueRef coro_mem_ptr = LLVMBuildLoad2(builder, hdl_ptr_type, coro_mem, "");
657 LLVMTypeRef mem_ptr_type = LLVMPointerType(LLVMInt8TypeInContext(gallivm->context), 0);
658 LLVMTypeRef free_type = LLVMFunctionType(LLVMVoidTypeInContext(gallivm->context), &mem_ptr_type, 1, 0);
659 LLVMBuildCall2(gallivm->builder, free_type, gallivm->coro_free_hook, &coro_mem_ptr, 1, "");
660
661 LLVMBuildRetVoid(builder);
662
663 /* This is stage (b) - generate the compute shader code inside the coroutine. */
664 context_ptr = LLVMGetParam(coro, CS_ARG_CONTEXT);
665 resources_ptr = LLVMGetParam(coro, CS_ARG_RESOURCES);
666 grid_x_arg = LLVMGetParam(coro, CS_ARG_GRID_X);
667 grid_y_arg = LLVMGetParam(coro, CS_ARG_GRID_Y);
668 grid_z_arg = LLVMGetParam(coro, CS_ARG_GRID_Z);
669 grid_size_x_arg = LLVMGetParam(coro, CS_ARG_GRID_SIZE_X);
670 grid_size_y_arg = LLVMGetParam(coro, CS_ARG_GRID_SIZE_Y);
671 grid_size_z_arg = LLVMGetParam(coro, CS_ARG_GRID_SIZE_Z);
672 work_dim_arg = LLVMGetParam(coro, CS_ARG_WORK_DIM);
673 draw_id_arg = LLVMGetParam(coro, CS_ARG_DRAW_ID);
674 io_ptr = LLVMGetParam(coro, CS_ARG_VERTEX_DATA);
675 thread_data_ptr = LLVMGetParam(coro, CS_ARG_PER_THREAD_DATA);
676 num_subgroup_loop = LLVMGetParam(coro, CS_ARG_CORO_SUBGROUP_COUNT);
677 partials = LLVMGetParam(coro, CS_ARG_CORO_PARTIALS);
678 block_x_size_arg = LLVMGetParam(coro, CS_ARG_CORO_BLOCK_X_SIZE);
679 block_y_size_arg = LLVMGetParam(coro, CS_ARG_CORO_BLOCK_Y_SIZE);
680 block_z_size_arg = LLVMGetParam(coro, CS_ARG_CORO_BLOCK_Z_SIZE);
681 LLVMValueRef subgroup_id = LLVMGetParam(coro, CS_ARG_CORO_IDX);
682 coro_mem = LLVMGetParam(coro, CS_ARG_CORO_MEM);
683 if (is_mesh)
684 output_array = LLVMGetParam(coro, CS_ARG_CORO_OUTPUTS);
685 block = LLVMAppendBasicBlockInContext(gallivm->context, coro, "entry");
686 LLVMPositionBuilderAtEnd(builder, block);
687 {
688 LLVMValueRef consts_ptr;
689 LLVMValueRef ssbo_ptr;
690 LLVMValueRef shared_ptr;
691 LLVMValueRef payload_ptr;
692 LLVMValueRef kernel_args_ptr;
693 struct lp_build_mask_context mask;
694 struct lp_bld_tgsi_system_values system_values;
695
696 memset(&system_values, 0, sizeof(system_values));
697 consts_ptr = lp_jit_resources_constants(gallivm, variant->jit_resources_type, resources_ptr);
698 ssbo_ptr = lp_jit_resources_ssbos(gallivm, variant->jit_resources_type, resources_ptr);
699 kernel_args_ptr = lp_jit_cs_context_kernel_args(gallivm,
700 variant->jit_cs_context_type,
701 context_ptr);
702
703 shared_ptr = lp_jit_cs_thread_data_shared(gallivm,
704 variant->jit_cs_thread_data_type,
705 thread_data_ptr);
706 payload_ptr = lp_jit_cs_thread_data_payload(gallivm,
707 variant->jit_cs_thread_data_type,
708 thread_data_ptr);
709
710 /* these are coroutine entrypoint necessities */
711 LLVMValueRef coro_id = lp_build_coro_id(gallivm);
712 LLVMValueRef coro_entry = lp_build_coro_alloc_mem_array(gallivm, coro_mem, subgroup_id, num_subgroup_loop);
713 LLVMTypeRef mem_ptr_type = LLVMInt8TypeInContext(gallivm->context);
714 LLVMValueRef alloced_ptr = LLVMBuildLoad2(gallivm->builder, hdl_ptr_type, coro_mem, "");
715 alloced_ptr = LLVMBuildGEP2(gallivm->builder, mem_ptr_type, alloced_ptr, &coro_entry, 1, "");
716 LLVMValueRef coro_hdl = lp_build_coro_begin(gallivm, coro_id, alloced_ptr);
717 LLVMValueRef has_partials = LLVMBuildICmp(gallivm->builder, LLVMIntNE, partials, lp_build_const_int32(gallivm, 0), "");
718
719 struct lp_build_context bld;
720 lp_build_context_init(&bld, gallivm, lp_uint_type(cs_type));
721
722 LLVMValueRef base_val = LLVMBuildMul(gallivm->builder, subgroup_id, vec_length, "");
723 LLVMValueRef invocation_indices[LP_MAX_VECTOR_LENGTH];
724 for (i = 0; i < cs_type.length; i++)
725 invocation_indices[i] = LLVMBuildAdd(gallivm->builder, base_val, lp_build_const_int32(gallivm, i), "");
726 LLVMValueRef invocation_index = lp_build_gather_values(gallivm, invocation_indices, cs_type.length);
727
728 LLVMValueRef block_x_size_vec = lp_build_broadcast_scalar(&bld, block_x_size_arg);
729 LLVMValueRef block_y_size_vec = lp_build_broadcast_scalar(&bld, block_y_size_arg);
730
731 system_values.thread_id[0] = LLVMBuildURem(gallivm->builder, invocation_index, block_x_size_vec, "");
732 system_values.thread_id[1] = LLVMBuildUDiv(gallivm->builder, invocation_index, block_x_size_vec, "");
733 system_values.thread_id[1] = LLVMBuildURem(gallivm->builder, system_values.thread_id[1], block_y_size_vec, "");
734 system_values.thread_id[2] = LLVMBuildUDiv(gallivm->builder, invocation_index, block_x_size_vec, "");
735 system_values.thread_id[2] = LLVMBuildUDiv(gallivm->builder, system_values.thread_id[2], block_y_size_vec, "");
736
737 system_values.block_id[0] = grid_x_arg;
738 system_values.block_id[1] = grid_y_arg;
739 system_values.block_id[2] = grid_z_arg;
740
741 system_values.grid_size[0] = grid_size_x_arg;
742 system_values.grid_size[1] = grid_size_y_arg;
743 system_values.grid_size[2] = grid_size_z_arg;
744
745 system_values.work_dim = work_dim_arg;
746 system_values.draw_id = draw_id_arg;
747
748 system_values.subgroup_id = subgroup_id;
749 system_values.num_subgroups = num_subgroup_loop;
750
751 system_values.block_size[0] = block_x_size_arg;
752 system_values.block_size[1] = block_y_size_arg;
753 system_values.block_size[2] = block_z_size_arg;
754
755 LLVMValueRef last_loop = LLVMBuildICmp(gallivm->builder, LLVMIntEQ, subgroup_id, LLVMBuildSub(gallivm->builder, num_subgroup_loop, lp_build_const_int32(gallivm, 1), ""), "");
756 LLVMValueRef use_partial_mask = LLVMBuildAnd(gallivm->builder, last_loop, has_partials, "");
757 struct lp_build_if_state if_state;
758 LLVMTypeRef mask_type = LLVMVectorType(int32_type, cs_type.length);
759 LLVMValueRef mask_val = lp_build_alloca(gallivm, mask_type, "mask");
760 LLVMValueRef full_mask_val = lp_build_const_int_vec(gallivm, cs_type, ~0);
761 LLVMBuildStore(gallivm->builder, full_mask_val, mask_val);
762
763 lp_build_if(&if_state, gallivm, use_partial_mask);
764 struct lp_build_loop_state mask_loop_state;
765 lp_build_loop_begin(&mask_loop_state, gallivm, partials);
766 LLVMValueRef tmask_val = LLVMBuildLoad2(gallivm->builder, mask_type, mask_val, "");
767 tmask_val = LLVMBuildInsertElement(gallivm->builder, tmask_val, lp_build_const_int32(gallivm, 0), mask_loop_state.counter, "");
768 LLVMBuildStore(gallivm->builder, tmask_val, mask_val);
769 lp_build_loop_end_cond(&mask_loop_state, vec_length, NULL, LLVMIntUGE);
770 lp_build_endif(&if_state);
771
772 mask_val = LLVMBuildLoad2(gallivm->builder, mask_type, mask_val, "");
773 lp_build_mask_begin(&mask, gallivm, cs_type, mask_val);
774
775 struct lp_build_coro_suspend_info coro_info;
776
777 LLVMBasicBlockRef sus_block = LLVMAppendBasicBlockInContext(gallivm->context, coro, "suspend");
778 LLVMBasicBlockRef clean_block = LLVMAppendBasicBlockInContext(gallivm->context, coro, "cleanup");
779
780 coro_info.suspend = sus_block;
781 coro_info.cleanup = clean_block;
782
783 if (is_mesh) {
784 LLVMValueRef vertex_count = lp_build_alloca(gallivm, LLVMInt32TypeInContext(gallivm->context), "vertex_count");
785 LLVMValueRef primitive_count = lp_build_alloca(gallivm, LLVMInt32TypeInContext(gallivm->context), "prim_count");
786 mesh_iface.base.emit_store_output = lp_mesh_llvm_emit_store_output;
787 mesh_iface.base.emit_vertex_and_primitive_count = lp_mesh_emit_vertex_and_primitive_count;
788 mesh_iface.vertex_count = vertex_count;
789 mesh_iface.prim_count = primitive_count;
790 mesh_iface.outputs = output_array;
791 }
792
793 struct lp_build_tgsi_params params;
794 memset(¶ms, 0, sizeof(params));
795
796 params.type = cs_type;
797 params.mask = &mask;
798 params.consts_ptr = consts_ptr;
799 params.system_values = &system_values;
800 params.context_type = variant->jit_cs_context_type;
801 params.context_ptr = context_ptr;
802 params.resources_type = variant->jit_resources_type;
803 params.resources_ptr = resources_ptr;
804 params.sampler = sampler;
805 params.ssbo_ptr = ssbo_ptr;
806 params.image = image;
807 params.shared_ptr = shared_ptr;
808 params.payload_ptr = payload_ptr;
809 params.coro = &coro_info;
810 params.kernel_args = kernel_args_ptr;
811 params.aniso_filter_table = lp_jit_resources_aniso_filter_table(gallivm,
812 variant->jit_resources_type,
813 resources_ptr);
814 params.mesh_iface = &mesh_iface.base;
815
816 params.current_func = NULL;
817 params.fns = fns;
818 lp_build_nir_soa_func(gallivm, nir,
819 nir_shader_get_entrypoint(nir),
820 ¶ms, NULL);
821
822 if (is_mesh) {
823 LLVMTypeRef i32t = LLVMInt32TypeInContext(gallivm->context);
824 LLVMValueRef clipmask = lp_build_const_int_vec(gallivm,
825 lp_int_type(cs_type), 0);
826
827 struct lp_build_if_state iter0state;
828 LLVMValueRef is_iter0 = LLVMBuildICmp(gallivm->builder, LLVMIntEQ, subgroup_id,
829 lp_build_const_int32(gallivm, 0), "");
830 LLVMValueRef vertex_count = LLVMBuildLoad2(gallivm->builder, i32t, mesh_iface.vertex_count, "");
831 LLVMValueRef prim_count = LLVMBuildLoad2(gallivm->builder, i32t, mesh_iface.prim_count, "");
832
833 LLVMValueRef vert_count_ptr, prim_count_ptr;
834 LLVMValueRef indices = lp_build_const_int32(gallivm, 1);
835 vert_count_ptr = LLVMBuildGEP2(gallivm->builder, i32t, io_ptr, &indices, 1, "");
836 indices = lp_build_const_int32(gallivm, 2);
837 prim_count_ptr = LLVMBuildGEP2(gallivm->builder, i32t, io_ptr, &indices, 1, "");
838
839 lp_build_if(&iter0state, gallivm, is_iter0);
840 LLVMBuildStore(gallivm->builder, vertex_count, vert_count_ptr);
841 LLVMBuildStore(gallivm->builder, prim_count, prim_count_ptr);
842 lp_build_endif(&iter0state);
843
844 LLVMBasicBlockRef resume = lp_build_insert_new_block(gallivm, "resume");
845
846 lp_build_coro_suspend_switch(gallivm, params.coro, resume, false);
847 LLVMPositionBuilderAtEnd(gallivm->builder, resume);
848
849 vertex_count = LLVMBuildLoad2(gallivm->builder, i32t, vert_count_ptr, "");
850 prim_count = LLVMBuildLoad2(gallivm->builder, i32t, prim_count_ptr, "");
851
852 int per_prim_count = util_bitcount64(nir->info.per_primitive_outputs);
853 int out_count = util_bitcount64(nir->info.outputs_written);
854 int per_vert_count = out_count - per_prim_count;
855 int vsize = (sizeof(struct vertex_header) + per_vert_count * 4 * sizeof(float)) * 8;
856 int psize = (per_prim_count * 4 * sizeof(float)) * 8;
857 struct lp_build_loop_state vertex_loop_state;
858
859 lp_build_loop_begin(&vertex_loop_state, gallivm,
860 lp_build_const_int32(gallivm, 0));
861 LLVMValueRef io;
862 io = LLVMBuildPtrToInt(gallivm->builder, io_ptr, LLVMInt64TypeInContext(gallivm->context), "");
863 io = LLVMBuildAdd(builder, io, LLVMBuildZExt(builder, LLVMBuildMul(builder, vertex_loop_state.counter, lp_build_const_int32(gallivm, vsize), ""), LLVMInt64TypeInContext(gallivm->context), ""), "");
864 io = LLVMBuildIntToPtr(gallivm->builder, io, LLVMPointerType(LLVMVoidTypeInContext(gallivm->context), 0), "");
865 mesh_convert_to_aos(gallivm, shader->base.ir.nir, true, variant->jit_vertex_header_type,
866 io, output_array, clipmask,
867 vertex_loop_state.counter, lp_elem_type(cs_type), -1, false);
868 lp_build_loop_end_cond(&vertex_loop_state,
869 vertex_count,
870 NULL, LLVMIntUGE);
871
872 struct lp_build_loop_state prim_loop_state;
873 lp_build_loop_begin(&prim_loop_state, gallivm,
874 lp_build_const_int32(gallivm, 0));
875 io = LLVMBuildPtrToInt(gallivm->builder, io_ptr, LLVMInt64TypeInContext(gallivm->context), "");
876 LLVMValueRef prim_offset = LLVMBuildMul(builder, prim_loop_state.counter, lp_build_const_int32(gallivm, psize), "");
877 prim_offset = LLVMBuildAdd(builder, prim_offset, lp_build_const_int32(gallivm, vsize * (nir->info.mesh.max_vertices_out + 8)), "");
878 io = LLVMBuildAdd(builder, io, LLVMBuildZExt(builder, prim_offset, LLVMInt64TypeInContext(gallivm->context), ""), "");
879 io = LLVMBuildIntToPtr(gallivm->builder, io, LLVMPointerType(LLVMVoidTypeInContext(gallivm->context), 0), "");
880 mesh_convert_to_aos(gallivm, shader->base.ir.nir, false, variant->jit_prim_type,
881 io, output_array, clipmask,
882 prim_loop_state.counter, lp_elem_type(cs_type), -1, false);
883 lp_build_loop_end_cond(&prim_loop_state,
884 prim_count,
885 NULL, LLVMIntUGE);
886 }
887
888 mask_val = lp_build_mask_end(&mask);
889
890 lp_build_coro_suspend_switch(gallivm, &coro_info, NULL, true);
891 LLVMPositionBuilderAtEnd(builder, clean_block);
892
893 LLVMBuildBr(builder, sus_block);
894 LLVMPositionBuilderAtEnd(builder, sus_block);
895
896 lp_build_coro_end(gallivm, coro_hdl);
897 LLVMBuildRet(builder, coro_hdl);
898 }
899
900 lp_bld_llvm_sampler_soa_destroy(sampler);
901 lp_bld_llvm_image_soa_destroy(image);
902 _mesa_hash_table_destroy(fns, NULL);
903
904 gallivm_verify_function(gallivm, coro);
905 gallivm_verify_function(gallivm, function);
906 }
907
908
909 static void *
llvmpipe_create_compute_state(struct pipe_context * pipe,const struct pipe_compute_state * templ)910 llvmpipe_create_compute_state(struct pipe_context *pipe,
911 const struct pipe_compute_state *templ)
912 {
913 struct lp_compute_shader *shader = CALLOC_STRUCT(lp_compute_shader);
914 struct nir_shader *nir = NULL;
915 if (!shader)
916 return NULL;
917
918 shader->no = cs_no++;
919
920 shader->base.type = PIPE_SHADER_IR_NIR;
921
922 if (templ->ir_type == PIPE_SHADER_IR_TGSI) {
923 shader->base.ir.nir = tgsi_to_nir(templ->prog, pipe->screen, false);
924 } else if (templ->ir_type == PIPE_SHADER_IR_NIR_SERIALIZED) {
925 struct blob_reader reader;
926 const struct pipe_binary_program_header *hdr = templ->prog;
927
928 blob_reader_init(&reader, hdr->blob, hdr->num_bytes);
929 shader->base.ir.nir = nir_deserialize(NULL, pipe->screen->get_compiler_options(pipe->screen, PIPE_SHADER_IR_NIR, PIPE_SHADER_COMPUTE), &reader);
930
931 pipe->screen->finalize_nir(pipe->screen, shader->base.ir.nir);
932 } else if (templ->ir_type == PIPE_SHADER_IR_NIR) {
933 shader->base.ir.nir = (struct nir_shader *)templ->prog;
934 }
935
936 nir = (struct nir_shader *)shader->base.ir.nir;
937 shader->req_local_mem += nir->info.shared_size;
938 shader->zero_initialize_shared_memory = nir->info.zero_initialize_shared_memory;
939
940 llvmpipe_register_shader(pipe, &shader->base);
941
942 list_inithead(&shader->variants.list);
943
944 int nr_samplers = BITSET_LAST_BIT(nir->info.samplers_used);
945 int nr_sampler_views = BITSET_LAST_BIT(nir->info.textures_used);
946 int nr_images = BITSET_LAST_BIT(nir->info.images_used);
947 shader->variant_key_size = lp_cs_variant_key_size(MAX2(nr_samplers, nr_sampler_views), nr_images);
948
949 return shader;
950 }
951
952
953 static void
llvmpipe_bind_compute_state(struct pipe_context * pipe,void * cs)954 llvmpipe_bind_compute_state(struct pipe_context *pipe,
955 void *cs)
956 {
957 struct llvmpipe_context *llvmpipe = llvmpipe_context(pipe);
958
959 if (llvmpipe->cs == cs)
960 return;
961
962 llvmpipe->cs = (struct lp_compute_shader *)cs;
963 llvmpipe->cs_dirty |= LP_CSNEW_CS;
964 }
965
966 static void
llvmpipe_get_compute_state_info(struct pipe_context * pipe,void * cs,struct pipe_compute_state_object_info * info)967 llvmpipe_get_compute_state_info(struct pipe_context *pipe, void *cs,
968 struct pipe_compute_state_object_info *info)
969 {
970 struct lp_compute_shader* shader = cs;
971 struct nir_shader* nir = shader->base.ir.nir;
972
973 info->max_threads = 1024;
974 info->simd_sizes = lp_native_vector_width / 32;
975 info->preferred_simd_size = info->simd_sizes;
976 // TODO: this is a bad estimate, but not much we can do without actually compiling the shaders
977 info->private_memory = nir->scratch_size;
978 }
979
980
981 /**
982 * Remove shader variant from two lists: the shader's variant list
983 * and the context's variant list.
984 */
985 static void
llvmpipe_remove_cs_shader_variant(struct llvmpipe_context * lp,struct lp_compute_shader_variant * variant)986 llvmpipe_remove_cs_shader_variant(struct llvmpipe_context *lp,
987 struct lp_compute_shader_variant *variant)
988 {
989 if ((LP_DEBUG & DEBUG_CS) || (gallivm_debug & GALLIVM_DEBUG_IR)) {
990 debug_printf("llvmpipe: del cs #%u var %u v created %u v cached %u "
991 "v total cached %u inst %u total inst %u\n",
992 variant->shader->no, variant->no,
993 variant->shader->variants_created,
994 variant->shader->variants_cached,
995 lp->nr_cs_variants, variant->nr_instrs, lp->nr_cs_instrs);
996 }
997
998 gallivm_destroy(variant->gallivm);
999
1000 /* remove from shader's list */
1001 list_del(&variant->list_item_local.list);
1002 variant->shader->variants_cached--;
1003
1004 /* remove from context's list */
1005 list_del(&variant->list_item_global.list);
1006 lp->nr_cs_variants--;
1007 lp->nr_cs_instrs -= variant->nr_instrs;
1008
1009 FREE(variant);
1010 }
1011
1012
1013 static void
llvmpipe_delete_compute_state(struct pipe_context * pipe,void * cs)1014 llvmpipe_delete_compute_state(struct pipe_context *pipe,
1015 void *cs)
1016 {
1017 struct llvmpipe_context *llvmpipe = llvmpipe_context(pipe);
1018 struct lp_compute_shader *shader = cs;
1019 struct lp_cs_variant_list_item *li, *next;
1020
1021 if (llvmpipe->cs == cs)
1022 llvmpipe->cs = NULL;
1023 for (unsigned i = 0; i < shader->max_global_buffers; i++)
1024 pipe_resource_reference(&shader->global_buffers[i], NULL);
1025 FREE(shader->global_buffers);
1026
1027 /* Delete all the variants */
1028 LIST_FOR_EACH_ENTRY_SAFE(li, next, &shader->variants.list, list) {
1029 llvmpipe_remove_cs_shader_variant(llvmpipe, li->base);
1030 }
1031 ralloc_free(shader->base.ir.nir);
1032 FREE(shader);
1033 }
1034
1035
1036 static struct lp_compute_shader_variant_key *
make_variant_key(struct llvmpipe_context * lp,struct lp_compute_shader * shader,enum pipe_shader_type sh_type,char * store)1037 make_variant_key(struct llvmpipe_context *lp,
1038 struct lp_compute_shader *shader,
1039 enum pipe_shader_type sh_type,
1040 char *store)
1041 {
1042 struct lp_compute_shader_variant_key *key =
1043 (struct lp_compute_shader_variant_key *)store;
1044 memset(key, 0, sizeof(*key));
1045
1046 struct nir_shader *nir = (struct nir_shader *)shader->base.ir.nir;
1047 /* This value will be the same for all the variants of a given shader:
1048 */
1049 key->nr_samplers = BITSET_LAST_BIT(nir->info.samplers_used);
1050 key->nr_sampler_views = BITSET_LAST_BIT(nir->info.textures_used);
1051 struct lp_sampler_static_state *cs_sampler;
1052
1053 cs_sampler = lp_cs_variant_key_samplers(key);
1054
1055 memset(cs_sampler, 0, MAX2(key->nr_samplers, key->nr_sampler_views) * sizeof *cs_sampler);
1056 for (unsigned i = 0; i < key->nr_samplers; ++i) {
1057 if (BITSET_TEST(nir->info.samplers_used, i)) {
1058 lp_sampler_static_sampler_state(&cs_sampler[i].sampler_state,
1059 lp->samplers[sh_type][i]);
1060 }
1061 }
1062
1063 /*
1064 * XXX If TGSI_FILE_SAMPLER_VIEW exists assume all texture opcodes
1065 * are dx10-style? Can't really have mixed opcodes, at least not
1066 * if we want to skip the holes here (without rescanning tgsi).
1067 */
1068 if (!BITSET_IS_EMPTY(nir->info.textures_used)) {
1069 for (unsigned i = 0; i < key->nr_sampler_views; ++i) {
1070 /*
1071 * Note sview may exceed what's representable by file_mask.
1072 * This will still work, the only downside is that not actually
1073 * used views may be included in the shader key.
1074 */
1075 if (BITSET_TEST(nir->info.textures_used, i)) {
1076 lp_sampler_static_texture_state(&cs_sampler[i].texture_state,
1077 lp->sampler_views[sh_type][i]);
1078 }
1079 }
1080 } else {
1081 key->nr_sampler_views = key->nr_samplers;
1082 for (unsigned i = 0; i < key->nr_sampler_views; ++i) {
1083 if (BITSET_TEST(nir->info.samplers_used, i)) {
1084 lp_sampler_static_texture_state(&cs_sampler[i].texture_state,
1085 lp->sampler_views[sh_type][i]);
1086 }
1087 }
1088 }
1089
1090 struct lp_image_static_state *lp_image;
1091 lp_image = lp_cs_variant_key_images(key);
1092 key->nr_images = BITSET_LAST_BIT(nir->info.images_used);
1093
1094 if (key->nr_images)
1095 memset(lp_image, 0,
1096 key->nr_images * sizeof *lp_image);
1097 for (unsigned i = 0; i < key->nr_images; ++i) {
1098 if (BITSET_TEST(nir->info.images_used, i)) {
1099 lp_sampler_static_texture_state_image(&lp_image[i].image_state,
1100 &lp->images[sh_type][i]);
1101 }
1102 }
1103 return key;
1104 }
1105
1106
1107 static void
dump_cs_variant_key(const struct lp_compute_shader_variant_key * key)1108 dump_cs_variant_key(const struct lp_compute_shader_variant_key *key)
1109 {
1110 int i;
1111 debug_printf("cs variant %p:\n", (void *) key);
1112
1113 for (i = 0; i < key->nr_samplers; ++i) {
1114 const struct lp_sampler_static_state *samplers = lp_cs_variant_key_samplers(key);
1115 const struct lp_static_sampler_state *sampler = &samplers[i].sampler_state;
1116 debug_printf("sampler[%u] = \n", i);
1117 debug_printf(" .wrap = %s %s %s\n",
1118 util_str_tex_wrap(sampler->wrap_s, true),
1119 util_str_tex_wrap(sampler->wrap_t, true),
1120 util_str_tex_wrap(sampler->wrap_r, true));
1121 debug_printf(" .min_img_filter = %s\n",
1122 util_str_tex_filter(sampler->min_img_filter, true));
1123 debug_printf(" .min_mip_filter = %s\n",
1124 util_str_tex_mipfilter(sampler->min_mip_filter, true));
1125 debug_printf(" .mag_img_filter = %s\n",
1126 util_str_tex_filter(sampler->mag_img_filter, true));
1127 if (sampler->compare_mode != PIPE_TEX_COMPARE_NONE)
1128 debug_printf(" .compare_func = %s\n", util_str_func(sampler->compare_func, true));
1129 debug_printf(" .normalized_coords = %u\n", sampler->normalized_coords);
1130 debug_printf(" .min_max_lod_equal = %u\n", sampler->min_max_lod_equal);
1131 debug_printf(" .lod_bias_non_zero = %u\n", sampler->lod_bias_non_zero);
1132 debug_printf(" .apply_min_lod = %u\n", sampler->apply_min_lod);
1133 debug_printf(" .apply_max_lod = %u\n", sampler->apply_max_lod);
1134 debug_printf(" .aniso = %u\n", sampler->aniso);
1135 }
1136 for (i = 0; i < key->nr_sampler_views; ++i) {
1137 const struct lp_sampler_static_state *samplers = lp_cs_variant_key_samplers(key);
1138 const struct lp_static_texture_state *texture = &samplers[i].texture_state;
1139 debug_printf("texture[%u] = \n", i);
1140 debug_printf(" .format = %s\n",
1141 util_format_name(texture->format));
1142 debug_printf(" .target = %s\n",
1143 util_str_tex_target(texture->target, true));
1144 debug_printf(" .level_zero_only = %u\n",
1145 texture->level_zero_only);
1146 debug_printf(" .pot = %u %u %u\n",
1147 texture->pot_width,
1148 texture->pot_height,
1149 texture->pot_depth);
1150 }
1151 struct lp_image_static_state *images = lp_cs_variant_key_images(key);
1152 for (i = 0; i < key->nr_images; ++i) {
1153 const struct lp_static_texture_state *image = &images[i].image_state;
1154 debug_printf("image[%u] = \n", i);
1155 debug_printf(" .format = %s\n",
1156 util_format_name(image->format));
1157 debug_printf(" .target = %s\n",
1158 util_str_tex_target(image->target, true));
1159 debug_printf(" .level_zero_only = %u\n",
1160 image->level_zero_only);
1161 debug_printf(" .pot = %u %u %u\n",
1162 image->pot_width,
1163 image->pot_height,
1164 image->pot_depth);
1165 }
1166 }
1167
1168
1169 static void
lp_debug_cs_variant(const struct lp_compute_shader_variant * variant)1170 lp_debug_cs_variant(const struct lp_compute_shader_variant *variant)
1171 {
1172 debug_printf("llvmpipe: Compute shader #%u variant #%u:\n",
1173 variant->shader->no, variant->no);
1174 nir_print_shader(variant->shader->base.ir.nir, stderr);
1175 dump_cs_variant_key(&variant->key);
1176 debug_printf("\n");
1177 }
1178
1179
1180 static void
lp_cs_get_ir_cache_key(struct lp_compute_shader_variant * variant,unsigned char ir_sha1_cache_key[20])1181 lp_cs_get_ir_cache_key(struct lp_compute_shader_variant *variant,
1182 unsigned char ir_sha1_cache_key[20])
1183 {
1184 struct blob blob = { 0 };
1185 unsigned ir_size;
1186 void *ir_binary;
1187
1188 blob_init(&blob);
1189 nir_serialize(&blob, variant->shader->base.ir.nir, true);
1190 ir_binary = blob.data;
1191 ir_size = blob.size;
1192
1193 struct mesa_sha1 ctx;
1194 _mesa_sha1_init(&ctx);
1195 _mesa_sha1_update(&ctx, &variant->key, variant->shader->variant_key_size);
1196 _mesa_sha1_update(&ctx, ir_binary, ir_size);
1197 _mesa_sha1_final(&ctx, ir_sha1_cache_key);
1198
1199 blob_finish(&blob);
1200 }
1201
1202
1203 static struct lp_compute_shader_variant *
generate_variant(struct llvmpipe_context * lp,struct lp_compute_shader * shader,enum pipe_shader_type sh_type,const struct lp_compute_shader_variant_key * key)1204 generate_variant(struct llvmpipe_context *lp,
1205 struct lp_compute_shader *shader,
1206 enum pipe_shader_type sh_type,
1207 const struct lp_compute_shader_variant_key *key)
1208 {
1209 struct llvmpipe_screen *screen = llvmpipe_screen(lp->pipe.screen);
1210
1211 struct lp_compute_shader_variant *variant =
1212 MALLOC(sizeof *variant + shader->variant_key_size - sizeof variant->key);
1213 if (!variant)
1214 return NULL;
1215
1216 memset(variant, 0, sizeof(*variant));
1217
1218 char module_name[64];
1219 const char *shname = sh_type == PIPE_SHADER_MESH ? "ms" :
1220 (sh_type == PIPE_SHADER_TASK ? "ts" : "cs");
1221 snprintf(module_name, sizeof(module_name), "%s%u_variant%u",
1222 shname, shader->no, shader->variants_created);
1223
1224 variant->shader = shader;
1225 memcpy(&variant->key, key, shader->variant_key_size);
1226
1227 unsigned char ir_sha1_cache_key[20];
1228 struct lp_cached_code cached = { 0 };
1229 bool needs_caching = false;
1230
1231 lp_cs_get_ir_cache_key(variant, ir_sha1_cache_key);
1232
1233 lp_disk_cache_find_shader(screen, &cached, ir_sha1_cache_key);
1234 if (!cached.data_size)
1235 needs_caching = true;
1236
1237 variant->gallivm = gallivm_create(module_name, lp->context, &cached);
1238 if (!variant->gallivm) {
1239 FREE(variant);
1240 return NULL;
1241 }
1242
1243 variant->list_item_global.base = variant;
1244 variant->list_item_local.base = variant;
1245 variant->no = shader->variants_created++;
1246
1247 if ((LP_DEBUG & DEBUG_CS) || (gallivm_debug & GALLIVM_DEBUG_IR)) {
1248 lp_debug_cs_variant(variant);
1249 }
1250
1251 lp_jit_init_cs_types(variant);
1252
1253 if (sh_type == PIPE_SHADER_MESH) {
1254 struct nir_shader *nir = shader->base.ir.nir;
1255 int per_prim_count = util_bitcount64(nir->info.per_primitive_outputs);
1256 int out_count = util_bitcount64(nir->info.outputs_written);
1257 int per_vert_count = out_count - per_prim_count;
1258 variant->jit_vertex_header_type = lp_build_create_jit_vertex_header_type(variant->gallivm, per_vert_count);
1259 variant->jit_vertex_header_ptr_type = LLVMPointerType(variant->jit_vertex_header_type, 0);
1260 variant->jit_prim_type = LLVMArrayType(LLVMArrayType(LLVMFloatTypeInContext(variant->gallivm->context), 4), per_prim_count);
1261 }
1262
1263 generate_compute(lp, shader, variant);
1264
1265 gallivm_compile_module(variant->gallivm);
1266
1267 variant->nr_instrs += lp_build_count_ir_module(variant->gallivm->module);
1268
1269 variant->jit_function = (lp_jit_cs_func)
1270 gallivm_jit_function(variant->gallivm, variant->function);
1271
1272 if (needs_caching) {
1273 lp_disk_cache_insert_shader(screen, &cached, ir_sha1_cache_key);
1274 }
1275 gallivm_free_ir(variant->gallivm);
1276 return variant;
1277 }
1278
1279
1280 static void
lp_cs_ctx_set_cs_variant(struct lp_cs_context * csctx,struct lp_compute_shader_variant * variant)1281 lp_cs_ctx_set_cs_variant(struct lp_cs_context *csctx,
1282 struct lp_compute_shader_variant *variant)
1283 {
1284 csctx->cs.current.variant = variant;
1285 }
1286
1287
1288 static struct lp_compute_shader_variant *
llvmpipe_update_cs_variant(struct llvmpipe_context * lp,enum pipe_shader_type sh_type,struct lp_compute_shader * shader)1289 llvmpipe_update_cs_variant(struct llvmpipe_context *lp,
1290 enum pipe_shader_type sh_type,
1291 struct lp_compute_shader *shader)
1292 {
1293 char store[LP_CS_MAX_VARIANT_KEY_SIZE];
1294 struct lp_compute_shader_variant_key *key =
1295 make_variant_key(lp, shader, sh_type, store);
1296 struct lp_compute_shader_variant *variant = NULL;
1297 struct lp_cs_variant_list_item *li;
1298
1299 /* Search the variants for one which matches the key */
1300 LIST_FOR_EACH_ENTRY(li, &shader->variants.list, list) {
1301 if (memcmp(&li->base->key, key, shader->variant_key_size) == 0) {
1302 variant = li->base;
1303 break;
1304 }
1305 }
1306
1307 if (variant) {
1308 /* Move this variant to the head of the list to implement LRU
1309 * deletion of shader's when we have too many.
1310 */
1311 list_move_to(&variant->list_item_global.list,
1312 &lp->cs_variants_list.list);
1313 } else {
1314 /* variant not found, create it now */
1315
1316 if (LP_DEBUG & DEBUG_CS) {
1317 debug_printf("%u variants,\t%u instrs,\t%u instrs/variant\n",
1318 lp->nr_cs_variants,
1319 lp->nr_cs_instrs,
1320 lp->nr_cs_variants
1321 ? lp->nr_cs_instrs / lp->nr_cs_variants : 0);
1322 }
1323
1324 /* First, check if we've exceeded the max number of shader variants.
1325 * If so, free 6.25% of them (the least recently used ones).
1326 */
1327 unsigned variants_to_cull = lp->nr_cs_variants >= LP_MAX_SHADER_VARIANTS
1328 ? LP_MAX_SHADER_VARIANTS / 16 : 0;
1329
1330 if (variants_to_cull ||
1331 lp->nr_cs_instrs >= LP_MAX_SHADER_INSTRUCTIONS) {
1332 if (gallivm_debug & GALLIVM_DEBUG_PERF) {
1333 debug_printf("Evicting CS: %u cs variants,\t%u total variants,"
1334 "\t%u instrs,\t%u instrs/variant\n",
1335 shader->variants_cached,
1336 lp->nr_cs_variants, lp->nr_cs_instrs,
1337 lp->nr_cs_instrs / lp->nr_cs_variants);
1338 }
1339
1340 /*
1341 * We need to re-check lp->nr_cs_variants because an arbitrarily large
1342 * number of shader variants (potentially all of them) could be
1343 * pending for destruction on flush.
1344 */
1345 for (unsigned i = 0;
1346 i < variants_to_cull ||
1347 lp->nr_cs_instrs >= LP_MAX_SHADER_INSTRUCTIONS; i++) {
1348 struct lp_cs_variant_list_item *item;
1349 if (list_is_empty(&lp->cs_variants_list.list)) {
1350 break;
1351 }
1352 item = list_last_entry(&lp->cs_variants_list.list,
1353 struct lp_cs_variant_list_item, list);
1354 assert(item);
1355 assert(item->base);
1356 llvmpipe_remove_cs_shader_variant(lp, item->base);
1357 }
1358 }
1359
1360 /*
1361 * Generate the new variant.
1362 */
1363 int64_t t0, t1, dt;
1364 t0 = os_time_get();
1365 variant = generate_variant(lp, shader, sh_type, key);
1366 t1 = os_time_get();
1367 dt = t1 - t0;
1368 LP_COUNT_ADD(llvm_compile_time, dt);
1369 LP_COUNT_ADD(nr_llvm_compiles, 2); /* emit vs. omit in/out test */
1370
1371 /* Put the new variant into the list */
1372 if (variant) {
1373 list_add(&variant->list_item_local.list, &shader->variants.list);
1374 list_add(&variant->list_item_global.list, &lp->cs_variants_list.list);
1375 lp->nr_cs_variants++;
1376 lp->nr_cs_instrs += variant->nr_instrs;
1377 shader->variants_cached++;
1378 }
1379 }
1380 return variant;
1381 }
1382
1383 static void
llvmpipe_update_cs(struct llvmpipe_context * lp)1384 llvmpipe_update_cs(struct llvmpipe_context *lp)
1385 {
1386 struct lp_compute_shader_variant *variant;
1387 variant = llvmpipe_update_cs_variant(lp, PIPE_SHADER_COMPUTE, lp->cs);
1388 /* Bind this variant */
1389 lp_cs_ctx_set_cs_variant(lp->csctx, variant);
1390 }
1391
1392
1393 /**
1394 * Called during state validation when LP_CSNEW_SAMPLER_VIEW is set.
1395 */
1396 static void
lp_csctx_set_sampler_views(struct lp_cs_context * csctx,unsigned num,struct pipe_sampler_view ** views)1397 lp_csctx_set_sampler_views(struct lp_cs_context *csctx,
1398 unsigned num,
1399 struct pipe_sampler_view **views)
1400 {
1401 LP_DBG(DEBUG_SETUP, "%s\n", __func__);
1402
1403 assert(num <= PIPE_MAX_SHADER_SAMPLER_VIEWS);
1404
1405 const unsigned max_tex_num = MAX2(num, csctx->cs.current_tex_num);
1406
1407 for (unsigned i = 0; i < max_tex_num; i++) {
1408 struct pipe_sampler_view *view = i < num ? views[i] : NULL;
1409
1410 /* We are going to overwrite/unref the current texture further below. If
1411 * set, make sure to unmap its resource to avoid leaking previous
1412 * mapping. */
1413 if (csctx->cs.current_tex[i])
1414 llvmpipe_resource_unmap(csctx->cs.current_tex[i], 0, 0);
1415
1416 if (view) {
1417 struct pipe_resource *res = view->texture;
1418 struct lp_jit_texture *jit_tex;
1419 jit_tex = &csctx->cs.current.jit_resources.textures[i];
1420
1421 /* We're referencing the texture's internal data, so save a
1422 * reference to it.
1423 */
1424 pipe_resource_reference(&csctx->cs.current_tex[i], res);
1425
1426 lp_jit_texture_from_pipe(jit_tex, view);
1427 } else {
1428 pipe_resource_reference(&csctx->cs.current_tex[i], NULL);
1429 }
1430 }
1431 csctx->cs.current_tex_num = num;
1432 }
1433
1434
1435 /**
1436 * Called during state validation when LP_NEW_SAMPLER is set.
1437 */
1438 static void
lp_csctx_set_sampler_state(struct lp_cs_context * csctx,unsigned num,struct pipe_sampler_state ** samplers)1439 lp_csctx_set_sampler_state(struct lp_cs_context *csctx,
1440 unsigned num,
1441 struct pipe_sampler_state **samplers)
1442 {
1443 LP_DBG(DEBUG_SETUP, "%s\n", __func__);
1444
1445 assert(num <= PIPE_MAX_SAMPLERS);
1446
1447 for (unsigned i = 0; i < PIPE_MAX_SAMPLERS; i++) {
1448 const struct pipe_sampler_state *sampler = i < num ? samplers[i] : NULL;
1449
1450 if (sampler) {
1451 struct lp_jit_sampler *jit_sam;
1452 jit_sam = &csctx->cs.current.jit_resources.samplers[i];
1453
1454 jit_sam->min_lod = sampler->min_lod;
1455 jit_sam->max_lod = sampler->max_lod;
1456 jit_sam->lod_bias = sampler->lod_bias;
1457 jit_sam->max_aniso = sampler->max_anisotropy;
1458 COPY_4V(jit_sam->border_color, sampler->border_color.f);
1459 }
1460 }
1461 }
1462
1463
1464 static void
lp_csctx_set_cs_constants(struct lp_cs_context * csctx,unsigned num,struct pipe_constant_buffer * buffers)1465 lp_csctx_set_cs_constants(struct lp_cs_context *csctx,
1466 unsigned num,
1467 struct pipe_constant_buffer *buffers)
1468 {
1469 unsigned i;
1470
1471 LP_DBG(DEBUG_SETUP, "%s %p\n", __func__, (void *) buffers);
1472
1473 assert(num <= ARRAY_SIZE(csctx->constants));
1474
1475 for (i = 0; i < num; ++i) {
1476 util_copy_constant_buffer(&csctx->constants[i].current, &buffers[i], false);
1477 }
1478 for (; i < ARRAY_SIZE(csctx->constants); i++) {
1479 util_copy_constant_buffer(&csctx->constants[i].current, NULL, false);
1480 }
1481 }
1482
1483
1484 static void
lp_csctx_set_cs_ssbos(struct lp_cs_context * csctx,unsigned num,struct pipe_shader_buffer * buffers)1485 lp_csctx_set_cs_ssbos(struct lp_cs_context *csctx,
1486 unsigned num,
1487 struct pipe_shader_buffer *buffers)
1488 {
1489 int i;
1490 LP_DBG(DEBUG_SETUP, "%s %p\n", __func__, (void *)buffers);
1491
1492 assert (num <= ARRAY_SIZE(csctx->ssbos));
1493
1494 for (i = 0; i < num; ++i) {
1495 util_copy_shader_buffer(&csctx->ssbos[i].current, &buffers[i]);
1496 }
1497 for (; i < ARRAY_SIZE(csctx->ssbos); i++) {
1498 util_copy_shader_buffer(&csctx->ssbos[i].current, NULL);
1499 }
1500 }
1501
1502
1503 static void
lp_csctx_set_cs_images(struct lp_cs_context * csctx,unsigned num,struct pipe_image_view * images)1504 lp_csctx_set_cs_images(struct lp_cs_context *csctx,
1505 unsigned num,
1506 struct pipe_image_view *images)
1507 {
1508 unsigned i;
1509
1510 LP_DBG(DEBUG_SETUP, "%s %p\n", __func__, (void *) images);
1511
1512 assert(num <= ARRAY_SIZE(csctx->images));
1513
1514 for (i = 0; i < num; ++i) {
1515 struct pipe_image_view *image = &images[i];
1516 util_copy_image_view(&csctx->images[i].current, &images[i]);
1517
1518 struct pipe_resource *res = image->resource;
1519 struct llvmpipe_resource *lp_res = llvmpipe_resource(res);
1520 struct lp_jit_image *jit_image;
1521
1522 jit_image = &csctx->cs.current.jit_resources.images[i];
1523 if (!lp_res)
1524 continue;
1525
1526 lp_jit_image_from_pipe(jit_image, image);
1527 }
1528 for (; i < ARRAY_SIZE(csctx->images); i++) {
1529 util_copy_image_view(&csctx->images[i].current, NULL);
1530 }
1531 }
1532
1533
1534 static void
update_csctx_consts(struct llvmpipe_context * llvmpipe,struct lp_cs_context * csctx)1535 update_csctx_consts(struct llvmpipe_context *llvmpipe,
1536 struct lp_cs_context *csctx)
1537 {
1538 for (int i = 0; i < ARRAY_SIZE(csctx->constants); ++i) {
1539 lp_jit_buffer_from_pipe_const(&csctx->cs.current.jit_resources.constants[i],
1540 &csctx->constants[i].current, llvmpipe->pipe.screen);
1541 }
1542 }
1543
1544
1545 static void
update_csctx_ssbo(struct llvmpipe_context * llvmpipe,struct lp_cs_context * csctx)1546 update_csctx_ssbo(struct llvmpipe_context *llvmpipe,
1547 struct lp_cs_context *csctx)
1548 {
1549 for (int i = 0; i < ARRAY_SIZE(csctx->ssbos); ++i) {
1550 struct pipe_resource *buffer = csctx->ssbos[i].current.buffer;
1551 const uint8_t *current_data = NULL;
1552
1553 /* resource buffer */
1554 if (buffer)
1555 current_data = (uint8_t *) llvmpipe_resource_data(buffer);
1556 if (current_data) {
1557 current_data += csctx->ssbos[i].current.buffer_offset;
1558
1559 csctx->cs.current.jit_resources.ssbos[i].u = (const uint32_t *)current_data;
1560 csctx->cs.current.jit_resources.ssbos[i].num_elements = csctx->ssbos[i].current.buffer_size;
1561 } else {
1562 csctx->cs.current.jit_resources.ssbos[i].u = NULL;
1563 csctx->cs.current.jit_resources.ssbos[i].num_elements = 0;
1564 }
1565 }
1566 }
1567
1568
1569 static void
llvmpipe_cs_update_derived(struct llvmpipe_context * llvmpipe,const void * input)1570 llvmpipe_cs_update_derived(struct llvmpipe_context *llvmpipe, const void *input)
1571 {
1572 if (llvmpipe->cs_dirty & LP_CSNEW_CONSTANTS) {
1573 lp_csctx_set_cs_constants(llvmpipe->csctx,
1574 ARRAY_SIZE(llvmpipe->constants[PIPE_SHADER_COMPUTE]),
1575 llvmpipe->constants[PIPE_SHADER_COMPUTE]);
1576 update_csctx_consts(llvmpipe, llvmpipe->csctx);
1577 }
1578
1579 if (llvmpipe->cs_dirty & LP_CSNEW_SSBOS) {
1580 lp_csctx_set_cs_ssbos(llvmpipe->csctx,
1581 ARRAY_SIZE(llvmpipe->ssbos[PIPE_SHADER_COMPUTE]),
1582 llvmpipe->ssbos[PIPE_SHADER_COMPUTE]);
1583 update_csctx_ssbo(llvmpipe, llvmpipe->csctx);
1584 }
1585
1586 if (llvmpipe->cs_dirty & LP_CSNEW_SAMPLER_VIEW)
1587 lp_csctx_set_sampler_views(llvmpipe->csctx,
1588 llvmpipe->num_sampler_views[PIPE_SHADER_COMPUTE],
1589 llvmpipe->sampler_views[PIPE_SHADER_COMPUTE]);
1590
1591 if (llvmpipe->cs_dirty & LP_CSNEW_SAMPLER)
1592 lp_csctx_set_sampler_state(llvmpipe->csctx,
1593 llvmpipe->num_samplers[PIPE_SHADER_COMPUTE],
1594 llvmpipe->samplers[PIPE_SHADER_COMPUTE]);
1595
1596 if (llvmpipe->cs_dirty & LP_CSNEW_IMAGES)
1597 lp_csctx_set_cs_images(llvmpipe->csctx,
1598 ARRAY_SIZE(llvmpipe->images[PIPE_SHADER_COMPUTE]),
1599 llvmpipe->images[PIPE_SHADER_COMPUTE]);
1600
1601 struct lp_cs_context *csctx = llvmpipe->csctx;
1602 csctx->cs.current.jit_resources.aniso_filter_table = lp_build_sample_aniso_filter_table();
1603 if (input) {
1604 csctx->input = input;
1605 csctx->cs.current.jit_context.kernel_args = input;
1606 }
1607
1608 if (llvmpipe->cs_dirty & (LP_CSNEW_CS |
1609 LP_CSNEW_IMAGES |
1610 LP_CSNEW_SAMPLER_VIEW |
1611 LP_CSNEW_SAMPLER))
1612 llvmpipe_update_cs(llvmpipe);
1613
1614
1615 llvmpipe->cs_dirty = 0;
1616 }
1617
1618
1619 static void
cs_exec_fn(void * init_data,int iter_idx,struct lp_cs_local_mem * lmem)1620 cs_exec_fn(void *init_data, int iter_idx, struct lp_cs_local_mem *lmem)
1621 {
1622 struct lp_cs_job_info *job_info = init_data;
1623 struct lp_jit_cs_thread_data thread_data;
1624
1625 memset(&thread_data, 0, sizeof(thread_data));
1626
1627 if (lmem->local_size < job_info->req_local_mem) {
1628 lmem->local_mem_ptr = REALLOC(lmem->local_mem_ptr, lmem->local_size,
1629 job_info->req_local_mem);
1630 lmem->local_size = job_info->req_local_mem;
1631 }
1632 if (job_info->zero_initialize_shared_memory)
1633 memset(lmem->local_mem_ptr, 0, job_info->req_local_mem);
1634 thread_data.shared = lmem->local_mem_ptr;
1635
1636 thread_data.payload = job_info->payload;
1637
1638 unsigned grid_z, grid_y, grid_x;
1639
1640 if (job_info->use_iters) {
1641 grid_z = iter_idx / (job_info->iter_size[0] * job_info->iter_size[1]);
1642 grid_y = (iter_idx - (grid_z * (job_info->iter_size[0] * job_info->iter_size[1]))) / job_info->iter_size[0];
1643 grid_x = (iter_idx - (grid_z * (job_info->iter_size[0] * job_info->iter_size[1])) - (grid_y * job_info->iter_size[0]));
1644 } else {
1645 grid_z = iter_idx / (job_info->grid_size[0] * job_info->grid_size[1]);
1646 grid_y = (iter_idx - (grid_z * (job_info->grid_size[0] * job_info->grid_size[1]))) / job_info->grid_size[0];
1647 grid_x = (iter_idx - (grid_z * (job_info->grid_size[0] * job_info->grid_size[1])) - (grid_y * job_info->grid_size[0]));
1648 }
1649
1650 grid_z += job_info->grid_base[2];
1651 grid_y += job_info->grid_base[1];
1652 grid_x += job_info->grid_base[0];
1653 struct lp_compute_shader_variant *variant = job_info->current->variant;
1654
1655 void *io_ptr = NULL;
1656 if (job_info->io) {
1657 size_t io_offset = job_info->io_stride * iter_idx;
1658 io_ptr = (char *)job_info->io + io_offset;
1659 }
1660 if (thread_data.payload) {
1661 size_t payload_offset = job_info->payload_stride * iter_idx;
1662 thread_data.payload = (char *)thread_data.payload + payload_offset;
1663 }
1664 variant->jit_function(&job_info->current->jit_context,
1665 &job_info->current->jit_resources,
1666 job_info->block_size[0], job_info->block_size[1], job_info->block_size[2],
1667 grid_x, grid_y, grid_z,
1668 job_info->grid_size[0], job_info->grid_size[1], job_info->grid_size[2],
1669 job_info->work_dim, job_info->draw_id,
1670 io_ptr,
1671 &thread_data);
1672 }
1673
1674
1675 static void
fill_grid_size(struct pipe_context * pipe,int idx,const struct pipe_grid_info * info,uint32_t grid_size[3])1676 fill_grid_size(struct pipe_context *pipe,
1677 int idx,
1678 const struct pipe_grid_info *info,
1679 uint32_t grid_size[3])
1680 {
1681 struct pipe_transfer *transfer;
1682 uint32_t *params;
1683 if (!info->indirect) {
1684 grid_size[0] = info->grid[0];
1685 grid_size[1] = info->grid[1];
1686 grid_size[2] = info->grid[2];
1687 return;
1688 }
1689 params = pipe_buffer_map_range(pipe, info->indirect,
1690 (info->indirect_stride * idx) + info->indirect_offset,
1691 3 * sizeof(uint32_t),
1692 PIPE_MAP_READ,
1693 &transfer);
1694
1695 if (!transfer)
1696 return;
1697
1698 grid_size[0] = params[0];
1699 grid_size[1] = params[1];
1700 grid_size[2] = params[2];
1701 pipe_buffer_unmap(pipe, transfer);
1702 }
1703
1704
1705 static void
llvmpipe_launch_grid(struct pipe_context * pipe,const struct pipe_grid_info * info)1706 llvmpipe_launch_grid(struct pipe_context *pipe,
1707 const struct pipe_grid_info *info)
1708 {
1709 struct llvmpipe_context *llvmpipe = llvmpipe_context(pipe);
1710 struct llvmpipe_screen *screen = llvmpipe_screen(pipe->screen);
1711 struct lp_cs_job_info job_info;
1712
1713 if (!llvmpipe_check_render_cond(llvmpipe))
1714 return;
1715
1716 memset(&job_info, 0, sizeof(job_info));
1717
1718 llvmpipe_cs_update_derived(llvmpipe, info->input);
1719
1720 fill_grid_size(pipe, 0, info, job_info.grid_size);
1721
1722 job_info.grid_base[0] = info->grid_base[0];
1723 job_info.grid_base[1] = info->grid_base[1];
1724 job_info.grid_base[2] = info->grid_base[2];
1725 job_info.block_size[0] = info->block[0];
1726 job_info.block_size[1] = info->block[1];
1727 job_info.block_size[2] = info->block[2];
1728 job_info.work_dim = info->work_dim;
1729 job_info.req_local_mem = llvmpipe->cs->req_local_mem + info->variable_shared_mem;
1730 job_info.zero_initialize_shared_memory = llvmpipe->cs->zero_initialize_shared_memory;
1731 job_info.current = &llvmpipe->csctx->cs.current;
1732
1733 int num_tasks = job_info.grid_size[2] * job_info.grid_size[1] * job_info.grid_size[0];
1734 if (num_tasks) {
1735 struct lp_cs_tpool_task *task;
1736 mtx_lock(&screen->cs_mutex);
1737 task = lp_cs_tpool_queue_task(screen->cs_tpool, cs_exec_fn, &job_info, num_tasks);
1738 mtx_unlock(&screen->cs_mutex);
1739
1740 lp_cs_tpool_wait_for_task(screen->cs_tpool, &task);
1741 }
1742 if (!llvmpipe->queries_disabled)
1743 llvmpipe->pipeline_statistics.cs_invocations += num_tasks * info->block[0] * info->block[1] * info->block[2];
1744 }
1745
1746
1747 static void
llvmpipe_set_compute_resources(struct pipe_context * pipe,unsigned start,unsigned count,struct pipe_surface ** resources)1748 llvmpipe_set_compute_resources(struct pipe_context *pipe,
1749 unsigned start, unsigned count,
1750 struct pipe_surface **resources)
1751 {
1752 }
1753
1754
1755 static void
llvmpipe_set_global_binding(struct pipe_context * pipe,unsigned first,unsigned count,struct pipe_resource ** resources,uint32_t ** handles)1756 llvmpipe_set_global_binding(struct pipe_context *pipe,
1757 unsigned first, unsigned count,
1758 struct pipe_resource **resources,
1759 uint32_t **handles)
1760 {
1761 struct llvmpipe_context *llvmpipe = llvmpipe_context(pipe);
1762 struct lp_compute_shader *cs = llvmpipe->cs;
1763
1764 if (first + count > cs->max_global_buffers) {
1765 unsigned old_max = cs->max_global_buffers;
1766 cs->max_global_buffers = first + count;
1767 cs->global_buffers = realloc(cs->global_buffers,
1768 cs->max_global_buffers * sizeof(cs->global_buffers[0]));
1769 if (!cs->global_buffers) {
1770 return;
1771 }
1772
1773 memset(&cs->global_buffers[old_max], 0, (cs->max_global_buffers - old_max) * sizeof(cs->global_buffers[0]));
1774 }
1775
1776 if (!resources) {
1777 for (unsigned i = 0; i < count; i++)
1778 pipe_resource_reference(&cs->global_buffers[first + i], NULL);
1779 return;
1780 }
1781
1782 for (unsigned i = 0; i < count; i++) {
1783 uintptr_t va;
1784 uint32_t offset;
1785 pipe_resource_reference(&cs->global_buffers[first + i], resources[i]);
1786 struct llvmpipe_resource *lp_res = llvmpipe_resource(resources[i]);
1787 offset = *handles[i];
1788 va = (uintptr_t)((char *)lp_res->data + offset);
1789 memcpy(handles[i], &va, sizeof(va));
1790 }
1791 }
1792
1793
1794 void
llvmpipe_init_compute_funcs(struct llvmpipe_context * llvmpipe)1795 llvmpipe_init_compute_funcs(struct llvmpipe_context *llvmpipe)
1796 {
1797 llvmpipe->pipe.create_compute_state = llvmpipe_create_compute_state;
1798 llvmpipe->pipe.bind_compute_state = llvmpipe_bind_compute_state;
1799 llvmpipe->pipe.get_compute_state_info = llvmpipe_get_compute_state_info;
1800 llvmpipe->pipe.delete_compute_state = llvmpipe_delete_compute_state;
1801 llvmpipe->pipe.set_compute_resources = llvmpipe_set_compute_resources;
1802 llvmpipe->pipe.set_global_binding = llvmpipe_set_global_binding;
1803 llvmpipe->pipe.launch_grid = llvmpipe_launch_grid;
1804 }
1805
1806
1807 void
lp_csctx_destroy(struct lp_cs_context * csctx)1808 lp_csctx_destroy(struct lp_cs_context *csctx)
1809 {
1810 unsigned i;
1811 for (i = 0; i < ARRAY_SIZE(csctx->cs.current_tex); i++) {
1812 struct pipe_resource **res_ptr = &csctx->cs.current_tex[i];
1813 if (*res_ptr)
1814 llvmpipe_resource_unmap(*res_ptr, 0, 0);
1815 pipe_resource_reference(res_ptr, NULL);
1816 }
1817 for (i = 0; i < ARRAY_SIZE(csctx->constants); i++) {
1818 pipe_resource_reference(&csctx->constants[i].current.buffer, NULL);
1819 }
1820 for (i = 0; i < ARRAY_SIZE(csctx->ssbos); i++) {
1821 pipe_resource_reference(&csctx->ssbos[i].current.buffer, NULL);
1822 }
1823 for (i = 0; i < ARRAY_SIZE(csctx->images); i++) {
1824 pipe_resource_reference(&csctx->images[i].current.resource, NULL);
1825 }
1826 FREE(csctx);
1827 }
1828
1829
1830 struct lp_cs_context *
lp_csctx_create(struct pipe_context * pipe)1831 lp_csctx_create(struct pipe_context *pipe)
1832 {
1833 struct lp_cs_context *csctx = CALLOC_STRUCT(lp_cs_context);
1834 if (!csctx)
1835 return NULL;
1836
1837 csctx->pipe = pipe;
1838 return csctx;
1839 }
1840
1841 void
llvmpipe_update_task_shader(struct llvmpipe_context * lp)1842 llvmpipe_update_task_shader(struct llvmpipe_context *lp)
1843 {
1844 if (!lp->tss)
1845 return;
1846 struct lp_compute_shader_variant *variant = llvmpipe_update_cs_variant(lp, PIPE_SHADER_TASK, lp->tss);
1847 lp_cs_ctx_set_cs_variant(lp->task_ctx, variant);
1848 }
1849
1850 static void *
llvmpipe_create_ts_state(struct pipe_context * pipe,const struct pipe_shader_state * templ)1851 llvmpipe_create_ts_state(struct pipe_context *pipe,
1852 const struct pipe_shader_state *templ)
1853 {
1854 struct lp_compute_shader *shader = CALLOC_STRUCT(lp_compute_shader);
1855 if (!shader)
1856 return NULL;
1857
1858 llvmpipe_register_shader(pipe, templ);
1859
1860 shader->no = task_no++;
1861 shader->base.type = templ->type;
1862
1863 shader->base.ir.nir = templ->ir.nir;
1864 shader->req_local_mem += ((struct nir_shader *)shader->base.ir.nir)->info.shared_size;
1865 list_inithead(&shader->variants.list);
1866
1867 struct nir_shader *nir = shader->base.ir.nir;
1868 int nr_samplers = BITSET_LAST_BIT(nir->info.samplers_used);
1869 int nr_sampler_views = BITSET_LAST_BIT(nir->info.textures_used);
1870 int nr_images = BITSET_LAST_BIT(nir->info.images_used);
1871 shader->variant_key_size = lp_cs_variant_key_size(MAX2(nr_samplers, nr_sampler_views), nr_images);
1872 return shader;
1873 }
1874
1875
1876 static void
llvmpipe_bind_ts_state(struct pipe_context * pipe,void * _task)1877 llvmpipe_bind_ts_state(struct pipe_context *pipe, void *_task)
1878 {
1879 struct llvmpipe_context *llvmpipe = llvmpipe_context(pipe);
1880
1881 if (llvmpipe->tss == _task)
1882 return;
1883
1884 llvmpipe->tss = (struct lp_compute_shader *)_task;
1885 llvmpipe->dirty |= LP_NEW_TASK;
1886 }
1887
1888 static void
llvmpipe_delete_ts_state(struct pipe_context * pipe,void * _task)1889 llvmpipe_delete_ts_state(struct pipe_context *pipe, void *_task)
1890 {
1891 struct llvmpipe_context *llvmpipe = llvmpipe_context(pipe);
1892 struct lp_compute_shader *shader = _task;
1893 struct lp_cs_variant_list_item *li, *next;
1894
1895 /* Delete all the variants */
1896 LIST_FOR_EACH_ENTRY_SAFE(li, next, &shader->variants.list, list) {
1897 llvmpipe_remove_cs_shader_variant(llvmpipe, li->base);
1898 }
1899 ralloc_free(shader->base.ir.nir);
1900 FREE(shader);
1901 }
1902
1903 void
llvmpipe_init_task_funcs(struct llvmpipe_context * llvmpipe)1904 llvmpipe_init_task_funcs(struct llvmpipe_context *llvmpipe)
1905 {
1906 llvmpipe->pipe.create_ts_state = llvmpipe_create_ts_state;
1907 llvmpipe->pipe.bind_ts_state = llvmpipe_bind_ts_state;
1908 llvmpipe->pipe.delete_ts_state = llvmpipe_delete_ts_state;
1909 }
1910
1911 void
llvmpipe_update_mesh_shader(struct llvmpipe_context * lp)1912 llvmpipe_update_mesh_shader(struct llvmpipe_context *lp)
1913 {
1914 if (!lp->mhs)
1915 return;
1916 struct lp_compute_shader_variant *variant = llvmpipe_update_cs_variant(lp, PIPE_SHADER_MESH, lp->mhs);
1917 lp_cs_ctx_set_cs_variant(lp->mesh_ctx, variant);
1918 }
1919
1920 static void *
llvmpipe_create_ms_state(struct pipe_context * pipe,const struct pipe_shader_state * templ)1921 llvmpipe_create_ms_state(struct pipe_context *pipe,
1922 const struct pipe_shader_state *templ)
1923 {
1924 struct llvmpipe_context *llvmpipe = llvmpipe_context(pipe);
1925 struct lp_compute_shader *shader = CALLOC_STRUCT(lp_compute_shader);
1926 if (!shader)
1927 return NULL;
1928
1929 llvmpipe_register_shader(pipe, templ);
1930
1931 shader->no = mesh_no++;
1932 shader->base.type = templ->type;
1933
1934 shader->base.ir.nir = templ->ir.nir;
1935 shader->req_local_mem += ((struct nir_shader *)shader->base.ir.nir)->info.shared_size;
1936 list_inithead(&shader->variants.list);
1937
1938 shader->draw_mesh_data = draw_create_mesh_shader(llvmpipe->draw, templ);
1939 if (shader->draw_mesh_data == NULL) {
1940 FREE(shader);
1941 return NULL;
1942 }
1943
1944 struct nir_shader *nir = shader->base.ir.nir;
1945 int nr_samplers = BITSET_LAST_BIT(nir->info.samplers_used);
1946 int nr_sampler_views = BITSET_LAST_BIT(nir->info.textures_used);
1947 int nr_images = BITSET_LAST_BIT(nir->info.images_used);
1948 shader->variant_key_size = lp_cs_variant_key_size(MAX2(nr_samplers, nr_sampler_views), nr_images);
1949 return shader;
1950 }
1951
1952
1953 static void
llvmpipe_bind_ms_state(struct pipe_context * pipe,void * _mesh)1954 llvmpipe_bind_ms_state(struct pipe_context *pipe, void *_mesh)
1955 {
1956 struct llvmpipe_context *llvmpipe = llvmpipe_context(pipe);
1957
1958 if (llvmpipe->mhs == _mesh)
1959 return;
1960
1961 llvmpipe->mhs = (struct lp_compute_shader *)_mesh;
1962
1963 draw_bind_mesh_shader(llvmpipe->draw, _mesh ? llvmpipe->mhs->draw_mesh_data : NULL);
1964 llvmpipe->dirty |= LP_NEW_MESH;
1965 }
1966
1967
1968 static void
llvmpipe_delete_ms_state(struct pipe_context * pipe,void * _mesh)1969 llvmpipe_delete_ms_state(struct pipe_context *pipe, void *_mesh)
1970 {
1971 struct llvmpipe_context *llvmpipe = llvmpipe_context(pipe);
1972 struct lp_compute_shader *shader = _mesh;
1973 struct lp_cs_variant_list_item *li, *next;
1974
1975 /* Delete all the variants */
1976 LIST_FOR_EACH_ENTRY_SAFE(li, next, &shader->variants.list, list) {
1977 llvmpipe_remove_cs_shader_variant(llvmpipe, li->base);
1978 }
1979
1980 draw_delete_mesh_shader(llvmpipe->draw, shader->draw_mesh_data);
1981 ralloc_free(shader->base.ir.nir);
1982
1983 FREE(shader);
1984 }
1985
1986 static void
lp_mesh_call_draw(struct llvmpipe_context * lp,enum mesa_prim prim,int prim_out_idx,int cull_prim_idx,int task_idx,void * vbuf,size_t task_out_size,int vsize,int psize,int per_prim_count,size_t prim_offset)1987 lp_mesh_call_draw(struct llvmpipe_context *lp,
1988 enum mesa_prim prim,
1989 int prim_out_idx,
1990 int cull_prim_idx,
1991 int task_idx,
1992 void *vbuf, size_t task_out_size,
1993 int vsize, int psize, int per_prim_count,
1994 size_t prim_offset)
1995 {
1996 unsigned prim_len = mesa_vertices_per_prim(prim);
1997 uint32_t *ptr = (uint32_t *)((char *)vbuf + task_out_size * task_idx);
1998 uint32_t vertex_count = ptr[1];
1999 uint32_t prim_count = ptr[2];
2000
2001 if (!vertex_count || !prim_count)
2002 return;
2003
2004 struct draw_vertex_info vinfo;
2005 vinfo.verts = (struct vertex_header *)ptr;
2006 vinfo.vertex_size = vsize / 8;
2007 vinfo.stride = vsize;
2008 vinfo.count = vertex_count;
2009
2010 unsigned elts_size = prim_len * prim_count;
2011 unsigned short *elts = calloc(sizeof(uint16_t), elts_size);
2012 uint32_t *prim_lengths = calloc(prim_count, sizeof(uint32_t));
2013 int elts_idx = 0;
2014 char *prim_ptr = (char *)ptr + prim_offset;
2015 for (unsigned p = 0; p < prim_count; p++) {
2016 uint32_t *prim_idxs = (uint32_t *)(prim_ptr + p * psize + prim_out_idx * 4 * sizeof(float));
2017 for (unsigned elt = 0; elt < prim_len; elt++){
2018 elts[elts_idx++] = prim_idxs[elt];
2019 }
2020 prim_lengths[p] = prim_len;
2021 }
2022
2023 struct draw_prim_info prim_info = { 0 };
2024 prim_info.prim = prim;
2025 prim_info.linear = false;
2026 prim_info.elts = elts;
2027 prim_info.count = prim_count;
2028 prim_info.primitive_count = prim_count;
2029 prim_info.primitive_lengths = prim_lengths;
2030
2031 struct draw_vertex_info vert_out = { 0 };
2032 struct draw_prim_info prim_out = { 0 };
2033 draw_mesh_prim_run(lp->draw,
2034 per_prim_count,
2035 prim_ptr,
2036 cull_prim_idx,
2037 &prim_info,
2038 &vinfo,
2039 &prim_out,
2040 &vert_out);
2041 free(elts);
2042 free(prim_lengths);
2043
2044 draw_collect_primitives_generated(lp->draw,
2045 lp->active_primgen_queries &&
2046 !lp->queries_disabled);
2047 draw_mesh(lp->draw, &vert_out, &prim_out);
2048
2049 free(vert_out.verts);
2050 free(prim_out.primitive_lengths);
2051 }
2052
2053 static void
llvmpipe_draw_mesh_tasks(struct pipe_context * pipe,const struct pipe_grid_info * info)2054 llvmpipe_draw_mesh_tasks(struct pipe_context *pipe,
2055 const struct pipe_grid_info *info)
2056 {
2057 struct llvmpipe_context *lp = llvmpipe_context(pipe);
2058 struct llvmpipe_screen *screen = llvmpipe_screen(pipe->screen);
2059 struct lp_cs_job_info job_info;
2060
2061 if (!llvmpipe_check_render_cond(lp))
2062 return;
2063
2064 memset(&job_info, 0, sizeof(job_info));
2065 if (lp->dirty)
2066 llvmpipe_update_derived(lp);
2067
2068 unsigned draw_count = info->draw_count;
2069 if (info->indirect && info->indirect_draw_count) {
2070 struct pipe_transfer *dc_transfer;
2071 uint32_t *dc_param = pipe_buffer_map_range(pipe,
2072 info->indirect_draw_count,
2073 info->indirect_draw_count_offset,
2074 4, PIPE_MAP_READ, &dc_transfer);
2075 if (!dc_transfer) {
2076 debug_printf("%s: failed to map indirect draw count buffer\n", __func__);
2077 return;
2078 }
2079 if (dc_param[0] < draw_count)
2080 draw_count = dc_param[0];
2081 pipe_buffer_unmap(pipe, dc_transfer);
2082 }
2083
2084 struct nir_shader *mhs_shader = lp->mhs->base.ir.nir;
2085 int prim_out_idx = -1;
2086 int first_per_prim_idx = -1;
2087 int cull_prim_idx = -1;
2088 nir_foreach_shader_out_variable(var, mhs_shader) {
2089 if (var->data.per_primitive) {
2090 first_per_prim_idx = var->data.driver_location;
2091 break;
2092 }
2093 }
2094 nir_foreach_shader_out_variable(var, mhs_shader) {
2095 if (var->data.location == VARYING_SLOT_PRIMITIVE_INDICES) {
2096 prim_out_idx = var->data.driver_location;
2097 break;
2098 }
2099 }
2100 nir_foreach_shader_out_variable(var, mhs_shader) {
2101 if (var->data.location == VARYING_SLOT_CULL_PRIMITIVE) {
2102 cull_prim_idx = var->data.driver_location - first_per_prim_idx;
2103 break;
2104 }
2105 }
2106 int per_prim_count = util_bitcount64(mhs_shader->info.per_primitive_outputs);
2107 int out_count = util_bitcount64(mhs_shader->info.outputs_written);
2108 int per_vert_count = out_count - per_prim_count;
2109 int vsize = (sizeof(struct vertex_header) + per_vert_count * 4 * sizeof(float)) * 8;
2110 int psize = (per_prim_count * 4 * sizeof(float)) * 8;
2111 size_t prim_offset = vsize * (mhs_shader->info.mesh.max_vertices_out + 8);
2112 size_t task_out_size = prim_offset + psize * (mhs_shader->info.mesh.max_primitives_out + 8);
2113
2114 for (unsigned dr = 0; dr < draw_count; dr++) {
2115 fill_grid_size(pipe, dr, info, job_info.grid_size);
2116
2117 job_info.grid_base[0] = info->grid_base[0];
2118 job_info.grid_base[1] = info->grid_base[1];
2119 job_info.grid_base[2] = info->grid_base[2];
2120 job_info.block_size[0] = info->block[0];
2121 job_info.block_size[1] = info->block[1];
2122 job_info.block_size[2] = info->block[2];
2123
2124 void *payload = NULL;
2125 size_t payload_stride = 0;
2126 int num_tasks = job_info.grid_size[2] * job_info.grid_size[1] * job_info.grid_size[0];
2127 int num_mesh_invocs = 1;
2128 if (lp->tss) {
2129 struct nir_shader *tsk_shader = lp->tss->base.ir.nir;
2130 payload_stride = tsk_shader->info.task_payload_size + 3 * sizeof(uint32_t);
2131
2132 payload = calloc(num_tasks, payload_stride);
2133
2134 job_info.use_iters = false;
2135 job_info.payload = payload;
2136 job_info.payload_stride = payload_stride;
2137 job_info.work_dim = info->work_dim;
2138 job_info.draw_id = dr;
2139 job_info.req_local_mem = lp->tss->req_local_mem + info->variable_shared_mem;
2140 job_info.current = &lp->task_ctx->cs.current;
2141
2142 if (num_tasks) {
2143 struct lp_cs_tpool_task *task;
2144 mtx_lock(&screen->cs_mutex);
2145 task = lp_cs_tpool_queue_task(screen->cs_tpool, cs_exec_fn, &job_info, num_tasks);
2146 mtx_unlock(&screen->cs_mutex);
2147
2148 lp_cs_tpool_wait_for_task(screen->cs_tpool, &task);
2149 }
2150 if (!lp->queries_disabled)
2151 lp->pipeline_statistics.ts_invocations += num_tasks * info->block[0] * info->block[1] * info->block[2];
2152 num_mesh_invocs = num_tasks;
2153 }
2154
2155 for (unsigned i = 0; i < num_mesh_invocs; i++) {
2156 if (payload) {
2157 void *this_payload = (char *)payload + (payload_stride * i);
2158 uint32_t *payload_grid = (uint32_t *)this_payload;
2159 assert(lp->tss);
2160 job_info.grid_size[0] = payload_grid[0];
2161 job_info.grid_size[1] = payload_grid[1];
2162 job_info.grid_size[2] = payload_grid[2];
2163 job_info.payload = this_payload;
2164 job_info.block_size[0] = mhs_shader->info.workgroup_size[0];
2165 job_info.block_size[1] = mhs_shader->info.workgroup_size[1];
2166 job_info.block_size[2] = mhs_shader->info.workgroup_size[2];
2167 }
2168
2169 job_info.req_local_mem = lp->mhs->req_local_mem + info->variable_shared_mem;
2170 job_info.current = &lp->mesh_ctx->cs.current;
2171 job_info.payload_stride = 0;
2172 job_info.draw_id = dr;
2173 job_info.io_stride = task_out_size;
2174
2175 uint32_t job_strides[3] = { job_info.grid_size[0], job_info.grid_size[1], job_info.grid_size[2] };
2176 uint32_t total_grid[3] = { job_info.grid_size[0], job_info.grid_size[1], job_info.grid_size[2] };
2177 const unsigned int max_tasks = 4096;
2178 /* limit how large memory allocation can get for vbuf */
2179 for (unsigned g = 0; g < 3; g++) {
2180 if (job_strides[g] > max_tasks) {
2181 job_strides[g] = max_tasks;
2182 }
2183 }
2184
2185 for (unsigned grid_z = 0; grid_z < total_grid[2]; grid_z += job_strides[2]) {
2186 int this_z = MIN2(total_grid[2] - grid_z, max_tasks);
2187 job_info.grid_base[2] = grid_z;
2188 for (unsigned grid_y = 0; grid_y < total_grid[1]; grid_y += job_strides[1]) {
2189 int this_y = MIN2(total_grid[1] - grid_y, max_tasks);
2190 job_info.grid_base[1] = grid_y;
2191 for (unsigned grid_x = 0; grid_x < total_grid[0]; grid_x += job_strides[0]) {
2192 int this_x = MIN2(total_grid[0] - grid_x, max_tasks);
2193 job_info.grid_base[0] = grid_x;
2194 num_tasks = this_x * this_y * this_z;
2195
2196 job_info.iter_size[0] = this_x;
2197 job_info.iter_size[1] = this_y;
2198 job_info.iter_size[2] = this_z;
2199 job_info.use_iters = true;
2200
2201 void *vbuf = CALLOC(num_tasks, task_out_size);
2202 if (!vbuf)
2203 return;
2204
2205 job_info.io = vbuf;
2206 if (num_tasks) {
2207 struct lp_cs_tpool_task *task;
2208 mtx_lock(&screen->cs_mutex);
2209 task = lp_cs_tpool_queue_task(screen->cs_tpool, cs_exec_fn, &job_info, num_tasks);
2210 mtx_unlock(&screen->cs_mutex);
2211
2212 lp_cs_tpool_wait_for_task(screen->cs_tpool, &task);
2213 }
2214 if (!lp->queries_disabled)
2215 lp->pipeline_statistics.ms_invocations += num_tasks * job_info.block_size[0] * job_info.block_size[1] * job_info.block_size[2];
2216
2217 for (unsigned t = 0; t < num_tasks; t++)
2218 lp_mesh_call_draw(lp,
2219 mhs_shader->info.mesh.primitive_type,
2220 prim_out_idx - first_per_prim_idx,
2221 cull_prim_idx, t, vbuf, task_out_size,
2222 vsize, psize, per_prim_count, prim_offset);
2223 free(vbuf);
2224 }
2225 }
2226 }
2227 }
2228 free(payload);
2229 }
2230 draw_flush(lp->draw);
2231 }
2232
2233 void
llvmpipe_init_mesh_funcs(struct llvmpipe_context * llvmpipe)2234 llvmpipe_init_mesh_funcs(struct llvmpipe_context *llvmpipe)
2235 {
2236 llvmpipe->pipe.create_ms_state = llvmpipe_create_ms_state;
2237 llvmpipe->pipe.bind_ms_state = llvmpipe_bind_ms_state;
2238 llvmpipe->pipe.delete_ms_state = llvmpipe_delete_ms_state;
2239
2240 llvmpipe->pipe.draw_mesh_tasks = llvmpipe_draw_mesh_tasks;
2241 }
2242
2243 void
llvmpipe_task_update_derived(struct llvmpipe_context * llvmpipe)2244 llvmpipe_task_update_derived(struct llvmpipe_context *llvmpipe)
2245 {
2246 if (llvmpipe->dirty & LP_NEW_TASK_CONSTANTS) {
2247 lp_csctx_set_cs_constants(llvmpipe->task_ctx,
2248 ARRAY_SIZE(llvmpipe->constants[PIPE_SHADER_TASK]),
2249 llvmpipe->constants[PIPE_SHADER_TASK]);
2250 update_csctx_consts(llvmpipe, llvmpipe->task_ctx);
2251 }
2252
2253 if (llvmpipe->dirty & LP_NEW_TASK_SSBOS) {
2254 lp_csctx_set_cs_ssbos(llvmpipe->task_ctx,
2255 ARRAY_SIZE(llvmpipe->ssbos[PIPE_SHADER_TASK]),
2256 llvmpipe->ssbos[PIPE_SHADER_TASK]);
2257 update_csctx_ssbo(llvmpipe, llvmpipe->task_ctx);
2258 }
2259
2260 if (llvmpipe->dirty & LP_NEW_TASK_SAMPLER_VIEW)
2261 lp_csctx_set_sampler_views(llvmpipe->task_ctx,
2262 llvmpipe->num_sampler_views[PIPE_SHADER_TASK],
2263 llvmpipe->sampler_views[PIPE_SHADER_TASK]);
2264
2265 if (llvmpipe->dirty & LP_NEW_TASK_SAMPLER)
2266 lp_csctx_set_sampler_state(llvmpipe->task_ctx,
2267 llvmpipe->num_samplers[PIPE_SHADER_TASK],
2268 llvmpipe->samplers[PIPE_SHADER_TASK]);
2269
2270 if (llvmpipe->dirty & LP_NEW_TASK_IMAGES)
2271 lp_csctx_set_cs_images(llvmpipe->task_ctx,
2272 ARRAY_SIZE(llvmpipe->images[PIPE_SHADER_TASK]),
2273 llvmpipe->images[PIPE_SHADER_TASK]);
2274
2275 struct lp_cs_context *csctx = llvmpipe->task_ctx;
2276 csctx->cs.current.jit_resources.aniso_filter_table = lp_build_sample_aniso_filter_table();
2277 }
2278
2279 void
llvmpipe_mesh_update_derived(struct llvmpipe_context * llvmpipe)2280 llvmpipe_mesh_update_derived(struct llvmpipe_context *llvmpipe)
2281 {
2282 if (llvmpipe->dirty & LP_NEW_MESH_CONSTANTS) {
2283 lp_csctx_set_cs_constants(llvmpipe->mesh_ctx,
2284 ARRAY_SIZE(llvmpipe->constants[PIPE_SHADER_MESH]),
2285 llvmpipe->constants[PIPE_SHADER_MESH]);
2286 update_csctx_consts(llvmpipe, llvmpipe->mesh_ctx);
2287 }
2288
2289 if (llvmpipe->dirty & LP_NEW_MESH_SSBOS) {
2290 lp_csctx_set_cs_ssbos(llvmpipe->mesh_ctx,
2291 ARRAY_SIZE(llvmpipe->ssbos[PIPE_SHADER_MESH]),
2292 llvmpipe->ssbos[PIPE_SHADER_MESH]);
2293 update_csctx_ssbo(llvmpipe, llvmpipe->mesh_ctx);
2294 }
2295
2296 if (llvmpipe->dirty & LP_NEW_MESH_SAMPLER_VIEW)
2297 lp_csctx_set_sampler_views(llvmpipe->mesh_ctx,
2298 llvmpipe->num_sampler_views[PIPE_SHADER_MESH],
2299 llvmpipe->sampler_views[PIPE_SHADER_MESH]);
2300
2301 if (llvmpipe->dirty & LP_NEW_MESH_SAMPLER)
2302 lp_csctx_set_sampler_state(llvmpipe->mesh_ctx,
2303 llvmpipe->num_samplers[PIPE_SHADER_MESH],
2304 llvmpipe->samplers[PIPE_SHADER_MESH]);
2305
2306 if (llvmpipe->dirty & LP_NEW_MESH_IMAGES)
2307 lp_csctx_set_cs_images(llvmpipe->mesh_ctx,
2308 ARRAY_SIZE(llvmpipe->images[PIPE_SHADER_MESH]),
2309 llvmpipe->images[PIPE_SHADER_MESH]);
2310
2311 struct lp_cs_context *csctx = llvmpipe->mesh_ctx;
2312 csctx->cs.current.jit_resources.aniso_filter_table = lp_build_sample_aniso_filter_table();
2313 }
2314