1 /*
2 * Copyright 2018 Collabora Ltd.
3 *
4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the "Software"),
6 * to deal in the Software without restriction, including without limitation
7 * on the rights to use, copy, modify, merge, publish, distribute, sub
8 * license, and/or sell copies of the Software, and to permit persons to whom
9 * the Software is furnished to do so, subject to the following conditions:
10 *
11 * The above copyright notice and this permission notice (including the next
12 * paragraph) shall be included in all copies or substantial portions of the
13 * Software.
14 *
15 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
16 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
17 * FITNESS FOR A PARTICULAR PURPOSE AND NON-INFRINGEMENT. IN NO EVENT SHALL
18 * THE AUTHOR(S) AND/OR THEIR SUPPLIERS BE LIABLE FOR ANY CLAIM,
19 * DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
20 * OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
21 * USE OR OTHER DEALINGS IN THE SOFTWARE.
22 */
23
24 #include "nir_to_spirv.h"
25 #include "spirv_builder.h"
26
27 #include "nir.h"
28 #include "pipe/p_state.h"
29 #include "util/u_math.h"
30 #include "util/u_memory.h"
31 #include "util/hash_table.h"
32
33 #define SLOT_UNSET ((unsigned char) -1)
34
35 struct ntv_context {
36 void *mem_ctx;
37
38 /* SPIR-V 1.4 and later requires entrypoints to list all global
39 * variables in the interface.
40 */
41 bool spirv_1_4_interfaces;
42
43 bool explicit_lod; //whether to set lod=0 for texture()
44
45 struct spirv_builder builder;
46 nir_shader *nir;
47
48 struct hash_table *glsl_types;
49 struct hash_table *bo_struct_types;
50 struct hash_table *bo_array_types;
51
52 SpvId GLSL_std_450;
53
54 gl_shader_stage stage;
55 const struct zink_shader_info *sinfo;
56
57 SpvId ubos[PIPE_MAX_CONSTANT_BUFFERS][5]; //8, 16, 32, unused, 64
58 nir_variable *ubo_vars[PIPE_MAX_CONSTANT_BUFFERS];
59
60 SpvId ssbos[5]; //8, 16, 32, unused, 64
61 nir_variable *ssbo_vars;
62
63 SpvId images[PIPE_MAX_SHADER_IMAGES];
64 struct hash_table image_types;
65 SpvId samplers[PIPE_MAX_SHADER_SAMPLER_VIEWS];
66 SpvId bindless_samplers[2];
67 nir_variable *sampler_var[PIPE_MAX_SHADER_SAMPLER_VIEWS]; /* driver_location -> variable */
68 nir_variable *bindless_sampler_var[2];
69 unsigned last_sampler;
70 unsigned bindless_set_idx;
71 nir_variable *image_var[PIPE_MAX_SHADER_IMAGES]; /* driver_location -> variable */
72
73 SpvId entry_ifaces[PIPE_MAX_SHADER_INPUTS * 4 + PIPE_MAX_SHADER_OUTPUTS * 4];
74 size_t num_entry_ifaces;
75
76 SpvId *defs;
77 nir_alu_type *def_types;
78 SpvId *resident_defs;
79 size_t num_defs;
80
81 struct hash_table *vars; /* nir_variable -> SpvId */
82
83 const SpvId *block_ids;
84 size_t num_blocks;
85 bool block_started;
86 SpvId loop_break, loop_cont;
87
88 SpvId shared_block_var[5]; //8, 16, 32, unused, 64
89 SpvId shared_block_arr_type[5]; //8, 16, 32, unused, 64
90 SpvId scratch_block_var[5]; //8, 16, 32, unused, 64
91
92 SpvId front_face_var, instance_id_var, vertex_id_var,
93 primitive_id_var, invocation_id_var, // geometry
94 sample_mask_type, sample_id_var, sample_pos_var, sample_mask_in_var,
95 tess_patch_vertices_in, tess_coord_var, // tess
96 push_const_var, point_coord_var,
97 workgroup_id_var, num_workgroups_var,
98 local_invocation_id_var, global_invocation_id_var,
99 local_invocation_index_var, helper_invocation_var,
100 local_group_size_var, view_index_var,
101 base_vertex_var, base_instance_var, draw_id_var;
102
103 SpvId shared_mem_size;
104
105 SpvId subgroup_eq_mask_var,
106 subgroup_ge_mask_var,
107 subgroup_gt_mask_var,
108 subgroup_id_var,
109 subgroup_invocation_var,
110 subgroup_le_mask_var,
111 subgroup_lt_mask_var,
112 subgroup_size_var;
113
114 SpvId discard_func;
115 SpvId float_array_type[2];
116
117 bool use_gl_per_vertex;
118 SpvId per_vertex_block_type[2];
119 SpvId per_vertex_in;
120 SpvId per_vertex_out;
121 SpvId float_array_1_type;
122 int gl_in_num_vertices[3];
123 int gl_out_num_vertices;
124 };
125
126 static SpvId
127 get_fvec_constant(struct ntv_context *ctx, unsigned bit_size,
128 unsigned num_components, double value);
129
130 static SpvId
131 get_ivec_constant(struct ntv_context *ctx, unsigned bit_size,
132 unsigned num_components, int64_t value);
133
134 static SpvId
135 emit_unop(struct ntv_context *ctx, SpvOp op, SpvId type, SpvId src);
136
137 static SpvId
138 emit_binop(struct ntv_context *ctx, SpvOp op, SpvId type,
139 SpvId src0, SpvId src1);
140
141 static SpvId
142 emit_triop(struct ntv_context *ctx, SpvOp op, SpvId type,
143 SpvId src0, SpvId src1, SpvId src2);
144
145 static bool
alu_op_is_typeless(nir_op op)146 alu_op_is_typeless(nir_op op)
147 {
148 switch (op) {
149 case nir_op_mov:
150 case nir_op_vec16:
151 case nir_op_vec2:
152 case nir_op_vec3:
153 case nir_op_vec4:
154 case nir_op_vec5:
155 case nir_op_vec8:
156 case nir_op_bcsel:
157 return true;
158 default:
159 break;
160 }
161 return false;
162 }
163
164 static nir_alu_type
get_nir_alu_type(const struct glsl_type * type)165 get_nir_alu_type(const struct glsl_type *type)
166 {
167 return nir_alu_type_get_base_type(nir_get_nir_type_for_glsl_base_type(glsl_get_base_type(glsl_without_array_or_matrix(type))));
168 }
169
170 static nir_alu_type
171 infer_nir_alu_type_from_uses_ssa(nir_def *ssa);
172
173 static nir_alu_type
infer_nir_alu_type_from_use(nir_src * src)174 infer_nir_alu_type_from_use(nir_src *src)
175 {
176 nir_instr *instr = nir_src_parent_instr(src);
177 nir_alu_type atype = nir_type_invalid;
178 switch (instr->type) {
179 case nir_instr_type_alu: {
180 nir_alu_instr *alu = nir_instr_as_alu(instr);
181 if (alu->op == nir_op_bcsel) {
182 if (nir_srcs_equal(alu->src[0].src, *src)) {
183 /* special case: the first src in bcsel is always bool */
184 return nir_type_bool;
185 }
186 }
187 /* ignore typeless ops */
188 if (alu_op_is_typeless(alu->op)) {
189 atype = infer_nir_alu_type_from_uses_ssa(&alu->def);
190 break;
191 }
192 for (unsigned i = 0; i < nir_op_infos[alu->op].num_inputs; i++) {
193 if (!nir_srcs_equal(alu->src[i].src, *src))
194 continue;
195 atype = nir_op_infos[alu->op].input_types[i];
196 break;
197 }
198 break;
199 }
200 case nir_instr_type_tex: {
201 nir_tex_instr *tex = nir_instr_as_tex(instr);
202 for (unsigned i = 0; i < tex->num_srcs; i++) {
203 if (!nir_srcs_equal(tex->src[i].src, *src))
204 continue;
205 switch (tex->src[i].src_type) {
206 case nir_tex_src_coord:
207 case nir_tex_src_lod:
208 if (tex->op == nir_texop_txf ||
209 tex->op == nir_texop_txf_ms ||
210 tex->op == nir_texop_txs)
211 atype = nir_type_int;
212 else
213 atype = nir_type_float;
214 break;
215 case nir_tex_src_projector:
216 case nir_tex_src_bias:
217 case nir_tex_src_min_lod:
218 case nir_tex_src_comparator:
219 case nir_tex_src_ddx:
220 case nir_tex_src_ddy:
221 atype = nir_type_float;
222 break;
223 case nir_tex_src_offset:
224 case nir_tex_src_ms_index:
225 case nir_tex_src_texture_offset:
226 case nir_tex_src_sampler_offset:
227 case nir_tex_src_sampler_handle:
228 case nir_tex_src_texture_handle:
229 atype = nir_type_int;
230 break;
231 default:
232 break;
233 }
234 break;
235 }
236 break;
237 }
238 case nir_instr_type_intrinsic: {
239 if (nir_instr_as_intrinsic(instr)->intrinsic == nir_intrinsic_load_deref) {
240 atype = get_nir_alu_type(nir_instr_as_deref(instr)->type);
241 } else if (nir_instr_as_intrinsic(instr)->intrinsic == nir_intrinsic_store_deref) {
242 atype = get_nir_alu_type(nir_src_as_deref(nir_instr_as_intrinsic(instr)->src[0])->type);
243 }
244 break;
245 }
246 default:
247 break;
248 }
249 return nir_alu_type_get_base_type(atype);
250 }
251
252 static nir_alu_type
infer_nir_alu_type_from_uses_ssa(nir_def * ssa)253 infer_nir_alu_type_from_uses_ssa(nir_def *ssa)
254 {
255 nir_alu_type atype = nir_type_invalid;
256 /* try to infer a type: if it's wrong then whatever, but at least we tried */
257 nir_foreach_use_including_if(src, ssa) {
258 if (nir_src_is_if(src))
259 return nir_type_bool;
260 atype = infer_nir_alu_type_from_use(src);
261 if (atype)
262 break;
263 }
264 return atype ? atype : nir_type_uint;
265 }
266
267 static SpvId
get_bvec_type(struct ntv_context * ctx,int num_components)268 get_bvec_type(struct ntv_context *ctx, int num_components)
269 {
270 SpvId bool_type = spirv_builder_type_bool(&ctx->builder);
271 if (num_components > 1)
272 return spirv_builder_type_vector(&ctx->builder, bool_type,
273 num_components);
274
275 assert(num_components == 1);
276 return bool_type;
277 }
278
279 static SpvId
find_image_type(struct ntv_context * ctx,nir_variable * var)280 find_image_type(struct ntv_context *ctx, nir_variable *var)
281 {
282 struct hash_entry *he = _mesa_hash_table_search(&ctx->image_types, var);
283 return he ? (intptr_t)he->data : 0;
284 }
285
286 static SpvScope
get_scope(mesa_scope scope)287 get_scope(mesa_scope scope)
288 {
289 SpvScope conv[] = {
290 [SCOPE_NONE] = 0,
291 [SCOPE_INVOCATION] = SpvScopeInvocation,
292 [SCOPE_SUBGROUP] = SpvScopeSubgroup,
293 [SCOPE_SHADER_CALL] = SpvScopeShaderCallKHR,
294 [SCOPE_WORKGROUP] = SpvScopeWorkgroup,
295 [SCOPE_QUEUE_FAMILY] = SpvScopeQueueFamily,
296 [SCOPE_DEVICE] = SpvScopeDevice,
297 };
298 return conv[scope];
299 }
300
301 static SpvId
block_label(struct ntv_context * ctx,nir_block * block)302 block_label(struct ntv_context *ctx, nir_block *block)
303 {
304 assert(block->index < ctx->num_blocks);
305 return ctx->block_ids[block->index];
306 }
307
308 static void
emit_access_decorations(struct ntv_context * ctx,nir_variable * var,SpvId var_id)309 emit_access_decorations(struct ntv_context *ctx, nir_variable *var, SpvId var_id)
310 {
311 u_foreach_bit(bit, var->data.access) {
312 switch (1 << bit) {
313 case ACCESS_COHERENT:
314 /* SpvDecorationCoherent can't be used with vulkan memory model */
315 break;
316 case ACCESS_RESTRICT:
317 spirv_builder_emit_decoration(&ctx->builder, var_id, SpvDecorationRestrict);
318 break;
319 case ACCESS_VOLATILE:
320 /* SpvDecorationVolatile can't be used with vulkan memory model */
321 break;
322 case ACCESS_NON_READABLE:
323 spirv_builder_emit_decoration(&ctx->builder, var_id, SpvDecorationNonReadable);
324 break;
325 case ACCESS_NON_WRITEABLE:
326 spirv_builder_emit_decoration(&ctx->builder, var_id, SpvDecorationNonWritable);
327 break;
328 case ACCESS_NON_UNIFORM:
329 spirv_builder_emit_decoration(&ctx->builder, var_id, SpvDecorationNonUniform);
330 break;
331 case ACCESS_CAN_REORDER:
332 case ACCESS_NON_TEMPORAL:
333 /* no equivalent */
334 break;
335 default:
336 unreachable("unknown access bit");
337 }
338 }
339 /* The Simple, GLSL, and Vulkan memory models can assume that aliasing is generally
340 * not present between the memory object declarations. Specifically, the consumer
341 * is free to assume aliasing is not present between memory object declarations,
342 * unless the memory object declarations explicitly indicate they alias.
343 * ...
344 * Applying Restrict is allowed, but has no effect.
345 * ...
346 * Only those memory object declarations decorated with Aliased or AliasedPointer may alias each other.
347 *
348 * - SPIRV 2.18.2 Aliasing
349 *
350 * thus if the variable isn't marked restrict, assume it may alias
351 */
352 if (!(var->data.access & ACCESS_RESTRICT))
353 spirv_builder_emit_decoration(&ctx->builder, var_id, SpvDecorationAliased);
354 }
355
356 static SpvOp
get_atomic_op(struct ntv_context * ctx,unsigned bit_size,nir_atomic_op op)357 get_atomic_op(struct ntv_context *ctx, unsigned bit_size, nir_atomic_op op)
358 {
359 switch (op) {
360 #define ATOMIC_FCAP(NAME) \
361 do {\
362 if (bit_size == 16) \
363 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityAtomicFloat16##NAME##EXT); \
364 if (bit_size == 32) \
365 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityAtomicFloat32##NAME##EXT); \
366 if (bit_size == 64) \
367 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityAtomicFloat64##NAME##EXT); \
368 } while (0)
369
370 case nir_atomic_op_fadd:
371 ATOMIC_FCAP(Add);
372 if (bit_size == 16)
373 spirv_builder_emit_extension(&ctx->builder, "SPV_EXT_shader_atomic_float16_add");
374 else
375 spirv_builder_emit_extension(&ctx->builder, "SPV_EXT_shader_atomic_float_add");
376 return SpvOpAtomicFAddEXT;
377 case nir_atomic_op_fmax:
378 ATOMIC_FCAP(MinMax);
379 spirv_builder_emit_extension(&ctx->builder, "SPV_EXT_shader_atomic_float_min_max");
380 return SpvOpAtomicFMaxEXT;
381 case nir_atomic_op_fmin:
382 ATOMIC_FCAP(MinMax);
383 spirv_builder_emit_extension(&ctx->builder, "SPV_EXT_shader_atomic_float_min_max");
384 return SpvOpAtomicFMinEXT;
385
386 case nir_atomic_op_iadd:
387 return SpvOpAtomicIAdd;
388 case nir_atomic_op_umin:
389 return SpvOpAtomicUMin;
390 case nir_atomic_op_imin:
391 return SpvOpAtomicSMin;
392 case nir_atomic_op_umax:
393 return SpvOpAtomicUMax;
394 case nir_atomic_op_imax:
395 return SpvOpAtomicSMax;
396 case nir_atomic_op_iand:
397 return SpvOpAtomicAnd;
398 case nir_atomic_op_ior:
399 return SpvOpAtomicOr;
400 case nir_atomic_op_ixor:
401 return SpvOpAtomicXor;
402 case nir_atomic_op_xchg:
403 return SpvOpAtomicExchange;
404 case nir_atomic_op_cmpxchg:
405 return SpvOpAtomicCompareExchange;
406 default:
407 debug_printf("%s - ", nir_intrinsic_infos[op].name);
408 unreachable("unhandled atomic op");
409 }
410 return 0;
411 }
412
413 static SpvId
emit_float_const(struct ntv_context * ctx,int bit_size,double value)414 emit_float_const(struct ntv_context *ctx, int bit_size, double value)
415 {
416 assert(bit_size == 16 || bit_size == 32 || bit_size == 64);
417 return spirv_builder_const_float(&ctx->builder, bit_size, value);
418 }
419
420 static SpvId
emit_uint_const(struct ntv_context * ctx,int bit_size,uint64_t value)421 emit_uint_const(struct ntv_context *ctx, int bit_size, uint64_t value)
422 {
423 assert(bit_size == 8 || bit_size == 16 || bit_size == 32 || bit_size == 64);
424 return spirv_builder_const_uint(&ctx->builder, bit_size, value);
425 }
426
427 static SpvId
emit_int_const(struct ntv_context * ctx,int bit_size,int64_t value)428 emit_int_const(struct ntv_context *ctx, int bit_size, int64_t value)
429 {
430 assert(bit_size == 8 || bit_size == 16 || bit_size == 32 || bit_size == 64);
431 return spirv_builder_const_int(&ctx->builder, bit_size, value);
432 }
433
434 static SpvId
get_fvec_type(struct ntv_context * ctx,unsigned bit_size,unsigned num_components)435 get_fvec_type(struct ntv_context *ctx, unsigned bit_size, unsigned num_components)
436 {
437 assert(bit_size == 16 || bit_size == 32 || bit_size == 64);
438
439 SpvId float_type = spirv_builder_type_float(&ctx->builder, bit_size);
440 if (num_components > 1)
441 return spirv_builder_type_vector(&ctx->builder, float_type,
442 num_components);
443
444 assert(num_components == 1);
445 return float_type;
446 }
447
448 static SpvId
get_ivec_type(struct ntv_context * ctx,unsigned bit_size,unsigned num_components)449 get_ivec_type(struct ntv_context *ctx, unsigned bit_size, unsigned num_components)
450 {
451 assert(bit_size == 8 || bit_size == 16 || bit_size == 32 || bit_size == 64);
452
453 SpvId int_type = spirv_builder_type_int(&ctx->builder, bit_size);
454 if (num_components > 1)
455 return spirv_builder_type_vector(&ctx->builder, int_type,
456 num_components);
457
458 assert(num_components == 1);
459 return int_type;
460 }
461
462 static SpvId
get_uvec_type(struct ntv_context * ctx,unsigned bit_size,unsigned num_components)463 get_uvec_type(struct ntv_context *ctx, unsigned bit_size, unsigned num_components)
464 {
465 assert(bit_size == 8 || bit_size == 16 || bit_size == 32 || bit_size == 64);
466
467 SpvId uint_type = spirv_builder_type_uint(&ctx->builder, bit_size);
468 if (num_components > 1)
469 return spirv_builder_type_vector(&ctx->builder, uint_type,
470 num_components);
471
472 assert(num_components == 1);
473 return uint_type;
474 }
475
476 static SpvId
get_alu_type(struct ntv_context * ctx,nir_alu_type type,unsigned num_components,unsigned bit_size)477 get_alu_type(struct ntv_context *ctx, nir_alu_type type, unsigned num_components, unsigned bit_size)
478 {
479 if (bit_size == 1)
480 return get_bvec_type(ctx, num_components);
481
482 type = nir_alu_type_get_base_type(type);
483 switch (nir_alu_type_get_base_type(type)) {
484 case nir_type_bool:
485 return get_bvec_type(ctx, num_components);
486
487 case nir_type_int:
488 return get_ivec_type(ctx, bit_size, num_components);
489
490 case nir_type_uint:
491 return get_uvec_type(ctx, bit_size, num_components);
492
493 case nir_type_float:
494 return get_fvec_type(ctx, bit_size, num_components);
495
496 default:
497 unreachable("unsupported nir_alu_type");
498 }
499 }
500
501 static SpvStorageClass
get_storage_class(struct nir_variable * var)502 get_storage_class(struct nir_variable *var)
503 {
504 switch (var->data.mode) {
505 case nir_var_function_temp:
506 return SpvStorageClassFunction;
507 case nir_var_mem_push_const:
508 return SpvStorageClassPushConstant;
509 case nir_var_shader_in:
510 return SpvStorageClassInput;
511 case nir_var_shader_out:
512 return SpvStorageClassOutput;
513 case nir_var_uniform:
514 case nir_var_image:
515 return SpvStorageClassUniformConstant;
516 case nir_var_mem_ubo:
517 return SpvStorageClassUniform;
518 case nir_var_mem_ssbo:
519 return SpvStorageClassStorageBuffer;
520 default:
521 unreachable("Unsupported nir_variable_mode");
522 }
523 return 0;
524 }
525
526 static SpvId
get_def_uvec_type(struct ntv_context * ctx,nir_def * def)527 get_def_uvec_type(struct ntv_context *ctx, nir_def *def)
528 {
529 unsigned bit_size = def->bit_size;
530 return get_uvec_type(ctx, bit_size, def->num_components);
531 }
532
533 static SpvId
get_glsl_basetype(struct ntv_context * ctx,enum glsl_base_type type)534 get_glsl_basetype(struct ntv_context *ctx, enum glsl_base_type type)
535 {
536 switch (type) {
537 case GLSL_TYPE_BOOL:
538 return spirv_builder_type_bool(&ctx->builder);
539
540 case GLSL_TYPE_FLOAT16:
541 return spirv_builder_type_float(&ctx->builder, 16);
542
543 case GLSL_TYPE_FLOAT:
544 return spirv_builder_type_float(&ctx->builder, 32);
545
546 case GLSL_TYPE_INT:
547 return spirv_builder_type_int(&ctx->builder, 32);
548
549 case GLSL_TYPE_UINT:
550 return spirv_builder_type_uint(&ctx->builder, 32);
551
552 case GLSL_TYPE_DOUBLE:
553 return spirv_builder_type_float(&ctx->builder, 64);
554
555 case GLSL_TYPE_INT64:
556 return spirv_builder_type_int(&ctx->builder, 64);
557
558 case GLSL_TYPE_UINT64:
559 return spirv_builder_type_uint(&ctx->builder, 64);
560
561 case GLSL_TYPE_UINT16:
562 return spirv_builder_type_uint(&ctx->builder, 16);
563 case GLSL_TYPE_INT16:
564 return spirv_builder_type_int(&ctx->builder, 16);
565 case GLSL_TYPE_INT8:
566 return spirv_builder_type_int(&ctx->builder, 8);
567 case GLSL_TYPE_UINT8:
568 return spirv_builder_type_uint(&ctx->builder, 8);
569
570 default:
571 unreachable("unknown GLSL type");
572 }
573 }
574
575 static SpvId
get_glsl_type(struct ntv_context * ctx,const struct glsl_type * type)576 get_glsl_type(struct ntv_context *ctx, const struct glsl_type *type)
577 {
578 assert(type);
579 if (glsl_type_is_scalar(type))
580 return get_glsl_basetype(ctx, glsl_get_base_type(type));
581
582 if (glsl_type_is_vector(type))
583 return spirv_builder_type_vector(&ctx->builder,
584 get_glsl_basetype(ctx, glsl_get_base_type(type)),
585 glsl_get_vector_elements(type));
586
587 if (glsl_type_is_matrix(type))
588 return spirv_builder_type_matrix(&ctx->builder,
589 spirv_builder_type_vector(&ctx->builder,
590 get_glsl_basetype(ctx, glsl_get_base_type(type)),
591 glsl_get_vector_elements(type)),
592 glsl_get_matrix_columns(type));
593
594 /* Aggregate types aren't cached in spirv_builder, so let's cache
595 * them here instead.
596 */
597
598 struct hash_entry *entry =
599 _mesa_hash_table_search(ctx->glsl_types, type);
600 if (entry)
601 return (SpvId)(uintptr_t)entry->data;
602
603 SpvId ret;
604 if (glsl_type_is_array(type)) {
605 SpvId element_type = get_glsl_type(ctx, glsl_get_array_element(type));
606 if (glsl_type_is_unsized_array(type))
607 ret = spirv_builder_type_runtime_array(&ctx->builder, element_type);
608 else
609 ret = spirv_builder_type_array(&ctx->builder,
610 element_type,
611 emit_uint_const(ctx, 32, glsl_get_length(type)));
612 uint32_t stride = glsl_get_explicit_stride(type);
613 if (!stride && glsl_type_is_scalar(glsl_get_array_element(type))) {
614 stride = MAX2(glsl_get_bit_size(glsl_get_array_element(type)) / 8, 1);
615 }
616 if (stride)
617 spirv_builder_emit_array_stride(&ctx->builder, ret, stride);
618 } else if (glsl_type_is_struct_or_ifc(type)) {
619 const unsigned length = glsl_get_length(type);
620
621 /* allocate some SpvId on the stack, falling back to the heap if the array is too long */
622 SpvId *types, types_stack[16];
623
624 if (length <= ARRAY_SIZE(types_stack)) {
625 types = types_stack;
626 } else {
627 types = ralloc_array_size(ctx->mem_ctx, sizeof(SpvId), length);
628 assert(types != NULL);
629 }
630
631 for (unsigned i = 0; i < glsl_get_length(type); i++)
632 types[i] = get_glsl_type(ctx, glsl_get_struct_field(type, i));
633 ret = spirv_builder_type_struct(&ctx->builder, types,
634 glsl_get_length(type));
635 for (unsigned i = 0; i < glsl_get_length(type); i++) {
636 int32_t offset = glsl_get_struct_field_offset(type, i);
637 if (offset >= 0)
638 spirv_builder_emit_member_offset(&ctx->builder, ret, i, offset);
639 }
640 } else
641 unreachable("Unhandled GLSL type");
642
643 _mesa_hash_table_insert(ctx->glsl_types, type, (void *)(uintptr_t)ret);
644 return ret;
645 }
646
647 static void
create_scratch_block(struct ntv_context * ctx,unsigned scratch_size,unsigned bit_size)648 create_scratch_block(struct ntv_context *ctx, unsigned scratch_size, unsigned bit_size)
649 {
650 unsigned idx = bit_size >> 4;
651 SpvId type = spirv_builder_type_uint(&ctx->builder, bit_size);
652 unsigned block_size = scratch_size / (bit_size / 8);
653 assert(block_size);
654 SpvId array = spirv_builder_type_array(&ctx->builder, type, emit_uint_const(ctx, 32, block_size));
655 spirv_builder_emit_array_stride(&ctx->builder, array, bit_size / 8);
656 SpvId ptr_type = spirv_builder_type_pointer(&ctx->builder,
657 SpvStorageClassPrivate,
658 array);
659 ctx->scratch_block_var[idx] = spirv_builder_emit_var(&ctx->builder, ptr_type, SpvStorageClassPrivate);
660 if (ctx->spirv_1_4_interfaces) {
661 assert(ctx->num_entry_ifaces < ARRAY_SIZE(ctx->entry_ifaces));
662 ctx->entry_ifaces[ctx->num_entry_ifaces++] = ctx->scratch_block_var[idx];
663 }
664 }
665
666 static SpvId
get_scratch_block(struct ntv_context * ctx,unsigned bit_size)667 get_scratch_block(struct ntv_context *ctx, unsigned bit_size)
668 {
669 unsigned idx = bit_size >> 4;
670 if (!ctx->scratch_block_var[idx])
671 create_scratch_block(ctx, ctx->nir->scratch_size, bit_size);
672 return ctx->scratch_block_var[idx];
673 }
674
675 static void
create_shared_block(struct ntv_context * ctx,unsigned bit_size)676 create_shared_block(struct ntv_context *ctx, unsigned bit_size)
677 {
678 unsigned idx = bit_size >> 4;
679 SpvId type = spirv_builder_type_uint(&ctx->builder, bit_size);
680 SpvId array;
681
682 assert(gl_shader_stage_is_compute(ctx->nir->info.stage));
683 if (ctx->nir->info.cs.has_variable_shared_mem) {
684 assert(ctx->shared_mem_size);
685 SpvId const_shared_size = emit_uint_const(ctx, 32, ctx->nir->info.shared_size);
686 SpvId shared_mem_size = spirv_builder_emit_triop(&ctx->builder, SpvOpSpecConstantOp, spirv_builder_type_uint(&ctx->builder, 32), SpvOpIAdd, const_shared_size, ctx->shared_mem_size);
687 shared_mem_size = spirv_builder_emit_triop(&ctx->builder, SpvOpSpecConstantOp, spirv_builder_type_uint(&ctx->builder, 32), SpvOpUDiv, shared_mem_size, emit_uint_const(ctx, 32, bit_size / 8));
688 array = spirv_builder_type_array(&ctx->builder, type, shared_mem_size);
689 } else {
690 unsigned block_size = ctx->nir->info.shared_size / (bit_size / 8);
691 assert(block_size);
692 array = spirv_builder_type_array(&ctx->builder, type, emit_uint_const(ctx, 32, block_size));
693 }
694
695 ctx->shared_block_arr_type[idx] = array;
696 spirv_builder_emit_array_stride(&ctx->builder, array, bit_size / 8);
697
698 /* Create wrapper struct for Block, Offset and Aliased decorations. */
699 SpvId block = spirv_builder_type_struct(&ctx->builder, &array, 1);
700
701 SpvId ptr_type = spirv_builder_type_pointer(&ctx->builder,
702 SpvStorageClassWorkgroup,
703 block);
704 ctx->shared_block_var[idx] = spirv_builder_emit_var(&ctx->builder, ptr_type, SpvStorageClassWorkgroup);
705 if (ctx->spirv_1_4_interfaces) {
706 assert(ctx->num_entry_ifaces < ARRAY_SIZE(ctx->entry_ifaces));
707 ctx->entry_ifaces[ctx->num_entry_ifaces++] = ctx->shared_block_var[idx];
708 }
709 /* Alias our shared memory blocks */
710 if (ctx->sinfo->have_workgroup_memory_explicit_layout) {
711 spirv_builder_emit_member_offset(&ctx->builder, block, 0, 0);
712 spirv_builder_emit_decoration(&ctx->builder, block, SpvDecorationBlock);
713 spirv_builder_emit_decoration(&ctx->builder, ctx->shared_block_var[idx], SpvDecorationAliased);
714 }
715 }
716
717 static SpvId
get_shared_block(struct ntv_context * ctx,unsigned bit_size)718 get_shared_block(struct ntv_context *ctx, unsigned bit_size)
719 {
720 unsigned idx = bit_size >> 4;
721 if (!ctx->shared_block_var[idx])
722 create_shared_block(ctx, bit_size);
723 if (ctx->sinfo->have_workgroup_memory_explicit_layout) {
724 spirv_builder_emit_extension(&ctx->builder, "SPV_KHR_workgroup_memory_explicit_layout");
725 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityWorkgroupMemoryExplicitLayoutKHR);
726 if (ctx->shared_block_var[0])
727 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityWorkgroupMemoryExplicitLayout8BitAccessKHR);
728 if (ctx->shared_block_var[1])
729 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityWorkgroupMemoryExplicitLayout16BitAccessKHR);
730 }
731
732 SpvId ptr_type = spirv_builder_type_pointer(&ctx->builder,
733 SpvStorageClassWorkgroup,
734 ctx->shared_block_arr_type[idx]);
735 SpvId zero = emit_uint_const(ctx, 32, 0);
736
737 return spirv_builder_emit_access_chain(&ctx->builder, ptr_type,
738 ctx->shared_block_var[idx], &zero, 1);
739 }
740
741 #define HANDLE_EMIT_BUILTIN(SLOT, BUILTIN) \
742 case VARYING_SLOT_##SLOT: \
743 spirv_builder_emit_builtin(&ctx->builder, var_id, SpvBuiltIn##BUILTIN); \
744 break
745
746
747 static SpvId
input_var_init(struct ntv_context * ctx,struct nir_variable * var)748 input_var_init(struct ntv_context *ctx, struct nir_variable *var)
749 {
750 SpvId var_type = get_glsl_type(ctx, var->type);
751 SpvStorageClass sc = get_storage_class(var);
752 if (sc == SpvStorageClassPushConstant)
753 spirv_builder_emit_decoration(&ctx->builder, var_type, SpvDecorationBlock);
754 SpvId pointer_type = spirv_builder_type_pointer(&ctx->builder,
755 sc, var_type);
756 SpvId var_id = spirv_builder_emit_var(&ctx->builder, pointer_type, sc);
757
758 if (var->name)
759 spirv_builder_emit_name(&ctx->builder, var_id, var->name);
760
761 if (var->data.mode == nir_var_mem_push_const) {
762 ctx->push_const_var = var_id;
763
764 if (ctx->spirv_1_4_interfaces) {
765 assert(ctx->num_entry_ifaces < ARRAY_SIZE(ctx->entry_ifaces));
766 ctx->entry_ifaces[ctx->num_entry_ifaces++] = var_id;
767 }
768 }
769 return var_id;
770 }
771
772 static void
emit_interpolation(struct ntv_context * ctx,SpvId var_id,enum glsl_interp_mode mode)773 emit_interpolation(struct ntv_context *ctx, SpvId var_id,
774 enum glsl_interp_mode mode)
775 {
776 switch (mode) {
777 case INTERP_MODE_NONE:
778 case INTERP_MODE_SMOOTH:
779 /* XXX spirv doesn't seem to have anything for this */
780 break;
781 case INTERP_MODE_FLAT:
782 spirv_builder_emit_decoration(&ctx->builder, var_id,
783 SpvDecorationFlat);
784 break;
785 case INTERP_MODE_EXPLICIT:
786 spirv_builder_emit_decoration(&ctx->builder, var_id,
787 SpvDecorationExplicitInterpAMD);
788 break;
789 case INTERP_MODE_NOPERSPECTIVE:
790 spirv_builder_emit_decoration(&ctx->builder, var_id,
791 SpvDecorationNoPerspective);
792 break;
793 default:
794 unreachable("unknown interpolation value");
795 }
796 }
797
798 enum per_vertex_member {
799 PV_MEMBER_POSITION = 0,
800 PV_MEMBER_POINT_SIZE = 1,
801 PV_MEMBER_CLIP_DISTANCE = 2,
802 PV_MEMBER_CULL_DISTANCE = 3,
803 PV_MEMBER_COUNT = 4
804 };
805
806 static inline bool
stage_has_per_vertex_input(gl_shader_stage stage)807 stage_has_per_vertex_input(gl_shader_stage stage)
808 {
809 return stage == MESA_SHADER_TESS_CTRL ||
810 stage == MESA_SHADER_TESS_EVAL ||
811 stage == MESA_SHADER_GEOMETRY;
812 }
813
814 static inline bool
stage_has_per_vertex_output(gl_shader_stage stage)815 stage_has_per_vertex_output(gl_shader_stage stage)
816 {
817 return stage == MESA_SHADER_VERTEX ||
818 stage == MESA_SHADER_TESS_CTRL ||
819 stage == MESA_SHADER_TESS_EVAL ||
820 stage == MESA_SHADER_GEOMETRY;
821 }
822
823 static inline bool
stage_has_per_vertex(struct ntv_context * ctx)824 stage_has_per_vertex(struct ntv_context *ctx)
825 {
826 return ctx->use_gl_per_vertex &&
827 (stage_has_per_vertex_input(ctx->stage) ||
828 stage_has_per_vertex_output(ctx->stage));
829 }
830
831 static bool
is_per_vertex_builtin(struct ntv_context * ctx,nir_variable * var)832 is_per_vertex_builtin(struct ntv_context *ctx, nir_variable *var)
833 {
834 if (var->data.mode != nir_var_shader_in &&
835 var->data.mode != nir_var_shader_out)
836 return false;
837
838 if (var->data.mode == nir_var_shader_in &&
839 ctx->stage == MESA_SHADER_VERTEX) {
840 return false;
841 }
842
843 if (ctx->stage == MESA_SHADER_FRAGMENT)
844 return false;
845
846 switch (var->data.location) {
847 case VARYING_SLOT_POS:
848 case VARYING_SLOT_PSIZ:
849 case VARYING_SLOT_CLIP_DIST0:
850 case VARYING_SLOT_CULL_DIST0:
851 return true;
852 default:
853 return false;
854 }
855 }
856
857 static enum per_vertex_member
location_to_per_vertex_member(unsigned location)858 location_to_per_vertex_member(unsigned location)
859 {
860 switch (location) {
861 case VARYING_SLOT_POS:
862 return PV_MEMBER_POSITION;
863 case VARYING_SLOT_PSIZ:
864 return PV_MEMBER_POINT_SIZE;
865 case VARYING_SLOT_CLIP_DIST0:
866 return PV_MEMBER_CLIP_DISTANCE;
867 case VARYING_SLOT_CULL_DIST0:
868 return PV_MEMBER_CULL_DISTANCE;
869 default:
870 unreachable("not a per-vertex builtin");
871 }
872 }
873
874 static SpvId
get_per_vertex_member_type(struct ntv_context * ctx,uint32_t member_idx)875 get_per_vertex_member_type(struct ntv_context *ctx, uint32_t member_idx)
876 {
877 switch (member_idx) {
878 case PV_MEMBER_POSITION:
879 return get_glsl_type(ctx, glsl_vec4_type());
880 case PV_MEMBER_POINT_SIZE:
881 return get_glsl_type(ctx, glsl_float_type());
882 case PV_MEMBER_CLIP_DISTANCE:
883 case PV_MEMBER_CULL_DISTANCE:
884 // Use cached type or create it once
885 if (ctx->float_array_1_type == 0) {
886 SpvId float_type = get_glsl_type(ctx, glsl_float_type());
887 SpvId array_size = spirv_builder_const_uint(&ctx->builder, 32, 5);
888 ctx->float_array_1_type = spirv_builder_type_array(&ctx->builder, float_type, array_size);
889 }
890 return ctx->float_array_1_type;
891 default:
892 unreachable("invalid per-vertex member");
893 }
894 }
895
896 static SpvId
create_per_vertex_block_type(struct ntv_context * ctx,bool in)897 create_per_vertex_block_type(struct ntv_context *ctx, bool in)
898 {
899 int index = in ? 0 : 1;
900 if (ctx->per_vertex_block_type[index] != 0)
901 return ctx->per_vertex_block_type[index];
902
903 SpvId vec4_type = get_glsl_type(ctx, glsl_vec4_type());
904 SpvId float_type = get_glsl_type(ctx, glsl_float_type());
905 SpvId float_array_type = get_per_vertex_member_type(ctx, PV_MEMBER_CLIP_DISTANCE);
906
907 SpvId member_types[PV_MEMBER_COUNT] = {
908 [PV_MEMBER_POSITION] = vec4_type,
909 [PV_MEMBER_POINT_SIZE] = float_type,
910 [PV_MEMBER_CLIP_DISTANCE] = float_array_type,
911 [PV_MEMBER_CULL_DISTANCE] = float_array_type // Same type as clip distance
912 };
913
914 SpvId block_type = spirv_builder_type_struct(&ctx->builder, member_types, PV_MEMBER_COUNT);
915
916 spirv_builder_emit_member_name(&ctx->builder, block_type, PV_MEMBER_POSITION, "gl_Position");
917 spirv_builder_emit_member_builtin(&ctx->builder, block_type, PV_MEMBER_POSITION,
918 SpvBuiltInPosition);
919
920 spirv_builder_emit_member_name(&ctx->builder, block_type, PV_MEMBER_POINT_SIZE, "gl_PointSize");
921 spirv_builder_emit_member_builtin(&ctx->builder, block_type, PV_MEMBER_POINT_SIZE,
922 SpvBuiltInPointSize);
923
924 spirv_builder_emit_member_name(&ctx->builder, block_type, PV_MEMBER_CLIP_DISTANCE, "gl_ClipDistance");
925 spirv_builder_emit_member_builtin(&ctx->builder, block_type, PV_MEMBER_CLIP_DISTANCE,
926 SpvBuiltInClipDistance);
927
928 spirv_builder_emit_member_name(&ctx->builder, block_type, PV_MEMBER_CULL_DISTANCE, "gl_CullDistance");
929 spirv_builder_emit_member_builtin(&ctx->builder, block_type, PV_MEMBER_CULL_DISTANCE,
930 SpvBuiltInCullDistance);
931
932 spirv_builder_emit_name(&ctx->builder, block_type, in ? "gl_PerVertex_0" : "gl_PerVertex");
933 spirv_builder_emit_decoration(&ctx->builder, block_type, SpvDecorationBlock);
934
935 ctx->per_vertex_block_type[index] = block_type;
936
937 return block_type;
938 }
939
940 static void
emit_input(struct ntv_context * ctx,struct nir_variable * var)941 emit_input(struct ntv_context *ctx, struct nir_variable *var)
942 {
943 if (stage_has_per_vertex(ctx) && is_per_vertex_builtin(ctx, var)) {
944 if (glsl_type_is_array(var->type))
945 ctx->gl_in_num_vertices[ctx->stage - 1] = glsl_get_length(var->type);
946 return;
947 }
948
949 SpvId var_id = input_var_init(ctx, var);
950 if (ctx->stage == MESA_SHADER_VERTEX)
951 spirv_builder_emit_location(&ctx->builder, var_id,
952 var->data.driver_location);
953 else if (ctx->stage == MESA_SHADER_FRAGMENT) {
954 switch (var->data.location) {
955 HANDLE_EMIT_BUILTIN(POS, FragCoord);
956 HANDLE_EMIT_BUILTIN(LAYER, Layer);
957 HANDLE_EMIT_BUILTIN(PRIMITIVE_ID, PrimitiveId);
958 HANDLE_EMIT_BUILTIN(CLIP_DIST0, ClipDistance);
959 HANDLE_EMIT_BUILTIN(CULL_DIST0, CullDistance);
960 HANDLE_EMIT_BUILTIN(VIEWPORT, ViewportIndex);
961 HANDLE_EMIT_BUILTIN(FACE, FrontFacing);
962
963 default:
964 spirv_builder_emit_location(&ctx->builder, var_id,
965 var->data.driver_location);
966 }
967 if (var->data.centroid)
968 spirv_builder_emit_decoration(&ctx->builder, var_id, SpvDecorationCentroid);
969 else if (var->data.sample)
970 spirv_builder_emit_decoration(&ctx->builder, var_id, SpvDecorationSample);
971 emit_interpolation(ctx, var_id, var->data.interpolation);
972 } else if (ctx->stage < MESA_SHADER_FRAGMENT) {
973 switch (var->data.location) {
974 HANDLE_EMIT_BUILTIN(POS, Position);
975 HANDLE_EMIT_BUILTIN(PSIZ, PointSize);
976 HANDLE_EMIT_BUILTIN(LAYER, Layer);
977 HANDLE_EMIT_BUILTIN(PRIMITIVE_ID, PrimitiveId);
978 HANDLE_EMIT_BUILTIN(CULL_DIST0, CullDistance);
979 HANDLE_EMIT_BUILTIN(VIEWPORT, ViewportIndex);
980 HANDLE_EMIT_BUILTIN(TESS_LEVEL_OUTER, TessLevelOuter);
981 HANDLE_EMIT_BUILTIN(TESS_LEVEL_INNER, TessLevelInner);
982
983 case VARYING_SLOT_CLIP_DIST0:
984 assert(glsl_type_is_array(var->type));
985 spirv_builder_emit_builtin(&ctx->builder, var_id, SpvBuiltInClipDistance);
986 break;
987
988 default:
989 spirv_builder_emit_location(&ctx->builder, var_id,
990 var->data.driver_location);
991 }
992 }
993
994 if (var->data.location_frac)
995 spirv_builder_emit_component(&ctx->builder, var_id,
996 var->data.location_frac);
997
998 if (var->data.patch)
999 spirv_builder_emit_decoration(&ctx->builder, var_id, SpvDecorationPatch);
1000
1001 _mesa_hash_table_insert(ctx->vars, var, (void *)(intptr_t)var_id);
1002
1003 assert(ctx->num_entry_ifaces < ARRAY_SIZE(ctx->entry_ifaces));
1004 ctx->entry_ifaces[ctx->num_entry_ifaces++] = var_id;
1005 }
1006
1007 static void
emit_output(struct ntv_context * ctx,struct nir_variable * var)1008 emit_output(struct ntv_context *ctx, struct nir_variable *var)
1009 {
1010 if (stage_has_per_vertex(ctx) && is_per_vertex_builtin(ctx, var)) {
1011 if (glsl_type_is_array(var->type))
1012 ctx->gl_out_num_vertices = glsl_get_length(var->type);
1013 return;
1014 }
1015
1016 SpvId var_type = get_glsl_type(ctx, var->type);
1017
1018 /* SampleMask is always an array in spirv */
1019 if (ctx->stage == MESA_SHADER_FRAGMENT && var->data.location == FRAG_RESULT_SAMPLE_MASK)
1020 ctx->sample_mask_type = var_type = spirv_builder_type_array(&ctx->builder, var_type, emit_uint_const(ctx, 32, 1));
1021 SpvId pointer_type = spirv_builder_type_pointer(&ctx->builder,
1022 SpvStorageClassOutput,
1023 var_type);
1024 SpvId var_id = spirv_builder_emit_var(&ctx->builder, pointer_type,
1025 SpvStorageClassOutput);
1026 if (var->name)
1027 spirv_builder_emit_name(&ctx->builder, var_id, var->name);
1028
1029 if (var->data.precision == GLSL_PRECISION_MEDIUM || var->data.precision == GLSL_PRECISION_LOW) {
1030 spirv_builder_emit_decoration(&ctx->builder, var_id,
1031 SpvDecorationRelaxedPrecision);
1032 }
1033
1034 if (ctx->stage != MESA_SHADER_FRAGMENT) {
1035 switch (var->data.location) {
1036 HANDLE_EMIT_BUILTIN(POS, Position);
1037 HANDLE_EMIT_BUILTIN(PSIZ, PointSize);
1038 HANDLE_EMIT_BUILTIN(LAYER, Layer);
1039 HANDLE_EMIT_BUILTIN(PRIMITIVE_ID, PrimitiveId);
1040 HANDLE_EMIT_BUILTIN(CLIP_DIST0, ClipDistance);
1041 HANDLE_EMIT_BUILTIN(CULL_DIST0, CullDistance);
1042 HANDLE_EMIT_BUILTIN(VIEWPORT, ViewportIndex);
1043 HANDLE_EMIT_BUILTIN(TESS_LEVEL_OUTER, TessLevelOuter);
1044 HANDLE_EMIT_BUILTIN(TESS_LEVEL_INNER, TessLevelInner);
1045
1046 default:
1047 /* non-xfb psiz output will have location -1 */
1048 if (var->data.location >= 0)
1049 spirv_builder_emit_location(&ctx->builder, var_id,
1050 var->data.driver_location);
1051 }
1052 emit_interpolation(ctx, var_id, var->data.interpolation);
1053 } else {
1054 if (var->data.location >= FRAG_RESULT_DATA0) {
1055 spirv_builder_emit_location(&ctx->builder, var_id,
1056 var->data.location - FRAG_RESULT_DATA0);
1057 spirv_builder_emit_index(&ctx->builder, var_id, var->data.index);
1058 } else {
1059 switch (var->data.location) {
1060 case FRAG_RESULT_COLOR:
1061 unreachable("gl_FragColor should be lowered by now");
1062
1063 case FRAG_RESULT_DEPTH:
1064 spirv_builder_emit_builtin(&ctx->builder, var_id, SpvBuiltInFragDepth);
1065 break;
1066
1067 case FRAG_RESULT_SAMPLE_MASK:
1068 spirv_builder_emit_builtin(&ctx->builder, var_id, SpvBuiltInSampleMask);
1069 break;
1070
1071 case FRAG_RESULT_STENCIL:
1072 spirv_builder_emit_builtin(&ctx->builder, var_id, SpvBuiltInFragStencilRefEXT);
1073 break;
1074
1075 default:
1076 spirv_builder_emit_location(&ctx->builder, var_id,
1077 var->data.location);
1078 spirv_builder_emit_index(&ctx->builder, var_id, var->data.index);
1079 }
1080 }
1081 if (var->data.sample)
1082 spirv_builder_emit_decoration(&ctx->builder, var_id, SpvDecorationSample);
1083 }
1084
1085 if (var->data.location_frac)
1086 spirv_builder_emit_component(&ctx->builder, var_id,
1087 var->data.location_frac);
1088
1089 if (var->data.patch)
1090 spirv_builder_emit_decoration(&ctx->builder, var_id, SpvDecorationPatch);
1091
1092 if (var->data.explicit_xfb_buffer && ctx->nir->xfb_info) {
1093 spirv_builder_emit_offset(&ctx->builder, var_id, var->data.offset);
1094 spirv_builder_emit_xfb_buffer(&ctx->builder, var_id, var->data.xfb.buffer);
1095 spirv_builder_emit_xfb_stride(&ctx->builder, var_id, var->data.xfb.stride);
1096 if (var->data.stream)
1097 spirv_builder_emit_stream(&ctx->builder, var_id, var->data.stream);
1098 }
1099
1100 _mesa_hash_table_insert(ctx->vars, var, (void *)(intptr_t)var_id);
1101
1102 assert(ctx->num_entry_ifaces < ARRAY_SIZE(ctx->entry_ifaces));
1103 ctx->entry_ifaces[ctx->num_entry_ifaces++] = var_id;
1104 }
1105
1106 static void
emit_shader_temp(struct ntv_context * ctx,struct nir_variable * var)1107 emit_shader_temp(struct ntv_context *ctx, struct nir_variable *var)
1108 {
1109 SpvId var_type = get_glsl_type(ctx, var->type);
1110
1111 SpvId pointer_type = spirv_builder_type_pointer(&ctx->builder,
1112 SpvStorageClassPrivate,
1113 var_type);
1114 SpvId var_id = spirv_builder_emit_var(&ctx->builder, pointer_type,
1115 SpvStorageClassPrivate);
1116 if (var->name)
1117 spirv_builder_emit_name(&ctx->builder, var_id, var->name);
1118
1119 _mesa_hash_table_insert(ctx->vars, var, (void *)(intptr_t)var_id);
1120
1121 assert(ctx->num_entry_ifaces < ARRAY_SIZE(ctx->entry_ifaces));
1122 ctx->entry_ifaces[ctx->num_entry_ifaces++] = var_id;
1123 }
1124
1125 static void
emit_temp(struct ntv_context * ctx,struct nir_variable * var)1126 emit_temp(struct ntv_context *ctx, struct nir_variable *var)
1127 {
1128 SpvId var_type = get_glsl_type(ctx, var->type);
1129
1130 SpvId pointer_type = spirv_builder_type_pointer(&ctx->builder,
1131 SpvStorageClassFunction,
1132 var_type);
1133 SpvId var_id = spirv_builder_emit_var(&ctx->builder, pointer_type,
1134 SpvStorageClassFunction);
1135 if (var->name)
1136 spirv_builder_emit_name(&ctx->builder, var_id, var->name);
1137
1138 _mesa_hash_table_insert(ctx->vars, var, (void *)(intptr_t)var_id);
1139 }
1140
1141 static SpvDim
type_to_dim(enum glsl_sampler_dim gdim,bool * is_ms)1142 type_to_dim(enum glsl_sampler_dim gdim, bool *is_ms)
1143 {
1144 *is_ms = false;
1145 switch (gdim) {
1146 case GLSL_SAMPLER_DIM_1D:
1147 return SpvDim1D;
1148 case GLSL_SAMPLER_DIM_2D:
1149 return SpvDim2D;
1150 case GLSL_SAMPLER_DIM_3D:
1151 return SpvDim3D;
1152 case GLSL_SAMPLER_DIM_CUBE:
1153 return SpvDimCube;
1154 case GLSL_SAMPLER_DIM_RECT:
1155 return SpvDim2D;
1156 case GLSL_SAMPLER_DIM_BUF:
1157 return SpvDimBuffer;
1158 case GLSL_SAMPLER_DIM_EXTERNAL:
1159 return SpvDim2D; /* seems dodgy... */
1160 case GLSL_SAMPLER_DIM_MS:
1161 *is_ms = true;
1162 return SpvDim2D;
1163 case GLSL_SAMPLER_DIM_SUBPASS_MS:
1164 *is_ms = true;
1165 return SpvDimSubpassData;
1166 case GLSL_SAMPLER_DIM_SUBPASS:
1167 return SpvDimSubpassData;
1168 default:
1169 fprintf(stderr, "unknown sampler type %d\n", gdim);
1170 break;
1171 }
1172 return SpvDim2D;
1173 }
1174
1175 static inline SpvImageFormat
get_shader_image_format(enum pipe_format format)1176 get_shader_image_format(enum pipe_format format)
1177 {
1178 switch (format) {
1179 case PIPE_FORMAT_R32G32B32A32_FLOAT:
1180 return SpvImageFormatRgba32f;
1181 case PIPE_FORMAT_R16G16B16A16_FLOAT:
1182 return SpvImageFormatRgba16f;
1183 case PIPE_FORMAT_R32_FLOAT:
1184 return SpvImageFormatR32f;
1185 case PIPE_FORMAT_R8G8B8A8_UNORM:
1186 return SpvImageFormatRgba8;
1187 case PIPE_FORMAT_R8G8B8A8_SNORM:
1188 return SpvImageFormatRgba8Snorm;
1189 case PIPE_FORMAT_R32G32B32A32_SINT:
1190 return SpvImageFormatRgba32i;
1191 case PIPE_FORMAT_R16G16B16A16_SINT:
1192 return SpvImageFormatRgba16i;
1193 case PIPE_FORMAT_R8G8B8A8_SINT:
1194 return SpvImageFormatRgba8i;
1195 case PIPE_FORMAT_R32_SINT:
1196 return SpvImageFormatR32i;
1197 case PIPE_FORMAT_R32G32B32A32_UINT:
1198 return SpvImageFormatRgba32ui;
1199 case PIPE_FORMAT_R16G16B16A16_UINT:
1200 return SpvImageFormatRgba16ui;
1201 case PIPE_FORMAT_R8G8B8A8_UINT:
1202 return SpvImageFormatRgba8ui;
1203 case PIPE_FORMAT_R32_UINT:
1204 return SpvImageFormatR32ui;
1205 default:
1206 return SpvImageFormatUnknown;
1207 }
1208 }
1209
1210 static inline SpvImageFormat
get_extended_image_format(enum pipe_format format)1211 get_extended_image_format(enum pipe_format format)
1212 {
1213 switch (format) {
1214 case PIPE_FORMAT_R32G32_FLOAT:
1215 return SpvImageFormatRg32f;
1216 case PIPE_FORMAT_R16G16_FLOAT:
1217 return SpvImageFormatRg16f;
1218 case PIPE_FORMAT_R11G11B10_FLOAT:
1219 return SpvImageFormatR11fG11fB10f;
1220 case PIPE_FORMAT_R16_FLOAT:
1221 return SpvImageFormatR16f;
1222 case PIPE_FORMAT_R16G16B16A16_UNORM:
1223 return SpvImageFormatRgba16;
1224 case PIPE_FORMAT_R10G10B10A2_UNORM:
1225 return SpvImageFormatRgb10A2;
1226 case PIPE_FORMAT_R16G16_UNORM:
1227 return SpvImageFormatRg16;
1228 case PIPE_FORMAT_R8G8_UNORM:
1229 return SpvImageFormatRg8;
1230 case PIPE_FORMAT_R16_UNORM:
1231 return SpvImageFormatR16;
1232 case PIPE_FORMAT_R8_UNORM:
1233 return SpvImageFormatR8;
1234 case PIPE_FORMAT_R16G16B16A16_SNORM:
1235 return SpvImageFormatRgba16Snorm;
1236 case PIPE_FORMAT_R16G16_SNORM:
1237 return SpvImageFormatRg16Snorm;
1238 case PIPE_FORMAT_R8G8_SNORM:
1239 return SpvImageFormatRg8Snorm;
1240 case PIPE_FORMAT_R16_SNORM:
1241 return SpvImageFormatR16Snorm;
1242 case PIPE_FORMAT_R8_SNORM:
1243 return SpvImageFormatR8Snorm;
1244 case PIPE_FORMAT_R32G32_SINT:
1245 return SpvImageFormatRg32i;
1246 case PIPE_FORMAT_R16G16_SINT:
1247 return SpvImageFormatRg16i;
1248 case PIPE_FORMAT_R8G8_SINT:
1249 return SpvImageFormatRg8i;
1250 case PIPE_FORMAT_R16_SINT:
1251 return SpvImageFormatR16i;
1252 case PIPE_FORMAT_R8_SINT:
1253 return SpvImageFormatR8i;
1254 case PIPE_FORMAT_R10G10B10A2_UINT:
1255 return SpvImageFormatRgb10a2ui;
1256 case PIPE_FORMAT_R32G32_UINT:
1257 return SpvImageFormatRg32ui;
1258 case PIPE_FORMAT_R16G16_UINT:
1259 return SpvImageFormatRg16ui;
1260 case PIPE_FORMAT_R8G8_UINT:
1261 return SpvImageFormatRg8ui;
1262 case PIPE_FORMAT_R16_UINT:
1263 return SpvImageFormatR16ui;
1264 case PIPE_FORMAT_R8_UINT:
1265 return SpvImageFormatR8ui;
1266
1267 default:
1268 return SpvImageFormatUnknown;
1269 }
1270 }
1271
1272 static inline SpvImageFormat
get_image_format(struct ntv_context * ctx,enum pipe_format format)1273 get_image_format(struct ntv_context *ctx, enum pipe_format format)
1274 {
1275 /* always supported */
1276 if (format == PIPE_FORMAT_NONE)
1277 return SpvImageFormatUnknown;
1278
1279 SpvImageFormat ret = get_shader_image_format(format);
1280 if (ret != SpvImageFormatUnknown) {
1281 /* requires the shader-cap, but we already emit that */
1282 return ret;
1283 }
1284
1285 ret = get_extended_image_format(format);
1286 assert(ret != SpvImageFormatUnknown);
1287 spirv_builder_emit_cap(&ctx->builder,
1288 SpvCapabilityStorageImageExtendedFormats);
1289 return ret;
1290 }
1291
1292 static SpvId
get_bare_image_type(struct ntv_context * ctx,struct nir_variable * var,bool is_sampler)1293 get_bare_image_type(struct ntv_context *ctx, struct nir_variable *var, bool is_sampler)
1294 {
1295 const struct glsl_type *type = glsl_without_array(var->type);
1296
1297 bool is_ms;
1298
1299 if (var->data.fb_fetch_output) {
1300 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityInputAttachment);
1301 } else if (!is_sampler && !var->data.image.format) {
1302 if (!(var->data.access & ACCESS_NON_WRITEABLE))
1303 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityStorageImageWriteWithoutFormat);
1304 if (!(var->data.access & ACCESS_NON_READABLE))
1305 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityStorageImageReadWithoutFormat);
1306 }
1307
1308 SpvDim dimension = type_to_dim(glsl_get_sampler_dim(type), &is_ms);
1309 if (dimension == SpvDim1D) {
1310 if (is_sampler)
1311 spirv_builder_emit_cap(&ctx->builder, SpvCapabilitySampled1D);
1312 else
1313 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityImage1D);
1314 }
1315 if (dimension == SpvDimBuffer) {
1316 if (is_sampler)
1317 spirv_builder_emit_cap(&ctx->builder, SpvCapabilitySampledBuffer);
1318 else
1319 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityImageBuffer);
1320 }
1321
1322 bool arrayed = glsl_sampler_type_is_array(type);
1323 if (dimension == SpvDimCube && arrayed)
1324 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityImageCubeArray);
1325 if (arrayed && !is_sampler && is_ms)
1326 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityImageMSArray);
1327
1328 SpvId result_type = get_glsl_basetype(ctx, glsl_get_sampler_result_type(type));
1329 return spirv_builder_type_image(&ctx->builder, result_type,
1330 dimension, false,
1331 arrayed,
1332 is_ms, is_sampler ? 1 : 2,
1333 get_image_format(ctx, var->data.image.format));
1334 }
1335
1336 static SpvId
get_image_type(struct ntv_context * ctx,struct nir_variable * var,bool is_sampler,bool is_buffer)1337 get_image_type(struct ntv_context *ctx, struct nir_variable *var,
1338 bool is_sampler, bool is_buffer)
1339 {
1340 SpvId image_type = get_bare_image_type(ctx, var, is_sampler);
1341 return is_sampler && ctx->stage != MESA_SHADER_KERNEL && !is_buffer ?
1342 spirv_builder_type_sampled_image(&ctx->builder, image_type) :
1343 image_type;
1344 }
1345
1346 static SpvId
emit_image(struct ntv_context * ctx,struct nir_variable * var,SpvId image_type)1347 emit_image(struct ntv_context *ctx, struct nir_variable *var, SpvId image_type)
1348 {
1349 if (var->data.bindless)
1350 return 0;
1351 const struct glsl_type *type = glsl_without_array(var->type);
1352
1353 bool is_sampler = glsl_type_is_sampler(type);
1354 bool is_buffer = glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_BUF;
1355 SpvId var_type = is_sampler && ctx->stage != MESA_SHADER_KERNEL && !is_buffer ?
1356 spirv_builder_type_sampled_image(&ctx->builder, image_type) : image_type;
1357
1358 bool mediump = (var->data.precision == GLSL_PRECISION_MEDIUM || var->data.precision == GLSL_PRECISION_LOW);
1359
1360 int index = var->data.driver_location;
1361 assert(!find_image_type(ctx, var));
1362
1363 if (glsl_type_is_array(var->type)) {
1364 var_type = spirv_builder_type_array(&ctx->builder, var_type,
1365 emit_uint_const(ctx, 32, glsl_get_aoa_size(var->type)));
1366 spirv_builder_emit_array_stride(&ctx->builder, var_type, sizeof(void*));
1367 }
1368 SpvId pointer_type = spirv_builder_type_pointer(&ctx->builder,
1369 SpvStorageClassUniformConstant,
1370 var_type);
1371
1372 SpvId var_id = spirv_builder_emit_var(&ctx->builder, pointer_type,
1373 SpvStorageClassUniformConstant);
1374
1375 if (mediump) {
1376 spirv_builder_emit_decoration(&ctx->builder, var_id,
1377 SpvDecorationRelaxedPrecision);
1378 }
1379
1380 if (var->name)
1381 spirv_builder_emit_name(&ctx->builder, var_id, var->name);
1382
1383 if (var->data.fb_fetch_output)
1384 spirv_builder_emit_input_attachment_index(&ctx->builder, var_id, var->data.index);
1385
1386 _mesa_hash_table_insert(ctx->vars, var, (void *)(intptr_t)var_id);
1387 if (is_sampler) {
1388 if (var->data.descriptor_set == ctx->bindless_set_idx) {
1389 assert(!ctx->bindless_samplers[index]);
1390 ctx->bindless_samplers[index] = var_id;
1391 } else {
1392 assert(!ctx->samplers[index]);
1393 ctx->samplers[index] = var_id;
1394 }
1395 } else {
1396 assert(!ctx->images[index]);
1397 ctx->images[index] = var_id;
1398 emit_access_decorations(ctx, var, var_id);
1399 }
1400 _mesa_hash_table_insert(&ctx->image_types, var, (void *)(intptr_t)image_type);
1401 if (ctx->spirv_1_4_interfaces) {
1402 assert(ctx->num_entry_ifaces < ARRAY_SIZE(ctx->entry_ifaces));
1403 ctx->entry_ifaces[ctx->num_entry_ifaces++] = var_id;
1404 }
1405
1406 spirv_builder_emit_descriptor_set(&ctx->builder, var_id, var->data.descriptor_set);
1407 spirv_builder_emit_binding(&ctx->builder, var_id, var->data.binding);
1408 return var_id;
1409 }
1410
1411 static void
emit_sampler(struct ntv_context * ctx,nir_variable * var)1412 emit_sampler(struct ntv_context *ctx, nir_variable *var)
1413 {
1414 SpvId type = spirv_builder_type_sampler(&ctx->builder);
1415 SpvId pointer_type = spirv_builder_type_pointer(&ctx->builder,
1416 SpvStorageClassUniformConstant,
1417 type);
1418
1419 SpvId var_id = spirv_builder_emit_var(&ctx->builder, pointer_type,
1420 SpvStorageClassUniformConstant);
1421 char buf[128];
1422 snprintf(buf, sizeof(buf), "sampler_%u", var->data.driver_location);
1423 spirv_builder_emit_name(&ctx->builder, var_id, buf);
1424 spirv_builder_emit_descriptor_set(&ctx->builder, var_id, var->data.descriptor_set);
1425 spirv_builder_emit_binding(&ctx->builder, var_id, var->data.driver_location);
1426 _mesa_hash_table_insert(ctx->vars, var, (void *)(intptr_t)var_id);
1427 if (ctx->spirv_1_4_interfaces) {
1428 assert(ctx->num_entry_ifaces < ARRAY_SIZE(ctx->entry_ifaces));
1429 ctx->entry_ifaces[ctx->num_entry_ifaces++] = var_id;
1430 }
1431 }
1432
1433 static SpvId
get_sized_uint_array_type(struct ntv_context * ctx,unsigned array_size,unsigned bitsize)1434 get_sized_uint_array_type(struct ntv_context *ctx, unsigned array_size, unsigned bitsize)
1435 {
1436 SpvId array_length = emit_uint_const(ctx, 32, array_size);
1437 SpvId array_type = spirv_builder_type_array(&ctx->builder, get_uvec_type(ctx, bitsize, 1),
1438 array_length);
1439 spirv_builder_emit_array_stride(&ctx->builder, array_type, bitsize / 8);
1440 return array_type;
1441 }
1442
1443 /* get array<struct(array_type <--this one)> */
1444 static SpvId
get_bo_array_type(struct ntv_context * ctx,struct nir_variable * var)1445 get_bo_array_type(struct ntv_context *ctx, struct nir_variable *var)
1446 {
1447 struct hash_entry *he = _mesa_hash_table_search(ctx->bo_array_types, var);
1448 if (he)
1449 return (SpvId)(uintptr_t)he->data;
1450 unsigned bitsize = glsl_get_bit_size(glsl_get_array_element(glsl_get_struct_field(glsl_without_array(var->type), 0)));
1451 assert(bitsize);
1452 SpvId array_type;
1453 const struct glsl_type *type = glsl_without_array(var->type);
1454 const struct glsl_type *first_type = glsl_get_struct_field(type, 0);
1455 if (!glsl_type_is_unsized_array(first_type)) {
1456 uint32_t array_size = glsl_get_length(first_type);
1457 assert(array_size);
1458 return get_sized_uint_array_type(ctx, array_size, bitsize);
1459 }
1460 SpvId uint_type = spirv_builder_type_uint(&ctx->builder, bitsize);
1461 array_type = spirv_builder_type_runtime_array(&ctx->builder, uint_type);
1462 spirv_builder_emit_array_stride(&ctx->builder, array_type, bitsize / 8);
1463 return array_type;
1464 }
1465
1466 /* get array<struct(array_type) <--this one> */
1467 static SpvId
get_bo_struct_type(struct ntv_context * ctx,struct nir_variable * var)1468 get_bo_struct_type(struct ntv_context *ctx, struct nir_variable *var)
1469 {
1470 struct hash_entry *he = _mesa_hash_table_search(ctx->bo_struct_types, var);
1471 if (he)
1472 return (SpvId)(uintptr_t)he->data;
1473 const struct glsl_type *bare_type = glsl_without_array(var->type);
1474 unsigned bitsize = glsl_get_bit_size(glsl_get_array_element(glsl_get_struct_field(bare_type, 0)));
1475 SpvId array_type = get_bo_array_type(ctx, var);
1476 _mesa_hash_table_insert(ctx->bo_array_types, var, (void *)(uintptr_t)array_type);
1477 bool ssbo = var->data.mode == nir_var_mem_ssbo;
1478
1479 // wrap UBO-array in a struct
1480 SpvId runtime_array = 0;
1481 if (ssbo && glsl_get_length(bare_type) > 1) {
1482 const struct glsl_type *last_member = glsl_get_struct_field(bare_type, glsl_get_length(bare_type) - 1);
1483 if (glsl_type_is_unsized_array(last_member)) {
1484 runtime_array = spirv_builder_type_runtime_array(&ctx->builder, get_uvec_type(ctx, bitsize, 1));
1485 spirv_builder_emit_array_stride(&ctx->builder, runtime_array, glsl_get_explicit_stride(last_member));
1486 }
1487 }
1488 SpvId types[] = {array_type, runtime_array};
1489 SpvId struct_type = spirv_builder_type_struct(&ctx->builder, types, 1 + !!runtime_array);
1490 if (var->name) {
1491 char struct_name[100];
1492 snprintf(struct_name, sizeof(struct_name), "struct_%s", var->name);
1493 spirv_builder_emit_name(&ctx->builder, struct_type, struct_name);
1494 }
1495
1496 spirv_builder_emit_decoration(&ctx->builder, struct_type,
1497 SpvDecorationBlock);
1498 spirv_builder_emit_member_offset(&ctx->builder, struct_type, 0, 0);
1499 if (runtime_array)
1500 spirv_builder_emit_member_offset(&ctx->builder, struct_type, 1, 0);
1501
1502 return struct_type;
1503 }
1504
1505 static void
emit_bo(struct ntv_context * ctx,struct nir_variable * var,bool aliased)1506 emit_bo(struct ntv_context *ctx, struct nir_variable *var, bool aliased)
1507 {
1508 unsigned bitsize = glsl_get_bit_size(glsl_get_array_element(glsl_get_struct_field(glsl_without_array(var->type), 0)));
1509 bool ssbo = var->data.mode == nir_var_mem_ssbo;
1510 SpvId struct_type = get_bo_struct_type(ctx, var);
1511 _mesa_hash_table_insert(ctx->bo_struct_types, var, (void *)(uintptr_t)struct_type);
1512 SpvId array_length = emit_uint_const(ctx, 32, glsl_get_length(var->type));
1513 SpvId array_type = spirv_builder_type_array(&ctx->builder, struct_type, array_length);
1514 SpvId pointer_type = spirv_builder_type_pointer(&ctx->builder,
1515 ssbo ? SpvStorageClassStorageBuffer : SpvStorageClassUniform,
1516 array_type);
1517 SpvId var_id = spirv_builder_emit_var(&ctx->builder, pointer_type,
1518 ssbo ? SpvStorageClassStorageBuffer : SpvStorageClassUniform);
1519 if (var->name)
1520 spirv_builder_emit_name(&ctx->builder, var_id, var->name);
1521
1522 if (aliased)
1523 spirv_builder_emit_decoration(&ctx->builder, var_id, SpvDecorationAliased);
1524
1525 unsigned idx = bitsize >> 4;
1526 assert(idx < ARRAY_SIZE(ctx->ssbos));
1527 if (ssbo) {
1528 assert(!ctx->ssbos[idx]);
1529 ctx->ssbos[idx] = var_id;
1530 if (bitsize == 32)
1531 ctx->ssbo_vars = var;
1532 } else {
1533 assert(!ctx->ubos[var->data.driver_location][idx]);
1534 ctx->ubos[var->data.driver_location][idx] = var_id;
1535 ctx->ubo_vars[var->data.driver_location] = var;
1536 }
1537 if (ctx->spirv_1_4_interfaces) {
1538 assert(ctx->num_entry_ifaces < ARRAY_SIZE(ctx->entry_ifaces));
1539 ctx->entry_ifaces[ctx->num_entry_ifaces++] = var_id;
1540 }
1541 _mesa_hash_table_insert(ctx->vars, var, (void *)(intptr_t)var_id);
1542
1543 spirv_builder_emit_descriptor_set(&ctx->builder, var_id, var->data.descriptor_set);
1544 spirv_builder_emit_binding(&ctx->builder, var_id, var->data.binding);
1545 }
1546
1547 static SpvId
get_vec_from_bit_size(struct ntv_context * ctx,uint32_t bit_size,uint32_t num_components)1548 get_vec_from_bit_size(struct ntv_context *ctx, uint32_t bit_size, uint32_t num_components)
1549 {
1550 if (bit_size == 1)
1551 return get_bvec_type(ctx, num_components);
1552 return get_uvec_type(ctx, bit_size, num_components);
1553 }
1554
1555 static SpvId
get_src_ssa(struct ntv_context * ctx,const nir_def * ssa,nir_alu_type * atype)1556 get_src_ssa(struct ntv_context *ctx, const nir_def *ssa, nir_alu_type *atype)
1557 {
1558 assert(ssa->index < ctx->num_defs);
1559 assert(ctx->defs[ssa->index] != 0);
1560 *atype = ctx->def_types[ssa->index];
1561 return ctx->defs[ssa->index];
1562 }
1563
1564 static void
init_reg(struct ntv_context * ctx,nir_intrinsic_instr * decl,nir_alu_type atype)1565 init_reg(struct ntv_context *ctx, nir_intrinsic_instr *decl, nir_alu_type atype)
1566 {
1567 unsigned index = decl->def.index;
1568 unsigned num_components = nir_intrinsic_num_components(decl);
1569 unsigned bit_size = nir_intrinsic_bit_size(decl);
1570
1571 if (ctx->defs[index])
1572 return;
1573
1574 SpvId type = get_alu_type(ctx, atype, num_components, bit_size);
1575 SpvId pointer_type = spirv_builder_type_pointer(&ctx->builder,
1576 SpvStorageClassFunction,
1577 type);
1578 SpvId var = spirv_builder_emit_var(&ctx->builder, pointer_type,
1579 SpvStorageClassFunction);
1580
1581 ctx->defs[index] = var;
1582 ctx->def_types[index] = nir_alu_type_get_base_type(atype);
1583 }
1584
1585 static SpvId
get_src(struct ntv_context * ctx,nir_src * src,nir_alu_type * atype)1586 get_src(struct ntv_context *ctx, nir_src *src, nir_alu_type *atype)
1587 {
1588 return get_src_ssa(ctx, src->ssa, atype);
1589 }
1590
1591 static SpvId
get_alu_src_raw(struct ntv_context * ctx,nir_alu_instr * alu,unsigned src,nir_alu_type * atype)1592 get_alu_src_raw(struct ntv_context *ctx, nir_alu_instr *alu, unsigned src, nir_alu_type *atype)
1593 {
1594 SpvId def = get_src(ctx, &alu->src[src].src, atype);
1595
1596 unsigned used_channels = 0;
1597 bool need_swizzle = false;
1598 for (unsigned i = 0; i < NIR_MAX_VEC_COMPONENTS; i++) {
1599 if (!nir_alu_instr_channel_used(alu, src, i))
1600 continue;
1601
1602 used_channels++;
1603
1604 if (alu->src[src].swizzle[i] != i)
1605 need_swizzle = true;
1606 }
1607 assert(used_channels != 0);
1608
1609 unsigned live_channels = nir_src_num_components(alu->src[src].src);
1610 if (used_channels != live_channels)
1611 need_swizzle = true;
1612
1613 if (!need_swizzle)
1614 return def;
1615
1616 int bit_size = nir_src_bit_size(alu->src[src].src);
1617 SpvId raw_type = get_alu_type(ctx, *atype, 1, bit_size);
1618
1619 if (used_channels == 1) {
1620 uint32_t indices[] = { alu->src[src].swizzle[0] };
1621 return spirv_builder_emit_composite_extract(&ctx->builder, raw_type,
1622 def, indices,
1623 ARRAY_SIZE(indices));
1624 } else if (live_channels == 1) {
1625 SpvId raw_vec_type = spirv_builder_type_vector(&ctx->builder,
1626 raw_type,
1627 used_channels);
1628
1629 SpvId constituents[NIR_MAX_VEC_COMPONENTS] = {0};
1630 for (unsigned i = 0; i < used_channels; ++i)
1631 constituents[i] = def;
1632
1633 return spirv_builder_emit_composite_construct(&ctx->builder,
1634 raw_vec_type,
1635 constituents,
1636 used_channels);
1637 } else {
1638 SpvId raw_vec_type = spirv_builder_type_vector(&ctx->builder,
1639 raw_type,
1640 used_channels);
1641
1642 uint32_t components[NIR_MAX_VEC_COMPONENTS] = {0};
1643 size_t num_components = 0;
1644 for (unsigned i = 0; i < NIR_MAX_VEC_COMPONENTS; i++) {
1645 if (!nir_alu_instr_channel_used(alu, src, i))
1646 continue;
1647
1648 components[num_components++] = alu->src[src].swizzle[i];
1649 }
1650
1651 return spirv_builder_emit_vector_shuffle(&ctx->builder, raw_vec_type,
1652 def, def, components,
1653 num_components);
1654 }
1655 }
1656
1657 static SpvId
emit_select(struct ntv_context * ctx,SpvId type,SpvId cond,SpvId if_true,SpvId if_false)1658 emit_select(struct ntv_context *ctx, SpvId type, SpvId cond,
1659 SpvId if_true, SpvId if_false)
1660 {
1661 return emit_triop(ctx, SpvOpSelect, type, cond, if_true, if_false);
1662 }
1663
1664 static SpvId
emit_bitcast(struct ntv_context * ctx,SpvId type,SpvId value)1665 emit_bitcast(struct ntv_context *ctx, SpvId type, SpvId value)
1666 {
1667 return emit_unop(ctx, SpvOpBitcast, type, value);
1668 }
1669
1670 static SpvId
bitcast_to_uvec(struct ntv_context * ctx,SpvId value,unsigned bit_size,unsigned num_components)1671 bitcast_to_uvec(struct ntv_context *ctx, SpvId value, unsigned bit_size,
1672 unsigned num_components)
1673 {
1674 SpvId type = get_uvec_type(ctx, bit_size, num_components);
1675 return emit_bitcast(ctx, type, value);
1676 }
1677
1678 static SpvId
bitcast_to_ivec(struct ntv_context * ctx,SpvId value,unsigned bit_size,unsigned num_components)1679 bitcast_to_ivec(struct ntv_context *ctx, SpvId value, unsigned bit_size,
1680 unsigned num_components)
1681 {
1682 SpvId type = get_ivec_type(ctx, bit_size, num_components);
1683 return emit_bitcast(ctx, type, value);
1684 }
1685
1686 static SpvId
bitcast_to_fvec(struct ntv_context * ctx,SpvId value,unsigned bit_size,unsigned num_components)1687 bitcast_to_fvec(struct ntv_context *ctx, SpvId value, unsigned bit_size,
1688 unsigned num_components)
1689 {
1690 SpvId type = get_fvec_type(ctx, bit_size, num_components);
1691 return emit_bitcast(ctx, type, value);
1692 }
1693
1694 static SpvId
cast_src_to_type(struct ntv_context * ctx,SpvId value,nir_src src,nir_alu_type atype)1695 cast_src_to_type(struct ntv_context *ctx, SpvId value, nir_src src, nir_alu_type atype)
1696 {
1697 atype = nir_alu_type_get_base_type(atype);
1698 unsigned num_components = nir_src_num_components(src);
1699 unsigned bit_size = nir_src_bit_size(src);
1700 return emit_bitcast(ctx, get_alu_type(ctx, atype, num_components, bit_size), value);
1701 }
1702
1703 static void
store_def(struct ntv_context * ctx,unsigned def_index,SpvId result,nir_alu_type type)1704 store_def(struct ntv_context *ctx, unsigned def_index, SpvId result, nir_alu_type type)
1705 {
1706 assert(result != 0);
1707 assert(def_index < ctx->num_defs);
1708 ctx->def_types[def_index] = nir_alu_type_get_base_type(type);
1709 ctx->defs[def_index] = result;
1710 }
1711
1712 static SpvId
emit_unop(struct ntv_context * ctx,SpvOp op,SpvId type,SpvId src)1713 emit_unop(struct ntv_context *ctx, SpvOp op, SpvId type, SpvId src)
1714 {
1715 return spirv_builder_emit_unop(&ctx->builder, op, type, src);
1716 }
1717
1718 static SpvId
emit_atomic(struct ntv_context * ctx,SpvId op,SpvId type,SpvId src0,SpvId src1,SpvId src2)1719 emit_atomic(struct ntv_context *ctx, SpvId op, SpvId type, SpvId src0, SpvId src1, SpvId src2)
1720 {
1721 if (op == SpvOpAtomicLoad)
1722 return spirv_builder_emit_triop(&ctx->builder, op, type, src0, emit_uint_const(ctx, 32, SpvScopeDevice),
1723 emit_uint_const(ctx, 32, 0));
1724 if (op == SpvOpAtomicCompareExchange)
1725 return spirv_builder_emit_hexop(&ctx->builder, op, type, src0, emit_uint_const(ctx, 32, SpvScopeDevice),
1726 emit_uint_const(ctx, 32, 0),
1727 emit_uint_const(ctx, 32, 0),
1728 /* these params are intentionally swapped */
1729 src2, src1);
1730
1731 return spirv_builder_emit_quadop(&ctx->builder, op, type, src0, emit_uint_const(ctx, 32, SpvScopeDevice),
1732 emit_uint_const(ctx, 32, 0), src1);
1733 }
1734
1735 static SpvId
emit_binop(struct ntv_context * ctx,SpvOp op,SpvId type,SpvId src0,SpvId src1)1736 emit_binop(struct ntv_context *ctx, SpvOp op, SpvId type,
1737 SpvId src0, SpvId src1)
1738 {
1739 return spirv_builder_emit_binop(&ctx->builder, op, type, src0, src1);
1740 }
1741
1742 static SpvId
emit_triop(struct ntv_context * ctx,SpvOp op,SpvId type,SpvId src0,SpvId src1,SpvId src2)1743 emit_triop(struct ntv_context *ctx, SpvOp op, SpvId type,
1744 SpvId src0, SpvId src1, SpvId src2)
1745 {
1746 return spirv_builder_emit_triop(&ctx->builder, op, type, src0, src1, src2);
1747 }
1748
1749 static SpvId
emit_builtin_unop(struct ntv_context * ctx,enum GLSLstd450 op,SpvId type,SpvId src)1750 emit_builtin_unop(struct ntv_context *ctx, enum GLSLstd450 op, SpvId type,
1751 SpvId src)
1752 {
1753 SpvId args[] = { src };
1754 return spirv_builder_emit_ext_inst(&ctx->builder, type, ctx->GLSL_std_450,
1755 op, args, ARRAY_SIZE(args));
1756 }
1757
1758 static SpvId
emit_builtin_binop(struct ntv_context * ctx,enum GLSLstd450 op,SpvId type,SpvId src0,SpvId src1)1759 emit_builtin_binop(struct ntv_context *ctx, enum GLSLstd450 op, SpvId type,
1760 SpvId src0, SpvId src1)
1761 {
1762 SpvId args[] = { src0, src1 };
1763 return spirv_builder_emit_ext_inst(&ctx->builder, type, ctx->GLSL_std_450,
1764 op, args, ARRAY_SIZE(args));
1765 }
1766
1767 static SpvId
emit_builtin_triop(struct ntv_context * ctx,enum GLSLstd450 op,SpvId type,SpvId src0,SpvId src1,SpvId src2)1768 emit_builtin_triop(struct ntv_context *ctx, enum GLSLstd450 op, SpvId type,
1769 SpvId src0, SpvId src1, SpvId src2)
1770 {
1771 SpvId args[] = { src0, src1, src2 };
1772 return spirv_builder_emit_ext_inst(&ctx->builder, type, ctx->GLSL_std_450,
1773 op, args, ARRAY_SIZE(args));
1774 }
1775
1776 static SpvId
get_fvec_constant(struct ntv_context * ctx,unsigned bit_size,unsigned num_components,double value)1777 get_fvec_constant(struct ntv_context *ctx, unsigned bit_size,
1778 unsigned num_components, double value)
1779 {
1780 assert(bit_size == 16 || bit_size == 32 || bit_size == 64);
1781
1782 SpvId result = emit_float_const(ctx, bit_size, value);
1783 if (num_components == 1)
1784 return result;
1785
1786 assert(num_components > 1);
1787 SpvId components[NIR_MAX_VEC_COMPONENTS];
1788 for (int i = 0; i < num_components; i++)
1789 components[i] = result;
1790
1791 SpvId type = get_fvec_type(ctx, bit_size, num_components);
1792 return spirv_builder_const_composite(&ctx->builder, type, components,
1793 num_components);
1794 }
1795
1796 static SpvId
get_ivec_constant(struct ntv_context * ctx,unsigned bit_size,unsigned num_components,int64_t value)1797 get_ivec_constant(struct ntv_context *ctx, unsigned bit_size,
1798 unsigned num_components, int64_t value)
1799 {
1800 assert(bit_size == 8 || bit_size == 16 || bit_size == 32 || bit_size == 64);
1801
1802 SpvId result = emit_int_const(ctx, bit_size, value);
1803 if (num_components == 1)
1804 return result;
1805
1806 assert(num_components > 1);
1807 SpvId components[NIR_MAX_VEC_COMPONENTS];
1808 for (int i = 0; i < num_components; i++)
1809 components[i] = result;
1810
1811 SpvId type = get_ivec_type(ctx, bit_size, num_components);
1812 return spirv_builder_const_composite(&ctx->builder, type, components,
1813 num_components);
1814 }
1815
1816 static inline unsigned
alu_instr_src_components(const nir_alu_instr * instr,unsigned src)1817 alu_instr_src_components(const nir_alu_instr *instr, unsigned src)
1818 {
1819 if (nir_op_infos[instr->op].input_sizes[src] > 0)
1820 return nir_op_infos[instr->op].input_sizes[src];
1821
1822 return instr->def.num_components;
1823 }
1824
1825 static SpvId
get_alu_src(struct ntv_context * ctx,nir_alu_instr * alu,unsigned src,SpvId * raw_value,nir_alu_type * atype)1826 get_alu_src(struct ntv_context *ctx, nir_alu_instr *alu, unsigned src, SpvId *raw_value, nir_alu_type *atype)
1827 {
1828 *raw_value = get_alu_src_raw(ctx, alu, src, atype);
1829
1830 unsigned num_components = alu_instr_src_components(alu, src);
1831 unsigned bit_size = nir_src_bit_size(alu->src[src].src);
1832 nir_alu_type type = alu_op_is_typeless(alu->op) ? *atype : nir_op_infos[alu->op].input_types[src];
1833 type = nir_alu_type_get_base_type(type);
1834 if (type == *atype)
1835 return *raw_value;
1836
1837 if (bit_size == 1)
1838 return *raw_value;
1839 else {
1840 switch (nir_alu_type_get_base_type(type)) {
1841 case nir_type_bool:
1842 unreachable("bool should have bit-size 1");
1843
1844 case nir_type_int:
1845 return bitcast_to_ivec(ctx, *raw_value, bit_size, num_components);
1846
1847 case nir_type_uint:
1848 return bitcast_to_uvec(ctx, *raw_value, bit_size, num_components);
1849
1850 case nir_type_float:
1851 return bitcast_to_fvec(ctx, *raw_value, bit_size, num_components);
1852
1853 default:
1854 unreachable("unknown nir_alu_type");
1855 }
1856 }
1857 }
1858
1859 static void
store_alu_result(struct ntv_context * ctx,nir_alu_instr * alu,SpvId result,nir_alu_type atype)1860 store_alu_result(struct ntv_context *ctx, nir_alu_instr *alu, SpvId result, nir_alu_type atype)
1861 {
1862 store_def(ctx, alu->def.index, result, atype);
1863 }
1864
1865 static SpvId
get_def_type(struct ntv_context * ctx,nir_def * def,nir_alu_type type)1866 get_def_type(struct ntv_context *ctx, nir_def *def, nir_alu_type type)
1867 {
1868 return get_alu_type(ctx, type, def->num_components, def->bit_size);
1869 }
1870
1871 static void
emit_alu(struct ntv_context * ctx,nir_alu_instr * alu)1872 emit_alu(struct ntv_context *ctx, nir_alu_instr *alu)
1873 {
1874 bool is_bcsel = alu->op == nir_op_bcsel;
1875 nir_alu_type stype[NIR_MAX_VEC_COMPONENTS] = {0};
1876 SpvId src[NIR_MAX_VEC_COMPONENTS];
1877 SpvId raw_src[NIR_MAX_VEC_COMPONENTS];
1878 for (unsigned i = 0; i < nir_op_infos[alu->op].num_inputs; i++)
1879 src[i] = get_alu_src(ctx, alu, i, &raw_src[i], &stype[i]);
1880
1881 nir_alu_type typeless_type = stype[is_bcsel];
1882 if (nir_op_infos[alu->op].num_inputs > 1 &&
1883 alu_op_is_typeless(alu->op) &&
1884 nir_src_bit_size(alu->src[is_bcsel].src) != 1) {
1885 unsigned uint_count = 0;
1886 unsigned int_count = 0;
1887 unsigned float_count = 0;
1888 for (unsigned i = is_bcsel; i < nir_op_infos[alu->op].num_inputs; i++) {
1889 if (stype[i] == nir_type_bool)
1890 break;
1891 switch (stype[i]) {
1892 case nir_type_uint:
1893 uint_count++;
1894 break;
1895 case nir_type_int:
1896 int_count++;
1897 break;
1898 case nir_type_float:
1899 float_count++;
1900 break;
1901 default:
1902 unreachable("this shouldn't happen");
1903 }
1904 }
1905 if (uint_count > int_count && uint_count > float_count)
1906 typeless_type = nir_type_uint;
1907 else if (int_count > uint_count && int_count > float_count)
1908 typeless_type = nir_type_int;
1909 else if (float_count > uint_count && float_count > int_count)
1910 typeless_type = nir_type_float;
1911 else if (float_count == uint_count || uint_count == int_count)
1912 typeless_type = nir_type_uint;
1913 else if (float_count == int_count)
1914 typeless_type = nir_type_float;
1915 else
1916 typeless_type = nir_type_uint;
1917 assert(typeless_type != nir_type_bool);
1918 for (unsigned i = is_bcsel; i < nir_op_infos[alu->op].num_inputs; i++) {
1919 unsigned num_components = alu_instr_src_components(alu, i);
1920 unsigned bit_size = nir_src_bit_size(alu->src[i].src);
1921 SpvId type = get_alu_type(ctx, typeless_type, num_components, bit_size);
1922 if (stype[i] != typeless_type) {
1923 src[i] = emit_bitcast(ctx, type, src[i]);
1924 }
1925 }
1926 }
1927
1928 unsigned bit_size = alu->def.bit_size;
1929 unsigned num_components = alu->def.num_components;
1930 nir_alu_type atype = bit_size == 1 ?
1931 nir_type_bool :
1932 (alu_op_is_typeless(alu->op) ? typeless_type : nir_op_infos[alu->op].output_type);
1933 SpvId dest_type = get_def_type(ctx, &alu->def, atype);
1934
1935 SpvId result = 0;
1936 switch (alu->op) {
1937 case nir_op_mov:
1938 assert(nir_op_infos[alu->op].num_inputs == 1);
1939 result = src[0];
1940 break;
1941
1942 #define UNOP(nir_op, spirv_op) \
1943 case nir_op: \
1944 assert(nir_op_infos[alu->op].num_inputs == 1); \
1945 result = emit_unop(ctx, spirv_op, dest_type, src[0]); \
1946 break;
1947
1948 UNOP(nir_op_ineg, SpvOpSNegate)
1949 UNOP(nir_op_fneg, SpvOpFNegate)
1950 UNOP(nir_op_f2i8, SpvOpConvertFToS)
1951 UNOP(nir_op_f2u8, SpvOpConvertFToU)
1952 UNOP(nir_op_f2i16, SpvOpConvertFToS)
1953 UNOP(nir_op_f2u16, SpvOpConvertFToU)
1954 UNOP(nir_op_f2i32, SpvOpConvertFToS)
1955 UNOP(nir_op_f2u32, SpvOpConvertFToU)
1956 UNOP(nir_op_i2f16, SpvOpConvertSToF)
1957 UNOP(nir_op_i2f32, SpvOpConvertSToF)
1958 UNOP(nir_op_u2f16, SpvOpConvertUToF)
1959 UNOP(nir_op_u2f32, SpvOpConvertUToF)
1960 UNOP(nir_op_i2i8, SpvOpSConvert)
1961 UNOP(nir_op_i2i16, SpvOpSConvert)
1962 UNOP(nir_op_i2i32, SpvOpSConvert)
1963 UNOP(nir_op_u2u8, SpvOpUConvert)
1964 UNOP(nir_op_u2u16, SpvOpUConvert)
1965 UNOP(nir_op_u2u32, SpvOpUConvert)
1966 UNOP(nir_op_f2f16, SpvOpFConvert)
1967 UNOP(nir_op_f2f32, SpvOpFConvert)
1968 UNOP(nir_op_f2i64, SpvOpConvertFToS)
1969 UNOP(nir_op_f2u64, SpvOpConvertFToU)
1970 UNOP(nir_op_u2f64, SpvOpConvertUToF)
1971 UNOP(nir_op_i2f64, SpvOpConvertSToF)
1972 UNOP(nir_op_i2i64, SpvOpSConvert)
1973 UNOP(nir_op_u2u64, SpvOpUConvert)
1974 UNOP(nir_op_f2f64, SpvOpFConvert)
1975 UNOP(nir_op_bitfield_reverse, SpvOpBitReverse)
1976 UNOP(nir_op_bit_count, SpvOpBitCount)
1977 #undef UNOP
1978
1979 case nir_op_f2f16_rtz:
1980 assert(nir_op_infos[alu->op].num_inputs == 1);
1981 result = emit_unop(ctx, SpvOpFConvert, dest_type, src[0]);
1982 spirv_builder_emit_rounding_mode(&ctx->builder, result, SpvFPRoundingModeRTZ);
1983 break;
1984
1985 case nir_op_inot:
1986 if (bit_size == 1)
1987 result = emit_unop(ctx, SpvOpLogicalNot, dest_type, src[0]);
1988 else
1989 result = emit_unop(ctx, SpvOpNot, dest_type, src[0]);
1990 break;
1991
1992 case nir_op_b2i8:
1993 case nir_op_b2i16:
1994 case nir_op_b2i32:
1995 case nir_op_b2i64:
1996 assert(nir_op_infos[alu->op].num_inputs == 1);
1997 result = emit_select(ctx, dest_type, src[0],
1998 get_ivec_constant(ctx, bit_size, num_components, 1),
1999 get_ivec_constant(ctx, bit_size, num_components, 0));
2000 break;
2001
2002 case nir_op_b2f16:
2003 case nir_op_b2f32:
2004 case nir_op_b2f64:
2005 assert(nir_op_infos[alu->op].num_inputs == 1);
2006 result = emit_select(ctx, dest_type, src[0],
2007 get_fvec_constant(ctx, bit_size, num_components, 1),
2008 get_fvec_constant(ctx, bit_size, num_components, 0));
2009 break;
2010
2011 case nir_op_uclz:
2012 assert(nir_op_infos[alu->op].num_inputs == 1);
2013 result = emit_unop(ctx, SpvOpUCountLeadingZerosINTEL, dest_type, src[0]);
2014 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityIntegerFunctions2INTEL);
2015 spirv_builder_emit_extension(&ctx->builder, "SPV_INTEL_shader_integer_functions2");
2016 break;
2017 #define BUILTIN_UNOP(nir_op, spirv_op) \
2018 case nir_op: \
2019 assert(nir_op_infos[alu->op].num_inputs == 1); \
2020 result = emit_builtin_unop(ctx, spirv_op, dest_type, src[0]); \
2021 break;
2022
2023 #define BUILTIN_UNOPF(nir_op, spirv_op) \
2024 case nir_op: \
2025 assert(nir_op_infos[alu->op].num_inputs == 1); \
2026 result = emit_builtin_unop(ctx, spirv_op, get_def_type(ctx, &alu->def, nir_type_float), src[0]); \
2027 atype = nir_type_float; \
2028 break;
2029
2030 BUILTIN_UNOP(nir_op_iabs, GLSLstd450SAbs)
2031 BUILTIN_UNOP(nir_op_fabs, GLSLstd450FAbs)
2032 BUILTIN_UNOP(nir_op_fsqrt, GLSLstd450Sqrt)
2033 BUILTIN_UNOP(nir_op_frsq, GLSLstd450InverseSqrt)
2034 BUILTIN_UNOP(nir_op_flog2, GLSLstd450Log2)
2035 BUILTIN_UNOP(nir_op_fexp2, GLSLstd450Exp2)
2036 BUILTIN_UNOP(nir_op_ffract, GLSLstd450Fract)
2037 BUILTIN_UNOP(nir_op_ffloor, GLSLstd450Floor)
2038 BUILTIN_UNOP(nir_op_fceil, GLSLstd450Ceil)
2039 BUILTIN_UNOP(nir_op_ftrunc, GLSLstd450Trunc)
2040 BUILTIN_UNOP(nir_op_fround_even, GLSLstd450RoundEven)
2041 BUILTIN_UNOP(nir_op_fsign, GLSLstd450FSign)
2042 BUILTIN_UNOP(nir_op_isign, GLSLstd450SSign)
2043 BUILTIN_UNOP(nir_op_fsin, GLSLstd450Sin)
2044 BUILTIN_UNOP(nir_op_fcos, GLSLstd450Cos)
2045 BUILTIN_UNOP(nir_op_ufind_msb, GLSLstd450FindUMsb)
2046 BUILTIN_UNOP(nir_op_find_lsb, GLSLstd450FindILsb)
2047 BUILTIN_UNOP(nir_op_ifind_msb, GLSLstd450FindSMsb)
2048
2049 case nir_op_pack_half_2x16:
2050 assert(nir_op_infos[alu->op].num_inputs == 1);
2051 result = emit_builtin_unop(ctx, GLSLstd450PackHalf2x16, get_def_type(ctx, &alu->def, nir_type_uint), src[0]);
2052 break;
2053
2054 BUILTIN_UNOPF(nir_op_unpack_half_2x16, GLSLstd450UnpackHalf2x16)
2055 #undef BUILTIN_UNOP
2056 #undef BUILTIN_UNOPF
2057
2058 case nir_op_frcp:
2059 assert(nir_op_infos[alu->op].num_inputs == 1);
2060 result = emit_binop(ctx, SpvOpFDiv, dest_type,
2061 get_fvec_constant(ctx, bit_size, num_components, 1),
2062 src[0]);
2063 break;
2064
2065
2066 #define BINOP(nir_op, spirv_op) \
2067 case nir_op: \
2068 assert(nir_op_infos[alu->op].num_inputs == 2); \
2069 result = emit_binop(ctx, spirv_op, dest_type, src[0], src[1]); \
2070 break;
2071
2072 BINOP(nir_op_iadd, SpvOpIAdd)
2073 BINOP(nir_op_isub, SpvOpISub)
2074 BINOP(nir_op_imul, SpvOpIMul)
2075 BINOP(nir_op_idiv, SpvOpSDiv)
2076 BINOP(nir_op_udiv, SpvOpUDiv)
2077 BINOP(nir_op_umod, SpvOpUMod)
2078 BINOP(nir_op_imod, SpvOpSMod)
2079 BINOP(nir_op_irem, SpvOpSRem)
2080 BINOP(nir_op_fadd, SpvOpFAdd)
2081 BINOP(nir_op_fsub, SpvOpFSub)
2082 BINOP(nir_op_fmul, SpvOpFMul)
2083 BINOP(nir_op_fdiv, SpvOpFDiv)
2084 BINOP(nir_op_fmod, SpvOpFMod)
2085 BINOP(nir_op_ilt, SpvOpSLessThan)
2086 BINOP(nir_op_ige, SpvOpSGreaterThanEqual)
2087 BINOP(nir_op_ult, SpvOpULessThan)
2088 BINOP(nir_op_uge, SpvOpUGreaterThanEqual)
2089 BINOP(nir_op_flt, SpvOpFOrdLessThan)
2090 BINOP(nir_op_fge, SpvOpFOrdGreaterThanEqual)
2091 BINOP(nir_op_frem, SpvOpFRem)
2092 #undef BINOP
2093
2094 #define BINOP_LOG(nir_op, spv_op, spv_log_op) \
2095 case nir_op: \
2096 assert(nir_op_infos[alu->op].num_inputs == 2); \
2097 if (nir_src_bit_size(alu->src[0].src) == 1) \
2098 result = emit_binop(ctx, spv_log_op, dest_type, src[0], src[1]); \
2099 else \
2100 result = emit_binop(ctx, spv_op, dest_type, src[0], src[1]); \
2101 break;
2102
2103 BINOP_LOG(nir_op_iand, SpvOpBitwiseAnd, SpvOpLogicalAnd)
2104 BINOP_LOG(nir_op_ior, SpvOpBitwiseOr, SpvOpLogicalOr)
2105 BINOP_LOG(nir_op_ieq, SpvOpIEqual, SpvOpLogicalEqual)
2106 BINOP_LOG(nir_op_ine, SpvOpINotEqual, SpvOpLogicalNotEqual)
2107 BINOP_LOG(nir_op_ixor, SpvOpBitwiseXor, SpvOpLogicalNotEqual)
2108 #undef BINOP_LOG
2109
2110 #define BINOP_SHIFT(nir_op, spirv_op) \
2111 case nir_op: { \
2112 assert(nir_op_infos[alu->op].num_inputs == 2); \
2113 int shift_bit_size = nir_src_bit_size(alu->src[1].src); \
2114 nir_alu_type shift_nir_type = nir_alu_type_get_base_type(nir_op_infos[alu->op].input_types[1]); \
2115 SpvId shift_type = get_alu_type(ctx, shift_nir_type, num_components, shift_bit_size); \
2116 SpvId shift_mask = get_ivec_constant(ctx, shift_bit_size, num_components, bit_size - 1); \
2117 SpvId shift_count = emit_binop(ctx, SpvOpBitwiseAnd, shift_type, src[1], shift_mask); \
2118 result = emit_binop(ctx, spirv_op, dest_type, src[0], shift_count); \
2119 break; \
2120 }
2121
2122 BINOP_SHIFT(nir_op_ishl, SpvOpShiftLeftLogical)
2123 BINOP_SHIFT(nir_op_ishr, SpvOpShiftRightArithmetic)
2124 BINOP_SHIFT(nir_op_ushr, SpvOpShiftRightLogical)
2125 #undef BINOP_SHIFT
2126
2127 #define BUILTIN_BINOP(nir_op, spirv_op) \
2128 case nir_op: \
2129 assert(nir_op_infos[alu->op].num_inputs == 2); \
2130 result = emit_builtin_binop(ctx, spirv_op, dest_type, src[0], src[1]); \
2131 break;
2132
2133 BUILTIN_BINOP(nir_op_fmin, GLSLstd450FMin)
2134 BUILTIN_BINOP(nir_op_fmax, GLSLstd450FMax)
2135 BUILTIN_BINOP(nir_op_imin, GLSLstd450SMin)
2136 BUILTIN_BINOP(nir_op_imax, GLSLstd450SMax)
2137 BUILTIN_BINOP(nir_op_umin, GLSLstd450UMin)
2138 BUILTIN_BINOP(nir_op_umax, GLSLstd450UMax)
2139 BUILTIN_BINOP(nir_op_ldexp, GLSLstd450Ldexp)
2140 BUILTIN_BINOP(nir_op_fpow, GLSLstd450Pow)
2141 #undef BUILTIN_BINOP
2142
2143 #define INTEL_BINOP(nir_op, spirv_op) \
2144 case nir_op: \
2145 assert(nir_op_infos[alu->op].num_inputs == 2); \
2146 result = emit_binop(ctx, spirv_op, dest_type, src[0], src[1]); \
2147 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityIntegerFunctions2INTEL); \
2148 spirv_builder_emit_extension(&ctx->builder, "SPV_INTEL_shader_integer_functions2"); \
2149 break;
2150
2151 INTEL_BINOP(nir_op_uabs_isub, SpvOpAbsISubINTEL)
2152 INTEL_BINOP(nir_op_uabs_usub, SpvOpAbsUSubINTEL)
2153 INTEL_BINOP(nir_op_iadd_sat, SpvOpIAddSatINTEL)
2154 INTEL_BINOP(nir_op_uadd_sat, SpvOpUAddSatINTEL)
2155 INTEL_BINOP(nir_op_ihadd, SpvOpIAverageINTEL)
2156 INTEL_BINOP(nir_op_uhadd, SpvOpUAverageINTEL)
2157 INTEL_BINOP(nir_op_irhadd, SpvOpIAverageRoundedINTEL)
2158 INTEL_BINOP(nir_op_urhadd, SpvOpUAverageRoundedINTEL)
2159 INTEL_BINOP(nir_op_isub_sat, SpvOpISubSatINTEL)
2160 INTEL_BINOP(nir_op_usub_sat, SpvOpUSubSatINTEL)
2161 INTEL_BINOP(nir_op_imul_32x16, SpvOpIMul32x16INTEL)
2162 INTEL_BINOP(nir_op_umul_32x16, SpvOpUMul32x16INTEL)
2163 #undef INTEL_BINOP
2164
2165 case nir_op_fdot2:
2166 case nir_op_fdot3:
2167 case nir_op_fdot4:
2168 assert(nir_op_infos[alu->op].num_inputs == 2);
2169 result = emit_binop(ctx, SpvOpDot, dest_type, src[0], src[1]);
2170 break;
2171
2172 case nir_op_fdph:
2173 case nir_op_seq:
2174 case nir_op_sne:
2175 case nir_op_slt:
2176 case nir_op_sge:
2177 unreachable("should already be lowered away");
2178
2179 case nir_op_fneu:
2180 assert(nir_op_infos[alu->op].num_inputs == 2);
2181 if (raw_src[0] == raw_src[1])
2182 result = emit_unop(ctx, SpvOpIsNan, dest_type, src[0]);
2183 else
2184 result = emit_binop(ctx, SpvOpFUnordNotEqual, dest_type, src[0], src[1]);
2185 break;
2186
2187 case nir_op_feq:
2188 assert(nir_op_infos[alu->op].num_inputs == 2);
2189 if (raw_src[0] == raw_src[1])
2190 result = emit_unop(ctx, SpvOpLogicalNot, dest_type,
2191 emit_unop(ctx, SpvOpIsNan, dest_type, src[0]));
2192 else
2193 result = emit_binop(ctx, SpvOpFOrdEqual, dest_type, src[0], src[1]);
2194 break;
2195
2196 case nir_op_flrp:
2197 assert(nir_op_infos[alu->op].num_inputs == 3);
2198 result = emit_builtin_triop(ctx, GLSLstd450FMix, dest_type,
2199 src[0], src[1], src[2]);
2200 break;
2201
2202 case nir_op_bcsel:
2203 assert(nir_op_infos[alu->op].num_inputs == 3);
2204 result = emit_select(ctx, dest_type, src[0], src[1], src[2]);
2205 break;
2206
2207 case nir_op_pack_half_2x16_split: {
2208 SpvId fvec = spirv_builder_emit_composite_construct(&ctx->builder, get_fvec_type(ctx, 32, 2),
2209 src, 2);
2210 result = emit_builtin_unop(ctx, GLSLstd450PackHalf2x16, dest_type, fvec);
2211 break;
2212 }
2213 case nir_op_vec2:
2214 case nir_op_vec3:
2215 case nir_op_vec4: {
2216 int num_inputs = nir_op_infos[alu->op].num_inputs;
2217 assert(2 <= num_inputs && num_inputs <= 4);
2218 result = spirv_builder_emit_composite_construct(&ctx->builder, dest_type,
2219 src, num_inputs);
2220 }
2221 break;
2222
2223 case nir_op_ubitfield_extract:
2224 assert(nir_op_infos[alu->op].num_inputs == 3);
2225 result = emit_triop(ctx, SpvOpBitFieldUExtract, dest_type, src[0], src[1], src[2]);
2226 break;
2227
2228 case nir_op_ibitfield_extract:
2229 assert(nir_op_infos[alu->op].num_inputs == 3);
2230 result = emit_triop(ctx, SpvOpBitFieldSExtract, dest_type, src[0], src[1], src[2]);
2231 break;
2232
2233 case nir_op_bitfield_insert:
2234 assert(nir_op_infos[alu->op].num_inputs == 4);
2235 result = spirv_builder_emit_quadop(&ctx->builder, SpvOpBitFieldInsert, dest_type, src[0], src[1], src[2], src[3]);
2236 break;
2237
2238 /* those are all simple bitcasts, we could do better, but it doesn't matter */
2239 case nir_op_pack_32_4x8:
2240 case nir_op_pack_32_2x16:
2241 case nir_op_pack_64_2x32:
2242 case nir_op_pack_64_4x16:
2243 case nir_op_unpack_32_4x8:
2244 case nir_op_unpack_32_2x16:
2245 case nir_op_unpack_64_2x32:
2246 case nir_op_unpack_64_4x16: {
2247 result = emit_bitcast(ctx, dest_type, src[0]);
2248 break;
2249 }
2250
2251 case nir_op_pack_32_2x16_split:
2252 case nir_op_pack_64_2x32_split: {
2253 nir_alu_type type = nir_alu_type_get_base_type(nir_op_infos[alu->op].input_types[0]);
2254 if (num_components <= 2) {
2255 SpvId components[] = {src[0], src[1]};
2256 SpvId vec_type = get_alu_type(ctx, type, num_components * 2, nir_src_bit_size(alu->src[0].src));
2257 result = spirv_builder_emit_composite_construct(&ctx->builder, vec_type, components, 2);
2258 result = emit_bitcast(ctx, dest_type, result);
2259 } else {
2260 SpvId components[NIR_MAX_VEC_COMPONENTS];
2261 SpvId conv_type = get_alu_type(ctx, type, 1, nir_src_bit_size(alu->src[0].src));
2262 SpvId vec_type = get_alu_type(ctx, type, 2, nir_src_bit_size(alu->src[0].src));
2263 SpvId dest_scalar_type = get_alu_type(ctx, nir_op_infos[alu->op].output_type, 1, bit_size);
2264 for (unsigned i = 0; i < nir_src_num_components(alu->src[0].src); i++) {
2265 SpvId conv[2];
2266 conv[0] = spirv_builder_emit_composite_extract(&ctx->builder, conv_type, src[0], &i, 1);
2267 conv[1] = spirv_builder_emit_composite_extract(&ctx->builder, conv_type, src[1], &i, 1);
2268 SpvId vec = spirv_builder_emit_composite_construct(&ctx->builder, vec_type, conv, 2);
2269 components[i] = emit_bitcast(ctx, dest_scalar_type, vec);
2270 }
2271 result = spirv_builder_emit_composite_construct(&ctx->builder, dest_type, components, num_components);
2272 }
2273 break;
2274 }
2275
2276 case nir_op_unpack_32_2x16_split_x:
2277 case nir_op_unpack_64_2x32_split_x: {
2278 nir_alu_type type = nir_alu_type_get_base_type(nir_op_infos[alu->op].input_types[0]);
2279 SpvId vec_type = get_alu_type(ctx, type, 2, bit_size);
2280 unsigned idx = 0;
2281 if (num_components == 1) {
2282 SpvId vec = emit_bitcast(ctx, vec_type, src[0]);
2283 result = spirv_builder_emit_composite_extract(&ctx->builder, dest_type, vec, &idx, 1);
2284 } else {
2285 SpvId components[NIR_MAX_VEC_COMPONENTS];
2286 for (unsigned i = 0; i < nir_src_num_components(alu->src[0].src); i++) {
2287 SpvId conv = spirv_builder_emit_composite_extract(&ctx->builder, get_alu_type(ctx, type, 1, nir_src_bit_size(alu->src[0].src)), src[0], &i, 1);
2288 conv = emit_bitcast(ctx, vec_type, conv);
2289 SpvId conv_type = get_alu_type(ctx, type, 1, bit_size);
2290 components[i] = spirv_builder_emit_composite_extract(&ctx->builder, conv_type, conv, &idx, 1);
2291 }
2292 result = spirv_builder_emit_composite_construct(&ctx->builder, dest_type, components, num_components);
2293 }
2294 break;
2295 }
2296
2297 case nir_op_unpack_32_2x16_split_y:
2298 case nir_op_unpack_64_2x32_split_y: {
2299 nir_alu_type type = nir_alu_type_get_base_type(nir_op_infos[alu->op].input_types[0]);
2300 SpvId vec_type = get_alu_type(ctx, type, 2, bit_size);
2301 unsigned idx = 1;
2302 if (num_components == 1) {
2303 SpvId vec = emit_bitcast(ctx, vec_type, src[0]);
2304 result = spirv_builder_emit_composite_extract(&ctx->builder, dest_type, vec, &idx, 1);
2305 } else {
2306 SpvId components[NIR_MAX_VEC_COMPONENTS];
2307 for (unsigned i = 0; i < nir_src_num_components(alu->src[0].src); i++) {
2308 SpvId conv = spirv_builder_emit_composite_extract(&ctx->builder, get_alu_type(ctx, type, 1, nir_src_bit_size(alu->src[0].src)), src[0], &i, 1);
2309 conv = emit_bitcast(ctx, vec_type, conv);
2310 SpvId conv_type = get_alu_type(ctx, type, 1, bit_size);
2311 components[i] = spirv_builder_emit_composite_extract(&ctx->builder, conv_type, conv, &idx, 1);
2312 }
2313 result = spirv_builder_emit_composite_construct(&ctx->builder, dest_type, components, num_components);
2314 }
2315 break;
2316 }
2317
2318 default:
2319 fprintf(stderr, "emit_alu: not implemented (%s)\n",
2320 nir_op_infos[alu->op].name);
2321
2322 unreachable("unsupported opcode");
2323 return;
2324 }
2325 if (alu->exact)
2326 spirv_builder_emit_decoration(&ctx->builder, result, SpvDecorationNoContraction);
2327
2328 store_alu_result(ctx, alu, result, atype);
2329 }
2330
2331 static void
emit_load_const(struct ntv_context * ctx,nir_load_const_instr * load_const)2332 emit_load_const(struct ntv_context *ctx, nir_load_const_instr *load_const)
2333 {
2334 unsigned bit_size = load_const->def.bit_size;
2335 unsigned num_components = load_const->def.num_components;
2336
2337 SpvId components[NIR_MAX_VEC_COMPONENTS];
2338 nir_alu_type atype;
2339 if (bit_size == 1) {
2340 atype = nir_type_bool;
2341 for (int i = 0; i < num_components; i++)
2342 components[i] = spirv_builder_const_bool(&ctx->builder,
2343 load_const->value[i].b);
2344 } else {
2345 if (ctx->sinfo->broken_arbitary_type_const)
2346 atype = nir_type_uint;
2347 else
2348 atype = infer_nir_alu_type_from_uses_ssa(&load_const->def);
2349 for (int i = 0; i < num_components; i++) {
2350 switch (atype) {
2351 case nir_type_uint: {
2352 uint64_t tmp = nir_const_value_as_uint(load_const->value[i], bit_size);
2353 components[i] = emit_uint_const(ctx, bit_size, tmp);
2354 break;
2355 }
2356 case nir_type_int: {
2357 int64_t tmp = nir_const_value_as_int(load_const->value[i], bit_size);
2358 components[i] = emit_int_const(ctx, bit_size, tmp);
2359 break;
2360 }
2361 case nir_type_float: {
2362 double tmp = nir_const_value_as_float(load_const->value[i], bit_size);
2363 components[i] = emit_float_const(ctx, bit_size, tmp);
2364 break;
2365 }
2366 default:
2367 unreachable("this shouldn't happen!");
2368 }
2369 }
2370 }
2371
2372 if (num_components > 1) {
2373 SpvId type = get_alu_type(ctx, atype, num_components, bit_size);
2374 SpvId value = spirv_builder_const_composite(&ctx->builder,
2375 type, components,
2376 num_components);
2377 store_def(ctx, load_const->def.index, value, atype);
2378 } else {
2379 assert(num_components == 1);
2380 store_def(ctx, load_const->def.index, components[0], atype);
2381 }
2382 }
2383
2384 static void
emit_discard(struct ntv_context * ctx,nir_intrinsic_instr * intr)2385 emit_discard(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2386 {
2387 assert(ctx->discard_func);
2388 SpvId type_void = spirv_builder_type_void(&ctx->builder);
2389 spirv_builder_function_call(&ctx->builder, type_void,
2390 ctx->discard_func, NULL, 0);
2391 }
2392
2393 static void
emit_load_deref(struct ntv_context * ctx,nir_intrinsic_instr * intr)2394 emit_load_deref(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2395 {
2396 nir_alu_type atype;
2397 SpvId ptr = get_src(ctx, intr->src, &atype);
2398
2399 nir_deref_instr *deref = nir_src_as_deref(intr->src[0]);
2400 SpvId type;
2401 if (glsl_type_is_image(deref->type)) {
2402 nir_variable *var = nir_deref_instr_get_variable(deref);
2403 const struct glsl_type *gtype = glsl_without_array(var->type);
2404 type = get_image_type(ctx, var,
2405 glsl_type_is_sampler(gtype),
2406 glsl_get_sampler_dim(gtype) == GLSL_SAMPLER_DIM_BUF);
2407 atype = nir_get_nir_type_for_glsl_base_type(glsl_get_sampler_result_type(gtype));
2408 } else {
2409 type = get_glsl_type(ctx, deref->type);
2410 atype = get_nir_alu_type(deref->type);
2411 }
2412 SpvId result;
2413
2414 if (nir_intrinsic_access(intr) & ACCESS_COHERENT)
2415 result = emit_atomic(ctx, SpvOpAtomicLoad, type, ptr, 0, 0);
2416 else
2417 result = spirv_builder_emit_load(&ctx->builder, type, ptr);
2418 store_def(ctx, intr->def.index, result, atype);
2419 }
2420
2421 static void
emit_store_deref(struct ntv_context * ctx,nir_intrinsic_instr * intr)2422 emit_store_deref(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2423 {
2424 nir_alu_type ptype, stype;
2425 SpvId ptr = get_src(ctx, &intr->src[0], &ptype);
2426 SpvId src = get_src(ctx, &intr->src[1], &stype);
2427
2428 const struct glsl_type *gtype = nir_src_as_deref(intr->src[0])->type;
2429 SpvId type = get_glsl_type(ctx, gtype);
2430 nir_variable *var = nir_intrinsic_get_var(intr, 0);
2431 unsigned wrmask = nir_intrinsic_write_mask(intr);
2432 if (!glsl_type_is_scalar(gtype) &&
2433 wrmask != BITFIELD_MASK(glsl_type_is_array(gtype) ? glsl_get_aoa_size(gtype) : glsl_get_vector_elements(gtype))) {
2434 /* no idea what we do if this fails */
2435 assert(glsl_type_is_array(gtype) || glsl_type_is_vector(gtype));
2436
2437 /* this is a partial write, so we have to loop and do a per-component write */
2438 SpvId result_type;
2439 SpvId member_type;
2440 if (glsl_type_is_vector(gtype)) {
2441 result_type = get_glsl_basetype(ctx, glsl_get_base_type(gtype));
2442 member_type = get_alu_type(ctx, stype, 1, glsl_get_bit_size(gtype));
2443 } else
2444 member_type = result_type = get_glsl_type(ctx, glsl_get_array_element(gtype));
2445 SpvId ptr_type = spirv_builder_type_pointer(&ctx->builder,
2446 get_storage_class(var),
2447 result_type);
2448 for (unsigned i = 0; i < 4; i++)
2449 if (wrmask & BITFIELD_BIT(i)) {
2450 SpvId idx = emit_uint_const(ctx, 32, i);
2451 SpvId val = spirv_builder_emit_composite_extract(&ctx->builder, member_type, src, &i, 1);
2452 if (stype != ptype)
2453 val = emit_bitcast(ctx, result_type, val);
2454 SpvId member = spirv_builder_emit_access_chain(&ctx->builder, ptr_type,
2455 ptr, &idx, 1);
2456 spirv_builder_emit_store(&ctx->builder, member, val);
2457 }
2458 return;
2459
2460 }
2461 SpvId result;
2462 if (ctx->stage == MESA_SHADER_FRAGMENT &&
2463 var->data.mode == nir_var_shader_out &&
2464 var->data.location == FRAG_RESULT_SAMPLE_MASK) {
2465 src = emit_bitcast(ctx, type, src);
2466 /* SampleMask is always an array in spirv, so we need to construct it into one */
2467 result = spirv_builder_emit_composite_construct(&ctx->builder, ctx->sample_mask_type, &src, 1);
2468 } else {
2469 if (ptype == stype)
2470 result = src;
2471 else
2472 result = emit_bitcast(ctx, type, src);
2473 }
2474 if (nir_intrinsic_access(intr) & ACCESS_COHERENT)
2475 spirv_builder_emit_atomic_store(&ctx->builder, ptr, SpvScopeDevice, 0, result);
2476 else
2477 spirv_builder_emit_store(&ctx->builder, ptr, result);
2478 }
2479
2480 static void
emit_load_shared(struct ntv_context * ctx,nir_intrinsic_instr * intr)2481 emit_load_shared(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2482 {
2483 SpvId dest_type = get_def_type(ctx, &intr->def, nir_type_uint);
2484 unsigned num_components = intr->def.num_components;
2485 unsigned bit_size = intr->def.bit_size;
2486 SpvId uint_type = get_uvec_type(ctx, bit_size, 1);
2487 SpvId ptr_type = spirv_builder_type_pointer(&ctx->builder,
2488 SpvStorageClassWorkgroup,
2489 uint_type);
2490 nir_alu_type atype;
2491 SpvId offset = get_src(ctx, &intr->src[0], &atype);
2492 if (atype == nir_type_float)
2493 offset = bitcast_to_uvec(ctx, offset, nir_src_bit_size(intr->src[0]), 1);
2494 SpvId constituents[NIR_MAX_VEC_COMPONENTS];
2495 SpvId shared_block = get_shared_block(ctx, bit_size);
2496 /* need to convert array -> vec */
2497 for (unsigned i = 0; i < num_components; i++) {
2498 SpvId member = spirv_builder_emit_access_chain(&ctx->builder, ptr_type,
2499 shared_block, &offset, 1);
2500 constituents[i] = spirv_builder_emit_load(&ctx->builder, uint_type, member);
2501 offset = emit_binop(ctx, SpvOpIAdd, spirv_builder_type_uint(&ctx->builder, 32), offset, emit_uint_const(ctx, 32, 1));
2502 }
2503 SpvId result;
2504 if (num_components > 1)
2505 result = spirv_builder_emit_composite_construct(&ctx->builder, dest_type, constituents, num_components);
2506 else
2507 result = constituents[0];
2508 store_def(ctx, intr->def.index, result, nir_type_uint);
2509 }
2510
2511 static void
emit_store_shared(struct ntv_context * ctx,nir_intrinsic_instr * intr)2512 emit_store_shared(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2513 {
2514 nir_alu_type atype;
2515 SpvId src = get_src(ctx, &intr->src[0], &atype);
2516
2517 unsigned wrmask = nir_intrinsic_write_mask(intr);
2518 unsigned bit_size = nir_src_bit_size(intr->src[0]);
2519 SpvId uint_type = get_uvec_type(ctx, bit_size, 1);
2520 SpvId ptr_type = spirv_builder_type_pointer(&ctx->builder,
2521 SpvStorageClassWorkgroup,
2522 uint_type);
2523 nir_alu_type otype;
2524 SpvId offset = get_src(ctx, &intr->src[1], &otype);
2525 if (otype == nir_type_float)
2526 offset = bitcast_to_uvec(ctx, offset, nir_src_bit_size(intr->src[0]), 1);
2527 SpvId shared_block = get_shared_block(ctx, bit_size);
2528 /* this is a partial write, so we have to loop and do a per-component write */
2529 u_foreach_bit(i, wrmask) {
2530 SpvId shared_offset = emit_binop(ctx, SpvOpIAdd, spirv_builder_type_uint(&ctx->builder, 32), offset, emit_uint_const(ctx, 32, i));
2531 SpvId val = src;
2532 if (nir_src_num_components(intr->src[0]) != 1)
2533 val = spirv_builder_emit_composite_extract(&ctx->builder, uint_type, src, &i, 1);
2534 if (atype != nir_type_uint)
2535 val = emit_bitcast(ctx, get_alu_type(ctx, nir_type_uint, 1, bit_size), val);
2536 SpvId member = spirv_builder_emit_access_chain(&ctx->builder, ptr_type,
2537 shared_block, &shared_offset, 1);
2538 spirv_builder_emit_store(&ctx->builder, member, val);
2539 }
2540 }
2541
2542 static void
emit_load_scratch(struct ntv_context * ctx,nir_intrinsic_instr * intr)2543 emit_load_scratch(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2544 {
2545 SpvId dest_type = get_def_type(ctx, &intr->def, nir_type_uint);
2546 unsigned num_components = intr->def.num_components;
2547 unsigned bit_size = intr->def.bit_size;
2548 SpvId uint_type = get_uvec_type(ctx, bit_size, 1);
2549 SpvId ptr_type = spirv_builder_type_pointer(&ctx->builder,
2550 SpvStorageClassPrivate,
2551 uint_type);
2552 nir_alu_type atype;
2553 SpvId offset = get_src(ctx, &intr->src[0], &atype);
2554 if (atype != nir_type_uint)
2555 offset = bitcast_to_uvec(ctx, offset, nir_src_bit_size(intr->src[0]), 1);
2556 SpvId constituents[NIR_MAX_VEC_COMPONENTS];
2557 SpvId scratch_block = get_scratch_block(ctx, bit_size);
2558 /* need to convert array -> vec */
2559 for (unsigned i = 0; i < num_components; i++) {
2560 SpvId member = spirv_builder_emit_access_chain(&ctx->builder, ptr_type,
2561 scratch_block, &offset, 1);
2562 constituents[i] = spirv_builder_emit_load(&ctx->builder, uint_type, member);
2563 offset = emit_binop(ctx, SpvOpIAdd, spirv_builder_type_uint(&ctx->builder, 32), offset, emit_uint_const(ctx, 32, 1));
2564 }
2565 SpvId result;
2566 if (num_components > 1)
2567 result = spirv_builder_emit_composite_construct(&ctx->builder, dest_type, constituents, num_components);
2568 else
2569 result = constituents[0];
2570 store_def(ctx, intr->def.index, result, nir_type_uint);
2571 }
2572
2573 static void
emit_store_scratch(struct ntv_context * ctx,nir_intrinsic_instr * intr)2574 emit_store_scratch(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2575 {
2576 nir_alu_type atype;
2577 SpvId src = get_src(ctx, &intr->src[0], &atype);
2578
2579 unsigned wrmask = nir_intrinsic_write_mask(intr);
2580 unsigned bit_size = nir_src_bit_size(intr->src[0]);
2581 SpvId uint_type = get_uvec_type(ctx, bit_size, 1);
2582 SpvId ptr_type = spirv_builder_type_pointer(&ctx->builder,
2583 SpvStorageClassPrivate,
2584 uint_type);
2585 nir_alu_type otype;
2586 SpvId offset = get_src(ctx, &intr->src[1], &otype);
2587 if (otype != nir_type_uint)
2588 offset = bitcast_to_uvec(ctx, offset, nir_src_bit_size(intr->src[1]), 1);
2589 SpvId scratch_block = get_scratch_block(ctx, bit_size);
2590 /* this is a partial write, so we have to loop and do a per-component write */
2591 u_foreach_bit(i, wrmask) {
2592 SpvId scratch_offset = emit_binop(ctx, SpvOpIAdd, spirv_builder_type_uint(&ctx->builder, 32), offset, emit_uint_const(ctx, 32, i));
2593 SpvId val = src;
2594 if (nir_src_num_components(intr->src[0]) != 1)
2595 val = spirv_builder_emit_composite_extract(&ctx->builder, uint_type, src, &i, 1);
2596 if (atype != nir_type_uint)
2597 val = emit_bitcast(ctx, get_alu_type(ctx, nir_type_uint, 1, bit_size), val);
2598 SpvId member = spirv_builder_emit_access_chain(&ctx->builder, ptr_type,
2599 scratch_block, &scratch_offset, 1);
2600 spirv_builder_emit_store(&ctx->builder, member, val);
2601 }
2602 }
2603
2604 static void
emit_load_push_const(struct ntv_context * ctx,nir_intrinsic_instr * intr)2605 emit_load_push_const(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2606 {
2607 SpvId uint_type = get_uvec_type(ctx, 32, 1);
2608 SpvId load_type = get_uvec_type(ctx, 32, 1);
2609
2610 /* number of components being loaded */
2611 unsigned num_components = intr->def.num_components;
2612 SpvId constituents[NIR_MAX_VEC_COMPONENTS * 2];
2613 SpvId result;
2614
2615 /* destination type for the load */
2616 SpvId type = get_def_uvec_type(ctx, &intr->def);
2617 SpvId one = emit_uint_const(ctx, 32, 1);
2618
2619 /* we grab a single array member at a time, so it's a pointer to a uint */
2620 SpvId pointer_type = spirv_builder_type_pointer(&ctx->builder,
2621 SpvStorageClassPushConstant,
2622 load_type);
2623
2624 nir_alu_type atype;
2625 SpvId member = get_src(ctx, &intr->src[0], &atype);
2626 if (atype == nir_type_float)
2627 member = bitcast_to_uvec(ctx, member, nir_src_bit_size(intr->src[0]), 1);
2628 /* reuse the offset from ZINK_PUSH_CONST_OFFSET */
2629 SpvId offset = emit_uint_const(ctx, 32, nir_intrinsic_component(intr));
2630 /* OpAccessChain takes an array of indices that drill into a hierarchy based on the type:
2631 * index 0 is accessing 'base'
2632 * index 1 is accessing 'base[index 1]'
2633 *
2634 */
2635 for (unsigned i = 0; i < num_components; i++) {
2636 SpvId indices[2] = { member, offset };
2637 SpvId ptr = spirv_builder_emit_access_chain(&ctx->builder, pointer_type,
2638 ctx->push_const_var, indices,
2639 ARRAY_SIZE(indices));
2640 /* load a single value into the constituents array */
2641 constituents[i] = spirv_builder_emit_load(&ctx->builder, load_type, ptr);
2642 /* increment to the next vec4 member index for the next load */
2643 offset = emit_binop(ctx, SpvOpIAdd, uint_type, offset, one);
2644 }
2645
2646 /* if loading more than 1 value, reassemble the results into the desired type,
2647 * otherwise just use the loaded result
2648 */
2649 if (num_components > 1) {
2650 result = spirv_builder_emit_composite_construct(&ctx->builder,
2651 type,
2652 constituents,
2653 num_components);
2654 } else
2655 result = constituents[0];
2656
2657 store_def(ctx, intr->def.index, result, nir_type_uint);
2658 }
2659
2660 static void
emit_load_global(struct ntv_context * ctx,nir_intrinsic_instr * intr)2661 emit_load_global(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2662 {
2663 bool coherent = ctx->sinfo->have_vulkan_memory_model && nir_intrinsic_access(intr) & ACCESS_COHERENT;
2664 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityPhysicalStorageBufferAddresses);
2665 SpvId dest_type = get_def_type(ctx, &intr->def, nir_type_uint);
2666 SpvId pointer_type = spirv_builder_type_pointer(&ctx->builder,
2667 SpvStorageClassPhysicalStorageBuffer,
2668 dest_type);
2669 nir_alu_type atype;
2670 SpvId ptr = emit_bitcast(ctx, pointer_type, get_src(ctx, &intr->src[0], &atype));
2671 SpvId result = spirv_builder_emit_load_aligned(&ctx->builder, dest_type, ptr, intr->def.bit_size / 8, coherent);
2672 store_def(ctx, intr->def.index, result, nir_type_uint);
2673 }
2674
2675 static void
emit_store_global(struct ntv_context * ctx,nir_intrinsic_instr * intr)2676 emit_store_global(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2677 {
2678 bool coherent = ctx->sinfo->have_vulkan_memory_model && nir_intrinsic_access(intr) & ACCESS_COHERENT;
2679 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityPhysicalStorageBufferAddresses);
2680 unsigned bit_size = nir_src_bit_size(intr->src[0]);
2681 SpvId dest_type = get_uvec_type(ctx, bit_size, 1);
2682 SpvId pointer_type = spirv_builder_type_pointer(&ctx->builder,
2683 SpvStorageClassPhysicalStorageBuffer,
2684 dest_type);
2685 nir_alu_type atype;
2686 SpvId param = get_src(ctx, &intr->src[0], &atype);
2687 if (atype != nir_type_uint)
2688 param = emit_bitcast(ctx, dest_type, param);
2689 SpvId ptr = emit_bitcast(ctx, pointer_type, get_src(ctx, &intr->src[1], &atype));
2690 spirv_builder_emit_store_aligned(&ctx->builder, ptr, param, bit_size / 8, coherent);
2691 }
2692
2693 static void
emit_load_reg(struct ntv_context * ctx,nir_intrinsic_instr * intr)2694 emit_load_reg(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2695 {
2696 assert(nir_intrinsic_base(intr) == 0 && "no array registers");
2697
2698 nir_intrinsic_instr *decl = nir_reg_get_decl(intr->src[0].ssa);
2699 unsigned num_components = nir_intrinsic_num_components(decl);
2700 unsigned bit_size = nir_intrinsic_bit_size(decl);
2701 unsigned index = decl->def.index;
2702 assert(index < ctx->num_defs);
2703
2704 init_reg(ctx, decl, nir_type_uint);
2705 assert(ctx->defs[index] != 0);
2706
2707 nir_alu_type atype = ctx->def_types[index];
2708 SpvId var = ctx->defs[index];
2709 SpvId type = get_alu_type(ctx, atype, num_components, bit_size);
2710 SpvId result = spirv_builder_emit_load(&ctx->builder, type, var);
2711 store_def(ctx, intr->def.index, result, atype);
2712 }
2713
2714 static void
emit_store_reg(struct ntv_context * ctx,nir_intrinsic_instr * intr)2715 emit_store_reg(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2716 {
2717 nir_alu_type atype;
2718 SpvId param = get_src(ctx, &intr->src[0], &atype);
2719
2720 nir_intrinsic_instr *decl = nir_reg_get_decl(intr->src[1].ssa);
2721 unsigned index = decl->def.index;
2722 unsigned num_components = nir_intrinsic_num_components(decl);
2723 unsigned bit_size = nir_intrinsic_bit_size(decl);
2724
2725 atype = nir_alu_type_get_base_type(atype);
2726 init_reg(ctx, decl, atype);
2727 SpvId var = ctx->defs[index];
2728 nir_alu_type vtype = ctx->def_types[index];
2729 if (atype != vtype) {
2730 assert(vtype != nir_type_bool);
2731 param = emit_bitcast(ctx, get_alu_type(ctx, vtype, num_components, bit_size), param);
2732 }
2733 assert(var);
2734 spirv_builder_emit_store(&ctx->builder, var, param);
2735 }
2736
2737 static SpvId
create_builtin_var(struct ntv_context * ctx,SpvId var_type,SpvStorageClass storage_class,const char * name,SpvBuiltIn builtin)2738 create_builtin_var(struct ntv_context *ctx, SpvId var_type,
2739 SpvStorageClass storage_class,
2740 const char *name, SpvBuiltIn builtin)
2741 {
2742 SpvId pointer_type = spirv_builder_type_pointer(&ctx->builder,
2743 storage_class,
2744 var_type);
2745 SpvId var = spirv_builder_emit_var(&ctx->builder, pointer_type,
2746 storage_class);
2747 spirv_builder_emit_name(&ctx->builder, var, name);
2748 spirv_builder_emit_builtin(&ctx->builder, var, builtin);
2749
2750 if (ctx->stage == MESA_SHADER_FRAGMENT) {
2751 switch (builtin) {
2752 case SpvBuiltInSampleId:
2753 case SpvBuiltInSubgroupLocalInvocationId:
2754 spirv_builder_emit_decoration(&ctx->builder, var, SpvDecorationFlat);
2755 break;
2756 default:
2757 break;
2758 }
2759 }
2760
2761 assert(ctx->num_entry_ifaces < ARRAY_SIZE(ctx->entry_ifaces));
2762 ctx->entry_ifaces[ctx->num_entry_ifaces++] = var;
2763 return var;
2764 }
2765
2766 static void
emit_load_front_face(struct ntv_context * ctx,nir_intrinsic_instr * intr)2767 emit_load_front_face(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2768 {
2769 SpvId var_type = spirv_builder_type_bool(&ctx->builder);
2770 if (!ctx->front_face_var)
2771 ctx->front_face_var = create_builtin_var(ctx, var_type,
2772 SpvStorageClassInput,
2773 "gl_FrontFacing",
2774 SpvBuiltInFrontFacing);
2775
2776 SpvId result = spirv_builder_emit_load(&ctx->builder, var_type,
2777 ctx->front_face_var);
2778 assert(1 == intr->def.num_components);
2779 store_def(ctx, intr->def.index, result, nir_type_bool);
2780 }
2781
2782 static void
emit_load_view_index(struct ntv_context * ctx,nir_intrinsic_instr * intr)2783 emit_load_view_index(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2784 {
2785 SpvId var_type = spirv_builder_type_uint(&ctx->builder, 32);
2786 spirv_builder_emit_extension(&ctx->builder, "SPV_KHR_multiview");
2787 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityMultiView);
2788 if (!ctx->view_index_var)
2789 ctx->view_index_var = create_builtin_var(ctx, var_type,
2790 SpvStorageClassInput,
2791 "gl_ViewIndex",
2792 SpvBuiltInViewIndex);
2793
2794 SpvId result = spirv_builder_emit_load(&ctx->builder, var_type,
2795 ctx->view_index_var);
2796 assert(1 == intr->def.num_components);
2797 store_def(ctx, intr->def.index, result, nir_type_uint);
2798 }
2799
2800 static void
emit_load_uint_input(struct ntv_context * ctx,nir_intrinsic_instr * intr,SpvId * var_id,const char * var_name,SpvBuiltIn builtin)2801 emit_load_uint_input(struct ntv_context *ctx, nir_intrinsic_instr *intr, SpvId *var_id, const char *var_name, SpvBuiltIn builtin)
2802 {
2803 SpvId var_type = spirv_builder_type_uint(&ctx->builder, 32);
2804 if (!*var_id) {
2805 if (builtin == SpvBuiltInSampleMask) {
2806 /* gl_SampleMaskIn is an array[1] in spirv... */
2807 var_type = spirv_builder_type_array(&ctx->builder, var_type, emit_uint_const(ctx, 32, 1));
2808 spirv_builder_emit_array_stride(&ctx->builder, var_type, sizeof(uint32_t));
2809 }
2810 *var_id = create_builtin_var(ctx, var_type,
2811 SpvStorageClassInput,
2812 var_name,
2813 builtin);
2814 }
2815
2816 SpvId load_var = *var_id;
2817 if (builtin == SpvBuiltInSampleMask) {
2818 SpvId zero = emit_uint_const(ctx, 32, 0);
2819 var_type = spirv_builder_type_uint(&ctx->builder, 32);
2820 SpvId pointer_type = spirv_builder_type_pointer(&ctx->builder,
2821 SpvStorageClassInput,
2822 var_type);
2823 load_var = spirv_builder_emit_access_chain(&ctx->builder, pointer_type, load_var, &zero, 1);
2824 }
2825
2826 SpvId result = spirv_builder_emit_load(&ctx->builder, var_type, load_var);
2827 assert(1 == intr->def.num_components);
2828 store_def(ctx, intr->def.index, result, nir_type_uint);
2829 }
2830
2831 static void
emit_load_vec_input(struct ntv_context * ctx,nir_intrinsic_instr * intr,SpvId * var_id,const char * var_name,SpvBuiltIn builtin,nir_alu_type type)2832 emit_load_vec_input(struct ntv_context *ctx, nir_intrinsic_instr *intr, SpvId *var_id, const char *var_name, SpvBuiltIn builtin, nir_alu_type type)
2833 {
2834 SpvId var_type;
2835
2836 switch (type) {
2837 case nir_type_bool:
2838 var_type = get_bvec_type(ctx, intr->def.num_components);
2839 break;
2840 case nir_type_int:
2841 var_type = get_ivec_type(ctx, intr->def.bit_size,
2842 intr->def.num_components);
2843 break;
2844 case nir_type_uint:
2845 var_type = get_uvec_type(ctx, intr->def.bit_size,
2846 intr->def.num_components);
2847 break;
2848 case nir_type_float:
2849 var_type = get_fvec_type(ctx, intr->def.bit_size,
2850 intr->def.num_components);
2851 break;
2852 default:
2853 unreachable("unknown type passed");
2854 }
2855 if (!*var_id)
2856 *var_id = create_builtin_var(ctx, var_type,
2857 SpvStorageClassInput,
2858 var_name,
2859 builtin);
2860
2861 SpvId result = spirv_builder_emit_load(&ctx->builder, var_type, *var_id);
2862 store_def(ctx, intr->def.index, result, type);
2863 }
2864
2865 static void
emit_interpolate(struct ntv_context * ctx,nir_intrinsic_instr * intr)2866 emit_interpolate(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2867 {
2868 SpvId op;
2869 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityInterpolationFunction);
2870 SpvId src1 = 0;
2871 nir_alu_type atype;
2872 switch (intr->intrinsic) {
2873 case nir_intrinsic_interp_deref_at_centroid:
2874 op = GLSLstd450InterpolateAtCentroid;
2875 break;
2876 case nir_intrinsic_interp_deref_at_sample:
2877 op = GLSLstd450InterpolateAtSample;
2878 src1 = get_src(ctx, &intr->src[1], &atype);
2879 if (atype != nir_type_int)
2880 src1 = emit_bitcast(ctx, get_ivec_type(ctx, 32, 1), src1);
2881 break;
2882 case nir_intrinsic_interp_deref_at_offset:
2883 op = GLSLstd450InterpolateAtOffset;
2884 src1 = get_src(ctx, &intr->src[1], &atype);
2885 /*
2886 The offset operand must be a vector of 2 components of 32-bit floating-point type.
2887 - InterpolateAtOffset spec
2888 */
2889 if (atype != nir_type_float)
2890 src1 = emit_bitcast(ctx, get_fvec_type(ctx, 32, 2), src1);
2891 break;
2892 default:
2893 unreachable("unknown interp op");
2894 }
2895 nir_alu_type ptype;
2896 SpvId ptr = get_src(ctx, &intr->src[0], &ptype);
2897 SpvId result;
2898 const struct glsl_type *gtype = nir_src_as_deref(intr->src[0])->type;
2899 assert(glsl_get_vector_elements(gtype) == intr->num_components);
2900 assert(ptype == get_nir_alu_type(gtype));
2901 if (intr->intrinsic == nir_intrinsic_interp_deref_at_centroid)
2902 result = emit_builtin_unop(ctx, op, get_glsl_type(ctx, gtype), ptr);
2903 else
2904 result = emit_builtin_binop(ctx, op, get_glsl_type(ctx, gtype), ptr, src1);
2905 store_def(ctx, intr->def.index, result, ptype);
2906 }
2907
2908 static void
handle_atomic_op(struct ntv_context * ctx,nir_intrinsic_instr * intr,SpvId ptr,SpvId param,SpvId param2,nir_alu_type type)2909 handle_atomic_op(struct ntv_context *ctx, nir_intrinsic_instr *intr, SpvId ptr, SpvId param, SpvId param2, nir_alu_type type)
2910 {
2911 SpvId dest_type = get_def_type(ctx, &intr->def, type);
2912 SpvId result = emit_atomic(ctx,
2913 get_atomic_op(ctx, intr->def.bit_size, nir_intrinsic_atomic_op(intr)),
2914 dest_type, ptr, param, param2);
2915 assert(result);
2916 store_def(ctx, intr->def.index, result, type);
2917 }
2918
2919 static void
emit_deref_atomic_intrinsic(struct ntv_context * ctx,nir_intrinsic_instr * intr)2920 emit_deref_atomic_intrinsic(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2921 {
2922 nir_alu_type atype;
2923 nir_alu_type ret_type = nir_atomic_op_type(nir_intrinsic_atomic_op(intr)) == nir_type_float ? nir_type_float : nir_type_uint;
2924 SpvId ptr = get_src(ctx, &intr->src[0], &atype);
2925 if (atype != ret_type && ret_type == nir_type_float) {
2926 unsigned bit_size = nir_src_bit_size(intr->src[0]);
2927 SpvId *float_array_type = &ctx->float_array_type[bit_size == 32 ? 0 : 1];
2928 if (!*float_array_type) {
2929 *float_array_type = spirv_builder_type_pointer(&ctx->builder, SpvStorageClassStorageBuffer,
2930 spirv_builder_type_float(&ctx->builder, bit_size));
2931 }
2932 ptr = emit_unop(ctx, SpvOpBitcast, *float_array_type, ptr);
2933 }
2934
2935 SpvId param = get_src(ctx, &intr->src[1], &atype);
2936 if (atype != ret_type)
2937 param = cast_src_to_type(ctx, param, intr->src[1], ret_type);
2938
2939 SpvId param2 = 0;
2940
2941 if (nir_src_bit_size(intr->src[1]) == 64)
2942 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityInt64Atomics);
2943
2944 if (intr->intrinsic == nir_intrinsic_deref_atomic_swap) {
2945 param2 = get_src(ctx, &intr->src[2], &atype);
2946 if (atype != ret_type)
2947 param2 = cast_src_to_type(ctx, param2, intr->src[2], ret_type);
2948 }
2949
2950 handle_atomic_op(ctx, intr, ptr, param, param2, ret_type);
2951 }
2952
2953 static void
emit_shared_atomic_intrinsic(struct ntv_context * ctx,nir_intrinsic_instr * intr)2954 emit_shared_atomic_intrinsic(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2955 {
2956 unsigned bit_size = nir_src_bit_size(intr->src[1]);
2957 SpvId dest_type = get_def_type(ctx, &intr->def, nir_type_uint);
2958 nir_alu_type atype;
2959 nir_alu_type ret_type = nir_atomic_op_type(nir_intrinsic_atomic_op(intr)) == nir_type_float ? nir_type_float : nir_type_uint;
2960 SpvId param = get_src(ctx, &intr->src[1], &atype);
2961 if (atype != ret_type)
2962 param = cast_src_to_type(ctx, param, intr->src[1], ret_type);
2963
2964 SpvId pointer_type = spirv_builder_type_pointer(&ctx->builder,
2965 SpvStorageClassWorkgroup,
2966 dest_type);
2967 SpvId offset = get_src(ctx, &intr->src[0], &atype);
2968 if (atype != nir_type_uint)
2969 offset = cast_src_to_type(ctx, offset, intr->src[0], nir_type_uint);
2970 offset = emit_binop(ctx, SpvOpUDiv, get_uvec_type(ctx, 32, 1), offset, emit_uint_const(ctx, 32, bit_size / 8));
2971 SpvId shared_block = get_shared_block(ctx, bit_size);
2972 SpvId ptr = spirv_builder_emit_access_chain(&ctx->builder, pointer_type,
2973 shared_block, &offset, 1);
2974 if (nir_src_bit_size(intr->src[1]) == 64)
2975 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityInt64Atomics);
2976 SpvId param2 = 0;
2977
2978 if (intr->intrinsic == nir_intrinsic_shared_atomic_swap) {
2979 param2 = get_src(ctx, &intr->src[2], &atype);
2980 if (atype != ret_type)
2981 param2 = cast_src_to_type(ctx, param2, intr->src[2], ret_type);
2982 }
2983
2984 handle_atomic_op(ctx, intr, ptr, param, param2, ret_type);
2985 }
2986
2987 static void
emit_global_atomic_intrinsic(struct ntv_context * ctx,nir_intrinsic_instr * intr)2988 emit_global_atomic_intrinsic(struct ntv_context *ctx, nir_intrinsic_instr *intr)
2989 {
2990 unsigned bit_size = nir_src_bit_size(intr->src[1]);
2991 SpvId dest_type = get_def_type(ctx, &intr->def, nir_type_uint);
2992 nir_alu_type atype;
2993 nir_alu_type ret_type = nir_atomic_op_type(nir_intrinsic_atomic_op(intr)) == nir_type_float ? nir_type_float : nir_type_uint;
2994 SpvId param = get_src(ctx, &intr->src[1], &atype);
2995
2996 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityPhysicalStorageBufferAddresses);
2997 SpvId pointer_type = spirv_builder_type_pointer(&ctx->builder,
2998 SpvStorageClassPhysicalStorageBuffer,
2999 dest_type);
3000 SpvId ptr = emit_bitcast(ctx, pointer_type, get_src(ctx, &intr->src[0], &atype));
3001
3002 if (bit_size == 64)
3003 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityInt64Atomics);
3004 SpvId param2 = 0;
3005
3006 if (intr->intrinsic == nir_intrinsic_global_atomic_swap)
3007 param2 = get_src(ctx, &intr->src[2], &atype);
3008
3009 handle_atomic_op(ctx, intr, ptr, param, param2, ret_type);
3010 }
3011
3012 static void
emit_get_ssbo_size(struct ntv_context * ctx,nir_intrinsic_instr * intr)3013 emit_get_ssbo_size(struct ntv_context *ctx, nir_intrinsic_instr *intr)
3014 {
3015 SpvId uint_type = get_uvec_type(ctx, 32, 1);
3016 nir_variable *var = ctx->ssbo_vars;
3017 const struct glsl_type *bare_type = glsl_without_array(var->type);
3018 unsigned last_member_idx = glsl_get_length(bare_type) - 1;
3019 SpvId pointer_type = spirv_builder_type_pointer(&ctx->builder,
3020 SpvStorageClassStorageBuffer,
3021 get_bo_struct_type(ctx, var));
3022 nir_alu_type atype;
3023 SpvId bo = get_src(ctx, &intr->src[0], &atype);
3024 if (atype == nir_type_float)
3025 bo = bitcast_to_uvec(ctx, bo, nir_src_bit_size(intr->src[0]), 1);
3026 SpvId indices[] = { bo };
3027 SpvId ptr = spirv_builder_emit_access_chain(&ctx->builder, pointer_type,
3028 ctx->ssbos[2], indices,
3029 ARRAY_SIZE(indices));
3030 SpvId result = spirv_builder_emit_binop(&ctx->builder, SpvOpArrayLength, uint_type,
3031 ptr, last_member_idx);
3032 /* this is going to be converted by nir to:
3033
3034 length = (buffer_size - offset) / stride
3035
3036 * so we need to un-convert it to avoid having the calculation performed twice
3037 */
3038 const struct glsl_type *last_member = glsl_get_struct_field(bare_type, last_member_idx);
3039 /* multiply by stride */
3040 result = emit_binop(ctx, SpvOpIMul, uint_type, result, emit_uint_const(ctx, 32, glsl_get_explicit_stride(last_member)));
3041 /* get total ssbo size by adding offset */
3042 result = emit_binop(ctx, SpvOpIAdd, uint_type, result,
3043 emit_uint_const(ctx, 32,
3044 glsl_get_struct_field_offset(bare_type, last_member_idx)));
3045 store_def(ctx, intr->def.index, result, nir_type_uint);
3046 }
3047
3048 static SpvId
get_image_coords(struct ntv_context * ctx,const struct glsl_type * type,nir_src * src)3049 get_image_coords(struct ntv_context *ctx, const struct glsl_type *type, nir_src *src)
3050 {
3051 uint32_t num_coords = glsl_get_sampler_coordinate_components(type);
3052 uint32_t src_components = nir_src_num_components(*src);
3053
3054 nir_alu_type atype;
3055 SpvId spv = get_src(ctx, src, &atype);
3056 if (num_coords == src_components)
3057 return spv;
3058
3059 /* need to extract the coord dimensions that the image can use */
3060 SpvId vec_type = get_alu_type(ctx, atype, num_coords, 32);
3061 if (num_coords == 1)
3062 return spirv_builder_emit_vector_extract(&ctx->builder, vec_type, spv, 0);
3063 uint32_t constituents[4];
3064 SpvId zero = atype == nir_type_uint ? emit_uint_const(ctx, nir_src_bit_size(*src), 0) : emit_float_const(ctx, nir_src_bit_size(*src), 0);
3065 assert(num_coords < ARRAY_SIZE(constituents));
3066 for (unsigned i = 0; i < num_coords; i++)
3067 constituents[i] = i < src_components ? i : zero;
3068 return spirv_builder_emit_vector_shuffle(&ctx->builder, vec_type, spv, spv, constituents, num_coords);
3069 }
3070
3071 static void
emit_image_deref_store(struct ntv_context * ctx,nir_intrinsic_instr * intr)3072 emit_image_deref_store(struct ntv_context *ctx, nir_intrinsic_instr *intr)
3073 {
3074 nir_alu_type atype;
3075 SpvId img_var = get_src(ctx, &intr->src[0], &atype);
3076 nir_deref_instr *deref = nir_src_as_deref(intr->src[0]);
3077 nir_variable *var = nir_deref_instr_get_variable(deref);
3078 SpvId img_type = find_image_type(ctx, var);
3079 const struct glsl_type *type = glsl_without_array(var->type);
3080 SpvId base_type = get_glsl_basetype(ctx, glsl_get_sampler_result_type(type));
3081 SpvId img = spirv_builder_emit_load(&ctx->builder, img_type, img_var);
3082 SpvId coord = get_image_coords(ctx, type, &intr->src[1]);
3083 SpvId texel = get_src(ctx, &intr->src[3], &atype);
3084 /* texel type must match image type */
3085 if (atype != nir_get_nir_type_for_glsl_base_type(glsl_get_sampler_result_type(type)))
3086 texel = emit_bitcast(ctx,
3087 spirv_builder_type_vector(&ctx->builder, base_type, 4),
3088 texel);
3089 bool use_sample = glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_MS ||
3090 glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_SUBPASS_MS;
3091 SpvId sample = use_sample ? get_src(ctx, &intr->src[2], &atype) : 0;
3092 assert(nir_src_bit_size(intr->src[3]) == glsl_base_type_bit_size(glsl_get_sampler_result_type(type)));
3093 spirv_builder_emit_image_write(&ctx->builder, img, coord, texel, 0, sample, 0);
3094 }
3095
3096 static SpvId
extract_sparse_load(struct ntv_context * ctx,SpvId result,SpvId dest_type,nir_def * def)3097 extract_sparse_load(struct ntv_context *ctx, SpvId result, SpvId dest_type, nir_def *def)
3098 {
3099 /* Result Type must be an OpTypeStruct with two members.
3100 * The first member’s type must be an integer type scalar.
3101 * It holds a Residency Code that can be passed to OpImageSparseTexelsResident
3102 * - OpImageSparseRead spec
3103 */
3104 uint32_t idx = 0;
3105 SpvId resident = spirv_builder_emit_composite_extract(&ctx->builder, spirv_builder_type_uint(&ctx->builder, 32), result, &idx, 1);
3106 idx = 1;
3107 /* normal vec4 return */
3108 if (def->num_components == 4)
3109 result = spirv_builder_emit_composite_extract(&ctx->builder, dest_type, result, &idx, 1);
3110 else {
3111 /* shadow */
3112 assert(def->num_components == 1);
3113 SpvId type = spirv_builder_type_float(&ctx->builder, def->bit_size);
3114 SpvId val[2];
3115 /* pad to 2 components: the upcoming is_sparse_texels_resident instr will always use the
3116 * separate residency value, but the shader still expects this return to be a vec2,
3117 * so give it a vec2
3118 */
3119 val[0] = spirv_builder_emit_composite_extract(&ctx->builder, type, result, &idx, 1);
3120 val[1] = emit_float_const(ctx, def->bit_size, 0);
3121 result = spirv_builder_emit_composite_construct(&ctx->builder, get_fvec_type(ctx, def->bit_size, 2), val, 2);
3122 }
3123 assert(resident != 0);
3124 assert(def->index < ctx->num_defs);
3125 ctx->resident_defs[def->index] = resident;
3126 return result;
3127 }
3128
3129 static void
emit_image_deref_load(struct ntv_context * ctx,nir_intrinsic_instr * intr)3130 emit_image_deref_load(struct ntv_context *ctx, nir_intrinsic_instr *intr)
3131 {
3132 bool sparse = intr->intrinsic == nir_intrinsic_image_deref_sparse_load;
3133 nir_alu_type atype;
3134 SpvId img_var = get_src(ctx, &intr->src[0], &atype);
3135 nir_deref_instr *deref = nir_src_as_deref(intr->src[0]);
3136 nir_variable *var = nir_deref_instr_get_variable(deref);
3137 bool mediump = (var->data.precision == GLSL_PRECISION_MEDIUM || var->data.precision == GLSL_PRECISION_LOW);
3138 SpvId img_type = find_image_type(ctx, var);
3139 const struct glsl_type *type = glsl_without_array(var->type);
3140 SpvId base_type = get_glsl_basetype(ctx, glsl_get_sampler_result_type(type));
3141 SpvId img = spirv_builder_emit_load(&ctx->builder, img_type, img_var);
3142 SpvId coord = get_image_coords(ctx, type, &intr->src[1]);
3143 bool use_sample = glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_MS ||
3144 glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_SUBPASS_MS;
3145 SpvId sample = use_sample ? get_src(ctx, &intr->src[2], &atype) : 0;
3146 SpvId dest_type = spirv_builder_type_vector(&ctx->builder, base_type,
3147 intr->def.num_components);
3148 SpvId result = spirv_builder_emit_image_read(&ctx->builder,
3149 dest_type,
3150 img, coord, 0, sample, 0, sparse);
3151 if (sparse)
3152 result = extract_sparse_load(ctx, result, dest_type, &intr->def);
3153
3154 if (!sparse && mediump) {
3155 spirv_builder_emit_decoration(&ctx->builder, result,
3156 SpvDecorationRelaxedPrecision);
3157 }
3158
3159 store_def(ctx, intr->def.index, result, nir_get_nir_type_for_glsl_base_type(glsl_get_sampler_result_type(type)));
3160 }
3161
3162 static void
emit_image_deref_size(struct ntv_context * ctx,nir_intrinsic_instr * intr)3163 emit_image_deref_size(struct ntv_context *ctx, nir_intrinsic_instr *intr)
3164 {
3165 nir_alu_type atype;
3166 SpvId img_var = get_src(ctx, &intr->src[0], &atype);
3167 nir_deref_instr *deref = nir_src_as_deref(intr->src[0]);
3168 nir_variable *var = nir_deref_instr_get_variable(deref);
3169 SpvId img_type = find_image_type(ctx, var);
3170 const struct glsl_type *type = glsl_without_array(var->type);
3171 SpvId img = spirv_builder_emit_load(&ctx->builder, img_type, img_var);
3172 unsigned num_components = glsl_get_sampler_coordinate_components(type);
3173 /* SPIRV requires 2 components for non-array cube size */
3174 if (glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_CUBE && !glsl_sampler_type_is_array(type))
3175 num_components = 2;
3176
3177 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityImageQuery);
3178 SpvId result = spirv_builder_emit_image_query_size(&ctx->builder, get_uvec_type(ctx, 32, num_components), img, 0);
3179 store_def(ctx, intr->def.index, result, nir_type_uint);
3180 }
3181
3182 static void
emit_image_deref_samples(struct ntv_context * ctx,nir_intrinsic_instr * intr)3183 emit_image_deref_samples(struct ntv_context *ctx, nir_intrinsic_instr *intr)
3184 {
3185 nir_alu_type atype;
3186 SpvId img_var = get_src(ctx, &intr->src[0], &atype);
3187 nir_deref_instr *deref = nir_src_as_deref(intr->src[0]);
3188 nir_variable *var = nir_deref_instr_get_variable(deref);
3189 SpvId img_type = find_image_type(ctx, var);
3190 SpvId img = spirv_builder_emit_load(&ctx->builder, img_type, img_var);
3191
3192 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityImageQuery);
3193 SpvId result = spirv_builder_emit_unop(&ctx->builder, SpvOpImageQuerySamples, get_def_type(ctx, &intr->def, nir_type_uint), img);
3194 store_def(ctx, intr->def.index, result, nir_type_uint);
3195 }
3196
3197 static void
emit_image_intrinsic(struct ntv_context * ctx,nir_intrinsic_instr * intr)3198 emit_image_intrinsic(struct ntv_context *ctx, nir_intrinsic_instr *intr)
3199 {
3200 nir_alu_type atype, ptype;
3201 SpvId param = get_src(ctx, &intr->src[3], &ptype);
3202 SpvId img_var = get_src(ctx, &intr->src[0], &atype);
3203 nir_deref_instr *deref = nir_src_as_deref(intr->src[0]);
3204 nir_variable *var = nir_deref_instr_get_variable(deref);
3205 const struct glsl_type *type = glsl_without_array(var->type);
3206 bool is_ms;
3207 type_to_dim(glsl_get_sampler_dim(type), &is_ms);
3208 SpvId sample = is_ms ? get_src(ctx, &intr->src[2], &atype) : emit_uint_const(ctx, 32, 0);
3209 SpvId coord = get_image_coords(ctx, type, &intr->src[1]);
3210 enum glsl_base_type glsl_result_type = glsl_get_sampler_result_type(type);
3211 SpvId base_type = get_glsl_basetype(ctx, glsl_result_type);
3212 SpvId texel = spirv_builder_emit_image_texel_pointer(&ctx->builder, base_type, img_var, coord, sample);
3213 SpvId param2 = 0;
3214
3215 /* The type of Value must be the same as Result Type.
3216 * The type of the value pointed to by Pointer must be the same as Result Type.
3217 */
3218 nir_alu_type ntype = nir_get_nir_type_for_glsl_base_type(glsl_result_type);
3219 if (ptype != ntype) {
3220 SpvId cast_type = get_def_type(ctx, &intr->def, ntype);
3221 param = emit_bitcast(ctx, cast_type, param);
3222 }
3223
3224 if (intr->intrinsic == nir_intrinsic_image_deref_atomic_swap) {
3225 param2 = get_src(ctx, &intr->src[4], &ptype);
3226 if (ptype != ntype) {
3227 SpvId cast_type = get_def_type(ctx, &intr->def, ntype);
3228 param2 = emit_bitcast(ctx, cast_type, param2);
3229 }
3230 }
3231
3232 handle_atomic_op(ctx, intr, texel, param, param2, ntype);
3233 }
3234
3235 static void
emit_ballot(struct ntv_context * ctx,nir_intrinsic_instr * intr)3236 emit_ballot(struct ntv_context *ctx, nir_intrinsic_instr *intr)
3237 {
3238 spirv_builder_emit_cap(&ctx->builder, SpvCapabilitySubgroupBallotKHR);
3239 spirv_builder_emit_extension(&ctx->builder, "SPV_KHR_shader_ballot");
3240 SpvId type = get_def_uvec_type(ctx, &intr->def);
3241 nir_alu_type atype;
3242 SpvId result = emit_unop(ctx, SpvOpSubgroupBallotKHR, type, get_src(ctx, &intr->src[0], &atype));
3243 store_def(ctx, intr->def.index, result, nir_type_uint);
3244 }
3245
3246 static void
emit_read_first_invocation(struct ntv_context * ctx,nir_intrinsic_instr * intr)3247 emit_read_first_invocation(struct ntv_context *ctx, nir_intrinsic_instr *intr)
3248 {
3249 spirv_builder_emit_cap(&ctx->builder, SpvCapabilitySubgroupBallotKHR);
3250 spirv_builder_emit_extension(&ctx->builder, "SPV_KHR_shader_ballot");
3251 nir_alu_type atype;
3252 SpvId src = get_src(ctx, &intr->src[0], &atype);
3253 SpvId type = get_def_type(ctx, &intr->def, atype);
3254 SpvId result = emit_unop(ctx, SpvOpSubgroupFirstInvocationKHR, type, src);
3255 store_def(ctx, intr->def.index, result, atype);
3256 }
3257
3258 static void
emit_read_invocation(struct ntv_context * ctx,nir_intrinsic_instr * intr)3259 emit_read_invocation(struct ntv_context *ctx, nir_intrinsic_instr *intr)
3260 {
3261 spirv_builder_emit_cap(&ctx->builder, SpvCapabilitySubgroupBallotKHR);
3262 spirv_builder_emit_extension(&ctx->builder, "SPV_KHR_shader_ballot");
3263 nir_alu_type atype, itype;
3264 SpvId src = get_src(ctx, &intr->src[0], &atype);
3265 SpvId type = get_def_type(ctx, &intr->def, atype);
3266 SpvId result = emit_binop(ctx, SpvOpSubgroupReadInvocationKHR, type,
3267 src,
3268 get_src(ctx, &intr->src[1], &itype));
3269 store_def(ctx, intr->def.index, result, atype);
3270 }
3271
3272 static void
emit_shader_clock(struct ntv_context * ctx,nir_intrinsic_instr * intr)3273 emit_shader_clock(struct ntv_context *ctx, nir_intrinsic_instr *intr)
3274 {
3275 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityShaderClockKHR);
3276 spirv_builder_emit_extension(&ctx->builder, "SPV_KHR_shader_clock");
3277
3278 SpvScope scope = get_scope(nir_intrinsic_memory_scope(intr));
3279 SpvId type = get_def_type(ctx, &intr->def, nir_type_uint);
3280 SpvId result = spirv_builder_emit_unop_const(&ctx->builder, SpvOpReadClockKHR, type, scope);
3281 store_def(ctx, intr->def.index, result, nir_type_uint);
3282 }
3283
3284 static void
emit_is_sparse_texels_resident(struct ntv_context * ctx,nir_intrinsic_instr * intr)3285 emit_is_sparse_texels_resident(struct ntv_context *ctx, nir_intrinsic_instr *intr)
3286 {
3287 spirv_builder_emit_cap(&ctx->builder, SpvCapabilitySparseResidency);
3288
3289 SpvId type = get_def_type(ctx, &intr->def, nir_type_uint);
3290
3291 unsigned index = intr->src[0].ssa->index;
3292 assert(index < ctx->num_defs);
3293 assert(ctx->resident_defs[index] != 0);
3294 SpvId resident = ctx->resident_defs[index];
3295
3296 SpvId result = spirv_builder_emit_unop(&ctx->builder, SpvOpImageSparseTexelsResident, type, resident);
3297 store_def(ctx, intr->def.index, result, nir_type_uint);
3298 }
3299
3300 static void
emit_vote(struct ntv_context * ctx,nir_intrinsic_instr * intr)3301 emit_vote(struct ntv_context *ctx, nir_intrinsic_instr *intr)
3302 {
3303 SpvOp op;
3304
3305 switch (intr->intrinsic) {
3306 case nir_intrinsic_vote_all:
3307 op = SpvOpGroupNonUniformAll;
3308 break;
3309 case nir_intrinsic_vote_any:
3310 op = SpvOpGroupNonUniformAny;
3311 break;
3312 case nir_intrinsic_vote_ieq:
3313 case nir_intrinsic_vote_feq:
3314 op = SpvOpGroupNonUniformAllEqual;
3315 break;
3316 default:
3317 unreachable("unknown vote intrinsic");
3318 }
3319 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityGroupNonUniformVote);
3320 nir_alu_type atype;
3321 SpvId result = spirv_builder_emit_vote(&ctx->builder, op, get_src(ctx, &intr->src[0], &atype));
3322 store_def(ctx, intr->def.index, result, nir_type_bool);
3323 }
3324
3325 static void
emit_is_helper_invocation(struct ntv_context * ctx,nir_intrinsic_instr * intr)3326 emit_is_helper_invocation(struct ntv_context *ctx, nir_intrinsic_instr *intr)
3327 {
3328 spirv_builder_emit_extension(&ctx->builder,
3329 "SPV_EXT_demote_to_helper_invocation");
3330 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityDemoteToHelperInvocation);
3331 SpvId result = spirv_is_helper_invocation(&ctx->builder);
3332 store_def(ctx, intr->def.index, result, nir_type_bool);
3333 }
3334
3335 static void
emit_barrier(struct ntv_context * ctx,nir_intrinsic_instr * intr)3336 emit_barrier(struct ntv_context *ctx, nir_intrinsic_instr *intr)
3337 {
3338 SpvScope scope = get_scope(nir_intrinsic_execution_scope(intr));
3339 SpvScope mem_scope = get_scope(nir_intrinsic_memory_scope(intr));
3340 SpvMemorySemanticsMask semantics = 0;
3341
3342 if (nir_intrinsic_memory_scope(intr) != SCOPE_NONE) {
3343 nir_variable_mode modes = nir_intrinsic_memory_modes(intr);
3344
3345 if (modes & nir_var_image)
3346 semantics |= SpvMemorySemanticsImageMemoryMask;
3347
3348 if (modes & nir_var_mem_shared)
3349 semantics |= SpvMemorySemanticsWorkgroupMemoryMask;
3350
3351 if (modes & (nir_var_mem_ssbo | nir_var_mem_global))
3352 semantics |= SpvMemorySemanticsUniformMemoryMask;
3353
3354 if (modes & nir_var_mem_global)
3355 semantics |= SpvMemorySemanticsCrossWorkgroupMemoryMask;
3356
3357 if (modes & (nir_var_shader_out | nir_var_mem_task_payload))
3358 semantics |= SpvMemorySemanticsOutputMemoryMask;
3359
3360 if (!modes)
3361 semantics = SpvMemorySemanticsWorkgroupMemoryMask |
3362 SpvMemorySemanticsUniformMemoryMask |
3363 SpvMemorySemanticsImageMemoryMask |
3364 SpvMemorySemanticsCrossWorkgroupMemoryMask;
3365 semantics |= SpvMemorySemanticsAcquireReleaseMask;
3366 }
3367
3368 if (nir_intrinsic_execution_scope(intr) != SCOPE_NONE)
3369 spirv_builder_emit_control_barrier(&ctx->builder, scope, mem_scope, semantics);
3370 else
3371 spirv_builder_emit_memory_barrier(&ctx->builder, mem_scope, semantics);
3372 }
3373
3374 static void
emit_derivative(struct ntv_context * ctx,nir_intrinsic_instr * intr)3375 emit_derivative(struct ntv_context *ctx, nir_intrinsic_instr *intr)
3376 {
3377 SpvOp op;
3378 switch (intr->intrinsic) {
3379 case nir_intrinsic_ddx:
3380 op = SpvOpDPdx;
3381 break;
3382 case nir_intrinsic_ddy:
3383 op = SpvOpDPdy;
3384 break;
3385 case nir_intrinsic_ddx_fine:
3386 op = SpvOpDPdxFine;
3387 break;
3388 case nir_intrinsic_ddy_fine:
3389 op = SpvOpDPdyFine;
3390 break;
3391 case nir_intrinsic_ddx_coarse:
3392 op = SpvOpDPdxCoarse;
3393 break;
3394 case nir_intrinsic_ddy_coarse:
3395 op = SpvOpDPdyCoarse;
3396 break;
3397 default:
3398 unreachable("invalid ddx/ddy");
3399 }
3400
3401 if (op != SpvOpDPdx && op != SpvOpDPdy)
3402 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityDerivativeControl);
3403
3404 SpvId type = get_fvec_type(ctx, intr->def.bit_size, intr->def.num_components);
3405
3406 nir_alu_type atype;
3407 SpvId value = get_src(ctx, &intr->src[0], &atype);
3408 if (atype != nir_type_float)
3409 value = emit_bitcast(ctx, type, value);
3410
3411 SpvId result = emit_unop(ctx, op, type, value);
3412 store_def(ctx, intr->def.index, result, nir_type_float);
3413 }
3414
3415 static void
emit_intrinsic(struct ntv_context * ctx,nir_intrinsic_instr * intr)3416 emit_intrinsic(struct ntv_context *ctx, nir_intrinsic_instr *intr)
3417 {
3418 switch (intr->intrinsic) {
3419 case nir_intrinsic_decl_reg:
3420 /* Nothing to do */
3421 break;
3422
3423 case nir_intrinsic_load_reg:
3424 emit_load_reg(ctx, intr);
3425 break;
3426
3427 case nir_intrinsic_store_reg:
3428 emit_store_reg(ctx, intr);
3429 break;
3430
3431 case nir_intrinsic_terminate:
3432 emit_discard(ctx, intr);
3433 break;
3434
3435 case nir_intrinsic_demote:
3436 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityDemoteToHelperInvocation);
3437 spirv_builder_emit_demote(&ctx->builder);
3438 break;
3439
3440 case nir_intrinsic_load_deref:
3441 emit_load_deref(ctx, intr);
3442 break;
3443
3444 case nir_intrinsic_store_deref:
3445 emit_store_deref(ctx, intr);
3446 break;
3447
3448 case nir_intrinsic_load_push_constant_zink:
3449 emit_load_push_const(ctx, intr);
3450 break;
3451
3452 case nir_intrinsic_load_global:
3453 case nir_intrinsic_load_global_constant:
3454 emit_load_global(ctx, intr);
3455 break;
3456
3457 case nir_intrinsic_store_global:
3458 emit_store_global(ctx, intr);
3459 break;
3460
3461 case nir_intrinsic_load_front_face:
3462 emit_load_front_face(ctx, intr);
3463 break;
3464
3465 case nir_intrinsic_load_view_index:
3466 emit_load_view_index(ctx, intr);
3467 break;
3468
3469 case nir_intrinsic_load_base_instance:
3470 emit_load_uint_input(ctx, intr, &ctx->base_instance_var, "gl_BaseInstance", SpvBuiltInBaseInstance);
3471 break;
3472
3473 case nir_intrinsic_load_instance_id:
3474 emit_load_uint_input(ctx, intr, &ctx->instance_id_var, "gl_InstanceId", SpvBuiltInInstanceIndex);
3475 break;
3476
3477 case nir_intrinsic_load_base_vertex:
3478 emit_load_uint_input(ctx, intr, &ctx->base_vertex_var, "gl_BaseVertex", SpvBuiltInBaseVertex);
3479 break;
3480
3481 case nir_intrinsic_load_draw_id:
3482 emit_load_uint_input(ctx, intr, &ctx->draw_id_var, "gl_DrawID", SpvBuiltInDrawIndex);
3483 break;
3484
3485 case nir_intrinsic_load_vertex_id:
3486 emit_load_uint_input(ctx, intr, &ctx->vertex_id_var, "gl_VertexId", SpvBuiltInVertexIndex);
3487 break;
3488
3489 case nir_intrinsic_load_primitive_id:
3490 emit_load_uint_input(ctx, intr, &ctx->primitive_id_var, "gl_PrimitiveIdIn", SpvBuiltInPrimitiveId);
3491 break;
3492
3493 case nir_intrinsic_load_invocation_id:
3494 emit_load_uint_input(ctx, intr, &ctx->invocation_id_var, "gl_InvocationId", SpvBuiltInInvocationId);
3495 break;
3496
3497 case nir_intrinsic_load_sample_id:
3498 spirv_builder_emit_cap(&ctx->builder, SpvCapabilitySampleRateShading);
3499 emit_load_uint_input(ctx, intr, &ctx->sample_id_var, "gl_SampleId", SpvBuiltInSampleId);
3500 break;
3501
3502 case nir_intrinsic_load_point_coord_maybe_flipped:
3503 case nir_intrinsic_load_point_coord:
3504 emit_load_vec_input(ctx, intr, &ctx->point_coord_var, "gl_PointCoord", SpvBuiltInPointCoord, nir_type_float);
3505 break;
3506
3507 case nir_intrinsic_load_sample_pos:
3508 emit_load_vec_input(ctx, intr, &ctx->sample_pos_var, "gl_SamplePosition", SpvBuiltInSamplePosition, nir_type_float);
3509 break;
3510
3511 case nir_intrinsic_load_sample_mask_in:
3512 emit_load_uint_input(ctx, intr, &ctx->sample_mask_in_var, "gl_SampleMaskIn", SpvBuiltInSampleMask);
3513 break;
3514
3515 case nir_intrinsic_emit_vertex:
3516 if (ctx->nir->info.gs.vertices_out) //skip vertex emission if !vertices_out
3517 spirv_builder_emit_vertex(&ctx->builder, nir_intrinsic_stream_id(intr),
3518 ctx->nir->info.stage == MESA_SHADER_GEOMETRY && util_bitcount(ctx->nir->info.gs.active_stream_mask) > 1);
3519 break;
3520
3521 case nir_intrinsic_end_primitive:
3522 spirv_builder_end_primitive(&ctx->builder, nir_intrinsic_stream_id(intr),
3523 ctx->nir->info.stage == MESA_SHADER_GEOMETRY && util_bitcount(ctx->nir->info.gs.active_stream_mask) > 1);
3524 break;
3525
3526 case nir_intrinsic_load_helper_invocation:
3527 emit_load_vec_input(ctx, intr, &ctx->helper_invocation_var, "gl_HelperInvocation", SpvBuiltInHelperInvocation, nir_type_bool);
3528 break;
3529
3530 case nir_intrinsic_load_patch_vertices_in:
3531 emit_load_vec_input(ctx, intr, &ctx->tess_patch_vertices_in, "gl_PatchVerticesIn",
3532 SpvBuiltInPatchVertices, nir_type_int);
3533 break;
3534
3535 case nir_intrinsic_load_tess_coord:
3536 emit_load_vec_input(ctx, intr, &ctx->tess_coord_var, "gl_TessCoord",
3537 SpvBuiltInTessCoord, nir_type_float);
3538 break;
3539
3540 case nir_intrinsic_barrier:
3541 emit_barrier(ctx, intr);
3542 break;
3543
3544 case nir_intrinsic_interp_deref_at_centroid:
3545 case nir_intrinsic_interp_deref_at_sample:
3546 case nir_intrinsic_interp_deref_at_offset:
3547 emit_interpolate(ctx, intr);
3548 break;
3549
3550 case nir_intrinsic_deref_atomic:
3551 case nir_intrinsic_deref_atomic_swap:
3552 emit_deref_atomic_intrinsic(ctx, intr);
3553 break;
3554
3555 case nir_intrinsic_shared_atomic:
3556 case nir_intrinsic_shared_atomic_swap:
3557 emit_shared_atomic_intrinsic(ctx, intr);
3558 break;
3559
3560 case nir_intrinsic_global_atomic:
3561 case nir_intrinsic_global_atomic_swap:
3562 emit_global_atomic_intrinsic(ctx, intr);
3563 break;
3564
3565 case nir_intrinsic_begin_invocation_interlock:
3566 case nir_intrinsic_end_invocation_interlock:
3567 spirv_builder_emit_interlock(&ctx->builder, intr->intrinsic == nir_intrinsic_end_invocation_interlock);
3568 break;
3569
3570 case nir_intrinsic_get_ssbo_size:
3571 emit_get_ssbo_size(ctx, intr);
3572 break;
3573
3574 case nir_intrinsic_image_deref_store:
3575 emit_image_deref_store(ctx, intr);
3576 break;
3577
3578 case nir_intrinsic_image_deref_sparse_load:
3579 case nir_intrinsic_image_deref_load:
3580 emit_image_deref_load(ctx, intr);
3581 break;
3582
3583 case nir_intrinsic_image_deref_size:
3584 emit_image_deref_size(ctx, intr);
3585 break;
3586
3587 case nir_intrinsic_image_deref_samples:
3588 emit_image_deref_samples(ctx, intr);
3589 break;
3590
3591 case nir_intrinsic_image_deref_atomic:
3592 case nir_intrinsic_image_deref_atomic_swap:
3593 emit_image_intrinsic(ctx, intr);
3594 break;
3595
3596 case nir_intrinsic_load_workgroup_id:
3597 emit_load_vec_input(ctx, intr, &ctx->workgroup_id_var, "gl_WorkGroupID", SpvBuiltInWorkgroupId, nir_type_uint);
3598 break;
3599
3600 case nir_intrinsic_load_num_workgroups:
3601 emit_load_vec_input(ctx, intr, &ctx->num_workgroups_var, "gl_NumWorkGroups", SpvBuiltInNumWorkgroups, nir_type_uint);
3602 break;
3603
3604 case nir_intrinsic_load_local_invocation_id:
3605 emit_load_vec_input(ctx, intr, &ctx->local_invocation_id_var, "gl_LocalInvocationID", SpvBuiltInLocalInvocationId, nir_type_uint);
3606 break;
3607
3608 case nir_intrinsic_load_global_invocation_id:
3609 emit_load_vec_input(ctx, intr, &ctx->global_invocation_id_var, "gl_GlobalInvocationID", SpvBuiltInGlobalInvocationId, nir_type_uint);
3610 break;
3611
3612 case nir_intrinsic_load_local_invocation_index:
3613 emit_load_uint_input(ctx, intr, &ctx->local_invocation_index_var, "gl_LocalInvocationIndex", SpvBuiltInLocalInvocationIndex);
3614 break;
3615
3616 #define LOAD_SHADER_BALLOT(lowercase, camelcase) \
3617 case nir_intrinsic_load_##lowercase: \
3618 emit_load_uint_input(ctx, intr, &ctx->lowercase##_var, "gl_"#camelcase, SpvBuiltIn##camelcase); \
3619 break
3620
3621 LOAD_SHADER_BALLOT(subgroup_id, SubgroupId);
3622 LOAD_SHADER_BALLOT(subgroup_eq_mask, SubgroupEqMask);
3623 LOAD_SHADER_BALLOT(subgroup_ge_mask, SubgroupGeMask);
3624 LOAD_SHADER_BALLOT(subgroup_invocation, SubgroupLocalInvocationId);
3625 LOAD_SHADER_BALLOT(subgroup_le_mask, SubgroupLeMask);
3626 LOAD_SHADER_BALLOT(subgroup_lt_mask, SubgroupLtMask);
3627 LOAD_SHADER_BALLOT(subgroup_size, SubgroupSize);
3628
3629 case nir_intrinsic_ballot:
3630 emit_ballot(ctx, intr);
3631 break;
3632
3633 case nir_intrinsic_read_first_invocation:
3634 emit_read_first_invocation(ctx, intr);
3635 break;
3636
3637 case nir_intrinsic_read_invocation:
3638 emit_read_invocation(ctx, intr);
3639 break;
3640
3641 case nir_intrinsic_load_workgroup_size:
3642 assert(ctx->local_group_size_var);
3643 store_def(ctx, intr->def.index, ctx->local_group_size_var, nir_type_uint);
3644 break;
3645
3646 case nir_intrinsic_load_shared:
3647 emit_load_shared(ctx, intr);
3648 break;
3649
3650 case nir_intrinsic_store_shared:
3651 emit_store_shared(ctx, intr);
3652 break;
3653
3654 case nir_intrinsic_load_scratch:
3655 emit_load_scratch(ctx, intr);
3656 break;
3657
3658 case nir_intrinsic_store_scratch:
3659 emit_store_scratch(ctx, intr);
3660 break;
3661
3662 case nir_intrinsic_shader_clock:
3663 emit_shader_clock(ctx, intr);
3664 break;
3665
3666 case nir_intrinsic_vote_all:
3667 case nir_intrinsic_vote_any:
3668 case nir_intrinsic_vote_ieq:
3669 case nir_intrinsic_vote_feq:
3670 emit_vote(ctx, intr);
3671 break;
3672
3673 case nir_intrinsic_is_sparse_resident_zink:
3674 emit_is_sparse_texels_resident(ctx, intr);
3675 break;
3676
3677 case nir_intrinsic_is_helper_invocation:
3678 emit_is_helper_invocation(ctx, intr);
3679 break;
3680
3681 case nir_intrinsic_ddx:
3682 case nir_intrinsic_ddy:
3683 case nir_intrinsic_ddx_fine:
3684 case nir_intrinsic_ddy_fine:
3685 case nir_intrinsic_ddx_coarse:
3686 case nir_intrinsic_ddy_coarse:
3687 emit_derivative(ctx, intr);
3688 break;
3689
3690 default:
3691 fprintf(stderr, "emit_intrinsic: not implemented (%s)\n",
3692 nir_intrinsic_infos[intr->intrinsic].name);
3693 unreachable("unsupported intrinsic");
3694 }
3695 }
3696
3697 static void
emit_undef(struct ntv_context * ctx,nir_undef_instr * undef)3698 emit_undef(struct ntv_context *ctx, nir_undef_instr *undef)
3699 {
3700 SpvId type = undef->def.bit_size == 1 ? get_bvec_type(ctx, undef->def.num_components) :
3701 get_uvec_type(ctx, undef->def.bit_size,
3702 undef->def.num_components);
3703
3704 store_def(ctx, undef->def.index,
3705 spirv_builder_emit_undef(&ctx->builder, type),
3706 undef->def.bit_size == 1 ? nir_type_bool : nir_type_uint);
3707 }
3708
3709 static SpvId
get_src_float(struct ntv_context * ctx,nir_src * src)3710 get_src_float(struct ntv_context *ctx, nir_src *src)
3711 {
3712 nir_alu_type atype;
3713 SpvId def = get_src(ctx, src, &atype);
3714 if (atype == nir_type_float)
3715 return def;
3716 unsigned num_components = nir_src_num_components(*src);
3717 unsigned bit_size = nir_src_bit_size(*src);
3718 return bitcast_to_fvec(ctx, def, bit_size, num_components);
3719 }
3720
3721 static SpvId
get_src_int(struct ntv_context * ctx,nir_src * src)3722 get_src_int(struct ntv_context *ctx, nir_src *src)
3723 {
3724 nir_alu_type atype;
3725 SpvId def = get_src(ctx, src, &atype);
3726 if (atype == nir_type_int)
3727 return def;
3728 unsigned num_components = nir_src_num_components(*src);
3729 unsigned bit_size = nir_src_bit_size(*src);
3730 return bitcast_to_ivec(ctx, def, bit_size, num_components);
3731 }
3732
3733 static inline bool
tex_instr_is_lod_allowed(nir_tex_instr * tex)3734 tex_instr_is_lod_allowed(nir_tex_instr *tex)
3735 {
3736 /* This can only be used with an OpTypeImage that has a Dim operand of 1D, 2D, 3D, or Cube
3737 * - SPIR-V: 3.14. Image Operands
3738 */
3739
3740 return (tex->sampler_dim == GLSL_SAMPLER_DIM_1D ||
3741 tex->sampler_dim == GLSL_SAMPLER_DIM_2D ||
3742 tex->sampler_dim == GLSL_SAMPLER_DIM_3D ||
3743 tex->sampler_dim == GLSL_SAMPLER_DIM_CUBE ||
3744 /* External images are interpreted as 2D in type_to_dim,
3745 * so LOD is allowed */
3746 tex->sampler_dim == GLSL_SAMPLER_DIM_EXTERNAL ||
3747 /* RECT will always become 2D, so this is fine */
3748 tex->sampler_dim == GLSL_SAMPLER_DIM_RECT);
3749 }
3750
3751 static nir_variable *
get_tex_srcs(struct ntv_context * ctx,nir_tex_instr * tex,nir_variable ** bindless_var,unsigned * coord_components,struct spriv_tex_src * tex_src)3752 get_tex_srcs(struct ntv_context *ctx, nir_tex_instr *tex,
3753 nir_variable **bindless_var, unsigned *coord_components,
3754 struct spriv_tex_src *tex_src)
3755 {
3756 nir_variable *var = NULL;
3757 nir_alu_type atype;
3758 tex_src->sparse = tex->is_sparse;
3759 for (unsigned i = 0; i < tex->num_srcs; i++) {
3760 nir_const_value *cv;
3761 switch (tex->src[i].src_type) {
3762 case nir_tex_src_texture_deref:
3763 var = nir_deref_instr_get_variable(nir_instr_as_deref(tex->src[i].src.ssa->parent_instr));
3764 tex_src->src = get_src(ctx, &tex->src[i].src, &atype);
3765 break;
3766 case nir_tex_src_sampler_deref:
3767 tex_src->cl_sampler = get_src(ctx, &tex->src[i].src, &atype);
3768 break;
3769
3770 case nir_tex_src_coord:
3771 if (tex->op == nir_texop_txf ||
3772 tex->op == nir_texop_txf_ms)
3773 tex_src->coord = get_src_int(ctx, &tex->src[i].src);
3774 else
3775 tex_src->coord = get_src_float(ctx, &tex->src[i].src);
3776 *coord_components = nir_src_num_components(tex->src[i].src);
3777 break;
3778
3779 case nir_tex_src_projector:
3780 assert(nir_src_num_components(tex->src[i].src) == 1);
3781 tex_src->proj = get_src_float(ctx, &tex->src[i].src);
3782 assert(tex_src->proj != 0);
3783 break;
3784
3785 case nir_tex_src_offset:
3786 cv = nir_src_as_const_value(tex->src[i].src);
3787 if (cv) {
3788 unsigned bit_size = nir_src_bit_size(tex->src[i].src);
3789 unsigned num_components = nir_src_num_components(tex->src[i].src);
3790
3791 SpvId components[NIR_MAX_VEC_COMPONENTS];
3792 for (int j = 0; j < num_components; ++j) {
3793 int64_t tmp = nir_const_value_as_int(cv[j], bit_size);
3794 components[j] = emit_int_const(ctx, bit_size, tmp);
3795 }
3796
3797 if (num_components > 1) {
3798 SpvId type = get_ivec_type(ctx, bit_size, num_components);
3799 tex_src->const_offset = spirv_builder_const_composite(&ctx->builder,
3800 type,
3801 components,
3802 num_components);
3803 } else
3804 tex_src->const_offset = components[0];
3805 } else
3806 tex_src->offset = get_src_int(ctx, &tex->src[i].src);
3807 break;
3808
3809 case nir_tex_src_bias:
3810 assert(tex->op == nir_texop_txb);
3811 tex_src->bias = get_src_float(ctx, &tex->src[i].src);
3812 assert(tex_src->bias != 0);
3813 break;
3814
3815 case nir_tex_src_min_lod:
3816 assert(nir_src_num_components(tex->src[i].src) == 1);
3817 tex_src->min_lod = get_src_float(ctx, &tex->src[i].src);
3818 assert(tex_src->min_lod != 0);
3819 break;
3820
3821 case nir_tex_src_lod:
3822 assert(nir_src_num_components(tex->src[i].src) == 1);
3823 if (tex->op == nir_texop_txf ||
3824 tex->op == nir_texop_txf_ms ||
3825 tex->op == nir_texop_txs)
3826 tex_src->lod = get_src_int(ctx, &tex->src[i].src);
3827 else
3828 tex_src->lod = get_src_float(ctx, &tex->src[i].src);
3829 assert(tex_src->lod != 0);
3830 break;
3831
3832 case nir_tex_src_ms_index:
3833 assert(nir_src_num_components(tex->src[i].src) == 1);
3834 tex_src->sample = get_src_int(ctx, &tex->src[i].src);
3835 break;
3836
3837 case nir_tex_src_comparator:
3838 assert(nir_src_num_components(tex->src[i].src) == 1);
3839 tex_src->dref = get_src_float(ctx, &tex->src[i].src);
3840 assert(tex_src->dref != 0);
3841 break;
3842
3843 case nir_tex_src_ddx:
3844 tex_src->dx = get_src_float(ctx, &tex->src[i].src);
3845 assert(tex_src->dx != 0);
3846 break;
3847
3848 case nir_tex_src_ddy:
3849 tex_src->dy = get_src_float(ctx, &tex->src[i].src);
3850 assert(tex_src->dy != 0);
3851 break;
3852
3853 case nir_tex_src_texture_offset:
3854 tex_src->tex_offset = get_src_int(ctx, &tex->src[i].src);
3855 break;
3856
3857 case nir_tex_src_sampler_offset:
3858 case nir_tex_src_sampler_handle:
3859 /* don't care */
3860 break;
3861
3862 case nir_tex_src_texture_handle:
3863 tex_src->src = get_src(ctx, &tex->src[i].src, &atype);
3864 var = *bindless_var = nir_deref_instr_get_variable(nir_src_as_deref(tex->src[i].src));
3865 break;
3866
3867 default:
3868 fprintf(stderr, "texture source: %d\n", tex->src[i].src_type);
3869 unreachable("unknown texture source");
3870 }
3871 }
3872 return var;
3873 }
3874
3875 static SpvId
get_texture_load(struct ntv_context * ctx,SpvId sampler_id,nir_tex_instr * tex,SpvId cl_sampler,SpvId image_type,SpvId sampled_type)3876 get_texture_load(struct ntv_context *ctx, SpvId sampler_id, nir_tex_instr *tex,
3877 SpvId cl_sampler, SpvId image_type, SpvId sampled_type)
3878 {
3879 if (ctx->stage == MESA_SHADER_KERNEL) {
3880 SpvId image_load = spirv_builder_emit_load(&ctx->builder, image_type, sampler_id);
3881 if (nir_tex_instr_need_sampler(tex)) {
3882 SpvId sampler_load = spirv_builder_emit_load(&ctx->builder, spirv_builder_type_sampler(&ctx->builder),
3883 cl_sampler);
3884 return spirv_builder_emit_sampled_image(&ctx->builder, sampled_type, image_load, sampler_load);
3885 } else {
3886 return image_load;
3887 }
3888 } else {
3889 return spirv_builder_emit_load(&ctx->builder, sampled_type, sampler_id);
3890 }
3891 }
3892
3893 static SpvId
get_texop_dest_type(struct ntv_context * ctx,const nir_tex_instr * tex)3894 get_texop_dest_type(struct ntv_context *ctx, const nir_tex_instr *tex)
3895 {
3896 SpvId actual_dest_type;
3897 unsigned num_components = tex->def.num_components;
3898 switch (nir_alu_type_get_base_type(tex->dest_type)) {
3899 case nir_type_int:
3900 actual_dest_type = get_ivec_type(ctx, 32, num_components);
3901 break;
3902
3903 case nir_type_uint:
3904 actual_dest_type = get_uvec_type(ctx, 32, num_components);
3905 break;
3906
3907 case nir_type_float:
3908 actual_dest_type = get_fvec_type(ctx, 32, num_components);
3909 break;
3910
3911 default:
3912 unreachable("unexpected nir_alu_type");
3913 }
3914
3915 return actual_dest_type;
3916 }
3917
3918 static void
move_tex_proj_to_coord(struct ntv_context * ctx,unsigned coord_components,struct spriv_tex_src * tex_src)3919 move_tex_proj_to_coord(struct ntv_context *ctx, unsigned coord_components, struct spriv_tex_src *tex_src)
3920 {
3921 SpvId constituents[NIR_MAX_VEC_COMPONENTS + 1];
3922 if (coord_components == 1)
3923 constituents[0] = tex_src->coord;
3924 else {
3925 assert(coord_components > 1);
3926 SpvId float_type = spirv_builder_type_float(&ctx->builder, 32);
3927 for (uint32_t i = 0; i < coord_components; ++i)
3928 constituents[i] = spirv_builder_emit_composite_extract(&ctx->builder,
3929 float_type,
3930 tex_src->coord,
3931 &i, 1);
3932 }
3933
3934 constituents[coord_components++] = tex_src->proj;
3935
3936 SpvId vec_type = get_fvec_type(ctx, 32, coord_components);
3937 tex_src->coord = spirv_builder_emit_composite_construct(&ctx->builder,
3938 vec_type,
3939 constituents,
3940 coord_components);
3941 }
3942
3943 static SpvId
get_tex_image_to_load(struct ntv_context * ctx,SpvId image_type,bool is_buffer,SpvId load)3944 get_tex_image_to_load( struct ntv_context *ctx, SpvId image_type, bool is_buffer, SpvId load)
3945 {
3946 return is_buffer || ctx->stage == MESA_SHADER_KERNEL ?
3947 load :
3948 spirv_builder_emit_image(&ctx->builder, image_type, load);
3949 }
3950
3951 static SpvId
emit_tex_readop(struct ntv_context * ctx,nir_variable * bindless_var,SpvId load,struct spriv_tex_src * tex_src,SpvId dest_type,bool is_buffer,nir_variable * var,SpvId image_type,nir_tex_instr * tex)3952 emit_tex_readop(struct ntv_context *ctx, nir_variable *bindless_var, SpvId load,
3953 struct spriv_tex_src *tex_src, SpvId dest_type, bool is_buffer,
3954 nir_variable *var, SpvId image_type, nir_tex_instr *tex)
3955 {
3956 SpvId actual_dest_type = get_texop_dest_type(ctx, tex);
3957
3958 SpvId result;
3959 if (tex_src->offset)
3960 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityImageGatherExtended);
3961 if (tex_src->min_lod)
3962 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityMinLod);
3963 if (tex->op == nir_texop_txf ||
3964 tex->op == nir_texop_txf_ms ||
3965 tex->op == nir_texop_tg4) {
3966 SpvId image = get_tex_image_to_load(ctx, image_type, is_buffer, load);
3967
3968 if (tex->op == nir_texop_tg4) {
3969 if (tex_src->const_offset)
3970 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityImageGatherExtended);
3971 result = spirv_builder_emit_image_gather(&ctx->builder, actual_dest_type,
3972 load, tex_src, emit_uint_const(ctx, 32, tex->component));
3973 actual_dest_type = dest_type;
3974 } else {
3975 assert(tex->op == nir_texop_txf_ms || !tex_src->sample);
3976 bool is_ms;
3977 type_to_dim(glsl_get_sampler_dim(glsl_without_array(var->type)), &is_ms);
3978 assert(is_ms || !tex_src->sample);
3979 result = spirv_builder_emit_image_fetch(&ctx->builder, actual_dest_type,
3980 image, tex_src);
3981 }
3982 } else {
3983 if (tex->op == nir_texop_txl)
3984 tex_src->min_lod = 0;
3985 result = spirv_builder_emit_image_sample(&ctx->builder,
3986 actual_dest_type, load,
3987 tex_src);
3988 }
3989
3990 if (!bindless_var && (var->data.precision == GLSL_PRECISION_MEDIUM || var->data.precision == GLSL_PRECISION_LOW)) {
3991 spirv_builder_emit_decoration(&ctx->builder, result,
3992 SpvDecorationRelaxedPrecision);
3993 }
3994
3995 if (tex->is_sparse)
3996 result = extract_sparse_load(ctx, result, actual_dest_type, &tex->def);
3997
3998 if (tex->def.bit_size != 32) {
3999 /* convert FP32 to FP16 */
4000 result = emit_unop(ctx, SpvOpFConvert, dest_type, result);
4001 }
4002
4003 return result;
4004 }
4005
4006 static void
emit_tex(struct ntv_context * ctx,nir_tex_instr * tex)4007 emit_tex(struct ntv_context *ctx, nir_tex_instr *tex)
4008 {
4009 assert(tex->op == nir_texop_tex ||
4010 tex->op == nir_texop_txb ||
4011 tex->op == nir_texop_txl ||
4012 tex->op == nir_texop_txd ||
4013 tex->op == nir_texop_txf ||
4014 tex->op == nir_texop_txf_ms ||
4015 tex->op == nir_texop_txs ||
4016 tex->op == nir_texop_lod ||
4017 tex->op == nir_texop_tg4 ||
4018 tex->op == nir_texop_texture_samples ||
4019 tex->op == nir_texop_query_levels);
4020
4021 struct spriv_tex_src tex_src = {0};
4022 unsigned coord_components = 0;
4023 nir_variable *bindless_var = NULL;
4024 nir_variable *var = get_tex_srcs(ctx, tex, &bindless_var, &coord_components, &tex_src);
4025
4026 assert(var);
4027 SpvId image_type = find_image_type(ctx, var);
4028 assert(image_type);
4029
4030 bool is_buffer = glsl_get_sampler_dim(glsl_without_array(var->type)) ==
4031 GLSL_SAMPLER_DIM_BUF;
4032 SpvId sampled_type = is_buffer ? image_type :
4033 spirv_builder_type_sampled_image(&ctx->builder, image_type);
4034 assert(sampled_type);
4035
4036 SpvId sampler_id = tex_src.src;
4037 if (tex_src.tex_offset) {
4038 SpvId ptr = spirv_builder_type_pointer(&ctx->builder, SpvStorageClassUniformConstant, sampled_type);
4039 sampler_id = spirv_builder_emit_access_chain(&ctx->builder, ptr, sampler_id, &tex_src.tex_offset, 1);
4040 }
4041
4042 SpvId load = get_texture_load(ctx, sampler_id, tex, tex_src.cl_sampler, image_type, sampled_type);
4043
4044 if (tex->is_sparse)
4045 tex->def.num_components--;
4046 SpvId dest_type = get_def_type(ctx, &tex->def, tex->dest_type);
4047
4048 if (nir_tex_instr_is_query(tex))
4049 spirv_builder_emit_cap(&ctx->builder, SpvCapabilityImageQuery);
4050
4051 if (!tex_instr_is_lod_allowed(tex))
4052 tex_src.lod = 0;
4053 else if (ctx->stage != MESA_SHADER_FRAGMENT &&
4054 tex->op == nir_texop_tex && ctx->explicit_lod && !tex_src.lod)
4055 tex_src.lod = emit_float_const(ctx, 32, 0.0);
4056
4057 if (tex_src.proj && coord_components > 0)
4058 move_tex_proj_to_coord(ctx, coord_components, &tex_src);
4059
4060 SpvId result = 0;
4061
4062 switch (tex->op) {
4063 case nir_texop_txs: {
4064 SpvId image = get_tex_image_to_load(ctx, image_type, is_buffer, load);
4065 /* Its Dim operand must be one of 1D, 2D, 3D, or Cube
4066 * - OpImageQuerySizeLod specification
4067 *
4068 * Additionally, if its Dim is 1D, 2D, 3D, or Cube,
4069 * it must also have either an MS of 1 or a Sampled of 0 or 2.
4070 * - OpImageQuerySize specification
4071 *
4072 * all spirv samplers use these types
4073 */
4074 if (!tex_src.lod && tex_instr_is_lod_allowed(tex))
4075 tex_src.lod = emit_uint_const(ctx, 32, 0);
4076 result = spirv_builder_emit_image_query_size(&ctx->builder,
4077 dest_type, image,
4078 tex_src.lod);
4079 break;
4080 }
4081 case nir_texop_query_levels: {
4082 SpvId image = get_tex_image_to_load(ctx, image_type, is_buffer, load);
4083 result = spirv_builder_emit_image_query_levels(&ctx->builder,
4084 dest_type, image);
4085 break;
4086 }
4087 case nir_texop_texture_samples: {
4088 SpvId image = get_tex_image_to_load(ctx, image_type, is_buffer, load);
4089 result = spirv_builder_emit_unop(&ctx->builder, SpvOpImageQuerySamples,
4090 dest_type, image);
4091 break;
4092 }
4093 case nir_texop_lod: {
4094 result = spirv_builder_emit_image_query_lod(&ctx->builder,
4095 dest_type, load,
4096 tex_src.coord);
4097 break;
4098 }
4099 default:
4100 result = emit_tex_readop(ctx, bindless_var, load, &tex_src,
4101 dest_type, is_buffer, var, image_type, tex);
4102 break;
4103 }
4104
4105 store_def(ctx, tex->def.index, result, tex->dest_type);
4106
4107 if (tex->is_sparse)
4108 tex->def.num_components++;
4109 }
4110
4111 static void
start_block(struct ntv_context * ctx,SpvId label)4112 start_block(struct ntv_context *ctx, SpvId label)
4113 {
4114 /* terminate previous block if needed */
4115 if (ctx->block_started)
4116 spirv_builder_emit_branch(&ctx->builder, label);
4117
4118 /* start new block */
4119 spirv_builder_label(&ctx->builder, label);
4120 ctx->block_started = true;
4121 }
4122
4123 static void
branch(struct ntv_context * ctx,SpvId label)4124 branch(struct ntv_context *ctx, SpvId label)
4125 {
4126 assert(ctx->block_started);
4127 spirv_builder_emit_branch(&ctx->builder, label);
4128 ctx->block_started = false;
4129 }
4130
4131 static void
branch_conditional(struct ntv_context * ctx,SpvId condition,SpvId then_id,SpvId else_id)4132 branch_conditional(struct ntv_context *ctx, SpvId condition, SpvId then_id,
4133 SpvId else_id)
4134 {
4135 assert(ctx->block_started);
4136 spirv_builder_emit_branch_conditional(&ctx->builder, condition,
4137 then_id, else_id);
4138 ctx->block_started = false;
4139 }
4140
4141 static void
emit_jump(struct ntv_context * ctx,nir_jump_instr * jump)4142 emit_jump(struct ntv_context *ctx, nir_jump_instr *jump)
4143 {
4144 switch (jump->type) {
4145 case nir_jump_break:
4146 assert(ctx->loop_break);
4147 branch(ctx, ctx->loop_break);
4148 break;
4149
4150 case nir_jump_continue:
4151 assert(ctx->loop_cont);
4152 branch(ctx, ctx->loop_cont);
4153 break;
4154
4155 default:
4156 unreachable("Unsupported jump type\n");
4157 }
4158 }
4159
4160 // Detect if this is a gl_PerVertex access chain and build the complete access
4161 static bool
try_emit_gl_per_vertex_access(struct ntv_context * ctx,nir_deref_instr * deref)4162 try_emit_gl_per_vertex_access(struct ntv_context *ctx, nir_deref_instr *deref)
4163 {
4164 if (!stage_has_per_vertex(ctx))
4165 return false;
4166
4167 // Walk up the dereference chain to find components
4168 nir_deref_instr *cur = deref;
4169 nir_deref_instr* array_derefs[2] = {NULL, NULL};
4170 int num_array_derefs = 0;
4171 nir_deref_instr *struct_deref = NULL;
4172 nir_variable *root_var = NULL;
4173
4174 // Walk up the chain: current -> parent -> parent...
4175 while (cur) {
4176 switch (cur->deref_type) {
4177 case nir_deref_type_var:
4178 root_var = cur->var;
4179 cur = NULL; // End of chain
4180 break;
4181 case nir_deref_type_array:
4182 if (num_array_derefs >= 2)
4183 return false;
4184 array_derefs[num_array_derefs++] = cur;
4185 cur = nir_deref_instr_parent(cur);
4186 break;
4187 case nir_deref_type_struct:
4188 if (!struct_deref) struct_deref = cur;
4189 cur = nir_deref_instr_parent(cur);
4190 break;
4191 default:
4192 return false; // Unsupported dereference type
4193 }
4194 }
4195
4196 // Check if this is a gl_PerVertex builtin access
4197 if (!root_var || !is_per_vertex_builtin(ctx, root_var))
4198 return false;
4199
4200 // Validate stage compatibility
4201 bool is_output = (root_var->data.mode == nir_var_shader_out);
4202 if (is_output && !stage_has_per_vertex_output(ctx->stage))
4203 return false;
4204 if (!is_output && !stage_has_per_vertex_input(ctx->stage))
4205 return false;
4206
4207 // Build the SPIR-V access chain
4208 SpvId base = is_output ? ctx->per_vertex_out : ctx->per_vertex_in;
4209 SpvId indices[3] = {0, 0, 0}; // Max: [array_index, struct_member, ...]
4210 int num_indices = 0;
4211
4212 // Add array index if present (for gl_in[i] or gl_out[i])
4213 bool is_access_per_vertex_array = is_output ? (ctx->stage == MESA_SHADER_TESS_CTRL) : true;
4214 if (is_access_per_vertex_array && num_array_derefs > 0 && array_derefs[num_array_derefs - 1]) {
4215 nir_deref_instr* array_deref = array_derefs[num_array_derefs - 1];
4216 nir_alu_type itype;
4217 SpvId array_index = get_src(ctx, &array_deref->arr.index, &itype);
4218 if (itype == nir_type_float)
4219 array_index = emit_bitcast(ctx, get_uvec_type(ctx, 32, 1), array_index);
4220 indices[num_indices++] = array_index;
4221 }
4222
4223 // Add struct member index (always present for gl_PerVertex)
4224 enum per_vertex_member member = location_to_per_vertex_member(root_var->data.location);
4225 SpvId member_index = spirv_builder_const_uint(&ctx->builder, 32, member);
4226 indices[num_indices++] = member_index;
4227
4228 // Add array index if present (for gl_ClipDistance[i])
4229 bool is_access_clip_distance = is_access_per_vertex_array ? num_array_derefs > 1 : num_array_derefs > 0;
4230 if (is_access_clip_distance && array_derefs[0]) {
4231 nir_deref_instr* array_deref = array_derefs[0];
4232 nir_alu_type itype;
4233 SpvId array_index = get_src(ctx, &array_deref->arr.index, &itype);
4234 if (itype == nir_type_float)
4235 array_index = emit_bitcast(ctx, get_uvec_type(ctx, 32, 1), array_index);
4236 indices[num_indices++] = array_index;
4237 }
4238
4239 // Determine result type and storage class
4240 SpvId member_type = get_per_vertex_member_type(ctx, member);
4241 if (is_access_per_vertex_array) {
4242 member_type = num_array_derefs <= 0 ? ctx->per_vertex_block_type[is_output ? 1 : 0] :
4243 (num_array_derefs > 1 ? get_glsl_type(ctx, glsl_float_type()) : member_type);
4244 } else {
4245 member_type = num_array_derefs > 0 ? get_glsl_type(ctx, glsl_float_type()) : member_type;
4246 }
4247
4248 SpvStorageClass storage_class = get_storage_class(root_var);
4249 SpvId ptr_type = spirv_builder_type_pointer(&ctx->builder, storage_class, member_type);
4250
4251 // Generate the access chain
4252 SpvId result = spirv_builder_emit_access_chain(&ctx->builder, ptr_type, base, indices, num_indices);
4253 store_def(ctx, deref->def.index, result, get_nir_alu_type(deref->type));
4254
4255 return true;
4256 }
4257
4258 static void
emit_deref_var(struct ntv_context * ctx,nir_deref_instr * deref)4259 emit_deref_var(struct ntv_context *ctx, nir_deref_instr *deref)
4260 {
4261 assert(deref->deref_type == nir_deref_type_var);
4262
4263 struct hash_entry *he = _mesa_hash_table_search(ctx->vars, deref->var);
4264 assert(he);
4265 SpvId result = (SpvId)(intptr_t)he->data;
4266 store_def(ctx, deref->def.index, result, get_nir_alu_type(deref->type));
4267 }
4268
4269 static void
emit_deref_array(struct ntv_context * ctx,nir_deref_instr * deref)4270 emit_deref_array(struct ntv_context *ctx, nir_deref_instr *deref)
4271 {
4272 assert(deref->deref_type == nir_deref_type_array);
4273 nir_variable *var = nir_deref_instr_get_variable(deref);
4274
4275 if (!nir_src_is_always_uniform(deref->arr.index)) {
4276 if (deref->modes & nir_var_mem_ubo)
4277 spirv_builder_emit_cap(&ctx->builder,
4278 SpvCapabilityUniformBufferArrayDynamicIndexing);
4279
4280 if (deref->modes & nir_var_mem_ssbo)
4281 spirv_builder_emit_cap(&ctx->builder,
4282 SpvCapabilityStorageBufferArrayDynamicIndexing);
4283
4284 if (deref->modes & (nir_var_uniform | nir_var_image)) {
4285 const struct glsl_type *type = glsl_without_array(var->type);
4286 assert(glsl_type_is_sampler(type) || glsl_type_is_image(type));
4287
4288 if (glsl_type_is_sampler(type))
4289 spirv_builder_emit_cap(&ctx->builder,
4290 SpvCapabilitySampledImageArrayDynamicIndexing);
4291 else
4292 spirv_builder_emit_cap(&ctx->builder,
4293 SpvCapabilityStorageImageArrayDynamicIndexing);
4294 }
4295 }
4296
4297 SpvStorageClass storage_class = get_storage_class(var);
4298 SpvId type;
4299 nir_alu_type atype = nir_type_uint;
4300
4301 SpvId base = get_src(ctx, &deref->parent, &atype);
4302
4303 switch (var->data.mode) {
4304
4305 case nir_var_mem_ubo:
4306 case nir_var_mem_ssbo:
4307 base = get_src(ctx, &deref->parent, &atype);
4308 /* this is either the array<buffers> deref or the array<uint> deref */
4309 if (glsl_type_is_struct_or_ifc(deref->type)) {
4310 /* array<buffers> */
4311 type = get_bo_struct_type(ctx, var);
4312 break;
4313 }
4314 /* array<uint> */
4315 FALLTHROUGH;
4316 case nir_var_function_temp:
4317 case nir_var_shader_in:
4318 case nir_var_shader_out:
4319 base = get_src(ctx, &deref->parent, &atype);
4320 type = get_glsl_type(ctx, deref->type);
4321 break;
4322
4323 case nir_var_uniform:
4324 case nir_var_image: {
4325 base = get_src(ctx, &deref->parent, &atype);
4326 const struct glsl_type *gtype = glsl_without_array(deref->type);
4327 type = get_image_type(ctx, var,
4328 glsl_type_is_sampler(gtype),
4329 glsl_get_sampler_dim(gtype) == GLSL_SAMPLER_DIM_BUF);
4330 break;
4331 }
4332
4333 default:
4334 unreachable("Unsupported nir_variable_mode\n");
4335 }
4336
4337 nir_alu_type itype;
4338 SpvId index = get_src(ctx, &deref->arr.index, &itype);
4339 if (itype == nir_type_float)
4340 index = emit_bitcast(ctx, get_uvec_type(ctx, 32, 1), index);
4341
4342 SpvId ptr_type = spirv_builder_type_pointer(&ctx->builder,
4343 storage_class,
4344 type);
4345
4346 SpvId result = spirv_builder_emit_access_chain(&ctx->builder,
4347 ptr_type,
4348 base,
4349 &index, 1);
4350 /* uint is a bit of a lie here, it's really just an opaque type */
4351 store_def(ctx, deref->def.index, result, get_nir_alu_type(deref->type));
4352 }
4353
4354 static void
emit_deref_struct(struct ntv_context * ctx,nir_deref_instr * deref)4355 emit_deref_struct(struct ntv_context *ctx, nir_deref_instr *deref)
4356 {
4357 assert(deref->deref_type == nir_deref_type_struct);
4358 nir_variable *var = nir_deref_instr_get_variable(deref);
4359
4360 SpvStorageClass storage_class = get_storage_class(var);
4361
4362 SpvId index = emit_uint_const(ctx, 32, deref->strct.index);
4363 SpvId type = (var->data.mode & (nir_var_mem_ubo | nir_var_mem_ssbo)) ?
4364 get_bo_array_type(ctx, var) :
4365 get_glsl_type(ctx, deref->type);
4366
4367 SpvId ptr_type = spirv_builder_type_pointer(&ctx->builder,
4368 storage_class,
4369 type);
4370
4371 nir_alu_type atype;
4372 SpvId result = spirv_builder_emit_access_chain(&ctx->builder,
4373 ptr_type,
4374 get_src(ctx, &deref->parent, &atype),
4375 &index, 1);
4376 /* uint is a bit of a lie here, it's really just an opaque type */
4377 store_def(ctx, deref->def.index, result, get_nir_alu_type(deref->type));
4378 }
4379
4380 static void
emit_deref(struct ntv_context * ctx,nir_deref_instr * deref)4381 emit_deref(struct ntv_context *ctx, nir_deref_instr *deref)
4382 {
4383 if (try_emit_gl_per_vertex_access(ctx, deref))
4384 return;
4385
4386 switch (deref->deref_type) {
4387 case nir_deref_type_var:
4388 emit_deref_var(ctx, deref);
4389 break;
4390
4391 case nir_deref_type_array:
4392 emit_deref_array(ctx, deref);
4393 break;
4394
4395 case nir_deref_type_struct:
4396 emit_deref_struct(ctx, deref);
4397 break;
4398
4399 default:
4400 unreachable("unexpected deref_type");
4401 }
4402 }
4403
4404 static void
emit_block(struct ntv_context * ctx,struct nir_block * block)4405 emit_block(struct ntv_context *ctx, struct nir_block *block)
4406 {
4407 start_block(ctx, block_label(ctx, block));
4408 nir_foreach_instr(instr, block) {
4409 switch (instr->type) {
4410 case nir_instr_type_alu:
4411 emit_alu(ctx, nir_instr_as_alu(instr));
4412 break;
4413 case nir_instr_type_intrinsic:
4414 emit_intrinsic(ctx, nir_instr_as_intrinsic(instr));
4415 break;
4416 case nir_instr_type_load_const:
4417 emit_load_const(ctx, nir_instr_as_load_const(instr));
4418 break;
4419 case nir_instr_type_undef:
4420 emit_undef(ctx, nir_instr_as_undef(instr));
4421 break;
4422 case nir_instr_type_tex:
4423 emit_tex(ctx, nir_instr_as_tex(instr));
4424 break;
4425 case nir_instr_type_phi:
4426 unreachable("nir_instr_type_phi not supported");
4427 break;
4428 case nir_instr_type_jump:
4429 emit_jump(ctx, nir_instr_as_jump(instr));
4430 break;
4431 case nir_instr_type_call:
4432 unreachable("nir_instr_type_call not supported");
4433 break;
4434 case nir_instr_type_parallel_copy:
4435 unreachable("nir_instr_type_parallel_copy not supported");
4436 break;
4437 case nir_instr_type_deref:
4438 emit_deref(ctx, nir_instr_as_deref(instr));
4439 break;
4440 case nir_instr_type_debug_info:
4441 unreachable("nir_instr_type_debug_info not supported");
4442 break;
4443 }
4444 }
4445 }
4446
4447 static void
4448 emit_cf_list(struct ntv_context *ctx, struct exec_list *list);
4449
4450 static SpvId
get_src_bool(struct ntv_context * ctx,nir_src * src)4451 get_src_bool(struct ntv_context *ctx, nir_src *src)
4452 {
4453 assert(nir_src_bit_size(*src) == 1);
4454 nir_alu_type atype;
4455 return get_src(ctx, src, &atype);
4456 }
4457
4458 static void
emit_if(struct ntv_context * ctx,nir_if * if_stmt)4459 emit_if(struct ntv_context *ctx, nir_if *if_stmt)
4460 {
4461 SpvId condition = get_src_bool(ctx, &if_stmt->condition);
4462
4463 SpvId header_id = spirv_builder_new_id(&ctx->builder);
4464 SpvId then_id = block_label(ctx, nir_if_first_then_block(if_stmt));
4465 SpvId endif_id = spirv_builder_new_id(&ctx->builder);
4466 SpvId else_id = endif_id;
4467
4468 bool has_else = !exec_list_is_empty(&if_stmt->else_list);
4469 if (has_else) {
4470 assert(nir_if_first_else_block(if_stmt)->index < ctx->num_blocks);
4471 else_id = block_label(ctx, nir_if_first_else_block(if_stmt));
4472 }
4473
4474 /* create a header-block */
4475 start_block(ctx, header_id);
4476 spirv_builder_emit_selection_merge(&ctx->builder, endif_id,
4477 SpvSelectionControlMaskNone);
4478 branch_conditional(ctx, condition, then_id, else_id);
4479
4480 emit_cf_list(ctx, &if_stmt->then_list);
4481
4482 if (has_else) {
4483 if (ctx->block_started)
4484 branch(ctx, endif_id);
4485
4486 emit_cf_list(ctx, &if_stmt->else_list);
4487 }
4488
4489 start_block(ctx, endif_id);
4490 }
4491
4492 static void
emit_loop(struct ntv_context * ctx,nir_loop * loop)4493 emit_loop(struct ntv_context *ctx, nir_loop *loop)
4494 {
4495 assert(!nir_loop_has_continue_construct(loop));
4496 SpvId header_id = spirv_builder_new_id(&ctx->builder);
4497 SpvId begin_id = block_label(ctx, nir_loop_first_block(loop));
4498 SpvId break_id = spirv_builder_new_id(&ctx->builder);
4499 SpvId cont_id = spirv_builder_new_id(&ctx->builder);
4500
4501 /* create a header-block */
4502 start_block(ctx, header_id);
4503 spirv_builder_loop_merge(&ctx->builder, break_id, cont_id, SpvLoopControlMaskNone);
4504 branch(ctx, begin_id);
4505
4506 SpvId save_break = ctx->loop_break;
4507 SpvId save_cont = ctx->loop_cont;
4508 ctx->loop_break = break_id;
4509 ctx->loop_cont = cont_id;
4510
4511 emit_cf_list(ctx, &loop->body);
4512
4513 ctx->loop_break = save_break;
4514 ctx->loop_cont = save_cont;
4515
4516 /* loop->body may have already ended our block */
4517 if (ctx->block_started)
4518 branch(ctx, cont_id);
4519 start_block(ctx, cont_id);
4520 branch(ctx, header_id);
4521
4522 start_block(ctx, break_id);
4523 }
4524
4525 static void
emit_cf_list(struct ntv_context * ctx,struct exec_list * list)4526 emit_cf_list(struct ntv_context *ctx, struct exec_list *list)
4527 {
4528 foreach_list_typed(nir_cf_node, node, node, list) {
4529 switch (node->type) {
4530 case nir_cf_node_block:
4531 emit_block(ctx, nir_cf_node_as_block(node));
4532 break;
4533
4534 case nir_cf_node_if:
4535 emit_if(ctx, nir_cf_node_as_if(node));
4536 break;
4537
4538 case nir_cf_node_loop:
4539 emit_loop(ctx, nir_cf_node_as_loop(node));
4540 break;
4541
4542 case nir_cf_node_function:
4543 unreachable("nir_cf_node_function not supported");
4544 break;
4545 }
4546 }
4547 }
4548
4549 static SpvExecutionMode
get_input_prim_type_mode(enum mesa_prim type)4550 get_input_prim_type_mode(enum mesa_prim type)
4551 {
4552 switch (type) {
4553 case MESA_PRIM_POINTS:
4554 return SpvExecutionModeInputPoints;
4555 case MESA_PRIM_LINES:
4556 case MESA_PRIM_LINE_LOOP:
4557 case MESA_PRIM_LINE_STRIP:
4558 return SpvExecutionModeInputLines;
4559 case MESA_PRIM_TRIANGLE_STRIP:
4560 case MESA_PRIM_TRIANGLES:
4561 case MESA_PRIM_TRIANGLE_FAN:
4562 return SpvExecutionModeTriangles;
4563 case MESA_PRIM_QUADS:
4564 case MESA_PRIM_QUAD_STRIP:
4565 return SpvExecutionModeQuads;
4566 break;
4567 case MESA_PRIM_POLYGON:
4568 unreachable("handle polygons in gs");
4569 break;
4570 case MESA_PRIM_LINES_ADJACENCY:
4571 case MESA_PRIM_LINE_STRIP_ADJACENCY:
4572 return SpvExecutionModeInputLinesAdjacency;
4573 case MESA_PRIM_TRIANGLES_ADJACENCY:
4574 case MESA_PRIM_TRIANGLE_STRIP_ADJACENCY:
4575 return SpvExecutionModeInputTrianglesAdjacency;
4576 break;
4577 default:
4578 debug_printf("unknown geometry shader input mode %u\n", type);
4579 unreachable("error!");
4580 break;
4581 }
4582
4583 return 0;
4584 }
4585 static SpvExecutionMode
get_output_prim_type_mode(enum mesa_prim type)4586 get_output_prim_type_mode(enum mesa_prim type)
4587 {
4588 switch (type) {
4589 case MESA_PRIM_POINTS:
4590 return SpvExecutionModeOutputPoints;
4591 case MESA_PRIM_LINES:
4592 case MESA_PRIM_LINE_LOOP:
4593 unreachable("MESA_PRIM_LINES/LINE_LOOP passed as gs output");
4594 break;
4595 case MESA_PRIM_LINE_STRIP:
4596 return SpvExecutionModeOutputLineStrip;
4597 case MESA_PRIM_TRIANGLE_STRIP:
4598 return SpvExecutionModeOutputTriangleStrip;
4599 case MESA_PRIM_TRIANGLES:
4600 case MESA_PRIM_TRIANGLE_FAN: //FIXME: not sure if right for output
4601 return SpvExecutionModeTriangles;
4602 case MESA_PRIM_QUADS:
4603 case MESA_PRIM_QUAD_STRIP:
4604 return SpvExecutionModeQuads;
4605 case MESA_PRIM_POLYGON:
4606 unreachable("handle polygons in gs");
4607 break;
4608 case MESA_PRIM_LINES_ADJACENCY:
4609 case MESA_PRIM_LINE_STRIP_ADJACENCY:
4610 unreachable("handle line adjacency in gs");
4611 break;
4612 case MESA_PRIM_TRIANGLES_ADJACENCY:
4613 case MESA_PRIM_TRIANGLE_STRIP_ADJACENCY:
4614 unreachable("handle triangle adjacency in gs");
4615 break;
4616 default:
4617 debug_printf("unknown geometry shader output mode %u\n", type);
4618 unreachable("error!");
4619 break;
4620 }
4621
4622 return 0;
4623 }
4624
4625 static SpvExecutionMode
get_depth_layout_mode(enum gl_frag_depth_layout depth_layout)4626 get_depth_layout_mode(enum gl_frag_depth_layout depth_layout)
4627 {
4628 switch (depth_layout) {
4629 case FRAG_DEPTH_LAYOUT_NONE:
4630 case FRAG_DEPTH_LAYOUT_ANY:
4631 return SpvExecutionModeDepthReplacing;
4632 case FRAG_DEPTH_LAYOUT_GREATER:
4633 return SpvExecutionModeDepthGreater;
4634 case FRAG_DEPTH_LAYOUT_LESS:
4635 return SpvExecutionModeDepthLess;
4636 case FRAG_DEPTH_LAYOUT_UNCHANGED:
4637 return SpvExecutionModeDepthUnchanged;
4638 default:
4639 unreachable("unexpected depth layout");
4640 }
4641 }
4642
4643 static SpvExecutionMode
get_primitive_mode(enum tess_primitive_mode primitive_mode)4644 get_primitive_mode(enum tess_primitive_mode primitive_mode)
4645 {
4646 switch (primitive_mode) {
4647 case TESS_PRIMITIVE_TRIANGLES: return SpvExecutionModeTriangles;
4648 case TESS_PRIMITIVE_QUADS: return SpvExecutionModeQuads;
4649 case TESS_PRIMITIVE_ISOLINES: return SpvExecutionModeIsolines;
4650 default:
4651 unreachable("unknown tess prim type!");
4652 }
4653 }
4654
4655 static SpvExecutionMode
get_spacing(enum gl_tess_spacing spacing)4656 get_spacing(enum gl_tess_spacing spacing)
4657 {
4658 switch (spacing) {
4659 case TESS_SPACING_EQUAL:
4660 return SpvExecutionModeSpacingEqual;
4661 case TESS_SPACING_FRACTIONAL_ODD:
4662 return SpvExecutionModeSpacingFractionalOdd;
4663 case TESS_SPACING_FRACTIONAL_EVEN:
4664 return SpvExecutionModeSpacingFractionalEven;
4665 default:
4666 unreachable("unknown tess spacing!");
4667 }
4668 }
4669
4670 static void
setup_per_vertex_blocks(struct ntv_context * ctx)4671 setup_per_vertex_blocks(struct ntv_context *ctx)
4672 {
4673 SpvId block_type_in = create_per_vertex_block_type(ctx, true);
4674 SpvId block_type_out = create_per_vertex_block_type(ctx, false);
4675
4676 if (stage_has_per_vertex_input(ctx->stage)) {
4677 SpvId input_type = block_type_in;
4678
4679 if (ctx->stage == MESA_SHADER_GEOMETRY ||
4680 ctx->stage == MESA_SHADER_TESS_CTRL ||
4681 ctx->stage == MESA_SHADER_TESS_EVAL) {
4682 // gl_in[]: array of gl_PerVertex
4683 uint32_t input_vertices = MAX2(ctx->gl_in_num_vertices[ctx->stage - 1], 1);
4684
4685 SpvId array_size = spirv_builder_const_uint(&ctx->builder, 32, input_vertices);
4686 input_type = spirv_builder_type_array(&ctx->builder, block_type_in, array_size);
4687 }
4688
4689 SpvId ptr_type = spirv_builder_type_pointer(&ctx->builder, SpvStorageClassInput, input_type);
4690 SpvId var_id = spirv_builder_emit_var(&ctx->builder, ptr_type, SpvStorageClassInput);
4691 spirv_builder_emit_name(&ctx->builder, var_id, "gl_in");
4692 ctx->per_vertex_in = var_id;
4693 ctx->entry_ifaces[ctx->num_entry_ifaces++] = var_id;
4694 }
4695
4696 if (stage_has_per_vertex_output(ctx->stage)) {
4697 SpvId output_type = block_type_out;
4698
4699 // Tessellation control shaders need gl_out[] as an array
4700 if (ctx->stage == MESA_SHADER_TESS_CTRL) {
4701 uint32_t output_vertices = MAX2(ctx->gl_out_num_vertices, 1);
4702 SpvId array_size = spirv_builder_const_uint(&ctx->builder, 32, output_vertices);
4703 output_type = spirv_builder_type_array(&ctx->builder, block_type_out, array_size);
4704 }
4705
4706 SpvId ptr_type = spirv_builder_type_pointer(&ctx->builder, SpvStorageClassOutput, output_type);
4707 SpvId var_id = spirv_builder_emit_var(&ctx->builder, ptr_type, SpvStorageClassOutput);
4708 spirv_builder_emit_name(&ctx->builder, var_id, "gl_out");
4709
4710 ctx->per_vertex_out = var_id;
4711 ctx->entry_ifaces[ctx->num_entry_ifaces++] = var_id;
4712 }
4713 }
4714
4715 struct spirv_shader *
nir_to_spirv(struct nir_shader * s,const struct zink_shader_info * sinfo,const struct zink_screen * screen)4716 nir_to_spirv(struct nir_shader *s, const struct zink_shader_info *sinfo, const struct zink_screen *screen)
4717 {
4718 const uint32_t spirv_version = screen->spirv_version;
4719 struct spirv_shader *ret = NULL;
4720
4721 struct ntv_context ctx = {0};
4722 ctx.mem_ctx = ralloc_context(NULL);
4723 ctx.builder.name_syms = _mesa_set_create(NULL, _mesa_hash_string, _mesa_key_string_equal);
4724 ctx.builder.name_syms_index = 0;
4725 ctx.nir = s;
4726 ctx.builder.mem_ctx = ctx.mem_ctx;
4727 assert(spirv_version >= SPIRV_VERSION(1, 0));
4728 ctx.spirv_1_4_interfaces = spirv_version >= SPIRV_VERSION(1, 4);
4729
4730 ctx.bindless_set_idx = sinfo->bindless_set_idx;
4731 ctx.glsl_types = _mesa_pointer_hash_table_create(ctx.mem_ctx);
4732 ctx.bo_array_types = _mesa_pointer_hash_table_create(ctx.mem_ctx);
4733 ctx.bo_struct_types = _mesa_pointer_hash_table_create(ctx.mem_ctx);
4734 if (!ctx.glsl_types || !ctx.bo_array_types || !ctx.bo_struct_types ||
4735 !_mesa_hash_table_init(&ctx.image_types, ctx.mem_ctx, _mesa_hash_pointer, _mesa_key_pointer_equal))
4736 goto fail;
4737
4738 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityShader);
4739
4740 switch (s->info.stage) {
4741 case MESA_SHADER_FRAGMENT:
4742 if (s->info.fs.uses_sample_shading)
4743 spirv_builder_emit_cap(&ctx.builder, SpvCapabilitySampleRateShading);
4744 if (s->info.fs.uses_discard && spirv_version < SPIRV_VERSION(1, 6) &&
4745 screen->info.have_EXT_shader_demote_to_helper_invocation)
4746 spirv_builder_emit_extension(&ctx.builder,
4747 "SPV_EXT_demote_to_helper_invocation");
4748 break;
4749
4750 case MESA_SHADER_VERTEX:
4751 if (BITSET_TEST(s->info.system_values_read, SYSTEM_VALUE_INSTANCE_ID) ||
4752 BITSET_TEST(s->info.system_values_read, SYSTEM_VALUE_DRAW_ID) ||
4753 BITSET_TEST(s->info.system_values_read, SYSTEM_VALUE_BASE_INSTANCE) ||
4754 BITSET_TEST(s->info.system_values_read, SYSTEM_VALUE_BASE_VERTEX)) {
4755 if (spirv_version < SPIRV_VERSION(1, 3))
4756 spirv_builder_emit_extension(&ctx.builder, "SPV_KHR_shader_draw_parameters");
4757 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityDrawParameters);
4758 }
4759 break;
4760
4761 case MESA_SHADER_TESS_CTRL:
4762 case MESA_SHADER_TESS_EVAL:
4763 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityTessellation);
4764 /* TODO: check features for this */
4765 if (s->info.outputs_written & BITFIELD64_BIT(VARYING_SLOT_PSIZ))
4766 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityTessellationPointSize);
4767 break;
4768
4769 case MESA_SHADER_GEOMETRY:
4770 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityGeometry);
4771 if (s->info.outputs_written & BITFIELD64_BIT(VARYING_SLOT_PSIZ))
4772 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityGeometryPointSize);
4773 break;
4774
4775 default: ;
4776 }
4777
4778 if (s->info.stage < MESA_SHADER_GEOMETRY) {
4779 if (s->info.outputs_written & BITFIELD64_BIT(VARYING_SLOT_LAYER) ||
4780 s->info.inputs_read & BITFIELD64_BIT(VARYING_SLOT_LAYER)) {
4781 if (spirv_version >= SPIRV_VERSION(1, 5))
4782 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityShaderLayer);
4783 else {
4784 spirv_builder_emit_extension(&ctx.builder, "SPV_EXT_shader_viewport_index_layer");
4785 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityShaderViewportIndexLayerEXT);
4786 }
4787 }
4788 } else if (s->info.stage == MESA_SHADER_FRAGMENT) {
4789 /* incredibly, this is legal and intended.
4790 * https://github.com/KhronosGroup/SPIRV-Registry/issues/95
4791 */
4792 if (s->info.inputs_read & (BITFIELD64_BIT(VARYING_SLOT_LAYER) |
4793 BITFIELD64_BIT(VARYING_SLOT_PRIMITIVE_ID)))
4794 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityGeometry);
4795 }
4796
4797 if (s->info.num_ssbos && spirv_version < SPIRV_VERSION(1, 1))
4798 spirv_builder_emit_extension(&ctx.builder, "SPV_KHR_storage_buffer_storage_class");
4799
4800 if (s->info.stage < MESA_SHADER_FRAGMENT &&
4801 s->info.outputs_written & BITFIELD64_BIT(VARYING_SLOT_VIEWPORT)) {
4802 if (s->info.stage < MESA_SHADER_GEOMETRY)
4803 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityShaderViewportIndex);
4804 else
4805 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityMultiViewport);
4806 }
4807
4808 if (s->info.stage > MESA_SHADER_VERTEX &&
4809 s->info.inputs_read & BITFIELD64_BIT(VARYING_SLOT_VIEWPORT)) {
4810 if (s->info.stage < MESA_SHADER_GEOMETRY)
4811 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityShaderViewportIndex);
4812 else
4813 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityMultiViewport);
4814 }
4815
4816 ctx.stage = s->info.stage;
4817 ctx.sinfo = sinfo;
4818 ctx.GLSL_std_450 = spirv_builder_import(&ctx.builder, "GLSL.std.450");
4819 ctx.explicit_lod = true;
4820 spirv_builder_emit_source(&ctx.builder, SpvSourceLanguageUnknown, 0);
4821
4822 SpvAddressingModel model = SpvAddressingModelLogical;
4823 if (gl_shader_stage_is_compute(s->info.stage)) {
4824 if (s->info.cs.ptr_size == 32)
4825 model = SpvAddressingModelPhysical32;
4826 else if (s->info.cs.ptr_size == 64) {
4827 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityPhysicalStorageBufferAddresses);
4828 model = SpvAddressingModelPhysicalStorageBuffer64;
4829 } else
4830 model = SpvAddressingModelLogical;
4831 }
4832
4833 if (ctx.sinfo->have_vulkan_memory_model) {
4834 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityVulkanMemoryModel);
4835 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityVulkanMemoryModelDeviceScope);
4836 spirv_builder_emit_mem_model(&ctx.builder, model,
4837 SpvMemoryModelVulkan);
4838 } else {
4839 spirv_builder_emit_mem_model(&ctx.builder, model,
4840 SpvMemoryModelGLSL450);
4841 }
4842
4843 if (s->info.stage == MESA_SHADER_FRAGMENT &&
4844 s->info.outputs_written & BITFIELD64_BIT(FRAG_RESULT_STENCIL)) {
4845 spirv_builder_emit_extension(&ctx.builder, "SPV_EXT_shader_stencil_export");
4846 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityStencilExportEXT);
4847 }
4848
4849 SpvExecutionModel exec_model;
4850 switch (s->info.stage) {
4851 case MESA_SHADER_VERTEX:
4852 exec_model = SpvExecutionModelVertex;
4853 break;
4854 case MESA_SHADER_TESS_CTRL:
4855 exec_model = SpvExecutionModelTessellationControl;
4856 break;
4857 case MESA_SHADER_TESS_EVAL:
4858 exec_model = SpvExecutionModelTessellationEvaluation;
4859 break;
4860 case MESA_SHADER_GEOMETRY:
4861 exec_model = SpvExecutionModelGeometry;
4862 break;
4863 case MESA_SHADER_FRAGMENT:
4864 exec_model = SpvExecutionModelFragment;
4865 break;
4866 case MESA_SHADER_COMPUTE:
4867 case MESA_SHADER_KERNEL:
4868 exec_model = SpvExecutionModelGLCompute;
4869 break;
4870 default:
4871 unreachable("invalid stage");
4872 }
4873
4874 SpvId type_void = spirv_builder_type_void(&ctx.builder);
4875 SpvId type_void_func = spirv_builder_type_function(&ctx.builder, type_void,
4876 NULL, 0);
4877 SpvId entry_point = spirv_builder_new_id(&ctx.builder);
4878 spirv_builder_emit_name(&ctx.builder, entry_point, "main");
4879
4880 ctx.vars = _mesa_hash_table_create(ctx.mem_ctx, _mesa_hash_pointer,
4881 _mesa_key_pointer_equal);
4882
4883 // TODO: set this on a per driver basis
4884 ctx.use_gl_per_vertex = true;
4885
4886 nir_foreach_variable_with_modes(var, s, nir_var_mem_push_const)
4887 input_var_init(&ctx, var);
4888
4889 nir_foreach_shader_in_variable(var, s)
4890 emit_input(&ctx, var);
4891
4892 int max_output = 0;
4893 nir_foreach_shader_out_variable(var, s) {
4894 /* ignore SPIR-V built-ins, tagged with a sentinel value */
4895 if (var->data.driver_location != UINT_MAX) {
4896 assert(var->data.driver_location < INT_MAX);
4897 unsigned extent = glsl_count_attribute_slots(var->type, false);
4898 max_output = MAX2(max_output, (int)var->data.driver_location + extent);
4899 }
4900 emit_output(&ctx, var);
4901 }
4902
4903 if (ctx.use_gl_per_vertex)
4904 setup_per_vertex_blocks(&ctx);
4905
4906 uint32_t tcs_vertices_out_word = 0;
4907
4908 unsigned ubo_counter[2] = {0};
4909 nir_foreach_variable_with_modes(var, s, nir_var_mem_ubo)
4910 ubo_counter[var->data.driver_location != 0]++;
4911 nir_foreach_variable_with_modes(var, s, nir_var_mem_ubo)
4912 emit_bo(&ctx, var, ubo_counter[var->data.driver_location != 0] > 1);
4913
4914 unsigned ssbo_counter = 0;
4915 nir_foreach_variable_with_modes(var, s, nir_var_mem_ssbo)
4916 ssbo_counter++;
4917 nir_foreach_variable_with_modes(var, s, nir_var_mem_ssbo)
4918 emit_bo(&ctx, var, ssbo_counter > 1);
4919
4920 nir_foreach_variable_with_modes(var, s, nir_var_image)
4921 ctx.image_var[var->data.driver_location] = var;
4922 nir_foreach_variable_with_modes(var, s, nir_var_uniform) {
4923 if (glsl_type_is_sampler(glsl_without_array(var->type))) {
4924 if (var->data.descriptor_set == ctx.bindless_set_idx)
4925 ctx.bindless_sampler_var[var->data.driver_location] = var;
4926 else
4927 ctx.sampler_var[var->data.driver_location] = var;
4928 ctx.last_sampler = MAX2(ctx.last_sampler, var->data.driver_location);
4929 }
4930 }
4931 nir_foreach_variable_with_modes(var, s, nir_var_image | nir_var_uniform) {
4932 const struct glsl_type *type = glsl_without_array(var->type);
4933 if (glsl_type_is_bare_sampler(type))
4934 emit_sampler(&ctx, var);
4935 else if (glsl_type_is_sampler(type))
4936 emit_image(&ctx, var, get_bare_image_type(&ctx, var, true));
4937 else if (glsl_type_is_image(type))
4938 emit_image(&ctx, var, get_bare_image_type(&ctx, var, false));
4939 }
4940
4941 if (sinfo->float_controls.flush_denorms) {
4942 unsigned execution_mode = s->info.float_controls_execution_mode;
4943 bool flush_16_bit = nir_is_denorm_flush_to_zero(execution_mode, 16);
4944 bool flush_32_bit = nir_is_denorm_flush_to_zero(execution_mode, 32);
4945 bool flush_64_bit = nir_is_denorm_flush_to_zero(execution_mode, 64);
4946 bool preserve_16_bit = nir_is_denorm_preserve(execution_mode, 16);
4947 bool preserve_32_bit = nir_is_denorm_preserve(execution_mode, 32);
4948 bool preserve_64_bit = nir_is_denorm_preserve(execution_mode, 64);
4949 bool emit_cap_flush = false;
4950 bool emit_cap_preserve = false;
4951
4952 if (!sinfo->float_controls.denorms_all_independence) {
4953 bool flush = flush_16_bit && flush_64_bit;
4954 bool preserve = preserve_16_bit && preserve_64_bit;
4955
4956 if (!sinfo->float_controls.denorms_32_bit_independence) {
4957 flush = flush && flush_32_bit;
4958 preserve = preserve && preserve_32_bit;
4959
4960 flush_32_bit = flush;
4961 preserve_32_bit = preserve;
4962 }
4963
4964 flush_16_bit = flush;
4965 flush_64_bit = flush;
4966 preserve_16_bit = preserve;
4967 preserve_64_bit = preserve;
4968 }
4969
4970 if (flush_16_bit && sinfo->float_controls.flush_denorms & BITFIELD_BIT(0)) {
4971 emit_cap_flush = true;
4972 spirv_builder_emit_exec_mode_literal(&ctx.builder, entry_point,
4973 SpvExecutionModeDenormFlushToZero, 16);
4974 }
4975 if (flush_32_bit && sinfo->float_controls.flush_denorms & BITFIELD_BIT(1)) {
4976 emit_cap_flush = true;
4977 spirv_builder_emit_exec_mode_literal(&ctx.builder, entry_point,
4978 SpvExecutionModeDenormFlushToZero, 32);
4979 }
4980 if (flush_64_bit && sinfo->float_controls.flush_denorms & BITFIELD_BIT(2)) {
4981 emit_cap_flush = true;
4982 spirv_builder_emit_exec_mode_literal(&ctx.builder, entry_point,
4983 SpvExecutionModeDenormFlushToZero, 64);
4984 }
4985
4986 if (preserve_16_bit && sinfo->float_controls.preserve_denorms & BITFIELD_BIT(0)) {
4987 emit_cap_preserve = true;
4988 spirv_builder_emit_exec_mode_literal(&ctx.builder, entry_point,
4989 SpvExecutionModeDenormPreserve, 16);
4990 }
4991 if (preserve_32_bit && sinfo->float_controls.preserve_denorms & BITFIELD_BIT(1)) {
4992 emit_cap_preserve = true;
4993 spirv_builder_emit_exec_mode_literal(&ctx.builder, entry_point,
4994 SpvExecutionModeDenormPreserve, 32);
4995 }
4996 if (preserve_64_bit && sinfo->float_controls.preserve_denorms & BITFIELD_BIT(2)) {
4997 emit_cap_preserve = true;
4998 spirv_builder_emit_exec_mode_literal(&ctx.builder, entry_point,
4999 SpvExecutionModeDenormPreserve, 64);
5000 }
5001
5002 if (emit_cap_flush)
5003 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityDenormFlushToZero);
5004 if (emit_cap_preserve)
5005 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityDenormPreserve);
5006 }
5007
5008 switch (s->info.stage) {
5009 case MESA_SHADER_FRAGMENT:
5010 spirv_builder_emit_exec_mode(&ctx.builder, entry_point,
5011 SpvExecutionModeOriginUpperLeft);
5012 if (s->info.outputs_written & BITFIELD64_BIT(FRAG_RESULT_DEPTH))
5013 spirv_builder_emit_exec_mode(&ctx.builder, entry_point,
5014 get_depth_layout_mode(s->info.fs.depth_layout));
5015 if (s->info.outputs_written & BITFIELD64_BIT(FRAG_RESULT_STENCIL))
5016 spirv_builder_emit_exec_mode(&ctx.builder, entry_point,
5017 SpvExecutionModeStencilRefReplacingEXT);
5018 if (s->info.fs.early_fragment_tests)
5019 spirv_builder_emit_exec_mode(&ctx.builder, entry_point,
5020 SpvExecutionModeEarlyFragmentTests);
5021 if (s->info.fs.post_depth_coverage) {
5022 spirv_builder_emit_extension(&ctx.builder, "SPV_KHR_post_depth_coverage");
5023 spirv_builder_emit_cap(&ctx.builder, SpvCapabilitySampleMaskPostDepthCoverage);
5024 spirv_builder_emit_exec_mode(&ctx.builder, entry_point,
5025 SpvExecutionModePostDepthCoverage);
5026 }
5027
5028 if (s->info.fs.pixel_interlock_ordered || s->info.fs.pixel_interlock_unordered ||
5029 s->info.fs.sample_interlock_ordered || s->info.fs.sample_interlock_unordered)
5030 spirv_builder_emit_extension(&ctx.builder, "SPV_EXT_fragment_shader_interlock");
5031 if (s->info.fs.pixel_interlock_ordered || s->info.fs.pixel_interlock_unordered)
5032 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityFragmentShaderPixelInterlockEXT);
5033 if (s->info.fs.sample_interlock_ordered || s->info.fs.sample_interlock_unordered)
5034 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityFragmentShaderSampleInterlockEXT);
5035 if (s->info.fs.pixel_interlock_ordered)
5036 spirv_builder_emit_exec_mode(&ctx.builder, entry_point, SpvExecutionModePixelInterlockOrderedEXT);
5037 if (s->info.fs.pixel_interlock_unordered)
5038 spirv_builder_emit_exec_mode(&ctx.builder, entry_point, SpvExecutionModePixelInterlockUnorderedEXT);
5039 if (s->info.fs.sample_interlock_ordered)
5040 spirv_builder_emit_exec_mode(&ctx.builder, entry_point, SpvExecutionModeSampleInterlockOrderedEXT);
5041 if (s->info.fs.sample_interlock_unordered)
5042 spirv_builder_emit_exec_mode(&ctx.builder, entry_point, SpvExecutionModeSampleInterlockUnorderedEXT);
5043 break;
5044 case MESA_SHADER_TESS_CTRL:
5045 tcs_vertices_out_word = spirv_builder_emit_exec_mode_literal(&ctx.builder, entry_point,
5046 SpvExecutionModeOutputVertices,
5047 s->info.tess.tcs_vertices_out);
5048 break;
5049 case MESA_SHADER_TESS_EVAL:
5050 spirv_builder_emit_exec_mode(&ctx.builder, entry_point,
5051 get_primitive_mode(s->info.tess._primitive_mode));
5052 spirv_builder_emit_exec_mode(&ctx.builder, entry_point,
5053 s->info.tess.ccw ? SpvExecutionModeVertexOrderCcw
5054 : SpvExecutionModeVertexOrderCw);
5055 spirv_builder_emit_exec_mode(&ctx.builder, entry_point,
5056 get_spacing(s->info.tess.spacing));
5057 if (s->info.tess.point_mode)
5058 spirv_builder_emit_exec_mode(&ctx.builder, entry_point, SpvExecutionModePointMode);
5059 break;
5060 case MESA_SHADER_GEOMETRY:
5061 spirv_builder_emit_exec_mode(&ctx.builder, entry_point,
5062 get_input_prim_type_mode(s->info.gs.input_primitive));
5063 spirv_builder_emit_exec_mode(&ctx.builder, entry_point,
5064 get_output_prim_type_mode(s->info.gs.output_primitive));
5065 spirv_builder_emit_exec_mode_literal(&ctx.builder, entry_point,
5066 SpvExecutionModeInvocations,
5067 s->info.gs.invocations);
5068 spirv_builder_emit_exec_mode_literal(&ctx.builder, entry_point,
5069 SpvExecutionModeOutputVertices,
5070 MAX2(s->info.gs.vertices_out, 1));
5071 break;
5072 case MESA_SHADER_KERNEL:
5073 case MESA_SHADER_COMPUTE:
5074 if (s->info.workgroup_size[0] || s->info.workgroup_size[1] || s->info.workgroup_size[2])
5075 spirv_builder_emit_exec_mode_literal3(&ctx.builder, entry_point, SpvExecutionModeLocalSize,
5076 (uint32_t[3]){(uint32_t)s->info.workgroup_size[0], (uint32_t)s->info.workgroup_size[1],
5077 (uint32_t)s->info.workgroup_size[2]});
5078 else {
5079 SpvId sizes[3];
5080 uint32_t ids[] = {ZINK_WORKGROUP_SIZE_X, ZINK_WORKGROUP_SIZE_Y, ZINK_WORKGROUP_SIZE_Z};
5081 const char *names[] = {"x", "y", "z"};
5082 for (int i = 0; i < 3; i ++) {
5083 sizes[i] = spirv_builder_spec_const_uint(&ctx.builder, 32);
5084 spirv_builder_emit_specid(&ctx.builder, sizes[i], ids[i]);
5085 spirv_builder_emit_name(&ctx.builder, sizes[i], names[i]);
5086 }
5087 SpvId var_type = get_uvec_type(&ctx, 32, 3);
5088 // Even when using LocalSizeId this need to be initialized for nir_intrinsic_load_workgroup_size
5089 ctx.local_group_size_var = spirv_builder_spec_const_composite(&ctx.builder, var_type, sizes, 3);
5090 spirv_builder_emit_name(&ctx.builder, ctx.local_group_size_var, "gl_LocalGroupSizeARB");
5091
5092 /* WorkgroupSize is deprecated in SPIR-V 1.6 */
5093 if (spirv_version >= SPIRV_VERSION(1, 6)) {
5094 spirv_builder_emit_exec_mode_id3(&ctx.builder, entry_point,
5095 SpvExecutionModeLocalSizeId,
5096 sizes);
5097 } else {
5098 spirv_builder_emit_builtin(&ctx.builder, ctx.local_group_size_var, SpvBuiltInWorkgroupSize);
5099 }
5100 }
5101 if (s->info.cs.has_variable_shared_mem) {
5102 ctx.shared_mem_size = spirv_builder_spec_const_uint(&ctx.builder, 32);
5103 spirv_builder_emit_specid(&ctx.builder, ctx.shared_mem_size, ZINK_VARIABLE_SHARED_MEM);
5104 spirv_builder_emit_name(&ctx.builder, ctx.shared_mem_size, "variable_shared_mem");
5105 }
5106 if (s->info.derivative_group) {
5107 SpvCapability caps[] = { 0, SpvCapabilityComputeDerivativeGroupQuadsNV, SpvCapabilityComputeDerivativeGroupLinearNV };
5108 SpvExecutionMode modes[] = { 0, SpvExecutionModeDerivativeGroupQuadsNV, SpvExecutionModeDerivativeGroupLinearNV };
5109 spirv_builder_emit_extension(&ctx.builder, "SPV_NV_compute_shader_derivatives");
5110 spirv_builder_emit_cap(&ctx.builder, caps[s->info.derivative_group]);
5111 spirv_builder_emit_exec_mode(&ctx.builder, entry_point, modes[s->info.derivative_group]);
5112 ctx.explicit_lod = false;
5113 }
5114 break;
5115 default:
5116 break;
5117 }
5118 if (BITSET_TEST_RANGE(s->info.system_values_read, SYSTEM_VALUE_SUBGROUP_SIZE, SYSTEM_VALUE_SUBGROUP_LT_MASK)) {
5119 spirv_builder_emit_cap(&ctx.builder, SpvCapabilitySubgroupBallotKHR);
5120 spirv_builder_emit_extension(&ctx.builder, "SPV_KHR_shader_ballot");
5121 }
5122 if (s->info.has_transform_feedback_varyings && s->info.stage != MESA_SHADER_FRAGMENT) {
5123 spirv_builder_emit_cap(&ctx.builder, SpvCapabilityTransformFeedback);
5124 spirv_builder_emit_exec_mode(&ctx.builder, entry_point,
5125 SpvExecutionModeXfb);
5126 }
5127
5128 if (s->info.stage == MESA_SHADER_FRAGMENT && s->info.fs.uses_discard) {
5129 ctx.discard_func = spirv_builder_new_id(&ctx.builder);
5130 spirv_builder_emit_name(&ctx.builder, ctx.discard_func, "discard");
5131 spirv_builder_function(&ctx.builder, ctx.discard_func, type_void,
5132 SpvFunctionControlMaskNone,
5133 type_void_func);
5134 SpvId label = spirv_builder_new_id(&ctx.builder);
5135 spirv_builder_label(&ctx.builder, label);
5136
5137 /* kill is deprecated in SPIR-V 1.6, use terminate instead */
5138 if (spirv_version >= SPIRV_VERSION(1, 6))
5139 spirv_builder_emit_terminate(&ctx.builder);
5140 else
5141 spirv_builder_emit_kill(&ctx.builder);
5142
5143 spirv_builder_function_end(&ctx.builder);
5144 }
5145
5146 spirv_builder_function(&ctx.builder, entry_point, type_void,
5147 SpvFunctionControlMaskNone,
5148 type_void_func);
5149
5150 nir_function_impl *entry = nir_shader_get_entrypoint(s);
5151 nir_metadata_require(entry, nir_metadata_block_index);
5152
5153 ctx.defs = rzalloc_array_size(ctx.mem_ctx,
5154 sizeof(SpvId), entry->ssa_alloc);
5155 ctx.def_types = ralloc_array_size(ctx.mem_ctx,
5156 sizeof(nir_alu_type), entry->ssa_alloc);
5157 if (!ctx.defs || !ctx.def_types)
5158 goto fail;
5159 if (sinfo->have_sparse) {
5160 spirv_builder_emit_cap(&ctx.builder, SpvCapabilitySparseResidency);
5161 /* this could be huge, so only alloc if needed since it's extremely unlikely to
5162 * ever be used by anything except cts
5163 */
5164 ctx.resident_defs = rzalloc_array_size(ctx.mem_ctx,
5165 sizeof(SpvId), entry->ssa_alloc);
5166 if (!ctx.resident_defs)
5167 goto fail;
5168 }
5169 ctx.num_defs = entry->ssa_alloc;
5170
5171 SpvId *block_ids = ralloc_array_size(ctx.mem_ctx,
5172 sizeof(SpvId), entry->num_blocks);
5173 if (!block_ids)
5174 goto fail;
5175
5176 for (int i = 0; i < entry->num_blocks; ++i)
5177 block_ids[i] = spirv_builder_new_id(&ctx.builder);
5178
5179 ctx.block_ids = block_ids;
5180 ctx.num_blocks = entry->num_blocks;
5181
5182 /* emit a block only for the variable declarations */
5183 start_block(&ctx, spirv_builder_new_id(&ctx.builder));
5184 spirv_builder_begin_local_vars(&ctx.builder);
5185
5186 nir_foreach_reg_decl(reg, entry) {
5187 if (nir_intrinsic_bit_size(reg) == 1)
5188 init_reg(&ctx, reg, nir_type_bool);
5189 }
5190
5191 nir_foreach_variable_with_modes(var, s, nir_var_shader_temp)
5192 emit_shader_temp(&ctx, var);
5193
5194 nir_foreach_function_temp_variable(var, entry)
5195 emit_temp(&ctx, var);
5196
5197
5198 emit_cf_list(&ctx, &entry->body);
5199
5200 spirv_builder_return(&ctx.builder); // doesn't belong here, but whatevz
5201 spirv_builder_function_end(&ctx.builder);
5202
5203 spirv_builder_emit_entry_point(&ctx.builder, exec_model, entry_point,
5204 "main", ctx.entry_ifaces,
5205 ctx.num_entry_ifaces);
5206
5207 size_t num_words = spirv_builder_get_num_words(&ctx.builder);
5208
5209 ret = ralloc(NULL, struct spirv_shader);
5210 if (!ret)
5211 goto fail;
5212
5213 ret->words = ralloc_size(ret, sizeof(uint32_t) * num_words);
5214 if (!ret->words)
5215 goto fail;
5216
5217 ret->num_words = spirv_builder_get_words(&ctx.builder, ret->words, num_words, spirv_version, &tcs_vertices_out_word);
5218 ret->tcs_vertices_out_word = tcs_vertices_out_word;
5219 assert(ret->num_words == num_words);
5220
5221 ralloc_free(ctx.mem_ctx);
5222
5223 return ret;
5224
5225 fail:
5226 ralloc_free(ctx.mem_ctx);
5227
5228 if (ret)
5229 spirv_shader_delete(ret);
5230
5231 return NULL;
5232 }
5233
5234 void
spirv_shader_delete(struct spirv_shader * s)5235 spirv_shader_delete(struct spirv_shader *s)
5236 {
5237 ralloc_free(s);
5238 }
5239