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