2
* Copyright © 2015 Intel Corporation
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:
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
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
24
* Jason Ekstrand (jason@jlekstrand.net)
28
#ifndef _VTN_PRIVATE_H_
29
#define _VTN_PRIVATE_H_
34
#include "nir/nir_builder.h"
35
#include "util/u_dynarray.h"
36
#include "nir_spirv.h"
38
#include "vtn_generator_ids.h"
41
struct vtn_decoration;
43
/* setjmp/longjmp is broken on MinGW: https://sourceforge.net/p/mingw-w64/bugs/406/ */
44
#if defined(__MINGW32__) && !defined(_UCRT)
45
#define vtn_setjmp __builtin_setjmp
46
#define vtn_longjmp __builtin_longjmp
48
#define vtn_setjmp setjmp
49
#define vtn_longjmp longjmp
52
void vtn_log(struct vtn_builder *b, enum nir_spirv_debug_level level,
53
size_t spirv_offset, const char *message);
55
void vtn_logf(struct vtn_builder *b, enum nir_spirv_debug_level level,
56
size_t spirv_offset, const char *fmt, ...) PRINTFLIKE(4, 5);
58
#define vtn_info(...) vtn_logf(b, NIR_SPIRV_DEBUG_LEVEL_INFO, 0, __VA_ARGS__)
60
void _vtn_warn(struct vtn_builder *b, const char *file, unsigned line,
61
const char *fmt, ...) PRINTFLIKE(4, 5);
62
#define vtn_warn(...) _vtn_warn(b, __FILE__, __LINE__, __VA_ARGS__)
64
void _vtn_err(struct vtn_builder *b, const char *file, unsigned line,
65
const char *fmt, ...) PRINTFLIKE(4, 5);
66
#define vtn_err(...) _vtn_err(b, __FILE__, __LINE__, __VA_ARGS__)
68
/** Fail SPIR-V parsing
70
* This function logs an error and then bails out of the shader compile using
71
* longjmp. This being safe relies on two things:
73
* 1) We must guarantee that setjmp is called after allocating the builder
74
* and setting up b->debug (so that logging works) but before before any
75
* errors have a chance to occur.
77
* 2) While doing the SPIR-V -> NIR conversion, we need to be careful to
78
* ensure that all heap allocations happen through ralloc and are parented
79
* to the builder. This way they will get properly cleaned up on error.
81
* 3) We must ensure that _vtn_fail is never called while a mutex lock or a
82
* reference to any other resource is held with the exception of ralloc
83
* objects which are parented to the builder.
85
* So long as these two things continue to hold, we can easily longjmp back to
86
* spirv_to_nir(), clean up the builder, and return NULL.
89
_vtn_fail(struct vtn_builder *b, const char *file, unsigned line,
90
const char *fmt, ...) PRINTFLIKE(4, 5);
92
#define vtn_fail(...) _vtn_fail(b, __FILE__, __LINE__, __VA_ARGS__)
94
/** Fail if the given expression evaluates to true */
95
#define vtn_fail_if(expr, ...) \
98
vtn_fail(__VA_ARGS__); \
101
#define _vtn_fail_with(t, msg, v) \
102
vtn_fail("%s: %s (%u)\n", msg, spirv_ ## t ## _to_string(v), v)
104
#define vtn_fail_with_decoration(msg, v) _vtn_fail_with(decoration, msg, v)
105
#define vtn_fail_with_opcode(msg, v) _vtn_fail_with(op, msg, v)
107
/** Assert that a condition is true and, if it isn't, vtn_fail
109
* This macro is transitional only and should not be used in new code. Use
110
* vtn_fail_if and provide a real message instead.
112
#define vtn_assert(expr) \
115
vtn_fail("%s", #expr); \
118
enum vtn_value_type {
119
vtn_value_type_invalid = 0,
120
vtn_value_type_undef,
121
vtn_value_type_string,
122
vtn_value_type_decoration_group,
124
vtn_value_type_constant,
125
vtn_value_type_pointer,
126
vtn_value_type_function,
127
vtn_value_type_block,
129
vtn_value_type_extension,
130
vtn_value_type_image_pointer,
133
enum vtn_branch_type {
134
vtn_branch_type_none,
135
vtn_branch_type_if_merge,
136
vtn_branch_type_switch_break,
137
vtn_branch_type_switch_fallthrough,
138
vtn_branch_type_loop_break,
139
vtn_branch_type_loop_continue,
140
vtn_branch_type_loop_back_edge,
141
vtn_branch_type_discard,
142
vtn_branch_type_terminate_invocation,
143
vtn_branch_type_ignore_intersection,
144
vtn_branch_type_terminate_ray,
145
vtn_branch_type_return,
148
enum vtn_cf_node_type {
149
vtn_cf_node_type_block,
151
vtn_cf_node_type_loop,
152
vtn_cf_node_type_case,
153
vtn_cf_node_type_switch,
154
vtn_cf_node_type_function,
158
struct list_head link;
159
struct vtn_cf_node *parent;
160
enum vtn_cf_node_type type;
164
struct vtn_cf_node node;
166
/* The main body of the loop */
167
struct list_head body;
169
/* The "continue" part of the loop. This gets executed after the body
170
* and is where you go when you hit a continue.
172
struct list_head cont_body;
174
struct vtn_block *header_block;
175
struct vtn_block *cont_block;
176
struct vtn_block *break_block;
178
SpvLoopControlMask control;
182
struct vtn_cf_node node;
184
enum vtn_branch_type then_type;
185
struct list_head then_body;
187
enum vtn_branch_type else_type;
188
struct list_head else_body;
190
struct vtn_block *header_block;
191
struct vtn_block *merge_block;
193
SpvSelectionControlMask control;
197
struct vtn_cf_node node;
199
struct vtn_block *block;
201
enum vtn_branch_type type;
202
struct list_head body;
204
/* The fallthrough case, if any */
205
struct vtn_case *fallthrough;
207
/* The uint32_t values that map to this case */
208
struct util_dynarray values;
210
/* True if this is the default case */
213
/* Initialized to false; used when sorting the list of cases */
218
struct vtn_cf_node node;
222
struct list_head cases;
224
struct vtn_block *break_block;
228
struct vtn_cf_node node;
230
/** A pointer to the label instruction */
231
const uint32_t *label;
233
/** A pointer to the merge instruction (or NULL if non exists) */
234
const uint32_t *merge;
236
/** A pointer to the branch instruction that ends this block */
237
const uint32_t *branch;
239
enum vtn_branch_type branch_type;
241
/* The CF node for which this is a merge target
243
* The SPIR-V spec requires that any given block can be the merge target
244
* for at most one merge instruction. If this block is a merge target,
245
* this points back to the block containing that merge instruction.
247
struct vtn_cf_node *merge_cf_node;
249
/** Points to the loop that this block starts (if it starts a loop) */
250
struct vtn_loop *loop;
252
/** Points to the switch case started by this block (if any) */
253
struct vtn_case *switch_case;
255
/** Every block ends in a nop intrinsic so that we can find it again */
256
nir_intrinsic_instr *end_nop;
258
/** attached nir_block */
259
struct nir_block *block;
262
struct vtn_function {
263
struct vtn_cf_node node;
265
struct vtn_type *type;
270
nir_function *nir_func;
271
struct vtn_block *start_block;
273
struct list_head body;
277
SpvLinkageType linkage;
278
SpvFunctionControlMask control;
281
#define VTN_DECL_CF_NODE_CAST(_type) \
282
static inline struct vtn_##_type * \
283
vtn_cf_node_as_##_type(struct vtn_cf_node *node) \
285
assert(node->type == vtn_cf_node_type_##_type); \
286
return (struct vtn_##_type *)node; \
289
VTN_DECL_CF_NODE_CAST(block)
290
VTN_DECL_CF_NODE_CAST(loop)
291
VTN_DECL_CF_NODE_CAST(if)
292
VTN_DECL_CF_NODE_CAST(case)
293
VTN_DECL_CF_NODE_CAST(switch)
294
VTN_DECL_CF_NODE_CAST(function)
296
#define vtn_foreach_cf_node(node, cf_list) \
297
list_for_each_entry(struct vtn_cf_node, node, cf_list, link)
299
typedef bool (*vtn_instruction_handler)(struct vtn_builder *, SpvOp,
300
const uint32_t *, unsigned);
302
void vtn_build_cfg(struct vtn_builder *b, const uint32_t *words,
303
const uint32_t *end);
304
void vtn_function_emit(struct vtn_builder *b, struct vtn_function *func,
305
vtn_instruction_handler instruction_handler);
306
void vtn_handle_function_call(struct vtn_builder *b, SpvOp opcode,
307
const uint32_t *w, unsigned count);
310
vtn_foreach_instruction(struct vtn_builder *b, const uint32_t *start,
311
const uint32_t *end, vtn_instruction_handler handler);
313
struct vtn_ssa_value {
316
struct vtn_ssa_value **elems;
319
/* For matrices, if this is non-NULL, then this value is actually the
320
* transpose of some other value. The value that `transposed` points to
321
* always dominates this value.
323
struct vtn_ssa_value *transposed;
325
const struct glsl_type *type;
330
vtn_base_type_scalar,
331
vtn_base_type_vector,
332
vtn_base_type_matrix,
334
vtn_base_type_struct,
335
vtn_base_type_pointer,
337
vtn_base_type_sampler,
338
vtn_base_type_sampled_image,
339
vtn_base_type_accel_struct,
340
vtn_base_type_ray_query,
341
vtn_base_type_function,
346
enum vtn_base_type base_type;
348
const struct glsl_type *type;
350
/* The SPIR-V id of the given type. */
353
/* Specifies the length of complex types.
355
* For Workgroup pointers, this is the size of the referenced type.
359
/* for arrays, matrices and pointers, the array stride */
362
/* Access qualifiers */
363
enum gl_access_qualifier access;
366
/* Members for scalar, vector, and array-like types */
368
/* for arrays, the vtn_type for the elements of the array */
369
struct vtn_type *array_element;
371
/* for matrices, whether the matrix is stored row-major */
374
/* Whether this type, or a parent type, has been decorated as a
379
/* Which built-in to use */
383
/* Members for struct types */
385
/* for structures, the vtn_type for each member */
386
struct vtn_type **members;
388
/* for structs, the offset of each member */
391
/* for structs, whether it was decorated as a "non-SSBO-like" block */
394
/* for structs, whether it was decorated as an "SSBO-like" block */
397
/* for structs with block == true, whether this is a builtin block
398
* (i.e. a block that contains only builtins).
400
bool builtin_block:1;
402
/* for structs and unions it specifies the minimum alignment of the
403
* members. 0 means packed.
405
* Set by CPacked and Alignment Decorations in kernels.
410
/* Members for pointer types */
412
/* For pointers, the vtn_type for dereferenced type */
413
struct vtn_type *deref;
415
/* Storage class for pointers */
416
SpvStorageClass storage_class;
418
/* Required alignment for pointers */
422
/* Members for image types */
424
/* GLSL image type for this type. This is not to be confused with
425
* vtn_type::type which is actually going to be the GLSL type for a
426
* pointer to an image, likely a uint32_t.
428
const struct glsl_type *glsl_image;
430
/* Image format for image_load_store type images */
431
unsigned image_format;
433
/* Access qualifier for storage images */
434
SpvAccessQualifier access_qualifier;
437
/* Members for sampled image types */
439
/* For sampled images, the image type */
440
struct vtn_type *image;
443
/* Members for function types */
445
/* For functions, the vtn_type for each parameter */
446
struct vtn_type **params;
448
/* Return type for functions */
449
struct vtn_type *return_type;
454
bool vtn_type_contains_block(struct vtn_builder *b, struct vtn_type *type);
456
bool vtn_types_compatible(struct vtn_builder *b,
457
struct vtn_type *t1, struct vtn_type *t2);
459
struct vtn_type *vtn_type_without_array(struct vtn_type *type);
463
enum vtn_access_mode {
465
vtn_access_mode_literal,
468
struct vtn_access_link {
469
enum vtn_access_mode mode;
473
struct vtn_access_chain {
476
/** Whether or not to treat the base pointer as an array. This is only
477
* true if this access chain came from an OpPtrAccessChain.
481
/* Access qualifiers */
482
enum gl_access_qualifier access;
484
/** Struct elements and array offsets.
486
* This is an array of 1 so that it can conveniently be created on the
487
* stack but the real length is given by the length field.
489
struct vtn_access_link link[1];
492
enum vtn_variable_mode {
493
vtn_variable_mode_function,
494
vtn_variable_mode_private,
495
vtn_variable_mode_uniform,
496
vtn_variable_mode_atomic_counter,
497
vtn_variable_mode_ubo,
498
vtn_variable_mode_ssbo,
499
vtn_variable_mode_phys_ssbo,
500
vtn_variable_mode_push_constant,
501
vtn_variable_mode_workgroup,
502
vtn_variable_mode_cross_workgroup,
503
vtn_variable_mode_task_payload,
504
vtn_variable_mode_generic,
505
vtn_variable_mode_constant,
506
vtn_variable_mode_input,
507
vtn_variable_mode_output,
508
vtn_variable_mode_image,
509
vtn_variable_mode_accel_struct,
510
vtn_variable_mode_call_data,
511
vtn_variable_mode_call_data_in,
512
vtn_variable_mode_ray_payload,
513
vtn_variable_mode_ray_payload_in,
514
vtn_variable_mode_hit_attrib,
515
vtn_variable_mode_shader_record,
519
/** The variable mode for the referenced data */
520
enum vtn_variable_mode mode;
522
/** The dereferenced type of this pointer */
523
struct vtn_type *type;
525
/** The pointer type of this pointer
527
* This may be NULL for some temporary pointers constructed as part of a
528
* large load, store, or copy. It MUST be valid for all pointers which are
529
* stored as SPIR-V SSA values.
531
struct vtn_type *ptr_type;
533
/** The referenced variable, if known
535
* This field may be NULL if the pointer uses a (block_index, offset) pair
536
* instead of an access chain or if the access chain starts at a deref.
538
struct vtn_variable *var;
540
/** The NIR deref corresponding to this pointer */
541
nir_deref_instr *deref;
543
/** A (block_index, offset) pair representing a UBO or SSBO position. */
544
struct nir_ssa_def *block_index;
545
struct nir_ssa_def *offset;
547
/* Access qualifiers */
548
enum gl_access_qualifier access;
551
struct vtn_variable {
552
enum vtn_variable_mode mode;
554
struct vtn_type *type;
556
unsigned descriptor_set;
558
bool explicit_binding;
560
unsigned input_attachment_index;
564
/* If the variable is a struct with a location set on it then this will be
565
* stored here. This will be used to calculate locations for members that
566
* don’t have their own explicit location.
571
* In some early released versions of GLSLang, it implemented all function
572
* calls by making copies of all parameters into temporary variables and
573
* passing those variables into the function. It even did so for samplers
574
* and images which violates the SPIR-V spec. Unfortunately, two games
575
* (Talos Principle and Doom) shipped with this old version of GLSLang and
576
* also happen to pass samplers into functions. Talos Principle received
577
* an update fairly shortly after release with an updated GLSLang. Doom,
578
* on the other hand, has never received an update so we need to work
579
* around this GLSLang issue in SPIR-V -> NIR. Hopefully, we can drop this
580
* hack at some point in the future.
582
struct vtn_pointer *copy_prop_sampler;
584
/* Access qualifiers. */
585
enum gl_access_qualifier access;
588
const struct glsl_type *
589
vtn_type_get_nir_type(struct vtn_builder *b, struct vtn_type *type,
590
enum vtn_variable_mode mode);
592
struct vtn_image_pointer {
593
nir_deref_instr *image;
600
enum vtn_value_type value_type;
602
/* Workaround for https://gitlab.freedesktop.org/mesa/mesa/-/issues/3406
603
* Only set for OpImage / OpSampledImage. Note that this is in addition
604
* the existence of a NonUniform decoration on this value.*/
605
uint32_t propagated_non_uniform : 1;
607
/* Valid for vtn_value_type_constant to indicate the value is OpConstantNull. */
608
bool is_null_constant:1;
610
/* Valid when all the members of the value are undef. */
611
bool is_undef_constant:1;
614
struct vtn_decoration *decoration;
615
struct vtn_type *type;
618
nir_constant *constant;
619
struct vtn_pointer *pointer;
620
struct vtn_image_pointer *image;
621
struct vtn_function *func;
622
struct vtn_block *block;
623
struct vtn_ssa_value *ssa;
624
vtn_instruction_handler ext_handler;
628
#define VTN_DEC_DECORATION -1
629
#define VTN_DEC_EXECUTION_MODE -2
630
#define VTN_DEC_STRUCT_MEMBER_NAME0 -3
631
#define VTN_DEC_STRUCT_MEMBER0 0
633
struct vtn_decoration {
634
struct vtn_decoration *next;
636
/* Different kinds of decorations are stored in a value,
637
the scope defines what decoration it refers to:
639
- VTN_DEC_DECORATION:
640
decoration associated with the value
641
- VTN_DEC_EXECUTION_MODE:
642
an execution mode associated with an entrypoint value
643
- VTN_DEC_STRUCT_MEMBER0 + m:
644
decoration associated with member m of a struct value
645
- VTN_DEC_STRUCT_MEMBER_NAME0 - m:
646
name of m'th member of a struct value
650
uint32_t num_operands;
651
const uint32_t *operands;
652
struct vtn_value *group;
655
SpvDecoration decoration;
656
SpvExecutionMode exec_mode;
657
const char *member_name;
664
/* Used by vtn_fail to jump back to the beginning of SPIR-V compilation */
667
const uint32_t *spirv;
668
size_t spirv_word_count;
672
struct spirv_to_nir_options *options;
673
struct vtn_block *block;
675
/* Current offset, file, line, and column. Useful for debugging. Set
676
* automatically by vtn_foreach_instruction.
683
* In SPIR-V, constants are global, whereas in NIR, the load_const
684
* instruction we use is per-function. So while we parse each function, we
685
* keep a hash table of constants we've resolved to nir_ssa_value's so
686
* far, and we lazily resolve them when we see them used in a function.
688
struct hash_table *const_table;
691
* Map from phi instructions (pointer to the start of the instruction)
692
* to the variable corresponding to it.
694
struct hash_table *phi_table;
696
/* In Vulkan, when lowering some modes variable access, the derefs of the
697
* variables are replaced with a resource index intrinsics, leaving the
698
* variable hanging. This set keeps track of them so they can be filtered
699
* (and not removed) in nir_remove_dead_variables.
701
struct set *vars_used_indirectly;
703
unsigned num_specializations;
704
struct nir_spirv_specialization *specializations;
706
unsigned value_id_bound;
707
struct vtn_value *values;
709
/* Information on the origin of the SPIR-V */
710
enum vtn_generator generator_id;
711
SpvSourceLanguage source_lang;
713
/* True if we need to fix up CS OpControlBarrier */
714
bool wa_glslang_cs_barrier;
716
/* True if we need to ignore undef initializers */
717
bool wa_llvm_spirv_ignore_workgroup_initializer;
719
/* Workaround discard bugs in HLSL -> SPIR-V compilers */
720
bool uses_demote_to_helper_invocation;
721
bool convert_discard_to_demote;
723
gl_shader_stage entry_point_stage;
724
const char *entry_point_name;
725
struct vtn_value *entry_point;
726
struct vtn_value *workgroup_size_builtin;
727
bool variable_pointers;
729
uint32_t *interface_ids;
730
size_t interface_ids_count;
732
struct vtn_function *func;
733
struct list_head functions;
735
/* Current function parameter index */
736
unsigned func_param_idx;
738
/* false by default, set to true by the ContractionOff execution mode */
741
/* when a physical memory model is choosen */
744
/* memory model specified by OpMemoryModel */
749
vtn_string_literal(struct vtn_builder *b, const uint32_t *words,
750
unsigned word_count, unsigned *words_used);
753
vtn_pointer_to_ssa(struct vtn_builder *b, struct vtn_pointer *ptr);
755
vtn_pointer_from_ssa(struct vtn_builder *b, nir_ssa_def *ssa,
756
struct vtn_type *ptr_type);
758
struct vtn_ssa_value *
759
vtn_const_ssa_value(struct vtn_builder *b, nir_constant *constant,
760
const struct glsl_type *type);
762
static inline struct vtn_value *
763
vtn_untyped_value(struct vtn_builder *b, uint32_t value_id)
765
vtn_fail_if(value_id >= b->value_id_bound,
766
"SPIR-V id %u is out-of-bounds", value_id);
767
return &b->values[value_id];
770
static inline uint32_t
771
vtn_id_for_value(struct vtn_builder *b, struct vtn_value *value)
773
vtn_fail_if(value <= b->values, "vtn_value pointer outside the range of valid values");
774
uint32_t value_id = value - b->values;
775
vtn_fail_if(value_id >= b->value_id_bound, "vtn_value pointer outside the range of valid values");
779
/* Consider not using this function directly and instead use
780
* vtn_push_ssa/vtn_push_pointer so that appropriate applying of
781
* decorations is handled by common code.
783
static inline struct vtn_value *
784
vtn_push_value(struct vtn_builder *b, uint32_t value_id,
785
enum vtn_value_type value_type)
787
struct vtn_value *val = vtn_untyped_value(b, value_id);
789
vtn_fail_if(value_type == vtn_value_type_ssa,
790
"Do not call vtn_push_value for value_type_ssa. Use "
791
"vtn_push_ssa_value instead.");
793
vtn_fail_if(val->value_type != vtn_value_type_invalid,
794
"SPIR-V id %u has already been written by another instruction",
797
val->value_type = value_type;
799
return &b->values[value_id];
802
static inline struct vtn_value *
803
vtn_value(struct vtn_builder *b, uint32_t value_id,
804
enum vtn_value_type value_type)
806
struct vtn_value *val = vtn_untyped_value(b, value_id);
807
vtn_fail_if(val->value_type != value_type,
808
"SPIR-V id %u is the wrong kind of value", value_id);
812
static inline struct vtn_value *
813
vtn_pointer_value(struct vtn_builder *b, uint32_t value_id)
815
struct vtn_value *val = vtn_untyped_value(b, value_id);
816
vtn_fail_if(val->value_type != vtn_value_type_pointer &&
817
!val->is_null_constant,
818
"SPIR-V id %u is the wrong kind of value", value_id);
822
static inline struct vtn_pointer *
823
vtn_value_to_pointer(struct vtn_builder *b, struct vtn_value *value)
825
if (value->is_null_constant) {
826
vtn_assert(glsl_type_is_vector_or_scalar(value->type->type));
827
nir_ssa_def *const_ssa =
828
vtn_const_ssa_value(b, value->constant, value->type->type)->def;
829
return vtn_pointer_from_ssa(b, const_ssa, value->type);
831
vtn_assert(value->value_type == vtn_value_type_pointer);
832
return value->pointer;
835
static inline struct vtn_pointer *
836
vtn_pointer(struct vtn_builder *b, uint32_t value_id)
838
return vtn_value_to_pointer(b, vtn_pointer_value(b, value_id));
842
vtn_set_instruction_result_type(struct vtn_builder *b, SpvOp opcode,
843
const uint32_t *w, unsigned count);
845
static inline uint64_t
846
vtn_constant_uint(struct vtn_builder *b, uint32_t value_id)
848
struct vtn_value *val = vtn_value(b, value_id, vtn_value_type_constant);
850
vtn_fail_if(val->type->base_type != vtn_base_type_scalar ||
851
!glsl_type_is_integer(val->type->type),
852
"Expected id %u to be an integer constant", value_id);
854
switch (glsl_get_bit_size(val->type->type)) {
855
case 8: return val->constant->values[0].u8;
856
case 16: return val->constant->values[0].u16;
857
case 32: return val->constant->values[0].u32;
858
case 64: return val->constant->values[0].u64;
859
default: unreachable("Invalid bit size");
863
static inline int64_t
864
vtn_constant_int(struct vtn_builder *b, uint32_t value_id)
866
struct vtn_value *val = vtn_value(b, value_id, vtn_value_type_constant);
868
vtn_fail_if(val->type->base_type != vtn_base_type_scalar ||
869
!glsl_type_is_integer(val->type->type),
870
"Expected id %u to be an integer constant", value_id);
872
switch (glsl_get_bit_size(val->type->type)) {
873
case 8: return val->constant->values[0].i8;
874
case 16: return val->constant->values[0].i16;
875
case 32: return val->constant->values[0].i32;
876
case 64: return val->constant->values[0].i64;
877
default: unreachable("Invalid bit size");
881
static inline struct vtn_type *
882
vtn_get_value_type(struct vtn_builder *b, uint32_t value_id)
884
struct vtn_value *val = vtn_untyped_value(b, value_id);
885
vtn_fail_if(val->type == NULL, "Value %u does not have a type", value_id);
889
static inline struct vtn_type *
890
vtn_get_type(struct vtn_builder *b, uint32_t value_id)
892
return vtn_value(b, value_id, vtn_value_type_type)->type;
895
struct vtn_ssa_value *vtn_ssa_value(struct vtn_builder *b, uint32_t value_id);
896
struct vtn_value *vtn_push_ssa_value(struct vtn_builder *b, uint32_t value_id,
897
struct vtn_ssa_value *ssa);
899
nir_ssa_def *vtn_get_nir_ssa(struct vtn_builder *b, uint32_t value_id);
900
struct vtn_value *vtn_push_nir_ssa(struct vtn_builder *b, uint32_t value_id,
903
struct vtn_value *vtn_push_pointer(struct vtn_builder *b,
905
struct vtn_pointer *ptr);
907
struct vtn_sampled_image {
908
nir_deref_instr *image;
909
nir_deref_instr *sampler;
912
nir_ssa_def *vtn_sampled_image_to_nir_ssa(struct vtn_builder *b,
913
struct vtn_sampled_image si);
916
vtn_copy_value(struct vtn_builder *b, uint32_t src_value_id,
917
uint32_t dst_value_id);
919
struct vtn_ssa_value *vtn_create_ssa_value(struct vtn_builder *b,
920
const struct glsl_type *type);
922
struct vtn_ssa_value *vtn_ssa_transpose(struct vtn_builder *b,
923
struct vtn_ssa_value *src);
925
nir_deref_instr *vtn_nir_deref(struct vtn_builder *b, uint32_t id);
927
nir_deref_instr *vtn_pointer_to_deref(struct vtn_builder *b,
928
struct vtn_pointer *ptr);
930
vtn_pointer_to_offset(struct vtn_builder *b, struct vtn_pointer *ptr,
931
nir_ssa_def **index_out);
934
vtn_get_call_payload_for_location(struct vtn_builder *b, uint32_t location_id);
936
struct vtn_ssa_value *
937
vtn_local_load(struct vtn_builder *b, nir_deref_instr *src,
938
enum gl_access_qualifier access);
940
void vtn_local_store(struct vtn_builder *b, struct vtn_ssa_value *src,
941
nir_deref_instr *dest,
942
enum gl_access_qualifier access);
944
struct vtn_ssa_value *
945
vtn_variable_load(struct vtn_builder *b, struct vtn_pointer *src,
946
enum gl_access_qualifier access);
948
void vtn_variable_store(struct vtn_builder *b, struct vtn_ssa_value *src,
949
struct vtn_pointer *dest, enum gl_access_qualifier access);
951
void vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
952
const uint32_t *w, unsigned count);
955
typedef void (*vtn_decoration_foreach_cb)(struct vtn_builder *,
958
const struct vtn_decoration *,
961
void vtn_foreach_decoration(struct vtn_builder *b, struct vtn_value *value,
962
vtn_decoration_foreach_cb cb, void *data);
964
typedef void (*vtn_execution_mode_foreach_cb)(struct vtn_builder *,
966
const struct vtn_decoration *,
969
void vtn_foreach_execution_mode(struct vtn_builder *b, struct vtn_value *value,
970
vtn_execution_mode_foreach_cb cb, void *data);
972
nir_op vtn_nir_alu_op_for_spirv_opcode(struct vtn_builder *b,
973
SpvOp opcode, bool *swap, bool *exact,
974
unsigned src_bit_size, unsigned dst_bit_size);
976
void vtn_handle_alu(struct vtn_builder *b, SpvOp opcode,
977
const uint32_t *w, unsigned count);
979
void vtn_handle_integer_dot(struct vtn_builder *b, SpvOp opcode,
980
const uint32_t *w, unsigned count);
982
void vtn_handle_bitcast(struct vtn_builder *b, const uint32_t *w,
985
void vtn_handle_no_contraction(struct vtn_builder *b, struct vtn_value *val);
987
void vtn_handle_subgroup(struct vtn_builder *b, SpvOp opcode,
988
const uint32_t *w, unsigned count);
990
bool vtn_handle_glsl450_instruction(struct vtn_builder *b, SpvOp ext_opcode,
991
const uint32_t *words, unsigned count);
993
bool vtn_handle_opencl_instruction(struct vtn_builder *b, SpvOp ext_opcode,
994
const uint32_t *words, unsigned count);
995
bool vtn_handle_opencl_core_instruction(struct vtn_builder *b, SpvOp opcode,
996
const uint32_t *w, unsigned count);
998
struct vtn_builder* vtn_create_builder(const uint32_t *words, size_t word_count,
999
gl_shader_stage stage, const char *entry_point_name,
1000
const struct spirv_to_nir_options *options);
1002
void vtn_handle_entry_point(struct vtn_builder *b, const uint32_t *w,
1005
void vtn_handle_decoration(struct vtn_builder *b, SpvOp opcode,
1006
const uint32_t *w, unsigned count);
1008
enum vtn_variable_mode vtn_storage_class_to_mode(struct vtn_builder *b,
1009
SpvStorageClass class,
1010
struct vtn_type *interface_type,
1011
nir_variable_mode *nir_mode_out);
1013
nir_address_format vtn_mode_to_address_format(struct vtn_builder *b,
1014
enum vtn_variable_mode);
1016
nir_rounding_mode vtn_rounding_mode_to_nir(struct vtn_builder *b,
1017
SpvFPRoundingMode mode);
1019
static inline uint32_t
1020
vtn_align_u32(uint32_t v, uint32_t a)
1022
assert(a != 0 && a == (a & -((int32_t) a)));
1023
return (v + a - 1) & ~(a - 1);
1026
static inline uint64_t
1027
vtn_u64_literal(const uint32_t *w)
1029
return (uint64_t)w[1] << 32 | w[0];
1032
bool vtn_handle_amd_gcn_shader_instruction(struct vtn_builder *b, SpvOp ext_opcode,
1033
const uint32_t *words, unsigned count);
1035
bool vtn_handle_amd_shader_ballot_instruction(struct vtn_builder *b, SpvOp ext_opcode,
1036
const uint32_t *w, unsigned count);
1038
bool vtn_handle_amd_shader_trinary_minmax_instruction(struct vtn_builder *b, SpvOp ext_opcode,
1039
const uint32_t *words, unsigned count);
1041
bool vtn_handle_amd_shader_explicit_vertex_parameter_instruction(struct vtn_builder *b,
1043
const uint32_t *words,
1046
SpvMemorySemanticsMask vtn_mode_to_memory_semantics(enum vtn_variable_mode mode);
1048
void vtn_emit_memory_barrier(struct vtn_builder *b, SpvScope scope,
1049
SpvMemorySemanticsMask semantics);
1052
cmp_uint32_t(const void *pa, const void *pb)
1054
uint32_t a = *((const uint32_t *)pa);
1055
uint32_t b = *((const uint32_t *)pb);
1063
#endif /* _VTN_PRIVATE_H_ */
added
removed
src/compiler/spirv/vtn_private.h
