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 pub const fn lang_version(&self) -> (u8, u8) {
122 self.physical_layout.lang_version()
123 }
124
125 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 let fresh = Writer {
145 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 id_gen,
156 void_type,
157 gl450_ext_inst_id,
158
159 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 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 let selected = match self.capabilities_available {
210 None => first,
211 Some(ref available) => {
212 match capabilities
213 .iter()
214 .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 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 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 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 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 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 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 (
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 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 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 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 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 if self.physical_layout.version >= 0x10400 {
944 iface.varying_ids.push(gv.var_id);
945 }
946 }
947
948 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 gv.access_id = gv.var_id;
979 };
980 }
981 }
982
983 self.global_variables[handle] = gv;
985 }
986
987 let mut context = BlockContext {
990 ir_module,
991 ir_function,
992 fun_info: info,
993 function: &mut function,
994 cached: core::mem::take(&mut self.saved_cached),
996
997 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 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 context.cache_expression_value(handle, &mut prelude)?;
1053 }
1054 crate::Expression::Access { base, .. }
1055 | crate::Expression::AccessIndex { base, .. } => {
1056 *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 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 Instruction::execution_mode(function_id, mode, &[])
1119 .to_words(&mut self.logical_layout.execution_modes);
1120 Ok(())
1121 }
1122
1123 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 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 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 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 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 self.request_type_capabilities(&ty.inner)?;
1398 let id = if let Some(local) = self.localtype_from_inner(&ty.inner) {
1399 match self.lookup_type.entry(LookupType::Local(local)) {
1403 Entry::Occupied(e) => *e.get(),
1405
1406 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 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 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 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 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 (class == spirv::StorageClass::Input && stage == crate::ShaderStage::Vertex) ||
1981 (class == spirv::StorageClass::Output && stage == crate::ShaderStage::Fragment);
1985
1986 if !no_decorations {
1987 match interpolation {
1988 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 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 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 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 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 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 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 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 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 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 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 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 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 if let crate::TypeInner::Struct { ref members, .. } = ty.inner {
2310 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 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 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 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 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 for (handle, _) in ir_module.types.iter() {
2521 self.write_type_declaration_arena(ir_module, handle)?;
2522 }
2523
2524 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 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 for (handle, var) in ir_module.global_variables.iter() {
2544 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 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 !ep_info.dominates_global_use(info) {
2568 log::info!("Skip function {:?}", ir_function.name);
2569 continue;
2570 }
2571
2572 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 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 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 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 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 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}