naga/back/spv/
writer.rs

1use alloc::{string::String, vec, vec::Vec};
2
3use hashbrown::hash_map::Entry;
4use spirv::Word;
5
6use super::{
7    block::DebugInfoInner,
8    helpers::{contains_builtin, global_needs_wrapper, map_storage_class},
9    Block, BlockContext, CachedConstant, CachedExpressions, DebugInfo, EntryPointContext, Error,
10    Function, FunctionArgument, GlobalVariable, IdGenerator, Instruction, LocalImageType,
11    LocalType, LocalVariable, LogicalLayout, LookupFunctionType, LookupType, NumericType, Options,
12    PhysicalLayout, PipelineOptions, ResultMember, Writer, WriterFlags, BITS_PER_BYTE,
13};
14use crate::{
15    arena::{Handle, HandleVec, UniqueArena},
16    back::spv::{BindingInfo, WrappedFunction},
17    path_like::PathLike,
18    proc::{Alignment, TypeResolution},
19    valid::{FunctionInfo, ModuleInfo},
20};
21
22struct FunctionInterface<'a> {
23    varying_ids: &'a mut Vec<Word>,
24    stage: crate::ShaderStage,
25}
26
27impl Function {
28    pub(super) fn to_words(&self, sink: &mut impl Extend<Word>) {
29        self.signature.as_ref().unwrap().to_words(sink);
30        for argument in self.parameters.iter() {
31            argument.instruction.to_words(sink);
32        }
33        for (index, block) in self.blocks.iter().enumerate() {
34            Instruction::label(block.label_id).to_words(sink);
35            if index == 0 {
36                for local_var in self.variables.values() {
37                    local_var.instruction.to_words(sink);
38                }
39                for local_var in self.force_loop_bounding_vars.iter() {
40                    local_var.instruction.to_words(sink);
41                }
42                for internal_var in self.spilled_composites.values() {
43                    internal_var.instruction.to_words(sink);
44                }
45            }
46            for instruction in block.body.iter() {
47                instruction.to_words(sink);
48            }
49        }
50        Instruction::function_end().to_words(sink);
51    }
52}
53
54impl Writer {
55    pub fn new(options: &Options) -> Result<Self, Error> {
56        let (major, minor) = options.lang_version;
57        if major != 1 {
58            return Err(Error::UnsupportedVersion(major, minor));
59        }
60
61        let mut capabilities_used = crate::FastIndexSet::default();
62        capabilities_used.insert(spirv::Capability::Shader);
63
64        let mut id_gen = IdGenerator::default();
65        let gl450_ext_inst_id = id_gen.next();
66        let void_type = id_gen.next();
67
68        Ok(Writer {
69            physical_layout: PhysicalLayout::new(major, minor),
70            logical_layout: LogicalLayout::default(),
71            id_gen,
72            capabilities_available: options.capabilities.clone(),
73            capabilities_used,
74            extensions_used: crate::FastIndexSet::default(),
75            debugs: vec![],
76            annotations: vec![],
77            flags: options.flags,
78            bounds_check_policies: options.bounds_check_policies,
79            zero_initialize_workgroup_memory: options.zero_initialize_workgroup_memory,
80            force_loop_bounding: options.force_loop_bounding,
81            use_storage_input_output_16: options.use_storage_input_output_16,
82            void_type,
83            lookup_type: crate::FastHashMap::default(),
84            lookup_function: crate::FastHashMap::default(),
85            lookup_function_type: crate::FastHashMap::default(),
86            wrapped_functions: crate::FastHashMap::default(),
87            constant_ids: HandleVec::new(),
88            cached_constants: crate::FastHashMap::default(),
89            global_variables: HandleVec::new(),
90            binding_map: options.binding_map.clone(),
91            saved_cached: CachedExpressions::default(),
92            gl450_ext_inst_id,
93            temp_list: Vec::new(),
94            ray_get_committed_intersection_function: None,
95            ray_get_candidate_intersection_function: None,
96            io_f16_polyfills: super::f16_polyfill::F16IoPolyfill::new(
97                options.use_storage_input_output_16,
98            ),
99        })
100    }
101
102    /// Returns `(major, minor)` of the SPIR-V language version.
103    pub const fn lang_version(&self) -> (u8, u8) {
104        self.physical_layout.lang_version()
105    }
106
107    /// Reset `Writer` to its initial state, retaining any allocations.
108    ///
109    /// Why not just implement `Recyclable` for `Writer`? By design,
110    /// `Recyclable::recycle` requires ownership of the value, not just
111    /// `&mut`; see the trait documentation. But we need to use this method
112    /// from functions like `Writer::write`, which only have `&mut Writer`.
113    /// Workarounds include unsafe code (`core::ptr::read`, then `write`, ugh)
114    /// or something like a `Default` impl that returns an oddly-initialized
115    /// `Writer`, which is worse.
116    fn reset(&mut self) {
117        use super::recyclable::Recyclable;
118        use core::mem::take;
119
120        let mut id_gen = IdGenerator::default();
121        let gl450_ext_inst_id = id_gen.next();
122        let void_type = id_gen.next();
123
124        // Every field of the old writer that is not determined by the `Options`
125        // passed to `Writer::new` should be reset somehow.
126        let fresh = Writer {
127            // Copied from the old Writer:
128            flags: self.flags,
129            bounds_check_policies: self.bounds_check_policies,
130            zero_initialize_workgroup_memory: self.zero_initialize_workgroup_memory,
131            force_loop_bounding: self.force_loop_bounding,
132            use_storage_input_output_16: self.use_storage_input_output_16,
133            capabilities_available: take(&mut self.capabilities_available),
134            binding_map: take(&mut self.binding_map),
135
136            // Initialized afresh:
137            id_gen,
138            void_type,
139            gl450_ext_inst_id,
140
141            // Recycled:
142            capabilities_used: take(&mut self.capabilities_used).recycle(),
143            extensions_used: take(&mut self.extensions_used).recycle(),
144            physical_layout: self.physical_layout.clone().recycle(),
145            logical_layout: take(&mut self.logical_layout).recycle(),
146            debugs: take(&mut self.debugs).recycle(),
147            annotations: take(&mut self.annotations).recycle(),
148            lookup_type: take(&mut self.lookup_type).recycle(),
149            lookup_function: take(&mut self.lookup_function).recycle(),
150            lookup_function_type: take(&mut self.lookup_function_type).recycle(),
151            wrapped_functions: take(&mut self.wrapped_functions).recycle(),
152            constant_ids: take(&mut self.constant_ids).recycle(),
153            cached_constants: take(&mut self.cached_constants).recycle(),
154            global_variables: take(&mut self.global_variables).recycle(),
155            saved_cached: take(&mut self.saved_cached).recycle(),
156            temp_list: take(&mut self.temp_list).recycle(),
157            ray_get_candidate_intersection_function: None,
158            ray_get_committed_intersection_function: None,
159            io_f16_polyfills: take(&mut self.io_f16_polyfills).recycle(),
160        };
161
162        *self = fresh;
163
164        self.capabilities_used.insert(spirv::Capability::Shader);
165    }
166
167    /// Indicate that the code requires any one of the listed capabilities.
168    ///
169    /// If nothing in `capabilities` appears in the available capabilities
170    /// specified in the [`Options`] from which this `Writer` was created,
171    /// return an error. The `what` string is used in the error message to
172    /// explain what provoked the requirement. (If no available capabilities were
173    /// given, assume everything is available.)
174    ///
175    /// The first acceptable capability will be added to this `Writer`'s
176    /// [`capabilities_used`] table, and an `OpCapability` emitted for it in the
177    /// result. For this reason, more specific capabilities should be listed
178    /// before more general.
179    ///
180    /// [`capabilities_used`]: Writer::capabilities_used
181    pub(super) fn require_any(
182        &mut self,
183        what: &'static str,
184        capabilities: &[spirv::Capability],
185    ) -> Result<(), Error> {
186        match *capabilities {
187            [] => Ok(()),
188            [first, ..] => {
189                // Find the first acceptable capability, or return an error if
190                // there is none.
191                let selected = match self.capabilities_available {
192                    None => first,
193                    Some(ref available) => {
194                        match capabilities
195                            .iter()
196                            // need explicit type for hashbrown::HashSet::contains fn call to keep rustc happy
197                            .find(|cap| available.contains::<spirv::Capability>(cap))
198                        {
199                            Some(&cap) => cap,
200                            None => {
201                                return Err(Error::MissingCapabilities(what, capabilities.to_vec()))
202                            }
203                        }
204                    }
205                };
206                self.capabilities_used.insert(selected);
207                Ok(())
208            }
209        }
210    }
211
212    /// Indicate that the code requires all of the listed capabilities.
213    ///
214    /// If all entries of `capabilities` appear in the available capabilities
215    /// specified in the [`Options`] from which this `Writer` was created
216    /// (including the case where [`Options::capabilities`] is `None`), add
217    /// them all to this `Writer`'s [`capabilities_used`] table, and return
218    /// `Ok(())`. If at least one of the listed capabilities is not available,
219    /// do not add anything to the `capabilities_used` table, and return the
220    /// first unavailable requested capability, wrapped in `Err()`.
221    ///
222    /// This method is does not return an [`enum@Error`] in case of failure
223    /// because it may be used in cases where the caller can recover (e.g.,
224    /// with a polyfill) if the requested capabilities are not available. In
225    /// this case, it would be unnecessary work to find *all* the unavailable
226    /// requested capabilities, and to allocate a `Vec` for them, just so we
227    /// could return an [`Error::MissingCapabilities`]).
228    ///
229    /// [`capabilities_used`]: Writer::capabilities_used
230    pub(super) fn require_all(
231        &mut self,
232        capabilities: &[spirv::Capability],
233    ) -> Result<(), spirv::Capability> {
234        if let Some(ref available) = self.capabilities_available {
235            for requested in capabilities {
236                if !available.contains(requested) {
237                    return Err(*requested);
238                }
239            }
240        }
241
242        for requested in capabilities {
243            self.capabilities_used.insert(*requested);
244        }
245
246        Ok(())
247    }
248
249    /// Indicate that the code uses the given extension.
250    pub(super) fn use_extension(&mut self, extension: &'static str) {
251        self.extensions_used.insert(extension);
252    }
253
254    pub(super) fn get_type_id(&mut self, lookup_ty: LookupType) -> Word {
255        match self.lookup_type.entry(lookup_ty) {
256            Entry::Occupied(e) => *e.get(),
257            Entry::Vacant(e) => {
258                let local = match lookup_ty {
259                    LookupType::Handle(_handle) => unreachable!("Handles are populated at start"),
260                    LookupType::Local(local) => local,
261                };
262
263                let id = self.id_gen.next();
264                e.insert(id);
265                self.write_type_declaration_local(id, local);
266                id
267            }
268        }
269    }
270
271    pub(super) fn get_handle_type_id(&mut self, handle: Handle<crate::Type>) -> Word {
272        self.get_type_id(LookupType::Handle(handle))
273    }
274
275    pub(super) fn get_expression_lookup_type(&mut self, tr: &TypeResolution) -> LookupType {
276        match *tr {
277            TypeResolution::Handle(ty_handle) => LookupType::Handle(ty_handle),
278            TypeResolution::Value(ref inner) => {
279                let inner_local_type = self.localtype_from_inner(inner).unwrap();
280                LookupType::Local(inner_local_type)
281            }
282        }
283    }
284
285    pub(super) fn get_expression_type_id(&mut self, tr: &TypeResolution) -> Word {
286        let lookup_ty = self.get_expression_lookup_type(tr);
287        self.get_type_id(lookup_ty)
288    }
289
290    pub(super) fn get_localtype_id(&mut self, local: LocalType) -> Word {
291        self.get_type_id(LookupType::Local(local))
292    }
293
294    pub(super) fn get_pointer_type_id(&mut self, base: Word, class: spirv::StorageClass) -> Word {
295        self.get_type_id(LookupType::Local(LocalType::Pointer { base, class }))
296    }
297
298    pub(super) fn get_handle_pointer_type_id(
299        &mut self,
300        base: Handle<crate::Type>,
301        class: spirv::StorageClass,
302    ) -> Word {
303        let base_id = self.get_handle_type_id(base);
304        self.get_pointer_type_id(base_id, class)
305    }
306
307    pub(super) fn get_ray_query_pointer_id(&mut self) -> Word {
308        let rq_id = self.get_type_id(LookupType::Local(LocalType::RayQuery));
309        self.get_pointer_type_id(rq_id, spirv::StorageClass::Function)
310    }
311
312    /// Return a SPIR-V type for a pointer to `resolution`.
313    ///
314    /// The given `resolution` must be one that we can represent
315    /// either as a `LocalType::Pointer` or `LocalType::LocalPointer`.
316    pub(super) fn get_resolution_pointer_id(
317        &mut self,
318        resolution: &TypeResolution,
319        class: spirv::StorageClass,
320    ) -> Word {
321        let resolution_type_id = self.get_expression_type_id(resolution);
322        self.get_pointer_type_id(resolution_type_id, class)
323    }
324
325    pub(super) fn get_numeric_type_id(&mut self, numeric: NumericType) -> Word {
326        self.get_type_id(LocalType::Numeric(numeric).into())
327    }
328
329    pub(super) fn get_u32_type_id(&mut self) -> Word {
330        self.get_numeric_type_id(NumericType::Scalar(crate::Scalar::U32))
331    }
332
333    pub(super) fn get_f32_type_id(&mut self) -> Word {
334        self.get_numeric_type_id(NumericType::Scalar(crate::Scalar::F32))
335    }
336
337    pub(super) fn get_vec2u_type_id(&mut self) -> Word {
338        self.get_numeric_type_id(NumericType::Vector {
339            size: crate::VectorSize::Bi,
340            scalar: crate::Scalar::U32,
341        })
342    }
343
344    pub(super) fn get_vec2f_type_id(&mut self) -> Word {
345        self.get_numeric_type_id(NumericType::Vector {
346            size: crate::VectorSize::Bi,
347            scalar: crate::Scalar::F32,
348        })
349    }
350
351    pub(super) fn get_vec3u_type_id(&mut self) -> Word {
352        self.get_numeric_type_id(NumericType::Vector {
353            size: crate::VectorSize::Tri,
354            scalar: crate::Scalar::U32,
355        })
356    }
357
358    pub(super) fn get_f32_pointer_type_id(&mut self, class: spirv::StorageClass) -> Word {
359        let f32_id = self.get_f32_type_id();
360        self.get_pointer_type_id(f32_id, class)
361    }
362
363    pub(super) fn get_vec2u_pointer_type_id(&mut self, class: spirv::StorageClass) -> Word {
364        let vec2u_id = self.get_numeric_type_id(NumericType::Vector {
365            size: crate::VectorSize::Bi,
366            scalar: crate::Scalar::U32,
367        });
368        self.get_pointer_type_id(vec2u_id, class)
369    }
370
371    pub(super) fn get_vec3u_pointer_type_id(&mut self, class: spirv::StorageClass) -> Word {
372        let vec3u_id = self.get_numeric_type_id(NumericType::Vector {
373            size: crate::VectorSize::Tri,
374            scalar: crate::Scalar::U32,
375        });
376        self.get_pointer_type_id(vec3u_id, class)
377    }
378
379    pub(super) fn get_bool_type_id(&mut self) -> Word {
380        self.get_numeric_type_id(NumericType::Scalar(crate::Scalar::BOOL))
381    }
382
383    pub(super) fn get_vec2_bool_type_id(&mut self) -> Word {
384        self.get_numeric_type_id(NumericType::Vector {
385            size: crate::VectorSize::Bi,
386            scalar: crate::Scalar::BOOL,
387        })
388    }
389
390    pub(super) fn get_vec3_bool_type_id(&mut self) -> Word {
391        self.get_numeric_type_id(NumericType::Vector {
392            size: crate::VectorSize::Tri,
393            scalar: crate::Scalar::BOOL,
394        })
395    }
396
397    pub(super) fn decorate(&mut self, id: Word, decoration: spirv::Decoration, operands: &[Word]) {
398        self.annotations
399            .push(Instruction::decorate(id, decoration, operands));
400    }
401
402    /// Return `inner` as a `LocalType`, if that's possible.
403    ///
404    /// If `inner` can be represented as a `LocalType`, return
405    /// `Some(local_type)`.
406    ///
407    /// Otherwise, return `None`. In this case, the type must always be looked
408    /// up using a `LookupType::Handle`.
409    fn localtype_from_inner(&mut self, inner: &crate::TypeInner) -> Option<LocalType> {
410        Some(match *inner {
411            crate::TypeInner::Scalar(_)
412            | crate::TypeInner::Atomic(_)
413            | crate::TypeInner::Vector { .. }
414            | crate::TypeInner::Matrix { .. } => {
415                // We expect `NumericType::from_inner` to handle all
416                // these cases, so unwrap.
417                LocalType::Numeric(NumericType::from_inner(inner).unwrap())
418            }
419            crate::TypeInner::Pointer { base, space } => {
420                let base_type_id = self.get_handle_type_id(base);
421                LocalType::Pointer {
422                    base: base_type_id,
423                    class: map_storage_class(space),
424                }
425            }
426            crate::TypeInner::ValuePointer {
427                size,
428                scalar,
429                space,
430            } => {
431                let base_numeric_type = match size {
432                    Some(size) => NumericType::Vector { size, scalar },
433                    None => NumericType::Scalar(scalar),
434                };
435                LocalType::Pointer {
436                    base: self.get_numeric_type_id(base_numeric_type),
437                    class: map_storage_class(space),
438                }
439            }
440            crate::TypeInner::Image {
441                dim,
442                arrayed,
443                class,
444            } => LocalType::Image(LocalImageType::from_inner(dim, arrayed, class)),
445            crate::TypeInner::Sampler { comparison: _ } => LocalType::Sampler,
446            crate::TypeInner::AccelerationStructure { .. } => LocalType::AccelerationStructure,
447            crate::TypeInner::RayQuery { .. } => LocalType::RayQuery,
448            crate::TypeInner::Array { .. }
449            | crate::TypeInner::Struct { .. }
450            | crate::TypeInner::BindingArray { .. } => return None,
451        })
452    }
453
454    /// Emits code for any wrapper functions required by the expressions in ir_function.
455    /// The IDs of any emitted functions will be stored in [`Self::wrapped_functions`].
456    fn write_wrapped_functions(
457        &mut self,
458        ir_function: &crate::Function,
459        info: &FunctionInfo,
460        ir_module: &crate::Module,
461    ) -> Result<(), Error> {
462        log::trace!("Generating wrapped functions for {:?}", ir_function.name);
463
464        for (expr_handle, expr) in ir_function.expressions.iter() {
465            match *expr {
466                crate::Expression::Binary { op, left, right } => {
467                    let expr_ty_inner = info[expr_handle].ty.inner_with(&ir_module.types);
468                    if let Some(expr_ty) = NumericType::from_inner(expr_ty_inner) {
469                        match (op, expr_ty.scalar().kind) {
470                            // Division and modulo are undefined behaviour when the
471                            // dividend is the minimum representable value and the divisor
472                            // is negative one, or when the divisor is zero. These wrapped
473                            // functions override the divisor to one in these cases,
474                            // matching the WGSL spec.
475                            (
476                                crate::BinaryOperator::Divide | crate::BinaryOperator::Modulo,
477                                crate::ScalarKind::Sint | crate::ScalarKind::Uint,
478                            ) => {
479                                self.write_wrapped_binary_op(
480                                    op,
481                                    expr_ty,
482                                    &info[left].ty,
483                                    &info[right].ty,
484                                )?;
485                            }
486                            _ => {}
487                        }
488                    }
489                }
490                _ => {}
491            }
492        }
493
494        Ok(())
495    }
496
497    /// Write a SPIR-V function that performs the operator `op` with Naga IR semantics.
498    ///
499    /// Define a function that performs an integer division or modulo operation,
500    /// except that using a divisor of zero or causing signed overflow with a
501    /// divisor of -1 returns the numerator unchanged, rather than exhibiting
502    /// undefined behavior.
503    ///
504    /// Store the generated function's id in the [`wrapped_functions`] table.
505    ///
506    /// The operator `op` must be either [`Divide`] or [`Modulo`].
507    ///
508    /// # Panics
509    ///
510    /// The `return_type`, `left_type` or `right_type` arguments must all be
511    /// integer scalars or vectors. If not, this function panics.
512    ///
513    /// [`wrapped_functions`]: Writer::wrapped_functions
514    /// [`Divide`]: crate::BinaryOperator::Divide
515    /// [`Modulo`]: crate::BinaryOperator::Modulo
516    fn write_wrapped_binary_op(
517        &mut self,
518        op: crate::BinaryOperator,
519        return_type: NumericType,
520        left_type: &TypeResolution,
521        right_type: &TypeResolution,
522    ) -> Result<(), Error> {
523        let return_type_id = self.get_localtype_id(LocalType::Numeric(return_type));
524        let left_type_id = self.get_expression_type_id(left_type);
525        let right_type_id = self.get_expression_type_id(right_type);
526
527        // Check if we've already emitted this function.
528        let wrapped = WrappedFunction::BinaryOp {
529            op,
530            left_type_id,
531            right_type_id,
532        };
533        let function_id = match self.wrapped_functions.entry(wrapped) {
534            Entry::Occupied(_) => return Ok(()),
535            Entry::Vacant(e) => *e.insert(self.id_gen.next()),
536        };
537
538        let scalar = return_type.scalar();
539
540        if self.flags.contains(WriterFlags::DEBUG) {
541            let function_name = match op {
542                crate::BinaryOperator::Divide => "naga_div",
543                crate::BinaryOperator::Modulo => "naga_mod",
544                _ => unreachable!(),
545            };
546            self.debugs
547                .push(Instruction::name(function_id, function_name));
548        }
549        let mut function = Function::default();
550
551        let function_type_id = self.get_function_type(LookupFunctionType {
552            parameter_type_ids: vec![left_type_id, right_type_id],
553            return_type_id,
554        });
555        function.signature = Some(Instruction::function(
556            return_type_id,
557            function_id,
558            spirv::FunctionControl::empty(),
559            function_type_id,
560        ));
561
562        let lhs_id = self.id_gen.next();
563        let rhs_id = self.id_gen.next();
564        if self.flags.contains(WriterFlags::DEBUG) {
565            self.debugs.push(Instruction::name(lhs_id, "lhs"));
566            self.debugs.push(Instruction::name(rhs_id, "rhs"));
567        }
568        let left_par = Instruction::function_parameter(left_type_id, lhs_id);
569        let right_par = Instruction::function_parameter(right_type_id, rhs_id);
570        for instruction in [left_par, right_par] {
571            function.parameters.push(FunctionArgument {
572                instruction,
573                handle_id: 0,
574            });
575        }
576
577        let label_id = self.id_gen.next();
578        let mut block = Block::new(label_id);
579
580        let bool_type = return_type.with_scalar(crate::Scalar::BOOL);
581        let bool_type_id = self.get_numeric_type_id(bool_type);
582
583        let maybe_splat_const = |writer: &mut Self, const_id| match return_type {
584            NumericType::Scalar(_) => const_id,
585            NumericType::Vector { size, .. } => {
586                let constituent_ids = [const_id; crate::VectorSize::MAX];
587                writer.get_constant_composite(
588                    LookupType::Local(LocalType::Numeric(return_type)),
589                    &constituent_ids[..size as usize],
590                )
591            }
592            NumericType::Matrix { .. } => unreachable!(),
593        };
594
595        let const_zero_id = self.get_constant_scalar_with(0, scalar)?;
596        let composite_zero_id = maybe_splat_const(self, const_zero_id);
597        let rhs_eq_zero_id = self.id_gen.next();
598        block.body.push(Instruction::binary(
599            spirv::Op::IEqual,
600            bool_type_id,
601            rhs_eq_zero_id,
602            rhs_id,
603            composite_zero_id,
604        ));
605        let divisor_selector_id = match scalar.kind {
606            crate::ScalarKind::Sint => {
607                let (const_min_id, const_neg_one_id) = match scalar.width {
608                    4 => Ok((
609                        self.get_constant_scalar(crate::Literal::I32(i32::MIN)),
610                        self.get_constant_scalar(crate::Literal::I32(-1i32)),
611                    )),
612                    8 => Ok((
613                        self.get_constant_scalar(crate::Literal::I64(i64::MIN)),
614                        self.get_constant_scalar(crate::Literal::I64(-1i64)),
615                    )),
616                    _ => Err(Error::Validation("Unexpected scalar width")),
617                }?;
618                let composite_min_id = maybe_splat_const(self, const_min_id);
619                let composite_neg_one_id = maybe_splat_const(self, const_neg_one_id);
620
621                let lhs_eq_int_min_id = self.id_gen.next();
622                block.body.push(Instruction::binary(
623                    spirv::Op::IEqual,
624                    bool_type_id,
625                    lhs_eq_int_min_id,
626                    lhs_id,
627                    composite_min_id,
628                ));
629                let rhs_eq_neg_one_id = self.id_gen.next();
630                block.body.push(Instruction::binary(
631                    spirv::Op::IEqual,
632                    bool_type_id,
633                    rhs_eq_neg_one_id,
634                    rhs_id,
635                    composite_neg_one_id,
636                ));
637                let lhs_eq_int_min_and_rhs_eq_neg_one_id = self.id_gen.next();
638                block.body.push(Instruction::binary(
639                    spirv::Op::LogicalAnd,
640                    bool_type_id,
641                    lhs_eq_int_min_and_rhs_eq_neg_one_id,
642                    lhs_eq_int_min_id,
643                    rhs_eq_neg_one_id,
644                ));
645                let rhs_eq_zero_or_lhs_eq_int_min_and_rhs_eq_neg_one_id = self.id_gen.next();
646                block.body.push(Instruction::binary(
647                    spirv::Op::LogicalOr,
648                    bool_type_id,
649                    rhs_eq_zero_or_lhs_eq_int_min_and_rhs_eq_neg_one_id,
650                    rhs_eq_zero_id,
651                    lhs_eq_int_min_and_rhs_eq_neg_one_id,
652                ));
653                rhs_eq_zero_or_lhs_eq_int_min_and_rhs_eq_neg_one_id
654            }
655            crate::ScalarKind::Uint => rhs_eq_zero_id,
656            _ => unreachable!(),
657        };
658
659        let const_one_id = self.get_constant_scalar_with(1, scalar)?;
660        let composite_one_id = maybe_splat_const(self, const_one_id);
661        let divisor_id = self.id_gen.next();
662        block.body.push(Instruction::select(
663            right_type_id,
664            divisor_id,
665            divisor_selector_id,
666            composite_one_id,
667            rhs_id,
668        ));
669        let op = match (op, scalar.kind) {
670            (crate::BinaryOperator::Divide, crate::ScalarKind::Sint) => spirv::Op::SDiv,
671            (crate::BinaryOperator::Divide, crate::ScalarKind::Uint) => spirv::Op::UDiv,
672            (crate::BinaryOperator::Modulo, crate::ScalarKind::Sint) => spirv::Op::SRem,
673            (crate::BinaryOperator::Modulo, crate::ScalarKind::Uint) => spirv::Op::UMod,
674            _ => unreachable!(),
675        };
676        let return_id = self.id_gen.next();
677        block.body.push(Instruction::binary(
678            op,
679            return_type_id,
680            return_id,
681            lhs_id,
682            divisor_id,
683        ));
684
685        function.consume(block, Instruction::return_value(return_id));
686        function.to_words(&mut self.logical_layout.function_definitions);
687        Ok(())
688    }
689
690    fn write_function(
691        &mut self,
692        ir_function: &crate::Function,
693        info: &FunctionInfo,
694        ir_module: &crate::Module,
695        mut interface: Option<FunctionInterface>,
696        debug_info: &Option<DebugInfoInner>,
697    ) -> Result<Word, Error> {
698        self.write_wrapped_functions(ir_function, info, ir_module)?;
699
700        log::trace!("Generating code for {:?}", ir_function.name);
701        let mut function = Function::default();
702
703        let prelude_id = self.id_gen.next();
704        let mut prelude = Block::new(prelude_id);
705        let mut ep_context = EntryPointContext {
706            argument_ids: Vec::new(),
707            results: Vec::new(),
708        };
709
710        let mut local_invocation_id = None;
711
712        let mut parameter_type_ids = Vec::with_capacity(ir_function.arguments.len());
713        for argument in ir_function.arguments.iter() {
714            let class = spirv::StorageClass::Input;
715            let handle_ty = ir_module.types[argument.ty].inner.is_handle();
716            let argument_type_id = if handle_ty {
717                self.get_handle_pointer_type_id(argument.ty, spirv::StorageClass::UniformConstant)
718            } else {
719                self.get_handle_type_id(argument.ty)
720            };
721
722            if let Some(ref mut iface) = interface {
723                let id = if let Some(ref binding) = argument.binding {
724                    let name = argument.name.as_deref();
725
726                    let varying_id = self.write_varying(
727                        ir_module,
728                        iface.stage,
729                        class,
730                        name,
731                        argument.ty,
732                        binding,
733                    )?;
734                    iface.varying_ids.push(varying_id);
735                    let id = self.load_io_with_f16_polyfill(
736                        &mut prelude.body,
737                        varying_id,
738                        argument_type_id,
739                    );
740
741                    if binding == &crate::Binding::BuiltIn(crate::BuiltIn::LocalInvocationId) {
742                        local_invocation_id = Some(id);
743                    }
744
745                    id
746                } else if let crate::TypeInner::Struct { ref members, .. } =
747                    ir_module.types[argument.ty].inner
748                {
749                    let struct_id = self.id_gen.next();
750                    let mut constituent_ids = Vec::with_capacity(members.len());
751                    for member in members {
752                        let type_id = self.get_handle_type_id(member.ty);
753                        let name = member.name.as_deref();
754                        let binding = member.binding.as_ref().unwrap();
755                        let varying_id = self.write_varying(
756                            ir_module,
757                            iface.stage,
758                            class,
759                            name,
760                            member.ty,
761                            binding,
762                        )?;
763                        iface.varying_ids.push(varying_id);
764                        let id =
765                            self.load_io_with_f16_polyfill(&mut prelude.body, varying_id, type_id);
766                        constituent_ids.push(id);
767
768                        if binding == &crate::Binding::BuiltIn(crate::BuiltIn::LocalInvocationId) {
769                            local_invocation_id = Some(id);
770                        }
771                    }
772                    prelude.body.push(Instruction::composite_construct(
773                        argument_type_id,
774                        struct_id,
775                        &constituent_ids,
776                    ));
777                    struct_id
778                } else {
779                    unreachable!("Missing argument binding on an entry point");
780                };
781                ep_context.argument_ids.push(id);
782            } else {
783                let argument_id = self.id_gen.next();
784                let instruction = Instruction::function_parameter(argument_type_id, argument_id);
785                if self.flags.contains(WriterFlags::DEBUG) {
786                    if let Some(ref name) = argument.name {
787                        self.debugs.push(Instruction::name(argument_id, name));
788                    }
789                }
790                function.parameters.push(FunctionArgument {
791                    instruction,
792                    handle_id: if handle_ty {
793                        let id = self.id_gen.next();
794                        prelude.body.push(Instruction::load(
795                            self.get_handle_type_id(argument.ty),
796                            id,
797                            argument_id,
798                            None,
799                        ));
800                        id
801                    } else {
802                        0
803                    },
804                });
805                parameter_type_ids.push(argument_type_id);
806            };
807        }
808
809        let return_type_id = match ir_function.result {
810            Some(ref result) => {
811                if let Some(ref mut iface) = interface {
812                    let mut has_point_size = false;
813                    let class = spirv::StorageClass::Output;
814                    if let Some(ref binding) = result.binding {
815                        has_point_size |=
816                            *binding == crate::Binding::BuiltIn(crate::BuiltIn::PointSize);
817                        let type_id = self.get_handle_type_id(result.ty);
818                        let varying_id = self.write_varying(
819                            ir_module,
820                            iface.stage,
821                            class,
822                            None,
823                            result.ty,
824                            binding,
825                        )?;
826                        iface.varying_ids.push(varying_id);
827                        ep_context.results.push(ResultMember {
828                            id: varying_id,
829                            type_id,
830                            built_in: binding.to_built_in(),
831                        });
832                    } else if let crate::TypeInner::Struct { ref members, .. } =
833                        ir_module.types[result.ty].inner
834                    {
835                        for member in members {
836                            let type_id = self.get_handle_type_id(member.ty);
837                            let name = member.name.as_deref();
838                            let binding = member.binding.as_ref().unwrap();
839                            has_point_size |=
840                                *binding == crate::Binding::BuiltIn(crate::BuiltIn::PointSize);
841                            let varying_id = self.write_varying(
842                                ir_module,
843                                iface.stage,
844                                class,
845                                name,
846                                member.ty,
847                                binding,
848                            )?;
849                            iface.varying_ids.push(varying_id);
850                            ep_context.results.push(ResultMember {
851                                id: varying_id,
852                                type_id,
853                                built_in: binding.to_built_in(),
854                            });
855                        }
856                    } else {
857                        unreachable!("Missing result binding on an entry point");
858                    }
859
860                    if self.flags.contains(WriterFlags::FORCE_POINT_SIZE)
861                        && iface.stage == crate::ShaderStage::Vertex
862                        && !has_point_size
863                    {
864                        // add point size artificially
865                        let varying_id = self.id_gen.next();
866                        let pointer_type_id = self.get_f32_pointer_type_id(class);
867                        Instruction::variable(pointer_type_id, varying_id, class, None)
868                            .to_words(&mut self.logical_layout.declarations);
869                        self.decorate(
870                            varying_id,
871                            spirv::Decoration::BuiltIn,
872                            &[spirv::BuiltIn::PointSize as u32],
873                        );
874                        iface.varying_ids.push(varying_id);
875
876                        let default_value_id = self.get_constant_scalar(crate::Literal::F32(1.0));
877                        prelude
878                            .body
879                            .push(Instruction::store(varying_id, default_value_id, None));
880                    }
881                    self.void_type
882                } else {
883                    self.get_handle_type_id(result.ty)
884                }
885            }
886            None => self.void_type,
887        };
888
889        let lookup_function_type = LookupFunctionType {
890            parameter_type_ids,
891            return_type_id,
892        };
893
894        let function_id = self.id_gen.next();
895        if self.flags.contains(WriterFlags::DEBUG) {
896            if let Some(ref name) = ir_function.name {
897                self.debugs.push(Instruction::name(function_id, name));
898            }
899        }
900
901        let function_type = self.get_function_type(lookup_function_type);
902        function.signature = Some(Instruction::function(
903            return_type_id,
904            function_id,
905            spirv::FunctionControl::empty(),
906            function_type,
907        ));
908
909        if interface.is_some() {
910            function.entry_point_context = Some(ep_context);
911        }
912
913        // fill up the `GlobalVariable::access_id`
914        for gv in self.global_variables.iter_mut() {
915            gv.reset_for_function();
916        }
917        for (handle, var) in ir_module.global_variables.iter() {
918            if info[handle].is_empty() {
919                continue;
920            }
921
922            let mut gv = self.global_variables[handle].clone();
923            if let Some(ref mut iface) = interface {
924                // Have to include global variables in the interface
925                if self.physical_layout.version >= 0x10400 {
926                    iface.varying_ids.push(gv.var_id);
927                }
928            }
929
930            // Handle globals are pre-emitted and should be loaded automatically.
931            //
932            // Any that are binding arrays we skip as we cannot load the array, we must load the result after indexing.
933            match ir_module.types[var.ty].inner {
934                crate::TypeInner::BindingArray { .. } => {
935                    gv.access_id = gv.var_id;
936                }
937                _ => {
938                    if var.space == crate::AddressSpace::Handle {
939                        let var_type_id = self.get_handle_type_id(var.ty);
940                        let id = self.id_gen.next();
941                        prelude
942                            .body
943                            .push(Instruction::load(var_type_id, id, gv.var_id, None));
944                        gv.access_id = gv.var_id;
945                        gv.handle_id = id;
946                    } else if global_needs_wrapper(ir_module, var) {
947                        let class = map_storage_class(var.space);
948                        let pointer_type_id = self.get_handle_pointer_type_id(var.ty, class);
949                        let index_id = self.get_index_constant(0);
950                        let id = self.id_gen.next();
951                        prelude.body.push(Instruction::access_chain(
952                            pointer_type_id,
953                            id,
954                            gv.var_id,
955                            &[index_id],
956                        ));
957                        gv.access_id = id;
958                    } else {
959                        // by default, the variable ID is accessed as is
960                        gv.access_id = gv.var_id;
961                    };
962                }
963            }
964
965            // work around borrow checking in the presence of `self.xxx()` calls
966            self.global_variables[handle] = gv;
967        }
968
969        // Create a `BlockContext` for generating SPIR-V for the function's
970        // body.
971        let mut context = BlockContext {
972            ir_module,
973            ir_function,
974            fun_info: info,
975            function: &mut function,
976            // Re-use the cached expression table from prior functions.
977            cached: core::mem::take(&mut self.saved_cached),
978
979            // Steal the Writer's temp list for a bit.
980            temp_list: core::mem::take(&mut self.temp_list),
981            force_loop_bounding: self.force_loop_bounding,
982            writer: self,
983            expression_constness: super::ExpressionConstnessTracker::from_arena(
984                &ir_function.expressions,
985            ),
986        };
987
988        // fill up the pre-emitted and const expressions
989        context.cached.reset(ir_function.expressions.len());
990        for (handle, expr) in ir_function.expressions.iter() {
991            if (expr.needs_pre_emit() && !matches!(*expr, crate::Expression::LocalVariable(_)))
992                || context.expression_constness.is_const(handle)
993            {
994                context.cache_expression_value(handle, &mut prelude)?;
995            }
996        }
997
998        for (handle, variable) in ir_function.local_variables.iter() {
999            let id = context.gen_id();
1000
1001            if context.writer.flags.contains(WriterFlags::DEBUG) {
1002                if let Some(ref name) = variable.name {
1003                    context.writer.debugs.push(Instruction::name(id, name));
1004                }
1005            }
1006
1007            let init_word = variable.init.map(|constant| context.cached[constant]);
1008            let pointer_type_id = context
1009                .writer
1010                .get_handle_pointer_type_id(variable.ty, spirv::StorageClass::Function);
1011            let instruction = Instruction::variable(
1012                pointer_type_id,
1013                id,
1014                spirv::StorageClass::Function,
1015                init_word.or_else(|| match ir_module.types[variable.ty].inner {
1016                    crate::TypeInner::RayQuery { .. } => None,
1017                    _ => {
1018                        let type_id = context.get_handle_type_id(variable.ty);
1019                        Some(context.writer.write_constant_null(type_id))
1020                    }
1021                }),
1022            );
1023            context
1024                .function
1025                .variables
1026                .insert(handle, LocalVariable { id, instruction });
1027        }
1028
1029        for (handle, expr) in ir_function.expressions.iter() {
1030            match *expr {
1031                crate::Expression::LocalVariable(_) => {
1032                    // Cache the `OpVariable` instruction we generated above as
1033                    // the value of this expression.
1034                    context.cache_expression_value(handle, &mut prelude)?;
1035                }
1036                crate::Expression::Access { base, .. }
1037                | crate::Expression::AccessIndex { base, .. } => {
1038                    // Count references to `base` by `Access` and `AccessIndex`
1039                    // instructions. See `access_uses` for details.
1040                    *context.function.access_uses.entry(base).or_insert(0) += 1;
1041                }
1042                _ => {}
1043            }
1044        }
1045
1046        let next_id = context.gen_id();
1047
1048        context
1049            .function
1050            .consume(prelude, Instruction::branch(next_id));
1051
1052        let workgroup_vars_init_exit_block_id =
1053            match (context.writer.zero_initialize_workgroup_memory, interface) {
1054                (
1055                    super::ZeroInitializeWorkgroupMemoryMode::Polyfill,
1056                    Some(
1057                        ref mut interface @ FunctionInterface {
1058                            stage: crate::ShaderStage::Compute,
1059                            ..
1060                        },
1061                    ),
1062                ) => context.writer.generate_workgroup_vars_init_block(
1063                    next_id,
1064                    ir_module,
1065                    info,
1066                    local_invocation_id,
1067                    interface,
1068                    context.function,
1069                ),
1070                _ => None,
1071            };
1072
1073        let main_id = if let Some(exit_id) = workgroup_vars_init_exit_block_id {
1074            exit_id
1075        } else {
1076            next_id
1077        };
1078
1079        context.write_function_body(main_id, debug_info.as_ref())?;
1080
1081        // Consume the `BlockContext`, ending its borrows and letting the
1082        // `Writer` steal back its cached expression table and temp_list.
1083        let BlockContext {
1084            cached, temp_list, ..
1085        } = context;
1086        self.saved_cached = cached;
1087        self.temp_list = temp_list;
1088
1089        function.to_words(&mut self.logical_layout.function_definitions);
1090
1091        Ok(function_id)
1092    }
1093
1094    fn write_execution_mode(
1095        &mut self,
1096        function_id: Word,
1097        mode: spirv::ExecutionMode,
1098    ) -> Result<(), Error> {
1099        //self.check(mode.required_capabilities())?;
1100        Instruction::execution_mode(function_id, mode, &[])
1101            .to_words(&mut self.logical_layout.execution_modes);
1102        Ok(())
1103    }
1104
1105    // TODO Move to instructions module
1106    fn write_entry_point(
1107        &mut self,
1108        entry_point: &crate::EntryPoint,
1109        info: &FunctionInfo,
1110        ir_module: &crate::Module,
1111        debug_info: &Option<DebugInfoInner>,
1112    ) -> Result<Instruction, Error> {
1113        let mut interface_ids = Vec::new();
1114        let function_id = self.write_function(
1115            &entry_point.function,
1116            info,
1117            ir_module,
1118            Some(FunctionInterface {
1119                varying_ids: &mut interface_ids,
1120                stage: entry_point.stage,
1121            }),
1122            debug_info,
1123        )?;
1124
1125        let exec_model = match entry_point.stage {
1126            crate::ShaderStage::Vertex => spirv::ExecutionModel::Vertex,
1127            crate::ShaderStage::Fragment => {
1128                self.write_execution_mode(function_id, spirv::ExecutionMode::OriginUpperLeft)?;
1129                match entry_point.early_depth_test {
1130                    Some(crate::EarlyDepthTest::Force) => {
1131                        self.write_execution_mode(
1132                            function_id,
1133                            spirv::ExecutionMode::EarlyFragmentTests,
1134                        )?;
1135                    }
1136                    Some(crate::EarlyDepthTest::Allow { conservative }) => {
1137                        // TODO: Consider emitting EarlyAndLateFragmentTestsAMD here, if available.
1138                        // https://github.khronos.org/SPIRV-Registry/extensions/AMD/SPV_AMD_shader_early_and_late_fragment_tests.html
1139                        // This permits early depth tests even if the shader writes to a storage
1140                        // binding
1141                        match conservative {
1142                            crate::ConservativeDepth::GreaterEqual => self.write_execution_mode(
1143                                function_id,
1144                                spirv::ExecutionMode::DepthGreater,
1145                            )?,
1146                            crate::ConservativeDepth::LessEqual => self.write_execution_mode(
1147                                function_id,
1148                                spirv::ExecutionMode::DepthLess,
1149                            )?,
1150                            crate::ConservativeDepth::Unchanged => self.write_execution_mode(
1151                                function_id,
1152                                spirv::ExecutionMode::DepthUnchanged,
1153                            )?,
1154                        }
1155                    }
1156                    None => {}
1157                }
1158                if let Some(ref result) = entry_point.function.result {
1159                    if contains_builtin(
1160                        result.binding.as_ref(),
1161                        result.ty,
1162                        &ir_module.types,
1163                        crate::BuiltIn::FragDepth,
1164                    ) {
1165                        self.write_execution_mode(
1166                            function_id,
1167                            spirv::ExecutionMode::DepthReplacing,
1168                        )?;
1169                    }
1170                }
1171                spirv::ExecutionModel::Fragment
1172            }
1173            crate::ShaderStage::Compute => {
1174                let execution_mode = spirv::ExecutionMode::LocalSize;
1175                //self.check(execution_mode.required_capabilities())?;
1176                Instruction::execution_mode(
1177                    function_id,
1178                    execution_mode,
1179                    &entry_point.workgroup_size,
1180                )
1181                .to_words(&mut self.logical_layout.execution_modes);
1182                spirv::ExecutionModel::GLCompute
1183            }
1184            crate::ShaderStage::Task | crate::ShaderStage::Mesh => unreachable!(),
1185        };
1186        //self.check(exec_model.required_capabilities())?;
1187
1188        Ok(Instruction::entry_point(
1189            exec_model,
1190            function_id,
1191            &entry_point.name,
1192            interface_ids.as_slice(),
1193        ))
1194    }
1195
1196    fn make_scalar(&mut self, id: Word, scalar: crate::Scalar) -> Instruction {
1197        use crate::ScalarKind as Sk;
1198
1199        let bits = (scalar.width * BITS_PER_BYTE) as u32;
1200        match scalar.kind {
1201            Sk::Sint | Sk::Uint => {
1202                let signedness = if scalar.kind == Sk::Sint {
1203                    super::instructions::Signedness::Signed
1204                } else {
1205                    super::instructions::Signedness::Unsigned
1206                };
1207                let cap = match bits {
1208                    8 => Some(spirv::Capability::Int8),
1209                    16 => Some(spirv::Capability::Int16),
1210                    64 => Some(spirv::Capability::Int64),
1211                    _ => None,
1212                };
1213                if let Some(cap) = cap {
1214                    self.capabilities_used.insert(cap);
1215                }
1216                Instruction::type_int(id, bits, signedness)
1217            }
1218            Sk::Float => {
1219                if bits == 64 {
1220                    self.capabilities_used.insert(spirv::Capability::Float64);
1221                }
1222                if bits == 16 {
1223                    self.capabilities_used.insert(spirv::Capability::Float16);
1224                    self.capabilities_used
1225                        .insert(spirv::Capability::StorageBuffer16BitAccess);
1226                    self.capabilities_used
1227                        .insert(spirv::Capability::UniformAndStorageBuffer16BitAccess);
1228                    if self.use_storage_input_output_16 {
1229                        self.capabilities_used
1230                            .insert(spirv::Capability::StorageInputOutput16);
1231                    }
1232                }
1233                Instruction::type_float(id, bits)
1234            }
1235            Sk::Bool => Instruction::type_bool(id),
1236            Sk::AbstractInt | Sk::AbstractFloat => {
1237                unreachable!("abstract types should never reach the backend");
1238            }
1239        }
1240    }
1241
1242    fn request_type_capabilities(&mut self, inner: &crate::TypeInner) -> Result<(), Error> {
1243        match *inner {
1244            crate::TypeInner::Image {
1245                dim,
1246                arrayed,
1247                class,
1248            } => {
1249                let sampled = match class {
1250                    crate::ImageClass::Sampled { .. } => true,
1251                    crate::ImageClass::Depth { .. } => true,
1252                    crate::ImageClass::Storage { format, .. } => {
1253                        self.request_image_format_capabilities(format.into())?;
1254                        false
1255                    }
1256                    crate::ImageClass::External => unimplemented!(),
1257                };
1258
1259                match dim {
1260                    crate::ImageDimension::D1 => {
1261                        if sampled {
1262                            self.require_any("sampled 1D images", &[spirv::Capability::Sampled1D])?;
1263                        } else {
1264                            self.require_any("1D storage images", &[spirv::Capability::Image1D])?;
1265                        }
1266                    }
1267                    crate::ImageDimension::Cube if arrayed => {
1268                        if sampled {
1269                            self.require_any(
1270                                "sampled cube array images",
1271                                &[spirv::Capability::SampledCubeArray],
1272                            )?;
1273                        } else {
1274                            self.require_any(
1275                                "cube array storage images",
1276                                &[spirv::Capability::ImageCubeArray],
1277                            )?;
1278                        }
1279                    }
1280                    _ => {}
1281                }
1282            }
1283            crate::TypeInner::AccelerationStructure { .. } => {
1284                self.require_any("Acceleration Structure", &[spirv::Capability::RayQueryKHR])?;
1285            }
1286            crate::TypeInner::RayQuery { .. } => {
1287                self.require_any("Ray Query", &[spirv::Capability::RayQueryKHR])?;
1288            }
1289            crate::TypeInner::Atomic(crate::Scalar { width: 8, kind: _ }) => {
1290                self.require_any("64 bit integer atomics", &[spirv::Capability::Int64Atomics])?;
1291            }
1292            crate::TypeInner::Atomic(crate::Scalar {
1293                width: 4,
1294                kind: crate::ScalarKind::Float,
1295            }) => {
1296                self.require_any(
1297                    "32 bit floating-point atomics",
1298                    &[spirv::Capability::AtomicFloat32AddEXT],
1299                )?;
1300                self.use_extension("SPV_EXT_shader_atomic_float_add");
1301            }
1302            // 16 bit floating-point support requires Float16 capability
1303            crate::TypeInner::Matrix {
1304                scalar: crate::Scalar::F16,
1305                ..
1306            }
1307            | crate::TypeInner::Vector {
1308                scalar: crate::Scalar::F16,
1309                ..
1310            }
1311            | crate::TypeInner::Scalar(crate::Scalar::F16) => {
1312                self.require_any("16 bit floating-point", &[spirv::Capability::Float16])?;
1313                self.use_extension("SPV_KHR_16bit_storage");
1314            }
1315            _ => {}
1316        }
1317        Ok(())
1318    }
1319
1320    fn write_numeric_type_declaration_local(&mut self, id: Word, numeric: NumericType) {
1321        let instruction = match numeric {
1322            NumericType::Scalar(scalar) => self.make_scalar(id, scalar),
1323            NumericType::Vector { size, scalar } => {
1324                let scalar_id = self.get_numeric_type_id(NumericType::Scalar(scalar));
1325                Instruction::type_vector(id, scalar_id, size)
1326            }
1327            NumericType::Matrix {
1328                columns,
1329                rows,
1330                scalar,
1331            } => {
1332                let column_id =
1333                    self.get_numeric_type_id(NumericType::Vector { size: rows, scalar });
1334                Instruction::type_matrix(id, column_id, columns)
1335            }
1336        };
1337
1338        instruction.to_words(&mut self.logical_layout.declarations);
1339    }
1340
1341    fn write_type_declaration_local(&mut self, id: Word, local_ty: LocalType) {
1342        let instruction = match local_ty {
1343            LocalType::Numeric(numeric) => {
1344                self.write_numeric_type_declaration_local(id, numeric);
1345                return;
1346            }
1347            LocalType::Pointer { base, class } => Instruction::type_pointer(id, class, base),
1348            LocalType::Image(image) => {
1349                let local_type = LocalType::Numeric(NumericType::Scalar(image.sampled_type));
1350                let type_id = self.get_localtype_id(local_type);
1351                Instruction::type_image(id, type_id, image.dim, image.flags, image.image_format)
1352            }
1353            LocalType::Sampler => Instruction::type_sampler(id),
1354            LocalType::SampledImage { image_type_id } => {
1355                Instruction::type_sampled_image(id, image_type_id)
1356            }
1357            LocalType::BindingArray { base, size } => {
1358                let inner_ty = self.get_handle_type_id(base);
1359                let scalar_id = self.get_constant_scalar(crate::Literal::U32(size));
1360                Instruction::type_array(id, inner_ty, scalar_id)
1361            }
1362            LocalType::AccelerationStructure => Instruction::type_acceleration_structure(id),
1363            LocalType::RayQuery => Instruction::type_ray_query(id),
1364        };
1365
1366        instruction.to_words(&mut self.logical_layout.declarations);
1367    }
1368
1369    fn write_type_declaration_arena(
1370        &mut self,
1371        module: &crate::Module,
1372        handle: Handle<crate::Type>,
1373    ) -> Result<Word, Error> {
1374        let ty = &module.types[handle];
1375        // If it's a type that needs SPIR-V capabilities, request them now.
1376        // This needs to happen regardless of the LocalType lookup succeeding,
1377        // because some types which map to the same LocalType have different
1378        // capability requirements. See https://github.com/gfx-rs/wgpu/issues/5569
1379        self.request_type_capabilities(&ty.inner)?;
1380        let id = if let Some(local) = self.localtype_from_inner(&ty.inner) {
1381            // This type can be represented as a `LocalType`, so check if we've
1382            // already written an instruction for it. If not, do so now, with
1383            // `write_type_declaration_local`.
1384            match self.lookup_type.entry(LookupType::Local(local)) {
1385                // We already have an id for this `LocalType`.
1386                Entry::Occupied(e) => *e.get(),
1387
1388                // It's a type we haven't seen before.
1389                Entry::Vacant(e) => {
1390                    let id = self.id_gen.next();
1391                    e.insert(id);
1392
1393                    self.write_type_declaration_local(id, local);
1394
1395                    id
1396                }
1397            }
1398        } else {
1399            use spirv::Decoration;
1400
1401            let id = self.id_gen.next();
1402            let instruction = match ty.inner {
1403                crate::TypeInner::Array { base, size, stride } => {
1404                    self.decorate(id, Decoration::ArrayStride, &[stride]);
1405
1406                    let type_id = self.get_handle_type_id(base);
1407                    match size.resolve(module.to_ctx())? {
1408                        crate::proc::IndexableLength::Known(length) => {
1409                            let length_id = self.get_index_constant(length);
1410                            Instruction::type_array(id, type_id, length_id)
1411                        }
1412                        crate::proc::IndexableLength::Dynamic => {
1413                            Instruction::type_runtime_array(id, type_id)
1414                        }
1415                    }
1416                }
1417                crate::TypeInner::BindingArray { base, size } => {
1418                    let type_id = self.get_handle_type_id(base);
1419                    match size.resolve(module.to_ctx())? {
1420                        crate::proc::IndexableLength::Known(length) => {
1421                            let length_id = self.get_index_constant(length);
1422                            Instruction::type_array(id, type_id, length_id)
1423                        }
1424                        crate::proc::IndexableLength::Dynamic => {
1425                            Instruction::type_runtime_array(id, type_id)
1426                        }
1427                    }
1428                }
1429                crate::TypeInner::Struct {
1430                    ref members,
1431                    span: _,
1432                } => {
1433                    let mut has_runtime_array = false;
1434                    let mut member_ids = Vec::with_capacity(members.len());
1435                    for (index, member) in members.iter().enumerate() {
1436                        let member_ty = &module.types[member.ty];
1437                        match member_ty.inner {
1438                            crate::TypeInner::Array {
1439                                base: _,
1440                                size: crate::ArraySize::Dynamic,
1441                                stride: _,
1442                            } => {
1443                                has_runtime_array = true;
1444                            }
1445                            _ => (),
1446                        }
1447                        self.decorate_struct_member(id, index, member, &module.types)?;
1448                        let member_id = self.get_handle_type_id(member.ty);
1449                        member_ids.push(member_id);
1450                    }
1451                    if has_runtime_array {
1452                        self.decorate(id, Decoration::Block, &[]);
1453                    }
1454                    Instruction::type_struct(id, member_ids.as_slice())
1455                }
1456
1457                // These all have TypeLocal representations, so they should have been
1458                // handled by `write_type_declaration_local` above.
1459                crate::TypeInner::Scalar(_)
1460                | crate::TypeInner::Atomic(_)
1461                | crate::TypeInner::Vector { .. }
1462                | crate::TypeInner::Matrix { .. }
1463                | crate::TypeInner::Pointer { .. }
1464                | crate::TypeInner::ValuePointer { .. }
1465                | crate::TypeInner::Image { .. }
1466                | crate::TypeInner::Sampler { .. }
1467                | crate::TypeInner::AccelerationStructure { .. }
1468                | crate::TypeInner::RayQuery { .. } => unreachable!(),
1469            };
1470
1471            instruction.to_words(&mut self.logical_layout.declarations);
1472            id
1473        };
1474
1475        // Add this handle as a new alias for that type.
1476        self.lookup_type.insert(LookupType::Handle(handle), id);
1477
1478        if self.flags.contains(WriterFlags::DEBUG) {
1479            if let Some(ref name) = ty.name {
1480                self.debugs.push(Instruction::name(id, name));
1481            }
1482        }
1483
1484        Ok(id)
1485    }
1486
1487    fn request_image_format_capabilities(
1488        &mut self,
1489        format: spirv::ImageFormat,
1490    ) -> Result<(), Error> {
1491        use spirv::ImageFormat as If;
1492        match format {
1493            If::Rg32f
1494            | If::Rg16f
1495            | If::R11fG11fB10f
1496            | If::R16f
1497            | If::Rgba16
1498            | If::Rgb10A2
1499            | If::Rg16
1500            | If::Rg8
1501            | If::R16
1502            | If::R8
1503            | If::Rgba16Snorm
1504            | If::Rg16Snorm
1505            | If::Rg8Snorm
1506            | If::R16Snorm
1507            | If::R8Snorm
1508            | If::Rg32i
1509            | If::Rg16i
1510            | If::Rg8i
1511            | If::R16i
1512            | If::R8i
1513            | If::Rgb10a2ui
1514            | If::Rg32ui
1515            | If::Rg16ui
1516            | If::Rg8ui
1517            | If::R16ui
1518            | If::R8ui => self.require_any(
1519                "storage image format",
1520                &[spirv::Capability::StorageImageExtendedFormats],
1521            ),
1522            If::R64ui | If::R64i => {
1523                self.use_extension("SPV_EXT_shader_image_int64");
1524                self.require_any(
1525                    "64-bit integer storage image format",
1526                    &[spirv::Capability::Int64ImageEXT],
1527                )
1528            }
1529            If::Unknown
1530            | If::Rgba32f
1531            | If::Rgba16f
1532            | If::R32f
1533            | If::Rgba8
1534            | If::Rgba8Snorm
1535            | If::Rgba32i
1536            | If::Rgba16i
1537            | If::Rgba8i
1538            | If::R32i
1539            | If::Rgba32ui
1540            | If::Rgba16ui
1541            | If::Rgba8ui
1542            | If::R32ui => Ok(()),
1543        }
1544    }
1545
1546    pub(super) fn get_index_constant(&mut self, index: Word) -> Word {
1547        self.get_constant_scalar(crate::Literal::U32(index))
1548    }
1549
1550    pub(super) fn get_constant_scalar_with(
1551        &mut self,
1552        value: u8,
1553        scalar: crate::Scalar,
1554    ) -> Result<Word, Error> {
1555        Ok(
1556            self.get_constant_scalar(crate::Literal::new(value, scalar).ok_or(
1557                Error::Validation("Unexpected kind and/or width for Literal"),
1558            )?),
1559        )
1560    }
1561
1562    pub(super) fn get_constant_scalar(&mut self, value: crate::Literal) -> Word {
1563        let scalar = CachedConstant::Literal(value.into());
1564        if let Some(&id) = self.cached_constants.get(&scalar) {
1565            return id;
1566        }
1567        let id = self.id_gen.next();
1568        self.write_constant_scalar(id, &value, None);
1569        self.cached_constants.insert(scalar, id);
1570        id
1571    }
1572
1573    fn write_constant_scalar(
1574        &mut self,
1575        id: Word,
1576        value: &crate::Literal,
1577        debug_name: Option<&String>,
1578    ) {
1579        if self.flags.contains(WriterFlags::DEBUG) {
1580            if let Some(name) = debug_name {
1581                self.debugs.push(Instruction::name(id, name));
1582            }
1583        }
1584        let type_id = self.get_numeric_type_id(NumericType::Scalar(value.scalar()));
1585        let instruction = match *value {
1586            crate::Literal::F64(value) => {
1587                let bits = value.to_bits();
1588                Instruction::constant_64bit(type_id, id, bits as u32, (bits >> 32) as u32)
1589            }
1590            crate::Literal::F32(value) => Instruction::constant_32bit(type_id, id, value.to_bits()),
1591            crate::Literal::F16(value) => {
1592                let low = value.to_bits();
1593                Instruction::constant_16bit(type_id, id, low as u32)
1594            }
1595            crate::Literal::U32(value) => Instruction::constant_32bit(type_id, id, value),
1596            crate::Literal::I32(value) => Instruction::constant_32bit(type_id, id, value as u32),
1597            crate::Literal::U64(value) => {
1598                Instruction::constant_64bit(type_id, id, value as u32, (value >> 32) as u32)
1599            }
1600            crate::Literal::I64(value) => {
1601                Instruction::constant_64bit(type_id, id, value as u32, (value >> 32) as u32)
1602            }
1603            crate::Literal::Bool(true) => Instruction::constant_true(type_id, id),
1604            crate::Literal::Bool(false) => Instruction::constant_false(type_id, id),
1605            crate::Literal::AbstractInt(_) | crate::Literal::AbstractFloat(_) => {
1606                unreachable!("Abstract types should not appear in IR presented to backends");
1607            }
1608        };
1609
1610        instruction.to_words(&mut self.logical_layout.declarations);
1611    }
1612
1613    pub(super) fn get_constant_composite(
1614        &mut self,
1615        ty: LookupType,
1616        constituent_ids: &[Word],
1617    ) -> Word {
1618        let composite = CachedConstant::Composite {
1619            ty,
1620            constituent_ids: constituent_ids.to_vec(),
1621        };
1622        if let Some(&id) = self.cached_constants.get(&composite) {
1623            return id;
1624        }
1625        let id = self.id_gen.next();
1626        self.write_constant_composite(id, ty, constituent_ids, None);
1627        self.cached_constants.insert(composite, id);
1628        id
1629    }
1630
1631    fn write_constant_composite(
1632        &mut self,
1633        id: Word,
1634        ty: LookupType,
1635        constituent_ids: &[Word],
1636        debug_name: Option<&String>,
1637    ) {
1638        if self.flags.contains(WriterFlags::DEBUG) {
1639            if let Some(name) = debug_name {
1640                self.debugs.push(Instruction::name(id, name));
1641            }
1642        }
1643        let type_id = self.get_type_id(ty);
1644        Instruction::constant_composite(type_id, id, constituent_ids)
1645            .to_words(&mut self.logical_layout.declarations);
1646    }
1647
1648    pub(super) fn get_constant_null(&mut self, type_id: Word) -> Word {
1649        let null = CachedConstant::ZeroValue(type_id);
1650        if let Some(&id) = self.cached_constants.get(&null) {
1651            return id;
1652        }
1653        let id = self.write_constant_null(type_id);
1654        self.cached_constants.insert(null, id);
1655        id
1656    }
1657
1658    pub(super) fn write_constant_null(&mut self, type_id: Word) -> Word {
1659        let null_id = self.id_gen.next();
1660        Instruction::constant_null(type_id, null_id)
1661            .to_words(&mut self.logical_layout.declarations);
1662        null_id
1663    }
1664
1665    fn write_constant_expr(
1666        &mut self,
1667        handle: Handle<crate::Expression>,
1668        ir_module: &crate::Module,
1669        mod_info: &ModuleInfo,
1670    ) -> Result<Word, Error> {
1671        let id = match ir_module.global_expressions[handle] {
1672            crate::Expression::Literal(literal) => self.get_constant_scalar(literal),
1673            crate::Expression::Constant(constant) => {
1674                let constant = &ir_module.constants[constant];
1675                self.constant_ids[constant.init]
1676            }
1677            crate::Expression::ZeroValue(ty) => {
1678                let type_id = self.get_handle_type_id(ty);
1679                self.get_constant_null(type_id)
1680            }
1681            crate::Expression::Compose { ty, ref components } => {
1682                let component_ids: Vec<_> = crate::proc::flatten_compose(
1683                    ty,
1684                    components,
1685                    &ir_module.global_expressions,
1686                    &ir_module.types,
1687                )
1688                .map(|component| self.constant_ids[component])
1689                .collect();
1690                self.get_constant_composite(LookupType::Handle(ty), component_ids.as_slice())
1691            }
1692            crate::Expression::Splat { size, value } => {
1693                let value_id = self.constant_ids[value];
1694                let component_ids = &[value_id; 4][..size as usize];
1695
1696                let ty = self.get_expression_lookup_type(&mod_info[handle]);
1697
1698                self.get_constant_composite(ty, component_ids)
1699            }
1700            _ => {
1701                return Err(Error::Override);
1702            }
1703        };
1704
1705        self.constant_ids[handle] = id;
1706
1707        Ok(id)
1708    }
1709
1710    pub(super) fn write_control_barrier(&mut self, flags: crate::Barrier, block: &mut Block) {
1711        let memory_scope = if flags.contains(crate::Barrier::STORAGE) {
1712            spirv::Scope::Device
1713        } else if flags.contains(crate::Barrier::SUB_GROUP) {
1714            spirv::Scope::Subgroup
1715        } else {
1716            spirv::Scope::Workgroup
1717        };
1718        let mut semantics = spirv::MemorySemantics::ACQUIRE_RELEASE;
1719        semantics.set(
1720            spirv::MemorySemantics::UNIFORM_MEMORY,
1721            flags.contains(crate::Barrier::STORAGE),
1722        );
1723        semantics.set(
1724            spirv::MemorySemantics::WORKGROUP_MEMORY,
1725            flags.contains(crate::Barrier::WORK_GROUP),
1726        );
1727        semantics.set(
1728            spirv::MemorySemantics::SUBGROUP_MEMORY,
1729            flags.contains(crate::Barrier::SUB_GROUP),
1730        );
1731        semantics.set(
1732            spirv::MemorySemantics::IMAGE_MEMORY,
1733            flags.contains(crate::Barrier::TEXTURE),
1734        );
1735        let exec_scope_id = if flags.contains(crate::Barrier::SUB_GROUP) {
1736            self.get_index_constant(spirv::Scope::Subgroup as u32)
1737        } else {
1738            self.get_index_constant(spirv::Scope::Workgroup as u32)
1739        };
1740        let mem_scope_id = self.get_index_constant(memory_scope as u32);
1741        let semantics_id = self.get_index_constant(semantics.bits());
1742        block.body.push(Instruction::control_barrier(
1743            exec_scope_id,
1744            mem_scope_id,
1745            semantics_id,
1746        ));
1747    }
1748
1749    pub(super) fn write_memory_barrier(&mut self, flags: crate::Barrier, block: &mut Block) {
1750        let mut semantics = spirv::MemorySemantics::ACQUIRE_RELEASE;
1751        semantics.set(
1752            spirv::MemorySemantics::UNIFORM_MEMORY,
1753            flags.contains(crate::Barrier::STORAGE),
1754        );
1755        semantics.set(
1756            spirv::MemorySemantics::WORKGROUP_MEMORY,
1757            flags.contains(crate::Barrier::WORK_GROUP),
1758        );
1759        semantics.set(
1760            spirv::MemorySemantics::SUBGROUP_MEMORY,
1761            flags.contains(crate::Barrier::SUB_GROUP),
1762        );
1763        semantics.set(
1764            spirv::MemorySemantics::IMAGE_MEMORY,
1765            flags.contains(crate::Barrier::TEXTURE),
1766        );
1767        let mem_scope_id = if flags.contains(crate::Barrier::STORAGE) {
1768            self.get_index_constant(spirv::Scope::Device as u32)
1769        } else if flags.contains(crate::Barrier::SUB_GROUP) {
1770            self.get_index_constant(spirv::Scope::Subgroup as u32)
1771        } else {
1772            self.get_index_constant(spirv::Scope::Workgroup as u32)
1773        };
1774        let semantics_id = self.get_index_constant(semantics.bits());
1775        block
1776            .body
1777            .push(Instruction::memory_barrier(mem_scope_id, semantics_id));
1778    }
1779
1780    fn generate_workgroup_vars_init_block(
1781        &mut self,
1782        entry_id: Word,
1783        ir_module: &crate::Module,
1784        info: &FunctionInfo,
1785        local_invocation_id: Option<Word>,
1786        interface: &mut FunctionInterface,
1787        function: &mut Function,
1788    ) -> Option<Word> {
1789        let body = ir_module
1790            .global_variables
1791            .iter()
1792            .filter(|&(handle, var)| {
1793                !info[handle].is_empty() && var.space == crate::AddressSpace::WorkGroup
1794            })
1795            .map(|(handle, var)| {
1796                // It's safe to use `var_id` here, not `access_id`, because only
1797                // variables in the `Uniform` and `StorageBuffer` address spaces
1798                // get wrapped, and we're initializing `WorkGroup` variables.
1799                let var_id = self.global_variables[handle].var_id;
1800                let var_type_id = self.get_handle_type_id(var.ty);
1801                let init_word = self.get_constant_null(var_type_id);
1802                Instruction::store(var_id, init_word, None)
1803            })
1804            .collect::<Vec<_>>();
1805
1806        if body.is_empty() {
1807            return None;
1808        }
1809
1810        let uint3_type_id = self.get_vec3u_type_id();
1811
1812        let mut pre_if_block = Block::new(entry_id);
1813
1814        let local_invocation_id = if let Some(local_invocation_id) = local_invocation_id {
1815            local_invocation_id
1816        } else {
1817            let varying_id = self.id_gen.next();
1818            let class = spirv::StorageClass::Input;
1819            let pointer_type_id = self.get_vec3u_pointer_type_id(class);
1820
1821            Instruction::variable(pointer_type_id, varying_id, class, None)
1822                .to_words(&mut self.logical_layout.declarations);
1823
1824            self.decorate(
1825                varying_id,
1826                spirv::Decoration::BuiltIn,
1827                &[spirv::BuiltIn::LocalInvocationId as u32],
1828            );
1829
1830            interface.varying_ids.push(varying_id);
1831            let id = self.id_gen.next();
1832            pre_if_block
1833                .body
1834                .push(Instruction::load(uint3_type_id, id, varying_id, None));
1835
1836            id
1837        };
1838
1839        let zero_id = self.get_constant_null(uint3_type_id);
1840        let bool3_type_id = self.get_vec3_bool_type_id();
1841
1842        let eq_id = self.id_gen.next();
1843        pre_if_block.body.push(Instruction::binary(
1844            spirv::Op::IEqual,
1845            bool3_type_id,
1846            eq_id,
1847            local_invocation_id,
1848            zero_id,
1849        ));
1850
1851        let condition_id = self.id_gen.next();
1852        let bool_type_id = self.get_bool_type_id();
1853        pre_if_block.body.push(Instruction::relational(
1854            spirv::Op::All,
1855            bool_type_id,
1856            condition_id,
1857            eq_id,
1858        ));
1859
1860        let merge_id = self.id_gen.next();
1861        pre_if_block.body.push(Instruction::selection_merge(
1862            merge_id,
1863            spirv::SelectionControl::NONE,
1864        ));
1865
1866        let accept_id = self.id_gen.next();
1867        function.consume(
1868            pre_if_block,
1869            Instruction::branch_conditional(condition_id, accept_id, merge_id),
1870        );
1871
1872        let accept_block = Block {
1873            label_id: accept_id,
1874            body,
1875        };
1876        function.consume(accept_block, Instruction::branch(merge_id));
1877
1878        let mut post_if_block = Block::new(merge_id);
1879
1880        self.write_control_barrier(crate::Barrier::WORK_GROUP, &mut post_if_block);
1881
1882        let next_id = self.id_gen.next();
1883        function.consume(post_if_block, Instruction::branch(next_id));
1884        Some(next_id)
1885    }
1886
1887    /// Generate an `OpVariable` for one value in an [`EntryPoint`]'s IO interface.
1888    ///
1889    /// The [`Binding`]s of the arguments and result of an [`EntryPoint`]'s
1890    /// [`Function`] describe a SPIR-V shader interface. In SPIR-V, the
1891    /// interface is represented by global variables in the `Input` and `Output`
1892    /// storage classes, with decorations indicating which builtin or location
1893    /// each variable corresponds to.
1894    ///
1895    /// This function emits a single global `OpVariable` for a single value from
1896    /// the interface, and adds appropriate decorations to indicate which
1897    /// builtin or location it represents, how it should be interpolated, and so
1898    /// on. The `class` argument gives the variable's SPIR-V storage class,
1899    /// which should be either [`Input`] or [`Output`].
1900    ///
1901    /// [`Binding`]: crate::Binding
1902    /// [`Function`]: crate::Function
1903    /// [`EntryPoint`]: crate::EntryPoint
1904    /// [`Input`]: spirv::StorageClass::Input
1905    /// [`Output`]: spirv::StorageClass::Output
1906    fn write_varying(
1907        &mut self,
1908        ir_module: &crate::Module,
1909        stage: crate::ShaderStage,
1910        class: spirv::StorageClass,
1911        debug_name: Option<&str>,
1912        ty: Handle<crate::Type>,
1913        binding: &crate::Binding,
1914    ) -> Result<Word, Error> {
1915        use crate::TypeInner;
1916
1917        let id = self.id_gen.next();
1918        let ty_inner = &ir_module.types[ty].inner;
1919        let needs_polyfill = self.needs_f16_polyfill(ty_inner);
1920
1921        let pointer_type_id = if needs_polyfill {
1922            let f32_value_local =
1923                super::f16_polyfill::F16IoPolyfill::create_polyfill_type(ty_inner)
1924                    .expect("needs_polyfill returned true but create_polyfill_type returned None");
1925
1926            let f32_type_id = self.get_localtype_id(f32_value_local);
1927            let ptr_id = self.get_pointer_type_id(f32_type_id, class);
1928            self.io_f16_polyfills.register_io_var(id, f32_type_id);
1929
1930            ptr_id
1931        } else {
1932            self.get_handle_pointer_type_id(ty, class)
1933        };
1934
1935        Instruction::variable(pointer_type_id, id, class, None)
1936            .to_words(&mut self.logical_layout.declarations);
1937
1938        if self
1939            .flags
1940            .contains(WriterFlags::DEBUG | WriterFlags::LABEL_VARYINGS)
1941        {
1942            if let Some(name) = debug_name {
1943                self.debugs.push(Instruction::name(id, name));
1944            }
1945        }
1946
1947        use spirv::{BuiltIn, Decoration};
1948
1949        match *binding {
1950            crate::Binding::Location {
1951                location,
1952                interpolation,
1953                sampling,
1954                blend_src,
1955            } => {
1956                self.decorate(id, Decoration::Location, &[location]);
1957
1958                let no_decorations =
1959                    // VUID-StandaloneSpirv-Flat-06202
1960                    // > The Flat, NoPerspective, Sample, and Centroid decorations
1961                    // > must not be used on variables with the Input storage class in a vertex shader
1962                    (class == spirv::StorageClass::Input && stage == crate::ShaderStage::Vertex) ||
1963                    // VUID-StandaloneSpirv-Flat-06201
1964                    // > The Flat, NoPerspective, Sample, and Centroid decorations
1965                    // > must not be used on variables with the Output storage class in a fragment shader
1966                    (class == spirv::StorageClass::Output && stage == crate::ShaderStage::Fragment);
1967
1968                if !no_decorations {
1969                    match interpolation {
1970                        // Perspective-correct interpolation is the default in SPIR-V.
1971                        None | Some(crate::Interpolation::Perspective) => (),
1972                        Some(crate::Interpolation::Flat) => {
1973                            self.decorate(id, Decoration::Flat, &[]);
1974                        }
1975                        Some(crate::Interpolation::Linear) => {
1976                            self.decorate(id, Decoration::NoPerspective, &[]);
1977                        }
1978                    }
1979                    match sampling {
1980                        // Center sampling is the default in SPIR-V.
1981                        None
1982                        | Some(
1983                            crate::Sampling::Center
1984                            | crate::Sampling::First
1985                            | crate::Sampling::Either,
1986                        ) => (),
1987                        Some(crate::Sampling::Centroid) => {
1988                            self.decorate(id, Decoration::Centroid, &[]);
1989                        }
1990                        Some(crate::Sampling::Sample) => {
1991                            self.require_any(
1992                                "per-sample interpolation",
1993                                &[spirv::Capability::SampleRateShading],
1994                            )?;
1995                            self.decorate(id, Decoration::Sample, &[]);
1996                        }
1997                    }
1998                }
1999                if let Some(blend_src) = blend_src {
2000                    self.decorate(id, Decoration::Index, &[blend_src]);
2001                }
2002            }
2003            crate::Binding::BuiltIn(built_in) => {
2004                use crate::BuiltIn as Bi;
2005                let built_in = match built_in {
2006                    Bi::Position { invariant } => {
2007                        if invariant {
2008                            self.decorate(id, Decoration::Invariant, &[]);
2009                        }
2010
2011                        if class == spirv::StorageClass::Output {
2012                            BuiltIn::Position
2013                        } else {
2014                            BuiltIn::FragCoord
2015                        }
2016                    }
2017                    Bi::ViewIndex => {
2018                        self.require_any("`view_index` built-in", &[spirv::Capability::MultiView])?;
2019                        BuiltIn::ViewIndex
2020                    }
2021                    // vertex
2022                    Bi::BaseInstance => BuiltIn::BaseInstance,
2023                    Bi::BaseVertex => BuiltIn::BaseVertex,
2024                    Bi::ClipDistance => {
2025                        self.require_any(
2026                            "`clip_distance` built-in",
2027                            &[spirv::Capability::ClipDistance],
2028                        )?;
2029                        BuiltIn::ClipDistance
2030                    }
2031                    Bi::CullDistance => {
2032                        self.require_any(
2033                            "`cull_distance` built-in",
2034                            &[spirv::Capability::CullDistance],
2035                        )?;
2036                        BuiltIn::CullDistance
2037                    }
2038                    Bi::InstanceIndex => BuiltIn::InstanceIndex,
2039                    Bi::PointSize => BuiltIn::PointSize,
2040                    Bi::VertexIndex => BuiltIn::VertexIndex,
2041                    Bi::DrawID => BuiltIn::DrawIndex,
2042                    // fragment
2043                    Bi::FragDepth => BuiltIn::FragDepth,
2044                    Bi::PointCoord => BuiltIn::PointCoord,
2045                    Bi::FrontFacing => BuiltIn::FrontFacing,
2046                    Bi::PrimitiveIndex => {
2047                        self.require_any(
2048                            "`primitive_index` built-in",
2049                            &[spirv::Capability::Geometry],
2050                        )?;
2051                        BuiltIn::PrimitiveId
2052                    }
2053                    Bi::SampleIndex => {
2054                        self.require_any(
2055                            "`sample_index` built-in",
2056                            &[spirv::Capability::SampleRateShading],
2057                        )?;
2058
2059                        BuiltIn::SampleId
2060                    }
2061                    Bi::SampleMask => BuiltIn::SampleMask,
2062                    // compute
2063                    Bi::GlobalInvocationId => BuiltIn::GlobalInvocationId,
2064                    Bi::LocalInvocationId => BuiltIn::LocalInvocationId,
2065                    Bi::LocalInvocationIndex => BuiltIn::LocalInvocationIndex,
2066                    Bi::WorkGroupId => BuiltIn::WorkgroupId,
2067                    Bi::WorkGroupSize => BuiltIn::WorkgroupSize,
2068                    Bi::NumWorkGroups => BuiltIn::NumWorkgroups,
2069                    // Subgroup
2070                    Bi::NumSubgroups => {
2071                        self.require_any(
2072                            "`num_subgroups` built-in",
2073                            &[spirv::Capability::GroupNonUniform],
2074                        )?;
2075                        BuiltIn::NumSubgroups
2076                    }
2077                    Bi::SubgroupId => {
2078                        self.require_any(
2079                            "`subgroup_id` built-in",
2080                            &[spirv::Capability::GroupNonUniform],
2081                        )?;
2082                        BuiltIn::SubgroupId
2083                    }
2084                    Bi::SubgroupSize => {
2085                        self.require_any(
2086                            "`subgroup_size` built-in",
2087                            &[
2088                                spirv::Capability::GroupNonUniform,
2089                                spirv::Capability::SubgroupBallotKHR,
2090                            ],
2091                        )?;
2092                        BuiltIn::SubgroupSize
2093                    }
2094                    Bi::SubgroupInvocationId => {
2095                        self.require_any(
2096                            "`subgroup_invocation_id` built-in",
2097                            &[
2098                                spirv::Capability::GroupNonUniform,
2099                                spirv::Capability::SubgroupBallotKHR,
2100                            ],
2101                        )?;
2102                        BuiltIn::SubgroupLocalInvocationId
2103                    }
2104                };
2105
2106                self.decorate(id, Decoration::BuiltIn, &[built_in as u32]);
2107
2108                use crate::ScalarKind as Sk;
2109
2110                // Per the Vulkan spec, `VUID-StandaloneSpirv-Flat-04744`:
2111                //
2112                // > Any variable with integer or double-precision floating-
2113                // > point type and with Input storage class in a fragment
2114                // > shader, must be decorated Flat
2115                if class == spirv::StorageClass::Input && stage == crate::ShaderStage::Fragment {
2116                    let is_flat = match ir_module.types[ty].inner {
2117                        TypeInner::Scalar(scalar) | TypeInner::Vector { scalar, .. } => match scalar
2118                            .kind
2119                        {
2120                            Sk::Uint | Sk::Sint | Sk::Bool => true,
2121                            Sk::Float => false,
2122                            Sk::AbstractInt | Sk::AbstractFloat => {
2123                                return Err(Error::Validation(
2124                                    "Abstract types should not appear in IR presented to backends",
2125                                ))
2126                            }
2127                        },
2128                        _ => false,
2129                    };
2130
2131                    if is_flat {
2132                        self.decorate(id, Decoration::Flat, &[]);
2133                    }
2134                }
2135            }
2136        }
2137
2138        Ok(id)
2139    }
2140
2141    /// Load an IO variable, converting from `f32` to `f16` if polyfill is active.
2142    /// Returns the id of the loaded value matching `target_type_id`.
2143    pub(super) fn load_io_with_f16_polyfill(
2144        &mut self,
2145        body: &mut Vec<Instruction>,
2146        varying_id: Word,
2147        target_type_id: Word,
2148    ) -> Word {
2149        let tmp = self.id_gen.next();
2150        if let Some(f32_ty) = self.io_f16_polyfills.get_f32_io_type(varying_id) {
2151            body.push(Instruction::load(f32_ty, tmp, varying_id, None));
2152            let converted = self.id_gen.next();
2153            super::f16_polyfill::F16IoPolyfill::emit_f32_to_f16_conversion(
2154                tmp,
2155                target_type_id,
2156                converted,
2157                body,
2158            );
2159            converted
2160        } else {
2161            body.push(Instruction::load(target_type_id, tmp, varying_id, None));
2162            tmp
2163        }
2164    }
2165
2166    /// Store an IO variable, converting from `f16` to `f32` if polyfill is active.
2167    pub(super) fn store_io_with_f16_polyfill(
2168        &mut self,
2169        body: &mut Vec<Instruction>,
2170        varying_id: Word,
2171        value_id: Word,
2172    ) {
2173        if let Some(f32_ty) = self.io_f16_polyfills.get_f32_io_type(varying_id) {
2174            let converted = self.id_gen.next();
2175            super::f16_polyfill::F16IoPolyfill::emit_f16_to_f32_conversion(
2176                value_id, f32_ty, converted, body,
2177            );
2178            body.push(Instruction::store(varying_id, converted, None));
2179        } else {
2180            body.push(Instruction::store(varying_id, value_id, None));
2181        }
2182    }
2183
2184    fn write_global_variable(
2185        &mut self,
2186        ir_module: &crate::Module,
2187        global_variable: &crate::GlobalVariable,
2188    ) -> Result<Word, Error> {
2189        use spirv::Decoration;
2190
2191        let id = self.id_gen.next();
2192        let class = map_storage_class(global_variable.space);
2193
2194        //self.check(class.required_capabilities())?;
2195
2196        if self.flags.contains(WriterFlags::DEBUG) {
2197            if let Some(ref name) = global_variable.name {
2198                self.debugs.push(Instruction::name(id, name));
2199            }
2200        }
2201
2202        let storage_access = match global_variable.space {
2203            crate::AddressSpace::Storage { access } => Some(access),
2204            _ => match ir_module.types[global_variable.ty].inner {
2205                crate::TypeInner::Image {
2206                    class: crate::ImageClass::Storage { access, .. },
2207                    ..
2208                } => Some(access),
2209                _ => None,
2210            },
2211        };
2212        if let Some(storage_access) = storage_access {
2213            if !storage_access.contains(crate::StorageAccess::LOAD) {
2214                self.decorate(id, Decoration::NonReadable, &[]);
2215            }
2216            if !storage_access.contains(crate::StorageAccess::STORE) {
2217                self.decorate(id, Decoration::NonWritable, &[]);
2218            }
2219        }
2220
2221        // Note: we should be able to substitute `binding_array<Foo, 0>`,
2222        // but there is still code that tries to register the pre-substituted type,
2223        // and it is failing on 0.
2224        let mut substitute_inner_type_lookup = None;
2225        if let Some(ref res_binding) = global_variable.binding {
2226            self.decorate(id, Decoration::DescriptorSet, &[res_binding.group]);
2227            self.decorate(id, Decoration::Binding, &[res_binding.binding]);
2228
2229            if let Some(&BindingInfo {
2230                binding_array_size: Some(remapped_binding_array_size),
2231            }) = self.binding_map.get(res_binding)
2232            {
2233                if let crate::TypeInner::BindingArray { base, .. } =
2234                    ir_module.types[global_variable.ty].inner
2235                {
2236                    let binding_array_type_id =
2237                        self.get_type_id(LookupType::Local(LocalType::BindingArray {
2238                            base,
2239                            size: remapped_binding_array_size,
2240                        }));
2241                    substitute_inner_type_lookup = Some(LookupType::Local(LocalType::Pointer {
2242                        base: binding_array_type_id,
2243                        class,
2244                    }));
2245                }
2246            }
2247        };
2248
2249        let init_word = global_variable
2250            .init
2251            .map(|constant| self.constant_ids[constant]);
2252        let inner_type_id = self.get_type_id(
2253            substitute_inner_type_lookup.unwrap_or(LookupType::Handle(global_variable.ty)),
2254        );
2255
2256        // generate the wrapping structure if needed
2257        let pointer_type_id = if global_needs_wrapper(ir_module, global_variable) {
2258            let wrapper_type_id = self.id_gen.next();
2259
2260            self.decorate(wrapper_type_id, Decoration::Block, &[]);
2261            let member = crate::StructMember {
2262                name: None,
2263                ty: global_variable.ty,
2264                binding: None,
2265                offset: 0,
2266            };
2267            self.decorate_struct_member(wrapper_type_id, 0, &member, &ir_module.types)?;
2268
2269            Instruction::type_struct(wrapper_type_id, &[inner_type_id])
2270                .to_words(&mut self.logical_layout.declarations);
2271
2272            let pointer_type_id = self.id_gen.next();
2273            Instruction::type_pointer(pointer_type_id, class, wrapper_type_id)
2274                .to_words(&mut self.logical_layout.declarations);
2275
2276            pointer_type_id
2277        } else {
2278            // This is a global variable in the Storage address space. The only
2279            // way it could have `global_needs_wrapper() == false` is if it has
2280            // a runtime-sized or binding array.
2281            // Runtime-sized arrays were decorated when iterating through struct content.
2282            // Now binding arrays require Block decorating.
2283            if let crate::AddressSpace::Storage { .. } = global_variable.space {
2284                match ir_module.types[global_variable.ty].inner {
2285                    crate::TypeInner::BindingArray { base, .. } => {
2286                        let ty = &ir_module.types[base];
2287                        let mut should_decorate = true;
2288                        // Check if the type has a runtime array.
2289                        // A normal runtime array gets validated out,
2290                        // so only structs can be with runtime arrays
2291                        if let crate::TypeInner::Struct { ref members, .. } = ty.inner {
2292                            // only the last member in a struct can be dynamically sized
2293                            if let Some(last_member) = members.last() {
2294                                if let &crate::TypeInner::Array {
2295                                    size: crate::ArraySize::Dynamic,
2296                                    ..
2297                                } = &ir_module.types[last_member.ty].inner
2298                                {
2299                                    should_decorate = false;
2300                                }
2301                            }
2302                        }
2303                        if should_decorate {
2304                            let decorated_id = self.get_handle_type_id(base);
2305                            self.decorate(decorated_id, Decoration::Block, &[]);
2306                        }
2307                    }
2308                    _ => (),
2309                };
2310            }
2311            if substitute_inner_type_lookup.is_some() {
2312                inner_type_id
2313            } else {
2314                self.get_handle_pointer_type_id(global_variable.ty, class)
2315            }
2316        };
2317
2318        let init_word = match (global_variable.space, self.zero_initialize_workgroup_memory) {
2319            (crate::AddressSpace::Private, _)
2320            | (crate::AddressSpace::WorkGroup, super::ZeroInitializeWorkgroupMemoryMode::Native) => {
2321                init_word.or_else(|| Some(self.get_constant_null(inner_type_id)))
2322            }
2323            _ => init_word,
2324        };
2325
2326        Instruction::variable(pointer_type_id, id, class, init_word)
2327            .to_words(&mut self.logical_layout.declarations);
2328        Ok(id)
2329    }
2330
2331    /// Write the necessary decorations for a struct member.
2332    ///
2333    /// Emit decorations for the `index`'th member of the struct type
2334    /// designated by `struct_id`, described by `member`.
2335    fn decorate_struct_member(
2336        &mut self,
2337        struct_id: Word,
2338        index: usize,
2339        member: &crate::StructMember,
2340        arena: &UniqueArena<crate::Type>,
2341    ) -> Result<(), Error> {
2342        use spirv::Decoration;
2343
2344        self.annotations.push(Instruction::member_decorate(
2345            struct_id,
2346            index as u32,
2347            Decoration::Offset,
2348            &[member.offset],
2349        ));
2350
2351        if self.flags.contains(WriterFlags::DEBUG) {
2352            if let Some(ref name) = member.name {
2353                self.debugs
2354                    .push(Instruction::member_name(struct_id, index as u32, name));
2355            }
2356        }
2357
2358        // Matrices and (potentially nested) arrays of matrices both require decorations,
2359        // so "see through" any arrays to determine if they're needed.
2360        let mut member_array_subty_inner = &arena[member.ty].inner;
2361        while let crate::TypeInner::Array { base, .. } = *member_array_subty_inner {
2362            member_array_subty_inner = &arena[base].inner;
2363        }
2364
2365        if let crate::TypeInner::Matrix {
2366            columns: _,
2367            rows,
2368            scalar,
2369        } = *member_array_subty_inner
2370        {
2371            let byte_stride = Alignment::from(rows) * scalar.width as u32;
2372            self.annotations.push(Instruction::member_decorate(
2373                struct_id,
2374                index as u32,
2375                Decoration::ColMajor,
2376                &[],
2377            ));
2378            self.annotations.push(Instruction::member_decorate(
2379                struct_id,
2380                index as u32,
2381                Decoration::MatrixStride,
2382                &[byte_stride],
2383            ));
2384        }
2385
2386        Ok(())
2387    }
2388
2389    pub(super) fn get_function_type(&mut self, lookup_function_type: LookupFunctionType) -> Word {
2390        match self
2391            .lookup_function_type
2392            .entry(lookup_function_type.clone())
2393        {
2394            Entry::Occupied(e) => *e.get(),
2395            Entry::Vacant(_) => {
2396                let id = self.id_gen.next();
2397                let instruction = Instruction::type_function(
2398                    id,
2399                    lookup_function_type.return_type_id,
2400                    &lookup_function_type.parameter_type_ids,
2401                );
2402                instruction.to_words(&mut self.logical_layout.declarations);
2403                self.lookup_function_type.insert(lookup_function_type, id);
2404                id
2405            }
2406        }
2407    }
2408
2409    fn write_physical_layout(&mut self) {
2410        self.physical_layout.bound = self.id_gen.0 + 1;
2411    }
2412
2413    fn write_logical_layout(
2414        &mut self,
2415        ir_module: &crate::Module,
2416        mod_info: &ModuleInfo,
2417        ep_index: Option<usize>,
2418        debug_info: &Option<DebugInfo>,
2419    ) -> Result<(), Error> {
2420        fn has_view_index_check(
2421            ir_module: &crate::Module,
2422            binding: Option<&crate::Binding>,
2423            ty: Handle<crate::Type>,
2424        ) -> bool {
2425            match ir_module.types[ty].inner {
2426                crate::TypeInner::Struct { ref members, .. } => members.iter().any(|member| {
2427                    has_view_index_check(ir_module, member.binding.as_ref(), member.ty)
2428                }),
2429                _ => binding == Some(&crate::Binding::BuiltIn(crate::BuiltIn::ViewIndex)),
2430            }
2431        }
2432
2433        let has_storage_buffers =
2434            ir_module
2435                .global_variables
2436                .iter()
2437                .any(|(_, var)| match var.space {
2438                    crate::AddressSpace::Storage { .. } => true,
2439                    _ => false,
2440                });
2441        let has_view_index = ir_module
2442            .entry_points
2443            .iter()
2444            .flat_map(|entry| entry.function.arguments.iter())
2445            .any(|arg| has_view_index_check(ir_module, arg.binding.as_ref(), arg.ty));
2446        let mut has_ray_query = ir_module.special_types.ray_desc.is_some()
2447            | ir_module.special_types.ray_intersection.is_some();
2448        let has_vertex_return = ir_module.special_types.ray_vertex_return.is_some();
2449
2450        for (_, &crate::Type { ref inner, .. }) in ir_module.types.iter() {
2451            // spirv does not know whether these have vertex return - that is done by us
2452            if let &crate::TypeInner::AccelerationStructure { .. }
2453            | &crate::TypeInner::RayQuery { .. } = inner
2454            {
2455                has_ray_query = true
2456            }
2457        }
2458
2459        if self.physical_layout.version < 0x10300 && has_storage_buffers {
2460            // enable the storage buffer class on < SPV-1.3
2461            Instruction::extension("SPV_KHR_storage_buffer_storage_class")
2462                .to_words(&mut self.logical_layout.extensions);
2463        }
2464        if has_view_index {
2465            Instruction::extension("SPV_KHR_multiview")
2466                .to_words(&mut self.logical_layout.extensions)
2467        }
2468        if has_ray_query {
2469            Instruction::extension("SPV_KHR_ray_query")
2470                .to_words(&mut self.logical_layout.extensions)
2471        }
2472        if has_vertex_return {
2473            Instruction::extension("SPV_KHR_ray_tracing_position_fetch")
2474                .to_words(&mut self.logical_layout.extensions);
2475        }
2476        Instruction::type_void(self.void_type).to_words(&mut self.logical_layout.declarations);
2477        Instruction::ext_inst_import(self.gl450_ext_inst_id, "GLSL.std.450")
2478            .to_words(&mut self.logical_layout.ext_inst_imports);
2479
2480        let mut debug_info_inner = None;
2481        if self.flags.contains(WriterFlags::DEBUG) {
2482            if let Some(debug_info) = debug_info.as_ref() {
2483                let source_file_id = self.id_gen.next();
2484                self.debugs.push(Instruction::string(
2485                    &debug_info.file_name.to_string_lossy(),
2486                    source_file_id,
2487                ));
2488
2489                debug_info_inner = Some(DebugInfoInner {
2490                    source_code: debug_info.source_code,
2491                    source_file_id,
2492                });
2493                self.debugs.append(&mut Instruction::source_auto_continued(
2494                    debug_info.language,
2495                    0,
2496                    &debug_info_inner,
2497                ));
2498            }
2499        }
2500
2501        // write all types
2502        for (handle, _) in ir_module.types.iter() {
2503            self.write_type_declaration_arena(ir_module, handle)?;
2504        }
2505
2506        // write all const-expressions as constants
2507        self.constant_ids
2508            .resize(ir_module.global_expressions.len(), 0);
2509        for (handle, _) in ir_module.global_expressions.iter() {
2510            self.write_constant_expr(handle, ir_module, mod_info)?;
2511        }
2512        debug_assert!(self.constant_ids.iter().all(|&id| id != 0));
2513
2514        // write the name of constants on their respective const-expression initializer
2515        if self.flags.contains(WriterFlags::DEBUG) {
2516            for (_, constant) in ir_module.constants.iter() {
2517                if let Some(ref name) = constant.name {
2518                    let id = self.constant_ids[constant.init];
2519                    self.debugs.push(Instruction::name(id, name));
2520                }
2521            }
2522        }
2523
2524        // write all global variables
2525        for (handle, var) in ir_module.global_variables.iter() {
2526            // If a single entry point was specified, only write `OpVariable` instructions
2527            // for the globals it actually uses. Emit dummies for the others,
2528            // to preserve the indices in `global_variables`.
2529            let gvar = match ep_index {
2530                Some(index) if mod_info.get_entry_point(index)[handle].is_empty() => {
2531                    GlobalVariable::dummy()
2532                }
2533                _ => {
2534                    let id = self.write_global_variable(ir_module, var)?;
2535                    GlobalVariable::new(id)
2536                }
2537            };
2538            self.global_variables.insert(handle, gvar);
2539        }
2540
2541        // write all functions
2542        for (handle, ir_function) in ir_module.functions.iter() {
2543            let info = &mod_info[handle];
2544            if let Some(index) = ep_index {
2545                let ep_info = mod_info.get_entry_point(index);
2546                // If this function uses globals that we omitted from the SPIR-V
2547                // because the entry point and its callees didn't use them,
2548                // then we must skip it.
2549                if !ep_info.dominates_global_use(info) {
2550                    log::info!("Skip function {:?}", ir_function.name);
2551                    continue;
2552                }
2553
2554                // Skip functions that that are not compatible with this entry point's stage.
2555                //
2556                // When validation is enabled, it rejects modules whose entry points try to call
2557                // incompatible functions, so if we got this far, then any functions incompatible
2558                // with our selected entry point must not be used.
2559                //
2560                // When validation is disabled, `fun_info.available_stages` is always just
2561                // `ShaderStages::all()`, so this will write all functions in the module, and
2562                // the downstream GLSL compiler will catch any problems.
2563                if !info.available_stages.contains(ep_info.available_stages) {
2564                    continue;
2565                }
2566            }
2567            let id = self.write_function(ir_function, info, ir_module, None, &debug_info_inner)?;
2568            self.lookup_function.insert(handle, id);
2569        }
2570
2571        // write all or one entry points
2572        for (index, ir_ep) in ir_module.entry_points.iter().enumerate() {
2573            if ep_index.is_some() && ep_index != Some(index) {
2574                continue;
2575            }
2576            let info = mod_info.get_entry_point(index);
2577            let ep_instruction =
2578                self.write_entry_point(ir_ep, info, ir_module, &debug_info_inner)?;
2579            ep_instruction.to_words(&mut self.logical_layout.entry_points);
2580        }
2581
2582        for capability in self.capabilities_used.iter() {
2583            Instruction::capability(*capability).to_words(&mut self.logical_layout.capabilities);
2584        }
2585        for extension in self.extensions_used.iter() {
2586            Instruction::extension(extension).to_words(&mut self.logical_layout.extensions);
2587        }
2588        if ir_module.entry_points.is_empty() {
2589            // SPIR-V doesn't like modules without entry points
2590            Instruction::capability(spirv::Capability::Linkage)
2591                .to_words(&mut self.logical_layout.capabilities);
2592        }
2593
2594        let addressing_model = spirv::AddressingModel::Logical;
2595        let memory_model = spirv::MemoryModel::GLSL450;
2596        //self.check(addressing_model.required_capabilities())?;
2597        //self.check(memory_model.required_capabilities())?;
2598
2599        Instruction::memory_model(addressing_model, memory_model)
2600            .to_words(&mut self.logical_layout.memory_model);
2601
2602        if self.flags.contains(WriterFlags::DEBUG) {
2603            for debug in self.debugs.iter() {
2604                debug.to_words(&mut self.logical_layout.debugs);
2605            }
2606        }
2607
2608        for annotation in self.annotations.iter() {
2609            annotation.to_words(&mut self.logical_layout.annotations);
2610        }
2611
2612        Ok(())
2613    }
2614
2615    pub fn write(
2616        &mut self,
2617        ir_module: &crate::Module,
2618        info: &ModuleInfo,
2619        pipeline_options: Option<&PipelineOptions>,
2620        debug_info: &Option<DebugInfo>,
2621        words: &mut Vec<Word>,
2622    ) -> Result<(), Error> {
2623        self.reset();
2624
2625        // Try to find the entry point and corresponding index
2626        let ep_index = match pipeline_options {
2627            Some(po) => {
2628                let index = ir_module
2629                    .entry_points
2630                    .iter()
2631                    .position(|ep| po.shader_stage == ep.stage && po.entry_point == ep.name)
2632                    .ok_or(Error::EntryPointNotFound)?;
2633                Some(index)
2634            }
2635            None => None,
2636        };
2637
2638        self.write_logical_layout(ir_module, info, ep_index, debug_info)?;
2639        self.write_physical_layout();
2640
2641        self.physical_layout.in_words(words);
2642        self.logical_layout.in_words(words);
2643        Ok(())
2644    }
2645
2646    /// Return the set of capabilities the last module written used.
2647    pub const fn get_capabilities_used(&self) -> &crate::FastIndexSet<spirv::Capability> {
2648        &self.capabilities_used
2649    }
2650
2651    pub fn decorate_non_uniform_binding_array_access(&mut self, id: Word) -> Result<(), Error> {
2652        self.require_any("NonUniformEXT", &[spirv::Capability::ShaderNonUniform])?;
2653        self.use_extension("SPV_EXT_descriptor_indexing");
2654        self.decorate(id, spirv::Decoration::NonUniform, &[]);
2655        Ok(())
2656    }
2657
2658    pub(super) fn needs_f16_polyfill(&self, ty_inner: &crate::TypeInner) -> bool {
2659        self.io_f16_polyfills.needs_polyfill(ty_inner)
2660    }
2661}
2662
2663#[test]
2664fn test_write_physical_layout() {
2665    let mut writer = Writer::new(&Options::default()).unwrap();
2666    assert_eq!(writer.physical_layout.bound, 0);
2667    writer.write_physical_layout();
2668    assert_eq!(writer.physical_layout.bound, 3);
2669}