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