1 /*
2 * Copyright © 2015 Intel Corporation
3 *
4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the "Software"),
6 * to deal in the Software without restriction, including without limitation
7 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
8 * and/or sell copies of the Software, and to permit persons to whom the
9 * Software is furnished to do so, subject to the following conditions:
10 *
11 * The above copyright notice and this permission notice (including the next
12 * paragraph) shall be included in all copies or substantial portions of the
13 * Software.
14 *
15 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
16 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
17 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
18 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
19 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
20 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
21 * IN THE SOFTWARE.
22 */
23
24 #ifndef _VTN_PRIVATE_H_
25 #define _VTN_PRIVATE_H_
26
27 #include <setjmp.h>
28
29 #include "nir/nir.h"
30 #include "nir/nir_builder.h"
31 #include "util/u_dynarray.h"
32 #include "nir_spirv.h"
33 #include "spirv.h"
34 #include "spirv_info.h"
35 #include "vtn_generator_ids.h"
36
37 extern uint32_t mesa_spirv_debug;
38
39 #define MESA_SPIRV_DEBUG(flag) unlikely(mesa_spirv_debug & (MESA_SPIRV_DEBUG_ ## flag))
40
41 #define MESA_SPIRV_DEBUG_STRUCTURED (1u << 0)
42 #define MESA_SPIRV_DEBUG_VALUES (1u << 1)
43 #define MESA_SPIRV_DEBUG_ASM (1u << 2)
44 #define MESA_SPIRV_DEBUG_COLOR (1u << 3)
45
46 struct vtn_builder;
47 struct vtn_decoration;
48
49 /* setjmp/longjmp is broken on MinGW: https://sourceforge.net/p/mingw-w64/bugs/406/ */
50 #if defined(__MINGW32__) && !defined(_UCRT)
51 #define vtn_setjmp __builtin_setjmp
52 #define vtn_longjmp __builtin_longjmp
53 #else
54 #define vtn_setjmp setjmp
55 #define vtn_longjmp longjmp
56 #endif
57
58 void vtn_log(struct vtn_builder *b, enum nir_spirv_debug_level level,
59 size_t spirv_offset, const char *message);
60
61 void vtn_logf(struct vtn_builder *b, enum nir_spirv_debug_level level,
62 size_t spirv_offset, const char *fmt, ...) PRINTFLIKE(4, 5);
63
64 #define vtn_info(...) vtn_logf(b, NIR_SPIRV_DEBUG_LEVEL_INFO, 0, __VA_ARGS__)
65
66 void _vtn_warn(struct vtn_builder *b, const char *file, unsigned line,
67 const char *fmt, ...) PRINTFLIKE(4, 5);
68 #define vtn_warn(...) _vtn_warn(b, __FILE__, __LINE__, __VA_ARGS__)
69
70 void _vtn_err(struct vtn_builder *b, const char *file, unsigned line,
71 const char *fmt, ...) PRINTFLIKE(4, 5);
72 #define vtn_err(...) _vtn_err(b, __FILE__, __LINE__, __VA_ARGS__)
73
74 /** Fail SPIR-V parsing
75 *
76 * This function logs an error and then bails out of the shader compile using
77 * longjmp. This being safe relies on two things:
78 *
79 * 1) We must guarantee that setjmp is called after allocating the builder
80 * and setting up b->debug (so that logging works) but before before any
81 * errors have a chance to occur.
82 *
83 * 2) While doing the SPIR-V -> NIR conversion, we need to be careful to
84 * ensure that all heap allocations happen through ralloc and are parented
85 * to the builder. This way they will get properly cleaned up on error.
86 *
87 * 3) We must ensure that _vtn_fail is never called while a mutex lock or a
88 * reference to any other resource is held with the exception of ralloc
89 * objects which are parented to the builder.
90 *
91 * So long as these two things continue to hold, we can easily longjmp back to
92 * spirv_to_nir(), clean up the builder, and return NULL.
93 */
94 NORETURN void
95 _vtn_fail(struct vtn_builder *b, const char *file, unsigned line,
96 const char *fmt, ...) PRINTFLIKE(4, 5);
97
98 #define vtn_fail(...) _vtn_fail(b, __FILE__, __LINE__, __VA_ARGS__)
99
100 /** Fail if the given expression evaluates to true */
101 #define vtn_fail_if(expr, ...) \
102 do { \
103 if (unlikely(expr)) \
104 vtn_fail(__VA_ARGS__); \
105 } while (0)
106
107 #define _vtn_fail_with(t, msg, v) \
108 vtn_fail("%s: %s (%u)\n", msg, spirv_ ## t ## _to_string(v), v)
109
110 #define vtn_fail_with_decoration(msg, v) _vtn_fail_with(decoration, msg, v)
111 #define vtn_fail_with_opcode(msg, v) _vtn_fail_with(op, msg, v)
112
113 /** Assert that a condition is true and, if it isn't, vtn_fail
114 *
115 * This macro is transitional only and should not be used in new code. Use
116 * vtn_fail_if and provide a real message instead.
117 */
118 #define vtn_assert(expr) \
119 do { \
120 if (!likely(expr)) \
121 vtn_fail("%s", #expr); \
122 } while (0)
123
124 /* These are used to allocate data that can be dropped at the end of
125 * the parsing. Any NIR data structure should keep using the ralloc,
126 * since they will outlive the parsing.
127 */
128 #define vtn_alloc(B, TYPE) linear_alloc(B->lin_ctx, TYPE)
129 #define vtn_zalloc(B, TYPE) linear_zalloc(B->lin_ctx, TYPE)
130 #define vtn_alloc_array(B, TYPE, ELEMS) linear_alloc_array(B->lin_ctx, TYPE, ELEMS)
131 #define vtn_zalloc_array(B, TYPE, ELEMS) linear_zalloc_array(B->lin_ctx, TYPE, ELEMS)
132 #define vtn_alloc_size(B, SIZE) linear_alloc_child(B->lin_ctx, SIZE)
133 #define vtn_zalloc_size(B, SIZE) linear_zalloc_child(B->lin_ctx, SIZE)
134
135 enum vtn_value_type {
136 vtn_value_type_invalid = 0,
137 vtn_value_type_undef,
138 vtn_value_type_string,
139 vtn_value_type_decoration_group,
140 vtn_value_type_type,
141 vtn_value_type_constant,
142 vtn_value_type_pointer,
143 vtn_value_type_function,
144 vtn_value_type_block,
145 vtn_value_type_ssa,
146 vtn_value_type_extension,
147 vtn_value_type_image_pointer,
148 };
149
150 const char *vtn_value_type_to_string(enum vtn_value_type t);
151
152 struct vtn_case {
153 struct list_head link;
154
155 struct vtn_block *block;
156
157 /* The uint32_t values that map to this case */
158 struct util_dynarray values;
159
160 /* True if this is the default case */
161 bool is_default;
162
163 /* Initialized to false; used when sorting the list of cases */
164 bool visited;
165 };
166
167 struct vtn_block {
168 struct list_head link;
169
170 /** A pointer to the label instruction */
171 const uint32_t *label;
172
173 /** A pointer to the merge instruction (or NULL if non exists) */
174 const uint32_t *merge;
175
176 /** A pointer to the branch instruction that ends this block */
177 const uint32_t *branch;
178
179 /** Points to the switch case started by this block (if any) */
180 struct vtn_case *switch_case;
181
182 /** Every block ends in a nop intrinsic so that we can find it again */
183 nir_intrinsic_instr *end_nop;
184
185 /** attached nir_block */
186 struct nir_block *block;
187
188 /* Inner-most construct that this block is part of. */
189 struct vtn_construct *parent;
190
191 /* Blocks that succeed this block. Used by structured control flow. */
192 struct vtn_successor *successors;
193 unsigned successors_count;
194
195 /* Position of this block in the structured post-order traversal. */
196 unsigned pos;
197
198 bool visited;
199 };
200
201 struct vtn_function {
202 struct list_head link;
203
204 struct vtn_type *type;
205
206 bool referenced;
207 bool emitted;
208
209 nir_function *nir_func;
210 struct vtn_block *start_block;
211
212 struct list_head body;
213
214 const uint32_t *end;
215
216 SpvLinkageType linkage;
217 SpvFunctionControlMask control;
218
219 unsigned block_count;
220
221 /* Ordering of blocks to be processed by structured control flow. See
222 * vtn_structured_cfg.c for details.
223 */
224 unsigned ordered_blocks_count;
225 struct vtn_block **ordered_blocks;
226
227 /* Structured control flow constructs. See struct vtn_construct. */
228 struct list_head constructs;
229 };
230
231 #define vtn_foreach_function(func, func_list) \
232 list_for_each_entry(struct vtn_function, func, func_list, link)
233
234 #define vtn_foreach_case(cse, case_list) \
235 list_for_each_entry(struct vtn_case, cse, case_list, link)
236
237 #define vtn_foreach_case_safe(cse, case_list) \
238 list_for_each_entry_safe(struct vtn_case, cse, case_list, link)
239
240 typedef bool (*vtn_instruction_handler)(struct vtn_builder *, SpvOp,
241 const uint32_t *, unsigned);
242
243 void vtn_build_cfg(struct vtn_builder *b, const uint32_t *words,
244 const uint32_t *end);
245 void vtn_function_emit(struct vtn_builder *b, struct vtn_function *func,
246 vtn_instruction_handler instruction_handler);
247 void vtn_handle_function_call(struct vtn_builder *b, SpvOp opcode,
248 const uint32_t *w, unsigned count);
249
250 bool vtn_cfg_handle_prepass_instruction(struct vtn_builder *b, SpvOp opcode,
251 const uint32_t *w, unsigned count);
252 void vtn_emit_cf_func_structured(struct vtn_builder *b, struct vtn_function *func,
253 vtn_instruction_handler handler);
254 bool vtn_handle_phis_first_pass(struct vtn_builder *b, SpvOp opcode,
255 const uint32_t *w, unsigned count);
256 void vtn_emit_ret_store(struct vtn_builder *b, const struct vtn_block *block);
257 void vtn_build_structured_cfg(struct vtn_builder *b, const uint32_t *words,
258 const uint32_t *end);
259
260 const uint32_t *
261 vtn_foreach_instruction(struct vtn_builder *b, const uint32_t *start,
262 const uint32_t *end, vtn_instruction_handler handler);
263
264 struct vtn_ssa_value {
265 bool is_variable;
266
267 union {
268 nir_def *def;
269 nir_variable *var;
270 struct vtn_ssa_value **elems;
271 };
272
273 /* For matrices, if this is non-NULL, then this value is actually the
274 * transpose of some other value. The value that `transposed` points to
275 * always dominates this value.
276 */
277 struct vtn_ssa_value *transposed;
278
279 const struct glsl_type *type;
280 };
281
282 enum vtn_base_type {
283 vtn_base_type_void,
284 vtn_base_type_scalar,
285 vtn_base_type_vector,
286 vtn_base_type_matrix,
287 vtn_base_type_array,
288 vtn_base_type_struct,
289 vtn_base_type_pointer,
290 vtn_base_type_image,
291 vtn_base_type_sampler,
292 vtn_base_type_sampled_image,
293 vtn_base_type_accel_struct,
294 vtn_base_type_ray_query,
295 vtn_base_type_function,
296 vtn_base_type_event,
297 vtn_base_type_cooperative_matrix,
298 };
299
300 struct vtn_type {
301 enum vtn_base_type base_type;
302
303 const struct glsl_type *type;
304
305 /* The SPIR-V id of the given type. */
306 uint32_t id;
307
308 /* Specifies the length of complex types.
309 *
310 * For Workgroup pointers, this is the size of the referenced type.
311 */
312 unsigned length;
313
314 /* for arrays, matrices and pointers, the array stride */
315 unsigned stride;
316
317 /* Access qualifiers */
318 enum gl_access_qualifier access;
319
320 union {
321 /* Members for scalar, vector, and array-like types */
322 struct {
323 /* for arrays, the vtn_type for the elements of the array */
324 struct vtn_type *array_element;
325
326 /* for matrices, whether the matrix is stored row-major */
327 bool row_major:1;
328
329 /* Whether this type, or a parent type, has been decorated as a
330 * builtin
331 */
332 bool is_builtin:1;
333
334 /* Which built-in to use */
335 SpvBuiltIn builtin;
336 };
337
338 /* Members for struct types */
339 struct {
340 /* for structures, the vtn_type for each member */
341 struct vtn_type **members;
342
343 /* for structs, the offset of each member */
344 unsigned *offsets;
345
346 /* for structs, whether it was decorated as a "non-SSBO-like" block */
347 bool block:1;
348
349 /* for structs, whether it was decorated as an "SSBO-like" block */
350 bool buffer_block:1;
351
352 /* for structs with block == true, whether this is a builtin block
353 * (i.e. a block that contains only builtins).
354 */
355 bool builtin_block:1;
356
357 /* for structs and unions it specifies the minimum alignment of the
358 * members. 0 means packed.
359 *
360 * Set by CPacked and Alignment Decorations in kernels.
361 */
362 bool packed:1;
363 };
364
365 /* Members for pointer types */
366 struct {
367 /* For pointers, the vtn_type of the object pointed to. */
368 struct vtn_type *pointed;
369
370 /* Storage class for pointers */
371 SpvStorageClass storage_class;
372
373 /* Required alignment for pointers */
374 uint32_t align;
375 };
376
377 /* Members for image types */
378 struct {
379 /* GLSL image type for this type. This is not to be confused with
380 * vtn_type::type which is actually going to be the GLSL type for a
381 * pointer to an image, likely a uint32_t.
382 */
383 const struct glsl_type *glsl_image;
384
385 /* Image format for image_load_store type images */
386 unsigned image_format;
387
388 /* Access qualifier for storage images */
389 SpvAccessQualifier access_qualifier;
390 };
391
392 /* Members for sampled image types */
393 struct {
394 /* For sampled images, the image type */
395 struct vtn_type *image;
396 };
397
398 /* Members for function types */
399 struct {
400 /* For functions, the vtn_type for each parameter */
401 struct vtn_type **params;
402
403 /* Return type for functions */
404 struct vtn_type *return_type;
405 };
406
407 /* Members for cooperative matrix types. */
408 struct {
409 struct glsl_cmat_description desc;
410 struct vtn_type *component_type;
411 };
412 };
413 };
414
415 bool vtn_type_contains_block(struct vtn_builder *b, struct vtn_type *type);
416
417 bool vtn_types_compatible(struct vtn_builder *b,
418 struct vtn_type *t1, struct vtn_type *t2);
419
420 struct vtn_type *vtn_type_without_array(struct vtn_type *type);
421
422 struct vtn_variable;
423
424 enum vtn_access_mode {
425 vtn_access_mode_id,
426 vtn_access_mode_literal,
427 };
428
429 struct vtn_access_link {
430 enum vtn_access_mode mode;
431 int64_t id;
432 };
433
434 struct vtn_access_chain {
435 uint32_t length;
436
437 /** Whether or not to treat the base pointer as an array. This is only
438 * true if this access chain came from an OpPtrAccessChain.
439 */
440 bool ptr_as_array;
441
442 /* Access qualifiers */
443 enum gl_access_qualifier access;
444
445 bool in_bounds;
446
447 /** Struct elements and array offsets.
448 *
449 * This is an array of 1 so that it can conveniently be created on the
450 * stack but the real length is given by the length field.
451 */
452 struct vtn_access_link link[1];
453 };
454
455 enum vtn_variable_mode {
456 vtn_variable_mode_function,
457 vtn_variable_mode_private,
458 vtn_variable_mode_uniform,
459 vtn_variable_mode_atomic_counter,
460 vtn_variable_mode_ubo,
461 vtn_variable_mode_ssbo,
462 vtn_variable_mode_phys_ssbo,
463 vtn_variable_mode_push_constant,
464 vtn_variable_mode_workgroup,
465 vtn_variable_mode_cross_workgroup,
466 vtn_variable_mode_task_payload,
467 vtn_variable_mode_generic,
468 vtn_variable_mode_constant,
469 vtn_variable_mode_input,
470 vtn_variable_mode_output,
471 vtn_variable_mode_image,
472 vtn_variable_mode_accel_struct,
473 vtn_variable_mode_call_data,
474 vtn_variable_mode_call_data_in,
475 vtn_variable_mode_ray_payload,
476 vtn_variable_mode_ray_payload_in,
477 vtn_variable_mode_hit_attrib,
478 vtn_variable_mode_shader_record,
479 vtn_variable_mode_node_payload,
480 };
481
482 struct vtn_pointer {
483 /** The variable mode for the referenced data */
484 enum vtn_variable_mode mode;
485
486 /** The pointer type of this pointer */
487 struct vtn_type *type;
488
489 /** The referenced variable, if known
490 *
491 * This field may be NULL if the pointer uses a (block_index, offset) pair
492 * instead of an access chain or if the access chain starts at a deref.
493 */
494 struct vtn_variable *var;
495
496 /** The NIR deref corresponding to this pointer */
497 nir_deref_instr *deref;
498
499 /** A (block_index, offset) pair representing a UBO or SSBO position. */
500 struct nir_def *block_index;
501 struct nir_def *offset;
502
503 /* Access qualifiers */
504 enum gl_access_qualifier access;
505 };
506
507 struct vtn_variable {
508 enum vtn_variable_mode mode;
509
510 struct vtn_type *type;
511
512 unsigned descriptor_set;
513 unsigned binding;
514 bool explicit_binding;
515 unsigned offset;
516 unsigned input_attachment_index;
517
518 nir_variable *var;
519
520 /* If the variable is a struct with a location set on it then this will be
521 * stored here. This will be used to calculate locations for members that
522 * don’t have their own explicit location.
523 */
524 int base_location;
525
526 /**
527 * In some early released versions of GLSLang, it implemented all function
528 * calls by making copies of all parameters into temporary variables and
529 * passing those variables into the function. It even did so for samplers
530 * and images which violates the SPIR-V spec. Unfortunately, two games
531 * (Talos Principle and Doom) shipped with this old version of GLSLang and
532 * also happen to pass samplers into functions. Talos Principle received
533 * an update fairly shortly after release with an updated GLSLang. Doom,
534 * on the other hand, has never received an update so we need to work
535 * around this GLSLang issue in SPIR-V -> NIR. Hopefully, we can drop this
536 * hack at some point in the future.
537 */
538 struct vtn_pointer *copy_prop_sampler;
539
540 /* Access qualifiers. */
541 enum gl_access_qualifier access;
542 };
543
544 const struct glsl_type *
545 vtn_type_get_nir_type(struct vtn_builder *b, struct vtn_type *type,
546 enum vtn_variable_mode mode);
547
548 mesa_scope
549 vtn_translate_scope(struct vtn_builder *b, SpvScope scope);
550
551 struct vtn_image_pointer {
552 nir_deref_instr *image;
553 nir_def *coord;
554 nir_def *sample;
555 nir_def *lod;
556 };
557
558 struct vtn_value {
559 enum vtn_value_type value_type;
560
561 /* Workaround for https://gitlab.freedesktop.org/mesa/mesa/-/issues/3406
562 * Only set for OpImage / OpSampledImage. Note that this is in addition
563 * the existence of a NonUniform decoration on this value.*/
564 uint32_t propagated_non_uniform : 1;
565
566 /* Valid for vtn_value_type_constant to indicate the value is OpConstantNull. */
567 bool is_null_constant:1;
568
569 /* Valid when all the members of the value are undef. */
570 bool is_undef_constant:1;
571
572 const char *name;
573 struct vtn_decoration *decoration;
574 struct vtn_type *type;
575 union {
576 const char *str;
577 nir_constant *constant;
578 struct vtn_pointer *pointer;
579 struct vtn_image_pointer *image;
580 struct vtn_function *func;
581 struct vtn_block *block;
582 struct vtn_ssa_value *ssa;
583 vtn_instruction_handler ext_handler;
584 };
585 };
586
587 #define VTN_DEC_DECORATION -1
588 #define VTN_DEC_EXECUTION_MODE -2
589 #define VTN_DEC_STRUCT_MEMBER_NAME0 -3
590 #define VTN_DEC_STRUCT_MEMBER0 0
591
592 struct vtn_decoration {
593 struct vtn_decoration *next;
594
595 /* Different kinds of decorations are stored in a value,
596 the scope defines what decoration it refers to:
597
598 - VTN_DEC_DECORATION:
599 decoration associated with the value
600 - VTN_DEC_EXECUTION_MODE:
601 an execution mode associated with an entrypoint value
602 - VTN_DEC_STRUCT_MEMBER0 + m:
603 decoration associated with member m of a struct value
604 - VTN_DEC_STRUCT_MEMBER_NAME0 - m:
605 name of m'th member of a struct value
606 */
607 int scope;
608
609 uint32_t num_operands;
610 const uint32_t *operands;
611 struct vtn_value *group;
612
613 union {
614 SpvDecoration decoration;
615 SpvExecutionMode exec_mode;
616 const char *member_name;
617 };
618 };
619
620 struct vtn_builder {
621 nir_builder nb;
622
623 linear_ctx *lin_ctx;
624
625 /* Used by vtn_fail to jump back to the beginning of SPIR-V compilation */
626 jmp_buf fail_jump;
627
628 const uint32_t *spirv;
629 size_t spirv_word_count;
630 uint32_t version;
631
632 nir_shader *shader;
633 struct spirv_to_nir_options *options;
634 struct vtn_block *block;
635
636 /* Current offset, file, line, and column. Useful for debugging. Set
637 * automatically by vtn_foreach_instruction.
638 */
639 size_t spirv_offset;
640 const char *file;
641 int line, col;
642
643 /*
644 * Map from phi instructions (pointer to the start of the instruction)
645 * to the variable corresponding to it.
646 */
647 struct hash_table *phi_table;
648
649 /* In Vulkan, when lowering some modes variable access, the derefs of the
650 * variables are replaced with a resource index intrinsics, leaving the
651 * variable hanging. This set keeps track of them so they can be filtered
652 * (and not removed) in nir_remove_dead_variables.
653 */
654 struct set *vars_used_indirectly;
655
656 unsigned num_specializations;
657 struct nir_spirv_specialization *specializations;
658
659 unsigned value_id_bound;
660 struct vtn_value *values;
661
662 /* Information on the origin of the SPIR-V */
663 enum vtn_generator generator_id;
664 SpvSourceLanguage source_lang;
665
666 struct spirv_capabilities supported_capabilities;
667 struct spirv_capabilities enabled_capabilities;
668
669 /* True if we need to fix up CS OpControlBarrier */
670 bool wa_glslang_cs_barrier;
671
672 /* True if we need to ignore undef initializers */
673 bool wa_llvm_spirv_ignore_workgroup_initializer;
674
675 /* True if we need to ignore OpReturn after OpEmitMeshTasksEXT. */
676 bool wa_ignore_return_after_emit_mesh_tasks;
677
678 /* Workaround discard bugs in HLSL -> SPIR-V compilers */
679 bool convert_discard_to_demote;
680
681 gl_shader_stage entry_point_stage;
682 const char *entry_point_name;
683 struct vtn_value *entry_point;
684 struct vtn_value *workgroup_size_builtin;
685
686 uint32_t *interface_ids;
687 size_t interface_ids_count;
688
689 struct vtn_function *func;
690 struct list_head functions;
691
692 struct hash_table *strings;
693
694 /* Current function parameter index */
695 unsigned func_param_idx;
696
697 /* false by default, set to true by the ContractionOff execution mode */
698 bool exact;
699
700 /* when a physical memory model is choosen */
701 bool physical_ptrs;
702
703 /* memory model specified by OpMemoryModel */
704 unsigned mem_model;
705 };
706
707 const char *
708 vtn_string_literal(struct vtn_builder *b, const uint32_t *words,
709 unsigned word_count, unsigned *words_used);
710
711 nir_def *
712 vtn_pointer_to_ssa(struct vtn_builder *b, struct vtn_pointer *ptr);
713 struct vtn_pointer *
714 vtn_pointer_from_ssa(struct vtn_builder *b, nir_def *ssa,
715 struct vtn_type *ptr_type);
716
717 struct vtn_ssa_value *
718 vtn_const_ssa_value(struct vtn_builder *b, nir_constant *constant,
719 const struct glsl_type *type);
720
721 static inline struct vtn_value *
vtn_untyped_value(struct vtn_builder * b,uint32_t value_id)722 vtn_untyped_value(struct vtn_builder *b, uint32_t value_id)
723 {
724 vtn_fail_if(value_id >= b->value_id_bound,
725 "SPIR-V id %u is out-of-bounds", value_id);
726 return &b->values[value_id];
727 }
728
729 void vtn_print_value(struct vtn_builder *b, struct vtn_value *val, FILE *f);
730 void vtn_dump_values(struct vtn_builder *b, FILE *f);
731
732 static inline uint32_t
vtn_id_for_value(struct vtn_builder * b,struct vtn_value * value)733 vtn_id_for_value(struct vtn_builder *b, struct vtn_value *value)
734 {
735 vtn_fail_if(value <= b->values, "vtn_value pointer outside the range of valid values");
736 uint32_t value_id = value - b->values;
737 vtn_fail_if(value_id >= b->value_id_bound, "vtn_value pointer outside the range of valid values");
738 return value_id;
739 }
740
741 /* Consider not using this function directly and instead use
742 * vtn_push_ssa/vtn_push_pointer so that appropriate applying of
743 * decorations is handled by common code.
744 */
745 static inline struct vtn_value *
vtn_push_value(struct vtn_builder * b,uint32_t value_id,enum vtn_value_type value_type)746 vtn_push_value(struct vtn_builder *b, uint32_t value_id,
747 enum vtn_value_type value_type)
748 {
749 struct vtn_value *val = vtn_untyped_value(b, value_id);
750
751 vtn_fail_if(value_type == vtn_value_type_ssa,
752 "Do not call vtn_push_value for value_type_ssa. Use "
753 "vtn_push_ssa_value instead.");
754
755 vtn_fail_if(val->value_type != vtn_value_type_invalid,
756 "SPIR-V id %u has already been written by another instruction",
757 value_id);
758
759 val->value_type = value_type;
760
761 return &b->values[value_id];
762 }
763
764 /* These separated fail functions exist so the helpers like vtn_value()
765 * can be inlined with minimal code size impact. This allows the failure
766 * handling to have more detailed output without harming callers.
767 */
768
769 void _vtn_fail_value_type_mismatch(struct vtn_builder *b, uint32_t value_id,
770 enum vtn_value_type value_type);
771 void _vtn_fail_value_not_pointer(struct vtn_builder *b, uint32_t value_id);
772
773 static inline struct vtn_value *
vtn_value(struct vtn_builder * b,uint32_t value_id,enum vtn_value_type value_type)774 vtn_value(struct vtn_builder *b, uint32_t value_id,
775 enum vtn_value_type value_type)
776 {
777 struct vtn_value *val = vtn_untyped_value(b, value_id);
778 if (unlikely(val->value_type != value_type))
779 _vtn_fail_value_type_mismatch(b, value_id, value_type);
780 return val;
781 }
782
783 static inline struct vtn_value *
vtn_pointer_value(struct vtn_builder * b,uint32_t value_id)784 vtn_pointer_value(struct vtn_builder *b, uint32_t value_id)
785 {
786 struct vtn_value *val = vtn_untyped_value(b, value_id);
787 if (unlikely(val->value_type != vtn_value_type_pointer &&
788 !val->is_null_constant))
789 _vtn_fail_value_not_pointer(b, value_id);
790 return val;
791 }
792
793 static inline struct vtn_pointer *
vtn_value_to_pointer(struct vtn_builder * b,struct vtn_value * value)794 vtn_value_to_pointer(struct vtn_builder *b, struct vtn_value *value)
795 {
796 if (value->is_null_constant) {
797 vtn_assert(glsl_type_is_vector_or_scalar(value->type->type));
798 nir_def *const_ssa =
799 vtn_const_ssa_value(b, value->constant, value->type->type)->def;
800 return vtn_pointer_from_ssa(b, const_ssa, value->type);
801 }
802 vtn_assert(value->value_type == vtn_value_type_pointer);
803 return value->pointer;
804 }
805
806 static inline struct vtn_pointer *
vtn_pointer(struct vtn_builder * b,uint32_t value_id)807 vtn_pointer(struct vtn_builder *b, uint32_t value_id)
808 {
809 return vtn_value_to_pointer(b, vtn_pointer_value(b, value_id));
810 }
811
812 bool
813 vtn_set_instruction_result_type(struct vtn_builder *b, SpvOp opcode,
814 const uint32_t *w, unsigned count);
815
816 static inline uint64_t
vtn_constant_uint(struct vtn_builder * b,uint32_t value_id)817 vtn_constant_uint(struct vtn_builder *b, uint32_t value_id)
818 {
819 struct vtn_value *val = vtn_value(b, value_id, vtn_value_type_constant);
820
821 vtn_fail_if(val->type->base_type != vtn_base_type_scalar ||
822 !glsl_type_is_integer(val->type->type),
823 "Expected id %u to be an integer constant", value_id);
824
825 switch (glsl_get_bit_size(val->type->type)) {
826 case 8: return val->constant->values[0].u8;
827 case 16: return val->constant->values[0].u16;
828 case 32: return val->constant->values[0].u32;
829 case 64: return val->constant->values[0].u64;
830 default: unreachable("Invalid bit size");
831 }
832 }
833
834 static inline int64_t
vtn_constant_int(struct vtn_builder * b,uint32_t value_id)835 vtn_constant_int(struct vtn_builder *b, uint32_t value_id)
836 {
837 struct vtn_value *val = vtn_value(b, value_id, vtn_value_type_constant);
838
839 vtn_fail_if(val->type->base_type != vtn_base_type_scalar ||
840 !glsl_type_is_integer(val->type->type),
841 "Expected id %u to be an integer constant", value_id);
842
843 switch (glsl_get_bit_size(val->type->type)) {
844 case 8: return val->constant->values[0].i8;
845 case 16: return val->constant->values[0].i16;
846 case 32: return val->constant->values[0].i32;
847 case 64: return val->constant->values[0].i64;
848 default: unreachable("Invalid bit size");
849 }
850 }
851
852 static inline struct vtn_type *
vtn_get_value_type(struct vtn_builder * b,uint32_t value_id)853 vtn_get_value_type(struct vtn_builder *b, uint32_t value_id)
854 {
855 struct vtn_value *val = vtn_untyped_value(b, value_id);
856 vtn_fail_if(val->type == NULL, "Value %u does not have a type", value_id);
857 return val->type;
858 }
859
860 static inline struct vtn_type *
vtn_get_type(struct vtn_builder * b,uint32_t value_id)861 vtn_get_type(struct vtn_builder *b, uint32_t value_id)
862 {
863 return vtn_value(b, value_id, vtn_value_type_type)->type;
864 }
865
866 static inline struct vtn_block *
vtn_block(struct vtn_builder * b,uint32_t value_id)867 vtn_block(struct vtn_builder *b, uint32_t value_id)
868 {
869 return vtn_value(b, value_id, vtn_value_type_block)->block;
870 }
871
872 struct vtn_ssa_value *vtn_ssa_value(struct vtn_builder *b, uint32_t value_id);
873 struct vtn_value *vtn_push_ssa_value(struct vtn_builder *b, uint32_t value_id,
874 struct vtn_ssa_value *ssa);
875
876 nir_def *vtn_get_nir_ssa(struct vtn_builder *b, uint32_t value_id);
877 struct vtn_value *vtn_push_nir_ssa(struct vtn_builder *b, uint32_t value_id,
878 nir_def *def);
879 nir_deref_instr *vtn_get_deref_for_id(struct vtn_builder *b, uint32_t value_id);
880 nir_deref_instr *vtn_get_deref_for_ssa_value(struct vtn_builder *b, struct vtn_ssa_value *ssa);
881 struct vtn_value *vtn_push_var_ssa(struct vtn_builder *b, uint32_t value_id,
882 nir_variable *var);
883
884 struct vtn_value *vtn_push_pointer(struct vtn_builder *b,
885 uint32_t value_id,
886 struct vtn_pointer *ptr);
887
888 struct vtn_sampled_image {
889 nir_deref_instr *image;
890 nir_deref_instr *sampler;
891 };
892
893 nir_def *vtn_sampled_image_to_nir_ssa(struct vtn_builder *b,
894 struct vtn_sampled_image si);
895
896 void
897 vtn_copy_value(struct vtn_builder *b, uint32_t src_value_id,
898 uint32_t dst_value_id);
899
900 struct vtn_ssa_value *vtn_create_ssa_value(struct vtn_builder *b,
901 const struct glsl_type *type);
902 void vtn_set_ssa_value_var(struct vtn_builder *b, struct vtn_ssa_value *ssa, nir_variable *var);
903
904 struct vtn_ssa_value *vtn_ssa_transpose(struct vtn_builder *b,
905 struct vtn_ssa_value *src);
906
907 nir_deref_instr *vtn_nir_deref(struct vtn_builder *b, uint32_t id);
908
909 nir_deref_instr *vtn_pointer_to_deref(struct vtn_builder *b,
910 struct vtn_pointer *ptr);
911 nir_def *
912 vtn_pointer_to_offset(struct vtn_builder *b, struct vtn_pointer *ptr,
913 nir_def **index_out);
914
915 nir_deref_instr *
916 vtn_get_call_payload_for_location(struct vtn_builder *b, uint32_t location_id);
917
918 struct vtn_ssa_value *
919 vtn_local_load(struct vtn_builder *b, nir_deref_instr *src,
920 enum gl_access_qualifier access);
921
922 void vtn_local_store(struct vtn_builder *b, struct vtn_ssa_value *src,
923 nir_deref_instr *dest,
924 enum gl_access_qualifier access);
925
926 struct vtn_ssa_value *
927 vtn_variable_load(struct vtn_builder *b, struct vtn_pointer *src,
928 enum gl_access_qualifier access);
929
930 void vtn_variable_store(struct vtn_builder *b, struct vtn_ssa_value *src,
931 struct vtn_pointer *dest, enum gl_access_qualifier access);
932
933 void vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
934 const uint32_t *w, unsigned count);
935
936
937 typedef void (*vtn_decoration_foreach_cb)(struct vtn_builder *,
938 struct vtn_value *,
939 int member,
940 const struct vtn_decoration *,
941 void *);
942
943 void vtn_foreach_decoration(struct vtn_builder *b, struct vtn_value *value,
944 vtn_decoration_foreach_cb cb, void *data);
945
946 typedef void (*vtn_execution_mode_foreach_cb)(struct vtn_builder *,
947 struct vtn_value *,
948 const struct vtn_decoration *,
949 void *);
950
951 void vtn_foreach_execution_mode(struct vtn_builder *b, struct vtn_value *value,
952 vtn_execution_mode_foreach_cb cb, void *data);
953
954 nir_op vtn_nir_alu_op_for_spirv_opcode(struct vtn_builder *b,
955 SpvOp opcode, bool *swap, bool *exact,
956 unsigned src_bit_size, unsigned dst_bit_size);
957
958 void vtn_handle_alu(struct vtn_builder *b, SpvOp opcode,
959 const uint32_t *w, unsigned count);
960
961 void vtn_handle_integer_dot(struct vtn_builder *b, SpvOp opcode,
962 const uint32_t *w, unsigned count);
963
964 void vtn_handle_bitcast(struct vtn_builder *b, const uint32_t *w,
965 unsigned count);
966
967 void vtn_handle_no_contraction(struct vtn_builder *b, struct vtn_value *val);
968
969 void vtn_handle_fp_fast_math(struct vtn_builder *b, struct vtn_value *val);
970
971 void vtn_handle_subgroup(struct vtn_builder *b, SpvOp opcode,
972 const uint32_t *w, unsigned count);
973
974 bool vtn_handle_glsl450_instruction(struct vtn_builder *b, SpvOp ext_opcode,
975 const uint32_t *words, unsigned count);
976
977 bool vtn_handle_opencl_instruction(struct vtn_builder *b, SpvOp ext_opcode,
978 const uint32_t *words, unsigned count);
979 bool vtn_handle_opencl_core_instruction(struct vtn_builder *b, SpvOp opcode,
980 const uint32_t *w, unsigned count);
981
982 struct vtn_builder* vtn_create_builder(const uint32_t *words, size_t word_count,
983 gl_shader_stage stage, const char *entry_point_name,
984 const struct spirv_to_nir_options *options);
985
986 void vtn_handle_entry_point(struct vtn_builder *b, const uint32_t *w,
987 unsigned count);
988
989 void vtn_handle_debug_text(struct vtn_builder *b, SpvOp opcode,
990 const uint32_t *w, unsigned count);
991
992 void vtn_handle_decoration(struct vtn_builder *b, SpvOp opcode,
993 const uint32_t *w, unsigned count);
994
995 enum vtn_variable_mode vtn_storage_class_to_mode(struct vtn_builder *b,
996 SpvStorageClass class,
997 struct vtn_type *interface_type,
998 nir_variable_mode *nir_mode_out);
999
1000 nir_address_format vtn_mode_to_address_format(struct vtn_builder *b,
1001 enum vtn_variable_mode);
1002
1003 nir_rounding_mode vtn_rounding_mode_to_nir(struct vtn_builder *b,
1004 SpvFPRoundingMode mode);
1005
1006 static inline uint32_t
vtn_align_u32(uint32_t v,uint32_t a)1007 vtn_align_u32(uint32_t v, uint32_t a)
1008 {
1009 assert(a != 0 && a == (a & -((int32_t) a)));
1010 return (v + a - 1) & ~(a - 1);
1011 }
1012
1013 static inline uint64_t
vtn_u64_literal(const uint32_t * w)1014 vtn_u64_literal(const uint32_t *w)
1015 {
1016 return (uint64_t)w[1] << 32 | w[0];
1017 }
1018
1019 bool vtn_handle_amd_gcn_shader_instruction(struct vtn_builder *b, SpvOp ext_opcode,
1020 const uint32_t *words, unsigned count);
1021
1022 bool vtn_handle_amd_shader_ballot_instruction(struct vtn_builder *b, SpvOp ext_opcode,
1023 const uint32_t *w, unsigned count);
1024
1025 bool vtn_handle_amd_shader_trinary_minmax_instruction(struct vtn_builder *b, SpvOp ext_opcode,
1026 const uint32_t *words, unsigned count);
1027
1028 bool vtn_handle_amd_shader_explicit_vertex_parameter_instruction(struct vtn_builder *b,
1029 SpvOp ext_opcode,
1030 const uint32_t *words,
1031 unsigned count);
1032
1033 SpvMemorySemanticsMask vtn_mode_to_memory_semantics(enum vtn_variable_mode mode);
1034
1035 void vtn_emit_memory_barrier(struct vtn_builder *b, SpvScope scope,
1036 SpvMemorySemanticsMask semantics);
1037
1038 bool vtn_value_is_relaxed_precision(struct vtn_builder *b, struct vtn_value *val);
1039 nir_def *
1040 vtn_mediump_downconvert(struct vtn_builder *b, enum glsl_base_type base_type, nir_def *def);
1041 struct vtn_ssa_value *
1042 vtn_mediump_downconvert_value(struct vtn_builder *b, struct vtn_ssa_value *src);
1043 void vtn_mediump_upconvert_value(struct vtn_builder *b, struct vtn_ssa_value *value);
1044
1045 static inline int
cmp_uint32_t(const void * pa,const void * pb)1046 cmp_uint32_t(const void *pa, const void *pb)
1047 {
1048 uint32_t a = *((const uint32_t *)pa);
1049 uint32_t b = *((const uint32_t *)pb);
1050 if (a < b)
1051 return -1;
1052 if (a > b)
1053 return 1;
1054 return 0;
1055 }
1056
1057 void
1058 vtn_parse_switch(struct vtn_builder *b,
1059 const uint32_t *branch,
1060 struct list_head *case_list);
1061
1062 bool vtn_get_mem_operands(struct vtn_builder *b, const uint32_t *w, unsigned count,
1063 unsigned *idx, SpvMemoryAccessMask *access, unsigned *alignment,
1064 SpvScope *dest_scope, SpvScope *src_scope);
1065 void vtn_emit_make_visible_barrier(struct vtn_builder *b, SpvMemoryAccessMask access,
1066 SpvScope scope, enum vtn_variable_mode mode);
1067 void vtn_emit_make_available_barrier(struct vtn_builder *b, SpvMemoryAccessMask access,
1068 SpvScope scope, enum vtn_variable_mode mode);
1069
1070
1071 void vtn_handle_cooperative_type(struct vtn_builder *b, struct vtn_value *val,
1072 SpvOp opcode, const uint32_t *w, unsigned count);
1073 void vtn_handle_cooperative_instruction(struct vtn_builder *b, SpvOp opcode,
1074 const uint32_t *w, unsigned count);
1075 void vtn_handle_cooperative_alu(struct vtn_builder *b, struct vtn_value *dest_val,
1076 const struct glsl_type *dest_type, SpvOp opcode,
1077 const uint32_t *w, unsigned count);
1078 struct vtn_ssa_value *vtn_cooperative_matrix_extract(struct vtn_builder *b, struct vtn_ssa_value *mat,
1079 const uint32_t *indices, unsigned num_indices);
1080 struct vtn_ssa_value *vtn_cooperative_matrix_insert(struct vtn_builder *b, struct vtn_ssa_value *mat,
1081 struct vtn_ssa_value *insert,
1082 const uint32_t *indices, unsigned num_indices);
1083 nir_deref_instr *vtn_create_cmat_temporary(struct vtn_builder *b,
1084 const struct glsl_type *t, const char *name);
1085
1086 gl_shader_stage vtn_stage_for_execution_model(SpvExecutionModel model);
1087
1088 #endif /* _VTN_PRIVATE_H_ */
1089