naga/back/glsl/mod.rs
1/*!
2Backend for [GLSL][glsl] (OpenGL Shading Language).
3
4The main structure is [`Writer`], it maintains internal state that is used
5to output a [`Module`](crate::Module) into glsl
6
7# Supported versions
8### Core
9- 330
10- 400
11- 410
12- 420
13- 430
14- 450
15
16### ES
17- 300
18- 310
19
20[glsl]: https://www.khronos.org/registry/OpenGL/index_gl.php
21*/
22
23// GLSL is mostly a superset of C but it also removes some parts of it this is a list of relevant
24// aspects for this backend.
25//
26// The most notable change is the introduction of the version preprocessor directive that must
27// always be the first line of a glsl file and is written as
28// `#version number profile`
29// `number` is the version itself (i.e. 300) and `profile` is the
30// shader profile we only support "core" and "es", the former is used in desktop applications and
31// the later is used in embedded contexts, mobile devices and browsers. Each one as it's own
32// versions (at the time of writing this the latest version for "core" is 460 and for "es" is 320)
33//
34// Other important preprocessor addition is the extension directive which is written as
35// `#extension name: behaviour`
36// Extensions provide increased features in a plugin fashion but they aren't required to be
37// supported hence why they are called extensions, that's why `behaviour` is used it specifies
38// whether the extension is strictly required or if it should only be enabled if needed. In our case
39// when we use extensions we set behaviour to `require` always.
40//
41// The only thing that glsl removes that makes a difference are pointers.
42//
43// Additions that are relevant for the backend are the discard keyword, the introduction of
44// vector, matrices, samplers, image types and functions that provide common shader operations
45
46pub use features::Features;
47
48use crate::{
49 back,
50 proc::{self, NameKey},
51 valid, Handle, ShaderStage, TypeInner,
52};
53use features::FeaturesManager;
54use std::{
55 cmp::Ordering,
56 fmt,
57 fmt::{Error as FmtError, Write},
58};
59use thiserror::Error;
60
61/// Contains the features related code and the features querying method
62mod features;
63/// Contains a constant with a slice of all the reserved keywords RESERVED_KEYWORDS
64mod keywords;
65
66/// List of supported `core` GLSL versions.
67pub const SUPPORTED_CORE_VERSIONS: &[u16] = &[330, 400, 410, 420, 430, 440, 450];
68/// List of supported `es` GLSL versions.
69pub const SUPPORTED_ES_VERSIONS: &[u16] = &[300, 310, 320];
70
71/// The suffix of the variable that will hold the calculated clamped level
72/// of detail for bounds checking in `ImageLoad`
73const CLAMPED_LOD_SUFFIX: &str = "_clamped_lod";
74
75/// Mapping between resources and bindings.
76pub type BindingMap = std::collections::BTreeMap<crate::ResourceBinding, u8>;
77
78impl crate::AtomicFunction {
79 const fn to_glsl(self) -> &'static str {
80 match self {
81 Self::Add | Self::Subtract => "Add",
82 Self::And => "And",
83 Self::InclusiveOr => "Or",
84 Self::ExclusiveOr => "Xor",
85 Self::Min => "Min",
86 Self::Max => "Max",
87 Self::Exchange { compare: None } => "Exchange",
88 Self::Exchange { compare: Some(_) } => "", //TODO
89 }
90 }
91}
92
93impl crate::AddressSpace {
94 const fn is_buffer(&self) -> bool {
95 match *self {
96 crate::AddressSpace::Uniform | crate::AddressSpace::Storage { .. } => true,
97 _ => false,
98 }
99 }
100
101 /// Whether a variable with this address space can be initialized
102 const fn initializable(&self) -> bool {
103 match *self {
104 crate::AddressSpace::Function | crate::AddressSpace::Private => true,
105 crate::AddressSpace::WorkGroup
106 | crate::AddressSpace::Uniform
107 | crate::AddressSpace::Storage { .. }
108 | crate::AddressSpace::Handle
109 | crate::AddressSpace::PushConstant => false,
110 }
111 }
112}
113
114/// A GLSL version.
115#[derive(Debug, Copy, Clone, PartialEq)]
116#[cfg_attr(feature = "serialize", derive(serde::Serialize))]
117#[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
118pub enum Version {
119 /// `core` GLSL.
120 Desktop(u16),
121 /// `es` GLSL.
122 Embedded { version: u16, is_webgl: bool },
123}
124
125impl Version {
126 /// Create a new gles version
127 pub const fn new_gles(version: u16) -> Self {
128 Self::Embedded {
129 version,
130 is_webgl: false,
131 }
132 }
133
134 /// Returns true if self is `Version::Embedded` (i.e. is a es version)
135 const fn is_es(&self) -> bool {
136 match *self {
137 Version::Desktop(_) => false,
138 Version::Embedded { .. } => true,
139 }
140 }
141
142 /// Returns true if targetting WebGL
143 const fn is_webgl(&self) -> bool {
144 match *self {
145 Version::Desktop(_) => false,
146 Version::Embedded { is_webgl, .. } => is_webgl,
147 }
148 }
149
150 /// Checks the list of currently supported versions and returns true if it contains the
151 /// specified version
152 ///
153 /// # Notes
154 /// As an invalid version number will never be added to the supported version list
155 /// so this also checks for version validity
156 fn is_supported(&self) -> bool {
157 match *self {
158 Version::Desktop(v) => SUPPORTED_CORE_VERSIONS.contains(&v),
159 Version::Embedded { version: v, .. } => SUPPORTED_ES_VERSIONS.contains(&v),
160 }
161 }
162
163 /// Checks if the version supports all of the explicit layouts:
164 /// - `location=` qualifiers for bindings
165 /// - `binding=` qualifiers for resources
166 ///
167 /// Note: `location=` for vertex inputs and fragment outputs is supported
168 /// unconditionally for GLES 300.
169 fn supports_explicit_locations(&self) -> bool {
170 *self >= Version::Desktop(410) || *self >= Version::new_gles(310)
171 }
172
173 fn supports_early_depth_test(&self) -> bool {
174 *self >= Version::Desktop(130) || *self >= Version::new_gles(310)
175 }
176
177 fn supports_std430_layout(&self) -> bool {
178 *self >= Version::Desktop(430) || *self >= Version::new_gles(310)
179 }
180
181 fn supports_fma_function(&self) -> bool {
182 *self >= Version::Desktop(400) || *self >= Version::new_gles(320)
183 }
184
185 fn supports_integer_functions(&self) -> bool {
186 *self >= Version::Desktop(400) || *self >= Version::new_gles(310)
187 }
188
189 fn supports_derivative_control(&self) -> bool {
190 *self >= Version::Desktop(450)
191 }
192}
193
194impl PartialOrd for Version {
195 fn partial_cmp(&self, other: &Self) -> Option<Ordering> {
196 match (*self, *other) {
197 (Version::Desktop(x), Version::Desktop(y)) => Some(x.cmp(&y)),
198 (Version::Embedded { version: x, .. }, Version::Embedded { version: y, .. }) => {
199 Some(x.cmp(&y))
200 }
201 _ => None,
202 }
203 }
204}
205
206impl fmt::Display for Version {
207 fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
208 match *self {
209 Version::Desktop(v) => write!(f, "{v} core"),
210 Version::Embedded { version: v, .. } => write!(f, "{v} es"),
211 }
212 }
213}
214
215bitflags::bitflags! {
216 /// Configuration flags for the [`Writer`].
217 #[cfg_attr(feature = "serialize", derive(serde::Serialize))]
218 #[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
219 #[derive(Clone, Copy, Debug, Eq, PartialEq)]
220 pub struct WriterFlags: u32 {
221 /// Flip output Y and extend Z from (0, 1) to (-1, 1).
222 const ADJUST_COORDINATE_SPACE = 0x1;
223 /// Supports GL_EXT_texture_shadow_lod on the host, which provides
224 /// additional functions on shadows and arrays of shadows.
225 const TEXTURE_SHADOW_LOD = 0x2;
226 /// Include unused global variables, constants and functions. By default the output will exclude
227 /// global variables that are not used in the specified entrypoint (including indirect use),
228 /// all constant declarations, and functions that use excluded global variables.
229 const INCLUDE_UNUSED_ITEMS = 0x4;
230 /// Emit `PointSize` output builtin to vertex shaders, which is
231 /// required for drawing with `PointList` topology.
232 ///
233 /// https://registry.khronos.org/OpenGL/specs/es/3.2/GLSL_ES_Specification_3.20.html#built-in-language-variables
234 /// The variable gl_PointSize is intended for a shader to write the size of the point to be rasterized. It is measured in pixels.
235 /// If gl_PointSize is not written to, its value is undefined in subsequent pipe stages.
236 const FORCE_POINT_SIZE = 0x10;
237 }
238}
239
240/// Configuration used in the [`Writer`].
241#[derive(Debug, Clone)]
242#[cfg_attr(feature = "serialize", derive(serde::Serialize))]
243#[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
244pub struct Options {
245 /// The GLSL version to be used.
246 pub version: Version,
247 /// Configuration flags for the [`Writer`].
248 pub writer_flags: WriterFlags,
249 /// Map of resources association to binding locations.
250 pub binding_map: BindingMap,
251 /// Should workgroup variables be zero initialized (by polyfilling)?
252 pub zero_initialize_workgroup_memory: bool,
253}
254
255impl Default for Options {
256 fn default() -> Self {
257 Options {
258 version: Version::new_gles(310),
259 writer_flags: WriterFlags::ADJUST_COORDINATE_SPACE,
260 binding_map: BindingMap::default(),
261 zero_initialize_workgroup_memory: true,
262 }
263 }
264}
265
266/// A subset of options meant to be changed per pipeline.
267#[derive(Debug, Clone, PartialEq, Eq, Hash)]
268#[cfg_attr(feature = "serialize", derive(serde::Serialize))]
269#[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
270pub struct PipelineOptions {
271 /// The stage of the entry point.
272 pub shader_stage: ShaderStage,
273 /// The name of the entry point.
274 ///
275 /// If no entry point that matches is found while creating a [`Writer`], a error will be thrown.
276 pub entry_point: String,
277 /// How many views to render to, if doing multiview rendering.
278 pub multiview: Option<std::num::NonZeroU32>,
279}
280
281/// Reflection info for texture mappings and uniforms.
282pub struct ReflectionInfo {
283 /// Mapping between texture names and variables/samplers.
284 pub texture_mapping: crate::FastHashMap<String, TextureMapping>,
285 /// Mapping between uniform variables and names.
286 pub uniforms: crate::FastHashMap<Handle<crate::GlobalVariable>, String>,
287}
288
289/// Mapping between a texture and its sampler, if it exists.
290///
291/// GLSL pre-Vulkan has no concept of separate textures and samplers. Instead, everything is a
292/// `gsamplerN` where `g` is the scalar type and `N` is the dimension. But naga uses separate textures
293/// and samplers in the IR, so the backend produces a [`FastHashMap`](crate::FastHashMap) with the texture name
294/// as a key and a [`TextureMapping`] as a value. This way, the user knows where to bind.
295///
296/// [`Storage`](crate::ImageClass::Storage) images produce `gimageN` and don't have an associated sampler,
297/// so the [`sampler`](Self::sampler) field will be [`None`].
298#[derive(Debug, Clone)]
299pub struct TextureMapping {
300 /// Handle to the image global variable.
301 pub texture: Handle<crate::GlobalVariable>,
302 /// Handle to the associated sampler global variable, if it exists.
303 pub sampler: Option<Handle<crate::GlobalVariable>>,
304}
305
306/// Helper structure that generates a number
307#[derive(Default)]
308struct IdGenerator(u32);
309
310impl IdGenerator {
311 /// Generates a number that's guaranteed to be unique for this `IdGenerator`
312 fn generate(&mut self) -> u32 {
313 // It's just an increasing number but it does the job
314 let ret = self.0;
315 self.0 += 1;
316 ret
317 }
318}
319
320/// Helper wrapper used to get a name for a varying
321///
322/// Varying have different naming schemes depending on their binding:
323/// - Varyings with builtin bindings get the from [`glsl_built_in`](glsl_built_in).
324/// - Varyings with location bindings are named `_S_location_X` where `S` is a
325/// prefix identifying which pipeline stage the varying connects, and `X` is
326/// the location.
327struct VaryingName<'a> {
328 binding: &'a crate::Binding,
329 stage: ShaderStage,
330 output: bool,
331 targetting_webgl: bool,
332}
333impl fmt::Display for VaryingName<'_> {
334 fn fmt(&self, f: &mut fmt::Formatter) -> fmt::Result {
335 match *self.binding {
336 crate::Binding::Location { location, .. } => {
337 let prefix = match (self.stage, self.output) {
338 (ShaderStage::Compute, _) => unreachable!(),
339 // pipeline to vertex
340 (ShaderStage::Vertex, false) => "p2vs",
341 // vertex to fragment
342 (ShaderStage::Vertex, true) | (ShaderStage::Fragment, false) => "vs2fs",
343 // fragment to pipeline
344 (ShaderStage::Fragment, true) => "fs2p",
345 };
346 write!(f, "_{prefix}_location{location}",)
347 }
348 crate::Binding::BuiltIn(built_in) => {
349 write!(
350 f,
351 "{}",
352 glsl_built_in(built_in, self.output, self.targetting_webgl)
353 )
354 }
355 }
356 }
357}
358
359impl ShaderStage {
360 const fn to_str(self) -> &'static str {
361 match self {
362 ShaderStage::Compute => "cs",
363 ShaderStage::Fragment => "fs",
364 ShaderStage::Vertex => "vs",
365 }
366 }
367}
368
369/// Shorthand result used internally by the backend
370type BackendResult<T = ()> = Result<T, Error>;
371
372/// A GLSL compilation error.
373#[derive(Debug, Error)]
374pub enum Error {
375 /// A error occurred while writing to the output.
376 #[error("Format error")]
377 FmtError(#[from] FmtError),
378 /// The specified [`Version`] doesn't have all required [`Features`].
379 ///
380 /// Contains the missing [`Features`].
381 #[error("The selected version doesn't support {0:?}")]
382 MissingFeatures(Features),
383 /// [`AddressSpace::PushConstant`](crate::AddressSpace::PushConstant) was used more than
384 /// once in the entry point, which isn't supported.
385 #[error("Multiple push constants aren't supported")]
386 MultiplePushConstants,
387 /// The specified [`Version`] isn't supported.
388 #[error("The specified version isn't supported")]
389 VersionNotSupported,
390 /// The entry point couldn't be found.
391 #[error("The requested entry point couldn't be found")]
392 EntryPointNotFound,
393 /// A call was made to an unsupported external.
394 #[error("A call was made to an unsupported external: {0}")]
395 UnsupportedExternal(String),
396 /// A scalar with an unsupported width was requested.
397 #[error("A scalar with an unsupported width was requested: {0:?} {1:?}")]
398 UnsupportedScalar(crate::ScalarKind, crate::Bytes),
399 /// A image was used with multiple samplers, which isn't supported.
400 #[error("A image was used with multiple samplers")]
401 ImageMultipleSamplers,
402 #[error("{0}")]
403 Custom(String),
404}
405
406/// Binary operation with a different logic on the GLSL side.
407enum BinaryOperation {
408 /// Vector comparison should use the function like `greaterThan()`, etc.
409 VectorCompare,
410 /// Vector component wise operation; used to polyfill unsupported ops like `|` and `&` for `bvecN`'s
411 VectorComponentWise,
412 /// GLSL `%` is SPIR-V `OpUMod/OpSMod` and `mod()` is `OpFMod`, but [`BinaryOperator::Modulo`](crate::BinaryOperator::Modulo) is `OpFRem`.
413 Modulo,
414 /// Any plain operation. No additional logic required.
415 Other,
416}
417
418/// Writer responsible for all code generation.
419pub struct Writer<'a, W> {
420 // Inputs
421 /// The module being written.
422 module: &'a crate::Module,
423 /// The module analysis.
424 info: &'a valid::ModuleInfo,
425 /// The output writer.
426 out: W,
427 /// User defined configuration to be used.
428 options: &'a Options,
429 /// The bound checking policies to be used
430 policies: proc::BoundsCheckPolicies,
431
432 // Internal State
433 /// Features manager used to store all the needed features and write them.
434 features: FeaturesManager,
435 namer: proc::Namer,
436 /// A map with all the names needed for writing the module
437 /// (generated by a [`Namer`](crate::proc::Namer)).
438 names: crate::FastHashMap<NameKey, String>,
439 /// A map with the names of global variables needed for reflections.
440 reflection_names_globals: crate::FastHashMap<Handle<crate::GlobalVariable>, String>,
441 /// The selected entry point.
442 entry_point: &'a crate::EntryPoint,
443 /// The index of the selected entry point.
444 entry_point_idx: proc::EntryPointIndex,
445 /// A generator for unique block numbers.
446 block_id: IdGenerator,
447 /// Set of expressions that have associated temporary variables.
448 named_expressions: crate::NamedExpressions,
449 /// Set of expressions that need to be baked to avoid unnecessary repetition in output
450 need_bake_expressions: back::NeedBakeExpressions,
451 /// How many views to render to, if doing multiview rendering.
452 multiview: Option<std::num::NonZeroU32>,
453}
454
455impl<'a, W: Write> Writer<'a, W> {
456 /// Creates a new [`Writer`] instance.
457 ///
458 /// # Errors
459 /// - If the version specified is invalid or supported.
460 /// - If the entry point couldn't be found in the module.
461 /// - If the version specified doesn't support some used features.
462 pub fn new(
463 out: W,
464 module: &'a crate::Module,
465 info: &'a valid::ModuleInfo,
466 options: &'a Options,
467 pipeline_options: &'a PipelineOptions,
468 policies: proc::BoundsCheckPolicies,
469 ) -> Result<Self, Error> {
470 // Check if the requested version is supported
471 if !options.version.is_supported() {
472 log::error!("Version {}", options.version);
473 return Err(Error::VersionNotSupported);
474 }
475
476 // Try to find the entry point and corresponding index
477 let ep_idx = module
478 .entry_points
479 .iter()
480 .position(|ep| {
481 pipeline_options.shader_stage == ep.stage && pipeline_options.entry_point == ep.name
482 })
483 .ok_or(Error::EntryPointNotFound)?;
484
485 // Generate a map with names required to write the module
486 let mut names = crate::FastHashMap::default();
487 let mut namer = proc::Namer::default();
488 namer.reset(
489 module,
490 keywords::RESERVED_KEYWORDS,
491 &[],
492 &[],
493 &["gl_"],
494 &mut names,
495 );
496
497 // Build the instance
498 let mut this = Self {
499 module,
500 info,
501 out,
502 options,
503 policies,
504
505 namer,
506 features: FeaturesManager::new(),
507 names,
508 reflection_names_globals: crate::FastHashMap::default(),
509 entry_point: &module.entry_points[ep_idx],
510 entry_point_idx: ep_idx as u16,
511 multiview: pipeline_options.multiview,
512 block_id: IdGenerator::default(),
513 named_expressions: Default::default(),
514 need_bake_expressions: Default::default(),
515 };
516
517 // Find all features required to print this module
518 this.collect_required_features()?;
519
520 Ok(this)
521 }
522
523 /// Writes the [`Module`](crate::Module) as glsl to the output
524 ///
525 /// # Notes
526 /// If an error occurs while writing, the output might have been written partially
527 ///
528 /// # Panics
529 /// Might panic if the module is invalid
530 pub fn write(&mut self) -> Result<ReflectionInfo, Error> {
531 // We use `writeln!(self.out)` throughout the write to add newlines
532 // to make the output more readable
533
534 let es = self.options.version.is_es();
535
536 // Write the version (It must be the first thing or it isn't a valid glsl output)
537 writeln!(self.out, "#version {}", self.options.version)?;
538 // Write all the needed extensions
539 //
540 // This used to be the last thing being written as it allowed to search for features while
541 // writing the module saving some loops but some older versions (420 or less) required the
542 // extensions to appear before being used, even though extensions are part of the
543 // preprocessor not the processor ¯\_(ツ)_/¯
544 self.features.write(self.options.version, &mut self.out)?;
545
546 // Write the additional extensions
547 if self
548 .options
549 .writer_flags
550 .contains(WriterFlags::TEXTURE_SHADOW_LOD)
551 {
552 // https://www.khronos.org/registry/OpenGL/extensions/EXT/EXT_texture_shadow_lod.txt
553 writeln!(self.out, "#extension GL_EXT_texture_shadow_lod : require")?;
554 }
555
556 // glsl es requires a precision to be specified for floats and ints
557 // TODO: Should this be user configurable?
558 if es {
559 writeln!(self.out)?;
560 writeln!(self.out, "precision highp float;")?;
561 writeln!(self.out, "precision highp int;")?;
562 writeln!(self.out)?;
563 }
564
565 if self.entry_point.stage == ShaderStage::Compute {
566 let workgroup_size = self.entry_point.workgroup_size;
567 writeln!(
568 self.out,
569 "layout(local_size_x = {}, local_size_y = {}, local_size_z = {}) in;",
570 workgroup_size[0], workgroup_size[1], workgroup_size[2]
571 )?;
572 writeln!(self.out)?;
573 }
574
575 // Enable early depth tests if needed
576 if let Some(depth_test) = self.entry_point.early_depth_test {
577 // If early depth test is supported for this version of GLSL
578 if self.options.version.supports_early_depth_test() {
579 writeln!(self.out, "layout(early_fragment_tests) in;")?;
580
581 if let Some(conservative) = depth_test.conservative {
582 use crate::ConservativeDepth as Cd;
583
584 let depth = match conservative {
585 Cd::GreaterEqual => "greater",
586 Cd::LessEqual => "less",
587 Cd::Unchanged => "unchanged",
588 };
589 writeln!(self.out, "layout (depth_{depth}) out float gl_FragDepth;")?;
590 }
591 writeln!(self.out)?;
592 } else {
593 log::warn!(
594 "Early depth testing is not supported for this version of GLSL: {}",
595 self.options.version
596 );
597 }
598 }
599
600 if self.entry_point.stage == ShaderStage::Vertex && self.options.version.is_webgl() {
601 if let Some(multiview) = self.multiview.as_ref() {
602 writeln!(self.out, "layout(num_views = {multiview}) in;")?;
603 writeln!(self.out)?;
604 }
605 }
606
607 // Write struct types.
608 //
609 // This are always ordered because the IR is structured in a way that
610 // you can't make a struct without adding all of its members first.
611 for (handle, ty) in self.module.types.iter() {
612 if let TypeInner::Struct { ref members, .. } = ty.inner {
613 // Structures ending with runtime-sized arrays can only be
614 // rendered as shader storage blocks in GLSL, not stand-alone
615 // struct types.
616 if !self.module.types[members.last().unwrap().ty]
617 .inner
618 .is_dynamically_sized(&self.module.types)
619 {
620 let name = &self.names[&NameKey::Type(handle)];
621 write!(self.out, "struct {name} ")?;
622 self.write_struct_body(handle, members)?;
623 writeln!(self.out, ";")?;
624 }
625 }
626 }
627
628 // Write all named constants
629 let mut constants = self
630 .module
631 .constants
632 .iter()
633 .filter(|&(_, c)| c.name.is_some())
634 .peekable();
635 while let Some((handle, _)) = constants.next() {
636 self.write_global_constant(handle)?;
637 // Add extra newline for readability on last iteration
638 if constants.peek().is_none() {
639 writeln!(self.out)?;
640 }
641 }
642
643 let ep_info = self.info.get_entry_point(self.entry_point_idx as usize);
644
645 // Write the globals
646 //
647 // Unless explicitly disabled with WriterFlags::INCLUDE_UNUSED_ITEMS,
648 // we filter all globals that aren't used by the selected entry point as they might be
649 // interfere with each other (i.e. two globals with the same location but different with
650 // different classes)
651 let include_unused = self
652 .options
653 .writer_flags
654 .contains(WriterFlags::INCLUDE_UNUSED_ITEMS);
655 for (handle, global) in self.module.global_variables.iter() {
656 let is_unused = ep_info[handle].is_empty();
657 if !include_unused && is_unused {
658 continue;
659 }
660
661 match self.module.types[global.ty].inner {
662 // We treat images separately because they might require
663 // writing the storage format
664 TypeInner::Image {
665 mut dim,
666 arrayed,
667 class,
668 } => {
669 // Gather the storage format if needed
670 let storage_format_access = match self.module.types[global.ty].inner {
671 TypeInner::Image {
672 class: crate::ImageClass::Storage { format, access },
673 ..
674 } => Some((format, access)),
675 _ => None,
676 };
677
678 if dim == crate::ImageDimension::D1 && es {
679 dim = crate::ImageDimension::D2
680 }
681
682 // Gether the location if needed
683 let layout_binding = if self.options.version.supports_explicit_locations() {
684 let br = global.binding.as_ref().unwrap();
685 self.options.binding_map.get(br).cloned()
686 } else {
687 None
688 };
689
690 // Write all the layout qualifiers
691 if layout_binding.is_some() || storage_format_access.is_some() {
692 write!(self.out, "layout(")?;
693 if let Some(binding) = layout_binding {
694 write!(self.out, "binding = {binding}")?;
695 }
696 if let Some((format, _)) = storage_format_access {
697 let format_str = glsl_storage_format(format);
698 let separator = match layout_binding {
699 Some(_) => ",",
700 None => "",
701 };
702 write!(self.out, "{separator}{format_str}")?;
703 }
704 write!(self.out, ") ")?;
705 }
706
707 if let Some((_, access)) = storage_format_access {
708 self.write_storage_access(access)?;
709 }
710
711 // All images in glsl are `uniform`
712 // The trailing space is important
713 write!(self.out, "uniform ")?;
714
715 // write the type
716 //
717 // This is way we need the leading space because `write_image_type` doesn't add
718 // any spaces at the beginning or end
719 self.write_image_type(dim, arrayed, class)?;
720
721 // Finally write the name and end the global with a `;`
722 // The leading space is important
723 let global_name = self.get_global_name(handle, global);
724 writeln!(self.out, " {global_name};")?;
725 writeln!(self.out)?;
726
727 self.reflection_names_globals.insert(handle, global_name);
728 }
729 // glsl has no concept of samplers so we just ignore it
730 TypeInner::Sampler { .. } => continue,
731 // All other globals are written by `write_global`
732 _ => {
733 self.write_global(handle, global)?;
734 // Add a newline (only for readability)
735 writeln!(self.out)?;
736 }
737 }
738 }
739
740 for arg in self.entry_point.function.arguments.iter() {
741 self.write_varying(arg.binding.as_ref(), arg.ty, false)?;
742 }
743 if let Some(ref result) = self.entry_point.function.result {
744 self.write_varying(result.binding.as_ref(), result.ty, true)?;
745 }
746 writeln!(self.out)?;
747
748 // Write all regular functions
749 for (handle, function) in self.module.functions.iter() {
750 // Check that the function doesn't use globals that aren't supported
751 // by the current entry point
752 if !include_unused && !ep_info.dominates_global_use(&self.info[handle]) {
753 continue;
754 }
755
756 let fun_info = &self.info[handle];
757
758 // Write the function
759 self.write_function(back::FunctionType::Function(handle), function, fun_info)?;
760
761 writeln!(self.out)?;
762 }
763
764 self.write_function(
765 back::FunctionType::EntryPoint(self.entry_point_idx),
766 &self.entry_point.function,
767 ep_info,
768 )?;
769
770 // Add newline at the end of file
771 writeln!(self.out)?;
772
773 // Collect all reflection info and return it to the user
774 self.collect_reflection_info()
775 }
776
777 fn write_array_size(
778 &mut self,
779 base: Handle<crate::Type>,
780 size: crate::ArraySize,
781 ) -> BackendResult {
782 write!(self.out, "[")?;
783
784 // Write the array size
785 // Writes nothing if `ArraySize::Dynamic`
786 match size {
787 crate::ArraySize::Constant(size) => {
788 write!(self.out, "{size}")?;
789 }
790 crate::ArraySize::Dynamic => (),
791 }
792
793 write!(self.out, "]")?;
794
795 if let TypeInner::Array {
796 base: next_base,
797 size: next_size,
798 ..
799 } = self.module.types[base].inner
800 {
801 self.write_array_size(next_base, next_size)?;
802 }
803
804 Ok(())
805 }
806
807 /// Helper method used to write value types
808 ///
809 /// # Notes
810 /// Adds no trailing or leading whitespace
811 fn write_value_type(&mut self, inner: &TypeInner) -> BackendResult {
812 match *inner {
813 // Scalars are simple we just get the full name from `glsl_scalar`
814 TypeInner::Scalar { kind, width }
815 | TypeInner::Atomic { kind, width }
816 | TypeInner::ValuePointer {
817 size: None,
818 kind,
819 width,
820 space: _,
821 } => write!(self.out, "{}", glsl_scalar(kind, width)?.full)?,
822 // Vectors are just `gvecN` where `g` is the scalar prefix and `N` is the vector size
823 TypeInner::Vector { size, kind, width }
824 | TypeInner::ValuePointer {
825 size: Some(size),
826 kind,
827 width,
828 space: _,
829 } => write!(
830 self.out,
831 "{}vec{}",
832 glsl_scalar(kind, width)?.prefix,
833 size as u8
834 )?,
835 // Matrices are written with `gmatMxN` where `g` is the scalar prefix (only floats and
836 // doubles are allowed), `M` is the columns count and `N` is the rows count
837 //
838 // glsl supports a matrix shorthand `gmatN` where `N` = `M` but it doesn't justify the
839 // extra branch to write matrices this way
840 TypeInner::Matrix {
841 columns,
842 rows,
843 width,
844 } => write!(
845 self.out,
846 "{}mat{}x{}",
847 glsl_scalar(crate::ScalarKind::Float, width)?.prefix,
848 columns as u8,
849 rows as u8
850 )?,
851 // GLSL arrays are written as `type name[size]`
852 // Here we only write the size of the array i.e. `[size]`
853 // Base `type` and `name` should be written outside
854 TypeInner::Array { base, size, .. } => self.write_array_size(base, size)?,
855 // Write all variants instead of `_` so that if new variants are added a
856 // no exhaustiveness error is thrown
857 TypeInner::Pointer { .. }
858 | TypeInner::Struct { .. }
859 | TypeInner::Image { .. }
860 | TypeInner::Sampler { .. }
861 | TypeInner::AccelerationStructure
862 | TypeInner::RayQuery
863 | TypeInner::BindingArray { .. } => {
864 return Err(Error::Custom(format!("Unable to write type {inner:?}")))
865 }
866 }
867
868 Ok(())
869 }
870
871 /// Helper method used to write non image/sampler types
872 ///
873 /// # Notes
874 /// Adds no trailing or leading whitespace
875 fn write_type(&mut self, ty: Handle<crate::Type>) -> BackendResult {
876 match self.module.types[ty].inner {
877 // glsl has no pointer types so just write types as normal and loads are skipped
878 TypeInner::Pointer { base, .. } => self.write_type(base),
879 // glsl structs are written as just the struct name
880 TypeInner::Struct { .. } => {
881 // Get the struct name
882 let name = &self.names[&NameKey::Type(ty)];
883 write!(self.out, "{name}")?;
884 Ok(())
885 }
886 // glsl array has the size separated from the base type
887 TypeInner::Array { base, .. } => self.write_type(base),
888 ref other => self.write_value_type(other),
889 }
890 }
891
892 /// Helper method to write a image type
893 ///
894 /// # Notes
895 /// Adds no leading or trailing whitespace
896 fn write_image_type(
897 &mut self,
898 dim: crate::ImageDimension,
899 arrayed: bool,
900 class: crate::ImageClass,
901 ) -> BackendResult {
902 // glsl images consist of four parts the scalar prefix, the image "type", the dimensions
903 // and modifiers
904 //
905 // There exists two image types
906 // - sampler - for sampled images
907 // - image - for storage images
908 //
909 // There are three possible modifiers that can be used together and must be written in
910 // this order to be valid
911 // - MS - used if it's a multisampled image
912 // - Array - used if it's an image array
913 // - Shadow - used if it's a depth image
914 use crate::ImageClass as Ic;
915
916 let (base, kind, ms, comparison) = match class {
917 Ic::Sampled { kind, multi: true } => ("sampler", kind, "MS", ""),
918 Ic::Sampled { kind, multi: false } => ("sampler", kind, "", ""),
919 Ic::Depth { multi: true } => ("sampler", crate::ScalarKind::Float, "MS", ""),
920 Ic::Depth { multi: false } => ("sampler", crate::ScalarKind::Float, "", "Shadow"),
921 Ic::Storage { format, .. } => ("image", format.into(), "", ""),
922 };
923
924 write!(
925 self.out,
926 "highp {}{}{}{}{}{}",
927 glsl_scalar(kind, 4)?.prefix,
928 base,
929 glsl_dimension(dim),
930 ms,
931 if arrayed { "Array" } else { "" },
932 comparison
933 )?;
934
935 Ok(())
936 }
937
938 /// Helper method used to write non images/sampler globals
939 ///
940 /// # Notes
941 /// Adds a newline
942 ///
943 /// # Panics
944 /// If the global has type sampler
945 fn write_global(
946 &mut self,
947 handle: Handle<crate::GlobalVariable>,
948 global: &crate::GlobalVariable,
949 ) -> BackendResult {
950 if self.options.version.supports_explicit_locations() {
951 if let Some(ref br) = global.binding {
952 match self.options.binding_map.get(br) {
953 Some(binding) => {
954 let layout = match global.space {
955 crate::AddressSpace::Storage { .. } => {
956 if self.options.version.supports_std430_layout() {
957 "std430, "
958 } else {
959 "std140, "
960 }
961 }
962 crate::AddressSpace::Uniform => "std140, ",
963 _ => "",
964 };
965 write!(self.out, "layout({layout}binding = {binding}) ")?
966 }
967 None => {
968 log::debug!("unassigned binding for {:?}", global.name);
969 if let crate::AddressSpace::Storage { .. } = global.space {
970 if self.options.version.supports_std430_layout() {
971 write!(self.out, "layout(std430) ")?
972 }
973 }
974 }
975 }
976 }
977 }
978
979 if let crate::AddressSpace::Storage { access } = global.space {
980 self.write_storage_access(access)?;
981 }
982
983 if let Some(storage_qualifier) = glsl_storage_qualifier(global.space) {
984 write!(self.out, "{storage_qualifier} ")?;
985 }
986
987 match global.space {
988 crate::AddressSpace::Private => {
989 self.write_simple_global(handle, global)?;
990 }
991 crate::AddressSpace::WorkGroup => {
992 self.write_simple_global(handle, global)?;
993 }
994 crate::AddressSpace::PushConstant => {
995 self.write_simple_global(handle, global)?;
996 }
997 crate::AddressSpace::Uniform => {
998 self.write_interface_block(handle, global)?;
999 }
1000 crate::AddressSpace::Storage { .. } => {
1001 self.write_interface_block(handle, global)?;
1002 }
1003 // A global variable in the `Function` address space is a
1004 // contradiction in terms.
1005 crate::AddressSpace::Function => unreachable!(),
1006 // Textures and samplers are handled directly in `Writer::write`.
1007 crate::AddressSpace::Handle => unreachable!(),
1008 }
1009
1010 Ok(())
1011 }
1012
1013 fn write_simple_global(
1014 &mut self,
1015 handle: Handle<crate::GlobalVariable>,
1016 global: &crate::GlobalVariable,
1017 ) -> BackendResult {
1018 self.write_type(global.ty)?;
1019 write!(self.out, " ")?;
1020 self.write_global_name(handle, global)?;
1021
1022 if let TypeInner::Array { base, size, .. } = self.module.types[global.ty].inner {
1023 self.write_array_size(base, size)?;
1024 }
1025
1026 if global.space.initializable() && is_value_init_supported(self.module, global.ty) {
1027 write!(self.out, " = ")?;
1028 if let Some(init) = global.init {
1029 self.write_const_expr(init)?;
1030 } else {
1031 self.write_zero_init_value(global.ty)?;
1032 }
1033 }
1034
1035 writeln!(self.out, ";")?;
1036
1037 if let crate::AddressSpace::PushConstant = global.space {
1038 let global_name = self.get_global_name(handle, global);
1039 self.reflection_names_globals.insert(handle, global_name);
1040 }
1041
1042 Ok(())
1043 }
1044
1045 /// Write an interface block for a single Naga global.
1046 ///
1047 /// Write `block_name { members }`. Since `block_name` must be unique
1048 /// between blocks and structs, we add `_block_ID` where `ID` is a
1049 /// `IdGenerator` generated number. Write `members` in the same way we write
1050 /// a struct's members.
1051 fn write_interface_block(
1052 &mut self,
1053 handle: Handle<crate::GlobalVariable>,
1054 global: &crate::GlobalVariable,
1055 ) -> BackendResult {
1056 // Write the block name, it's just the struct name appended with `_block_ID`
1057 let ty_name = &self.names[&NameKey::Type(global.ty)];
1058 let block_name = format!(
1059 "{}_block_{}{:?}",
1060 ty_name,
1061 self.block_id.generate(),
1062 self.entry_point.stage,
1063 );
1064 write!(self.out, "{block_name} ")?;
1065 self.reflection_names_globals.insert(handle, block_name);
1066
1067 match self.module.types[global.ty].inner {
1068 crate::TypeInner::Struct { ref members, .. }
1069 if self.module.types[members.last().unwrap().ty]
1070 .inner
1071 .is_dynamically_sized(&self.module.types) =>
1072 {
1073 // Structs with dynamically sized arrays must have their
1074 // members lifted up as members of the interface block. GLSL
1075 // can't write such struct types anyway.
1076 self.write_struct_body(global.ty, members)?;
1077 write!(self.out, " ")?;
1078 self.write_global_name(handle, global)?;
1079 }
1080 _ => {
1081 // A global of any other type is written as the sole member
1082 // of the interface block. Since the interface block is
1083 // anonymous, this becomes visible in the global scope.
1084 write!(self.out, "{{ ")?;
1085 self.write_type(global.ty)?;
1086 write!(self.out, " ")?;
1087 self.write_global_name(handle, global)?;
1088 if let TypeInner::Array { base, size, .. } = self.module.types[global.ty].inner {
1089 self.write_array_size(base, size)?;
1090 }
1091 write!(self.out, "; }}")?;
1092 }
1093 }
1094
1095 writeln!(self.out, ";")?;
1096
1097 Ok(())
1098 }
1099
1100 /// Helper method used to find which expressions of a given function require baking
1101 ///
1102 /// # Notes
1103 /// Clears `need_bake_expressions` set before adding to it
1104 fn update_expressions_to_bake(&mut self, func: &crate::Function, info: &valid::FunctionInfo) {
1105 use crate::Expression;
1106 self.need_bake_expressions.clear();
1107 for (fun_handle, expr) in func.expressions.iter() {
1108 let expr_info = &info[fun_handle];
1109 let min_ref_count = func.expressions[fun_handle].bake_ref_count();
1110 if min_ref_count <= expr_info.ref_count {
1111 self.need_bake_expressions.insert(fun_handle);
1112 }
1113
1114 let inner = expr_info.ty.inner_with(&self.module.types);
1115
1116 if let Expression::Math { fun, arg, arg1, .. } = *expr {
1117 match fun {
1118 crate::MathFunction::Dot => {
1119 // if the expression is a Dot product with integer arguments,
1120 // then the args needs baking as well
1121 if let TypeInner::Scalar { kind, .. } = *inner {
1122 match kind {
1123 crate::ScalarKind::Sint | crate::ScalarKind::Uint => {
1124 self.need_bake_expressions.insert(arg);
1125 self.need_bake_expressions.insert(arg1.unwrap());
1126 }
1127 _ => {}
1128 }
1129 }
1130 }
1131 crate::MathFunction::CountLeadingZeros => {
1132 if let Some(crate::ScalarKind::Sint) = inner.scalar_kind() {
1133 self.need_bake_expressions.insert(arg);
1134 }
1135 }
1136 _ => {}
1137 }
1138 }
1139 }
1140 }
1141
1142 /// Helper method used to get a name for a global
1143 ///
1144 /// Globals have different naming schemes depending on their binding:
1145 /// - Globals without bindings use the name from the [`Namer`](crate::proc::Namer)
1146 /// - Globals with resource binding are named `_group_X_binding_Y` where `X`
1147 /// is the group and `Y` is the binding
1148 fn get_global_name(
1149 &self,
1150 handle: Handle<crate::GlobalVariable>,
1151 global: &crate::GlobalVariable,
1152 ) -> String {
1153 match global.binding {
1154 Some(ref br) => {
1155 format!(
1156 "_group_{}_binding_{}_{}",
1157 br.group,
1158 br.binding,
1159 self.entry_point.stage.to_str()
1160 )
1161 }
1162 None => self.names[&NameKey::GlobalVariable(handle)].clone(),
1163 }
1164 }
1165
1166 /// Helper method used to write a name for a global without additional heap allocation
1167 fn write_global_name(
1168 &mut self,
1169 handle: Handle<crate::GlobalVariable>,
1170 global: &crate::GlobalVariable,
1171 ) -> BackendResult {
1172 match global.binding {
1173 Some(ref br) => write!(
1174 self.out,
1175 "_group_{}_binding_{}_{}",
1176 br.group,
1177 br.binding,
1178 self.entry_point.stage.to_str()
1179 )?,
1180 None => write!(
1181 self.out,
1182 "{}",
1183 &self.names[&NameKey::GlobalVariable(handle)]
1184 )?,
1185 }
1186
1187 Ok(())
1188 }
1189
1190 /// Write a GLSL global that will carry a Naga entry point's argument or return value.
1191 ///
1192 /// A Naga entry point's arguments and return value are rendered in GLSL as
1193 /// variables at global scope with the `in` and `out` storage qualifiers.
1194 /// The code we generate for `main` loads from all the `in` globals into
1195 /// appropriately named locals. Before it returns, `main` assigns the
1196 /// components of its return value into all the `out` globals.
1197 ///
1198 /// This function writes a declaration for one such GLSL global,
1199 /// representing a value passed into or returned from [`self.entry_point`]
1200 /// that has a [`Location`] binding. The global's name is generated based on
1201 /// the location index and the shader stages being connected; see
1202 /// [`VaryingName`]. This means we don't need to know the names of
1203 /// arguments, just their types and bindings.
1204 ///
1205 /// Emit nothing for entry point arguments or return values with [`BuiltIn`]
1206 /// bindings; `main` will read from or assign to the appropriate GLSL
1207 /// special variable; these are pre-declared. As an exception, we do declare
1208 /// `gl_Position` or `gl_FragCoord` with the `invariant` qualifier if
1209 /// needed.
1210 ///
1211 /// Use `output` together with [`self.entry_point.stage`] to determine which
1212 /// shader stages are being connected, and choose the `in` or `out` storage
1213 /// qualifier.
1214 ///
1215 /// [`self.entry_point`]: Writer::entry_point
1216 /// [`self.entry_point.stage`]: crate::EntryPoint::stage
1217 /// [`Location`]: crate::Binding::Location
1218 /// [`BuiltIn`]: crate::Binding::BuiltIn
1219 fn write_varying(
1220 &mut self,
1221 binding: Option<&crate::Binding>,
1222 ty: Handle<crate::Type>,
1223 output: bool,
1224 ) -> Result<(), Error> {
1225 // For a struct, emit a separate global for each member with a binding.
1226 if let crate::TypeInner::Struct { ref members, .. } = self.module.types[ty].inner {
1227 for member in members {
1228 self.write_varying(member.binding.as_ref(), member.ty, output)?;
1229 }
1230 return Ok(());
1231 }
1232
1233 let binding = match binding {
1234 None => return Ok(()),
1235 Some(binding) => binding,
1236 };
1237
1238 let (location, interpolation, sampling) = match *binding {
1239 crate::Binding::Location {
1240 location,
1241 interpolation,
1242 sampling,
1243 } => (location, interpolation, sampling),
1244 crate::Binding::BuiltIn(built_in) => {
1245 if let crate::BuiltIn::Position { invariant: true } = built_in {
1246 match (self.options.version, self.entry_point.stage) {
1247 (
1248 Version::Embedded {
1249 version: 300,
1250 is_webgl: true,
1251 },
1252 ShaderStage::Fragment,
1253 ) => {
1254 // `invariant gl_FragCoord` is not allowed in WebGL2 and possibly
1255 // OpenGL ES in general (waiting on confirmation).
1256 //
1257 // See https://github.com/KhronosGroup/WebGL/issues/3518
1258 }
1259 _ => {
1260 writeln!(
1261 self.out,
1262 "invariant {};",
1263 glsl_built_in(built_in, output, self.options.version.is_webgl())
1264 )?;
1265 }
1266 }
1267 }
1268 return Ok(());
1269 }
1270 };
1271
1272 // Write the interpolation modifier if needed
1273 //
1274 // We ignore all interpolation and auxiliary modifiers that aren't used in fragment
1275 // shaders' input globals or vertex shaders' output globals.
1276 let emit_interpolation_and_auxiliary = match self.entry_point.stage {
1277 ShaderStage::Vertex => output,
1278 ShaderStage::Fragment => !output,
1279 ShaderStage::Compute => false,
1280 };
1281
1282 // Write the I/O locations, if allowed
1283 if self.options.version.supports_explicit_locations() || !emit_interpolation_and_auxiliary {
1284 write!(self.out, "layout(location = {location}) ")?;
1285 }
1286
1287 // Write the interpolation qualifier.
1288 if let Some(interp) = interpolation {
1289 if emit_interpolation_and_auxiliary {
1290 write!(self.out, "{} ", glsl_interpolation(interp))?;
1291 }
1292 }
1293
1294 // Write the sampling auxiliary qualifier.
1295 //
1296 // Before GLSL 4.2, the `centroid` and `sample` qualifiers were required to appear
1297 // immediately before the `in` / `out` qualifier, so we'll just follow that rule
1298 // here, regardless of the version.
1299 if let Some(sampling) = sampling {
1300 if emit_interpolation_and_auxiliary {
1301 if let Some(qualifier) = glsl_sampling(sampling) {
1302 write!(self.out, "{qualifier} ")?;
1303 }
1304 }
1305 }
1306
1307 // Write the input/output qualifier.
1308 write!(self.out, "{} ", if output { "out" } else { "in" })?;
1309
1310 // Write the type
1311 // `write_type` adds no leading or trailing spaces
1312 self.write_type(ty)?;
1313
1314 // Finally write the global name and end the global with a `;` and a newline
1315 // Leading space is important
1316 let vname = VaryingName {
1317 binding: &crate::Binding::Location {
1318 location,
1319 interpolation: None,
1320 sampling: None,
1321 },
1322 stage: self.entry_point.stage,
1323 output,
1324 targetting_webgl: self.options.version.is_webgl(),
1325 };
1326 writeln!(self.out, " {vname};")?;
1327
1328 Ok(())
1329 }
1330
1331 /// Helper method used to write functions (both entry points and regular functions)
1332 ///
1333 /// # Notes
1334 /// Adds a newline
1335 fn write_function(
1336 &mut self,
1337 ty: back::FunctionType,
1338 func: &crate::Function,
1339 info: &valid::FunctionInfo,
1340 ) -> BackendResult {
1341 // Create a function context for the function being written
1342 let ctx = back::FunctionCtx {
1343 ty,
1344 info,
1345 expressions: &func.expressions,
1346 named_expressions: &func.named_expressions,
1347 };
1348
1349 self.named_expressions.clear();
1350 self.update_expressions_to_bake(func, info);
1351
1352 // Write the function header
1353 //
1354 // glsl headers are the same as in c:
1355 // `ret_type name(args)`
1356 // `ret_type` is the return type
1357 // `name` is the function name
1358 // `args` is a comma separated list of `type name`
1359 // | - `type` is the argument type
1360 // | - `name` is the argument name
1361
1362 // Start by writing the return type if any otherwise write void
1363 // This is the only place where `void` is a valid type
1364 // (though it's more a keyword than a type)
1365 if let back::FunctionType::EntryPoint(_) = ctx.ty {
1366 write!(self.out, "void")?;
1367 } else if let Some(ref result) = func.result {
1368 self.write_type(result.ty)?;
1369 if let TypeInner::Array { base, size, .. } = self.module.types[result.ty].inner {
1370 self.write_array_size(base, size)?
1371 }
1372 } else {
1373 write!(self.out, "void")?;
1374 }
1375
1376 // Write the function name and open parentheses for the argument list
1377 let function_name = match ctx.ty {
1378 back::FunctionType::Function(handle) => &self.names[&NameKey::Function(handle)],
1379 back::FunctionType::EntryPoint(_) => "main",
1380 };
1381 write!(self.out, " {function_name}(")?;
1382
1383 // Write the comma separated argument list
1384 //
1385 // We need access to `Self` here so we use the reference passed to the closure as an
1386 // argument instead of capturing as that would cause a borrow checker error
1387 let arguments = match ctx.ty {
1388 back::FunctionType::EntryPoint(_) => &[][..],
1389 back::FunctionType::Function(_) => &func.arguments,
1390 };
1391 let arguments: Vec<_> = arguments
1392 .iter()
1393 .enumerate()
1394 .filter(|&(_, arg)| match self.module.types[arg.ty].inner {
1395 TypeInner::Sampler { .. } => false,
1396 _ => true,
1397 })
1398 .collect();
1399 self.write_slice(&arguments, |this, _, &(i, arg)| {
1400 // Write the argument type
1401 match this.module.types[arg.ty].inner {
1402 // We treat images separately because they might require
1403 // writing the storage format
1404 TypeInner::Image {
1405 dim,
1406 arrayed,
1407 class,
1408 } => {
1409 // Write the storage format if needed
1410 if let TypeInner::Image {
1411 class: crate::ImageClass::Storage { format, .. },
1412 ..
1413 } = this.module.types[arg.ty].inner
1414 {
1415 write!(this.out, "layout({}) ", glsl_storage_format(format))?;
1416 }
1417
1418 // write the type
1419 //
1420 // This is way we need the leading space because `write_image_type` doesn't add
1421 // any spaces at the beginning or end
1422 this.write_image_type(dim, arrayed, class)?;
1423 }
1424 TypeInner::Pointer { base, .. } => {
1425 // write parameter qualifiers
1426 write!(this.out, "inout ")?;
1427 this.write_type(base)?;
1428 }
1429 // All other types are written by `write_type`
1430 _ => {
1431 this.write_type(arg.ty)?;
1432 }
1433 }
1434
1435 // Write the argument name
1436 // The leading space is important
1437 write!(this.out, " {}", &this.names[&ctx.argument_key(i as u32)])?;
1438
1439 // Write array size
1440 match this.module.types[arg.ty].inner {
1441 TypeInner::Array { base, size, .. } => {
1442 this.write_array_size(base, size)?;
1443 }
1444 TypeInner::Pointer { base, .. } => {
1445 if let TypeInner::Array { base, size, .. } = this.module.types[base].inner {
1446 this.write_array_size(base, size)?;
1447 }
1448 }
1449 _ => {}
1450 }
1451
1452 Ok(())
1453 })?;
1454
1455 // Close the parentheses and open braces to start the function body
1456 writeln!(self.out, ") {{")?;
1457
1458 if self.options.zero_initialize_workgroup_memory
1459 && ctx.ty.is_compute_entry_point(self.module)
1460 {
1461 self.write_workgroup_variables_initialization(&ctx)?;
1462 }
1463
1464 // Compose the function arguments from globals, in case of an entry point.
1465 if let back::FunctionType::EntryPoint(ep_index) = ctx.ty {
1466 let stage = self.module.entry_points[ep_index as usize].stage;
1467 for (index, arg) in func.arguments.iter().enumerate() {
1468 write!(self.out, "{}", back::INDENT)?;
1469 self.write_type(arg.ty)?;
1470 let name = &self.names[&NameKey::EntryPointArgument(ep_index, index as u32)];
1471 write!(self.out, " {name}")?;
1472 write!(self.out, " = ")?;
1473 match self.module.types[arg.ty].inner {
1474 crate::TypeInner::Struct { ref members, .. } => {
1475 self.write_type(arg.ty)?;
1476 write!(self.out, "(")?;
1477 for (index, member) in members.iter().enumerate() {
1478 let varying_name = VaryingName {
1479 binding: member.binding.as_ref().unwrap(),
1480 stage,
1481 output: false,
1482 targetting_webgl: self.options.version.is_webgl(),
1483 };
1484 if index != 0 {
1485 write!(self.out, ", ")?;
1486 }
1487 write!(self.out, "{varying_name}")?;
1488 }
1489 writeln!(self.out, ");")?;
1490 }
1491 _ => {
1492 let varying_name = VaryingName {
1493 binding: arg.binding.as_ref().unwrap(),
1494 stage,
1495 output: false,
1496 targetting_webgl: self.options.version.is_webgl(),
1497 };
1498 writeln!(self.out, "{varying_name};")?;
1499 }
1500 }
1501 }
1502 }
1503
1504 // Write all function locals
1505 // Locals are `type name (= init)?;` where the init part (including the =) are optional
1506 //
1507 // Always adds a newline
1508 for (handle, local) in func.local_variables.iter() {
1509 // Write indentation (only for readability) and the type
1510 // `write_type` adds no trailing space
1511 write!(self.out, "{}", back::INDENT)?;
1512 self.write_type(local.ty)?;
1513
1514 // Write the local name
1515 // The leading space is important
1516 write!(self.out, " {}", self.names[&ctx.name_key(handle)])?;
1517 // Write size for array type
1518 if let TypeInner::Array { base, size, .. } = self.module.types[local.ty].inner {
1519 self.write_array_size(base, size)?;
1520 }
1521 // Write the local initializer if needed
1522 if let Some(init) = local.init {
1523 // Put the equal signal only if there's a initializer
1524 // The leading and trailing spaces aren't needed but help with readability
1525 write!(self.out, " = ")?;
1526
1527 // Write the constant
1528 // `write_constant` adds no trailing or leading space/newline
1529 self.write_const_expr(init)?;
1530 } else if is_value_init_supported(self.module, local.ty) {
1531 write!(self.out, " = ")?;
1532 self.write_zero_init_value(local.ty)?;
1533 }
1534
1535 // Finish the local with `;` and add a newline (only for readability)
1536 writeln!(self.out, ";")?
1537 }
1538
1539 // Write the function body (statement list)
1540 for sta in func.body.iter() {
1541 // Write a statement, the indentation should always be 1 when writing the function body
1542 // `write_stmt` adds a newline
1543 self.write_stmt(sta, &ctx, back::Level(1))?;
1544 }
1545
1546 // Close braces and add a newline
1547 writeln!(self.out, "}}")?;
1548
1549 Ok(())
1550 }
1551
1552 fn write_workgroup_variables_initialization(
1553 &mut self,
1554 ctx: &back::FunctionCtx,
1555 ) -> BackendResult {
1556 let mut vars = self
1557 .module
1558 .global_variables
1559 .iter()
1560 .filter(|&(handle, var)| {
1561 !ctx.info[handle].is_empty() && var.space == crate::AddressSpace::WorkGroup
1562 })
1563 .peekable();
1564
1565 if vars.peek().is_some() {
1566 let level = back::Level(1);
1567
1568 writeln!(self.out, "{level}if (gl_LocalInvocationID == uvec3(0u)) {{")?;
1569
1570 for (handle, var) in vars {
1571 let name = &self.names[&NameKey::GlobalVariable(handle)];
1572 write!(self.out, "{}{} = ", level.next(), name)?;
1573 self.write_zero_init_value(var.ty)?;
1574 writeln!(self.out, ";")?;
1575 }
1576
1577 writeln!(self.out, "{level}}}")?;
1578 self.write_barrier(crate::Barrier::WORK_GROUP, level)?;
1579 }
1580
1581 Ok(())
1582 }
1583
1584 /// Write a list of comma separated `T` values using a writer function `F`.
1585 ///
1586 /// The writer function `F` receives a mutable reference to `self` that if needed won't cause
1587 /// borrow checker issues (using for example a closure with `self` will cause issues), the
1588 /// second argument is the 0 based index of the element on the list, and the last element is
1589 /// a reference to the element `T` being written
1590 ///
1591 /// # Notes
1592 /// - Adds no newlines or leading/trailing whitespace
1593 /// - The last element won't have a trailing `,`
1594 fn write_slice<T, F: FnMut(&mut Self, u32, &T) -> BackendResult>(
1595 &mut self,
1596 data: &[T],
1597 mut f: F,
1598 ) -> BackendResult {
1599 // Loop through `data` invoking `f` for each element
1600 for (index, item) in data.iter().enumerate() {
1601 if index != 0 {
1602 write!(self.out, ", ")?;
1603 }
1604 f(self, index as u32, item)?;
1605 }
1606
1607 Ok(())
1608 }
1609
1610 /// Helper method used to write global constants
1611 fn write_global_constant(&mut self, handle: Handle<crate::Constant>) -> BackendResult {
1612 write!(self.out, "const ")?;
1613 let constant = &self.module.constants[handle];
1614 self.write_type(constant.ty)?;
1615 let name = &self.names[&NameKey::Constant(handle)];
1616 write!(self.out, " {name}")?;
1617 if let TypeInner::Array { base, size, .. } = self.module.types[constant.ty].inner {
1618 self.write_array_size(base, size)?;
1619 }
1620 write!(self.out, " = ")?;
1621 self.write_const_expr(constant.init)?;
1622 writeln!(self.out, ";")?;
1623 Ok(())
1624 }
1625
1626 /// Helper method used to output a dot product as an arithmetic expression
1627 ///
1628 fn write_dot_product(
1629 &mut self,
1630 arg: Handle<crate::Expression>,
1631 arg1: Handle<crate::Expression>,
1632 size: usize,
1633 ctx: &back::FunctionCtx<'_>,
1634 ) -> BackendResult {
1635 // Write parantheses around the dot product expression to prevent operators
1636 // with different precedences from applying earlier.
1637 write!(self.out, "(")?;
1638
1639 // Cycle trough all the components of the vector
1640 for index in 0..size {
1641 let component = back::COMPONENTS[index];
1642 // Write the addition to the previous product
1643 // This will print an extra '+' at the beginning but that is fine in glsl
1644 write!(self.out, " + ")?;
1645 // Write the first vector expression, this expression is marked to be
1646 // cached so unless it can't be cached (for example, it's a Constant)
1647 // it shouldn't produce large expressions.
1648 self.write_expr(arg, ctx)?;
1649 // Access the current component on the first vector
1650 write!(self.out, ".{component} * ")?;
1651 // Write the second vector expression, this expression is marked to be
1652 // cached so unless it can't be cached (for example, it's a Constant)
1653 // it shouldn't produce large expressions.
1654 self.write_expr(arg1, ctx)?;
1655 // Access the current component on the second vector
1656 write!(self.out, ".{component}")?;
1657 }
1658
1659 write!(self.out, ")")?;
1660 Ok(())
1661 }
1662
1663 /// Helper method used to write structs
1664 ///
1665 /// # Notes
1666 /// Ends in a newline
1667 fn write_struct_body(
1668 &mut self,
1669 handle: Handle<crate::Type>,
1670 members: &[crate::StructMember],
1671 ) -> BackendResult {
1672 // glsl structs are written as in C
1673 // `struct name() { members };`
1674 // | `struct` is a keyword
1675 // | `name` is the struct name
1676 // | `members` is a semicolon separated list of `type name`
1677 // | `type` is the member type
1678 // | `name` is the member name
1679 writeln!(self.out, "{{")?;
1680
1681 for (idx, member) in members.iter().enumerate() {
1682 // The indentation is only for readability
1683 write!(self.out, "{}", back::INDENT)?;
1684
1685 match self.module.types[member.ty].inner {
1686 TypeInner::Array {
1687 base,
1688 size,
1689 stride: _,
1690 } => {
1691 self.write_type(base)?;
1692 write!(
1693 self.out,
1694 " {}",
1695 &self.names[&NameKey::StructMember(handle, idx as u32)]
1696 )?;
1697 // Write [size]
1698 self.write_array_size(base, size)?;
1699 // Newline is important
1700 writeln!(self.out, ";")?;
1701 }
1702 _ => {
1703 // Write the member type
1704 // Adds no trailing space
1705 self.write_type(member.ty)?;
1706
1707 // Write the member name and put a semicolon
1708 // The leading space is important
1709 // All members must have a semicolon even the last one
1710 writeln!(
1711 self.out,
1712 " {};",
1713 &self.names[&NameKey::StructMember(handle, idx as u32)]
1714 )?;
1715 }
1716 }
1717 }
1718
1719 write!(self.out, "}}")?;
1720 Ok(())
1721 }
1722
1723 /// Helper method used to write statements
1724 ///
1725 /// # Notes
1726 /// Always adds a newline
1727 fn write_stmt(
1728 &mut self,
1729 sta: &crate::Statement,
1730 ctx: &back::FunctionCtx,
1731 level: back::Level,
1732 ) -> BackendResult {
1733 use crate::Statement;
1734
1735 match *sta {
1736 // This is where we can generate intermediate constants for some expression types.
1737 Statement::Emit(ref range) => {
1738 for handle in range.clone() {
1739 let info = &ctx.info[handle];
1740 let ptr_class = info.ty.inner_with(&self.module.types).pointer_space();
1741 let expr_name = if ptr_class.is_some() {
1742 // GLSL can't save a pointer-valued expression in a variable,
1743 // but we shouldn't ever need to: they should never be named expressions,
1744 // and none of the expression types flagged by bake_ref_count can be pointer-valued.
1745 None
1746 } else if let Some(name) = ctx.named_expressions.get(&handle) {
1747 // Front end provides names for all variables at the start of writing.
1748 // But we write them to step by step. We need to recache them
1749 // Otherwise, we could accidentally write variable name instead of full expression.
1750 // Also, we use sanitized names! It defense backend from generating variable with name from reserved keywords.
1751 Some(self.namer.call(name))
1752 } else if self.need_bake_expressions.contains(&handle) {
1753 Some(format!("{}{}", back::BAKE_PREFIX, handle.index()))
1754 } else {
1755 None
1756 };
1757
1758 // If we are going to write an `ImageLoad` next and the target image
1759 // is sampled and we are using the `Restrict` policy for bounds
1760 // checking images we need to write a local holding the clamped lod.
1761 if let crate::Expression::ImageLoad {
1762 image,
1763 level: Some(level_expr),
1764 ..
1765 } = ctx.expressions[handle]
1766 {
1767 if let TypeInner::Image {
1768 class: crate::ImageClass::Sampled { .. },
1769 ..
1770 } = *ctx.info[image].ty.inner_with(&self.module.types)
1771 {
1772 if let proc::BoundsCheckPolicy::Restrict = self.policies.image_load {
1773 write!(self.out, "{level}")?;
1774 self.write_clamped_lod(ctx, handle, image, level_expr)?
1775 }
1776 }
1777 }
1778
1779 if let Some(name) = expr_name {
1780 write!(self.out, "{level}")?;
1781 self.write_named_expr(handle, name, handle, ctx)?;
1782 }
1783 }
1784 }
1785 // Blocks are simple we just need to write the block statements between braces
1786 // We could also just print the statements but this is more readable and maps more
1787 // closely to the IR
1788 Statement::Block(ref block) => {
1789 write!(self.out, "{level}")?;
1790 writeln!(self.out, "{{")?;
1791 for sta in block.iter() {
1792 // Increase the indentation to help with readability
1793 self.write_stmt(sta, ctx, level.next())?
1794 }
1795 writeln!(self.out, "{level}}}")?
1796 }
1797 // Ifs are written as in C:
1798 // ```
1799 // if(condition) {
1800 // accept
1801 // } else {
1802 // reject
1803 // }
1804 // ```
1805 Statement::If {
1806 condition,
1807 ref accept,
1808 ref reject,
1809 } => {
1810 write!(self.out, "{level}")?;
1811 write!(self.out, "if (")?;
1812 self.write_expr(condition, ctx)?;
1813 writeln!(self.out, ") {{")?;
1814
1815 for sta in accept {
1816 // Increase indentation to help with readability
1817 self.write_stmt(sta, ctx, level.next())?;
1818 }
1819
1820 // If there are no statements in the reject block we skip writing it
1821 // This is only for readability
1822 if !reject.is_empty() {
1823 writeln!(self.out, "{level}}} else {{")?;
1824
1825 for sta in reject {
1826 // Increase indentation to help with readability
1827 self.write_stmt(sta, ctx, level.next())?;
1828 }
1829 }
1830
1831 writeln!(self.out, "{level}}}")?
1832 }
1833 // Switch are written as in C:
1834 // ```
1835 // switch (selector) {
1836 // // Fallthrough
1837 // case label:
1838 // block
1839 // // Non fallthrough
1840 // case label:
1841 // block
1842 // break;
1843 // default:
1844 // block
1845 // }
1846 // ```
1847 // Where the `default` case happens isn't important but we put it last
1848 // so that we don't need to print a `break` for it
1849 Statement::Switch {
1850 selector,
1851 ref cases,
1852 } => {
1853 // Start the switch
1854 write!(self.out, "{level}")?;
1855 write!(self.out, "switch(")?;
1856 self.write_expr(selector, ctx)?;
1857 writeln!(self.out, ") {{")?;
1858
1859 // Write all cases
1860 let l2 = level.next();
1861 for case in cases {
1862 match case.value {
1863 crate::SwitchValue::I32(value) => write!(self.out, "{l2}case {value}:")?,
1864 crate::SwitchValue::U32(value) => write!(self.out, "{l2}case {value}u:")?,
1865 crate::SwitchValue::Default => write!(self.out, "{l2}default:")?,
1866 }
1867
1868 let write_block_braces = !(case.fall_through && case.body.is_empty());
1869 if write_block_braces {
1870 writeln!(self.out, " {{")?;
1871 } else {
1872 writeln!(self.out)?;
1873 }
1874
1875 for sta in case.body.iter() {
1876 self.write_stmt(sta, ctx, l2.next())?;
1877 }
1878
1879 if !case.fall_through && case.body.last().map_or(true, |s| !s.is_terminator()) {
1880 writeln!(self.out, "{}break;", l2.next())?;
1881 }
1882
1883 if write_block_braces {
1884 writeln!(self.out, "{l2}}}")?;
1885 }
1886 }
1887
1888 writeln!(self.out, "{level}}}")?
1889 }
1890 // Loops in naga IR are based on wgsl loops, glsl can emulate the behaviour by using a
1891 // while true loop and appending the continuing block to the body resulting on:
1892 // ```
1893 // bool loop_init = true;
1894 // while(true) {
1895 // if (!loop_init) { <continuing> }
1896 // loop_init = false;
1897 // <body>
1898 // }
1899 // ```
1900 Statement::Loop {
1901 ref body,
1902 ref continuing,
1903 break_if,
1904 } => {
1905 if !continuing.is_empty() || break_if.is_some() {
1906 let gate_name = self.namer.call("loop_init");
1907 writeln!(self.out, "{level}bool {gate_name} = true;")?;
1908 writeln!(self.out, "{level}while(true) {{")?;
1909 let l2 = level.next();
1910 let l3 = l2.next();
1911 writeln!(self.out, "{l2}if (!{gate_name}) {{")?;
1912 for sta in continuing {
1913 self.write_stmt(sta, ctx, l3)?;
1914 }
1915 if let Some(condition) = break_if {
1916 write!(self.out, "{l3}if (")?;
1917 self.write_expr(condition, ctx)?;
1918 writeln!(self.out, ") {{")?;
1919 writeln!(self.out, "{}break;", l3.next())?;
1920 writeln!(self.out, "{l3}}}")?;
1921 }
1922 writeln!(self.out, "{l2}}}")?;
1923 writeln!(self.out, "{}{} = false;", level.next(), gate_name)?;
1924 } else {
1925 writeln!(self.out, "{level}while(true) {{")?;
1926 }
1927 for sta in body {
1928 self.write_stmt(sta, ctx, level.next())?;
1929 }
1930 writeln!(self.out, "{level}}}")?
1931 }
1932 // Break, continue and return as written as in C
1933 // `break;`
1934 Statement::Break => {
1935 write!(self.out, "{level}")?;
1936 writeln!(self.out, "break;")?
1937 }
1938 // `continue;`
1939 Statement::Continue => {
1940 write!(self.out, "{level}")?;
1941 writeln!(self.out, "continue;")?
1942 }
1943 // `return expr;`, `expr` is optional
1944 Statement::Return { value } => {
1945 write!(self.out, "{level}")?;
1946 match ctx.ty {
1947 back::FunctionType::Function(_) => {
1948 write!(self.out, "return")?;
1949 // Write the expression to be returned if needed
1950 if let Some(expr) = value {
1951 write!(self.out, " ")?;
1952 self.write_expr(expr, ctx)?;
1953 }
1954 writeln!(self.out, ";")?;
1955 }
1956 back::FunctionType::EntryPoint(ep_index) => {
1957 let mut has_point_size = false;
1958 let ep = &self.module.entry_points[ep_index as usize];
1959 if let Some(ref result) = ep.function.result {
1960 let value = value.unwrap();
1961 match self.module.types[result.ty].inner {
1962 crate::TypeInner::Struct { ref members, .. } => {
1963 let temp_struct_name = match ctx.expressions[value] {
1964 crate::Expression::Compose { .. } => {
1965 let return_struct = "_tmp_return";
1966 write!(
1967 self.out,
1968 "{} {} = ",
1969 &self.names[&NameKey::Type(result.ty)],
1970 return_struct
1971 )?;
1972 self.write_expr(value, ctx)?;
1973 writeln!(self.out, ";")?;
1974 write!(self.out, "{level}")?;
1975 Some(return_struct)
1976 }
1977 _ => None,
1978 };
1979
1980 for (index, member) in members.iter().enumerate() {
1981 if let Some(crate::Binding::BuiltIn(
1982 crate::BuiltIn::PointSize,
1983 )) = member.binding
1984 {
1985 has_point_size = true;
1986 }
1987
1988 let varying_name = VaryingName {
1989 binding: member.binding.as_ref().unwrap(),
1990 stage: ep.stage,
1991 output: true,
1992 targetting_webgl: self.options.version.is_webgl(),
1993 };
1994 write!(self.out, "{varying_name} = ")?;
1995
1996 if let Some(struct_name) = temp_struct_name {
1997 write!(self.out, "{struct_name}")?;
1998 } else {
1999 self.write_expr(value, ctx)?;
2000 }
2001
2002 // Write field name
2003 writeln!(
2004 self.out,
2005 ".{};",
2006 &self.names
2007 [&NameKey::StructMember(result.ty, index as u32)]
2008 )?;
2009 write!(self.out, "{level}")?;
2010 }
2011 }
2012 _ => {
2013 let name = VaryingName {
2014 binding: result.binding.as_ref().unwrap(),
2015 stage: ep.stage,
2016 output: true,
2017 targetting_webgl: self.options.version.is_webgl(),
2018 };
2019 write!(self.out, "{name} = ")?;
2020 self.write_expr(value, ctx)?;
2021 writeln!(self.out, ";")?;
2022 write!(self.out, "{level}")?;
2023 }
2024 }
2025 }
2026
2027 let is_vertex_stage = self.module.entry_points[ep_index as usize].stage
2028 == ShaderStage::Vertex;
2029 if is_vertex_stage
2030 && self
2031 .options
2032 .writer_flags
2033 .contains(WriterFlags::ADJUST_COORDINATE_SPACE)
2034 {
2035 writeln!(
2036 self.out,
2037 "gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w);",
2038 )?;
2039 write!(self.out, "{level}")?;
2040 }
2041
2042 if is_vertex_stage
2043 && self
2044 .options
2045 .writer_flags
2046 .contains(WriterFlags::FORCE_POINT_SIZE)
2047 && !has_point_size
2048 {
2049 writeln!(self.out, "gl_PointSize = 1.0;")?;
2050 write!(self.out, "{level}")?;
2051 }
2052 writeln!(self.out, "return;")?;
2053 }
2054 }
2055 }
2056 // This is one of the places were glsl adds to the syntax of C in this case the discard
2057 // keyword which ceases all further processing in a fragment shader, it's called OpKill
2058 // in spir-v that's why it's called `Statement::Kill`
2059 Statement::Kill => writeln!(self.out, "{level}discard;")?,
2060 Statement::Barrier(flags) => {
2061 self.write_barrier(flags, level)?;
2062 }
2063 // Stores in glsl are just variable assignments written as `pointer = value;`
2064 Statement::Store { pointer, value } => {
2065 write!(self.out, "{level}")?;
2066 self.write_expr(pointer, ctx)?;
2067 write!(self.out, " = ")?;
2068 self.write_expr(value, ctx)?;
2069 writeln!(self.out, ";")?
2070 }
2071 Statement::WorkGroupUniformLoad { pointer, result } => {
2072 // GLSL doesn't have pointers, which means that this backend needs to ensure that
2073 // the actual "loading" is happening between the two barriers.
2074 // This is done in `Emit` by never emitting a variable name for pointer variables
2075 self.write_barrier(crate::Barrier::WORK_GROUP, level)?;
2076
2077 let result_name = format!("{}{}", back::BAKE_PREFIX, result.index());
2078 write!(self.out, "{level}")?;
2079 // Expressions cannot have side effects, so just writing the expression here is fine.
2080 self.write_named_expr(pointer, result_name, result, ctx)?;
2081
2082 self.write_barrier(crate::Barrier::WORK_GROUP, level)?;
2083 }
2084 // Stores a value into an image.
2085 Statement::ImageStore {
2086 image,
2087 coordinate,
2088 array_index,
2089 value,
2090 } => {
2091 write!(self.out, "{level}")?;
2092 self.write_image_store(ctx, image, coordinate, array_index, value)?
2093 }
2094 // A `Call` is written `name(arguments)` where `arguments` is a comma separated expressions list
2095 Statement::Call {
2096 function,
2097 ref arguments,
2098 result,
2099 } => {
2100 write!(self.out, "{level}")?;
2101 if let Some(expr) = result {
2102 let name = format!("{}{}", back::BAKE_PREFIX, expr.index());
2103 let result = self.module.functions[function].result.as_ref().unwrap();
2104 self.write_type(result.ty)?;
2105 write!(self.out, " {name}")?;
2106 if let TypeInner::Array { base, size, .. } = self.module.types[result.ty].inner
2107 {
2108 self.write_array_size(base, size)?
2109 }
2110 write!(self.out, " = ")?;
2111 self.named_expressions.insert(expr, name);
2112 }
2113 write!(self.out, "{}(", &self.names[&NameKey::Function(function)])?;
2114 let arguments: Vec<_> = arguments
2115 .iter()
2116 .enumerate()
2117 .filter_map(|(i, arg)| {
2118 let arg_ty = self.module.functions[function].arguments[i].ty;
2119 match self.module.types[arg_ty].inner {
2120 TypeInner::Sampler { .. } => None,
2121 _ => Some(*arg),
2122 }
2123 })
2124 .collect();
2125 self.write_slice(&arguments, |this, _, arg| this.write_expr(*arg, ctx))?;
2126 writeln!(self.out, ");")?
2127 }
2128 Statement::Atomic {
2129 pointer,
2130 ref fun,
2131 value,
2132 result,
2133 } => {
2134 write!(self.out, "{level}")?;
2135 let res_name = format!("{}{}", back::BAKE_PREFIX, result.index());
2136 let res_ty = ctx.info[result].ty.inner_with(&self.module.types);
2137 self.write_value_type(res_ty)?;
2138 write!(self.out, " {res_name} = ")?;
2139 self.named_expressions.insert(result, res_name);
2140
2141 let fun_str = fun.to_glsl();
2142 write!(self.out, "atomic{fun_str}(")?;
2143 self.write_expr(pointer, ctx)?;
2144 write!(self.out, ", ")?;
2145 // handle the special cases
2146 match *fun {
2147 crate::AtomicFunction::Subtract => {
2148 // we just wrote `InterlockedAdd`, so negate the argument
2149 write!(self.out, "-")?;
2150 }
2151 crate::AtomicFunction::Exchange { compare: Some(_) } => {
2152 return Err(Error::Custom(
2153 "atomic CompareExchange is not implemented".to_string(),
2154 ));
2155 }
2156 _ => {}
2157 }
2158 self.write_expr(value, ctx)?;
2159 writeln!(self.out, ");")?;
2160 }
2161 Statement::RayQuery { .. } => unreachable!(),
2162 }
2163
2164 Ok(())
2165 }
2166
2167 /// Write a const expression.
2168 ///
2169 /// Write `expr`, a handle to an [`Expression`] in the current [`Module`]'s
2170 /// constant expression arena, as GLSL expression.
2171 ///
2172 /// # Notes
2173 /// Adds no newlines or leading/trailing whitespace
2174 ///
2175 /// [`Expression`]: crate::Expression
2176 /// [`Module`]: crate::Module
2177 fn write_const_expr(&mut self, expr: Handle<crate::Expression>) -> BackendResult {
2178 self.write_possibly_const_expr(expr, &self.module.const_expressions, |writer, expr| {
2179 writer.write_const_expr(expr)
2180 })
2181 }
2182
2183 /// Write [`Expression`] variants that can occur in both runtime and const expressions.
2184 ///
2185 /// Write `expr`, a handle to an [`Expression`] in the arena `expressions`,
2186 /// as as GLSL expression. This must be one of the [`Expression`] variants
2187 /// that is allowed to occur in constant expressions.
2188 ///
2189 /// Use `write_expression` to write subexpressions.
2190 ///
2191 /// This is the common code for `write_expr`, which handles arbitrary
2192 /// runtime expressions, and `write_const_expr`, which only handles
2193 /// const-expressions. Each of those callers passes itself (essentially) as
2194 /// the `write_expression` callback, so that subexpressions are restricted
2195 /// to the appropriate variants.
2196 ///
2197 /// # Notes
2198 /// Adds no newlines or leading/trailing whitespace
2199 ///
2200 /// [`Expression`]: crate::Expression
2201 fn write_possibly_const_expr<E>(
2202 &mut self,
2203 expr: Handle<crate::Expression>,
2204 expressions: &crate::Arena<crate::Expression>,
2205 write_expression: E,
2206 ) -> BackendResult
2207 where
2208 E: Fn(&mut Self, Handle<crate::Expression>) -> BackendResult,
2209 {
2210 use crate::Expression;
2211
2212 match expressions[expr] {
2213 Expression::Literal(literal) => {
2214 match literal {
2215 // Floats are written using `Debug` instead of `Display` because it always appends the
2216 // decimal part even it's zero which is needed for a valid glsl float constant
2217 crate::Literal::F64(value) => write!(self.out, "{:?}LF", value)?,
2218 crate::Literal::F32(value) => write!(self.out, "{:?}", value)?,
2219 // Unsigned integers need a `u` at the end
2220 //
2221 // While `core` doesn't necessarily need it, it's allowed and since `es` needs it we
2222 // always write it as the extra branch wouldn't have any benefit in readability
2223 crate::Literal::U32(value) => write!(self.out, "{}u", value)?,
2224 crate::Literal::I32(value) => write!(self.out, "{}", value)?,
2225 crate::Literal::Bool(value) => write!(self.out, "{}", value)?,
2226 }
2227 }
2228 Expression::Constant(handle) => {
2229 let constant = &self.module.constants[handle];
2230 if constant.name.is_some() {
2231 write!(self.out, "{}", self.names[&NameKey::Constant(handle)])?;
2232 } else {
2233 self.write_const_expr(constant.init)?;
2234 }
2235 }
2236 Expression::ZeroValue(ty) => {
2237 self.write_zero_init_value(ty)?;
2238 }
2239 Expression::Compose { ty, ref components } => {
2240 self.write_type(ty)?;
2241
2242 if let TypeInner::Array { base, size, .. } = self.module.types[ty].inner {
2243 self.write_array_size(base, size)?;
2244 }
2245
2246 write!(self.out, "(")?;
2247 for (index, component) in components.iter().enumerate() {
2248 if index != 0 {
2249 write!(self.out, ", ")?;
2250 }
2251 write_expression(self, *component)?;
2252 }
2253 write!(self.out, ")")?
2254 }
2255 _ => unreachable!(),
2256 }
2257
2258 Ok(())
2259 }
2260
2261 /// Helper method to write expressions
2262 ///
2263 /// # Notes
2264 /// Doesn't add any newlines or leading/trailing spaces
2265 fn write_expr(
2266 &mut self,
2267 expr: Handle<crate::Expression>,
2268 ctx: &back::FunctionCtx<'_>,
2269 ) -> BackendResult {
2270 use crate::Expression;
2271
2272 if let Some(name) = self.named_expressions.get(&expr) {
2273 write!(self.out, "{name}")?;
2274 return Ok(());
2275 }
2276
2277 match ctx.expressions[expr] {
2278 Expression::Literal(_)
2279 | Expression::Constant(_)
2280 | Expression::ZeroValue(_)
2281 | Expression::Compose { .. } => {
2282 self.write_possibly_const_expr(expr, ctx.expressions, |writer, expr| {
2283 writer.write_expr(expr, ctx)
2284 })?;
2285 }
2286 // `Access` is applied to arrays, vectors and matrices and is written as indexing
2287 Expression::Access { base, index } => {
2288 self.write_expr(base, ctx)?;
2289 write!(self.out, "[")?;
2290 self.write_expr(index, ctx)?;
2291 write!(self.out, "]")?
2292 }
2293 // `AccessIndex` is the same as `Access` except that the index is a constant and it can
2294 // be applied to structs, in this case we need to find the name of the field at that
2295 // index and write `base.field_name`
2296 Expression::AccessIndex { base, index } => {
2297 self.write_expr(base, ctx)?;
2298
2299 let base_ty_res = &ctx.info[base].ty;
2300 let mut resolved = base_ty_res.inner_with(&self.module.types);
2301 let base_ty_handle = match *resolved {
2302 TypeInner::Pointer { base, space: _ } => {
2303 resolved = &self.module.types[base].inner;
2304 Some(base)
2305 }
2306 _ => base_ty_res.handle(),
2307 };
2308
2309 match *resolved {
2310 TypeInner::Vector { .. } => {
2311 // Write vector access as a swizzle
2312 write!(self.out, ".{}", back::COMPONENTS[index as usize])?
2313 }
2314 TypeInner::Matrix { .. }
2315 | TypeInner::Array { .. }
2316 | TypeInner::ValuePointer { .. } => write!(self.out, "[{index}]")?,
2317 TypeInner::Struct { .. } => {
2318 // This will never panic in case the type is a `Struct`, this is not true
2319 // for other types so we can only check while inside this match arm
2320 let ty = base_ty_handle.unwrap();
2321
2322 write!(
2323 self.out,
2324 ".{}",
2325 &self.names[&NameKey::StructMember(ty, index)]
2326 )?
2327 }
2328 ref other => return Err(Error::Custom(format!("Cannot index {other:?}"))),
2329 }
2330 }
2331 // `Splat` needs to actually write down a vector, it's not always inferred in GLSL.
2332 Expression::Splat { size: _, value } => {
2333 let resolved = ctx.info[expr].ty.inner_with(&self.module.types);
2334 self.write_value_type(resolved)?;
2335 write!(self.out, "(")?;
2336 self.write_expr(value, ctx)?;
2337 write!(self.out, ")")?
2338 }
2339 // `Swizzle` adds a few letters behind the dot.
2340 Expression::Swizzle {
2341 size,
2342 vector,
2343 pattern,
2344 } => {
2345 self.write_expr(vector, ctx)?;
2346 write!(self.out, ".")?;
2347 for &sc in pattern[..size as usize].iter() {
2348 self.out.write_char(back::COMPONENTS[sc as usize])?;
2349 }
2350 }
2351 // Function arguments are written as the argument name
2352 Expression::FunctionArgument(pos) => {
2353 write!(self.out, "{}", &self.names[&ctx.argument_key(pos)])?
2354 }
2355 // Global variables need some special work for their name but
2356 // `get_global_name` does the work for us
2357 Expression::GlobalVariable(handle) => {
2358 let global = &self.module.global_variables[handle];
2359 self.write_global_name(handle, global)?
2360 }
2361 // A local is written as it's name
2362 Expression::LocalVariable(handle) => {
2363 write!(self.out, "{}", self.names[&ctx.name_key(handle)])?
2364 }
2365 // glsl has no pointers so there's no load operation, just write the pointer expression
2366 Expression::Load { pointer } => self.write_expr(pointer, ctx)?,
2367 // `ImageSample` is a bit complicated compared to the rest of the IR.
2368 //
2369 // First there are three variations depending whether the sample level is explicitly set,
2370 // if it's automatic or it it's bias:
2371 // `texture(image, coordinate)` - Automatic sample level
2372 // `texture(image, coordinate, bias)` - Bias sample level
2373 // `textureLod(image, coordinate, level)` - Zero or Exact sample level
2374 //
2375 // Furthermore if `depth_ref` is some we need to append it to the coordinate vector
2376 Expression::ImageSample {
2377 image,
2378 sampler: _, //TODO?
2379 gather,
2380 coordinate,
2381 array_index,
2382 offset,
2383 level,
2384 depth_ref,
2385 } => {
2386 let dim = match *ctx.info[image].ty.inner_with(&self.module.types) {
2387 TypeInner::Image { dim, .. } => dim,
2388 _ => unreachable!(),
2389 };
2390
2391 if dim == crate::ImageDimension::Cube
2392 && array_index.is_some()
2393 && depth_ref.is_some()
2394 {
2395 match level {
2396 crate::SampleLevel::Zero
2397 | crate::SampleLevel::Exact(_)
2398 | crate::SampleLevel::Gradient { .. }
2399 | crate::SampleLevel::Bias(_) => {
2400 return Err(Error::Custom(String::from(
2401 "gsamplerCubeArrayShadow isn't supported in textureGrad, \
2402 textureLod or texture with bias",
2403 )))
2404 }
2405 crate::SampleLevel::Auto => {}
2406 }
2407 }
2408
2409 // textureLod on sampler2DArrayShadow and samplerCubeShadow does not exist in GLSL.
2410 // To emulate this, we will have to use textureGrad with a constant gradient of 0.
2411 let workaround_lod_array_shadow_as_grad = (array_index.is_some()
2412 || dim == crate::ImageDimension::Cube)
2413 && depth_ref.is_some()
2414 && gather.is_none()
2415 && !self
2416 .options
2417 .writer_flags
2418 .contains(WriterFlags::TEXTURE_SHADOW_LOD);
2419
2420 //Write the function to be used depending on the sample level
2421 let fun_name = match level {
2422 crate::SampleLevel::Zero if gather.is_some() => "textureGather",
2423 crate::SampleLevel::Auto | crate::SampleLevel::Bias(_) => "texture",
2424 crate::SampleLevel::Zero | crate::SampleLevel::Exact(_) => {
2425 if workaround_lod_array_shadow_as_grad {
2426 "textureGrad"
2427 } else {
2428 "textureLod"
2429 }
2430 }
2431 crate::SampleLevel::Gradient { .. } => "textureGrad",
2432 };
2433 let offset_name = match offset {
2434 Some(_) => "Offset",
2435 None => "",
2436 };
2437
2438 write!(self.out, "{fun_name}{offset_name}(")?;
2439
2440 // Write the image that will be used
2441 self.write_expr(image, ctx)?;
2442 // The space here isn't required but it helps with readability
2443 write!(self.out, ", ")?;
2444
2445 // We need to get the coordinates vector size to later build a vector that's `size + 1`
2446 // if `depth_ref` is some, if it isn't a vector we panic as that's not a valid expression
2447 let mut coord_dim = match *ctx.info[coordinate].ty.inner_with(&self.module.types) {
2448 TypeInner::Vector { size, .. } => size as u8,
2449 TypeInner::Scalar { .. } => 1,
2450 _ => unreachable!(),
2451 };
2452
2453 if array_index.is_some() {
2454 coord_dim += 1;
2455 }
2456 let merge_depth_ref = depth_ref.is_some() && gather.is_none() && coord_dim < 4;
2457 if merge_depth_ref {
2458 coord_dim += 1;
2459 }
2460
2461 let tex_1d_hack = dim == crate::ImageDimension::D1 && self.options.version.is_es();
2462 let is_vec = tex_1d_hack || coord_dim != 1;
2463 // Compose a new texture coordinates vector
2464 if is_vec {
2465 write!(self.out, "vec{}(", coord_dim + tex_1d_hack as u8)?;
2466 }
2467 self.write_expr(coordinate, ctx)?;
2468 if tex_1d_hack {
2469 write!(self.out, ", 0.0")?;
2470 }
2471 if let Some(expr) = array_index {
2472 write!(self.out, ", ")?;
2473 self.write_expr(expr, ctx)?;
2474 }
2475 if merge_depth_ref {
2476 write!(self.out, ", ")?;
2477 self.write_expr(depth_ref.unwrap(), ctx)?;
2478 }
2479 if is_vec {
2480 write!(self.out, ")")?;
2481 }
2482
2483 if let (Some(expr), false) = (depth_ref, merge_depth_ref) {
2484 write!(self.out, ", ")?;
2485 self.write_expr(expr, ctx)?;
2486 }
2487
2488 match level {
2489 // Auto needs no more arguments
2490 crate::SampleLevel::Auto => (),
2491 // Zero needs level set to 0
2492 crate::SampleLevel::Zero => {
2493 if workaround_lod_array_shadow_as_grad {
2494 let vec_dim = match dim {
2495 crate::ImageDimension::Cube => 3,
2496 _ => 2,
2497 };
2498 write!(self.out, ", vec{vec_dim}(0.0), vec{vec_dim}(0.0)")?;
2499 } else if gather.is_none() {
2500 write!(self.out, ", 0.0")?;
2501 }
2502 }
2503 // Exact and bias require another argument
2504 crate::SampleLevel::Exact(expr) => {
2505 if workaround_lod_array_shadow_as_grad {
2506 log::warn!("Unable to `textureLod` a shadow array, ignoring the LOD");
2507 write!(self.out, ", vec2(0,0), vec2(0,0)")?;
2508 } else {
2509 write!(self.out, ", ")?;
2510 self.write_expr(expr, ctx)?;
2511 }
2512 }
2513 crate::SampleLevel::Bias(_) => {
2514 // This needs to be done after the offset writing
2515 }
2516 crate::SampleLevel::Gradient { x, y } => {
2517 // If we are using sampler2D to replace sampler1D, we also
2518 // need to make sure to use vec2 gradients
2519 if tex_1d_hack {
2520 write!(self.out, ", vec2(")?;
2521 self.write_expr(x, ctx)?;
2522 write!(self.out, ", 0.0)")?;
2523 write!(self.out, ", vec2(")?;
2524 self.write_expr(y, ctx)?;
2525 write!(self.out, ", 0.0)")?;
2526 } else {
2527 write!(self.out, ", ")?;
2528 self.write_expr(x, ctx)?;
2529 write!(self.out, ", ")?;
2530 self.write_expr(y, ctx)?;
2531 }
2532 }
2533 }
2534
2535 if let Some(constant) = offset {
2536 write!(self.out, ", ")?;
2537 if tex_1d_hack {
2538 write!(self.out, "ivec2(")?;
2539 }
2540 self.write_const_expr(constant)?;
2541 if tex_1d_hack {
2542 write!(self.out, ", 0)")?;
2543 }
2544 }
2545
2546 // Bias is always the last argument
2547 if let crate::SampleLevel::Bias(expr) = level {
2548 write!(self.out, ", ")?;
2549 self.write_expr(expr, ctx)?;
2550 }
2551
2552 if let (Some(component), None) = (gather, depth_ref) {
2553 write!(self.out, ", {}", component as usize)?;
2554 }
2555
2556 // End the function
2557 write!(self.out, ")")?
2558 }
2559 Expression::ImageLoad {
2560 image,
2561 coordinate,
2562 array_index,
2563 sample,
2564 level,
2565 } => self.write_image_load(expr, ctx, image, coordinate, array_index, sample, level)?,
2566 // Query translates into one of the:
2567 // - textureSize/imageSize
2568 // - textureQueryLevels
2569 // - textureSamples/imageSamples
2570 Expression::ImageQuery { image, query } => {
2571 use crate::ImageClass;
2572
2573 // This will only panic if the module is invalid
2574 let (dim, class) = match *ctx.info[image].ty.inner_with(&self.module.types) {
2575 TypeInner::Image {
2576 dim,
2577 arrayed: _,
2578 class,
2579 } => (dim, class),
2580 _ => unreachable!(),
2581 };
2582 let components = match dim {
2583 crate::ImageDimension::D1 => 1,
2584 crate::ImageDimension::D2 => 2,
2585 crate::ImageDimension::D3 => 3,
2586 crate::ImageDimension::Cube => 2,
2587 };
2588
2589 if let crate::ImageQuery::Size { .. } = query {
2590 match components {
2591 1 => write!(self.out, "uint(")?,
2592 _ => write!(self.out, "uvec{components}(")?,
2593 }
2594 } else {
2595 write!(self.out, "uint(")?;
2596 }
2597
2598 match query {
2599 crate::ImageQuery::Size { level } => {
2600 match class {
2601 ImageClass::Sampled { multi, .. } | ImageClass::Depth { multi } => {
2602 write!(self.out, "textureSize(")?;
2603 self.write_expr(image, ctx)?;
2604 if let Some(expr) = level {
2605 let cast_to_int = matches!(
2606 *ctx.info[expr].ty.inner_with(&self.module.types),
2607 crate::TypeInner::Scalar {
2608 kind: crate::ScalarKind::Uint,
2609 ..
2610 }
2611 );
2612
2613 write!(self.out, ", ")?;
2614
2615 if cast_to_int {
2616 write!(self.out, "int(")?;
2617 }
2618
2619 self.write_expr(expr, ctx)?;
2620
2621 if cast_to_int {
2622 write!(self.out, ")")?;
2623 }
2624 } else if !multi {
2625 // All textureSize calls requires an lod argument
2626 // except for multisampled samplers
2627 write!(self.out, ", 0")?;
2628 }
2629 }
2630 ImageClass::Storage { .. } => {
2631 write!(self.out, "imageSize(")?;
2632 self.write_expr(image, ctx)?;
2633 }
2634 }
2635 write!(self.out, ")")?;
2636 if components != 1 || self.options.version.is_es() {
2637 write!(self.out, ".{}", &"xyz"[..components])?;
2638 }
2639 }
2640 crate::ImageQuery::NumLevels => {
2641 write!(self.out, "textureQueryLevels(",)?;
2642 self.write_expr(image, ctx)?;
2643 write!(self.out, ")",)?;
2644 }
2645 crate::ImageQuery::NumLayers => {
2646 let fun_name = match class {
2647 ImageClass::Sampled { .. } | ImageClass::Depth { .. } => "textureSize",
2648 ImageClass::Storage { .. } => "imageSize",
2649 };
2650 write!(self.out, "{fun_name}(")?;
2651 self.write_expr(image, ctx)?;
2652 // All textureSize calls requires an lod argument
2653 // except for multisampled samplers
2654 if class.is_multisampled() {
2655 write!(self.out, ", 0")?;
2656 }
2657 write!(self.out, ")")?;
2658 if components != 1 || self.options.version.is_es() {
2659 write!(self.out, ".{}", back::COMPONENTS[components])?;
2660 }
2661 }
2662 crate::ImageQuery::NumSamples => {
2663 let fun_name = match class {
2664 ImageClass::Sampled { .. } | ImageClass::Depth { .. } => {
2665 "textureSamples"
2666 }
2667 ImageClass::Storage { .. } => "imageSamples",
2668 };
2669 write!(self.out, "{fun_name}(")?;
2670 self.write_expr(image, ctx)?;
2671 write!(self.out, ")",)?;
2672 }
2673 }
2674
2675 write!(self.out, ")")?;
2676 }
2677 // `Unary` is pretty straightforward
2678 // "-" - for `Negate`
2679 // "~" - for `Not` if it's an integer
2680 // "!" - for `Not` if it's a boolean
2681 //
2682 // We also wrap the everything in parentheses to avoid precedence issues
2683 Expression::Unary { op, expr } => {
2684 use crate::{ScalarKind as Sk, UnaryOperator as Uo};
2685
2686 let ty = ctx.info[expr].ty.inner_with(&self.module.types);
2687
2688 match *ty {
2689 TypeInner::Vector { kind: Sk::Bool, .. } => {
2690 write!(self.out, "not(")?;
2691 }
2692 _ => {
2693 let operator = match op {
2694 Uo::Negate => "-",
2695 Uo::Not => match ty.scalar_kind() {
2696 Some(Sk::Sint) | Some(Sk::Uint) => "~",
2697 Some(Sk::Bool) => "!",
2698 ref other => {
2699 return Err(Error::Custom(format!(
2700 "Cannot apply not to type {other:?}"
2701 )))
2702 }
2703 },
2704 };
2705
2706 write!(self.out, "{operator}(")?;
2707 }
2708 }
2709
2710 self.write_expr(expr, ctx)?;
2711
2712 write!(self.out, ")")?
2713 }
2714 // `Binary` we just write `left op right`, except when dealing with
2715 // comparison operations on vectors as they are implemented with
2716 // builtin functions.
2717 // Once again we wrap everything in parentheses to avoid precedence issues
2718 Expression::Binary {
2719 mut op,
2720 left,
2721 right,
2722 } => {
2723 // Holds `Some(function_name)` if the binary operation is
2724 // implemented as a function call
2725 use crate::{BinaryOperator as Bo, ScalarKind as Sk, TypeInner as Ti};
2726
2727 let left_inner = ctx.info[left].ty.inner_with(&self.module.types);
2728 let right_inner = ctx.info[right].ty.inner_with(&self.module.types);
2729
2730 let function = match (left_inner, right_inner) {
2731 (&Ti::Vector { kind, .. }, &Ti::Vector { .. }) => match op {
2732 Bo::Less
2733 | Bo::LessEqual
2734 | Bo::Greater
2735 | Bo::GreaterEqual
2736 | Bo::Equal
2737 | Bo::NotEqual => BinaryOperation::VectorCompare,
2738 Bo::Modulo if kind == Sk::Float => BinaryOperation::Modulo,
2739 Bo::And if kind == Sk::Bool => {
2740 op = crate::BinaryOperator::LogicalAnd;
2741 BinaryOperation::VectorComponentWise
2742 }
2743 Bo::InclusiveOr if kind == Sk::Bool => {
2744 op = crate::BinaryOperator::LogicalOr;
2745 BinaryOperation::VectorComponentWise
2746 }
2747 _ => BinaryOperation::Other,
2748 },
2749 _ => match (left_inner.scalar_kind(), right_inner.scalar_kind()) {
2750 (Some(Sk::Float), _) | (_, Some(Sk::Float)) => match op {
2751 Bo::Modulo => BinaryOperation::Modulo,
2752 _ => BinaryOperation::Other,
2753 },
2754 (Some(Sk::Bool), Some(Sk::Bool)) => match op {
2755 Bo::InclusiveOr => {
2756 op = crate::BinaryOperator::LogicalOr;
2757 BinaryOperation::Other
2758 }
2759 Bo::And => {
2760 op = crate::BinaryOperator::LogicalAnd;
2761 BinaryOperation::Other
2762 }
2763 _ => BinaryOperation::Other,
2764 },
2765 _ => BinaryOperation::Other,
2766 },
2767 };
2768
2769 match function {
2770 BinaryOperation::VectorCompare => {
2771 let op_str = match op {
2772 Bo::Less => "lessThan(",
2773 Bo::LessEqual => "lessThanEqual(",
2774 Bo::Greater => "greaterThan(",
2775 Bo::GreaterEqual => "greaterThanEqual(",
2776 Bo::Equal => "equal(",
2777 Bo::NotEqual => "notEqual(",
2778 _ => unreachable!(),
2779 };
2780 write!(self.out, "{op_str}")?;
2781 self.write_expr(left, ctx)?;
2782 write!(self.out, ", ")?;
2783 self.write_expr(right, ctx)?;
2784 write!(self.out, ")")?;
2785 }
2786 BinaryOperation::VectorComponentWise => {
2787 self.write_value_type(left_inner)?;
2788 write!(self.out, "(")?;
2789
2790 let size = match *left_inner {
2791 Ti::Vector { size, .. } => size,
2792 _ => unreachable!(),
2793 };
2794
2795 for i in 0..size as usize {
2796 if i != 0 {
2797 write!(self.out, ", ")?;
2798 }
2799
2800 self.write_expr(left, ctx)?;
2801 write!(self.out, ".{}", back::COMPONENTS[i])?;
2802
2803 write!(self.out, " {} ", back::binary_operation_str(op))?;
2804
2805 self.write_expr(right, ctx)?;
2806 write!(self.out, ".{}", back::COMPONENTS[i])?;
2807 }
2808
2809 write!(self.out, ")")?;
2810 }
2811 // TODO: handle undefined behavior of BinaryOperator::Modulo
2812 //
2813 // sint:
2814 // if right == 0 return 0
2815 // if left == min(type_of(left)) && right == -1 return 0
2816 // if sign(left) == -1 || sign(right) == -1 return result as defined by WGSL
2817 //
2818 // uint:
2819 // if right == 0 return 0
2820 //
2821 // float:
2822 // if right == 0 return ? see https://github.com/gpuweb/gpuweb/issues/2798
2823 BinaryOperation::Modulo => {
2824 write!(self.out, "(")?;
2825
2826 // write `e1 - e2 * trunc(e1 / e2)`
2827 self.write_expr(left, ctx)?;
2828 write!(self.out, " - ")?;
2829 self.write_expr(right, ctx)?;
2830 write!(self.out, " * ")?;
2831 write!(self.out, "trunc(")?;
2832 self.write_expr(left, ctx)?;
2833 write!(self.out, " / ")?;
2834 self.write_expr(right, ctx)?;
2835 write!(self.out, ")")?;
2836
2837 write!(self.out, ")")?;
2838 }
2839 BinaryOperation::Other => {
2840 write!(self.out, "(")?;
2841
2842 self.write_expr(left, ctx)?;
2843 write!(self.out, " {} ", back::binary_operation_str(op))?;
2844 self.write_expr(right, ctx)?;
2845
2846 write!(self.out, ")")?;
2847 }
2848 }
2849 }
2850 // `Select` is written as `condition ? accept : reject`
2851 // We wrap everything in parentheses to avoid precedence issues
2852 Expression::Select {
2853 condition,
2854 accept,
2855 reject,
2856 } => {
2857 let cond_ty = ctx.info[condition].ty.inner_with(&self.module.types);
2858 let vec_select = if let TypeInner::Vector { .. } = *cond_ty {
2859 true
2860 } else {
2861 false
2862 };
2863
2864 // TODO: Boolean mix on desktop required GL_EXT_shader_integer_mix
2865 if vec_select {
2866 // Glsl defines that for mix when the condition is a boolean the first element
2867 // is picked if condition is false and the second if condition is true
2868 write!(self.out, "mix(")?;
2869 self.write_expr(reject, ctx)?;
2870 write!(self.out, ", ")?;
2871 self.write_expr(accept, ctx)?;
2872 write!(self.out, ", ")?;
2873 self.write_expr(condition, ctx)?;
2874 } else {
2875 write!(self.out, "(")?;
2876 self.write_expr(condition, ctx)?;
2877 write!(self.out, " ? ")?;
2878 self.write_expr(accept, ctx)?;
2879 write!(self.out, " : ")?;
2880 self.write_expr(reject, ctx)?;
2881 }
2882
2883 write!(self.out, ")")?
2884 }
2885 // `Derivative` is a function call to a glsl provided function
2886 Expression::Derivative { axis, ctrl, expr } => {
2887 use crate::{DerivativeAxis as Axis, DerivativeControl as Ctrl};
2888 let fun_name = if self.options.version.supports_derivative_control() {
2889 match (axis, ctrl) {
2890 (Axis::X, Ctrl::Coarse) => "dFdxCoarse",
2891 (Axis::X, Ctrl::Fine) => "dFdxFine",
2892 (Axis::X, Ctrl::None) => "dFdx",
2893 (Axis::Y, Ctrl::Coarse) => "dFdyCoarse",
2894 (Axis::Y, Ctrl::Fine) => "dFdyFine",
2895 (Axis::Y, Ctrl::None) => "dFdy",
2896 (Axis::Width, Ctrl::Coarse) => "fwidthCoarse",
2897 (Axis::Width, Ctrl::Fine) => "fwidthFine",
2898 (Axis::Width, Ctrl::None) => "fwidth",
2899 }
2900 } else {
2901 match axis {
2902 Axis::X => "dFdx",
2903 Axis::Y => "dFdy",
2904 Axis::Width => "fwidth",
2905 }
2906 };
2907 write!(self.out, "{fun_name}(")?;
2908 self.write_expr(expr, ctx)?;
2909 write!(self.out, ")")?
2910 }
2911 // `Relational` is a normal function call to some glsl provided functions
2912 Expression::Relational { fun, argument } => {
2913 use crate::RelationalFunction as Rf;
2914
2915 let fun_name = match fun {
2916 // There's no specific function for this but we can invert the result of `isinf`
2917 Rf::IsFinite => "!isinf",
2918 Rf::IsInf => "isinf",
2919 Rf::IsNan => "isnan",
2920 // There's also no function for this but we can invert `isnan`
2921 Rf::IsNormal => "!isnan",
2922 Rf::All => "all",
2923 Rf::Any => "any",
2924 };
2925 write!(self.out, "{fun_name}(")?;
2926
2927 self.write_expr(argument, ctx)?;
2928
2929 write!(self.out, ")")?
2930 }
2931 Expression::Math {
2932 fun,
2933 arg,
2934 arg1,
2935 arg2,
2936 arg3,
2937 } => {
2938 use crate::MathFunction as Mf;
2939
2940 let fun_name = match fun {
2941 // comparison
2942 Mf::Abs => "abs",
2943 Mf::Min => "min",
2944 Mf::Max => "max",
2945 Mf::Clamp => "clamp",
2946 Mf::Saturate => {
2947 write!(self.out, "clamp(")?;
2948
2949 self.write_expr(arg, ctx)?;
2950
2951 match *ctx.info[arg].ty.inner_with(&self.module.types) {
2952 crate::TypeInner::Vector { size, .. } => write!(
2953 self.out,
2954 ", vec{}(0.0), vec{0}(1.0)",
2955 back::vector_size_str(size)
2956 )?,
2957 _ => write!(self.out, ", 0.0, 1.0")?,
2958 }
2959
2960 write!(self.out, ")")?;
2961
2962 return Ok(());
2963 }
2964 // trigonometry
2965 Mf::Cos => "cos",
2966 Mf::Cosh => "cosh",
2967 Mf::Sin => "sin",
2968 Mf::Sinh => "sinh",
2969 Mf::Tan => "tan",
2970 Mf::Tanh => "tanh",
2971 Mf::Acos => "acos",
2972 Mf::Asin => "asin",
2973 Mf::Atan => "atan",
2974 Mf::Asinh => "asinh",
2975 Mf::Acosh => "acosh",
2976 Mf::Atanh => "atanh",
2977 Mf::Radians => "radians",
2978 Mf::Degrees => "degrees",
2979 // glsl doesn't have atan2 function
2980 // use two-argument variation of the atan function
2981 Mf::Atan2 => "atan",
2982 // decomposition
2983 Mf::Ceil => "ceil",
2984 Mf::Floor => "floor",
2985 Mf::Round => "roundEven",
2986 Mf::Fract => "fract",
2987 Mf::Trunc => "trunc",
2988 Mf::Modf => "modf",
2989 Mf::Frexp => "frexp",
2990 Mf::Ldexp => "ldexp",
2991 // exponent
2992 Mf::Exp => "exp",
2993 Mf::Exp2 => "exp2",
2994 Mf::Log => "log",
2995 Mf::Log2 => "log2",
2996 Mf::Pow => "pow",
2997 // geometry
2998 Mf::Dot => match *ctx.info[arg].ty.inner_with(&self.module.types) {
2999 crate::TypeInner::Vector {
3000 kind: crate::ScalarKind::Float,
3001 ..
3002 } => "dot",
3003 crate::TypeInner::Vector { size, .. } => {
3004 return self.write_dot_product(arg, arg1.unwrap(), size as usize, ctx)
3005 }
3006 _ => unreachable!(
3007 "Correct TypeInner for dot product should be already validated"
3008 ),
3009 },
3010 Mf::Outer => "outerProduct",
3011 Mf::Cross => "cross",
3012 Mf::Distance => "distance",
3013 Mf::Length => "length",
3014 Mf::Normalize => "normalize",
3015 Mf::FaceForward => "faceforward",
3016 Mf::Reflect => "reflect",
3017 Mf::Refract => "refract",
3018 // computational
3019 Mf::Sign => "sign",
3020 Mf::Fma => {
3021 if self.options.version.supports_fma_function() {
3022 // Use the fma function when available
3023 "fma"
3024 } else {
3025 // No fma support. Transform the function call into an arithmetic expression
3026 write!(self.out, "(")?;
3027
3028 self.write_expr(arg, ctx)?;
3029 write!(self.out, " * ")?;
3030
3031 let arg1 =
3032 arg1.ok_or_else(|| Error::Custom("Missing fma arg1".to_owned()))?;
3033 self.write_expr(arg1, ctx)?;
3034 write!(self.out, " + ")?;
3035
3036 let arg2 =
3037 arg2.ok_or_else(|| Error::Custom("Missing fma arg2".to_owned()))?;
3038 self.write_expr(arg2, ctx)?;
3039 write!(self.out, ")")?;
3040
3041 return Ok(());
3042 }
3043 }
3044 Mf::Mix => "mix",
3045 Mf::Step => "step",
3046 Mf::SmoothStep => "smoothstep",
3047 Mf::Sqrt => "sqrt",
3048 Mf::InverseSqrt => "inversesqrt",
3049 Mf::Inverse => "inverse",
3050 Mf::Transpose => "transpose",
3051 Mf::Determinant => "determinant",
3052 // bits
3053 Mf::CountTrailingZeros => {
3054 match *ctx.info[arg].ty.inner_with(&self.module.types) {
3055 crate::TypeInner::Vector { size, kind, .. } => {
3056 let s = back::vector_size_str(size);
3057 if let crate::ScalarKind::Uint = kind {
3058 write!(self.out, "min(uvec{s}(findLSB(")?;
3059 self.write_expr(arg, ctx)?;
3060 write!(self.out, ")), uvec{s}(32u))")?;
3061 } else {
3062 write!(self.out, "ivec{s}(min(uvec{s}(findLSB(")?;
3063 self.write_expr(arg, ctx)?;
3064 write!(self.out, ")), uvec{s}(32u)))")?;
3065 }
3066 }
3067 crate::TypeInner::Scalar { kind, .. } => {
3068 if let crate::ScalarKind::Uint = kind {
3069 write!(self.out, "min(uint(findLSB(")?;
3070 self.write_expr(arg, ctx)?;
3071 write!(self.out, ")), 32u)")?;
3072 } else {
3073 write!(self.out, "int(min(uint(findLSB(")?;
3074 self.write_expr(arg, ctx)?;
3075 write!(self.out, ")), 32u))")?;
3076 }
3077 }
3078 _ => unreachable!(),
3079 };
3080 return Ok(());
3081 }
3082 Mf::CountLeadingZeros => {
3083 if self.options.version.supports_integer_functions() {
3084 match *ctx.info[arg].ty.inner_with(&self.module.types) {
3085 crate::TypeInner::Vector { size, kind, .. } => {
3086 let s = back::vector_size_str(size);
3087
3088 if let crate::ScalarKind::Uint = kind {
3089 write!(self.out, "uvec{s}(ivec{s}(31) - findMSB(")?;
3090 self.write_expr(arg, ctx)?;
3091 write!(self.out, "))")?;
3092 } else {
3093 write!(self.out, "mix(ivec{s}(31) - findMSB(")?;
3094 self.write_expr(arg, ctx)?;
3095 write!(self.out, "), ivec{s}(0), lessThan(")?;
3096 self.write_expr(arg, ctx)?;
3097 write!(self.out, ", ivec{s}(0)))")?;
3098 }
3099 }
3100 crate::TypeInner::Scalar { kind, .. } => {
3101 if let crate::ScalarKind::Uint = kind {
3102 write!(self.out, "uint(31 - findMSB(")?;
3103 } else {
3104 write!(self.out, "(")?;
3105 self.write_expr(arg, ctx)?;
3106 write!(self.out, " < 0 ? 0 : 31 - findMSB(")?;
3107 }
3108
3109 self.write_expr(arg, ctx)?;
3110 write!(self.out, "))")?;
3111 }
3112 _ => unreachable!(),
3113 };
3114 } else {
3115 match *ctx.info[arg].ty.inner_with(&self.module.types) {
3116 crate::TypeInner::Vector { size, kind, .. } => {
3117 let s = back::vector_size_str(size);
3118
3119 if let crate::ScalarKind::Uint = kind {
3120 write!(self.out, "uvec{s}(")?;
3121 write!(self.out, "vec{s}(31.0) - floor(log2(vec{s}(")?;
3122 self.write_expr(arg, ctx)?;
3123 write!(self.out, ") + 0.5)))")?;
3124 } else {
3125 write!(self.out, "ivec{s}(")?;
3126 write!(self.out, "mix(vec{s}(31.0) - floor(log2(vec{s}(")?;
3127 self.write_expr(arg, ctx)?;
3128 write!(self.out, ") + 0.5)), ")?;
3129 write!(self.out, "vec{s}(0.0), lessThan(")?;
3130 self.write_expr(arg, ctx)?;
3131 write!(self.out, ", ivec{s}(0u))))")?;
3132 }
3133 }
3134 crate::TypeInner::Scalar { kind, .. } => {
3135 if let crate::ScalarKind::Uint = kind {
3136 write!(self.out, "uint(31.0 - floor(log2(float(")?;
3137 self.write_expr(arg, ctx)?;
3138 write!(self.out, ") + 0.5)))")?;
3139 } else {
3140 write!(self.out, "(")?;
3141 self.write_expr(arg, ctx)?;
3142 write!(self.out, " < 0 ? 0 : int(")?;
3143 write!(self.out, "31.0 - floor(log2(float(")?;
3144 self.write_expr(arg, ctx)?;
3145 write!(self.out, ") + 0.5))))")?;
3146 }
3147 }
3148 _ => unreachable!(),
3149 };
3150 }
3151
3152 return Ok(());
3153 }
3154 Mf::CountOneBits => "bitCount",
3155 Mf::ReverseBits => "bitfieldReverse",
3156 Mf::ExtractBits => "bitfieldExtract",
3157 Mf::InsertBits => "bitfieldInsert",
3158 Mf::FindLsb => "findLSB",
3159 Mf::FindMsb => "findMSB",
3160 // data packing
3161 Mf::Pack4x8snorm => "packSnorm4x8",
3162 Mf::Pack4x8unorm => "packUnorm4x8",
3163 Mf::Pack2x16snorm => "packSnorm2x16",
3164 Mf::Pack2x16unorm => "packUnorm2x16",
3165 Mf::Pack2x16float => "packHalf2x16",
3166 // data unpacking
3167 Mf::Unpack4x8snorm => "unpackSnorm4x8",
3168 Mf::Unpack4x8unorm => "unpackUnorm4x8",
3169 Mf::Unpack2x16snorm => "unpackSnorm2x16",
3170 Mf::Unpack2x16unorm => "unpackUnorm2x16",
3171 Mf::Unpack2x16float => "unpackHalf2x16",
3172 };
3173
3174 let extract_bits = fun == Mf::ExtractBits;
3175 let insert_bits = fun == Mf::InsertBits;
3176
3177 // Some GLSL functions always return signed integers (like findMSB),
3178 // so they need to be cast to uint if the argument is also an uint.
3179 let ret_might_need_int_to_uint =
3180 matches!(fun, Mf::FindLsb | Mf::FindMsb | Mf::CountOneBits | Mf::Abs);
3181
3182 // Some GLSL functions only accept signed integers (like abs),
3183 // so they need their argument cast from uint to int.
3184 let arg_might_need_uint_to_int = matches!(fun, Mf::Abs);
3185
3186 // Check if the argument is an unsigned integer and return the vector size
3187 // in case it's a vector
3188 let maybe_uint_size = match *ctx.info[arg].ty.inner_with(&self.module.types) {
3189 crate::TypeInner::Scalar {
3190 kind: crate::ScalarKind::Uint,
3191 ..
3192 } => Some(None),
3193 crate::TypeInner::Vector {
3194 kind: crate::ScalarKind::Uint,
3195 size,
3196 ..
3197 } => Some(Some(size)),
3198 _ => None,
3199 };
3200
3201 // Cast to uint if the function needs it
3202 if ret_might_need_int_to_uint {
3203 if let Some(maybe_size) = maybe_uint_size {
3204 match maybe_size {
3205 Some(size) => write!(self.out, "uvec{}(", size as u8)?,
3206 None => write!(self.out, "uint(")?,
3207 }
3208 }
3209 }
3210
3211 write!(self.out, "{fun_name}(")?;
3212
3213 // Cast to int if the function needs it
3214 if arg_might_need_uint_to_int {
3215 if let Some(maybe_size) = maybe_uint_size {
3216 match maybe_size {
3217 Some(size) => write!(self.out, "ivec{}(", size as u8)?,
3218 None => write!(self.out, "int(")?,
3219 }
3220 }
3221 }
3222
3223 self.write_expr(arg, ctx)?;
3224
3225 // Close the cast from uint to int
3226 if arg_might_need_uint_to_int && maybe_uint_size.is_some() {
3227 write!(self.out, ")")?
3228 }
3229
3230 if let Some(arg) = arg1 {
3231 write!(self.out, ", ")?;
3232 if extract_bits {
3233 write!(self.out, "int(")?;
3234 self.write_expr(arg, ctx)?;
3235 write!(self.out, ")")?;
3236 } else {
3237 self.write_expr(arg, ctx)?;
3238 }
3239 }
3240 if let Some(arg) = arg2 {
3241 write!(self.out, ", ")?;
3242 if extract_bits || insert_bits {
3243 write!(self.out, "int(")?;
3244 self.write_expr(arg, ctx)?;
3245 write!(self.out, ")")?;
3246 } else {
3247 self.write_expr(arg, ctx)?;
3248 }
3249 }
3250 if let Some(arg) = arg3 {
3251 write!(self.out, ", ")?;
3252 if insert_bits {
3253 write!(self.out, "int(")?;
3254 self.write_expr(arg, ctx)?;
3255 write!(self.out, ")")?;
3256 } else {
3257 self.write_expr(arg, ctx)?;
3258 }
3259 }
3260 write!(self.out, ")")?;
3261
3262 // Close the cast from int to uint
3263 if ret_might_need_int_to_uint && maybe_uint_size.is_some() {
3264 write!(self.out, ")")?
3265 }
3266 }
3267 // `As` is always a call.
3268 // If `convert` is true the function name is the type
3269 // Else the function name is one of the glsl provided bitcast functions
3270 Expression::As {
3271 expr,
3272 kind: target_kind,
3273 convert,
3274 } => {
3275 let inner = ctx.info[expr].ty.inner_with(&self.module.types);
3276 match convert {
3277 Some(width) => {
3278 // this is similar to `write_type`, but with the target kind
3279 let scalar = glsl_scalar(target_kind, width)?;
3280 match *inner {
3281 TypeInner::Matrix { columns, rows, .. } => write!(
3282 self.out,
3283 "{}mat{}x{}",
3284 scalar.prefix, columns as u8, rows as u8
3285 )?,
3286 TypeInner::Vector { size, .. } => {
3287 write!(self.out, "{}vec{}", scalar.prefix, size as u8)?
3288 }
3289 _ => write!(self.out, "{}", scalar.full)?,
3290 }
3291
3292 write!(self.out, "(")?;
3293 self.write_expr(expr, ctx)?;
3294 write!(self.out, ")")?
3295 }
3296 None => {
3297 use crate::ScalarKind as Sk;
3298
3299 let target_vector_type = match *inner {
3300 TypeInner::Vector { size, width, .. } => Some(TypeInner::Vector {
3301 size,
3302 width,
3303 kind: target_kind,
3304 }),
3305 _ => None,
3306 };
3307
3308 let source_kind = inner.scalar_kind().unwrap();
3309
3310 match (source_kind, target_kind, target_vector_type) {
3311 // No conversion needed
3312 (Sk::Sint, Sk::Sint, _)
3313 | (Sk::Uint, Sk::Uint, _)
3314 | (Sk::Float, Sk::Float, _)
3315 | (Sk::Bool, Sk::Bool, _) => {
3316 self.write_expr(expr, ctx)?;
3317 return Ok(());
3318 }
3319
3320 // Cast to/from floats
3321 (Sk::Float, Sk::Sint, _) => write!(self.out, "floatBitsToInt")?,
3322 (Sk::Float, Sk::Uint, _) => write!(self.out, "floatBitsToUint")?,
3323 (Sk::Sint, Sk::Float, _) => write!(self.out, "intBitsToFloat")?,
3324 (Sk::Uint, Sk::Float, _) => write!(self.out, "uintBitsToFloat")?,
3325
3326 // Cast between vector types
3327 (_, _, Some(vector)) => {
3328 self.write_value_type(&vector)?;
3329 }
3330
3331 // There is no way to bitcast between Uint/Sint in glsl. Use constructor conversion
3332 (Sk::Uint | Sk::Bool, Sk::Sint, None) => write!(self.out, "int")?,
3333 (Sk::Sint | Sk::Bool, Sk::Uint, None) => write!(self.out, "uint")?,
3334 (Sk::Bool, Sk::Float, None) => write!(self.out, "float")?,
3335 (Sk::Sint | Sk::Uint | Sk::Float, Sk::Bool, None) => {
3336 write!(self.out, "bool")?
3337 }
3338 };
3339
3340 write!(self.out, "(")?;
3341 self.write_expr(expr, ctx)?;
3342 write!(self.out, ")")?;
3343 }
3344 }
3345 }
3346 // These expressions never show up in `Emit`.
3347 Expression::CallResult(_)
3348 | Expression::AtomicResult { .. }
3349 | Expression::RayQueryProceedResult
3350 | Expression::WorkGroupUniformLoadResult { .. } => unreachable!(),
3351 // `ArrayLength` is written as `expr.length()` and we convert it to a uint
3352 Expression::ArrayLength(expr) => {
3353 write!(self.out, "uint(")?;
3354 self.write_expr(expr, ctx)?;
3355 write!(self.out, ".length())")?
3356 }
3357 // not supported yet
3358 Expression::RayQueryGetIntersection { .. } => unreachable!(),
3359 }
3360
3361 Ok(())
3362 }
3363
3364 /// Helper function to write the local holding the clamped lod
3365 fn write_clamped_lod(
3366 &mut self,
3367 ctx: &back::FunctionCtx,
3368 expr: Handle<crate::Expression>,
3369 image: Handle<crate::Expression>,
3370 level_expr: Handle<crate::Expression>,
3371 ) -> Result<(), Error> {
3372 // Define our local and start a call to `clamp`
3373 write!(
3374 self.out,
3375 "int {}{}{} = clamp(",
3376 back::BAKE_PREFIX,
3377 expr.index(),
3378 CLAMPED_LOD_SUFFIX
3379 )?;
3380 // Write the lod that will be clamped
3381 self.write_expr(level_expr, ctx)?;
3382 // Set the min value to 0 and start a call to `textureQueryLevels` to get
3383 // the maximum value
3384 write!(self.out, ", 0, textureQueryLevels(")?;
3385 // Write the target image as an argument to `textureQueryLevels`
3386 self.write_expr(image, ctx)?;
3387 // Close the call to `textureQueryLevels` subtract 1 from it since
3388 // the lod argument is 0 based, close the `clamp` call and end the
3389 // local declaration statement.
3390 writeln!(self.out, ") - 1);")?;
3391
3392 Ok(())
3393 }
3394
3395 // Helper method used to retrieve how many elements a coordinate vector
3396 // for the images operations need.
3397 fn get_coordinate_vector_size(&self, dim: crate::ImageDimension, arrayed: bool) -> u8 {
3398 // openGL es doesn't have 1D images so we need workaround it
3399 let tex_1d_hack = dim == crate::ImageDimension::D1 && self.options.version.is_es();
3400 // Get how many components the coordinate vector needs for the dimensions only
3401 let tex_coord_size = match dim {
3402 crate::ImageDimension::D1 => 1,
3403 crate::ImageDimension::D2 => 2,
3404 crate::ImageDimension::D3 => 3,
3405 crate::ImageDimension::Cube => 2,
3406 };
3407 // Calculate the true size of the coordinate vector by adding 1 for arrayed images
3408 // and another 1 if we need to workaround 1D images by making them 2D
3409 tex_coord_size + tex_1d_hack as u8 + arrayed as u8
3410 }
3411
3412 /// Helper method to write the coordinate vector for image operations
3413 fn write_texture_coord(
3414 &mut self,
3415 ctx: &back::FunctionCtx,
3416 vector_size: u8,
3417 coordinate: Handle<crate::Expression>,
3418 array_index: Option<Handle<crate::Expression>>,
3419 // Emulate 1D images as 2D for profiles that don't support it (glsl es)
3420 tex_1d_hack: bool,
3421 ) -> Result<(), Error> {
3422 match array_index {
3423 // If the image needs an array indice we need to add it to the end of our
3424 // coordinate vector, to do so we will use the `ivec(ivec, scalar)`
3425 // constructor notation (NOTE: the inner `ivec` can also be a scalar, this
3426 // is important for 1D arrayed images).
3427 Some(layer_expr) => {
3428 write!(self.out, "ivec{vector_size}(")?;
3429 self.write_expr(coordinate, ctx)?;
3430 write!(self.out, ", ")?;
3431 // If we are replacing sampler1D with sampler2D we also need
3432 // to add another zero to the coordinates vector for the y component
3433 if tex_1d_hack {
3434 write!(self.out, "0, ")?;
3435 }
3436 self.write_expr(layer_expr, ctx)?;
3437 write!(self.out, ")")?;
3438 }
3439 // Otherwise write just the expression (and the 1D hack if needed)
3440 None => {
3441 let uvec_size = match *ctx.info[coordinate].ty.inner_with(&self.module.types) {
3442 TypeInner::Scalar {
3443 kind: crate::ScalarKind::Uint,
3444 ..
3445 } => Some(None),
3446 TypeInner::Vector {
3447 size,
3448 kind: crate::ScalarKind::Uint,
3449 ..
3450 } => Some(Some(size as u32)),
3451 _ => None,
3452 };
3453 if tex_1d_hack {
3454 write!(self.out, "ivec2(")?;
3455 } else if uvec_size.is_some() {
3456 match uvec_size {
3457 Some(None) => write!(self.out, "int(")?,
3458 Some(Some(size)) => write!(self.out, "ivec{size}(")?,
3459 _ => {}
3460 }
3461 }
3462 self.write_expr(coordinate, ctx)?;
3463 if tex_1d_hack {
3464 write!(self.out, ", 0)")?;
3465 } else if uvec_size.is_some() {
3466 write!(self.out, ")")?;
3467 }
3468 }
3469 }
3470
3471 Ok(())
3472 }
3473
3474 /// Helper method to write the `ImageStore` statement
3475 fn write_image_store(
3476 &mut self,
3477 ctx: &back::FunctionCtx,
3478 image: Handle<crate::Expression>,
3479 coordinate: Handle<crate::Expression>,
3480 array_index: Option<Handle<crate::Expression>>,
3481 value: Handle<crate::Expression>,
3482 ) -> Result<(), Error> {
3483 use crate::ImageDimension as IDim;
3484
3485 // NOTE: openGL requires that `imageStore`s have no effets when the texel is invalid
3486 // so we don't need to generate bounds checks (OpenGL 4.2 Core §3.9.20)
3487
3488 // This will only panic if the module is invalid
3489 let dim = match *ctx.info[image].ty.inner_with(&self.module.types) {
3490 TypeInner::Image { dim, .. } => dim,
3491 _ => unreachable!(),
3492 };
3493
3494 // Begin our call to `imageStore`
3495 write!(self.out, "imageStore(")?;
3496 self.write_expr(image, ctx)?;
3497 // Separate the image argument from the coordinates
3498 write!(self.out, ", ")?;
3499
3500 // openGL es doesn't have 1D images so we need workaround it
3501 let tex_1d_hack = dim == IDim::D1 && self.options.version.is_es();
3502 // Write the coordinate vector
3503 self.write_texture_coord(
3504 ctx,
3505 // Get the size of the coordinate vector
3506 self.get_coordinate_vector_size(dim, array_index.is_some()),
3507 coordinate,
3508 array_index,
3509 tex_1d_hack,
3510 )?;
3511
3512 // Separate the coordinate from the value to write and write the expression
3513 // of the value to write.
3514 write!(self.out, ", ")?;
3515 self.write_expr(value, ctx)?;
3516 // End the call to `imageStore` and the statement.
3517 writeln!(self.out, ");")?;
3518
3519 Ok(())
3520 }
3521
3522 /// Helper method for writing an `ImageLoad` expression.
3523 #[allow(clippy::too_many_arguments)]
3524 fn write_image_load(
3525 &mut self,
3526 handle: Handle<crate::Expression>,
3527 ctx: &back::FunctionCtx,
3528 image: Handle<crate::Expression>,
3529 coordinate: Handle<crate::Expression>,
3530 array_index: Option<Handle<crate::Expression>>,
3531 sample: Option<Handle<crate::Expression>>,
3532 level: Option<Handle<crate::Expression>>,
3533 ) -> Result<(), Error> {
3534 use crate::ImageDimension as IDim;
3535
3536 // `ImageLoad` is a bit complicated.
3537 // There are two functions one for sampled
3538 // images another for storage images, the former uses `texelFetch` and the
3539 // latter uses `imageLoad`.
3540 //
3541 // Furthermore we have `level` which is always `Some` for sampled images
3542 // and `None` for storage images, so we end up with two functions:
3543 // - `texelFetch(image, coordinate, level)` for sampled images
3544 // - `imageLoad(image, coordinate)` for storage images
3545 //
3546 // Finally we also have to consider bounds checking, for storage images
3547 // this is easy since openGL requires that invalid texels always return
3548 // 0, for sampled images we need to either verify that all arguments are
3549 // in bounds (`ReadZeroSkipWrite`) or make them a valid texel (`Restrict`).
3550
3551 // This will only panic if the module is invalid
3552 let (dim, class) = match *ctx.info[image].ty.inner_with(&self.module.types) {
3553 TypeInner::Image {
3554 dim,
3555 arrayed: _,
3556 class,
3557 } => (dim, class),
3558 _ => unreachable!(),
3559 };
3560
3561 // Get the name of the function to be used for the load operation
3562 // and the policy to be used with it.
3563 let (fun_name, policy) = match class {
3564 // Sampled images inherit the policy from the user passed policies
3565 crate::ImageClass::Sampled { .. } => ("texelFetch", self.policies.image_load),
3566 crate::ImageClass::Storage { .. } => {
3567 // OpenGL ES 3.1 mentiones in Chapter "8.22 Texture Image Loads and Stores" that:
3568 // "Invalid image loads will return a vector where the value of R, G, and B components
3569 // is 0 and the value of the A component is undefined."
3570 //
3571 // OpenGL 4.2 Core mentiones in Chapter "3.9.20 Texture Image Loads and Stores" that:
3572 // "Invalid image loads will return zero."
3573 //
3574 // So, we only inject bounds checks for ES
3575 let policy = if self.options.version.is_es() {
3576 self.policies.image_load
3577 } else {
3578 proc::BoundsCheckPolicy::Unchecked
3579 };
3580 ("imageLoad", policy)
3581 }
3582 // TODO: Is there even a function for this?
3583 crate::ImageClass::Depth { multi: _ } => {
3584 return Err(Error::Custom(
3585 "WGSL `textureLoad` from depth textures is not supported in GLSL".to_string(),
3586 ))
3587 }
3588 };
3589
3590 // openGL es doesn't have 1D images so we need workaround it
3591 let tex_1d_hack = dim == IDim::D1 && self.options.version.is_es();
3592 // Get the size of the coordinate vector
3593 let vector_size = self.get_coordinate_vector_size(dim, array_index.is_some());
3594
3595 if let proc::BoundsCheckPolicy::ReadZeroSkipWrite = policy {
3596 // To write the bounds checks for `ReadZeroSkipWrite` we will use a
3597 // ternary operator since we are in the middle of an expression and
3598 // need to return a value.
3599 //
3600 // NOTE: glsl does short circuit when evaluating logical
3601 // expressions so we can be sure that after we test a
3602 // condition it will be true for the next ones
3603
3604 // Write parantheses around the ternary operator to prevent problems with
3605 // expressions emitted before or after it having more precedence
3606 write!(self.out, "(",)?;
3607
3608 // The lod check needs to precede the size check since we need
3609 // to use the lod to get the size of the image at that level.
3610 if let Some(level_expr) = level {
3611 self.write_expr(level_expr, ctx)?;
3612 write!(self.out, " < textureQueryLevels(",)?;
3613 self.write_expr(image, ctx)?;
3614 // Chain the next check
3615 write!(self.out, ") && ")?;
3616 }
3617
3618 // Check that the sample arguments doesn't exceed the number of samples
3619 if let Some(sample_expr) = sample {
3620 self.write_expr(sample_expr, ctx)?;
3621 write!(self.out, " < textureSamples(",)?;
3622 self.write_expr(image, ctx)?;
3623 // Chain the next check
3624 write!(self.out, ") && ")?;
3625 }
3626
3627 // We now need to write the size checks for the coordinates and array index
3628 // first we write the comparation function in case the image is 1D non arrayed
3629 // (and no 1D to 2D hack was needed) we are comparing scalars so the less than
3630 // operator will suffice, but otherwise we'll be comparing two vectors so we'll
3631 // need to use the `lessThan` function but it returns a vector of booleans (one
3632 // for each comparison) so we need to fold it all in one scalar boolean, since
3633 // we want all comparisons to pass we use the `all` function which will only
3634 // return `true` if all the elements of the boolean vector are also `true`.
3635 //
3636 // So we'll end with one of the following forms
3637 // - `coord < textureSize(image, lod)` for 1D images
3638 // - `all(lessThan(coord, textureSize(image, lod)))` for normal images
3639 // - `all(lessThan(ivec(coord, array_index), textureSize(image, lod)))`
3640 // for arrayed images
3641 // - `all(lessThan(coord, textureSize(image)))` for multi sampled images
3642
3643 if vector_size != 1 {
3644 write!(self.out, "all(lessThan(")?;
3645 }
3646
3647 // Write the coordinate vector
3648 self.write_texture_coord(ctx, vector_size, coordinate, array_index, tex_1d_hack)?;
3649
3650 if vector_size != 1 {
3651 // If we used the `lessThan` function we need to separate the
3652 // coordinates from the image size.
3653 write!(self.out, ", ")?;
3654 } else {
3655 // If we didn't use it (ie. 1D images) we perform the comparsion
3656 // using the less than operator.
3657 write!(self.out, " < ")?;
3658 }
3659
3660 // Call `textureSize` to get our image size
3661 write!(self.out, "textureSize(")?;
3662 self.write_expr(image, ctx)?;
3663 // `textureSize` uses the lod as a second argument for mipmapped images
3664 if let Some(level_expr) = level {
3665 // Separate the image from the lod
3666 write!(self.out, ", ")?;
3667 self.write_expr(level_expr, ctx)?;
3668 }
3669 // Close the `textureSize` call
3670 write!(self.out, ")")?;
3671
3672 if vector_size != 1 {
3673 // Close the `all` and `lessThan` calls
3674 write!(self.out, "))")?;
3675 }
3676
3677 // Finally end the condition part of the ternary operator
3678 write!(self.out, " ? ")?;
3679 }
3680
3681 // Begin the call to the function used to load the texel
3682 write!(self.out, "{fun_name}(")?;
3683 self.write_expr(image, ctx)?;
3684 write!(self.out, ", ")?;
3685
3686 // If we are using `Restrict` bounds checking we need to pass valid texel
3687 // coordinates, to do so we use the `clamp` function to get a value between
3688 // 0 and the image size - 1 (indexing begins at 0)
3689 if let proc::BoundsCheckPolicy::Restrict = policy {
3690 write!(self.out, "clamp(")?;
3691 }
3692
3693 // Write the coordinate vector
3694 self.write_texture_coord(ctx, vector_size, coordinate, array_index, tex_1d_hack)?;
3695
3696 // If we are using `Restrict` bounds checking we need to write the rest of the
3697 // clamp we initiated before writing the coordinates.
3698 if let proc::BoundsCheckPolicy::Restrict = policy {
3699 // Write the min value 0
3700 if vector_size == 1 {
3701 write!(self.out, ", 0")?;
3702 } else {
3703 write!(self.out, ", ivec{vector_size}(0)")?;
3704 }
3705 // Start the `textureSize` call to use as the max value.
3706 write!(self.out, ", textureSize(")?;
3707 self.write_expr(image, ctx)?;
3708 // If the image is mipmapped we need to add the lod argument to the
3709 // `textureSize` call, but this needs to be the clamped lod, this should
3710 // have been generated earlier and put in a local.
3711 if class.is_mipmapped() {
3712 write!(
3713 self.out,
3714 ", {}{}{}",
3715 back::BAKE_PREFIX,
3716 handle.index(),
3717 CLAMPED_LOD_SUFFIX
3718 )?;
3719 }
3720 // Close the `textureSize` call
3721 write!(self.out, ")")?;
3722
3723 // Subtract 1 from the `textureSize` call since the coordinates are zero based.
3724 if vector_size == 1 {
3725 write!(self.out, " - 1")?;
3726 } else {
3727 write!(self.out, " - ivec{vector_size}(1)")?;
3728 }
3729
3730 // Close the `clamp` call
3731 write!(self.out, ")")?;
3732
3733 // Add the clamped lod (if present) as the second argument to the
3734 // image load function.
3735 if level.is_some() {
3736 write!(
3737 self.out,
3738 ", {}{}{}",
3739 back::BAKE_PREFIX,
3740 handle.index(),
3741 CLAMPED_LOD_SUFFIX
3742 )?;
3743 }
3744
3745 // If a sample argument is needed we need to clamp it between 0 and
3746 // the number of samples the image has.
3747 if let Some(sample_expr) = sample {
3748 write!(self.out, ", clamp(")?;
3749 self.write_expr(sample_expr, ctx)?;
3750 // Set the min value to 0 and start the call to `textureSamples`
3751 write!(self.out, ", 0, textureSamples(")?;
3752 self.write_expr(image, ctx)?;
3753 // Close the `textureSamples` call, subtract 1 from it since the sample
3754 // argument is zero based, and close the `clamp` call
3755 writeln!(self.out, ") - 1)")?;
3756 }
3757 } else if let Some(sample_or_level) = sample.or(level) {
3758 // If no bounds checking is need just add the sample or level argument
3759 // after the coordinates
3760 write!(self.out, ", ")?;
3761 self.write_expr(sample_or_level, ctx)?;
3762 }
3763
3764 // Close the image load function.
3765 write!(self.out, ")")?;
3766
3767 // If we were using the `ReadZeroSkipWrite` policy we need to end the first branch
3768 // (which is taken if the condition is `true`) with a colon (`:`) and write the
3769 // second branch which is just a 0 value.
3770 if let proc::BoundsCheckPolicy::ReadZeroSkipWrite = policy {
3771 // Get the kind of the output value.
3772 let kind = match class {
3773 // Only sampled images can reach here since storage images
3774 // don't need bounds checks and depth images aren't implmented
3775 crate::ImageClass::Sampled { kind, .. } => kind,
3776 _ => unreachable!(),
3777 };
3778
3779 // End the first branch
3780 write!(self.out, " : ")?;
3781 // Write the 0 value
3782 write!(self.out, "{}vec4(", glsl_scalar(kind, 4)?.prefix,)?;
3783 self.write_zero_init_scalar(kind)?;
3784 // Close the zero value constructor
3785 write!(self.out, ")")?;
3786 // Close the parantheses surrounding our ternary
3787 write!(self.out, ")")?;
3788 }
3789
3790 Ok(())
3791 }
3792
3793 fn write_named_expr(
3794 &mut self,
3795 handle: Handle<crate::Expression>,
3796 name: String,
3797 // The expression which is being named.
3798 // Generally, this is the same as handle, except in WorkGroupUniformLoad
3799 named: Handle<crate::Expression>,
3800 ctx: &back::FunctionCtx,
3801 ) -> BackendResult {
3802 match ctx.info[named].ty {
3803 proc::TypeResolution::Handle(ty_handle) => match self.module.types[ty_handle].inner {
3804 TypeInner::Struct { .. } => {
3805 let ty_name = &self.names[&NameKey::Type(ty_handle)];
3806 write!(self.out, "{ty_name}")?;
3807 }
3808 _ => {
3809 self.write_type(ty_handle)?;
3810 }
3811 },
3812 proc::TypeResolution::Value(ref inner) => {
3813 self.write_value_type(inner)?;
3814 }
3815 }
3816
3817 let base_ty_res = &ctx.info[named].ty;
3818 let resolved = base_ty_res.inner_with(&self.module.types);
3819
3820 write!(self.out, " {name}")?;
3821 if let TypeInner::Array { base, size, .. } = *resolved {
3822 self.write_array_size(base, size)?;
3823 }
3824 write!(self.out, " = ")?;
3825 self.write_expr(handle, ctx)?;
3826 writeln!(self.out, ";")?;
3827 self.named_expressions.insert(named, name);
3828
3829 Ok(())
3830 }
3831
3832 /// Helper function that write string with default zero initialization for supported types
3833 fn write_zero_init_value(&mut self, ty: Handle<crate::Type>) -> BackendResult {
3834 let inner = &self.module.types[ty].inner;
3835 match *inner {
3836 TypeInner::Scalar { kind, .. } | TypeInner::Atomic { kind, .. } => {
3837 self.write_zero_init_scalar(kind)?;
3838 }
3839 TypeInner::Vector { kind, .. } => {
3840 self.write_value_type(inner)?;
3841 write!(self.out, "(")?;
3842 self.write_zero_init_scalar(kind)?;
3843 write!(self.out, ")")?;
3844 }
3845 TypeInner::Matrix { .. } => {
3846 self.write_value_type(inner)?;
3847 write!(self.out, "(")?;
3848 self.write_zero_init_scalar(crate::ScalarKind::Float)?;
3849 write!(self.out, ")")?;
3850 }
3851 TypeInner::Array { base, size, .. } => {
3852 let count = match size
3853 .to_indexable_length(self.module)
3854 .expect("Bad array size")
3855 {
3856 proc::IndexableLength::Known(count) => count,
3857 proc::IndexableLength::Dynamic => return Ok(()),
3858 };
3859 self.write_type(base)?;
3860 self.write_array_size(base, size)?;
3861 write!(self.out, "(")?;
3862 for _ in 1..count {
3863 self.write_zero_init_value(base)?;
3864 write!(self.out, ", ")?;
3865 }
3866 // write last parameter without comma and space
3867 self.write_zero_init_value(base)?;
3868 write!(self.out, ")")?;
3869 }
3870 TypeInner::Struct { ref members, .. } => {
3871 let name = &self.names[&NameKey::Type(ty)];
3872 write!(self.out, "{name}(")?;
3873 for (index, member) in members.iter().enumerate() {
3874 if index != 0 {
3875 write!(self.out, ", ")?;
3876 }
3877 self.write_zero_init_value(member.ty)?;
3878 }
3879 write!(self.out, ")")?;
3880 }
3881 _ => unreachable!(),
3882 }
3883
3884 Ok(())
3885 }
3886
3887 /// Helper function that write string with zero initialization for scalar
3888 fn write_zero_init_scalar(&mut self, kind: crate::ScalarKind) -> BackendResult {
3889 match kind {
3890 crate::ScalarKind::Bool => write!(self.out, "false")?,
3891 crate::ScalarKind::Uint => write!(self.out, "0u")?,
3892 crate::ScalarKind::Float => write!(self.out, "0.0")?,
3893 crate::ScalarKind::Sint => write!(self.out, "0")?,
3894 }
3895
3896 Ok(())
3897 }
3898
3899 /// Issue a memory barrier. Please note that to ensure visibility,
3900 /// OpenGL always requires a call to the `barrier()` function after a `memoryBarrier*()`
3901 fn write_barrier(&mut self, flags: crate::Barrier, level: back::Level) -> BackendResult {
3902 if flags.contains(crate::Barrier::STORAGE) {
3903 writeln!(self.out, "{level}memoryBarrierBuffer();")?;
3904 }
3905 if flags.contains(crate::Barrier::WORK_GROUP) {
3906 writeln!(self.out, "{level}memoryBarrierShared();")?;
3907 }
3908 writeln!(self.out, "{level}barrier();")?;
3909 Ok(())
3910 }
3911
3912 /// Helper function that return the glsl storage access string of [`StorageAccess`](crate::StorageAccess)
3913 ///
3914 /// glsl allows adding both `readonly` and `writeonly` but this means that
3915 /// they can only be used to query information about the resource which isn't what
3916 /// we want here so when storage access is both `LOAD` and `STORE` add no modifiers
3917 fn write_storage_access(&mut self, storage_access: crate::StorageAccess) -> BackendResult {
3918 if !storage_access.contains(crate::StorageAccess::STORE) {
3919 write!(self.out, "readonly ")?;
3920 }
3921 if !storage_access.contains(crate::StorageAccess::LOAD) {
3922 write!(self.out, "writeonly ")?;
3923 }
3924 Ok(())
3925 }
3926
3927 /// Helper method used to produce the reflection info that's returned to the user
3928 fn collect_reflection_info(&self) -> Result<ReflectionInfo, Error> {
3929 use std::collections::hash_map::Entry;
3930 let info = self.info.get_entry_point(self.entry_point_idx as usize);
3931 let mut texture_mapping = crate::FastHashMap::default();
3932 let mut uniforms = crate::FastHashMap::default();
3933
3934 for sampling in info.sampling_set.iter() {
3935 let tex_name = self.reflection_names_globals[&sampling.image].clone();
3936
3937 match texture_mapping.entry(tex_name) {
3938 Entry::Vacant(v) => {
3939 v.insert(TextureMapping {
3940 texture: sampling.image,
3941 sampler: Some(sampling.sampler),
3942 });
3943 }
3944 Entry::Occupied(e) => {
3945 if e.get().sampler != Some(sampling.sampler) {
3946 log::error!("Conflicting samplers for {}", e.key());
3947 return Err(Error::ImageMultipleSamplers);
3948 }
3949 }
3950 }
3951 }
3952
3953 for (handle, var) in self.module.global_variables.iter() {
3954 if info[handle].is_empty() {
3955 continue;
3956 }
3957 match self.module.types[var.ty].inner {
3958 crate::TypeInner::Image { .. } => {
3959 let tex_name = self.reflection_names_globals[&handle].clone();
3960 match texture_mapping.entry(tex_name) {
3961 Entry::Vacant(v) => {
3962 v.insert(TextureMapping {
3963 texture: handle,
3964 sampler: None,
3965 });
3966 }
3967 Entry::Occupied(_) => {
3968 // already used with a sampler, do nothing
3969 }
3970 }
3971 }
3972 _ => match var.space {
3973 crate::AddressSpace::Uniform | crate::AddressSpace::Storage { .. } => {
3974 let name = self.reflection_names_globals[&handle].clone();
3975 uniforms.insert(handle, name);
3976 }
3977 _ => (),
3978 },
3979 }
3980 }
3981
3982 Ok(ReflectionInfo {
3983 texture_mapping,
3984 uniforms,
3985 })
3986 }
3987}
3988
3989/// Structure returned by [`glsl_scalar`](glsl_scalar)
3990///
3991/// It contains both a prefix used in other types and the full type name
3992struct ScalarString<'a> {
3993 /// The prefix used to compose other types
3994 prefix: &'a str,
3995 /// The name of the scalar type
3996 full: &'a str,
3997}
3998
3999/// Helper function that returns scalar related strings
4000///
4001/// Check [`ScalarString`](ScalarString) for the information provided
4002///
4003/// # Errors
4004/// If a [`Float`](crate::ScalarKind::Float) with an width that isn't 4 or 8
4005const fn glsl_scalar(
4006 kind: crate::ScalarKind,
4007 width: crate::Bytes,
4008) -> Result<ScalarString<'static>, Error> {
4009 use crate::ScalarKind as Sk;
4010
4011 Ok(match kind {
4012 Sk::Sint => ScalarString {
4013 prefix: "i",
4014 full: "int",
4015 },
4016 Sk::Uint => ScalarString {
4017 prefix: "u",
4018 full: "uint",
4019 },
4020 Sk::Float => match width {
4021 4 => ScalarString {
4022 prefix: "",
4023 full: "float",
4024 },
4025 8 => ScalarString {
4026 prefix: "d",
4027 full: "double",
4028 },
4029 _ => return Err(Error::UnsupportedScalar(kind, width)),
4030 },
4031 Sk::Bool => ScalarString {
4032 prefix: "b",
4033 full: "bool",
4034 },
4035 })
4036}
4037
4038/// Helper function that returns the glsl variable name for a builtin
4039const fn glsl_built_in(
4040 built_in: crate::BuiltIn,
4041 output: bool,
4042 targetting_webgl: bool,
4043) -> &'static str {
4044 use crate::BuiltIn as Bi;
4045
4046 match built_in {
4047 Bi::Position { .. } => {
4048 if output {
4049 "gl_Position"
4050 } else {
4051 "gl_FragCoord"
4052 }
4053 }
4054 Bi::ViewIndex if targetting_webgl => "int(gl_ViewID_OVR)",
4055 Bi::ViewIndex => "gl_ViewIndex",
4056 // vertex
4057 Bi::BaseInstance => "uint(gl_BaseInstance)",
4058 Bi::BaseVertex => "uint(gl_BaseVertex)",
4059 Bi::ClipDistance => "gl_ClipDistance",
4060 Bi::CullDistance => "gl_CullDistance",
4061 Bi::InstanceIndex => "uint(gl_InstanceID)",
4062 Bi::PointSize => "gl_PointSize",
4063 Bi::VertexIndex => "uint(gl_VertexID)",
4064 // fragment
4065 Bi::FragDepth => "gl_FragDepth",
4066 Bi::PointCoord => "gl_PointCoord",
4067 Bi::FrontFacing => "gl_FrontFacing",
4068 Bi::PrimitiveIndex => "uint(gl_PrimitiveID)",
4069 Bi::SampleIndex => "gl_SampleID",
4070 Bi::SampleMask => {
4071 if output {
4072 "gl_SampleMask"
4073 } else {
4074 "gl_SampleMaskIn"
4075 }
4076 }
4077 // compute
4078 Bi::GlobalInvocationId => "gl_GlobalInvocationID",
4079 Bi::LocalInvocationId => "gl_LocalInvocationID",
4080 Bi::LocalInvocationIndex => "gl_LocalInvocationIndex",
4081 Bi::WorkGroupId => "gl_WorkGroupID",
4082 Bi::WorkGroupSize => "gl_WorkGroupSize",
4083 Bi::NumWorkGroups => "gl_NumWorkGroups",
4084 }
4085}
4086
4087/// Helper function that returns the string corresponding to the address space
4088const fn glsl_storage_qualifier(space: crate::AddressSpace) -> Option<&'static str> {
4089 use crate::AddressSpace as As;
4090
4091 match space {
4092 As::Function => None,
4093 As::Private => None,
4094 As::Storage { .. } => Some("buffer"),
4095 As::Uniform => Some("uniform"),
4096 As::Handle => Some("uniform"),
4097 As::WorkGroup => Some("shared"),
4098 As::PushConstant => Some("uniform"),
4099 }
4100}
4101
4102/// Helper function that returns the string corresponding to the glsl interpolation qualifier
4103const fn glsl_interpolation(interpolation: crate::Interpolation) -> &'static str {
4104 use crate::Interpolation as I;
4105
4106 match interpolation {
4107 I::Perspective => "smooth",
4108 I::Linear => "noperspective",
4109 I::Flat => "flat",
4110 }
4111}
4112
4113/// Return the GLSL auxiliary qualifier for the given sampling value.
4114const fn glsl_sampling(sampling: crate::Sampling) -> Option<&'static str> {
4115 use crate::Sampling as S;
4116
4117 match sampling {
4118 S::Center => None,
4119 S::Centroid => Some("centroid"),
4120 S::Sample => Some("sample"),
4121 }
4122}
4123
4124/// Helper function that returns the glsl dimension string of [`ImageDimension`](crate::ImageDimension)
4125const fn glsl_dimension(dim: crate::ImageDimension) -> &'static str {
4126 use crate::ImageDimension as IDim;
4127
4128 match dim {
4129 IDim::D1 => "1D",
4130 IDim::D2 => "2D",
4131 IDim::D3 => "3D",
4132 IDim::Cube => "Cube",
4133 }
4134}
4135
4136/// Helper function that returns the glsl storage format string of [`StorageFormat`](crate::StorageFormat)
4137const fn glsl_storage_format(format: crate::StorageFormat) -> &'static str {
4138 use crate::StorageFormat as Sf;
4139
4140 match format {
4141 Sf::R8Unorm => "r8",
4142 Sf::R8Snorm => "r8_snorm",
4143 Sf::R8Uint => "r8ui",
4144 Sf::R8Sint => "r8i",
4145 Sf::R16Uint => "r16ui",
4146 Sf::R16Sint => "r16i",
4147 Sf::R16Float => "r16f",
4148 Sf::Rg8Unorm => "rg8",
4149 Sf::Rg8Snorm => "rg8_snorm",
4150 Sf::Rg8Uint => "rg8ui",
4151 Sf::Rg8Sint => "rg8i",
4152 Sf::R32Uint => "r32ui",
4153 Sf::R32Sint => "r32i",
4154 Sf::R32Float => "r32f",
4155 Sf::Rg16Uint => "rg16ui",
4156 Sf::Rg16Sint => "rg16i",
4157 Sf::Rg16Float => "rg16f",
4158 Sf::Rgba8Unorm => "rgba8",
4159 Sf::Rgba8Snorm => "rgba8_snorm",
4160 Sf::Rgba8Uint => "rgba8ui",
4161 Sf::Rgba8Sint => "rgba8i",
4162 Sf::Rgb10a2Unorm => "rgb10_a2ui",
4163 Sf::Rg11b10Float => "r11f_g11f_b10f",
4164 Sf::Rg32Uint => "rg32ui",
4165 Sf::Rg32Sint => "rg32i",
4166 Sf::Rg32Float => "rg32f",
4167 Sf::Rgba16Uint => "rgba16ui",
4168 Sf::Rgba16Sint => "rgba16i",
4169 Sf::Rgba16Float => "rgba16f",
4170 Sf::Rgba32Uint => "rgba32ui",
4171 Sf::Rgba32Sint => "rgba32i",
4172 Sf::Rgba32Float => "rgba32f",
4173 Sf::R16Unorm => "r16",
4174 Sf::R16Snorm => "r16_snorm",
4175 Sf::Rg16Unorm => "rg16",
4176 Sf::Rg16Snorm => "rg16_snorm",
4177 Sf::Rgba16Unorm => "rgba16",
4178 Sf::Rgba16Snorm => "rgba16_snorm",
4179 }
4180}
4181
4182fn is_value_init_supported(module: &crate::Module, ty: Handle<crate::Type>) -> bool {
4183 match module.types[ty].inner {
4184 TypeInner::Scalar { .. } | TypeInner::Vector { .. } | TypeInner::Matrix { .. } => true,
4185 TypeInner::Array { base, size, .. } => {
4186 size != crate::ArraySize::Dynamic && is_value_init_supported(module, base)
4187 }
4188 TypeInner::Struct { ref members, .. } => members
4189 .iter()
4190 .all(|member| is_value_init_supported(module, member.ty)),
4191 _ => false,
4192 }
4193}