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