src/compiler/spirv/vtn_private.h

   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  * Authors:
  24  *    Jason Ekstrand (jason@jlekstrand.net)
  25  *
  26  */
  27
  28 #ifndef _VTN_PRIVATE_H_
  29 #define _VTN_PRIVATE_H_
  30
  31 #include <setjmp.h>
  32
  33 #include "nir/nir.h"
  34 #include "nir/nir_builder.h"
  35 #include "util/u_dynarray.h"
  36 #include "nir_spirv.h"
  37 #include "spirv.h"
  38
  39 struct vtn_builder;
  40 struct vtn_decoration;
  41
  42 void vtn_log(struct vtn_builder *b, enum nir_spirv_debug_level level,
  43              size_t spirv_offset, const char *message);
  44
  45 void vtn_logf(struct vtn_builder *b, enum nir_spirv_debug_level level,
  46               size_t spirv_offset, const char *fmt, ...) PRINTFLIKE(4, 5);
  47
  48 #define vtn_info(...) vtn_logf(b, NIR_SPIRV_DEBUG_LEVEL_INFO, 0, __VA_ARGS__)
  49
  50 void _vtn_warn(struct vtn_builder *b, const char *file, unsigned line,
  51                const char *fmt, ...) PRINTFLIKE(4, 5);
  52 #define vtn_warn(...) _vtn_warn(b, __FILE__, __LINE__, __VA_ARGS__)
  53
  54 void _vtn_err(struct vtn_builder *b, const char *file, unsigned line,
  55                const char *fmt, ...) PRINTFLIKE(4, 5);
  56 #define vtn_err(...) _vtn_err(b, __FILE__, __LINE__, __VA_ARGS__)
  57
  58 /** Fail SPIR-V parsing
  59  *
  60  * This function logs an error and then bails out of the shader compile using
  61  * longjmp.  This being safe relies on two things:
  62  *
  63  *  1) We must guarantee that setjmp is called after allocating the builder
  64  *     and setting up b->debug (so that logging works) but before before any
  65  *     errors have a chance to occur.
  66  *
  67  *  2) While doing the SPIR-V -> NIR conversion, we need to be careful to
  68  *     ensure that all heap allocations happen through ralloc and are parented
  69  *     to the builder.  This way they will get properly cleaned up on error.
  70  *
  71  *  3) We must ensure that _vtn_fail is never called while a mutex lock or a
  72  *     reference to any other resource is held with the exception of ralloc
  73  *     objects which are parented to the builder.
  74  *
  75  * So long as these two things continue to hold, we can easily longjmp back to
  76  * spirv_to_nir(), clean up the builder, and return NULL.
  77  */
  78 NORETURN void
  79 _vtn_fail(struct vtn_builder *b, const char *file, unsigned line,
  80              const char *fmt, ...) PRINTFLIKE(4, 5);
  81
  82 #define vtn_fail(...) _vtn_fail(b, __FILE__, __LINE__, __VA_ARGS__)
  83
  84 /** Fail if the given expression evaluates to true */
  85 #define vtn_fail_if(expr, ...) \
  86    do { \
  87       if (unlikely(expr)) \
  88          vtn_fail(__VA_ARGS__); \
  89    } while (0)
  90
  91 #define _vtn_fail_with(t, msg, v) \
  92    vtn_fail("%s: %s (%u)\n", msg, spirv_ ## t ## _to_string(v), v)
  93
  94 #define vtn_fail_with_decoration(msg, v) _vtn_fail_with(decoration, msg, v)
  95 #define vtn_fail_with_opcode(msg, v)     _vtn_fail_with(op, msg, v)
  96
  97 /** Assert that a condition is true and, if it isn't, vtn_fail
  98  *
  99  * This macro is transitional only and should not be used in new code.  Use
 100  * vtn_fail_if and provide a real message instead.
 101  */
 102 #define vtn_assert(expr) \
 103    do { \
 104       if (!likely(expr)) \
 105          vtn_fail("%s", #expr); \
 106    } while (0)
 107
 108 enum vtn_value_type {
 109    vtn_value_type_invalid = 0,
 110    vtn_value_type_undef,
 111    vtn_value_type_string,
 112    vtn_value_type_decoration_group,
 113    vtn_value_type_type,
 114    vtn_value_type_constant,
 115    vtn_value_type_pointer,
 116    vtn_value_type_function,
 117    vtn_value_type_block,
 118    vtn_value_type_ssa,
 119    vtn_value_type_extension,
 120    vtn_value_type_image_pointer,
 121    vtn_value_type_sampled_image,
 122 };
 123
 124 enum vtn_branch_type {
 125    vtn_branch_type_none,
 126    vtn_branch_type_if_merge,
 127    vtn_branch_type_switch_break,
 128    vtn_branch_type_switch_fallthrough,
 129    vtn_branch_type_loop_break,
 130    vtn_branch_type_loop_continue,
 131    vtn_branch_type_loop_back_edge,
 132    vtn_branch_type_discard,
 133    vtn_branch_type_return,
 134 };
 135
 136 enum vtn_cf_node_type {
 137    vtn_cf_node_type_block,
 138    vtn_cf_node_type_if,
 139    vtn_cf_node_type_loop,
 140    vtn_cf_node_type_case,
 141    vtn_cf_node_type_switch,
 142    vtn_cf_node_type_function,
 143 };
 144
 145 struct vtn_cf_node {
 146    struct list_head link;
 147    struct vtn_cf_node *parent;
 148    enum vtn_cf_node_type type;
 149 };
 150
 151 struct vtn_loop {
 152    struct vtn_cf_node node;
 153
 154    /* The main body of the loop */
 155    struct list_head body;
 156
 157    /* The "continue" part of the loop.  This gets executed after the body
 158     * and is where you go when you hit a continue.
 159     */
 160    struct list_head cont_body;
 161
 162    struct vtn_block *header_block;
 163    struct vtn_block *cont_block;
 164    struct vtn_block *break_block;
 165
 166    SpvLoopControlMask control;
 167 };
 168
 169 struct vtn_if {
 170    struct vtn_cf_node node;
 171
 172    uint32_t condition;
 173
 174    enum vtn_branch_type then_type;
 175    struct list_head then_body;
 176
 177    enum vtn_branch_type else_type;
 178    struct list_head else_body;
 179
 180    struct vtn_block *merge_block;
 181
 182    SpvSelectionControlMask control;
 183 };
 184
 185 struct vtn_case {
 186    struct vtn_cf_node node;
 187
 188    enum vtn_branch_type type;
 189    struct list_head body;
 190
 191    /* The fallthrough case, if any */
 192    struct vtn_case *fallthrough;
 193
 194    /* The uint32_t values that map to this case */
 195    struct util_dynarray values;
 196
 197    /* True if this is the default case */
 198    bool is_default;
 199
 200    /* Initialized to false; used when sorting the list of cases */
 201    bool visited;
 202 };
 203
 204 struct vtn_switch {
 205    struct vtn_cf_node node;
 206
 207    uint32_t selector;
 208
 209    struct list_head cases;
 210
 211    struct vtn_block *break_block;
 212 };
 213
 214 struct vtn_block {
 215    struct vtn_cf_node node;
 216
 217    /** A pointer to the label instruction */
 218    const uint32_t *label;
 219
 220    /** A pointer to the merge instruction (or NULL if non exists) */
 221    const uint32_t *merge;
 222
 223    /** A pointer to the branch instruction that ends this block */
 224    const uint32_t *branch;
 225
 226    enum vtn_branch_type branch_type;
 227
 228    /* The CF node for which this is a merge target
 229     *
 230     * The SPIR-V spec requires that any given block can be the merge target
 231     * for at most one merge instruction.  If this block is a merge target,
 232     * this points back to the block containing that merge instruction.
 233     */
 234    struct vtn_cf_node *merge_cf_node;
 235
 236    /** Points to the loop that this block starts (if it starts a loop) */
 237    struct vtn_loop *loop;
 238
 239    /** Points to the switch case started by this block (if any) */
 240    struct vtn_case *switch_case;
 241
 242    /** Every block ends in a nop intrinsic so that we can find it again */
 243    nir_intrinsic_instr *end_nop;
 244 };
 245
 246 struct vtn_function {
 247    struct vtn_cf_node node;
 248
 249    struct vtn_type *type;
 250
 251    bool referenced;
 252    bool emitted;
 253
 254    nir_function_impl *impl;
 255    struct vtn_block *start_block;
 256
 257    struct list_head body;
 258
 259    const uint32_t *end;
 260
 261    SpvFunctionControlMask control;
 262 };
 263
 264 #define VTN_DECL_CF_NODE_CAST(_type)               \
 265 static inline struct vtn_##_type *                 \
 266 vtn_cf_node_as_##_type(struct vtn_cf_node *node)   \
 267 {                                                  \
 268    assert(node->type == vtn_cf_node_type_##_type); \
 269    return (struct vtn_##_type *)node;              \
 270 }
 271
 272 VTN_DECL_CF_NODE_CAST(block)
 273 VTN_DECL_CF_NODE_CAST(loop)
 274 VTN_DECL_CF_NODE_CAST(if)
 275 VTN_DECL_CF_NODE_CAST(case)
 276 VTN_DECL_CF_NODE_CAST(switch)
 277 VTN_DECL_CF_NODE_CAST(function)
 278
 279 #define vtn_foreach_cf_node(node, cf_list) \
 280    list_for_each_entry(struct vtn_cf_node, node, cf_list, link)
 281
 282 typedef bool (*vtn_instruction_handler)(struct vtn_builder *, SpvOp,
 283                                         const uint32_t *, unsigned);
 284
 285 void vtn_build_cfg(struct vtn_builder *b, const uint32_t *words,
 286                    const uint32_t *end);
 287 void vtn_function_emit(struct vtn_builder *b, struct vtn_function *func,
 288                        vtn_instruction_handler instruction_handler);
 289 void vtn_handle_function_call(struct vtn_builder *b, SpvOp opcode,
 290                               const uint32_t *w, unsigned count);
 291
 292 const uint32_t *
 293 vtn_foreach_instruction(struct vtn_builder *b, const uint32_t *start,
 294                         const uint32_t *end, vtn_instruction_handler handler);
 295
 296 struct vtn_ssa_value {
 297    union {
 298       nir_ssa_def *def;
 299       struct vtn_ssa_value **elems;
 300    };
 301
 302    /* For matrices, if this is non-NULL, then this value is actually the
 303     * transpose of some other value.  The value that `transposed` points to
 304     * always dominates this value.
 305     */
 306    struct vtn_ssa_value *transposed;
 307
 308    const struct glsl_type *type;
 309 };
 310
 311 enum vtn_base_type {
 312    vtn_base_type_void,
 313    vtn_base_type_scalar,
 314    vtn_base_type_vector,
 315    vtn_base_type_matrix,
 316    vtn_base_type_array,
 317    vtn_base_type_struct,
 318    vtn_base_type_pointer,
 319    vtn_base_type_image,
 320    vtn_base_type_sampler,
 321    vtn_base_type_sampled_image,
 322    vtn_base_type_function,
 323 };
 324
 325 struct vtn_type {
 326    enum vtn_base_type base_type;
 327
 328    const struct glsl_type *type;
 329
 330    /* The SPIR-V id of the given type. */
 331    uint32_t id;
 332
 333    /* Specifies the length of complex types.
 334     *
 335     * For Workgroup pointers, this is the size of the referenced type.
 336     */
 337    unsigned length;
 338
 339    /* for arrays, matrices and pointers, the array stride */
 340    unsigned stride;
 341
 342    /* Access qualifiers */
 343    enum gl_access_qualifier access;
 344
 345    union {
 346       /* Members for scalar, vector, and array-like types */
 347       struct {
 348          /* for arrays, the vtn_type for the elements of the array */
 349          struct vtn_type *array_element;
 350
 351          /* for matrices, whether the matrix is stored row-major */
 352          bool row_major:1;
 353
 354          /* Whether this type, or a parent type, has been decorated as a
 355           * builtin
 356           */
 357          bool is_builtin:1;
 358
 359          /* Which built-in to use */
 360          SpvBuiltIn builtin;
 361       };
 362
 363       /* Members for struct types */
 364       struct {
 365          /* for structures, the vtn_type for each member */
 366          struct vtn_type **members;
 367
 368          /* for structs, the offset of each member */
 369          unsigned *offsets;
 370
 371          /* for structs, whether it was decorated as a "non-SSBO-like" block */
 372          bool block:1;
 373
 374          /* for structs, whether it was decorated as an "SSBO-like" block */
 375          bool buffer_block:1;
 376
 377          /* for structs with block == true, whether this is a builtin block
 378           * (i.e. a block that contains only builtins).
 379           */
 380          bool builtin_block:1;
 381
 382          /* for structs and unions it specifies the minimum alignment of the
 383           * members. 0 means packed.
 384           *
 385           * Set by CPacked and Alignment Decorations in kernels.
 386           */
 387          bool packed:1;
 388       };
 389
 390       /* Members for pointer types */
 391       struct {
 392          /* For pointers, the vtn_type for dereferenced type */
 393          struct vtn_type *deref;
 394
 395          /* Storage class for pointers */
 396          SpvStorageClass storage_class;
 397
 398          /* Required alignment for pointers */
 399          uint32_t align;
 400       };
 401
 402       /* Members for image types */
 403       struct {
 404          /* Image format for image_load_store type images */
 405          unsigned image_format;
 406
 407          /* Access qualifier for storage images */
 408          SpvAccessQualifier access_qualifier;
 409       };
 410
 411       /* Members for sampled image types */
 412       struct {
 413          /* For sampled images, the image type */
 414          struct vtn_type *image;
 415       };
 416
 417       /* Members for function types */
 418       struct {
 419          /* For functions, the vtn_type for each parameter */
 420          struct vtn_type **params;
 421
 422          /* Return type for functions */
 423          struct vtn_type *return_type;
 424       };
 425    };
 426 };
 427
 428 bool vtn_type_contains_block(struct vtn_builder *b, struct vtn_type *type);
 429
 430 bool vtn_types_compatible(struct vtn_builder *b,
 431                           struct vtn_type *t1, struct vtn_type *t2);
 432
 433 struct vtn_type *vtn_type_without_array(struct vtn_type *type);
 434
 435 struct vtn_variable;
 436
 437 enum vtn_access_mode {
 438    vtn_access_mode_id,
 439    vtn_access_mode_literal,
 440 };
 441
 442 struct vtn_access_link {
 443    enum vtn_access_mode mode;
 444    int64_t id;
 445 };
 446
 447 struct vtn_access_chain {
 448    uint32_t length;
 449
 450    /** Whether or not to treat the base pointer as an array.  This is only
 451     * true if this access chain came from an OpPtrAccessChain.
 452     */
 453    bool ptr_as_array;
 454
 455    /* Access qualifiers */
 456    enum gl_access_qualifier access;
 457
 458    /** Struct elements and array offsets.
 459     *
 460     * This is an array of 1 so that it can conveniently be created on the
 461     * stack but the real length is given by the length field.
 462     */
 463    struct vtn_access_link link[1];
 464 };
 465
 466 enum vtn_variable_mode {
 467    vtn_variable_mode_function,
 468    vtn_variable_mode_private,
 469    vtn_variable_mode_uniform,
 470    vtn_variable_mode_atomic_counter,
 471    vtn_variable_mode_ubo,
 472    vtn_variable_mode_ssbo,
 473    vtn_variable_mode_phys_ssbo,
 474    vtn_variable_mode_push_constant,
 475    vtn_variable_mode_workgroup,
 476    vtn_variable_mode_cross_workgroup,
 477    vtn_variable_mode_input,
 478    vtn_variable_mode_output,
 479    vtn_variable_mode_image,
 480 };
 481
 482 struct vtn_pointer {
 483    /** The variable mode for the referenced data */
 484    enum vtn_variable_mode mode;
 485
 486    /** The dereferenced type of this pointer */
 487    struct vtn_type *type;
 488
 489    /** The pointer type of this pointer
 490     *
 491     * This may be NULL for some temporary pointers constructed as part of a
 492     * large load, store, or copy.  It MUST be valid for all pointers which are
 493     * stored as SPIR-V SSA values.
 494     */
 495    struct vtn_type *ptr_type;
 496
 497    /** The referenced variable, if known
 498     *
 499     * This field may be NULL if the pointer uses a (block_index, offset) pair
 500     * instead of an access chain or if the access chain starts at a deref.
 501     */
 502    struct vtn_variable *var;
 503
 504    /** The NIR deref corresponding to this pointer */
 505    nir_deref_instr *deref;
 506
 507    /** A (block_index, offset) pair representing a UBO or SSBO position. */
 508    struct nir_ssa_def *block_index;
 509    struct nir_ssa_def *offset;
 510
 511    /* Access qualifiers */
 512    enum gl_access_qualifier access;
 513 };
 514
 515 bool vtn_mode_uses_ssa_offset(struct vtn_builder *b,
 516                               enum vtn_variable_mode mode);
 517
 518 static inline bool vtn_pointer_uses_ssa_offset(struct vtn_builder *b,
 519                                                struct vtn_pointer *ptr)
 520 {
 521    return vtn_mode_uses_ssa_offset(b, ptr->mode);
 522 }
 523
 524
 525 struct vtn_variable {
 526    enum vtn_variable_mode mode;
 527
 528    struct vtn_type *type;
 529
 530    unsigned descriptor_set;
 531    unsigned binding;
 532    bool explicit_binding;
 533    unsigned offset;
 534    unsigned input_attachment_index;
 535    bool patch;
 536
 537    nir_variable *var;
 538
 539    /* If the variable is a struct with a location set on it then this will be
 540     * stored here. This will be used to calculate locations for members that
 541     * don’t have their own explicit location.
 542     */
 543    int base_location;
 544
 545    int shared_location;
 546
 547    /**
 548     * In some early released versions of GLSLang, it implemented all function
 549     * calls by making copies of all parameters into temporary variables and
 550     * passing those variables into the function.  It even did so for samplers
 551     * and images which violates the SPIR-V spec.  Unfortunately, two games
 552     * (Talos Principle and Doom) shipped with this old version of GLSLang and
 553     * also happen to pass samplers into functions.  Talos Principle received
 554     * an update fairly shortly after release with an updated GLSLang.  Doom,
 555     * on the other hand, has never received an update so we need to work
 556     * around this GLSLang issue in SPIR-V -> NIR.  Hopefully, we can drop this
 557     * hack at some point in the future.
 558     */
 559    struct vtn_pointer *copy_prop_sampler;
 560
 561    /* Access qualifiers. */
 562    enum gl_access_qualifier access;
 563 };
 564
 565 struct vtn_image_pointer {
 566    struct vtn_pointer *image;
 567    nir_ssa_def *coord;
 568    nir_ssa_def *sample;
 569    nir_ssa_def *lod;
 570 };
 571
 572 struct vtn_sampled_image {
 573    struct vtn_pointer *image; /* Image or array of images */
 574    struct vtn_pointer *sampler; /* Sampler */
 575 };
 576
 577 struct vtn_value {
 578    enum vtn_value_type value_type;
 579    const char *name;
 580    struct vtn_decoration *decoration;
 581    struct vtn_type *type;
 582    union {
 583       char *str;
 584       nir_constant *constant;
 585       struct vtn_pointer *pointer;
 586       struct vtn_image_pointer *image;
 587       struct vtn_sampled_image *sampled_image;
 588       struct vtn_function *func;
 589       struct vtn_block *block;
 590       struct vtn_ssa_value *ssa;
 591       vtn_instruction_handler ext_handler;
 592    };
 593 };
 594
 595 #define VTN_DEC_DECORATION -1
 596 #define VTN_DEC_EXECUTION_MODE -2
 597 #define VTN_DEC_STRUCT_MEMBER0 0
 598
 599 struct vtn_decoration {
 600    struct vtn_decoration *next;
 601
 602    /* Specifies how to apply this decoration.  Negative values represent a
 603     * decoration or execution mode. (See the VTN_DEC_ #defines above.)
 604     * Non-negative values specify that it applies to a structure member.
 605     */
 606    int scope;
 607
 608    const uint32_t *operands;
 609    struct vtn_value *group;
 610
 611    union {
 612       SpvDecoration decoration;
 613       SpvExecutionMode exec_mode;
 614    };
 615 };
 616
 617 struct vtn_builder {
 618    nir_builder nb;
 619
 620    /* Used by vtn_fail to jump back to the beginning of SPIR-V compilation */
 621    jmp_buf fail_jump;
 622
 623    const uint32_t *spirv;
 624    size_t spirv_word_count;
 625
 626    nir_shader *shader;
 627    struct spirv_to_nir_options *options;
 628    struct vtn_block *block;
 629
 630    /* Current offset, file, line, and column.  Useful for debugging.  Set
 631     * automatically by vtn_foreach_instruction.
 632     */
 633    size_t spirv_offset;
 634    char *file;
 635    int line, col;
 636
 637    /*
 638     * In SPIR-V, constants are global, whereas in NIR, the load_const
 639     * instruction we use is per-function. So while we parse each function, we
 640     * keep a hash table of constants we've resolved to nir_ssa_value's so
 641     * far, and we lazily resolve them when we see them used in a function.
 642     */
 643    struct hash_table *const_table;
 644
 645    /*
 646     * Map from phi instructions (pointer to the start of the instruction)
 647     * to the variable corresponding to it.
 648     */
 649    struct hash_table *phi_table;
 650
 651    unsigned num_specializations;
 652    struct nir_spirv_specialization *specializations;
 653
 654    unsigned value_id_bound;
 655    struct vtn_value *values;
 656
 657    /* True if we should watch out for GLSLang issue #179 */
 658    bool wa_glslang_179;
 659
 660    /* True if we need to fix up CS OpControlBarrier */
 661    bool wa_glslang_cs_barrier;
 662
 663    gl_shader_stage entry_point_stage;
 664    const char *entry_point_name;
 665    struct vtn_value *entry_point;
 666    struct vtn_value *workgroup_size_builtin;
 667    bool variable_pointers;
 668
 669    struct vtn_function *func;
 670    struct list_head functions;
 671
 672    /* Current function parameter index */
 673    unsigned func_param_idx;
 674
 675    bool has_loop_continue;
 676
 677    /* false by default, set to true by the ContractionOff execution mode */
 678    bool exact;
 679
 680    /* when a physical memory model is choosen */
 681    bool physical_ptrs;
 682
 683    /* memory model specified by OpMemoryModel */
 684    unsigned mem_model;
 685 };
 686
 687 nir_ssa_def *
 688 vtn_pointer_to_ssa(struct vtn_builder *b, struct vtn_pointer *ptr);
 689 struct vtn_pointer *
 690 vtn_pointer_from_ssa(struct vtn_builder *b, nir_ssa_def *ssa,
 691                      struct vtn_type *ptr_type);
 692
 693 static inline struct vtn_value *
 694 vtn_untyped_value(struct vtn_builder *b, uint32_t value_id)
 695 {
 696    vtn_fail_if(value_id >= b->value_id_bound,
 697                "SPIR-V id %u is out-of-bounds", value_id);
 698    return &b->values[value_id];
 699 }
 700
 701 /* Consider not using this function directly and instead use
 702  * vtn_push_ssa/vtn_push_pointer so that appropriate applying of
 703  * decorations is handled by common code.
 704  */
 705 static inline struct vtn_value *
 706 vtn_push_value(struct vtn_builder *b, uint32_t value_id,
 707                enum vtn_value_type value_type)
 708 {
 709    struct vtn_value *val = vtn_untyped_value(b, value_id);
 710
 711    vtn_fail_if(value_type == vtn_value_type_ssa,
 712                "Do not call vtn_push_value for value_type_ssa.  Use "
 713                "vtn_push_ssa_value instead.");
 714
 715    vtn_fail_if(val->value_type != vtn_value_type_invalid,
 716                "SPIR-V id %u has already been written by another instruction",
 717                value_id);
 718
 719    val->value_type = value_type;
 720
 721    return &b->values[value_id];
 722 }
 723
 724 static inline struct vtn_value *
 725 vtn_value(struct vtn_builder *b, uint32_t value_id,
 726           enum vtn_value_type value_type)
 727 {
 728    struct vtn_value *val = vtn_untyped_value(b, value_id);
 729    vtn_fail_if(val->value_type != value_type,
 730                "SPIR-V id %u is the wrong kind of value", value_id);
 731    return val;
 732 }
 733
 734 bool
 735 vtn_set_instruction_result_type(struct vtn_builder *b, SpvOp opcode,
 736                                 const uint32_t *w, unsigned count);
 737
 738 static inline uint64_t
 739 vtn_constant_uint(struct vtn_builder *b, uint32_t value_id)
 740 {
 741    struct vtn_value *val = vtn_value(b, value_id, vtn_value_type_constant);
 742
 743    vtn_fail_if(val->type->base_type != vtn_base_type_scalar ||
 744                !glsl_type_is_integer(val->type->type),
 745                "Expected id %u to be an integer constant", value_id);
 746
 747    switch (glsl_get_bit_size(val->type->type)) {
 748    case 8:  return val->constant->values[0].u8;
 749    case 16: return val->constant->values[0].u16;
 750    case 32: return val->constant->values[0].u32;
 751    case 64: return val->constant->values[0].u64;
 752    default: unreachable("Invalid bit size");
 753    }
 754 }
 755
 756 static inline int64_t
 757 vtn_constant_int(struct vtn_builder *b, uint32_t value_id)
 758 {
 759    struct vtn_value *val = vtn_value(b, value_id, vtn_value_type_constant);
 760
 761    vtn_fail_if(val->type->base_type != vtn_base_type_scalar ||
 762                !glsl_type_is_integer(val->type->type),
 763                "Expected id %u to be an integer constant", value_id);
 764
 765    switch (glsl_get_bit_size(val->type->type)) {
 766    case 8:  return val->constant->values[0].i8;
 767    case 16: return val->constant->values[0].i16;
 768    case 32: return val->constant->values[0].i32;
 769    case 64: return val->constant->values[0].i64;
 770    default: unreachable("Invalid bit size");
 771    }
 772 }
 773
 774 static inline struct vtn_type *
 775 vtn_get_value_type(struct vtn_builder *b, uint32_t value_id)
 776 {
 777    struct vtn_value *val = vtn_untyped_value(b, value_id);
 778    vtn_fail_if(val->type == NULL, "Value %u does not have a type", value_id);
 779    return val->type;
 780 }
 781
 782 static inline struct vtn_type *
 783 vtn_get_type(struct vtn_builder *b, uint32_t value_id)
 784 {
 785    return vtn_value(b, value_id, vtn_value_type_type)->type;
 786 }
 787
 788 struct vtn_ssa_value *vtn_ssa_value(struct vtn_builder *b, uint32_t value_id);
 789 struct vtn_value *vtn_push_ssa_value(struct vtn_builder *b, uint32_t value_id,
 790                                      struct vtn_ssa_value *ssa);
 791
 792 nir_ssa_def *vtn_get_nir_ssa(struct vtn_builder *b, uint32_t value_id);
 793 struct vtn_value *vtn_push_nir_ssa(struct vtn_builder *b, uint32_t value_id,
 794                                    nir_ssa_def *def);
 795
 796 struct vtn_value *vtn_push_pointer(struct vtn_builder *b,
 797                                    uint32_t value_id,
 798                                    struct vtn_pointer *ptr);
 799
 800 void
 801 vtn_copy_value(struct vtn_builder *b, uint32_t src_value_id,
 802                uint32_t dst_value_id);
 803
 804 struct vtn_ssa_value *vtn_create_ssa_value(struct vtn_builder *b,
 805                                            const struct glsl_type *type);
 806
 807 struct vtn_ssa_value *vtn_ssa_transpose(struct vtn_builder *b,
 808                                         struct vtn_ssa_value *src);
 809
 810 nir_deref_instr *vtn_nir_deref(struct vtn_builder *b, uint32_t id);
 811
 812 nir_deref_instr *vtn_pointer_to_deref(struct vtn_builder *b,
 813                                       struct vtn_pointer *ptr);
 814 nir_ssa_def *
 815 vtn_pointer_to_offset(struct vtn_builder *b, struct vtn_pointer *ptr,
 816                       nir_ssa_def **index_out);
 817
 818 struct vtn_ssa_value *
 819 vtn_local_load(struct vtn_builder *b, nir_deref_instr *src,
 820                enum gl_access_qualifier access);
 821
 822 void vtn_local_store(struct vtn_builder *b, struct vtn_ssa_value *src,
 823                      nir_deref_instr *dest,
 824                      enum gl_access_qualifier access);
 825
 826 struct vtn_ssa_value *
 827 vtn_variable_load(struct vtn_builder *b, struct vtn_pointer *src);
 828
 829 void vtn_variable_store(struct vtn_builder *b, struct vtn_ssa_value *src,
 830                         struct vtn_pointer *dest);
 831
 832 void vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
 833                           const uint32_t *w, unsigned count);
 834
 835
 836 typedef void (*vtn_decoration_foreach_cb)(struct vtn_builder *,
 837                                           struct vtn_value *,
 838                                           int member,
 839                                           const struct vtn_decoration *,
 840                                           void *);
 841
 842 void vtn_foreach_decoration(struct vtn_builder *b, struct vtn_value *value,
 843                             vtn_decoration_foreach_cb cb, void *data);
 844
 845 typedef void (*vtn_execution_mode_foreach_cb)(struct vtn_builder *,
 846                                               struct vtn_value *,
 847                                               const struct vtn_decoration *,
 848                                               void *);
 849
 850 void vtn_foreach_execution_mode(struct vtn_builder *b, struct vtn_value *value,
 851                                 vtn_execution_mode_foreach_cb cb, void *data);
 852
 853 nir_op vtn_nir_alu_op_for_spirv_opcode(struct vtn_builder *b,
 854                                        SpvOp opcode, bool *swap,
 855                                        unsigned src_bit_size, unsigned dst_bit_size);
 856
 857 void vtn_handle_alu(struct vtn_builder *b, SpvOp opcode,
 858                     const uint32_t *w, unsigned count);
 859
 860 void vtn_handle_bitcast(struct vtn_builder *b, const uint32_t *w,
 861                         unsigned count);
 862
 863 void vtn_handle_subgroup(struct vtn_builder *b, SpvOp opcode,
 864                          const uint32_t *w, unsigned count);
 865
 866 bool vtn_handle_glsl450_instruction(struct vtn_builder *b, SpvOp ext_opcode,
 867                                     const uint32_t *words, unsigned count);
 868
 869 bool vtn_handle_opencl_instruction(struct vtn_builder *b, SpvOp ext_opcode,
 870                                    const uint32_t *words, unsigned count);
 871
 872 struct vtn_builder* vtn_create_builder(const uint32_t *words, size_t word_count,
 873                                        gl_shader_stage stage, const char *entry_point_name,
 874                                        const struct spirv_to_nir_options *options);
 875
 876 void vtn_handle_entry_point(struct vtn_builder *b, const uint32_t *w,
 877                             unsigned count);
 878
 879 void vtn_handle_decoration(struct vtn_builder *b, SpvOp opcode,
 880                            const uint32_t *w, unsigned count);
 881
 882 enum vtn_variable_mode vtn_storage_class_to_mode(struct vtn_builder *b,
 883                                                  SpvStorageClass class,
 884                                                  struct vtn_type *interface_type,
 885                                                  nir_variable_mode *nir_mode_out);
 886
 887 nir_address_format vtn_mode_to_address_format(struct vtn_builder *b,
 888                                               enum vtn_variable_mode);
 889
 890 static inline uint32_t
 891 vtn_align_u32(uint32_t v, uint32_t a)
 892 {
 893    assert(a != 0 && a == (a & -((int32_t) a)));
 894    return (v + a - 1) & ~(a - 1);
 895 }
 896
 897 static inline uint64_t
 898 vtn_u64_literal(const uint32_t *w)
 899 {
 900    return (uint64_t)w[1] << 32 | w[0];
 901 }
 902
 903 bool vtn_handle_amd_gcn_shader_instruction(struct vtn_builder *b, SpvOp ext_opcode,
 904                                            const uint32_t *words, unsigned count);
 905
 906 bool vtn_handle_amd_shader_ballot_instruction(struct vtn_builder *b, SpvOp ext_opcode,
 907                                               const uint32_t *w, unsigned count);
 908
 909 bool vtn_handle_amd_shader_trinary_minmax_instruction(struct vtn_builder *b, SpvOp ext_opcode,
 910                                                       const uint32_t *words, unsigned count);
 911
 912 bool vtn_handle_amd_shader_explicit_vertex_parameter_instruction(struct vtn_builder *b,
 913                                                                  SpvOp ext_opcode,
 914                                                                  const uint32_t *words,
 915                                                                  unsigned count);
 916
 917 SpvMemorySemanticsMask vtn_storage_class_to_memory_semantics(SpvStorageClass sc);
 918
 919 void vtn_emit_memory_barrier(struct vtn_builder *b, SpvScope scope,
 920                              SpvMemorySemanticsMask semantics);
 921
 922 #endif /* _VTN_PRIVATE_H_ */