1 //! OpenGL shading language backend
2 //!
3 //! The main structure is [`Writer`](Writer), it maintains internal state that is used
4 //! to output a [`Module`](crate::Module) into glsl
5 //!
6 //! # Supported versions
7 //! ### Core
8 //! - 330
9 //! - 400
10 //! - 410
11 //! - 420
12 //! - 430
13 //! - 450
14 //! - 460
15 //!
16 //! ### ES
17 //! - 300
18 //! - 310
19 //!
20
21 // GLSL is mostly a superset of C but it also removes some parts of it this is a list of relevant
22 // aspects for this backend.
23 //
24 // The most notable change is the introduction of the version preprocessor directive that must
25 // always be the first line of a glsl file and is written as
26 // `#version number profile`
27 // `number` is the version itself (i.e. 300) and `profile` is the
28 // shader profile we only support "core" and "es", the former is used in desktop applications and
29 // the later is used in embedded contexts, mobile devices and browsers. Each one as it's own
30 // versions (at the time of writing this the latest version for "core" is 460 and for "es" is 320)
31 //
32 // Other important preprocessor addition is the extension directive which is written as
33 // `#extension name: behaviour`
34 // Extensions provide increased features in a plugin fashion but they aren't required to be
35 // supported hence why they are called extensions, that's why `behaviour` is used it specifies
36 // wether the extension is strictly required or if it should only be enabled if needed. In our case
37 // when we use extensions we set behaviour to `require` always.
38 //
39 // The only thing that glsl removes that makes a difference are pointers.
40 //
41 // Additions that are relevant for the backend are the discard keyword, the introduction of
42 // vector, matrices, samplers, image types and functions that provide common shader operations
43
44 pub use features::Features;
45
46 use crate::{
47 proc::{EntryPointIndex, NameKey, Namer, TypeResolution},
48 valid::{FunctionInfo, ModuleInfo},
49 Arena, ArraySize, BinaryOperator, Binding, BuiltIn, Bytes, ConservativeDepth, Constant,
50 ConstantInner, DerivativeAxis, Expression, FastHashMap, Function, GlobalVariable, Handle,
51 ImageClass, Interpolation, LocalVariable, Module, RelationalFunction, Sampling, ScalarKind,
52 ScalarValue, ShaderStage, Statement, StorageAccess, StorageClass, StorageFormat, StructMember,
53 Type, TypeInner, UnaryOperator,
54 };
55 use features::FeaturesManager;
56 use std::{
57 cmp::Ordering,
58 fmt,
59 fmt::{Error as FmtError, Write},
60 };
61 use thiserror::Error;
62
63 /// Contains the features related code and the features querying method
64 mod features;
65 /// Contains a constant with a slice of all the reserved keywords RESERVED_KEYWORDS
66 mod keywords;
67
68 /// List of supported core glsl versions
69 pub const SUPPORTED_CORE_VERSIONS: &[u16] = &[330, 400, 410, 420, 430, 440, 450];
70 /// List of supported es glsl versions
71 pub const SUPPORTED_ES_VERSIONS: &[u16] = &[300, 310, 320];
72 const INDENT: &str = " ";
73
74 const COMPONENTS: &[char] = &['x', 'y', 'z', 'w'];
75
76 /// glsl version
77 #[derive(Debug, Copy, Clone, PartialEq)]
78 pub enum Version {
79 /// `core` glsl
80 Desktop(u16),
81 /// `es` glsl
82 Embedded(u16),
83 }
84
85 impl Version {
86 /// Returns true if self is `Version::Embedded` (i.e. is a es version)
is_es(&self) -> bool87 fn is_es(&self) -> bool {
88 match *self {
89 Version::Desktop(_) => false,
90 Version::Embedded(_) => true,
91 }
92 }
93
94 /// Checks the list of currently supported versions and returns true if it contains the
95 /// specified version
96 ///
97 /// # Notes
98 /// As an invalid version number will never be added to the supported version list
99 /// so this also checks for version validity
is_supported(&self) -> bool100 fn is_supported(&self) -> bool {
101 match *self {
102 Version::Desktop(v) => SUPPORTED_CORE_VERSIONS.contains(&v),
103 Version::Embedded(v) => SUPPORTED_ES_VERSIONS.contains(&v),
104 }
105 }
106
107 /// Checks if the version supports explicit `layout(location=)` qualifiers.
supports_explicit_locations(&self) -> bool108 fn supports_explicit_locations(&self) -> bool {
109 *self >= Version::Embedded(310) || *self >= Version::Desktop(410)
110 }
111 }
112
113 impl PartialOrd for Version {
partial_cmp(&self, other: &Self) -> Option<Ordering>114 fn partial_cmp(&self, other: &Self) -> Option<Ordering> {
115 match (*self, *other) {
116 (Version::Desktop(x), Version::Desktop(y)) => Some(x.cmp(&y)),
117 (Version::Embedded(x), Version::Embedded(y)) => Some(x.cmp(&y)),
118 _ => None,
119 }
120 }
121 }
122
123 impl fmt::Display for Version {
fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result124 fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
125 match *self {
126 Version::Desktop(v) => write!(f, "{} core", v),
127 Version::Embedded(v) => write!(f, "{} es", v),
128 }
129 }
130 }
131
132 /// Structure that contains the configuration used in the [`Writer`](Writer)
133 #[derive(Debug, Clone)]
134 pub struct Options {
135 /// The glsl version to be used
136 pub version: Version,
137 /// The stage of the entry point
138 pub shader_stage: ShaderStage,
139 /// The name of the entry point
140 ///
141 /// If no entry point that matches is found a error will be thrown while creating a new instance
142 /// of [`Writer`](struct.Writer.html)
143 pub entry_point: String,
144 }
145
146 impl Default for Options {
default() -> Self147 fn default() -> Self {
148 Options {
149 version: Version::Embedded(320),
150 shader_stage: ShaderStage::Compute,
151 entry_point: "main".to_string(),
152 }
153 }
154 }
155
156 /// Structure that contains a reflection info
157 pub struct ReflectionInfo {
158 pub texture_mapping: FastHashMap<String, TextureMapping>,
159 pub uniforms: FastHashMap<Handle<GlobalVariable>, String>,
160 }
161
162 /// Structure that connects a texture to a sampler or not
163 ///
164 /// glsl pre vulkan has no concept of separate textures and samplers instead everything is a
165 /// `gsamplerN` where `g` is the scalar type and `N` is the dimension, but naga uses separate textures
166 /// and samplers in the IR so the backend produces a [`HashMap`](crate::FastHashMap) with the texture name
167 /// as a key and a [`TextureMapping`](TextureMapping) as a value this way the user knows where to bind.
168 ///
169 /// [`Storage`](crate::ImageClass::Storage) images produce `gimageN` and don't have an associated sampler
170 /// so the [`sampler`](Self::sampler) field will be [`None`](std::option::Option::None)
171 #[derive(Debug, Clone)]
172 pub struct TextureMapping {
173 /// Handle to the image global variable
174 pub texture: Handle<GlobalVariable>,
175 /// Handle to the associated sampler global variable if it exists
176 pub sampler: Option<Handle<GlobalVariable>>,
177 }
178
179 /// Stores the current function type (either a regular function or an entry point)
180 ///
181 /// Also stores data needed to identify it (handle for a regular function or index for an entry point)
182 enum FunctionType {
183 /// A regular function and it's handle
184 Function(Handle<Function>),
185 /// A entry point and it's index
186 EntryPoint(EntryPointIndex),
187 }
188
189 /// Helper structure that stores data needed when writing the function
190 struct FunctionCtx<'a> {
191 /// The current function being written
192 func: FunctionType,
193 /// Analysis about the function
194 info: &'a FunctionInfo,
195 /// The expression arena of the current function being written
196 expressions: &'a Arena<Expression>,
197 }
198
199 impl<'a> FunctionCtx<'a> {
200 /// Helper method that generates a [`NameKey`](crate::proc::NameKey) for a local in the current function
name_key(&self, local: Handle<LocalVariable>) -> NameKey201 fn name_key(&self, local: Handle<LocalVariable>) -> NameKey {
202 match self.func {
203 FunctionType::Function(handle) => NameKey::FunctionLocal(handle, local),
204 FunctionType::EntryPoint(idx) => NameKey::EntryPointLocal(idx, local),
205 }
206 }
207
208 /// Helper method that generates a [`NameKey`](crate::proc::NameKey) for a function argument.
209 ///
210 /// # Panics
211 /// - If the function arguments are less or equal to `arg`
argument_key(&self, arg: u32) -> NameKey212 fn argument_key(&self, arg: u32) -> NameKey {
213 match self.func {
214 FunctionType::Function(handle) => NameKey::FunctionArgument(handle, arg),
215 FunctionType::EntryPoint(ep_index) => NameKey::EntryPointArgument(ep_index, arg),
216 }
217 }
218 }
219
220 /// Helper structure that generates a number
221 #[derive(Default)]
222 struct IdGenerator(u32);
223
224 impl IdGenerator {
225 /// Generates a number that's guaranteed to be unique for this `IdGenerator`
generate(&mut self) -> u32226 fn generate(&mut self) -> u32 {
227 // It's just an increasing number but it does the job
228 let ret = self.0;
229 self.0 += 1;
230 ret
231 }
232 }
233
234 /// Helper wrapper used to get a name for a varying
235 ///
236 /// Varying have different naming schemes depending on their binding:
237 /// - Varyings with builtin bindings get the from [`glsl_built_in`](glsl_built_in).
238 /// - Varyings with location bindings are named `_S_location_X` where `S` is a
239 /// prefix identifying which pipeline stage the varying connects, and `X` is
240 /// the location.
241 struct VaryingName<'a> {
242 binding: &'a Binding,
243 stage: ShaderStage,
244 output: bool,
245 }
246 impl fmt::Display for VaryingName<'_> {
fmt(&self, f: &mut fmt::Formatter) -> fmt::Result247 fn fmt(&self, f: &mut fmt::Formatter) -> fmt::Result {
248 match *self.binding {
249 Binding::Location { location, .. } => {
250 let prefix = match (self.stage, self.output) {
251 (ShaderStage::Compute, _) => unreachable!(),
252 // pipeline to vertex
253 (ShaderStage::Vertex, false) => "p2vs",
254 // vertex to fragment
255 (ShaderStage::Vertex, true) | (ShaderStage::Fragment, false) => "vs2fs",
256 // fragment to pipeline
257 (ShaderStage::Fragment, true) => "fs2p",
258 };
259 write!(f, "_{}_location{}", prefix, location,)
260 }
261 Binding::BuiltIn(built_in) => {
262 write!(f, "{}", glsl_built_in(built_in, self.output))
263 }
264 }
265 }
266 }
267
268 /// Shorthand result used internally by the backend
269 type BackendResult = Result<(), Error>;
270
271 /// A glsl compilation error.
272 #[derive(Debug, Error)]
273 pub enum Error {
274 /// A error occurred while writing to the output
275 #[error("Format error")]
276 FmtError(#[from] FmtError),
277 /// The specified [`Version`](Version) doesn't have all required [`Features`](super)
278 ///
279 /// Contains the missing [`Features`](Features)
280 #[error("The selected version doesn't support {0:?}")]
281 MissingFeatures(Features),
282 /// [`StorageClass::PushConstant`](crate::StorageClass::PushConstant) was used and isn't
283 /// supported in the glsl backend
284 #[error("Push constants aren't supported")]
285 PushConstantNotSupported,
286 /// The specified [`Version`](Version) isn't supported
287 #[error("The specified version isn't supported")]
288 VersionNotSupported,
289 /// The entry point couldn't be found
290 #[error("The requested entry point couldn't be found")]
291 EntryPointNotFound,
292 /// A call was made to an unsupported external
293 #[error("A call was made to an unsupported external: {0}")]
294 UnsupportedExternal(String),
295 /// A scalar with an unsupported width was requested
296 #[error("A scalar with an unsupported width was requested: {0:?} {1:?}")]
297 UnsupportedScalar(ScalarKind, Bytes),
298 /// [`Interpolation::Patch`](crate::Interpolation::Patch) isn't supported
299 #[error("Patch interpolation isn't supported")]
300 PatchInterpolationNotSupported,
301 /// A image was used with multiple samplers, this isn't supported
302 #[error("A image was used with multiple samplers")]
303 ImageMultipleSamplers,
304 #[error("{0}")]
305 Custom(String),
306 }
307
308 /// Main structure of the glsl backend responsible for all code generation
309 pub struct Writer<'a, W> {
310 // Inputs
311 /// The module being written
312 module: &'a Module,
313 /// The module analysis.
314 info: &'a ModuleInfo,
315 /// The output writer
316 out: W,
317 /// User defined configuration to be used
318 options: &'a Options,
319
320 // Internal State
321 /// Features manager used to store all the needed features and write them
322 features: FeaturesManager,
323 /// A map with all the names needed for writing the module
324 /// (generated by a [`Namer`](crate::proc::Namer))
325 names: FastHashMap<NameKey, String>,
326 /// A map with all the names needed for reflections
327 reflection_names: FastHashMap<Handle<Type>, String>,
328 /// The selected entry point
329 entry_point: &'a crate::EntryPoint,
330 /// The index of the selected entry point
331 entry_point_idx: EntryPointIndex,
332 /// Used to generate a unique number for blocks
333 block_id: IdGenerator,
334 /// Set of expressions that have associated temporary variables
335 cached_expressions: FastHashMap<Handle<Expression>, String>,
336 }
337
338 impl<'a, W: Write> Writer<'a, W> {
339 /// Creates a new [`Writer`](Writer) instance
340 ///
341 /// # Errors
342 /// - If the version specified isn't supported (or invalid)
343 /// - If the entry point couldn't be found on the module
344 /// - If the version specified doesn't support some used features
new( out: W, module: &'a Module, info: &'a ModuleInfo, options: &'a Options, ) -> Result<Self, Error>345 pub fn new(
346 out: W,
347 module: &'a Module,
348 info: &'a ModuleInfo,
349 options: &'a Options,
350 ) -> Result<Self, Error> {
351 // Check if the requested version is supported
352 if !options.version.is_supported() {
353 log::error!("Version {}", options.version);
354 return Err(Error::VersionNotSupported);
355 }
356
357 // Try to find the entry point and corresponding index
358 let ep_idx = module
359 .entry_points
360 .iter()
361 .position(|ep| options.shader_stage == ep.stage && options.entry_point == ep.name)
362 .ok_or(Error::EntryPointNotFound)?;
363
364 // Generate a map with names required to write the module
365 let mut names = FastHashMap::default();
366 Namer::default().reset(module, keywords::RESERVED_KEYWORDS, &["gl_"], &mut names);
367
368 // Build the instance
369 let mut this = Self {
370 module,
371 info,
372 out,
373 options,
374
375 features: FeaturesManager::new(),
376 names,
377 reflection_names: FastHashMap::default(),
378 entry_point: &module.entry_points[ep_idx],
379 entry_point_idx: ep_idx as u16,
380
381 block_id: IdGenerator::default(),
382 cached_expressions: FastHashMap::default(),
383 };
384
385 // Find all features required to print this module
386 this.collect_required_features()?;
387
388 Ok(this)
389 }
390
391 /// Writes the [`Module`](crate::Module) as glsl to the output
392 ///
393 /// # Notes
394 /// If an error occurs while writing, the output might have been written partially
395 ///
396 /// # Panics
397 /// Might panic if the module is invalid
write(&mut self) -> Result<ReflectionInfo, Error>398 pub fn write(&mut self) -> Result<ReflectionInfo, Error> {
399 // We use `writeln!(self.out)` throughout the write to add newlines
400 // to make the output more readable
401
402 let es = self.options.version.is_es();
403
404 // Write the version (It must be the first thing or it isn't a valid glsl output)
405 writeln!(self.out, "#version {}", self.options.version)?;
406 // Write all the needed extensions
407 //
408 // This used to be the last thing being written as it allowed to search for features while
409 // writing the module saving some loops but some older versions (420 or less) required the
410 // extensions to appear before being used, even though extensions are part of the
411 // preprocessor not the processor ¯\_(ツ)_/¯
412 self.features.write(self.options.version, &mut self.out)?;
413
414 // glsl es requires a precision to be specified for floats
415 // TODO: Should this be user configurable?
416 if es {
417 writeln!(self.out)?;
418 writeln!(self.out, "precision highp float;")?;
419 writeln!(self.out)?;
420 }
421
422 if self.options.shader_stage == ShaderStage::Compute {
423 let workgroup_size = self.entry_point.workgroup_size;
424 writeln!(
425 self.out,
426 "layout(local_size_x = {}, local_size_y = {}, local_size_z = {}) in;",
427 workgroup_size[0], workgroup_size[1], workgroup_size[2]
428 )?;
429 writeln!(self.out)?;
430 }
431
432 // Enable early depth tests if needed
433 if let Some(depth_test) = self.entry_point.early_depth_test {
434 writeln!(self.out, "layout(early_fragment_tests) in;")?;
435
436 if let Some(conservative) = depth_test.conservative {
437 writeln!(
438 self.out,
439 "layout (depth_{}) out float gl_FragDepth;",
440 match conservative {
441 ConservativeDepth::GreaterEqual => "greater",
442 ConservativeDepth::LessEqual => "less",
443 ConservativeDepth::Unchanged => "unchanged",
444 }
445 )?;
446 }
447 writeln!(self.out)?;
448 }
449
450 // Write all structs
451 //
452 // This are always ordered because of the IR is structured in a way that you can't make a
453 // struct without adding all of it's members first
454 for (handle, ty) in self.module.types.iter() {
455 if let TypeInner::Struct { ref members, .. } = ty.inner {
456 // No needed to write a struct that also should be written as a global variable
457 let is_global_struct = self
458 .module
459 .global_variables
460 .iter()
461 .any(|e| e.1.ty == handle);
462
463 if !is_global_struct {
464 self.write_struct(false, handle, members)?
465 }
466 }
467 }
468
469 let ep_info = self.info.get_entry_point(self.entry_point_idx as usize);
470
471 // Write the globals
472 //
473 // We filter all globals that aren't used by the selected entry point as they might be
474 // interfere with each other (i.e. two globals with the same location but different with
475 // different classes)
476 for (handle, global) in self.module.global_variables.iter() {
477 if ep_info[handle].is_empty() {
478 continue;
479 }
480
481 match self.module.types[global.ty].inner {
482 // We treat images separately because they might require
483 // writing the storage format
484 TypeInner::Image {
485 dim,
486 arrayed,
487 class,
488 } => {
489 // Write the storage format if needed
490 if let TypeInner::Image {
491 class: ImageClass::Storage(format),
492 ..
493 } = self.module.types[global.ty].inner
494 {
495 write!(self.out, "layout({}) ", glsl_storage_format(format))?;
496 }
497
498 if let Some(storage_access) = glsl_storage_access(global.storage_access) {
499 write!(self.out, "{} ", storage_access)?;
500 }
501
502 // All images in glsl are `uniform`
503 // The trailing space is important
504 write!(self.out, "uniform ")?;
505
506 // write the type
507 //
508 // This is way we need the leading space because `write_image_type` doesn't add
509 // any spaces at the beginning or end
510 self.write_image_type(dim, arrayed, class)?;
511
512 // Finally write the name and end the global with a `;`
513 // The leading space is important
514 let global_name = self.get_global_name(handle, global);
515 writeln!(self.out, " {};", global_name)?;
516 writeln!(self.out)?;
517
518 self.reflection_names.insert(global.ty, global_name);
519 }
520 // glsl has no concept of samplers so we just ignore it
521 TypeInner::Sampler { .. } => continue,
522 // All other globals are written by `write_global`
523 _ => self.write_global(handle, global)?,
524 }
525 }
526
527 for arg in self.entry_point.function.arguments.iter() {
528 self.write_varying(arg.binding.as_ref(), arg.ty, false)?;
529 }
530 if let Some(ref result) = self.entry_point.function.result {
531 self.write_varying(result.binding.as_ref(), result.ty, true)?;
532 }
533 writeln!(self.out)?;
534
535 // Write all regular functions
536 for (handle, function) in self.module.functions.iter() {
537 // Check that the function doesn't use globals that aren't supported
538 // by the current entry point
539 if !ep_info.dominates_global_use(&self.info[handle]) {
540 continue;
541 }
542
543 // We also `clone` to satisfy the borrow checker
544 let name = self.names[&NameKey::Function(handle)].clone();
545 let fun_info = &self.info[handle];
546
547 // Write the function
548 self.write_function(FunctionType::Function(handle), function, fun_info, &name)?;
549
550 writeln!(self.out)?;
551 }
552
553 self.write_function(
554 FunctionType::EntryPoint(self.entry_point_idx),
555 &self.entry_point.function,
556 ep_info,
557 "main",
558 )?;
559
560 // Add newline at the end of file
561 writeln!(self.out)?;
562
563 // Collect all relection info and return it to the user
564 self.collect_reflection_info()
565 }
566
567 /// Helper method used to write value types
568 ///
569 /// # Notes
570 /// Adds no trailing or leading whitespace
571 ///
572 /// # Panics
573 /// - If type is either a image, a sampler, a pointer, or a struct
574 /// - If it's an Array with a [`ArraySize::Constant`](crate::ArraySize::Constant) with a
575 /// constant that isn't [`Uint`](crate::ConstantInner::Uint)
write_value_type(&mut self, inner: &TypeInner) -> BackendResult576 fn write_value_type(&mut self, inner: &TypeInner) -> BackendResult {
577 match *inner {
578 // Scalars are simple we just get the full name from `glsl_scalar`
579 TypeInner::Scalar { kind, width }
580 | TypeInner::ValuePointer {
581 size: None,
582 kind,
583 width,
584 class: _,
585 } => write!(self.out, "{}", glsl_scalar(kind, width)?.full)?,
586 // Vectors are just `gvecN` where `g` is the scalar prefix and `N` is the vector size
587 TypeInner::Vector { size, kind, width }
588 | TypeInner::ValuePointer {
589 size: Some(size),
590 kind,
591 width,
592 class: _,
593 } => write!(
594 self.out,
595 "{}vec{}",
596 glsl_scalar(kind, width)?.prefix,
597 size as u8
598 )?,
599 // Matrices are written with `gmatMxN` where `g` is the scalar prefix (only floats and
600 // doubles are allowed), `M` is the columns count and `N` is the rows count
601 //
602 // glsl supports a matrix shorthand `gmatN` where `N` = `M` but it doesn't justify the
603 // extra branch to write matrices this way
604 TypeInner::Matrix {
605 columns,
606 rows,
607 width,
608 } => write!(
609 self.out,
610 "{}mat{}x{}",
611 glsl_scalar(ScalarKind::Float, width)?.prefix,
612 columns as u8,
613 rows as u8
614 )?,
615 // GLSL arrays are written as `type name[size]`
616 // Current code is written arrays only as `[size]`
617 // Base `type` and `name` should be written outside
618 TypeInner::Array { base: _, size, .. } => {
619 write!(self.out, "[")?;
620
621 // Write the array size
622 // Writes nothing if `ArraySize::Dynamic`
623 // Panics if `ArraySize::Constant` has a constant that isn't an uint
624 match size {
625 ArraySize::Constant(const_handle) => {
626 match self.module.constants[const_handle].inner {
627 ConstantInner::Scalar {
628 width: _,
629 value: ScalarValue::Uint(size),
630 } => write!(self.out, "{}", size)?,
631 _ => unreachable!(),
632 }
633 }
634 ArraySize::Dynamic => (),
635 }
636
637 write!(self.out, "]")?
638 }
639 // Panic if either Image, Sampler, Pointer, or a Struct is being written
640 //
641 // Write all variants instead of `_` so that if new variants are added a
642 // no exhaustiveness error is thrown
643 TypeInner::Pointer { .. }
644 | TypeInner::Struct { .. }
645 | TypeInner::Image { .. }
646 | TypeInner::Sampler { .. } => unreachable!(),
647 }
648
649 Ok(())
650 }
651
652 /// Helper method used to write non image/sampler types
653 ///
654 /// # Notes
655 /// Adds no trailing or leading whitespace
656 ///
657 /// # Panics
658 /// - If type is either a image or sampler
659 /// - If it's an Array with a [`ArraySize::Constant`](crate::ArraySize::Constant) with a
660 /// constant that isn't [`Uint`](crate::ConstantInner::Uint)
write_type(&mut self, ty: Handle<Type>) -> BackendResult661 fn write_type(&mut self, ty: Handle<Type>) -> BackendResult {
662 match self.module.types[ty].inner {
663 // glsl has no pointer types so just write types as normal and loads are skipped
664 TypeInner::Pointer { base, .. } => self.write_type(base),
665 TypeInner::Struct {
666 level: crate::StructLevel::Root,
667 ref members,
668 span: _,
669 } => self.write_struct(true, ty, members),
670 // glsl structs are written as just the struct name if it isn't a block
671 TypeInner::Struct { .. } => {
672 // Get the struct name
673 let name = &self.names[&NameKey::Type(ty)];
674 write!(self.out, "{}", name)?;
675 Ok(())
676 }
677 ref other => self.write_value_type(other),
678 }
679 }
680
681 /// Helper method to write a image type
682 ///
683 /// # Notes
684 /// Adds no leading or trailing whitespace
write_image_type( &mut self, dim: crate::ImageDimension, arrayed: bool, class: ImageClass, ) -> BackendResult685 fn write_image_type(
686 &mut self,
687 dim: crate::ImageDimension,
688 arrayed: bool,
689 class: ImageClass,
690 ) -> BackendResult {
691 // glsl images consist of four parts the scalar prefix, the image "type", the dimensions
692 // and modifiers
693 //
694 // There exists two image types
695 // - sampler - for sampled images
696 // - image - for storage images
697 //
698 // There are three possible modifiers that can be used together and must be written in
699 // this order to be valid
700 // - MS - used if it's a multisampled image
701 // - Array - used if it's an image array
702 // - Shadow - used if it's a depth image
703
704 let (base, kind, ms, comparison) = match class {
705 ImageClass::Sampled { kind, multi: true } => ("sampler", kind, "MS", ""),
706 ImageClass::Sampled { kind, multi: false } => ("sampler", kind, "", ""),
707 ImageClass::Depth => ("sampler", crate::ScalarKind::Float, "", "Shadow"),
708 ImageClass::Storage(format) => ("image", format.into(), "", ""),
709 };
710
711 write!(
712 self.out,
713 "highp {}{}{}{}{}{}",
714 glsl_scalar(kind, 4)?.prefix,
715 base,
716 glsl_dimension(dim),
717 ms,
718 if arrayed { "Array" } else { "" },
719 comparison
720 )?;
721
722 Ok(())
723 }
724
725 /// Helper method used to write non images/sampler globals
726 ///
727 /// # Notes
728 /// Adds a newline
729 ///
730 /// # Panics
731 /// If the global has type sampler
write_global( &mut self, handle: Handle<GlobalVariable>, global: &GlobalVariable, ) -> BackendResult732 fn write_global(
733 &mut self,
734 handle: Handle<GlobalVariable>,
735 global: &GlobalVariable,
736 ) -> BackendResult {
737 if let Some(storage_access) = glsl_storage_access(global.storage_access) {
738 write!(self.out, "{} ", storage_access)?;
739 }
740
741 // Write the storage class
742 // Trailing space is important
743 if let Some(storage_class) = glsl_storage_class(global.class) {
744 write!(self.out, "{} ", storage_class)?;
745 } else if let TypeInner::Struct {
746 level: crate::StructLevel::Root,
747 ..
748 } = self.module.types[global.ty].inner
749 {
750 write!(self.out, "struct ")?;
751 }
752
753 // Write the type
754 // `write_type` adds no leading or trailing spaces
755 self.write_type(global.ty)?;
756
757 // Finally write the global name and end the global with a `;` and a newline
758 // Leading space is important
759 let global_name = self.get_global_name(handle, global);
760 let global_str =
761 if let Some(default_value) = zero_init_value_str(&self.module.types[global.ty].inner) {
762 format!("{} = {}", global_name, default_value)
763 } else {
764 global_name
765 };
766 writeln!(self.out, " {};", global_str)?;
767 writeln!(self.out)?;
768
769 Ok(())
770 }
771
772 /// Helper method used to get a name for a global
773 ///
774 /// Globals have different naming schemes depending on their binding:
775 /// - Globals without bindings use the name from the [`Namer`](crate::proc::Namer)
776 /// - Globals with resource binding are named `_group_X_binding_Y` where `X`
777 /// is the group and `Y` is the binding
get_global_name(&self, handle: Handle<GlobalVariable>, global: &GlobalVariable) -> String778 fn get_global_name(&self, handle: Handle<GlobalVariable>, global: &GlobalVariable) -> String {
779 match global.binding {
780 Some(ref br) => {
781 format!("_group_{}_binding_{}", br.group, br.binding)
782 }
783 None => self.names[&NameKey::GlobalVariable(handle)].clone(),
784 }
785 }
786
787 /// Writes the varying declaration.
write_varying( &mut self, binding: Option<&Binding>, ty: Handle<Type>, output: bool, ) -> Result<(), Error>788 fn write_varying(
789 &mut self,
790 binding: Option<&Binding>,
791 ty: Handle<Type>,
792 output: bool,
793 ) -> Result<(), Error> {
794 match self.module.types[ty].inner {
795 crate::TypeInner::Struct { ref members, .. } => {
796 for member in members {
797 self.write_varying(member.binding.as_ref(), member.ty, output)?;
798 }
799 }
800 _ => {
801 let (location, interpolation, sampling) = match binding {
802 Some(&Binding::Location {
803 location,
804 interpolation,
805 sampling,
806 }) => (location, interpolation, sampling),
807 _ => return Ok(()),
808 };
809
810 // Write the interpolation modifier if needed
811 //
812 // We ignore all interpolation and auxiliary modifiers that aren't used in fragment
813 // shaders' input globals or vertex shaders' output globals.
814 let emit_interpolation_and_auxiliary = match self.options.shader_stage {
815 ShaderStage::Vertex => output,
816 ShaderStage::Fragment => !output,
817 _ => false,
818 };
819 if let Some(interp) = interpolation {
820 if emit_interpolation_and_auxiliary {
821 write!(self.out, "{} ", glsl_interpolation(interp))?;
822 }
823 }
824
825 // Write the storage class
826 if self.options.version.supports_explicit_locations() {
827 write!(self.out, "layout(location = {}) ", location)?;
828 }
829
830 // Write the sampling auxiliary qualifier.
831 //
832 // Before GLSL 4.2, the `centroid` and `sample` qualifiers were required to appear
833 // immediately before the `in` / `out` qualifier, so we'll just follow that rule
834 // here, regardless of the version.
835 if let Some(sampling) = sampling {
836 if emit_interpolation_and_auxiliary {
837 if let Some(qualifier) = glsl_sampling(sampling) {
838 write!(self.out, "{} ", qualifier)?;
839 }
840 }
841 }
842
843 // Write the input/output qualifier.
844 write!(self.out, "{} ", if output { "out" } else { "in" })?;
845
846 // Write the type
847 // `write_type` adds no leading or trailing spaces
848 self.write_type(ty)?;
849
850 // Finally write the global name and end the global with a `;` and a newline
851 // Leading space is important
852 let vname = VaryingName {
853 binding: &Binding::Location {
854 location,
855 interpolation: None,
856 sampling: None,
857 },
858 stage: self.entry_point.stage,
859 output,
860 };
861 writeln!(self.out, " {};", vname)?;
862 }
863 }
864 Ok(())
865 }
866
867 /// Helper method used to write functions (both entry points and regular functions)
868 ///
869 /// # Notes
870 /// Adds a newline
write_function( &mut self, ty: FunctionType, func: &Function, info: &FunctionInfo, name: &str, ) -> BackendResult871 fn write_function(
872 &mut self,
873 ty: FunctionType,
874 func: &Function,
875 info: &FunctionInfo,
876 name: &str,
877 ) -> BackendResult {
878 // Create a function context for the function being written
879 let ctx = FunctionCtx {
880 func: ty,
881 info,
882 expressions: &func.expressions,
883 };
884
885 self.cached_expressions.clear();
886
887 // Write the function header
888 //
889 // glsl headers are the same as in c:
890 // `ret_type name(args)`
891 // `ret_type` is the return type
892 // `name` is the function name
893 // `args` is a comma separated list of `type name`
894 // | - `type` is the argument type
895 // | - `name` is the argument name
896
897 // Start by writing the return type if any otherwise write void
898 // This is the only place where `void` is a valid type
899 // (though it's more a keyword than a type)
900 if let FunctionType::EntryPoint(_) = ctx.func {
901 write!(self.out, "void")?;
902 } else if let Some(ref result) = func.result {
903 self.write_type(result.ty)?;
904 } else {
905 write!(self.out, "void")?;
906 }
907
908 // Write the function name and open parentheses for the argument list
909 write!(self.out, " {}(", name)?;
910
911 // Write the comma separated argument list
912 //
913 // We need access to `Self` here so we use the reference passed to the closure as an
914 // argument instead of capturing as that would cause a borrow checker error
915 let arguments = match ctx.func {
916 FunctionType::EntryPoint(_) => &[][..],
917 FunctionType::Function(_) => &func.arguments,
918 };
919 self.write_slice(arguments, |this, i, arg| {
920 // Write the argument type
921 // `write_type` adds no trailing spaces
922 this.write_type(arg.ty)?;
923
924 // Write the argument name
925 // The leading space is important
926 write!(this.out, " {}", &this.names[&ctx.argument_key(i)])?;
927
928 Ok(())
929 })?;
930
931 // Close the parentheses and open braces to start the function body
932 writeln!(self.out, ") {{")?;
933
934 // Compose the function arguments from globals, in case of an entry point.
935 if let FunctionType::EntryPoint(ep_index) = ctx.func {
936 let stage = self.module.entry_points[ep_index as usize].stage;
937 for (index, arg) in func.arguments.iter().enumerate() {
938 write!(self.out, "{}", INDENT)?;
939 self.write_type(arg.ty)?;
940 let name = &self.names[&NameKey::EntryPointArgument(ep_index, index as u32)];
941 write!(self.out, " {}", name)?;
942 write!(self.out, " = ")?;
943 match self.module.types[arg.ty].inner {
944 crate::TypeInner::Struct { ref members, .. } => {
945 self.write_type(arg.ty)?;
946 write!(self.out, "(")?;
947 for (index, member) in members.iter().enumerate() {
948 let varying_name = VaryingName {
949 binding: member.binding.as_ref().unwrap(),
950 stage,
951 output: false,
952 };
953 if index != 0 {
954 write!(self.out, ", ")?;
955 }
956 write!(self.out, "{}", varying_name)?;
957 }
958 writeln!(self.out, ");")?;
959 }
960 _ => {
961 let varying_name = VaryingName {
962 binding: arg.binding.as_ref().unwrap(),
963 stage,
964 output: false,
965 };
966 writeln!(self.out, "{};", varying_name)?;
967 }
968 }
969 }
970 }
971
972 // Write all function locals
973 // Locals are `type name (= init)?;` where the init part (including the =) are optional
974 //
975 // Always adds a newline
976 for (handle, local) in func.local_variables.iter() {
977 // Write indentation (only for readability) and the type
978 // `write_type` adds no trailing space
979 write!(self.out, "{}", INDENT)?;
980 self.write_type(local.ty)?;
981
982 // Write the local name
983 // The leading space is important
984 write!(self.out, " {}", self.names[&ctx.name_key(handle)])?;
985
986 // Write the local initializer if needed
987 if let Some(init) = local.init {
988 // Put the equal signal only if there's a initializer
989 // The leading and trailing spaces aren't needed but help with readability
990 write!(self.out, " = ")?;
991
992 // Write the constant
993 // `write_constant` adds no trailing or leading space/newline
994 self.write_constant(&self.module.constants[init])?;
995 }
996
997 // Finish the local with `;` and add a newline (only for readability)
998 writeln!(self.out, ";")?
999 }
1000
1001 // Write the function body (statement list)
1002 for sta in func.body.iter() {
1003 // Write a statement, the indentation should always be 1 when writing the function body
1004 // `write_stmt` adds a newline
1005 self.write_stmt(sta, &ctx, 1)?;
1006 }
1007
1008 // Close braces and add a newline
1009 writeln!(self.out, "}}")?;
1010
1011 Ok(())
1012 }
1013
1014 /// Helper method that writes a list of comma separated `T` with a writer function `F`
1015 ///
1016 /// The writer function `F` receives a mutable reference to `self` that if needed won't cause
1017 /// borrow checker issues (using for example a closure with `self` will cause issues), the
1018 /// second argument is the 0 based index of the element on the list, and the last element is
1019 /// a reference to the element `T` being written
1020 ///
1021 /// # Notes
1022 /// - Adds no newlines or leading/trailing whitespace
1023 /// - The last element won't have a trailing `,`
write_slice<T, F: FnMut(&mut Self, u32, &T) -> BackendResult>( &mut self, data: &[T], mut f: F, ) -> BackendResult1024 fn write_slice<T, F: FnMut(&mut Self, u32, &T) -> BackendResult>(
1025 &mut self,
1026 data: &[T],
1027 mut f: F,
1028 ) -> BackendResult {
1029 // Loop trough `data` invoking `f` for each element
1030 for (i, item) in data.iter().enumerate() {
1031 f(self, i as u32, item)?;
1032
1033 // Only write a comma if isn't the last element
1034 if i != data.len().saturating_sub(1) {
1035 // The leading space is for readability only
1036 write!(self.out, ", ")?;
1037 }
1038 }
1039
1040 Ok(())
1041 }
1042
1043 /// Helper method used to write constants
1044 ///
1045 /// # Notes
1046 /// Adds no newlines or leading/trailing whitespace
write_constant(&mut self, constant: &Constant) -> BackendResult1047 fn write_constant(&mut self, constant: &Constant) -> BackendResult {
1048 match constant.inner {
1049 ConstantInner::Scalar {
1050 width: _,
1051 ref value,
1052 } => match *value {
1053 // Signed integers don't need anything special
1054 ScalarValue::Sint(int) => write!(self.out, "{}", int)?,
1055 // Unsigned integers need a `u` at the end
1056 //
1057 // While `core` doesn't necessarily need it, it's allowed and since `es` needs it we
1058 // always write it as the extra branch wouldn't have any benefit in readability
1059 ScalarValue::Uint(int) => write!(self.out, "{}u", int)?,
1060 // Floats are written using `Debug` instead of `Display` because it always appends the
1061 // decimal part even it's zero which is needed for a valid glsl float constant
1062 ScalarValue::Float(float) => write!(self.out, "{:?}", float)?,
1063 // Booleans are either `true` or `false` so nothing special needs to be done
1064 ScalarValue::Bool(boolean) => write!(self.out, "{}", boolean)?,
1065 },
1066 // Composite constant are created using the same syntax as compose
1067 // `type(components)` where `components` is a comma separated list of constants
1068 ConstantInner::Composite { ty, ref components } => {
1069 self.write_type(ty)?;
1070 write!(self.out, "(")?;
1071
1072 // Write the comma separated constants
1073 self.write_slice(components, |this, _, arg| {
1074 this.write_constant(&this.module.constants[*arg])
1075 })?;
1076
1077 write!(self.out, ")")?
1078 }
1079 }
1080
1081 Ok(())
1082 }
1083
1084 /// Helper method used to write structs
1085 ///
1086 /// # Notes
1087 /// Ends in a newline
write_struct( &mut self, block: bool, handle: Handle<Type>, members: &[StructMember], ) -> BackendResult1088 fn write_struct(
1089 &mut self,
1090 block: bool,
1091 handle: Handle<Type>,
1092 members: &[StructMember],
1093 ) -> BackendResult {
1094 // glsl structs are written as in C
1095 // `struct name() { members };`
1096 // | `struct` is a keyword
1097 // | `name` is the struct name
1098 // | `members` is a semicolon separated list of `type name`
1099 // | `type` is the member type
1100 // | `name` is the member name
1101 let name = &self.names[&NameKey::Type(handle)];
1102
1103 // If struct is a block we need to write `block_name { members }` where `block_name` must be
1104 // unique between blocks and structs so we add `_block_ID` where `ID` is a `IdGenerator`
1105 // generated number so it's unique and `members` are the same as in a struct
1106 if block {
1107 // Write the block name, it's just the struct name appended with `_block_ID`
1108 let block_name = format!("{}_block_{}", name, self.block_id.generate());
1109 writeln!(self.out, "{} {{", block_name)?;
1110
1111 self.reflection_names.insert(handle, block_name);
1112 } else {
1113 writeln!(self.out, "struct {} {{", name)?;
1114 }
1115
1116 for (idx, member) in members.iter().enumerate() {
1117 // The indentation is only for readability
1118 write!(self.out, "{}", INDENT)?;
1119
1120 match self.module.types[member.ty].inner {
1121 TypeInner::Array { base, .. } => {
1122 // GLSL arrays are written as `type name[size]`
1123 let ty_name = match self.module.types[base].inner {
1124 // Write scalar type by backend so as not to depend on the front-end implementation
1125 // Name returned from frontend can be generated (type1, float1, etc.)
1126 TypeInner::Scalar { kind, width } => glsl_scalar(kind, width)?.full,
1127 _ => &self.names[&NameKey::Type(base)],
1128 };
1129
1130 // Write `type` and `name`
1131 write!(self.out, "{}", ty_name)?;
1132 write!(
1133 self.out,
1134 " {}",
1135 &self.names[&NameKey::StructMember(handle, idx as u32)]
1136 )?;
1137 // Write [size]
1138 self.write_type(member.ty)?;
1139 // Newline is important
1140 writeln!(self.out, ";")?;
1141 }
1142 _ => {
1143 // Write the member type
1144 // Adds no trailing space
1145 self.write_type(member.ty)?;
1146
1147 // Write the member name and put a semicolon
1148 // The leading space is important
1149 // All members must have a semicolon even the last one
1150 writeln!(
1151 self.out,
1152 " {};",
1153 &self.names[&NameKey::StructMember(handle, idx as u32)]
1154 )?;
1155 }
1156 }
1157 }
1158
1159 write!(self.out, "}}")?;
1160
1161 if !block {
1162 writeln!(self.out, ";")?;
1163 // Add a newline for readability
1164 writeln!(self.out)?;
1165 }
1166
1167 Ok(())
1168 }
1169
1170 /// Helper method used to write statements
1171 ///
1172 /// # Notes
1173 /// Always adds a newline
write_stmt( &mut self, sta: &Statement, ctx: &FunctionCtx<'_>, indent: usize, ) -> BackendResult1174 fn write_stmt(
1175 &mut self,
1176 sta: &Statement,
1177 ctx: &FunctionCtx<'_>,
1178 indent: usize,
1179 ) -> BackendResult {
1180 match *sta {
1181 // This is where we can generate intermediate constants for some expression types.
1182 Statement::Emit(ref range) => {
1183 for handle in range.clone() {
1184 let min_ref_count = ctx.expressions[handle].bake_ref_count();
1185 if min_ref_count <= ctx.info[handle].ref_count {
1186 write!(self.out, "{}", INDENT.repeat(indent))?;
1187 match ctx.info[handle].ty {
1188 TypeResolution::Handle(ty_handle) => {
1189 match self.module.types[ty_handle].inner {
1190 TypeInner::Struct { .. } => {
1191 let ty_name = &self.names[&NameKey::Type(ty_handle)];
1192 write!(self.out, "{}", ty_name)?;
1193 }
1194 _ => {
1195 self.write_type(ty_handle)?;
1196 }
1197 }
1198 }
1199 TypeResolution::Value(ref inner) => {
1200 self.write_value_type(inner)?;
1201 }
1202 }
1203 let name = format!("_expr{}", handle.index());
1204 write!(self.out, " {} = ", name)?;
1205 self.write_expr(handle, ctx)?;
1206 writeln!(self.out, ";")?;
1207 self.cached_expressions.insert(handle, name);
1208 }
1209 }
1210 }
1211 // Blocks are simple we just need to write the block statements between braces
1212 // We could also just print the statements but this is more readable and maps more
1213 // closely to the IR
1214 Statement::Block(ref block) => {
1215 write!(self.out, "{}", INDENT.repeat(indent))?;
1216 writeln!(self.out, "{{")?;
1217 for sta in block.iter() {
1218 // Increase the indentation to help with readability
1219 self.write_stmt(sta, ctx, indent + 1)?
1220 }
1221 writeln!(self.out, "{}}}", INDENT.repeat(indent))?
1222 }
1223 // Ifs are written as in C:
1224 // ```
1225 // if(condition) {
1226 // accept
1227 // } else {
1228 // reject
1229 // }
1230 // ```
1231 Statement::If {
1232 condition,
1233 ref accept,
1234 ref reject,
1235 } => {
1236 write!(self.out, "{}", INDENT.repeat(indent))?;
1237 write!(self.out, "if(")?;
1238 self.write_expr(condition, ctx)?;
1239 writeln!(self.out, ") {{")?;
1240
1241 for sta in accept {
1242 // Increase indentation to help with readability
1243 self.write_stmt(sta, ctx, indent + 1)?;
1244 }
1245
1246 // If there are no statements in the reject block we skip writing it
1247 // This is only for readability
1248 if !reject.is_empty() {
1249 writeln!(self.out, "{}}} else {{", INDENT.repeat(indent))?;
1250
1251 for sta in reject {
1252 // Increase indentation to help with readability
1253 self.write_stmt(sta, ctx, indent + 1)?;
1254 }
1255 }
1256
1257 writeln!(self.out, "{}}}", INDENT.repeat(indent))?
1258 }
1259 // Switch are written as in C:
1260 // ```
1261 // switch (selector) {
1262 // // Fallthrough
1263 // case label:
1264 // block
1265 // // Non fallthrough
1266 // case label:
1267 // block
1268 // break;
1269 // default:
1270 // block
1271 // }
1272 // ```
1273 // Where the `default` case happens isn't important but we put it last
1274 // so that we don't need to print a `break` for it
1275 Statement::Switch {
1276 selector,
1277 ref cases,
1278 ref default,
1279 } => {
1280 // Start the switch
1281 write!(self.out, "{}", INDENT.repeat(indent))?;
1282 write!(self.out, "switch(")?;
1283 self.write_expr(selector, ctx)?;
1284 writeln!(self.out, ") {{")?;
1285
1286 // Write all cases
1287 for case in cases {
1288 writeln!(
1289 self.out,
1290 "{}case {}:",
1291 INDENT.repeat(indent + 1),
1292 case.value
1293 )?;
1294
1295 for sta in case.body.iter() {
1296 self.write_stmt(sta, ctx, indent + 2)?;
1297 }
1298
1299 // Write `break;` if the block isn't fallthrough
1300 if !case.fall_through {
1301 writeln!(self.out, "{}break;", INDENT.repeat(indent + 2))?;
1302 }
1303 }
1304
1305 // Only write the default block if the block isn't empty
1306 // Writing default without a block is valid but it's more readable this way
1307 if !default.is_empty() {
1308 writeln!(self.out, "{}default:", INDENT.repeat(indent + 1))?;
1309
1310 for sta in default {
1311 self.write_stmt(sta, ctx, indent + 2)?;
1312 }
1313 }
1314
1315 writeln!(self.out, "{}}}", INDENT.repeat(indent))?
1316 }
1317 // Loops in naga IR are based on wgsl loops, glsl can emulate the behaviour by using a
1318 // while true loop and appending the continuing block to the body resulting on:
1319 // ```
1320 // while(true) {
1321 // body
1322 // continuing
1323 // }
1324 // ```
1325 Statement::Loop {
1326 ref body,
1327 ref continuing,
1328 } => {
1329 write!(self.out, "{}", INDENT.repeat(indent))?;
1330 writeln!(self.out, "while(true) {{")?;
1331
1332 for sta in body.iter().chain(continuing.iter()) {
1333 self.write_stmt(sta, ctx, indent + 1)?;
1334 }
1335
1336 writeln!(self.out, "{}}}", INDENT.repeat(indent))?
1337 }
1338 // Break, continue and return as written as in C
1339 // `break;`
1340 Statement::Break => {
1341 write!(self.out, "{}", INDENT.repeat(indent))?;
1342 writeln!(self.out, "break;")?
1343 }
1344 // `continue;`
1345 Statement::Continue => {
1346 write!(self.out, "{}", INDENT.repeat(indent))?;
1347 writeln!(self.out, "continue;")?
1348 }
1349 // `return expr;`, `expr` is optional
1350 Statement::Return { value } => {
1351 write!(self.out, "{}", INDENT.repeat(indent))?;
1352 match ctx.func {
1353 FunctionType::Function(_) => {
1354 write!(self.out, "return")?;
1355 // Write the expression to be returned if needed
1356 if let Some(expr) = value {
1357 write!(self.out, " ")?;
1358 self.write_expr(expr, ctx)?;
1359 }
1360 writeln!(self.out, ";")?;
1361 }
1362 FunctionType::EntryPoint(ep_index) => {
1363 let ep = &self.module.entry_points[ep_index as usize];
1364 if let Some(ref result) = ep.function.result {
1365 let value = value.unwrap();
1366 match self.module.types[result.ty].inner {
1367 crate::TypeInner::Struct { ref members, .. } => {
1368 let (mut is_temp_struct_used, mut return_struct) = (false, "");
1369 if let Expression::Compose { .. } = ctx.expressions[value] {
1370 is_temp_struct_used = true;
1371 return_struct = "_tmp_return";
1372 write!(
1373 self.out,
1374 "{} {} = ",
1375 &self.names[&NameKey::Type(result.ty)],
1376 return_struct
1377 )?;
1378 self.write_expr(value, ctx)?;
1379 writeln!(self.out, ";")?;
1380 write!(self.out, "{}", INDENT.repeat(indent))?;
1381 }
1382 for (index, member) in members.iter().enumerate() {
1383 // TODO: handle builtin in better way
1384 if let Some(Binding::BuiltIn(builtin)) = member.binding {
1385 match builtin {
1386 crate::BuiltIn::ClipDistance
1387 | crate::BuiltIn::CullDistance
1388 | crate::BuiltIn::PointSize => {
1389 if self.options.version.is_es() {
1390 continue;
1391 }
1392 }
1393 _ => {}
1394 }
1395 }
1396
1397 let varying_name = VaryingName {
1398 binding: member.binding.as_ref().unwrap(),
1399 stage: ep.stage,
1400 output: true,
1401 };
1402 let field_name = self.names
1403 [&NameKey::StructMember(result.ty, index as u32)]
1404 .clone();
1405 write!(self.out, "{} = ", varying_name)?;
1406
1407 if !is_temp_struct_used {
1408 self.write_expr(value, ctx)?;
1409 }
1410
1411 writeln!(self.out, "{}.{};", return_struct, &field_name)?;
1412 write!(self.out, "{}", INDENT.repeat(indent))?;
1413 }
1414 }
1415 _ => {
1416 let name = VaryingName {
1417 binding: result.binding.as_ref().unwrap(),
1418 stage: ep.stage,
1419 output: true,
1420 };
1421 write!(self.out, "{} = ", name)?;
1422 self.write_expr(value, ctx)?;
1423 writeln!(self.out, ";")?;
1424 write!(self.out, "{}", INDENT.repeat(indent))?;
1425 }
1426 }
1427 }
1428 writeln!(self.out, "return;")?;
1429 }
1430 }
1431 }
1432 // This is one of the places were glsl adds to the syntax of C in this case the discard
1433 // keyword which ceases all further processing in a fragment shader, it's called OpKill
1434 // in spir-v that's why it's called `Statement::Kill`
1435 Statement::Kill => writeln!(self.out, "{}discard;", INDENT.repeat(indent))?,
1436 // Issue an execution or a memory barrier.
1437 Statement::Barrier(flags) => {
1438 if flags.is_empty() {
1439 writeln!(self.out, "{}barrier();", INDENT.repeat(indent))?;
1440 } else {
1441 writeln!(self.out, "{}groupMemoryBarrier();", INDENT.repeat(indent))?;
1442 }
1443 }
1444 // Stores in glsl are just variable assignments written as `pointer = value;`
1445 Statement::Store { pointer, value } => {
1446 write!(self.out, "{}", INDENT.repeat(indent))?;
1447 self.write_expr(pointer, ctx)?;
1448 write!(self.out, " = ")?;
1449 self.write_expr(value, ctx)?;
1450 writeln!(self.out, ";")?
1451 }
1452 // Stores a value into an image.
1453 Statement::ImageStore {
1454 image,
1455 coordinate,
1456 array_index,
1457 value,
1458 } => {
1459 write!(self.out, "{}", INDENT.repeat(indent))?;
1460 // This will only panic if the module is invalid
1461 let dim = match *ctx.info[image].ty.inner_with(&self.module.types) {
1462 TypeInner::Image { dim, .. } => dim,
1463 _ => unreachable!(),
1464 };
1465
1466 write!(self.out, "imageStore(")?;
1467 self.write_expr(image, ctx)?;
1468 write!(self.out, ", ")?;
1469 self.write_texture_coordinates(coordinate, array_index, dim, ctx)?;
1470 write!(self.out, ", ")?;
1471 self.write_expr(value, ctx)?;
1472 writeln!(self.out, ");")?;
1473 }
1474 // A `Call` is written `name(arguments)` where `arguments` is a comma separated expressions list
1475 Statement::Call {
1476 function,
1477 ref arguments,
1478 result,
1479 } => {
1480 write!(self.out, "{}", INDENT.repeat(indent))?;
1481 if let Some(expr) = result {
1482 let name = format!("_expr{}", expr.index());
1483 let result = self.module.functions[function].result.as_ref().unwrap();
1484 self.write_type(result.ty)?;
1485 write!(self.out, " {} = ", name)?;
1486 self.cached_expressions.insert(expr, name);
1487 }
1488 write!(self.out, "{}(", &self.names[&NameKey::Function(function)])?;
1489 self.write_slice(arguments, |this, _, arg| this.write_expr(*arg, ctx))?;
1490 writeln!(self.out, ");")?
1491 }
1492 }
1493
1494 Ok(())
1495 }
1496
1497 /// Helper method to write expressions
1498 ///
1499 /// # Notes
1500 /// Doesn't add any newlines or leading/trailing spaces
write_expr(&mut self, expr: Handle<Expression>, ctx: &FunctionCtx<'_>) -> BackendResult1501 fn write_expr(&mut self, expr: Handle<Expression>, ctx: &FunctionCtx<'_>) -> BackendResult {
1502 if let Some(name) = self.cached_expressions.get(&expr) {
1503 write!(self.out, "{}", name)?;
1504 return Ok(());
1505 }
1506 match ctx.expressions[expr] {
1507 // `Access` is applied to arrays, vectors and matrices and is written as indexing
1508 Expression::Access { base, index } => {
1509 self.write_expr(base, ctx)?;
1510 write!(self.out, "[")?;
1511 self.write_expr(index, ctx)?;
1512 write!(self.out, "]")?
1513 }
1514 // `AccessIndex` is the same as `Access` except that the index is a constant and it can
1515 // be applied to structs, in this case we need to find the name of the field at that
1516 // index and write `base.field_name`
1517 Expression::AccessIndex { base, index } => {
1518 self.write_expr(base, ctx)?;
1519
1520 let base_ty_res = &ctx.info[base].ty;
1521 let mut resolved = base_ty_res.inner_with(&self.module.types);
1522 let base_ty_handle = match *resolved {
1523 TypeInner::Pointer { base, class: _ } => {
1524 resolved = &self.module.types[base].inner;
1525 Some(base)
1526 }
1527 _ => base_ty_res.handle(),
1528 };
1529
1530 match *resolved {
1531 TypeInner::Vector { .. }
1532 | TypeInner::Matrix { .. }
1533 | TypeInner::Array { .. }
1534 | TypeInner::ValuePointer { .. } => write!(self.out, "[{}]", index)?,
1535 TypeInner::Struct { .. } => {
1536 // This will never panic in case the type is a `Struct`, this is not true
1537 // for other types so we can only check while inside this match arm
1538 let ty = base_ty_handle.unwrap();
1539
1540 write!(
1541 self.out,
1542 ".{}",
1543 &self.names[&NameKey::StructMember(ty, index)]
1544 )?
1545 }
1546 ref other => return Err(Error::Custom(format!("Cannot index {:?}", other))),
1547 }
1548 }
1549 // Constants are delegated to `write_constant`
1550 Expression::Constant(constant) => {
1551 self.write_constant(&self.module.constants[constant])?
1552 }
1553 // `Splat` needs to actually write down a vector, it's not always inferred in GLSL.
1554 Expression::Splat { size: _, value } => {
1555 let resolved = ctx.info[expr].ty.inner_with(&self.module.types);
1556 self.write_value_type(resolved)?;
1557 write!(self.out, "(")?;
1558 self.write_expr(value, ctx)?;
1559 write!(self.out, ")")?
1560 }
1561 // `Swizzle` adds a few letters behind the dot.
1562 Expression::Swizzle {
1563 size,
1564 vector,
1565 pattern,
1566 } => {
1567 self.write_expr(vector, ctx)?;
1568 write!(self.out, ".")?;
1569 for &sc in pattern[..size as usize].iter() {
1570 self.out.write_char(COMPONENTS[sc as usize])?;
1571 }
1572 }
1573 // `Compose` is pretty simple we just write `type(components)` where `components` is a
1574 // comma separated list of expressions
1575 Expression::Compose { ty, ref components } => {
1576 self.write_type(ty)?;
1577
1578 write!(self.out, "(")?;
1579 self.write_slice(components, |this, _, arg| this.write_expr(*arg, ctx))?;
1580 write!(self.out, ")")?
1581 }
1582 // Function arguments are written as the argument name
1583 Expression::FunctionArgument(pos) => {
1584 write!(self.out, "{}", &self.names[&ctx.argument_key(pos)])?
1585 }
1586 // Global variables need some special work for their name but
1587 // `get_global_name` does the work for us
1588 Expression::GlobalVariable(handle) => {
1589 let global = &self.module.global_variables[handle];
1590 write!(self.out, "{}", self.get_global_name(handle, global))?
1591 }
1592 // A local is written as it's name
1593 Expression::LocalVariable(handle) => {
1594 write!(self.out, "{}", self.names[&ctx.name_key(handle)])?
1595 }
1596 // glsl has no pointers so there's no load operation, just write the pointer expression
1597 Expression::Load { pointer } => self.write_expr(pointer, ctx)?,
1598 // `ImageSample` is a bit complicated compared to the rest of the IR.
1599 //
1600 // First there are three variations depending wether the sample level is explicitly set,
1601 // if it's automatic or it it's bias:
1602 // `texture(image, coordinate)` - Automatic sample level
1603 // `texture(image, coordinate, bias)` - Bias sample level
1604 // `textureLod(image, coordinate, level)` - Zero or Exact sample level
1605 //
1606 // Furthermore if `depth_ref` is some we need to append it to the coordinate vector
1607 Expression::ImageSample {
1608 image,
1609 sampler: _, //TODO
1610 coordinate,
1611 array_index,
1612 offset: _, //TODO
1613 level,
1614 depth_ref,
1615 } => {
1616 //TODO: handle MS
1617
1618 // textureLod on sampler2DArrayShadow and samplerCubeShadow does not exist in GLSL.
1619 // To emulate this, we will have to use textureGrad with a constant gradient of 0.
1620 let workaround_lod_array_shadow_as_grad =
1621 array_index.is_some() && depth_ref.is_some();
1622
1623 //Write the function to be used depending on the sample level
1624 let fun_name = match level {
1625 crate::SampleLevel::Auto | crate::SampleLevel::Bias(_) => "texture",
1626 crate::SampleLevel::Zero | crate::SampleLevel::Exact(_) => {
1627 if workaround_lod_array_shadow_as_grad {
1628 "textureGrad"
1629 } else {
1630 "textureLod"
1631 }
1632 }
1633 crate::SampleLevel::Gradient { .. } => "textureGrad",
1634 };
1635
1636 write!(self.out, "{}(", fun_name)?;
1637
1638 // Write the image that will be used
1639 self.write_expr(image, ctx)?;
1640 // The space here isn't required but it helps with readability
1641 write!(self.out, ", ")?;
1642
1643 // We need to get the coordinates vector size to later build a vector that's `size + 1`
1644 // if `depth_ref` is some, if it isn't a vector we panic as that's not a valid expression
1645 let size = match *ctx.info[coordinate].ty.inner_with(&self.module.types) {
1646 TypeInner::Vector { size, .. } => size,
1647 _ => unreachable!(),
1648 };
1649
1650 let mut coord_dim = size as u8;
1651 if array_index.is_some() {
1652 coord_dim += 1;
1653 }
1654 if depth_ref.is_some() {
1655 coord_dim += 1;
1656 }
1657
1658 // Compose a new texture coordinates vector
1659 write!(self.out, "vec{}(", coord_dim)?;
1660 self.write_expr(coordinate, ctx)?;
1661 if let Some(expr) = array_index {
1662 write!(self.out, ", ")?;
1663 self.write_expr(expr, ctx)?;
1664 }
1665 if let Some(expr) = depth_ref {
1666 write!(self.out, ", ")?;
1667 self.write_expr(expr, ctx)?;
1668 }
1669 write!(self.out, ")")?;
1670
1671 match level {
1672 // Auto needs no more arguments
1673 crate::SampleLevel::Auto => (),
1674 // Zero needs level set to 0
1675 crate::SampleLevel::Zero => {
1676 if workaround_lod_array_shadow_as_grad {
1677 write!(self.out, ", vec2(0, 0), vec2(0,0)")?;
1678 } else {
1679 write!(self.out, ", 0")?;
1680 }
1681 }
1682 // Exact and bias require another argument
1683 crate::SampleLevel::Exact(expr) => {
1684 if workaround_lod_array_shadow_as_grad {
1685 write!(self.out, ", vec2(0, 0), vec2(0,0)")?;
1686 } else {
1687 write!(self.out, ", ")?;
1688 self.write_expr(expr, ctx)?;
1689 }
1690 }
1691 crate::SampleLevel::Bias(expr) => {
1692 write!(self.out, ", ")?;
1693 self.write_expr(expr, ctx)?;
1694 }
1695 crate::SampleLevel::Gradient { x, y } => {
1696 write!(self.out, ", ")?;
1697 self.write_expr(x, ctx)?;
1698 write!(self.out, ", ")?;
1699 self.write_expr(y, ctx)?;
1700 }
1701 }
1702
1703 // End the function
1704 write!(self.out, ")")?
1705 }
1706 // `ImageLoad` is also a bit complicated.
1707 // There are two functions one for sampled
1708 // images another for storage images, the former uses `texelFetch` and the latter uses
1709 // `imageLoad`.
1710 // Furthermore we have `index` which is always `Some` for sampled images
1711 // and `None` for storage images, so we end up with two functions:
1712 // `texelFetch(image, coordinate, index)` - for sampled images
1713 // `imageLoad(image, coordinate)` - for storage images
1714 Expression::ImageLoad {
1715 image,
1716 coordinate,
1717 array_index,
1718 index,
1719 } => {
1720 // This will only panic if the module is invalid
1721 let (dim, class) = match *ctx.info[image].ty.inner_with(&self.module.types) {
1722 TypeInner::Image {
1723 dim,
1724 arrayed: _,
1725 class,
1726 } => (dim, class),
1727 _ => unreachable!(),
1728 };
1729
1730 let fun_name = match class {
1731 ImageClass::Sampled { .. } => "texelFetch",
1732 ImageClass::Storage(_) => "imageLoad",
1733 // TODO: Is there even a function for this?
1734 ImageClass::Depth => todo!(),
1735 };
1736
1737 write!(self.out, "{}(", fun_name)?;
1738 self.write_expr(image, ctx)?;
1739 write!(self.out, ", ")?;
1740 self.write_texture_coordinates(coordinate, array_index, dim, ctx)?;
1741
1742 if let Some(index_expr) = index {
1743 write!(self.out, ", ")?;
1744 self.write_expr(index_expr, ctx)?;
1745 }
1746 write!(self.out, ")")?;
1747 }
1748 // Query translates into one of the:
1749 // - textureSize/imageSize
1750 // - textureQueryLevels
1751 // - textureSamples/imageSamples
1752 Expression::ImageQuery { image, query } => {
1753 // This will only panic if the module is invalid
1754 let (dim, class) = match *ctx.info[image].ty.inner_with(&self.module.types) {
1755 TypeInner::Image {
1756 dim,
1757 arrayed: _,
1758 class,
1759 } => (dim, class),
1760 _ => unreachable!(),
1761 };
1762 let components = match dim {
1763 crate::ImageDimension::D1 => 1,
1764 crate::ImageDimension::D2 => 2,
1765 crate::ImageDimension::D3 => 3,
1766 crate::ImageDimension::Cube => 2,
1767 };
1768 match query {
1769 crate::ImageQuery::Size { level } => {
1770 match class {
1771 ImageClass::Sampled { .. } | ImageClass::Depth => {
1772 write!(self.out, "textureSize(")?;
1773 self.write_expr(image, ctx)?;
1774 write!(self.out, ",")?;
1775 if let Some(expr) = level {
1776 self.write_expr(expr, ctx)?;
1777 } else {
1778 write!(self.out, "0",)?;
1779 }
1780 }
1781 ImageClass::Storage(_) => {
1782 write!(self.out, "imageSize(")?;
1783 self.write_expr(image, ctx)?;
1784 }
1785 }
1786 write!(self.out, ").{}", &"xyz"[..components])?;
1787 }
1788 crate::ImageQuery::NumLevels => {
1789 write!(self.out, "textureQueryLevels(",)?;
1790 self.write_expr(image, ctx)?;
1791 write!(self.out, ")",)?;
1792 }
1793 crate::ImageQuery::NumLayers => {
1794 let fun_name = match class {
1795 ImageClass::Sampled { .. } | ImageClass::Depth => "textureSize",
1796 ImageClass::Storage(_) => "imageSize",
1797 };
1798 write!(self.out, "{}(", fun_name)?;
1799 self.write_expr(image, ctx)?;
1800 write!(self.out, ",0).{}", COMPONENTS[components])?;
1801 }
1802 crate::ImageQuery::NumSamples => {
1803 // assumes ARB_shader_texture_image_samples
1804 let fun_name = match class {
1805 ImageClass::Sampled { .. } | ImageClass::Depth => "textureSamples",
1806 ImageClass::Storage(_) => "imageSamples",
1807 };
1808 write!(self.out, "{}(", fun_name)?;
1809 self.write_expr(image, ctx)?;
1810 write!(self.out, ")",)?;
1811 }
1812 }
1813 return Err(Error::Custom("ImageQuery not implemented".to_string()));
1814 }
1815 // `Unary` is pretty straightforward
1816 // "-" - for `Negate`
1817 // "~" - for `Not` if it's an integer
1818 // "!" - for `Not` if it's a boolean
1819 //
1820 // We also wrap the everything in parentheses to avoid precedence issues
1821 Expression::Unary { op, expr } => {
1822 write!(
1823 self.out,
1824 "({} ",
1825 match op {
1826 UnaryOperator::Negate => "-",
1827 UnaryOperator::Not =>
1828 match *ctx.info[expr].ty.inner_with(&self.module.types) {
1829 TypeInner::Scalar {
1830 kind: ScalarKind::Sint,
1831 ..
1832 } => "~",
1833 TypeInner::Scalar {
1834 kind: ScalarKind::Uint,
1835 ..
1836 } => "~",
1837 TypeInner::Scalar {
1838 kind: ScalarKind::Bool,
1839 ..
1840 } => "!",
1841 ref other =>
1842 return Err(Error::Custom(format!(
1843 "Cannot apply not to type {:?}",
1844 other
1845 ))),
1846 },
1847 }
1848 )?;
1849
1850 self.write_expr(expr, ctx)?;
1851
1852 write!(self.out, ")")?
1853 }
1854 // `Binary` we just write `left op right`, except when dealing with
1855 // comparison operations on vectors as they are implemented with
1856 // builtin functions.
1857 // Once again we wrap everything in parentheses to avoid precedence issues
1858 Expression::Binary { op, left, right } => {
1859 // Holds `Some(function_name)` if the binary operation is
1860 // implemented as a function call
1861 let function = if let (&TypeInner::Vector { .. }, &TypeInner::Vector { .. }) = (
1862 ctx.info[left].ty.inner_with(&self.module.types),
1863 ctx.info[right].ty.inner_with(&self.module.types),
1864 ) {
1865 match op {
1866 BinaryOperator::Less => Some("lessThan"),
1867 BinaryOperator::LessEqual => Some("lessThanEqual"),
1868 BinaryOperator::Greater => Some("greaterThan"),
1869 BinaryOperator::GreaterEqual => Some("greaterThanEqual"),
1870 BinaryOperator::Equal => Some("equal"),
1871 BinaryOperator::NotEqual => Some("notEqual"),
1872 _ => None,
1873 }
1874 } else {
1875 None
1876 };
1877
1878 write!(self.out, "{}(", function.unwrap_or(""))?;
1879 self.write_expr(left, ctx)?;
1880
1881 if function.is_some() {
1882 write!(self.out, ",")?
1883 } else {
1884 write!(self.out, " {} ", super::binary_operation_str(op))?;
1885 }
1886
1887 self.write_expr(right, ctx)?;
1888
1889 write!(self.out, ")")?
1890 }
1891 // `Select` is written as `condition ? accept : reject`
1892 // We wrap everything in parentheses to avoid precedence issues
1893 Expression::Select {
1894 condition,
1895 accept,
1896 reject,
1897 } => {
1898 write!(self.out, "(")?;
1899 self.write_expr(condition, ctx)?;
1900 write!(self.out, " ? ")?;
1901 self.write_expr(accept, ctx)?;
1902 write!(self.out, " : ")?;
1903 self.write_expr(reject, ctx)?;
1904 write!(self.out, ")")?
1905 }
1906 // `Derivative` is a function call to a glsl provided function
1907 Expression::Derivative { axis, expr } => {
1908 write!(
1909 self.out,
1910 "{}(",
1911 match axis {
1912 DerivativeAxis::X => "dFdx",
1913 DerivativeAxis::Y => "dFdy",
1914 DerivativeAxis::Width => "fwidth",
1915 }
1916 )?;
1917 self.write_expr(expr, ctx)?;
1918 write!(self.out, ")")?
1919 }
1920 // `Relational` is a normal function call to some glsl provided functions
1921 Expression::Relational { fun, argument } => {
1922 let fun_name = match fun {
1923 // There's no specific function for this but we can invert the result of `isinf`
1924 RelationalFunction::IsFinite => "!isinf",
1925 RelationalFunction::IsInf => "isinf",
1926 RelationalFunction::IsNan => "isnan",
1927 // There's also no function for this but we can invert `isnan`
1928 RelationalFunction::IsNormal => "!isnan",
1929 RelationalFunction::All => "all",
1930 RelationalFunction::Any => "any",
1931 };
1932 write!(self.out, "{}(", fun_name)?;
1933
1934 self.write_expr(argument, ctx)?;
1935
1936 write!(self.out, ")")?
1937 }
1938 Expression::Math {
1939 fun,
1940 arg,
1941 arg1,
1942 arg2,
1943 } => {
1944 use crate::MathFunction as Mf;
1945
1946 let fun_name = match fun {
1947 // comparison
1948 Mf::Abs => "abs",
1949 Mf::Min => "min",
1950 Mf::Max => "max",
1951 Mf::Clamp => "clamp",
1952 // trigonometry
1953 Mf::Cos => "cos",
1954 Mf::Cosh => "cosh",
1955 Mf::Sin => "sin",
1956 Mf::Sinh => "sinh",
1957 Mf::Tan => "tan",
1958 Mf::Tanh => "tanh",
1959 Mf::Acos => "acos",
1960 Mf::Asin => "asin",
1961 Mf::Atan => "atan",
1962 // glsl doesn't have atan2 function
1963 // use two-argument variation of the atan function
1964 Mf::Atan2 => "atan",
1965 // decomposition
1966 Mf::Ceil => "ceil",
1967 Mf::Floor => "floor",
1968 Mf::Round => "round",
1969 Mf::Fract => "fract",
1970 Mf::Trunc => "trunc",
1971 Mf::Modf => "modf",
1972 Mf::Frexp => "frexp",
1973 Mf::Ldexp => "ldexp",
1974 // exponent
1975 Mf::Exp => "exp",
1976 Mf::Exp2 => "exp2",
1977 Mf::Log => "log",
1978 Mf::Log2 => "log2",
1979 Mf::Pow => "pow",
1980 // geometry
1981 Mf::Dot => "dot",
1982 Mf::Outer => "outerProduct",
1983 Mf::Cross => "cross",
1984 Mf::Distance => "distance",
1985 Mf::Length => "length",
1986 Mf::Normalize => "normalize",
1987 Mf::FaceForward => "faceforward",
1988 Mf::Reflect => "reflect",
1989 Mf::Refract => "refract",
1990 // computational
1991 Mf::Sign => "sign",
1992 Mf::Fma => "fma",
1993 Mf::Mix => "mix",
1994 Mf::Step => "step",
1995 Mf::SmoothStep => "smoothstep",
1996 Mf::Sqrt => "sqrt",
1997 Mf::InverseSqrt => "inversesqrt",
1998 Mf::Inverse => "inverse",
1999 Mf::Transpose => "transpose",
2000 Mf::Determinant => "determinant",
2001 // bits
2002 Mf::CountOneBits => "bitCount",
2003 Mf::ReverseBits => "bitfieldReverse",
2004 };
2005
2006 write!(self.out, "{}(", fun_name)?;
2007 self.write_expr(arg, ctx)?;
2008 if let Some(arg) = arg1 {
2009 write!(self.out, ", ")?;
2010 self.write_expr(arg, ctx)?;
2011 }
2012 if let Some(arg) = arg2 {
2013 write!(self.out, ", ")?;
2014 self.write_expr(arg, ctx)?;
2015 }
2016 write!(self.out, ")")?
2017 }
2018 // `As` is always a call.
2019 // If `convert` is true the function name is the type
2020 // Else the function name is one of the glsl provided bitcast functions
2021 Expression::As {
2022 expr,
2023 kind: target_kind,
2024 convert,
2025 } => {
2026 let inner = ctx.info[expr].ty.inner_with(&self.module.types);
2027 match convert {
2028 Some(width) => {
2029 // this is similar to `write_type`, but with the target kind
2030 let scalar = glsl_scalar(target_kind, width)?;
2031 match *inner {
2032 TypeInner::Vector { size, .. } => {
2033 write!(self.out, "{}vec{}", scalar.prefix, size as u8)?
2034 }
2035 _ => write!(self.out, "{}", scalar.full)?,
2036 }
2037 }
2038 None => {
2039 let source_kind = inner.scalar_kind().unwrap();
2040 write!(
2041 self.out,
2042 "{}",
2043 match (source_kind, target_kind) {
2044 (ScalarKind::Float, ScalarKind::Sint) => "floatBitsToInt",
2045 (ScalarKind::Float, ScalarKind::Uint) => "floatBitsToUInt",
2046 (ScalarKind::Sint, ScalarKind::Float) => "intBitsToFloat",
2047 (ScalarKind::Uint, ScalarKind::Float) => "uintBitsToFloat",
2048 // There is no way to bitcast between Uint/Sint in glsl. Use constructor conversion
2049 (ScalarKind::Uint, ScalarKind::Sint) => "int",
2050 (ScalarKind::Sint, ScalarKind::Uint) => "uint",
2051 _ => {
2052 return Err(Error::Custom(format!(
2053 "Cannot bitcast {:?} to {:?}",
2054 source_kind, target_kind
2055 )));
2056 }
2057 }
2058 )?;
2059 }
2060 }
2061
2062 write!(self.out, "(")?;
2063 self.write_expr(expr, ctx)?;
2064 write!(self.out, ")")?
2065 }
2066 Expression::Call(_function) => unreachable!(),
2067 // `ArrayLength` is written as `expr.length()` and we convert it to a uint
2068 Expression::ArrayLength(expr) => {
2069 write!(self.out, "uint(")?;
2070 self.write_expr(expr, ctx)?;
2071 write!(self.out, ".length())")?
2072 }
2073 }
2074
2075 Ok(())
2076 }
2077
write_texture_coordinates( &mut self, coordinate: Handle<Expression>, array_index: Option<Handle<Expression>>, dim: crate::ImageDimension, ctx: &FunctionCtx, ) -> Result<(), Error>2078 fn write_texture_coordinates(
2079 &mut self,
2080 coordinate: Handle<Expression>,
2081 array_index: Option<Handle<Expression>>,
2082 dim: crate::ImageDimension,
2083 ctx: &FunctionCtx,
2084 ) -> Result<(), Error> {
2085 match array_index {
2086 Some(layer_expr) => {
2087 let tex_coord_type = match dim {
2088 crate::ImageDimension::D1 => "ivec2",
2089 crate::ImageDimension::D2 => "ivec3",
2090 crate::ImageDimension::D3 => "ivec4",
2091 crate::ImageDimension::Cube => "ivec4",
2092 };
2093 write!(self.out, "{}(", tex_coord_type)?;
2094 self.write_expr(coordinate, ctx)?;
2095 write!(self.out, ", ")?;
2096 self.write_expr(layer_expr, ctx)?;
2097 write!(self.out, ")")?;
2098 }
2099 None => {
2100 self.write_expr(coordinate, ctx)?;
2101 }
2102 }
2103 Ok(())
2104 }
2105
2106 /// Helper method used to produce the reflection info that's returned to the user
2107 ///
2108 /// It takes an iterator of [`Function`](crate::Function) references instead of
2109 /// [`Handle`](crate::arena::Handle) because [`EntryPoint`](crate::EntryPoint) isn't in any
2110 /// [`Arena`](crate::arena::Arena) and we need to traverse it
collect_reflection_info(&self) -> Result<ReflectionInfo, Error>2111 fn collect_reflection_info(&self) -> Result<ReflectionInfo, Error> {
2112 use std::collections::hash_map::Entry;
2113 let info = self.info.get_entry_point(self.entry_point_idx as usize);
2114 let mut mappings = FastHashMap::default();
2115 let mut uniforms = FastHashMap::default();
2116
2117 for sampling in info.sampling_set.iter() {
2118 let global = self.module.global_variables[sampling.image].clone();
2119 let tex_name = self.reflection_names[&global.ty].clone();
2120
2121 match mappings.entry(tex_name) {
2122 Entry::Vacant(v) => {
2123 v.insert(TextureMapping {
2124 texture: sampling.image,
2125 sampler: Some(sampling.sampler),
2126 });
2127 }
2128 Entry::Occupied(e) => {
2129 if e.get().sampler != Some(sampling.sampler) {
2130 log::error!("Conflicting samplers for {}", e.key());
2131 return Err(Error::ImageMultipleSamplers);
2132 }
2133 }
2134 }
2135 }
2136
2137 for (handle, var) in self.module.global_variables.iter() {
2138 if info[handle].is_empty() {
2139 continue;
2140 }
2141 match self.module.types[var.ty].inner {
2142 crate::TypeInner::Struct { .. } => match var.class {
2143 StorageClass::Uniform | StorageClass::Storage => {
2144 let name = self.reflection_names[&var.ty].clone();
2145 uniforms.insert(handle, name);
2146 }
2147 _ => (),
2148 },
2149 _ => continue,
2150 }
2151 }
2152
2153 Ok(ReflectionInfo {
2154 texture_mapping: mappings,
2155 uniforms,
2156 })
2157 }
2158 }
2159
2160 /// Structure returned by [`glsl_scalar`](glsl_scalar)
2161 ///
2162 /// It contains both a prefix used in other types and the full type name
2163 struct ScalarString<'a> {
2164 /// The prefix used to compose other types
2165 prefix: &'a str,
2166 /// The name of the scalar type
2167 full: &'a str,
2168 }
2169
2170 /// Helper function that returns scalar related strings
2171 ///
2172 /// Check [`ScalarString`](ScalarString) for the information provided
2173 ///
2174 /// # Errors
2175 /// If a [`Float`](crate::ScalarKind::Float) with an width that isn't 4 or 8
glsl_scalar(kind: ScalarKind, width: Bytes) -> Result<ScalarString<'static>, Error>2176 fn glsl_scalar(kind: ScalarKind, width: Bytes) -> Result<ScalarString<'static>, Error> {
2177 Ok(match kind {
2178 ScalarKind::Sint => ScalarString {
2179 prefix: "i",
2180 full: "int",
2181 },
2182 ScalarKind::Uint => ScalarString {
2183 prefix: "u",
2184 full: "uint",
2185 },
2186 ScalarKind::Float => match width {
2187 4 => ScalarString {
2188 prefix: "",
2189 full: "float",
2190 },
2191 8 => ScalarString {
2192 prefix: "d",
2193 full: "double",
2194 },
2195 _ => return Err(Error::UnsupportedScalar(kind, width)),
2196 },
2197 ScalarKind::Bool => ScalarString {
2198 prefix: "b",
2199 full: "bool",
2200 },
2201 })
2202 }
2203
2204 /// Helper function that returns the glsl variable name for a builtin
glsl_built_in(built_in: BuiltIn, output: bool) -> &'static str2205 fn glsl_built_in(built_in: BuiltIn, output: bool) -> &'static str {
2206 match built_in {
2207 BuiltIn::Position => {
2208 if output {
2209 "gl_Position"
2210 } else {
2211 "gl_FragCoord"
2212 }
2213 }
2214 // vertex
2215 BuiltIn::BaseInstance => "uint(gl_BaseInstance)",
2216 BuiltIn::BaseVertex => "uint(gl_BaseVertex)",
2217 BuiltIn::ClipDistance => "gl_ClipDistance",
2218 BuiltIn::CullDistance => "gl_CullDistance",
2219 BuiltIn::InstanceIndex => "uint(gl_InstanceID)",
2220 BuiltIn::PointSize => "gl_PointSize",
2221 BuiltIn::VertexIndex => "uint(gl_VertexID)",
2222 // fragment
2223 BuiltIn::FragDepth => "gl_FragDepth",
2224 BuiltIn::FrontFacing => "gl_FrontFacing",
2225 BuiltIn::SampleIndex => "gl_SampleID",
2226 BuiltIn::SampleMask => {
2227 if output {
2228 "gl_SampleMask"
2229 } else {
2230 "gl_SampleMaskIn"
2231 }
2232 }
2233 // compute
2234 BuiltIn::GlobalInvocationId => "gl_GlobalInvocationID",
2235 BuiltIn::LocalInvocationId => "gl_LocalInvocationID",
2236 BuiltIn::LocalInvocationIndex => "gl_LocalInvocationIndex",
2237 BuiltIn::WorkGroupId => "gl_WorkGroupID",
2238 BuiltIn::WorkGroupSize => "gl_WorkGroupSize",
2239 }
2240 }
2241
2242 /// Helper function that returns the string corresponding to the storage class
glsl_storage_class(class: StorageClass) -> Option<&'static str>2243 fn glsl_storage_class(class: StorageClass) -> Option<&'static str> {
2244 match class {
2245 StorageClass::Function => None,
2246 StorageClass::Private => None,
2247 StorageClass::Storage => Some("buffer"),
2248 StorageClass::Uniform => Some("uniform"),
2249 StorageClass::Handle => Some("uniform"),
2250 StorageClass::WorkGroup => Some("shared"),
2251 StorageClass::PushConstant => None,
2252 }
2253 }
2254
2255 /// Helper function that returns the string corresponding to the glsl interpolation qualifier
glsl_interpolation(interpolation: Interpolation) -> &'static str2256 fn glsl_interpolation(interpolation: Interpolation) -> &'static str {
2257 match interpolation {
2258 Interpolation::Perspective => "smooth",
2259 Interpolation::Linear => "noperspective",
2260 Interpolation::Flat => "flat",
2261 }
2262 }
2263
2264 /// Return the GLSL auxiliary qualifier for the given sampling value.
glsl_sampling(sampling: Sampling) -> Option<&'static str>2265 fn glsl_sampling(sampling: Sampling) -> Option<&'static str> {
2266 match sampling {
2267 Sampling::Center => None,
2268 Sampling::Centroid => Some("centroid"),
2269 Sampling::Sample => Some("sample"),
2270 }
2271 }
2272
2273 /// Helper function that returns the glsl dimension string of [`ImageDimension`](crate::ImageDimension)
glsl_dimension(dim: crate::ImageDimension) -> &'static str2274 fn glsl_dimension(dim: crate::ImageDimension) -> &'static str {
2275 match dim {
2276 crate::ImageDimension::D1 => "1D",
2277 crate::ImageDimension::D2 => "2D",
2278 crate::ImageDimension::D3 => "3D",
2279 crate::ImageDimension::Cube => "Cube",
2280 }
2281 }
2282
2283 /// Helper function that returns the glsl storage format string of [`StorageFormat`](crate::StorageFormat)
glsl_storage_format(format: StorageFormat) -> &'static str2284 fn glsl_storage_format(format: StorageFormat) -> &'static str {
2285 match format {
2286 StorageFormat::R8Unorm => "r8",
2287 StorageFormat::R8Snorm => "r8_snorm",
2288 StorageFormat::R8Uint => "r8ui",
2289 StorageFormat::R8Sint => "r8i",
2290 StorageFormat::R16Uint => "r16ui",
2291 StorageFormat::R16Sint => "r16i",
2292 StorageFormat::R16Float => "r16f",
2293 StorageFormat::Rg8Unorm => "rg8",
2294 StorageFormat::Rg8Snorm => "rg8_snorm",
2295 StorageFormat::Rg8Uint => "rg8ui",
2296 StorageFormat::Rg8Sint => "rg8i",
2297 StorageFormat::R32Uint => "r32ui",
2298 StorageFormat::R32Sint => "r32i",
2299 StorageFormat::R32Float => "r32f",
2300 StorageFormat::Rg16Uint => "rg16ui",
2301 StorageFormat::Rg16Sint => "rg16i",
2302 StorageFormat::Rg16Float => "rg16f",
2303 StorageFormat::Rgba8Unorm => "rgba8ui",
2304 StorageFormat::Rgba8Snorm => "rgba8_snorm",
2305 StorageFormat::Rgba8Uint => "rgba8ui",
2306 StorageFormat::Rgba8Sint => "rgba8i",
2307 StorageFormat::Rgb10a2Unorm => "rgb10_a2ui",
2308 StorageFormat::Rg11b10Float => "r11f_g11f_b10f",
2309 StorageFormat::Rg32Uint => "rg32ui",
2310 StorageFormat::Rg32Sint => "rg32i",
2311 StorageFormat::Rg32Float => "rg32f",
2312 StorageFormat::Rgba16Uint => "rgba16ui",
2313 StorageFormat::Rgba16Sint => "rgba16i",
2314 StorageFormat::Rgba16Float => "rgba16f",
2315 StorageFormat::Rgba32Uint => "rgba32ui",
2316 StorageFormat::Rgba32Sint => "rgba32i",
2317 StorageFormat::Rgba32Float => "rgba32f",
2318 }
2319 }
2320
2321 /// Helper function that return the glsl storage access string of [`StorageAccess`](crate::StorageAccess)
2322 ///
2323 /// glsl allows adding both `readonly` and `writeonly` but this means that
2324 /// they can only be used to query information about the resource which isn't what
2325 /// we want here so when storage access is both `LOAD` and `STORE` add no modifiers
glsl_storage_access(storage_access: StorageAccess) -> Option<&'static str>2326 fn glsl_storage_access(storage_access: StorageAccess) -> Option<&'static str> {
2327 if storage_access == StorageAccess::LOAD {
2328 Some("readonly")
2329 } else if storage_access == StorageAccess::STORE {
2330 Some("writeonly")
2331 } else {
2332 None
2333 }
2334 }
2335
2336 /// Helper function that return string with default zero initialization for supported types
zero_init_value_str(inner: &TypeInner) -> Option<String>2337 fn zero_init_value_str(inner: &TypeInner) -> Option<String> {
2338 match *inner {
2339 TypeInner::Scalar { kind, .. } => match kind {
2340 ScalarKind::Bool => Some(String::from("false")),
2341 _ => Some(String::from("0")),
2342 },
2343 TypeInner::Vector { size, kind, width } => {
2344 if let Ok(scalar_string) = glsl_scalar(kind, width) {
2345 let vec_type = format!("{}vec{}", scalar_string.prefix, size as u8);
2346 match size {
2347 crate::VectorSize::Bi => Some(format!("{}(0, 0)", vec_type)),
2348 crate::VectorSize::Tri => Some(format!("{}(0, 0, 0)", vec_type)),
2349 crate::VectorSize::Quad => Some(format!("{}(0, 0, 0, 0)", vec_type)),
2350 }
2351 } else {
2352 None
2353 }
2354 }
2355 _ => None,
2356 }
2357 }
2358