1 use super::{
2 helpers::{contains_builtin, map_storage_class},
3 Instruction, LogicalLayout, Options, PhysicalLayout, WriterFlags,
4 };
5 use crate::{
6 arena::{Arena, Handle},
7 proc::TypeResolution,
8 valid::{FunctionInfo, ModuleInfo},
9 };
10 use spirv::Word;
11 use std::{collections::hash_map::Entry, ops};
12 use thiserror::Error;
13
14 const BITS_PER_BYTE: crate::Bytes = 8;
15
16 #[derive(Clone, Debug, Error)]
17 pub enum Error {
18 #[error("target SPIRV-{0}.{1} is not supported")]
19 UnsupportedVersion(u8, u8),
20 #[error("one of the required capabilities {0:?} is missing")]
21 MissingCapabilities(Vec<spirv::Capability>),
22 #[error("unimplemented {0}")]
23 FeatureNotImplemented(&'static str),
24 #[error("module is not validated properly: {0}")]
25 Validation(&'static str),
26 }
27
28 #[derive(Default)]
29 struct IdGenerator(Word);
30
31 impl IdGenerator {
next(&mut self) -> Word32 fn next(&mut self) -> Word {
33 self.0 += 1;
34 self.0
35 }
36 }
37
38 struct Block {
39 label_id: Word,
40 body: Vec<Instruction>,
41 termination: Option<Instruction>,
42 }
43
44 impl Block {
new(label_id: Word) -> Self45 fn new(label_id: Word) -> Self {
46 Block {
47 label_id,
48 body: Vec::new(),
49 termination: None,
50 }
51 }
52 }
53
54 struct LocalVariable {
55 id: Word,
56 instruction: Instruction,
57 }
58
59 struct ResultMember {
60 id: Word,
61 type_id: Word,
62 built_in: Option<crate::BuiltIn>,
63 }
64
65 struct EntryPointContext {
66 argument_ids: Vec<Word>,
67 results: Vec<ResultMember>,
68 }
69
70 #[derive(Default)]
71 struct Function {
72 signature: Option<Instruction>,
73 parameters: Vec<Instruction>,
74 variables: crate::FastHashMap<Handle<crate::LocalVariable>, LocalVariable>,
75 internal_variables: Vec<LocalVariable>,
76 blocks: Vec<Block>,
77 entry_point_context: Option<EntryPointContext>,
78 }
79
80 impl Function {
to_words(&self, sink: &mut impl Extend<Word>)81 fn to_words(&self, sink: &mut impl Extend<Word>) {
82 self.signature.as_ref().unwrap().to_words(sink);
83 for instruction in self.parameters.iter() {
84 instruction.to_words(sink);
85 }
86 for (index, block) in self.blocks.iter().enumerate() {
87 Instruction::label(block.label_id).to_words(sink);
88 if index == 0 {
89 for local_var in self.variables.values() {
90 local_var.instruction.to_words(sink);
91 }
92 for internal_var in self.internal_variables.iter() {
93 internal_var.instruction.to_words(sink);
94 }
95 }
96 for instruction in block.body.iter() {
97 instruction.to_words(sink);
98 }
99 block.termination.as_ref().unwrap().to_words(sink);
100 }
101 }
102
consume(&mut self, mut block: Block, termination: Instruction)103 fn consume(&mut self, mut block: Block, termination: Instruction) {
104 block.termination = Some(termination);
105 self.blocks.push(block);
106 }
107 }
108
109 #[derive(Debug, PartialEq, Hash, Eq, Copy, Clone)]
110 enum LocalType {
111 Value {
112 vector_size: Option<crate::VectorSize>,
113 kind: crate::ScalarKind,
114 width: crate::Bytes,
115 pointer_class: Option<spirv::StorageClass>,
116 },
117 Matrix {
118 columns: crate::VectorSize,
119 rows: crate::VectorSize,
120 width: crate::Bytes,
121 },
122 Pointer {
123 base: Handle<crate::Type>,
124 class: spirv::StorageClass,
125 },
126 Image {
127 dim: crate::ImageDimension,
128 arrayed: bool,
129 class: crate::ImageClass,
130 },
131 SampledImage {
132 image_type_id: Word,
133 },
134 Sampler,
135 }
136
137 impl PhysicalLayout {
make_local(&self, inner: &crate::TypeInner) -> Option<LocalType>138 fn make_local(&self, inner: &crate::TypeInner) -> Option<LocalType> {
139 Some(match *inner {
140 crate::TypeInner::Scalar { kind, width } => LocalType::Value {
141 vector_size: None,
142 kind,
143 width,
144 pointer_class: None,
145 },
146 crate::TypeInner::Vector { size, kind, width } => LocalType::Value {
147 vector_size: Some(size),
148 kind,
149 width,
150 pointer_class: None,
151 },
152 crate::TypeInner::Matrix {
153 columns,
154 rows,
155 width,
156 } => LocalType::Matrix {
157 columns,
158 rows,
159 width,
160 },
161 crate::TypeInner::Pointer { base, class } => LocalType::Pointer {
162 base,
163 class: map_storage_class(class),
164 },
165 crate::TypeInner::ValuePointer {
166 size,
167 kind,
168 width,
169 class,
170 } => LocalType::Value {
171 vector_size: size,
172 kind,
173 width,
174 pointer_class: Some(map_storage_class(class)),
175 },
176 crate::TypeInner::Image {
177 dim,
178 arrayed,
179 class,
180 } => LocalType::Image {
181 dim,
182 arrayed,
183 class,
184 },
185 crate::TypeInner::Sampler { comparison: _ } => LocalType::Sampler,
186 _ => return None,
187 })
188 }
189 }
190
191 #[derive(Debug, PartialEq, Hash, Eq, Copy, Clone)]
192 enum LookupType {
193 Handle(Handle<crate::Type>),
194 Local(LocalType),
195 }
196
197 impl From<LocalType> for LookupType {
from(local: LocalType) -> Self198 fn from(local: LocalType) -> Self {
199 Self::Local(local)
200 }
201 }
202
map_dim(dim: crate::ImageDimension) -> spirv::Dim203 fn map_dim(dim: crate::ImageDimension) -> spirv::Dim {
204 match dim {
205 crate::ImageDimension::D1 => spirv::Dim::Dim1D,
206 crate::ImageDimension::D2 => spirv::Dim::Dim2D,
207 crate::ImageDimension::D3 => spirv::Dim::Dim3D,
208 crate::ImageDimension::Cube => spirv::Dim::DimCube,
209 }
210 }
211
212 #[derive(Debug, PartialEq, Clone, Hash, Eq)]
213 struct LookupFunctionType {
214 parameter_type_ids: Vec<Word>,
215 return_type_id: Word,
216 }
217
218 #[derive(Debug)]
219 enum Dimension {
220 Scalar,
221 Vector,
222 Matrix,
223 }
224
get_dimension(type_inner: &crate::TypeInner) -> Dimension225 fn get_dimension(type_inner: &crate::TypeInner) -> Dimension {
226 match *type_inner {
227 crate::TypeInner::Scalar { .. } => Dimension::Scalar,
228 crate::TypeInner::Vector { .. } => Dimension::Vector,
229 crate::TypeInner::Matrix { .. } => Dimension::Matrix,
230 _ => unreachable!(),
231 }
232 }
233
234 #[derive(Clone, Copy, Default)]
235 struct LoopContext {
236 continuing_id: Option<Word>,
237 break_id: Option<Word>,
238 }
239
240 #[derive(Default)]
241 struct CachedExpressions {
242 ids: Vec<Word>,
243 }
244 impl CachedExpressions {
reset(&mut self, length: usize)245 fn reset(&mut self, length: usize) {
246 self.ids.clear();
247 self.ids.resize(length, 0);
248 }
249 }
250 impl ops::Index<Handle<crate::Expression>> for CachedExpressions {
251 type Output = Word;
index(&self, h: Handle<crate::Expression>) -> &Word252 fn index(&self, h: Handle<crate::Expression>) -> &Word {
253 let id = &self.ids[h.index()];
254 if *id == 0 {
255 unreachable!("Expression {:?} is not cached!", h);
256 }
257 id
258 }
259 }
260 impl ops::IndexMut<Handle<crate::Expression>> for CachedExpressions {
index_mut(&mut self, h: Handle<crate::Expression>) -> &mut Word261 fn index_mut(&mut self, h: Handle<crate::Expression>) -> &mut Word {
262 let id = &mut self.ids[h.index()];
263 if *id != 0 {
264 unreachable!("Expression {:?} is already cached!", h);
265 }
266 id
267 }
268 }
269
270 struct GlobalVariable {
271 /// Actual ID of the variable.
272 id: Word,
273 /// For `StorageClass::Handle` variables, this ID is recorded in the function
274 /// prelude block (and reset before every function) as `OpLoad` of the variable.
275 /// It is then used for all the global ops, such as `OpImageSample`.
276 handle_id: Word,
277 /// SPIR-V storage class.
278 class: spirv::StorageClass,
279 }
280
281 pub struct Writer {
282 physical_layout: PhysicalLayout,
283 logical_layout: LogicalLayout,
284 id_gen: IdGenerator,
285 capabilities: crate::FastHashSet<spirv::Capability>,
286 strict_capabilities: bool,
287 debugs: Vec<Instruction>,
288 annotations: Vec<Instruction>,
289 flags: WriterFlags,
290 void_type: u32,
291 //TODO: convert most of these into vectors, addressable by handle indices
292 lookup_type: crate::FastHashMap<LookupType, Word>,
293 lookup_function: crate::FastHashMap<Handle<crate::Function>, Word>,
294 lookup_function_type: crate::FastHashMap<LookupFunctionType, Word>,
295 lookup_function_call: crate::FastHashMap<Handle<crate::Expression>, Word>,
296 constant_ids: Vec<Word>,
297 cached_constants: crate::FastHashMap<(crate::ScalarValue, crate::Bytes), Word>,
298 global_variables: Vec<GlobalVariable>,
299 cached: CachedExpressions,
300 gl450_ext_inst_id: Word,
301 // Just a temporary list of SPIR-V ids
302 temp_list: Vec<Word>,
303 }
304
305 impl Writer {
new(options: &Options) -> Result<Self, Error>306 pub fn new(options: &Options) -> Result<Self, Error> {
307 let (major, minor) = options.lang_version;
308 if major != 1 {
309 return Err(Error::UnsupportedVersion(major, minor));
310 }
311 let raw_version = ((major as u32) << 16) | ((minor as u32) << 8);
312 let mut id_gen = IdGenerator::default();
313 let gl450_ext_inst_id = id_gen.next();
314 let void_type = id_gen.next();
315
316 Ok(Writer {
317 physical_layout: PhysicalLayout::new(raw_version),
318 logical_layout: LogicalLayout::default(),
319 id_gen,
320 capabilities: match options.capabilities {
321 Some(ref caps) => caps.clone(),
322 None => {
323 let mut caps = crate::FastHashSet::default();
324 caps.insert(spirv::Capability::Shader);
325 caps
326 }
327 },
328 strict_capabilities: options.capabilities.is_some(),
329 debugs: vec![],
330 annotations: vec![],
331 flags: options.flags,
332 void_type,
333 lookup_type: crate::FastHashMap::default(),
334 lookup_function: crate::FastHashMap::default(),
335 lookup_function_type: crate::FastHashMap::default(),
336 lookup_function_call: crate::FastHashMap::default(),
337 constant_ids: Vec::new(),
338 cached_constants: crate::FastHashMap::default(),
339 global_variables: Vec::new(),
340 cached: CachedExpressions::default(),
341 gl450_ext_inst_id,
342 temp_list: Vec::new(),
343 })
344 }
345
check(&mut self, capabilities: &[spirv::Capability]) -> Result<(), Error>346 fn check(&mut self, capabilities: &[spirv::Capability]) -> Result<(), Error> {
347 if self.strict_capabilities {
348 if capabilities.is_empty()
349 || capabilities
350 .iter()
351 .any(|cap| self.capabilities.contains(cap))
352 {
353 Ok(())
354 } else {
355 Err(Error::MissingCapabilities(capabilities.to_vec()))
356 }
357 } else {
358 self.capabilities.extend(capabilities);
359 Ok(())
360 }
361 }
362
get_type_id( &mut self, arena: &Arena<crate::Type>, lookup_ty: LookupType, ) -> Result<Word, Error>363 fn get_type_id(
364 &mut self,
365 arena: &Arena<crate::Type>,
366 lookup_ty: LookupType,
367 ) -> Result<Word, Error> {
368 if let Entry::Occupied(e) = self.lookup_type.entry(lookup_ty) {
369 Ok(*e.get())
370 } else {
371 match lookup_ty {
372 LookupType::Handle(_handle) => unreachable!("Handles are populated at start"),
373 LookupType::Local(local_ty) => self.write_type_declaration_local(arena, local_ty),
374 }
375 }
376 }
377
get_expression_type_id( &mut self, arena: &Arena<crate::Type>, tr: &TypeResolution, ) -> Result<Word, Error>378 fn get_expression_type_id(
379 &mut self,
380 arena: &Arena<crate::Type>,
381 tr: &TypeResolution,
382 ) -> Result<Word, Error> {
383 let lookup_ty = match *tr {
384 TypeResolution::Handle(ty_handle) => LookupType::Handle(ty_handle),
385 TypeResolution::Value(ref inner) => {
386 LookupType::Local(self.physical_layout.make_local(inner).unwrap())
387 }
388 };
389 self.get_type_id(arena, lookup_ty)
390 }
391
get_pointer_id( &mut self, arena: &Arena<crate::Type>, handle: Handle<crate::Type>, class: spirv::StorageClass, ) -> Result<Word, Error>392 fn get_pointer_id(
393 &mut self,
394 arena: &Arena<crate::Type>,
395 handle: Handle<crate::Type>,
396 class: spirv::StorageClass,
397 ) -> Result<Word, Error> {
398 let ty_id = self.get_type_id(arena, LookupType::Handle(handle))?;
399 if let crate::TypeInner::Pointer { .. } = arena[handle].inner {
400 return Ok(ty_id);
401 }
402 let lookup_type = LookupType::Local(LocalType::Pointer {
403 base: handle,
404 class,
405 });
406 Ok(if let Some(&id) = self.lookup_type.get(&lookup_type) {
407 id
408 } else {
409 let id = self.id_gen.next();
410 let instruction = Instruction::type_pointer(id, class, ty_id);
411 instruction.to_words(&mut self.logical_layout.declarations);
412 self.lookup_type.insert(lookup_type, id);
413 id
414 })
415 }
416
decorate(&mut self, id: Word, decoration: spirv::Decoration, operands: &[Word])417 fn decorate(&mut self, id: Word, decoration: spirv::Decoration, operands: &[Word]) {
418 self.annotations
419 .push(Instruction::decorate(id, decoration, operands));
420 }
421
write_function( &mut self, ir_function: &crate::Function, info: &FunctionInfo, ir_module: &crate::Module, mut varying_ids: Option<&mut Vec<Word>>, ) -> Result<Word, Error>422 fn write_function(
423 &mut self,
424 ir_function: &crate::Function,
425 info: &FunctionInfo,
426 ir_module: &crate::Module,
427 mut varying_ids: Option<&mut Vec<Word>>,
428 ) -> Result<Word, Error> {
429 let mut function = Function::default();
430
431 for (handle, variable) in ir_function.local_variables.iter() {
432 let id = self.id_gen.next();
433
434 if self.flags.contains(WriterFlags::DEBUG) {
435 if let Some(ref name) = variable.name {
436 self.debugs.push(Instruction::name(id, name));
437 }
438 }
439
440 let init_word = variable
441 .init
442 .map(|constant| self.constant_ids[constant.index()]);
443 let pointer_type_id =
444 self.get_pointer_id(&ir_module.types, variable.ty, spirv::StorageClass::Function)?;
445 let instruction = Instruction::variable(
446 pointer_type_id,
447 id,
448 spirv::StorageClass::Function,
449 init_word,
450 );
451 function
452 .variables
453 .insert(handle, LocalVariable { id, instruction });
454 }
455
456 let prelude_id = self.id_gen.next();
457 let mut prelude = Block::new(prelude_id);
458 let mut ep_context = EntryPointContext {
459 argument_ids: Vec::new(),
460 results: Vec::new(),
461 };
462
463 let mut parameter_type_ids = Vec::with_capacity(ir_function.arguments.len());
464 for argument in ir_function.arguments.iter() {
465 let class = spirv::StorageClass::Input;
466 let argument_type_id =
467 self.get_type_id(&ir_module.types, LookupType::Handle(argument.ty))?;
468 if let Some(ref mut list) = varying_ids {
469 let id = if let Some(ref binding) = argument.binding {
470 let name = argument.name.as_ref().map(AsRef::as_ref);
471 let varying_id =
472 self.write_varying(ir_module, class, name, argument.ty, binding)?;
473 list.push(varying_id);
474 let id = self.id_gen.next();
475 prelude
476 .body
477 .push(Instruction::load(argument_type_id, id, varying_id, None));
478 id
479 } else if let crate::TypeInner::Struct { ref members, .. } =
480 ir_module.types[argument.ty].inner
481 {
482 let struct_id = self.id_gen.next();
483 let mut constituent_ids = Vec::with_capacity(members.len());
484 for member in members {
485 let type_id =
486 self.get_type_id(&ir_module.types, LookupType::Handle(member.ty))?;
487 let name = member.name.as_ref().map(AsRef::as_ref);
488 let binding = member.binding.as_ref().unwrap();
489 let varying_id =
490 self.write_varying(ir_module, class, name, member.ty, binding)?;
491 list.push(varying_id);
492 let id = self.id_gen.next();
493 prelude
494 .body
495 .push(Instruction::load(type_id, id, varying_id, None));
496 constituent_ids.push(id);
497 }
498 prelude.body.push(Instruction::composite_construct(
499 argument_type_id,
500 struct_id,
501 &constituent_ids,
502 ));
503 struct_id
504 } else {
505 unreachable!("Missing argument binding on an entry point");
506 };
507 ep_context.argument_ids.push(id);
508 } else {
509 let id = self.id_gen.next();
510 let instruction = Instruction::function_parameter(argument_type_id, id);
511 function.parameters.push(instruction);
512 parameter_type_ids.push(argument_type_id);
513 };
514 }
515
516 let return_type_id = match ir_function.result {
517 Some(ref result) => {
518 if let Some(ref mut list) = varying_ids {
519 let class = spirv::StorageClass::Output;
520 if let Some(ref binding) = result.binding {
521 let type_id =
522 self.get_type_id(&ir_module.types, LookupType::Handle(result.ty))?;
523 let varying_id =
524 self.write_varying(ir_module, class, None, result.ty, binding)?;
525 list.push(varying_id);
526 ep_context.results.push(ResultMember {
527 id: varying_id,
528 type_id,
529 built_in: binding.to_built_in(),
530 });
531 } else if let crate::TypeInner::Struct { ref members, .. } =
532 ir_module.types[result.ty].inner
533 {
534 for member in members {
535 let type_id =
536 self.get_type_id(&ir_module.types, LookupType::Handle(member.ty))?;
537 let name = member.name.as_ref().map(AsRef::as_ref);
538 let binding = member.binding.as_ref().unwrap();
539 let varying_id =
540 self.write_varying(ir_module, class, name, member.ty, binding)?;
541 list.push(varying_id);
542 ep_context.results.push(ResultMember {
543 id: varying_id,
544 type_id,
545 built_in: binding.to_built_in(),
546 });
547 }
548 } else {
549 unreachable!("Missing result binding on an entry point");
550 }
551 self.void_type
552 } else {
553 self.get_type_id(&ir_module.types, LookupType::Handle(result.ty))?
554 }
555 }
556 None => self.void_type,
557 };
558
559 let lookup_function_type = LookupFunctionType {
560 parameter_type_ids,
561 return_type_id,
562 };
563
564 let function_id = self.id_gen.next();
565 if self.flags.contains(WriterFlags::DEBUG) {
566 if let Some(ref name) = ir_function.name {
567 self.debugs.push(Instruction::name(function_id, name));
568 }
569 }
570
571 let function_type = self.get_function_type(lookup_function_type);
572 function.signature = Some(Instruction::function(
573 return_type_id,
574 function_id,
575 spirv::FunctionControl::empty(),
576 function_type,
577 ));
578
579 if varying_ids.is_some() {
580 function.entry_point_context = Some(ep_context);
581 }
582
583 // fill up the `GlobalVariable::handle_id`
584 for gv in self.global_variables.iter_mut() {
585 gv.handle_id = 0;
586 }
587 for (handle, var) in ir_module.global_variables.iter() {
588 // Handle globals are pre-emitted and should be loaded automatically.
589 if info[handle].is_empty() || var.class != crate::StorageClass::Handle {
590 continue;
591 }
592 let id = self.id_gen.next();
593 let result_type_id = self.get_type_id(&ir_module.types, LookupType::Handle(var.ty))?;
594 let gv = &mut self.global_variables[handle.index()];
595 prelude
596 .body
597 .push(Instruction::load(result_type_id, id, gv.id, None));
598 gv.handle_id = id;
599 }
600 // fill up the pre-emitted expressions
601 self.cached.reset(ir_function.expressions.len());
602 for (handle, expr) in ir_function.expressions.iter() {
603 if expr.needs_pre_emit() {
604 self.cache_expression_value(
605 ir_module,
606 ir_function,
607 info,
608 handle,
609 &mut prelude,
610 &mut function,
611 )?;
612 }
613 }
614
615 let main_id = self.id_gen.next();
616 function.consume(prelude, Instruction::branch(main_id));
617 self.write_block(
618 main_id,
619 &ir_function.body,
620 ir_module,
621 ir_function,
622 info,
623 &mut function,
624 None,
625 LoopContext::default(),
626 )?;
627
628 function.to_words(&mut self.logical_layout.function_definitions);
629 Instruction::function_end().to_words(&mut self.logical_layout.function_definitions);
630
631 Ok(function_id)
632 }
633
write_execution_mode( &mut self, function_id: Word, mode: spirv::ExecutionMode, ) -> Result<(), Error>634 fn write_execution_mode(
635 &mut self,
636 function_id: Word,
637 mode: spirv::ExecutionMode,
638 ) -> Result<(), Error> {
639 self.check(mode.required_capabilities())?;
640 Instruction::execution_mode(function_id, mode, &[])
641 .to_words(&mut self.logical_layout.execution_modes);
642 Ok(())
643 }
644
645 // TODO Move to instructions module
write_entry_point( &mut self, entry_point: &crate::EntryPoint, info: &FunctionInfo, ir_module: &crate::Module, ) -> Result<Instruction, Error>646 fn write_entry_point(
647 &mut self,
648 entry_point: &crate::EntryPoint,
649 info: &FunctionInfo,
650 ir_module: &crate::Module,
651 ) -> Result<Instruction, Error> {
652 let mut interface_ids = Vec::new();
653 let function_id = self.write_function(
654 &entry_point.function,
655 info,
656 ir_module,
657 Some(&mut interface_ids),
658 )?;
659
660 let exec_model = match entry_point.stage {
661 crate::ShaderStage::Vertex => spirv::ExecutionModel::Vertex,
662 crate::ShaderStage::Fragment => {
663 self.write_execution_mode(function_id, spirv::ExecutionMode::OriginUpperLeft)?;
664 if let Some(ref result) = entry_point.function.result {
665 if contains_builtin(
666 result.binding.as_ref(),
667 result.ty,
668 &ir_module.types,
669 crate::BuiltIn::FragDepth,
670 ) {
671 self.write_execution_mode(
672 function_id,
673 spirv::ExecutionMode::DepthReplacing,
674 )?;
675 }
676 }
677 spirv::ExecutionModel::Fragment
678 }
679 crate::ShaderStage::Compute => {
680 let execution_mode = spirv::ExecutionMode::LocalSize;
681 self.check(execution_mode.required_capabilities())?;
682 Instruction::execution_mode(
683 function_id,
684 execution_mode,
685 &entry_point.workgroup_size,
686 )
687 .to_words(&mut self.logical_layout.execution_modes);
688 spirv::ExecutionModel::GLCompute
689 }
690 };
691 self.check(exec_model.required_capabilities())?;
692
693 Ok(Instruction::entry_point(
694 exec_model,
695 function_id,
696 &entry_point.name,
697 interface_ids.as_slice(),
698 ))
699 }
700
make_scalar( &mut self, id: Word, kind: crate::ScalarKind, width: crate::Bytes, ) -> Instruction701 fn make_scalar(
702 &mut self,
703 id: Word,
704 kind: crate::ScalarKind,
705 width: crate::Bytes,
706 ) -> Instruction {
707 use crate::ScalarKind as Sk;
708
709 let bits = (width * BITS_PER_BYTE) as u32;
710 match kind {
711 Sk::Sint | Sk::Uint => {
712 let signedness = if kind == Sk::Sint {
713 super::instructions::Signedness::Signed
714 } else {
715 super::instructions::Signedness::Unsigned
716 };
717 let cap = match bits {
718 8 => Some(spirv::Capability::Int8),
719 16 => Some(spirv::Capability::Int16),
720 64 => Some(spirv::Capability::Int64),
721 _ => None,
722 };
723 if let Some(cap) = cap {
724 self.capabilities.insert(cap);
725 }
726 Instruction::type_int(id, bits, signedness)
727 }
728 crate::ScalarKind::Float => {
729 if bits == 64 {
730 self.capabilities.insert(spirv::Capability::Float64);
731 }
732 Instruction::type_float(id, bits)
733 }
734 Sk::Bool => Instruction::type_bool(id),
735 }
736 }
737
write_type_declaration_local( &mut self, arena: &Arena<crate::Type>, local_ty: LocalType, ) -> Result<Word, Error>738 fn write_type_declaration_local(
739 &mut self,
740 arena: &Arena<crate::Type>,
741 local_ty: LocalType,
742 ) -> Result<Word, Error> {
743 let id = self.id_gen.next();
744 let instruction = match local_ty {
745 LocalType::Value {
746 vector_size: None,
747 kind,
748 width,
749 pointer_class: None,
750 } => self.make_scalar(id, kind, width),
751 LocalType::Value {
752 vector_size: Some(size),
753 kind,
754 width,
755 pointer_class: None,
756 } => {
757 let scalar_id = self.get_type_id(
758 arena,
759 LookupType::Local(LocalType::Value {
760 vector_size: None,
761 kind,
762 width,
763 pointer_class: None,
764 }),
765 )?;
766 Instruction::type_vector(id, scalar_id, size)
767 }
768 LocalType::Matrix {
769 columns,
770 rows,
771 width,
772 } => {
773 let vector_id = self.get_type_id(
774 arena,
775 LookupType::Local(LocalType::Value {
776 vector_size: Some(rows),
777 kind: crate::ScalarKind::Float,
778 width,
779 pointer_class: None,
780 }),
781 )?;
782 Instruction::type_matrix(id, vector_id, columns)
783 }
784 LocalType::Pointer { base, class } => {
785 let type_id = self.get_type_id(arena, LookupType::Handle(base))?;
786 Instruction::type_pointer(id, class, type_id)
787 }
788 LocalType::Value {
789 vector_size,
790 kind,
791 width,
792 pointer_class: Some(class),
793 } => {
794 let type_id = self.get_type_id(
795 arena,
796 LookupType::Local(LocalType::Value {
797 vector_size,
798 kind,
799 width,
800 pointer_class: None,
801 }),
802 )?;
803 Instruction::type_pointer(id, class, type_id)
804 }
805 // all the samplers and image types go through `write_type_declaration_arena`
806 LocalType::Image { .. } | LocalType::Sampler => unreachable!(),
807 LocalType::SampledImage { image_type_id } => {
808 Instruction::type_sampled_image(id, image_type_id)
809 }
810 };
811
812 self.lookup_type.insert(LookupType::Local(local_ty), id);
813 instruction.to_words(&mut self.logical_layout.declarations);
814 Ok(id)
815 }
816
write_type_declaration_arena( &mut self, arena: &Arena<crate::Type>, handle: Handle<crate::Type>, ) -> Result<Word, Error>817 fn write_type_declaration_arena(
818 &mut self,
819 arena: &Arena<crate::Type>,
820 handle: Handle<crate::Type>,
821 ) -> Result<Word, Error> {
822 let ty = &arena[handle];
823 let decorate_layout = true; //TODO?
824
825 let id = if let Some(local) = self.physical_layout.make_local(&ty.inner) {
826 match self.lookup_type.entry(LookupType::Local(local)) {
827 // if it's already known as local, re-use it
828 Entry::Occupied(e) => {
829 let id = *e.into_mut();
830 self.lookup_type.insert(LookupType::Handle(handle), id);
831 return Ok(id);
832 }
833 // also register the type as "local", to avoid duplication
834 Entry::Vacant(e) => {
835 let id = self.id_gen.next();
836 *e.insert(id)
837 }
838 }
839 } else {
840 self.id_gen.next()
841 };
842 self.lookup_type.insert(LookupType::Handle(handle), id);
843
844 if self.flags.contains(WriterFlags::DEBUG) {
845 if let Some(ref name) = ty.name {
846 self.debugs.push(Instruction::name(id, name));
847 }
848 }
849
850 use spirv::Decoration;
851
852 let instruction = match ty.inner {
853 crate::TypeInner::Scalar { kind, width } => self.make_scalar(id, kind, width),
854 crate::TypeInner::Vector { size, kind, width } => {
855 let scalar_id = self.get_type_id(
856 arena,
857 LookupType::Local(LocalType::Value {
858 vector_size: None,
859 kind,
860 width,
861 pointer_class: None,
862 }),
863 )?;
864 Instruction::type_vector(id, scalar_id, size)
865 }
866 crate::TypeInner::Matrix {
867 columns,
868 rows,
869 width,
870 } => {
871 let vector_id = self.get_type_id(
872 arena,
873 LookupType::Local(LocalType::Value {
874 vector_size: Some(rows),
875 kind: crate::ScalarKind::Float,
876 width,
877 pointer_class: None,
878 }),
879 )?;
880 Instruction::type_matrix(id, vector_id, columns)
881 }
882 crate::TypeInner::Image {
883 dim,
884 arrayed,
885 class,
886 } => {
887 let kind = match class {
888 crate::ImageClass::Sampled { kind, multi: _ } => kind,
889 crate::ImageClass::Depth => crate::ScalarKind::Float,
890 crate::ImageClass::Storage(format) => format.into(),
891 };
892 let local_type = LocalType::Value {
893 vector_size: None,
894 kind,
895 width: 4,
896 pointer_class: None,
897 };
898 let type_id = self.get_type_id(arena, LookupType::Local(local_type))?;
899 let dim = map_dim(dim);
900 self.check(dim.required_capabilities())?;
901 Instruction::type_image(id, type_id, dim, arrayed, class)
902 }
903 crate::TypeInner::Sampler { comparison: _ } => Instruction::type_sampler(id),
904 crate::TypeInner::Array { base, size, stride } => {
905 if decorate_layout {
906 self.decorate(id, Decoration::ArrayStride, &[stride]);
907 }
908
909 let type_id = self.get_type_id(arena, LookupType::Handle(base))?;
910 match size {
911 crate::ArraySize::Constant(const_handle) => {
912 let length_id = self.constant_ids[const_handle.index()];
913 Instruction::type_array(id, type_id, length_id)
914 }
915 crate::ArraySize::Dynamic => Instruction::type_runtime_array(id, type_id),
916 }
917 }
918 crate::TypeInner::Struct {
919 ref level,
920 ref members,
921 span: _,
922 } => {
923 if let crate::StructLevel::Root = *level {
924 self.decorate(id, Decoration::Block, &[]);
925 }
926
927 let mut member_ids = Vec::with_capacity(members.len());
928 for (index, member) in members.iter().enumerate() {
929 if decorate_layout {
930 self.annotations.push(Instruction::member_decorate(
931 id,
932 index as u32,
933 Decoration::Offset,
934 &[member.offset],
935 ));
936 }
937
938 if self.flags.contains(WriterFlags::DEBUG) {
939 if let Some(ref name) = member.name {
940 self.debugs
941 .push(Instruction::member_name(id, index as u32, name));
942 }
943 }
944
945 // The matrix decorations also go on arrays of matrices,
946 // so lets check this first.
947 let member_array_subty_inner = match arena[member.ty].inner {
948 crate::TypeInner::Array { base, .. } => &arena[base].inner,
949 ref other => other,
950 };
951 if let crate::TypeInner::Matrix {
952 columns,
953 rows: _,
954 width,
955 } = *member_array_subty_inner
956 {
957 let byte_stride = match columns {
958 crate::VectorSize::Bi => 2 * width,
959 crate::VectorSize::Tri | crate::VectorSize::Quad => 4 * width,
960 };
961 self.annotations.push(Instruction::member_decorate(
962 id,
963 index as u32,
964 Decoration::ColMajor,
965 &[],
966 ));
967 self.annotations.push(Instruction::member_decorate(
968 id,
969 index as u32,
970 Decoration::MatrixStride,
971 &[byte_stride as u32],
972 ));
973 }
974
975 let member_id = self.get_type_id(arena, LookupType::Handle(member.ty))?;
976 member_ids.push(member_id);
977 }
978 Instruction::type_struct(id, member_ids.as_slice())
979 }
980 crate::TypeInner::Pointer { base, class } => {
981 let type_id = self.get_type_id(arena, LookupType::Handle(base))?;
982 let raw_class = map_storage_class(class);
983 Instruction::type_pointer(id, raw_class, type_id)
984 }
985 crate::TypeInner::ValuePointer {
986 size,
987 kind,
988 width,
989 class,
990 } => {
991 let raw_class = map_storage_class(class);
992 let type_id = self.get_type_id(
993 arena,
994 LookupType::Local(LocalType::Value {
995 vector_size: size,
996 kind,
997 width,
998 pointer_class: None,
999 }),
1000 )?;
1001 Instruction::type_pointer(id, raw_class, type_id)
1002 }
1003 };
1004
1005 instruction.to_words(&mut self.logical_layout.declarations);
1006 Ok(id)
1007 }
1008
get_index_constant( &mut self, index: Word, types: &Arena<crate::Type>, ) -> Result<Word, Error>1009 fn get_index_constant(
1010 &mut self,
1011 index: Word,
1012 types: &Arena<crate::Type>,
1013 ) -> Result<Word, Error> {
1014 self.get_constant_scalar(crate::ScalarValue::Uint(index as _), 4, types)
1015 }
1016
get_constant_scalar( &mut self, value: crate::ScalarValue, width: crate::Bytes, types: &Arena<crate::Type>, ) -> Result<Word, Error>1017 fn get_constant_scalar(
1018 &mut self,
1019 value: crate::ScalarValue,
1020 width: crate::Bytes,
1021 types: &Arena<crate::Type>,
1022 ) -> Result<Word, Error> {
1023 if let Some(&id) = self.cached_constants.get(&(value, width)) {
1024 return Ok(id);
1025 }
1026 let id = self.id_gen.next();
1027 self.write_constant_scalar(id, &value, width, None, types)?;
1028 self.cached_constants.insert((value, width), id);
1029 Ok(id)
1030 }
1031
write_constant_scalar( &mut self, id: Word, value: &crate::ScalarValue, width: crate::Bytes, debug_name: Option<&String>, types: &Arena<crate::Type>, ) -> Result<(), Error>1032 fn write_constant_scalar(
1033 &mut self,
1034 id: Word,
1035 value: &crate::ScalarValue,
1036 width: crate::Bytes,
1037 debug_name: Option<&String>,
1038 types: &Arena<crate::Type>,
1039 ) -> Result<(), Error> {
1040 if self.flags.contains(WriterFlags::DEBUG) {
1041 if let Some(name) = debug_name {
1042 self.debugs.push(Instruction::name(id, name));
1043 }
1044 }
1045 let type_id = self.get_type_id(
1046 types,
1047 LookupType::Local(LocalType::Value {
1048 vector_size: None,
1049 kind: value.scalar_kind(),
1050 width,
1051 pointer_class: None,
1052 }),
1053 )?;
1054 let (solo, pair);
1055 let instruction = match *value {
1056 crate::ScalarValue::Sint(val) => {
1057 let words = match width {
1058 4 => {
1059 solo = [val as u32];
1060 &solo[..]
1061 }
1062 8 => {
1063 pair = [(val >> 32) as u32, val as u32];
1064 &pair
1065 }
1066 _ => unreachable!(),
1067 };
1068 Instruction::constant(type_id, id, words)
1069 }
1070 crate::ScalarValue::Uint(val) => {
1071 let words = match width {
1072 4 => {
1073 solo = [val as u32];
1074 &solo[..]
1075 }
1076 8 => {
1077 pair = [(val >> 32) as u32, val as u32];
1078 &pair
1079 }
1080 _ => unreachable!(),
1081 };
1082 Instruction::constant(type_id, id, words)
1083 }
1084 crate::ScalarValue::Float(val) => {
1085 let words = match width {
1086 4 => {
1087 solo = [(val as f32).to_bits()];
1088 &solo[..]
1089 }
1090 8 => {
1091 let bits = f64::to_bits(val);
1092 pair = [(bits >> 32) as u32, bits as u32];
1093 &pair
1094 }
1095 _ => unreachable!(),
1096 };
1097 Instruction::constant(type_id, id, words)
1098 }
1099 crate::ScalarValue::Bool(true) => Instruction::constant_true(type_id, id),
1100 crate::ScalarValue::Bool(false) => Instruction::constant_false(type_id, id),
1101 };
1102
1103 instruction.to_words(&mut self.logical_layout.declarations);
1104 Ok(())
1105 }
1106
write_constant_composite( &mut self, id: Word, ty: Handle<crate::Type>, components: &[Handle<crate::Constant>], types: &Arena<crate::Type>, ) -> Result<(), Error>1107 fn write_constant_composite(
1108 &mut self,
1109 id: Word,
1110 ty: Handle<crate::Type>,
1111 components: &[Handle<crate::Constant>],
1112 types: &Arena<crate::Type>,
1113 ) -> Result<(), Error> {
1114 let mut constituent_ids = Vec::with_capacity(components.len());
1115 for constituent in components.iter() {
1116 let constituent_id = self.constant_ids[constituent.index()];
1117 constituent_ids.push(constituent_id);
1118 }
1119
1120 let type_id = self.get_type_id(types, LookupType::Handle(ty))?;
1121 Instruction::constant_composite(type_id, id, constituent_ids.as_slice())
1122 .to_words(&mut self.logical_layout.declarations);
1123 Ok(())
1124 }
1125
write_varying( &mut self, ir_module: &crate::Module, class: spirv::StorageClass, debug_name: Option<&str>, ty: Handle<crate::Type>, binding: &crate::Binding, ) -> Result<Word, Error>1126 fn write_varying(
1127 &mut self,
1128 ir_module: &crate::Module,
1129 class: spirv::StorageClass,
1130 debug_name: Option<&str>,
1131 ty: Handle<crate::Type>,
1132 binding: &crate::Binding,
1133 ) -> Result<Word, Error> {
1134 let id = self.id_gen.next();
1135 let pointer_type_id = self.get_pointer_id(&ir_module.types, ty, class)?;
1136 Instruction::variable(pointer_type_id, id, class, None)
1137 .to_words(&mut self.logical_layout.declarations);
1138
1139 if self.flags.contains(WriterFlags::DEBUG) {
1140 if let Some(name) = debug_name {
1141 self.debugs.push(Instruction::name(id, name));
1142 }
1143 }
1144
1145 use spirv::{BuiltIn, Decoration};
1146
1147 match *binding {
1148 crate::Binding::Location {
1149 location,
1150 interpolation,
1151 sampling,
1152 } => {
1153 self.decorate(id, Decoration::Location, &[location]);
1154
1155 match interpolation {
1156 // Perspective-correct interpolation is the default in SPIR-V.
1157 None | Some(crate::Interpolation::Perspective) => (),
1158 Some(crate::Interpolation::Flat) => {
1159 self.decorate(id, Decoration::Flat, &[]);
1160 }
1161 Some(crate::Interpolation::Linear) => {
1162 self.decorate(id, Decoration::NoPerspective, &[]);
1163 }
1164 }
1165
1166 match sampling {
1167 // Center sampling is the default in SPIR-V.
1168 None | Some(crate::Sampling::Center) => (),
1169 Some(crate::Sampling::Centroid) => {
1170 self.decorate(id, Decoration::Centroid, &[]);
1171 }
1172 Some(crate::Sampling::Sample) => {
1173 self.decorate(id, Decoration::Sample, &[]);
1174 }
1175 }
1176 }
1177 crate::Binding::BuiltIn(built_in) => {
1178 use crate::BuiltIn as Bi;
1179 let built_in = match built_in {
1180 Bi::Position => {
1181 if class == spirv::StorageClass::Output {
1182 BuiltIn::Position
1183 } else {
1184 BuiltIn::FragCoord
1185 }
1186 }
1187 // vertex
1188 Bi::BaseInstance => BuiltIn::BaseInstance,
1189 Bi::BaseVertex => BuiltIn::BaseVertex,
1190 Bi::ClipDistance => BuiltIn::ClipDistance,
1191 Bi::CullDistance => BuiltIn::CullDistance,
1192 Bi::InstanceIndex => BuiltIn::InstanceIndex,
1193 Bi::PointSize => BuiltIn::PointSize,
1194 Bi::VertexIndex => BuiltIn::VertexIndex,
1195 // fragment
1196 Bi::FragDepth => BuiltIn::FragDepth,
1197 Bi::FrontFacing => BuiltIn::FrontFacing,
1198 Bi::SampleIndex => BuiltIn::SampleId,
1199 Bi::SampleMask => BuiltIn::SampleMask,
1200 // compute
1201 Bi::GlobalInvocationId => BuiltIn::GlobalInvocationId,
1202 Bi::LocalInvocationId => BuiltIn::LocalInvocationId,
1203 Bi::LocalInvocationIndex => BuiltIn::LocalInvocationIndex,
1204 Bi::WorkGroupId => BuiltIn::WorkgroupId,
1205 Bi::WorkGroupSize => BuiltIn::WorkgroupSize,
1206 };
1207
1208 self.decorate(id, Decoration::BuiltIn, &[built_in as u32]);
1209 }
1210 }
1211
1212 Ok(id)
1213 }
1214
write_global_variable( &mut self, ir_module: &crate::Module, global_variable: &crate::GlobalVariable, ) -> Result<(Instruction, Word, spirv::StorageClass), Error>1215 fn write_global_variable(
1216 &mut self,
1217 ir_module: &crate::Module,
1218 global_variable: &crate::GlobalVariable,
1219 ) -> Result<(Instruction, Word, spirv::StorageClass), Error> {
1220 let id = self.id_gen.next();
1221
1222 let class = map_storage_class(global_variable.class);
1223 self.check(class.required_capabilities())?;
1224
1225 let init_word = global_variable
1226 .init
1227 .map(|constant| self.constant_ids[constant.index()]);
1228 let pointer_type_id = self.get_pointer_id(&ir_module.types, global_variable.ty, class)?;
1229 let instruction = Instruction::variable(pointer_type_id, id, class, init_word);
1230
1231 if self.flags.contains(WriterFlags::DEBUG) {
1232 if let Some(ref name) = global_variable.name {
1233 self.debugs.push(Instruction::name(id, name));
1234 }
1235 }
1236
1237 use spirv::Decoration;
1238
1239 let access_decoration = match global_variable.storage_access {
1240 crate::StorageAccess::LOAD => Some(Decoration::NonWritable),
1241 crate::StorageAccess::STORE => Some(Decoration::NonReadable),
1242 _ => None,
1243 };
1244 if let Some(decoration) = access_decoration {
1245 self.decorate(id, decoration, &[]);
1246 }
1247
1248 if let Some(ref res_binding) = global_variable.binding {
1249 self.decorate(id, Decoration::DescriptorSet, &[res_binding.group]);
1250 self.decorate(id, Decoration::Binding, &[res_binding.binding]);
1251 }
1252
1253 // TODO Initializer is optional and not (yet) included in the IR
1254 Ok((instruction, id, class))
1255 }
1256
get_function_type(&mut self, lookup_function_type: LookupFunctionType) -> Word1257 fn get_function_type(&mut self, lookup_function_type: LookupFunctionType) -> Word {
1258 match self
1259 .lookup_function_type
1260 .entry(lookup_function_type.clone())
1261 {
1262 Entry::Occupied(e) => *e.get(),
1263 _ => {
1264 let id = self.id_gen.next();
1265 let instruction = Instruction::type_function(
1266 id,
1267 lookup_function_type.return_type_id,
1268 &lookup_function_type.parameter_type_ids,
1269 );
1270 instruction.to_words(&mut self.logical_layout.declarations);
1271 self.lookup_function_type.insert(lookup_function_type, id);
1272 id
1273 }
1274 }
1275 }
1276
write_texture_coordinates( &mut self, ir_module: &crate::Module, fun_info: &FunctionInfo, coordinates: Handle<crate::Expression>, array_index: Option<Handle<crate::Expression>>, block: &mut Block, ) -> Result<Word, Error>1277 fn write_texture_coordinates(
1278 &mut self,
1279 ir_module: &crate::Module,
1280 fun_info: &FunctionInfo,
1281 coordinates: Handle<crate::Expression>,
1282 array_index: Option<Handle<crate::Expression>>,
1283 block: &mut Block,
1284 ) -> Result<Word, Error> {
1285 let coordinate_id = self.cached[coordinates];
1286
1287 Ok(if let Some(array_index) = array_index {
1288 let coordinate_scalar_type_id = self.get_type_id(
1289 &ir_module.types,
1290 LookupType::Local(LocalType::Value {
1291 vector_size: None,
1292 kind: crate::ScalarKind::Float,
1293 width: 4,
1294 pointer_class: None,
1295 }),
1296 )?;
1297
1298 let mut constituent_ids = [0u32; 4];
1299 let size = match *fun_info[coordinates].ty.inner_with(&ir_module.types) {
1300 crate::TypeInner::Scalar { .. } => {
1301 constituent_ids[0] = coordinate_id;
1302 crate::VectorSize::Bi
1303 }
1304 crate::TypeInner::Vector { size, .. } => {
1305 for i in 0..size as u32 {
1306 let id = self.id_gen.next();
1307 constituent_ids[i as usize] = id;
1308 block.body.push(Instruction::composite_extract(
1309 coordinate_scalar_type_id,
1310 id,
1311 coordinate_id,
1312 &[i],
1313 ));
1314 }
1315 match size {
1316 crate::VectorSize::Bi => crate::VectorSize::Tri,
1317 crate::VectorSize::Tri => crate::VectorSize::Quad,
1318 crate::VectorSize::Quad => {
1319 return Err(Error::Validation("extending vec4 coordinate"));
1320 }
1321 }
1322 }
1323 ref other => {
1324 log::error!("wrong coordinate type {:?}", other);
1325 return Err(Error::Validation("coordinate type"));
1326 }
1327 };
1328
1329 let array_index_f32_id = self.id_gen.next();
1330 constituent_ids[size as usize - 1] = array_index_f32_id;
1331
1332 let array_index_u32_id = self.cached[array_index];
1333 let cast_instruction = Instruction::unary(
1334 spirv::Op::ConvertUToF,
1335 coordinate_scalar_type_id,
1336 array_index_f32_id,
1337 array_index_u32_id,
1338 );
1339 block.body.push(cast_instruction);
1340
1341 let extended_coordinate_type_id = self.get_type_id(
1342 &ir_module.types,
1343 LookupType::Local(LocalType::Value {
1344 vector_size: Some(size),
1345 kind: crate::ScalarKind::Float,
1346 width: 4,
1347 pointer_class: None,
1348 }),
1349 )?;
1350
1351 let id = self.id_gen.next();
1352 block.body.push(Instruction::composite_construct(
1353 extended_coordinate_type_id,
1354 id,
1355 &constituent_ids[..size as usize],
1356 ));
1357 id
1358 } else {
1359 coordinate_id
1360 })
1361 }
1362
1363 #[allow(clippy::too_many_arguments)]
promote_access_expression_to_variable( &mut self, ir_types: &Arena<crate::Type>, result_type_id: Word, container_id: Word, container_resolution: &TypeResolution, index_id: Word, element_ty: Handle<crate::Type>, block: &mut Block, ) -> Result<(Word, LocalVariable), Error>1364 fn promote_access_expression_to_variable(
1365 &mut self,
1366 ir_types: &Arena<crate::Type>,
1367 result_type_id: Word,
1368 container_id: Word,
1369 container_resolution: &TypeResolution,
1370 index_id: Word,
1371 element_ty: Handle<crate::Type>,
1372 block: &mut Block,
1373 ) -> Result<(Word, LocalVariable), Error> {
1374 let container_type_id = self.get_expression_type_id(ir_types, container_resolution)?;
1375 let pointer_type_id = self.id_gen.next();
1376 Instruction::type_pointer(
1377 pointer_type_id,
1378 spirv::StorageClass::Function,
1379 container_type_id,
1380 )
1381 .to_words(&mut self.logical_layout.declarations);
1382
1383 let variable = {
1384 let id = self.id_gen.next();
1385 LocalVariable {
1386 id,
1387 instruction: Instruction::variable(
1388 pointer_type_id,
1389 id,
1390 spirv::StorageClass::Function,
1391 None,
1392 ),
1393 }
1394 };
1395 block
1396 .body
1397 .push(Instruction::store(variable.id, container_id, None));
1398
1399 let element_pointer_id = self.id_gen.next();
1400 let element_pointer_type_id =
1401 self.get_pointer_id(ir_types, element_ty, spirv::StorageClass::Function)?;
1402 block.body.push(Instruction::access_chain(
1403 element_pointer_type_id,
1404 element_pointer_id,
1405 variable.id,
1406 &[index_id],
1407 ));
1408 let id = self.id_gen.next();
1409 block.body.push(Instruction::load(
1410 result_type_id,
1411 id,
1412 element_pointer_id,
1413 None,
1414 ));
1415
1416 Ok((id, variable))
1417 }
1418
is_intermediate( &self, expr_handle: Handle<crate::Expression>, ir_function: &crate::Function, ir_types: &Arena<crate::Type>, ) -> bool1419 fn is_intermediate(
1420 &self,
1421 expr_handle: Handle<crate::Expression>,
1422 ir_function: &crate::Function,
1423 ir_types: &Arena<crate::Type>,
1424 ) -> bool {
1425 match ir_function.expressions[expr_handle] {
1426 crate::Expression::GlobalVariable(_) | crate::Expression::LocalVariable(_) => true,
1427 crate::Expression::FunctionArgument(index) => {
1428 let arg = &ir_function.arguments[index as usize];
1429 match ir_types[arg.ty].inner {
1430 crate::TypeInner::Pointer { .. } | crate::TypeInner::ValuePointer { .. } => {
1431 true
1432 }
1433 _ => false,
1434 }
1435 }
1436 _ => self.cached.ids[expr_handle.index()] == 0,
1437 }
1438 }
1439
1440 /// Cache an expression for a value.
cache_expression_value( &mut self, ir_module: &crate::Module, ir_function: &crate::Function, fun_info: &FunctionInfo, expr_handle: Handle<crate::Expression>, block: &mut Block, function: &mut Function, ) -> Result<(), Error>1441 fn cache_expression_value(
1442 &mut self,
1443 ir_module: &crate::Module,
1444 ir_function: &crate::Function,
1445 fun_info: &FunctionInfo,
1446 expr_handle: Handle<crate::Expression>,
1447 block: &mut Block,
1448 function: &mut Function,
1449 ) -> Result<(), Error> {
1450 let result_type_id =
1451 self.get_expression_type_id(&ir_module.types, &fun_info[expr_handle].ty)?;
1452
1453 let id = match ir_function.expressions[expr_handle] {
1454 crate::Expression::Access { base, index: _ }
1455 if self.is_intermediate(base, ir_function, &ir_module.types) =>
1456 {
1457 0
1458 }
1459 crate::Expression::Access { base, index } => {
1460 let index_id = self.cached[index];
1461 let base_id = self.cached[base];
1462 match *fun_info[base].ty.inner_with(&ir_module.types) {
1463 crate::TypeInner::Vector { .. } => {
1464 let id = self.id_gen.next();
1465 block.body.push(Instruction::vector_extract_dynamic(
1466 result_type_id,
1467 id,
1468 base_id,
1469 index_id,
1470 ));
1471 id
1472 }
1473 crate::TypeInner::Array {
1474 base: ty_element, ..
1475 } => {
1476 let (id, variable) = self.promote_access_expression_to_variable(
1477 &ir_module.types,
1478 result_type_id,
1479 base_id,
1480 &fun_info[base].ty,
1481 index_id,
1482 ty_element,
1483 block,
1484 )?;
1485 function.internal_variables.push(variable);
1486 id
1487 }
1488 ref other => {
1489 log::error!(
1490 "Unable to access base {:?} of type {:?}",
1491 ir_function.expressions[base],
1492 other
1493 );
1494 return Err(Error::FeatureNotImplemented("access for type"));
1495 }
1496 }
1497 }
1498 crate::Expression::AccessIndex { base, index: _ }
1499 if self.is_intermediate(base, ir_function, &ir_module.types) =>
1500 {
1501 0
1502 }
1503 crate::Expression::AccessIndex { base, index } => {
1504 match *fun_info[base].ty.inner_with(&ir_module.types) {
1505 crate::TypeInner::Vector { .. }
1506 | crate::TypeInner::Matrix { .. }
1507 | crate::TypeInner::Array { .. }
1508 | crate::TypeInner::Struct { .. } => {
1509 let id = self.id_gen.next();
1510 let base_id = self.cached[base];
1511 block.body.push(Instruction::composite_extract(
1512 result_type_id,
1513 id,
1514 base_id,
1515 &[index],
1516 ));
1517 id
1518 }
1519 ref other => {
1520 log::error!("Unable to access index of {:?}", other);
1521 return Err(Error::FeatureNotImplemented("access index for type"));
1522 }
1523 }
1524 }
1525 crate::Expression::GlobalVariable(handle) => self.global_variables[handle.index()].id,
1526 crate::Expression::Constant(handle) => self.constant_ids[handle.index()],
1527 crate::Expression::Splat { size, value } => {
1528 let value_id = self.cached[value];
1529 self.temp_list.clear();
1530 self.temp_list.resize(size as usize, value_id);
1531
1532 let id = self.id_gen.next();
1533 block.body.push(Instruction::composite_construct(
1534 result_type_id,
1535 id,
1536 &self.temp_list,
1537 ));
1538 id
1539 }
1540 crate::Expression::Swizzle {
1541 size,
1542 vector,
1543 pattern,
1544 } => {
1545 let vector_id = self.cached[vector];
1546 self.temp_list.clear();
1547 for &sc in pattern[..size as usize].iter() {
1548 self.temp_list.push(sc as Word);
1549 }
1550 let id = self.id_gen.next();
1551 block.body.push(Instruction::vector_shuffle(
1552 result_type_id,
1553 id,
1554 vector_id,
1555 vector_id,
1556 &self.temp_list,
1557 ));
1558 id
1559 }
1560 crate::Expression::Compose {
1561 ty: _,
1562 ref components,
1563 } => {
1564 self.temp_list.clear();
1565 for &component in components {
1566 self.temp_list.push(self.cached[component]);
1567 }
1568
1569 let id = self.id_gen.next();
1570 block.body.push(Instruction::composite_construct(
1571 result_type_id,
1572 id,
1573 &self.temp_list,
1574 ));
1575 id
1576 }
1577 crate::Expression::Unary { op, expr } => {
1578 let id = self.id_gen.next();
1579 let expr_id = self.cached[expr];
1580 let expr_ty_inner = fun_info[expr].ty.inner_with(&ir_module.types);
1581
1582 let spirv_op = match op {
1583 crate::UnaryOperator::Negate => match expr_ty_inner.scalar_kind() {
1584 Some(crate::ScalarKind::Float) => spirv::Op::FNegate,
1585 Some(crate::ScalarKind::Sint) => spirv::Op::SNegate,
1586 Some(crate::ScalarKind::Bool) => spirv::Op::LogicalNot,
1587 Some(crate::ScalarKind::Uint) | None => {
1588 log::error!("Unable to negate {:?}", expr_ty_inner);
1589 return Err(Error::FeatureNotImplemented("negation"));
1590 }
1591 },
1592 crate::UnaryOperator::Not => match expr_ty_inner.scalar_kind() {
1593 Some(crate::ScalarKind::Bool) => spirv::Op::LogicalNot,
1594 _ => spirv::Op::Not,
1595 },
1596 };
1597
1598 block
1599 .body
1600 .push(Instruction::unary(spirv_op, result_type_id, id, expr_id));
1601 id
1602 }
1603 crate::Expression::Binary { op, left, right } => {
1604 let id = self.id_gen.next();
1605 let left_id = self.cached[left];
1606 let right_id = self.cached[right];
1607
1608 let left_ty_inner = fun_info[left].ty.inner_with(&ir_module.types);
1609 let right_ty_inner = fun_info[right].ty.inner_with(&ir_module.types);
1610
1611 let left_dimension = get_dimension(left_ty_inner);
1612 let right_dimension = get_dimension(right_ty_inner);
1613
1614 let mut preserve_order = true;
1615
1616 let spirv_op = match op {
1617 crate::BinaryOperator::Add => match *left_ty_inner {
1618 crate::TypeInner::Scalar { kind, .. }
1619 | crate::TypeInner::Vector { kind, .. } => match kind {
1620 crate::ScalarKind::Float => spirv::Op::FAdd,
1621 _ => spirv::Op::IAdd,
1622 },
1623 _ => unimplemented!(),
1624 },
1625 crate::BinaryOperator::Subtract => match *left_ty_inner {
1626 crate::TypeInner::Scalar { kind, .. }
1627 | crate::TypeInner::Vector { kind, .. } => match kind {
1628 crate::ScalarKind::Float => spirv::Op::FSub,
1629 _ => spirv::Op::ISub,
1630 },
1631 _ => unimplemented!(),
1632 },
1633 crate::BinaryOperator::Multiply => match (left_dimension, right_dimension) {
1634 (Dimension::Scalar, Dimension::Vector { .. }) => {
1635 preserve_order = false;
1636 spirv::Op::VectorTimesScalar
1637 }
1638 (Dimension::Vector, Dimension::Scalar { .. }) => {
1639 spirv::Op::VectorTimesScalar
1640 }
1641 (Dimension::Vector, Dimension::Matrix) => spirv::Op::VectorTimesMatrix,
1642 (Dimension::Matrix, Dimension::Scalar { .. }) => {
1643 spirv::Op::MatrixTimesScalar
1644 }
1645 (Dimension::Matrix, Dimension::Vector) => spirv::Op::MatrixTimesVector,
1646 (Dimension::Matrix, Dimension::Matrix) => spirv::Op::MatrixTimesMatrix,
1647 (Dimension::Vector, Dimension::Vector)
1648 | (Dimension::Scalar, Dimension::Scalar)
1649 if left_ty_inner.scalar_kind() == Some(crate::ScalarKind::Float) =>
1650 {
1651 spirv::Op::FMul
1652 }
1653 (Dimension::Vector, Dimension::Vector)
1654 | (Dimension::Scalar, Dimension::Scalar) => spirv::Op::IMul,
1655 other => unimplemented!("Mul {:?}", other),
1656 },
1657 crate::BinaryOperator::Divide => match left_ty_inner.scalar_kind() {
1658 Some(crate::ScalarKind::Sint) => spirv::Op::SDiv,
1659 Some(crate::ScalarKind::Uint) => spirv::Op::UDiv,
1660 Some(crate::ScalarKind::Float) => spirv::Op::FDiv,
1661 _ => unimplemented!(),
1662 },
1663 crate::BinaryOperator::Modulo => match left_ty_inner.scalar_kind() {
1664 Some(crate::ScalarKind::Sint) => spirv::Op::SMod,
1665 Some(crate::ScalarKind::Uint) => spirv::Op::UMod,
1666 Some(crate::ScalarKind::Float) => spirv::Op::FMod,
1667 _ => unimplemented!(),
1668 },
1669 crate::BinaryOperator::Equal => match left_ty_inner.scalar_kind() {
1670 Some(crate::ScalarKind::Sint) | Some(crate::ScalarKind::Uint) => {
1671 spirv::Op::IEqual
1672 }
1673 Some(crate::ScalarKind::Float) => spirv::Op::FOrdEqual,
1674 Some(crate::ScalarKind::Bool) => spirv::Op::LogicalEqual,
1675 _ => unimplemented!(),
1676 },
1677 crate::BinaryOperator::NotEqual => match left_ty_inner.scalar_kind() {
1678 Some(crate::ScalarKind::Sint) | Some(crate::ScalarKind::Uint) => {
1679 spirv::Op::INotEqual
1680 }
1681 Some(crate::ScalarKind::Float) => spirv::Op::FOrdNotEqual,
1682 Some(crate::ScalarKind::Bool) => spirv::Op::LogicalNotEqual,
1683 _ => unimplemented!(),
1684 },
1685 crate::BinaryOperator::Less => match left_ty_inner.scalar_kind() {
1686 Some(crate::ScalarKind::Sint) => spirv::Op::SLessThan,
1687 Some(crate::ScalarKind::Uint) => spirv::Op::ULessThan,
1688 Some(crate::ScalarKind::Float) => spirv::Op::FOrdLessThan,
1689 _ => unimplemented!(),
1690 },
1691 crate::BinaryOperator::LessEqual => match left_ty_inner.scalar_kind() {
1692 Some(crate::ScalarKind::Sint) => spirv::Op::SLessThanEqual,
1693 Some(crate::ScalarKind::Uint) => spirv::Op::ULessThanEqual,
1694 Some(crate::ScalarKind::Float) => spirv::Op::FOrdLessThanEqual,
1695 _ => unimplemented!(),
1696 },
1697 crate::BinaryOperator::Greater => match left_ty_inner.scalar_kind() {
1698 Some(crate::ScalarKind::Sint) => spirv::Op::SGreaterThan,
1699 Some(crate::ScalarKind::Uint) => spirv::Op::UGreaterThan,
1700 Some(crate::ScalarKind::Float) => spirv::Op::FOrdGreaterThan,
1701 _ => unimplemented!(),
1702 },
1703 crate::BinaryOperator::GreaterEqual => match left_ty_inner.scalar_kind() {
1704 Some(crate::ScalarKind::Sint) => spirv::Op::SGreaterThanEqual,
1705 Some(crate::ScalarKind::Uint) => spirv::Op::UGreaterThanEqual,
1706 Some(crate::ScalarKind::Float) => spirv::Op::FOrdGreaterThanEqual,
1707 _ => unimplemented!(),
1708 },
1709 crate::BinaryOperator::And => spirv::Op::BitwiseAnd,
1710 crate::BinaryOperator::ExclusiveOr => spirv::Op::BitwiseXor,
1711 crate::BinaryOperator::InclusiveOr => spirv::Op::BitwiseOr,
1712 crate::BinaryOperator::LogicalAnd => spirv::Op::LogicalAnd,
1713 crate::BinaryOperator::LogicalOr => spirv::Op::LogicalOr,
1714 crate::BinaryOperator::ShiftLeft => spirv::Op::ShiftLeftLogical,
1715 crate::BinaryOperator::ShiftRight => match left_ty_inner.scalar_kind() {
1716 Some(crate::ScalarKind::Sint) => spirv::Op::ShiftRightArithmetic,
1717 Some(crate::ScalarKind::Uint) => spirv::Op::ShiftRightLogical,
1718 _ => unimplemented!(),
1719 },
1720 };
1721
1722 block.body.push(Instruction::binary(
1723 spirv_op,
1724 result_type_id,
1725 id,
1726 if preserve_order { left_id } else { right_id },
1727 if preserve_order { right_id } else { left_id },
1728 ));
1729 id
1730 }
1731 crate::Expression::Math {
1732 fun,
1733 arg,
1734 arg1,
1735 arg2,
1736 } => {
1737 use crate::MathFunction as Mf;
1738 enum MathOp {
1739 Ext(spirv::GLOp),
1740 Custom(Instruction),
1741 }
1742
1743 let arg0_id = self.cached[arg];
1744 let arg_scalar_kind = fun_info[arg].ty.inner_with(&ir_module.types).scalar_kind();
1745 let arg1_id = match arg1 {
1746 Some(handle) => self.cached[handle],
1747 None => 0,
1748 };
1749 let arg2_id = match arg2 {
1750 Some(handle) => self.cached[handle],
1751 None => 0,
1752 };
1753
1754 let id = self.id_gen.next();
1755 let math_op = match fun {
1756 // comparison
1757 Mf::Abs => {
1758 match arg_scalar_kind {
1759 Some(crate::ScalarKind::Float) => MathOp::Ext(spirv::GLOp::FAbs),
1760 Some(crate::ScalarKind::Sint) => MathOp::Ext(spirv::GLOp::SAbs),
1761 Some(crate::ScalarKind::Uint) => {
1762 MathOp::Custom(Instruction::unary(
1763 spirv::Op::CopyObject, // do nothing
1764 result_type_id,
1765 id,
1766 arg0_id,
1767 ))
1768 }
1769 other => unimplemented!("Unexpected abs({:?})", other),
1770 }
1771 }
1772 Mf::Min => MathOp::Ext(match arg_scalar_kind {
1773 Some(crate::ScalarKind::Float) => spirv::GLOp::FMin,
1774 Some(crate::ScalarKind::Sint) => spirv::GLOp::SMin,
1775 Some(crate::ScalarKind::Uint) => spirv::GLOp::UMin,
1776 other => unimplemented!("Unexpected min({:?})", other),
1777 }),
1778 Mf::Max => MathOp::Ext(match arg_scalar_kind {
1779 Some(crate::ScalarKind::Float) => spirv::GLOp::FMax,
1780 Some(crate::ScalarKind::Sint) => spirv::GLOp::SMax,
1781 Some(crate::ScalarKind::Uint) => spirv::GLOp::UMax,
1782 other => unimplemented!("Unexpected max({:?})", other),
1783 }),
1784 Mf::Clamp => MathOp::Ext(match arg_scalar_kind {
1785 Some(crate::ScalarKind::Float) => spirv::GLOp::FClamp,
1786 Some(crate::ScalarKind::Sint) => spirv::GLOp::SClamp,
1787 Some(crate::ScalarKind::Uint) => spirv::GLOp::UClamp,
1788 other => unimplemented!("Unexpected max({:?})", other),
1789 }),
1790 // trigonometry
1791 Mf::Sin => MathOp::Ext(spirv::GLOp::Sin),
1792 Mf::Sinh => MathOp::Ext(spirv::GLOp::Sinh),
1793 Mf::Asin => MathOp::Ext(spirv::GLOp::Asin),
1794 Mf::Cos => MathOp::Ext(spirv::GLOp::Cos),
1795 Mf::Cosh => MathOp::Ext(spirv::GLOp::Cosh),
1796 Mf::Acos => MathOp::Ext(spirv::GLOp::Acos),
1797 Mf::Tan => MathOp::Ext(spirv::GLOp::Tan),
1798 Mf::Tanh => MathOp::Ext(spirv::GLOp::Tanh),
1799 Mf::Atan => MathOp::Ext(spirv::GLOp::Atan),
1800 Mf::Atan2 => MathOp::Ext(spirv::GLOp::Atan2),
1801 // decomposition
1802 Mf::Ceil => MathOp::Ext(spirv::GLOp::Ceil),
1803 Mf::Round => MathOp::Ext(spirv::GLOp::Round),
1804 Mf::Floor => MathOp::Ext(spirv::GLOp::Floor),
1805 Mf::Fract => MathOp::Ext(spirv::GLOp::Fract),
1806 Mf::Trunc => MathOp::Ext(spirv::GLOp::Trunc),
1807 Mf::Modf => MathOp::Ext(spirv::GLOp::Modf),
1808 Mf::Frexp => MathOp::Ext(spirv::GLOp::Frexp),
1809 Mf::Ldexp => MathOp::Ext(spirv::GLOp::Ldexp),
1810 // geometry
1811 Mf::Dot => MathOp::Custom(Instruction::binary(
1812 spirv::Op::Dot,
1813 result_type_id,
1814 id,
1815 arg0_id,
1816 arg1_id,
1817 )),
1818 Mf::Outer => MathOp::Custom(Instruction::binary(
1819 spirv::Op::OuterProduct,
1820 result_type_id,
1821 id,
1822 arg0_id,
1823 arg1_id,
1824 )),
1825 Mf::Cross => MathOp::Ext(spirv::GLOp::Cross),
1826 Mf::Distance => MathOp::Ext(spirv::GLOp::Distance),
1827 Mf::Length => MathOp::Ext(spirv::GLOp::Length),
1828 Mf::Normalize => MathOp::Ext(spirv::GLOp::Normalize),
1829 Mf::FaceForward => MathOp::Ext(spirv::GLOp::FaceForward),
1830 Mf::Reflect => MathOp::Ext(spirv::GLOp::Reflect),
1831 Mf::Refract => MathOp::Ext(spirv::GLOp::Refract),
1832 // exponent
1833 Mf::Exp => MathOp::Ext(spirv::GLOp::Exp),
1834 Mf::Exp2 => MathOp::Ext(spirv::GLOp::Exp2),
1835 Mf::Log => MathOp::Ext(spirv::GLOp::Log),
1836 Mf::Log2 => MathOp::Ext(spirv::GLOp::Log2),
1837 Mf::Pow => MathOp::Ext(spirv::GLOp::Pow),
1838 // computational
1839 Mf::Sign => MathOp::Ext(match arg_scalar_kind {
1840 Some(crate::ScalarKind::Float) => spirv::GLOp::FSign,
1841 Some(crate::ScalarKind::Sint) => spirv::GLOp::SSign,
1842 other => unimplemented!("Unexpected sign({:?})", other),
1843 }),
1844 Mf::Fma => MathOp::Ext(spirv::GLOp::Fma),
1845 Mf::Mix => MathOp::Ext(spirv::GLOp::FMix),
1846 Mf::Step => MathOp::Ext(spirv::GLOp::Step),
1847 Mf::SmoothStep => MathOp::Ext(spirv::GLOp::SmoothStep),
1848 Mf::Sqrt => MathOp::Ext(spirv::GLOp::Sqrt),
1849 Mf::InverseSqrt => MathOp::Ext(spirv::GLOp::InverseSqrt),
1850 Mf::Inverse => MathOp::Ext(spirv::GLOp::MatrixInverse),
1851 Mf::Transpose => MathOp::Custom(Instruction::unary(
1852 spirv::Op::Transpose,
1853 result_type_id,
1854 id,
1855 arg0_id,
1856 )),
1857 Mf::Determinant => MathOp::Ext(spirv::GLOp::Determinant),
1858 Mf::ReverseBits | Mf::CountOneBits => {
1859 log::error!("unimplemented math function {:?}", fun);
1860 return Err(Error::FeatureNotImplemented("math function"));
1861 }
1862 };
1863
1864 block.body.push(match math_op {
1865 MathOp::Ext(op) => Instruction::ext_inst(
1866 self.gl450_ext_inst_id,
1867 op,
1868 result_type_id,
1869 id,
1870 &[arg0_id, arg1_id, arg2_id][..fun.argument_count()],
1871 ),
1872 MathOp::Custom(inst) => inst,
1873 });
1874 id
1875 }
1876 crate::Expression::LocalVariable(variable) => function.variables[&variable].id,
1877 crate::Expression::Load { pointer } => {
1878 let (pointer_id, _) = self.write_expression_pointer(
1879 ir_module,
1880 ir_function,
1881 fun_info,
1882 pointer,
1883 block,
1884 function,
1885 )?;
1886
1887 let id = self.id_gen.next();
1888 block
1889 .body
1890 .push(Instruction::load(result_type_id, id, pointer_id, None));
1891 id
1892 }
1893 crate::Expression::FunctionArgument(index) => match function.entry_point_context {
1894 Some(ref context) => context.argument_ids[index as usize],
1895 None => function.parameters[index as usize].result_id.unwrap(),
1896 },
1897 crate::Expression::Call(_function) => self.lookup_function_call[&expr_handle],
1898 crate::Expression::As {
1899 expr,
1900 kind,
1901 convert,
1902 } => {
1903 let expr_id = self.cached[expr];
1904 let expr_kind = fun_info[expr]
1905 .ty
1906 .inner_with(&ir_module.types)
1907 .scalar_kind()
1908 .unwrap();
1909
1910 let op = match (expr_kind, kind) {
1911 _ if convert.is_none() => spirv::Op::Bitcast,
1912 (crate::ScalarKind::Float, crate::ScalarKind::Uint) => spirv::Op::ConvertFToU,
1913 (crate::ScalarKind::Float, crate::ScalarKind::Sint) => spirv::Op::ConvertFToS,
1914 (crate::ScalarKind::Float, crate::ScalarKind::Float) => spirv::Op::FConvert,
1915 (crate::ScalarKind::Sint, crate::ScalarKind::Float) => spirv::Op::ConvertSToF,
1916 (crate::ScalarKind::Sint, crate::ScalarKind::Sint) => spirv::Op::SConvert,
1917 (crate::ScalarKind::Uint, crate::ScalarKind::Float) => spirv::Op::ConvertUToF,
1918 (crate::ScalarKind::Uint, crate::ScalarKind::Uint) => spirv::Op::UConvert,
1919 // We assume it's either an identity cast, or int-uint.
1920 _ => spirv::Op::Bitcast,
1921 };
1922
1923 let id = self.id_gen.next();
1924 let instruction = Instruction::unary(op, result_type_id, id, expr_id);
1925 block.body.push(instruction);
1926 id
1927 }
1928 crate::Expression::ImageLoad {
1929 image,
1930 coordinate,
1931 array_index,
1932 index,
1933 } => {
1934 let image_id = self.get_expression_global(ir_function, image);
1935 let coordinate_id = self.write_texture_coordinates(
1936 ir_module,
1937 fun_info,
1938 coordinate,
1939 array_index,
1940 block,
1941 )?;
1942
1943 let id = self.id_gen.next();
1944
1945 let image_ty = fun_info[image].ty.inner_with(&ir_module.types);
1946 let mut instruction = match *image_ty {
1947 crate::TypeInner::Image {
1948 class: crate::ImageClass::Storage { .. },
1949 ..
1950 } => Instruction::image_read(result_type_id, id, image_id, coordinate_id),
1951 crate::TypeInner::Image {
1952 class: crate::ImageClass::Depth,
1953 ..
1954 } => {
1955 // Vulkan doesn't know about our `Depth` class, and it returns `vec4<f32>`,
1956 // so we need to grab the first component out of it.
1957 let load_result_type_id = self.get_type_id(
1958 &ir_module.types,
1959 LookupType::Local(LocalType::Value {
1960 vector_size: Some(crate::VectorSize::Quad),
1961 kind: crate::ScalarKind::Float,
1962 width: 4,
1963 pointer_class: None,
1964 }),
1965 )?;
1966 Instruction::image_fetch(load_result_type_id, id, image_id, coordinate_id)
1967 }
1968 _ => Instruction::image_fetch(result_type_id, id, image_id, coordinate_id),
1969 };
1970
1971 if let Some(index) = index {
1972 let index_id = self.cached[index];
1973 let image_ops = match *fun_info[image].ty.inner_with(&ir_module.types) {
1974 crate::TypeInner::Image {
1975 class: crate::ImageClass::Sampled { multi: true, .. },
1976 ..
1977 } => spirv::ImageOperands::SAMPLE,
1978 _ => spirv::ImageOperands::LOD,
1979 };
1980 instruction.add_operand(image_ops.bits());
1981 instruction.add_operand(index_id);
1982 }
1983
1984 let inst_type_id = instruction.type_id;
1985 block.body.push(instruction);
1986 if inst_type_id != Some(result_type_id) {
1987 let sub_id = self.id_gen.next();
1988 block.body.push(Instruction::composite_extract(
1989 result_type_id,
1990 sub_id,
1991 id,
1992 &[0],
1993 ));
1994 sub_id
1995 } else {
1996 id
1997 }
1998 }
1999 crate::Expression::ImageSample {
2000 image,
2001 sampler,
2002 coordinate,
2003 array_index,
2004 offset,
2005 level,
2006 depth_ref,
2007 } => {
2008 use super::instructions::SampleLod;
2009 // image
2010 let image_id = self.get_expression_global(ir_function, image);
2011 let image_type = fun_info[image].ty.handle().unwrap();
2012 // Vulkan doesn't know about our `Depth` class, and it returns `vec4<f32>`,
2013 // so we need to grab the first component out of it.
2014 let needs_sub_access = match ir_module.types[image_type].inner {
2015 crate::TypeInner::Image {
2016 class: crate::ImageClass::Depth,
2017 ..
2018 } => depth_ref.is_none(),
2019 _ => false,
2020 };
2021 let sample_result_type_id = if needs_sub_access {
2022 self.get_type_id(
2023 &ir_module.types,
2024 LookupType::Local(LocalType::Value {
2025 vector_size: Some(crate::VectorSize::Quad),
2026 kind: crate::ScalarKind::Float,
2027 width: 4,
2028 pointer_class: None,
2029 }),
2030 )?
2031 } else {
2032 result_type_id
2033 };
2034
2035 // OpTypeSampledImage
2036 let image_type_id =
2037 self.get_type_id(&ir_module.types, LookupType::Handle(image_type))?;
2038 let sampled_image_type_id = self.get_type_id(
2039 &ir_module.types,
2040 LookupType::Local(LocalType::SampledImage { image_type_id }),
2041 )?;
2042
2043 let sampler_id = self.get_expression_global(ir_function, sampler);
2044 let coordinate_id = self.write_texture_coordinates(
2045 ir_module,
2046 fun_info,
2047 coordinate,
2048 array_index,
2049 block,
2050 )?;
2051
2052 let sampled_image_id = self.id_gen.next();
2053 block.body.push(Instruction::sampled_image(
2054 sampled_image_type_id,
2055 sampled_image_id,
2056 image_id,
2057 sampler_id,
2058 ));
2059 let id = self.id_gen.next();
2060
2061 let depth_id = depth_ref.map(|handle| self.cached[handle]);
2062 let mut mask = spirv::ImageOperands::empty();
2063 mask.set(spirv::ImageOperands::CONST_OFFSET, offset.is_some());
2064
2065 let mut main_instruction = match level {
2066 crate::SampleLevel::Zero => {
2067 let mut inst = Instruction::image_sample(
2068 sample_result_type_id,
2069 id,
2070 SampleLod::Explicit,
2071 sampled_image_id,
2072 coordinate_id,
2073 depth_id,
2074 );
2075
2076 let zero_id = self.get_constant_scalar(
2077 crate::ScalarValue::Float(0.0),
2078 4,
2079 &ir_module.types,
2080 )?;
2081
2082 mask |= spirv::ImageOperands::LOD;
2083 inst.add_operand(mask.bits());
2084 inst.add_operand(zero_id);
2085
2086 inst
2087 }
2088 crate::SampleLevel::Auto => Instruction::image_sample(
2089 sample_result_type_id,
2090 id,
2091 SampleLod::Implicit,
2092 sampled_image_id,
2093 coordinate_id,
2094 depth_id,
2095 ),
2096 crate::SampleLevel::Exact(lod_handle) => {
2097 let mut inst = Instruction::image_sample(
2098 sample_result_type_id,
2099 id,
2100 SampleLod::Explicit,
2101 sampled_image_id,
2102 coordinate_id,
2103 depth_id,
2104 );
2105
2106 let lod_id = self.cached[lod_handle];
2107 mask |= spirv::ImageOperands::LOD;
2108 inst.add_operand(mask.bits());
2109 inst.add_operand(lod_id);
2110
2111 inst
2112 }
2113 crate::SampleLevel::Bias(bias_handle) => {
2114 let mut inst = Instruction::image_sample(
2115 sample_result_type_id,
2116 id,
2117 SampleLod::Implicit,
2118 sampled_image_id,
2119 coordinate_id,
2120 depth_id,
2121 );
2122
2123 let bias_id = self.cached[bias_handle];
2124 mask |= spirv::ImageOperands::BIAS;
2125 inst.add_operand(bias_id);
2126
2127 inst
2128 }
2129 crate::SampleLevel::Gradient { x, y } => {
2130 let mut inst = Instruction::image_sample(
2131 sample_result_type_id,
2132 id,
2133 SampleLod::Explicit,
2134 sampled_image_id,
2135 coordinate_id,
2136 depth_id,
2137 );
2138
2139 let x_id = self.cached[x];
2140 let y_id = self.cached[y];
2141 mask |= spirv::ImageOperands::GRAD;
2142 inst.add_operand(x_id);
2143 inst.add_operand(y_id);
2144
2145 inst
2146 }
2147 };
2148
2149 if let Some(offset_const) = offset {
2150 let offset_id = self.constant_ids[offset_const.index()];
2151 main_instruction.add_operand(offset_id);
2152 }
2153
2154 block.body.push(main_instruction);
2155
2156 if needs_sub_access {
2157 let sub_id = self.id_gen.next();
2158 block.body.push(Instruction::composite_extract(
2159 result_type_id,
2160 sub_id,
2161 id,
2162 &[0],
2163 ));
2164 sub_id
2165 } else {
2166 id
2167 }
2168 }
2169 crate::Expression::Select {
2170 condition,
2171 accept,
2172 reject,
2173 } => {
2174 let id = self.id_gen.next();
2175 let condition_id = self.cached[condition];
2176 let accept_id = self.cached[accept];
2177 let reject_id = self.cached[reject];
2178
2179 let instruction =
2180 Instruction::select(result_type_id, id, condition_id, accept_id, reject_id);
2181 block.body.push(instruction);
2182 id
2183 }
2184 crate::Expression::Derivative { axis, expr } => {
2185 use crate::DerivativeAxis as Da;
2186
2187 let id = self.id_gen.next();
2188 let expr_id = self.cached[expr];
2189 let op = match axis {
2190 Da::X => spirv::Op::DPdx,
2191 Da::Y => spirv::Op::DPdy,
2192 Da::Width => spirv::Op::Fwidth,
2193 };
2194 block
2195 .body
2196 .push(Instruction::derivative(op, result_type_id, id, expr_id));
2197 id
2198 }
2199 crate::Expression::ImageQuery { image, query } => {
2200 use crate::{ImageClass as Ic, ImageDimension as Id, ImageQuery as Iq};
2201
2202 let image_id = self.get_expression_global(ir_function, image);
2203 let image_type = fun_info[image].ty.handle().unwrap();
2204 let (dim, arrayed, class) = match ir_module.types[image_type].inner {
2205 crate::TypeInner::Image {
2206 dim,
2207 arrayed,
2208 class,
2209 } => (dim, arrayed, class),
2210 _ => {
2211 return Err(Error::Validation("image type"));
2212 }
2213 };
2214
2215 match query {
2216 Iq::Size { level } => {
2217 let dim_coords = match dim {
2218 Id::D1 => 1,
2219 Id::D2 | Id::Cube => 2,
2220 Id::D3 => 3,
2221 };
2222 let extended_size_type_id = {
2223 let array_coords = if arrayed { 1 } else { 0 };
2224 let vector_size = match dim_coords + array_coords {
2225 2 => Some(crate::VectorSize::Bi),
2226 3 => Some(crate::VectorSize::Tri),
2227 4 => Some(crate::VectorSize::Quad),
2228 _ => None,
2229 };
2230 self.get_type_id(
2231 &ir_module.types,
2232 LookupType::Local(LocalType::Value {
2233 vector_size,
2234 kind: crate::ScalarKind::Sint,
2235 width: 4,
2236 pointer_class: None,
2237 }),
2238 )?
2239 };
2240
2241 let (query_op, level_id) = match class {
2242 Ic::Storage(_) => (spirv::Op::ImageQuerySize, None),
2243 _ => {
2244 let level_id = match level {
2245 Some(expr) => self.cached[expr],
2246 None => self.get_index_constant(0, &ir_module.types)?,
2247 };
2248 (spirv::Op::ImageQuerySizeLod, Some(level_id))
2249 }
2250 };
2251 // The ID of the vector returned by SPIR-V, which contains the dimensions
2252 // as well as the layer count.
2253 let id_extended = self.id_gen.next();
2254 let mut inst = Instruction::image_query(
2255 query_op,
2256 extended_size_type_id,
2257 id_extended,
2258 image_id,
2259 );
2260 if let Some(expr_id) = level_id {
2261 inst.add_operand(expr_id);
2262 }
2263 block.body.push(inst);
2264
2265 if result_type_id != extended_size_type_id {
2266 let id = self.id_gen.next();
2267 let components = match dim {
2268 // always pick the first component, and duplicate it for all 3 dimensions
2269 Id::Cube => &[0u32, 0, 0][..],
2270 _ => &[0u32, 1, 2, 3][..dim_coords],
2271 };
2272 block.body.push(Instruction::vector_shuffle(
2273 result_type_id,
2274 id,
2275 id_extended,
2276 id_extended,
2277 components,
2278 ));
2279 id
2280 } else {
2281 id_extended
2282 }
2283 }
2284 Iq::NumLevels => {
2285 let id = self.id_gen.next();
2286 block.body.push(Instruction::image_query(
2287 spirv::Op::ImageQueryLevels,
2288 result_type_id,
2289 id,
2290 image_id,
2291 ));
2292 id
2293 }
2294 Iq::NumLayers => {
2295 let vec_size = match dim {
2296 Id::D1 => crate::VectorSize::Bi,
2297 Id::D2 | Id::Cube => crate::VectorSize::Tri,
2298 Id::D3 => crate::VectorSize::Quad,
2299 };
2300 let extended_size_type_id = self.get_type_id(
2301 &ir_module.types,
2302 LookupType::Local(LocalType::Value {
2303 vector_size: Some(vec_size),
2304 kind: crate::ScalarKind::Sint,
2305 width: 4,
2306 pointer_class: None,
2307 }),
2308 )?;
2309 let id_extended = self.id_gen.next();
2310 let mut inst = Instruction::image_query(
2311 spirv::Op::ImageQuerySizeLod,
2312 extended_size_type_id,
2313 id_extended,
2314 image_id,
2315 );
2316 inst.add_operand(self.get_index_constant(0, &ir_module.types)?);
2317 block.body.push(inst);
2318 let id = self.id_gen.next();
2319 block.body.push(Instruction::composite_extract(
2320 result_type_id,
2321 id,
2322 id_extended,
2323 &[vec_size as u32 - 1],
2324 ));
2325 id
2326 }
2327 Iq::NumSamples => {
2328 let id = self.id_gen.next();
2329 block.body.push(Instruction::image_query(
2330 spirv::Op::ImageQuerySamples,
2331 result_type_id,
2332 id,
2333 image_id,
2334 ));
2335 id
2336 }
2337 }
2338 }
2339 crate::Expression::Relational { fun, argument } => {
2340 use crate::RelationalFunction as Rf;
2341 let arg_id = self.cached[argument];
2342 let op = match fun {
2343 Rf::All => spirv::Op::All,
2344 Rf::Any => spirv::Op::Any,
2345 Rf::IsNan => spirv::Op::IsNan,
2346 Rf::IsInf => spirv::Op::IsInf,
2347 //TODO: these require Kernel capability
2348 Rf::IsFinite | Rf::IsNormal => {
2349 return Err(Error::FeatureNotImplemented("is finite/normal"))
2350 }
2351 };
2352 let id = self.id_gen.next();
2353 block
2354 .body
2355 .push(Instruction::relational(op, result_type_id, id, arg_id));
2356 id
2357 }
2358 crate::Expression::ArrayLength(expr) => {
2359 let (structure_id, member_idx) = match ir_function.expressions[expr] {
2360 crate::Expression::AccessIndex { base, .. } => {
2361 match ir_function.expressions[base] {
2362 crate::Expression::GlobalVariable(handle) => {
2363 let global = &ir_module.global_variables[handle];
2364 let last_idx = match ir_module.types[global.ty].inner {
2365 crate::TypeInner::Struct { ref members, .. } => {
2366 members.len() as u32 - 1
2367 }
2368 _ => return Err(Error::Validation("array length expression")),
2369 };
2370
2371 (self.global_variables[handle.index()].id, last_idx)
2372 }
2373 _ => return Err(Error::Validation("array length expression")),
2374 }
2375 }
2376 _ => return Err(Error::Validation("array length expression")),
2377 };
2378
2379 // let structure_id = self.get_expression_global(ir_function, global);
2380 let id = self.id_gen.next();
2381
2382 block.body.push(Instruction::array_length(
2383 result_type_id,
2384 id,
2385 structure_id,
2386 member_idx,
2387 ));
2388 id
2389 }
2390 };
2391
2392 self.cached[expr_handle] = id;
2393 Ok(())
2394 }
2395
2396 /// Write a left-hand-side expression, returning an `id` of the pointer.
write_expression_pointer<'a>( &mut self, ir_module: &'a crate::Module, ir_function: &crate::Function, fun_info: &FunctionInfo, mut expr_handle: Handle<crate::Expression>, block: &mut Block, function: &mut Function, ) -> Result<(Word, spirv::StorageClass), Error>2397 fn write_expression_pointer<'a>(
2398 &mut self,
2399 ir_module: &'a crate::Module,
2400 ir_function: &crate::Function,
2401 fun_info: &FunctionInfo,
2402 mut expr_handle: Handle<crate::Expression>,
2403 block: &mut Block,
2404 function: &mut Function,
2405 ) -> Result<(Word, spirv::StorageClass), Error> {
2406 let result_lookup_ty = match fun_info[expr_handle].ty {
2407 TypeResolution::Handle(ty_handle) => LookupType::Handle(ty_handle),
2408 TypeResolution::Value(ref inner) => {
2409 LookupType::Local(self.physical_layout.make_local(inner).unwrap())
2410 }
2411 };
2412 let result_type_id = self.get_type_id(&ir_module.types, result_lookup_ty)?;
2413
2414 self.temp_list.clear();
2415 let (root_id, class) = loop {
2416 expr_handle = match ir_function.expressions[expr_handle] {
2417 crate::Expression::Access { base, index } => {
2418 let index_id = self.cached[index];
2419 self.temp_list.push(index_id);
2420 base
2421 }
2422 crate::Expression::AccessIndex { base, index } => {
2423 let const_id = self.get_index_constant(index, &ir_module.types)?;
2424 self.temp_list.push(const_id);
2425 base
2426 }
2427 crate::Expression::GlobalVariable(handle) => {
2428 let gv = &self.global_variables[handle.index()];
2429 break (gv.id, gv.class);
2430 }
2431 crate::Expression::LocalVariable(variable) => {
2432 let local_var = &function.variables[&variable];
2433 break (local_var.id, spirv::StorageClass::Function);
2434 }
2435 crate::Expression::FunctionArgument(index) => {
2436 let id = function.parameters[index as usize].result_id.unwrap();
2437 break (id, spirv::StorageClass::Function);
2438 }
2439 ref other => unimplemented!("Unexpected pointer expression {:?}", other),
2440 }
2441 };
2442
2443 let id = if self.temp_list.is_empty() {
2444 root_id
2445 } else {
2446 self.temp_list.reverse();
2447 let id = self.id_gen.next();
2448 block.body.push(Instruction::access_chain(
2449 result_type_id,
2450 id,
2451 root_id,
2452 &self.temp_list,
2453 ));
2454 id
2455 };
2456 Ok((id, class))
2457 }
2458
get_expression_global( &self, ir_function: &crate::Function, expr_handle: Handle<crate::Expression>, ) -> Word2459 fn get_expression_global(
2460 &self,
2461 ir_function: &crate::Function,
2462 expr_handle: Handle<crate::Expression>,
2463 ) -> Word {
2464 match ir_function.expressions[expr_handle] {
2465 crate::Expression::GlobalVariable(handle) => {
2466 let id = self.global_variables[handle.index()].handle_id;
2467 if id == 0 {
2468 unreachable!("Global variable {:?} doesn't have a handle ID", handle);
2469 }
2470 id
2471 }
2472 ref other => unreachable!("Unexpected global expression {:?}", other),
2473 }
2474 }
2475
write_entry_point_return( &mut self, value_id: Word, ir_result: &crate::FunctionResult, type_arena: &Arena<crate::Type>, result_members: &[ResultMember], body: &mut Vec<Instruction>, ) -> Result<(), Error>2476 fn write_entry_point_return(
2477 &mut self,
2478 value_id: Word,
2479 ir_result: &crate::FunctionResult,
2480 type_arena: &Arena<crate::Type>,
2481 result_members: &[ResultMember],
2482 body: &mut Vec<Instruction>,
2483 ) -> Result<(), Error> {
2484 for (index, res_member) in result_members.iter().enumerate() {
2485 let member_value_id = match ir_result.binding {
2486 Some(_) => value_id,
2487 None => {
2488 let member_value_id = self.id_gen.next();
2489 body.push(Instruction::composite_extract(
2490 res_member.type_id,
2491 member_value_id,
2492 value_id,
2493 &[index as u32],
2494 ));
2495 member_value_id
2496 }
2497 };
2498
2499 body.push(Instruction::store(res_member.id, member_value_id, None));
2500
2501 // Flip Y coordinate to adjust for coordinate space difference
2502 // between SPIR-V and our IR.
2503 if self.flags.contains(WriterFlags::ADJUST_COORDINATE_SPACE)
2504 && res_member.built_in == Some(crate::BuiltIn::Position)
2505 {
2506 let access_id = self.id_gen.next();
2507 let float_ptr_type_id = self.get_type_id(
2508 type_arena,
2509 LookupType::Local(LocalType::Value {
2510 vector_size: None,
2511 kind: crate::ScalarKind::Float,
2512 width: 4,
2513 pointer_class: Some(spirv::StorageClass::Output),
2514 }),
2515 )?;
2516 let index_y_id = self.get_index_constant(1, type_arena)?;
2517 body.push(Instruction::access_chain(
2518 float_ptr_type_id,
2519 access_id,
2520 res_member.id,
2521 &[index_y_id],
2522 ));
2523
2524 let load_id = self.id_gen.next();
2525 let float_type_id = self.get_type_id(
2526 type_arena,
2527 LookupType::Local(LocalType::Value {
2528 vector_size: None,
2529 kind: crate::ScalarKind::Float,
2530 width: 4,
2531 pointer_class: None,
2532 }),
2533 )?;
2534 body.push(Instruction::load(float_type_id, load_id, access_id, None));
2535
2536 let neg_id = self.id_gen.next();
2537 body.push(Instruction::unary(
2538 spirv::Op::FNegate,
2539 float_type_id,
2540 neg_id,
2541 load_id,
2542 ));
2543 body.push(Instruction::store(access_id, neg_id, None));
2544 }
2545 }
2546 Ok(())
2547 }
2548
2549 //TODO: put most of these into a `BlockContext` structure!
2550 #[allow(clippy::too_many_arguments)]
write_block( &mut self, label_id: Word, statements: &[crate::Statement], ir_module: &crate::Module, ir_function: &crate::Function, fun_info: &FunctionInfo, function: &mut Function, exit_id: Option<Word>, loop_context: LoopContext, ) -> Result<(), Error>2551 fn write_block(
2552 &mut self,
2553 label_id: Word,
2554 statements: &[crate::Statement],
2555 ir_module: &crate::Module,
2556 ir_function: &crate::Function,
2557 fun_info: &FunctionInfo,
2558 function: &mut Function,
2559 exit_id: Option<Word>,
2560 loop_context: LoopContext,
2561 ) -> Result<(), Error> {
2562 let mut block = Block::new(label_id);
2563
2564 for statement in statements {
2565 if block.termination.is_some() {
2566 unimplemented!("No statements are expected after block termination");
2567 }
2568 match *statement {
2569 crate::Statement::Emit(ref range) => {
2570 for handle in range.clone() {
2571 self.cache_expression_value(
2572 ir_module,
2573 ir_function,
2574 fun_info,
2575 handle,
2576 &mut block,
2577 function,
2578 )?;
2579 }
2580 }
2581 crate::Statement::Block(ref block_statements) => {
2582 let scope_id = self.id_gen.next();
2583 function.consume(block, Instruction::branch(scope_id));
2584
2585 let merge_id = self.id_gen.next();
2586 self.write_block(
2587 scope_id,
2588 block_statements,
2589 ir_module,
2590 ir_function,
2591 fun_info,
2592 function,
2593 Some(merge_id),
2594 loop_context,
2595 )?;
2596
2597 block = Block::new(merge_id);
2598 }
2599 crate::Statement::If {
2600 condition,
2601 ref accept,
2602 ref reject,
2603 } => {
2604 let condition_id = self.cached[condition];
2605
2606 let merge_id = self.id_gen.next();
2607 block.body.push(Instruction::selection_merge(
2608 merge_id,
2609 spirv::SelectionControl::NONE,
2610 ));
2611
2612 let accept_id = if accept.is_empty() {
2613 None
2614 } else {
2615 Some(self.id_gen.next())
2616 };
2617 let reject_id = if reject.is_empty() {
2618 None
2619 } else {
2620 Some(self.id_gen.next())
2621 };
2622
2623 function.consume(
2624 block,
2625 Instruction::branch_conditional(
2626 condition_id,
2627 accept_id.unwrap_or(merge_id),
2628 reject_id.unwrap_or(merge_id),
2629 ),
2630 );
2631
2632 if let Some(block_id) = accept_id {
2633 self.write_block(
2634 block_id,
2635 accept,
2636 ir_module,
2637 ir_function,
2638 fun_info,
2639 function,
2640 Some(merge_id),
2641 loop_context,
2642 )?;
2643 }
2644 if let Some(block_id) = reject_id {
2645 self.write_block(
2646 block_id,
2647 reject,
2648 ir_module,
2649 ir_function,
2650 fun_info,
2651 function,
2652 Some(merge_id),
2653 loop_context,
2654 )?;
2655 }
2656
2657 block = Block::new(merge_id);
2658 }
2659 crate::Statement::Switch {
2660 selector,
2661 ref cases,
2662 ref default,
2663 } => {
2664 let selector_id = self.cached[selector];
2665
2666 let merge_id = self.id_gen.next();
2667 block.body.push(Instruction::selection_merge(
2668 merge_id,
2669 spirv::SelectionControl::NONE,
2670 ));
2671
2672 let default_id = self.id_gen.next();
2673 let raw_cases = cases
2674 .iter()
2675 .map(|c| super::instructions::Case {
2676 value: c.value as Word,
2677 label_id: self.id_gen.next(),
2678 })
2679 .collect::<Vec<_>>();
2680
2681 function.consume(
2682 block,
2683 Instruction::switch(selector_id, default_id, &raw_cases),
2684 );
2685
2686 for (i, (case, raw_case)) in cases.iter().zip(raw_cases.iter()).enumerate() {
2687 let case_finish_id = if case.fall_through {
2688 match raw_cases.get(i + 1) {
2689 Some(rc) => rc.label_id,
2690 None => default_id,
2691 }
2692 } else {
2693 merge_id
2694 };
2695 self.write_block(
2696 raw_case.label_id,
2697 &case.body,
2698 ir_module,
2699 ir_function,
2700 fun_info,
2701 function,
2702 Some(case_finish_id),
2703 LoopContext::default(),
2704 )?;
2705 }
2706
2707 self.write_block(
2708 default_id,
2709 default,
2710 ir_module,
2711 ir_function,
2712 fun_info,
2713 function,
2714 Some(merge_id),
2715 LoopContext::default(),
2716 )?;
2717
2718 block = Block::new(merge_id);
2719 }
2720 crate::Statement::Loop {
2721 ref body,
2722 ref continuing,
2723 } => {
2724 let preamble_id = self.id_gen.next();
2725 function.consume(block, Instruction::branch(preamble_id));
2726
2727 let merge_id = self.id_gen.next();
2728 let body_id = self.id_gen.next();
2729 let continuing_id = self.id_gen.next();
2730
2731 // SPIR-V requires the continuing to the `OpLoopMerge`,
2732 // so we have to start a new block with it.
2733 block = Block::new(preamble_id);
2734 block.body.push(Instruction::loop_merge(
2735 merge_id,
2736 continuing_id,
2737 spirv::SelectionControl::NONE,
2738 ));
2739 function.consume(block, Instruction::branch(body_id));
2740
2741 self.write_block(
2742 body_id,
2743 body,
2744 ir_module,
2745 ir_function,
2746 fun_info,
2747 function,
2748 Some(continuing_id),
2749 LoopContext {
2750 continuing_id: Some(continuing_id),
2751 break_id: Some(merge_id),
2752 },
2753 )?;
2754
2755 self.write_block(
2756 continuing_id,
2757 continuing,
2758 ir_module,
2759 ir_function,
2760 fun_info,
2761 function,
2762 Some(preamble_id),
2763 LoopContext {
2764 continuing_id: None,
2765 break_id: Some(merge_id),
2766 },
2767 )?;
2768
2769 block = Block::new(merge_id);
2770 }
2771 crate::Statement::Break => {
2772 block.termination = Some(Instruction::branch(loop_context.break_id.unwrap()));
2773 }
2774 crate::Statement::Continue => {
2775 block.termination =
2776 Some(Instruction::branch(loop_context.continuing_id.unwrap()));
2777 }
2778 crate::Statement::Return { value: Some(value) } => {
2779 let value_id = self.cached[value];
2780 let instruction = match function.entry_point_context {
2781 // If this is an entry point, and we need to return anything,
2782 // let's instead store the output variables and return `void`.
2783 Some(ref context) => {
2784 self.write_entry_point_return(
2785 value_id,
2786 ir_function.result.as_ref().unwrap(),
2787 &ir_module.types,
2788 &context.results,
2789 &mut block.body,
2790 )?;
2791 Instruction::return_void()
2792 }
2793 None => Instruction::return_value(value_id),
2794 };
2795 block.termination = Some(instruction);
2796 }
2797 crate::Statement::Return { value: None } => {
2798 block.termination = Some(Instruction::return_void());
2799 }
2800 crate::Statement::Kill => {
2801 block.termination = Some(Instruction::kill());
2802 }
2803 crate::Statement::Barrier(flags) => {
2804 let memory_scope = if flags.contains(crate::Barrier::STORAGE) {
2805 spirv::Scope::Device
2806 } else {
2807 spirv::Scope::Workgroup
2808 };
2809 let mut semantics = spirv::MemorySemantics::ACQUIRE_RELEASE;
2810 semantics.set(
2811 spirv::MemorySemantics::UNIFORM_MEMORY,
2812 flags.contains(crate::Barrier::STORAGE),
2813 );
2814 semantics.set(
2815 spirv::MemorySemantics::WORKGROUP_MEMORY,
2816 flags.contains(crate::Barrier::WORK_GROUP),
2817 );
2818 let exec_scope_id =
2819 self.get_index_constant(spirv::Scope::Workgroup as u32, &ir_module.types)?;
2820 let mem_scope_id =
2821 self.get_index_constant(memory_scope as u32, &ir_module.types)?;
2822 let semantics_id =
2823 self.get_index_constant(semantics.bits(), &ir_module.types)?;
2824 block.body.push(Instruction::control_barrier(
2825 exec_scope_id,
2826 mem_scope_id,
2827 semantics_id,
2828 ));
2829 }
2830 crate::Statement::Store { pointer, value } => {
2831 let (pointer_id, _) = self.write_expression_pointer(
2832 ir_module,
2833 ir_function,
2834 fun_info,
2835 pointer,
2836 &mut block,
2837 function,
2838 )?;
2839 let value_id = self.cached[value];
2840
2841 block
2842 .body
2843 .push(Instruction::store(pointer_id, value_id, None));
2844 }
2845 crate::Statement::ImageStore {
2846 image,
2847 coordinate,
2848 array_index,
2849 value,
2850 } => {
2851 let image_id = self.get_expression_global(ir_function, image);
2852 let coordinate_id = self.write_texture_coordinates(
2853 ir_module,
2854 fun_info,
2855 coordinate,
2856 array_index,
2857 &mut block,
2858 )?;
2859 let value_id = self.cached[value];
2860
2861 block
2862 .body
2863 .push(Instruction::image_write(image_id, coordinate_id, value_id));
2864 }
2865 crate::Statement::Call {
2866 function: local_function,
2867 ref arguments,
2868 result,
2869 } => {
2870 let id = self.id_gen.next();
2871 self.temp_list.clear();
2872 for &argument in arguments {
2873 self.temp_list.push(self.cached[argument]);
2874 }
2875
2876 let type_id = match result {
2877 Some(expr) => {
2878 self.cached[expr] = id;
2879 self.lookup_function_call.insert(expr, id);
2880 let ty_handle = ir_module.functions[local_function]
2881 .result
2882 .as_ref()
2883 .unwrap()
2884 .ty;
2885 self.get_type_id(&ir_module.types, LookupType::Handle(ty_handle))?
2886 }
2887 None => self.void_type,
2888 };
2889
2890 block.body.push(Instruction::function_call(
2891 type_id,
2892 id,
2893 self.lookup_function[&local_function],
2894 &self.temp_list,
2895 ));
2896 }
2897 }
2898 }
2899
2900 if block.termination.is_none() {
2901 block.termination = Some(match exit_id {
2902 Some(id) => Instruction::branch(id),
2903 // This can happen if the last branch had all the paths
2904 // leading out of the graph (i.e. returning).
2905 // Or it may be the end of the function.
2906 None => match ir_function.result {
2907 Some(ref result) if function.entry_point_context.is_none() => {
2908 // create a Null and return it
2909 let null_id = self.id_gen.next();
2910 let type_id =
2911 self.get_type_id(&ir_module.types, LookupType::Handle(result.ty))?;
2912 Instruction::constant_null(type_id, null_id)
2913 .to_words(&mut self.logical_layout.declarations);
2914 Instruction::return_value(null_id)
2915 }
2916 _ => Instruction::return_void(),
2917 },
2918 });
2919 }
2920
2921 function.blocks.push(block);
2922 Ok(())
2923 }
2924
write_physical_layout(&mut self)2925 fn write_physical_layout(&mut self) {
2926 self.physical_layout.bound = self.id_gen.0 + 1;
2927 }
2928
write_logical_layout( &mut self, ir_module: &crate::Module, mod_info: &ModuleInfo, ) -> Result<(), Error>2929 fn write_logical_layout(
2930 &mut self,
2931 ir_module: &crate::Module,
2932 mod_info: &ModuleInfo,
2933 ) -> Result<(), Error> {
2934 let has_storage_buffers = ir_module
2935 .global_variables
2936 .iter()
2937 .any(|(_, var)| var.class == crate::StorageClass::Storage);
2938 if self.physical_layout.version < 0x10300 && has_storage_buffers {
2939 // enable the storage buffer class on < SPV-1.3
2940 Instruction::extension("SPV_KHR_storage_buffer_storage_class")
2941 .to_words(&mut self.logical_layout.extensions);
2942 }
2943 Instruction::type_void(self.void_type).to_words(&mut self.logical_layout.declarations);
2944 Instruction::ext_inst_import(self.gl450_ext_inst_id, "GLSL.std.450")
2945 .to_words(&mut self.logical_layout.ext_inst_imports);
2946
2947 if self.flags.contains(WriterFlags::DEBUG) {
2948 self.debugs
2949 .push(Instruction::source(spirv::SourceLanguage::GLSL, 450));
2950 }
2951
2952 self.constant_ids.clear();
2953 self.constant_ids.resize(ir_module.constants.len(), 0);
2954 // first, output all the scalar constants
2955 for (handle, constant) in ir_module.constants.iter() {
2956 match constant.inner {
2957 crate::ConstantInner::Composite { .. } => continue,
2958 crate::ConstantInner::Scalar { width, ref value } => {
2959 self.constant_ids[handle.index()] = match constant.name {
2960 Some(ref name) => {
2961 let id = self.id_gen.next();
2962 self.write_constant_scalar(
2963 id,
2964 value,
2965 width,
2966 Some(name),
2967 &ir_module.types,
2968 )?;
2969 id
2970 }
2971 None => self.get_constant_scalar(*value, width, &ir_module.types)?,
2972 };
2973 }
2974 }
2975 }
2976
2977 // then all types, some of them may rely on constants and struct type set
2978 for (handle, _) in ir_module.types.iter() {
2979 self.write_type_declaration_arena(&ir_module.types, handle)?;
2980 }
2981
2982 // the all the composite constants, they rely on types
2983 for (handle, constant) in ir_module.constants.iter() {
2984 match constant.inner {
2985 crate::ConstantInner::Scalar { .. } => continue,
2986 crate::ConstantInner::Composite { ty, ref components } => {
2987 let id = self.id_gen.next();
2988 self.constant_ids[handle.index()] = id;
2989 if self.flags.contains(WriterFlags::DEBUG) {
2990 if let Some(ref name) = constant.name {
2991 self.debugs.push(Instruction::name(id, name));
2992 }
2993 }
2994 self.write_constant_composite(id, ty, components, &ir_module.types)?;
2995 }
2996 }
2997 }
2998 debug_assert_eq!(self.constant_ids.iter().position(|&id| id == 0), None);
2999
3000 // now write all globals
3001 self.global_variables.clear();
3002 for (_, var) in ir_module.global_variables.iter() {
3003 let (instruction, id, class) = self.write_global_variable(ir_module, var)?;
3004 instruction.to_words(&mut self.logical_layout.declarations);
3005 self.global_variables.push(GlobalVariable {
3006 id,
3007 handle_id: 0,
3008 class,
3009 });
3010 }
3011
3012 // all functions
3013 for (handle, ir_function) in ir_module.functions.iter() {
3014 let info = &mod_info[handle];
3015 let id = self.write_function(ir_function, info, ir_module, None)?;
3016 self.lookup_function.insert(handle, id);
3017 }
3018
3019 // and entry points
3020 for (ep_index, ir_ep) in ir_module.entry_points.iter().enumerate() {
3021 let info = mod_info.get_entry_point(ep_index);
3022 let ep_instruction = self.write_entry_point(ir_ep, info, ir_module)?;
3023 ep_instruction.to_words(&mut self.logical_layout.entry_points);
3024 }
3025
3026 for capability in self.capabilities.iter() {
3027 Instruction::capability(*capability).to_words(&mut self.logical_layout.capabilities);
3028 }
3029 if ir_module.entry_points.is_empty() {
3030 // SPIR-V doesn't like modules without entry points
3031 Instruction::capability(spirv::Capability::Linkage)
3032 .to_words(&mut self.logical_layout.capabilities);
3033 }
3034
3035 let addressing_model = spirv::AddressingModel::Logical;
3036 let memory_model = spirv::MemoryModel::GLSL450;
3037 self.check(addressing_model.required_capabilities())?;
3038 self.check(memory_model.required_capabilities())?;
3039
3040 Instruction::memory_model(addressing_model, memory_model)
3041 .to_words(&mut self.logical_layout.memory_model);
3042
3043 if self.flags.contains(WriterFlags::DEBUG) {
3044 for debug in self.debugs.iter() {
3045 debug.to_words(&mut self.logical_layout.debugs);
3046 }
3047 }
3048
3049 for annotation in self.annotations.iter() {
3050 annotation.to_words(&mut self.logical_layout.annotations);
3051 }
3052
3053 Ok(())
3054 }
3055
write( &mut self, ir_module: &crate::Module, info: &ModuleInfo, words: &mut Vec<Word>, ) -> Result<(), Error>3056 pub fn write(
3057 &mut self,
3058 ir_module: &crate::Module,
3059 info: &ModuleInfo,
3060 words: &mut Vec<Word>,
3061 ) -> Result<(), Error> {
3062 self.lookup_function.clear();
3063 self.lookup_function_type.clear();
3064 self.lookup_function_call.clear();
3065
3066 self.write_logical_layout(ir_module, info)?;
3067 self.write_physical_layout();
3068
3069 self.physical_layout.in_words(words);
3070 self.logical_layout.in_words(words);
3071 Ok(())
3072 }
3073 }
3074
3075 #[test]
test_write_physical_layout()3076 fn test_write_physical_layout() {
3077 let mut writer = Writer::new(&Options::default()).unwrap();
3078 assert_eq!(writer.physical_layout.bound, 0);
3079 writer.write_physical_layout();
3080 assert_eq!(writer.physical_layout.bound, 3);
3081 }
3082