1use alloc::{
2 format,
3 string::{String, ToString},
4 vec,
5 vec::Vec,
6};
7use core::fmt::Write;
8
9use super::Error;
10use super::ToWgslIfImplemented as _;
11use crate::{back::wgsl::polyfill::InversePolyfill, common::wgsl::TypeContext};
12use crate::{
13 back::{self, Baked},
14 common::{
15 self,
16 wgsl::{address_space_str, ToWgsl, TryToWgsl},
17 },
18 proc::{self, NameKey},
19 valid, Handle, Module, ShaderStage, TypeInner,
20};
21
22type BackendResult = Result<(), Error>;
24
25enum Attribute {
27 Binding(u32),
28 BuiltIn(crate::BuiltIn),
29 Group(u32),
30 Invariant,
31 Interpolate(Option<crate::Interpolation>, Option<crate::Sampling>),
32 Location(u32),
33 BlendSrc(u32),
34 Stage(ShaderStage),
35 WorkGroupSize([u32; 3]),
36}
37
38#[derive(Clone, Copy, Debug)]
51enum Indirection {
52 Ordinary,
58
59 Reference,
65}
66
67bitflags::bitflags! {
68 #[cfg_attr(feature = "serialize", derive(serde::Serialize))]
69 #[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
70 #[derive(Clone, Copy, Debug, Eq, PartialEq)]
71 pub struct WriterFlags: u32 {
72 const EXPLICIT_TYPES = 0x1;
74 }
75}
76
77pub struct Writer<W> {
78 out: W,
79 flags: WriterFlags,
80 names: crate::FastHashMap<NameKey, String>,
81 namer: proc::Namer,
82 named_expressions: crate::NamedExpressions,
83 required_polyfills: crate::FastIndexSet<InversePolyfill>,
84}
85
86impl<W: Write> Writer<W> {
87 pub fn new(out: W, flags: WriterFlags) -> Self {
88 Writer {
89 out,
90 flags,
91 names: crate::FastHashMap::default(),
92 namer: proc::Namer::default(),
93 named_expressions: crate::NamedExpressions::default(),
94 required_polyfills: crate::FastIndexSet::default(),
95 }
96 }
97
98 fn reset(&mut self, module: &Module) {
99 self.names.clear();
100 self.namer.reset(
101 module,
102 &crate::keywords::wgsl::RESERVED_SET,
103 &[],
105 &["__", "_naga"],
106 &mut self.names,
107 );
108 self.named_expressions.clear();
109 self.required_polyfills.clear();
110 }
111
112 fn is_builtin_wgsl_struct(&self, module: &Module, handle: Handle<crate::Type>) -> bool {
113 module
114 .special_types
115 .predeclared_types
116 .values()
117 .any(|t| *t == handle)
118 }
119
120 pub fn write(&mut self, module: &Module, info: &valid::ModuleInfo) -> BackendResult {
121 if !module.overrides.is_empty() {
122 return Err(Error::Unimplemented(
123 "Pipeline constants are not yet supported for this back-end".to_string(),
124 ));
125 }
126
127 self.reset(module);
128
129 self.write_enable_declarations(module)?;
131
132 for (handle, ty) in module.types.iter() {
134 if let TypeInner::Struct { ref members, .. } = ty.inner {
135 {
136 if !self.is_builtin_wgsl_struct(module, handle) {
137 self.write_struct(module, handle, members)?;
138 writeln!(self.out)?;
139 }
140 }
141 }
142 }
143
144 let mut constants = module
146 .constants
147 .iter()
148 .filter(|&(_, c)| c.name.is_some())
149 .peekable();
150 while let Some((handle, _)) = constants.next() {
151 self.write_global_constant(module, handle)?;
152 if constants.peek().is_none() {
154 writeln!(self.out)?;
155 }
156 }
157
158 for (ty, global) in module.global_variables.iter() {
160 self.write_global(module, global, ty)?;
161 }
162
163 if !module.global_variables.is_empty() {
164 writeln!(self.out)?;
166 }
167
168 for (handle, function) in module.functions.iter() {
170 let fun_info = &info[handle];
171
172 let func_ctx = back::FunctionCtx {
173 ty: back::FunctionType::Function(handle),
174 info: fun_info,
175 expressions: &function.expressions,
176 named_expressions: &function.named_expressions,
177 };
178
179 self.write_function(module, function, &func_ctx)?;
181
182 writeln!(self.out)?;
183 }
184
185 for (index, ep) in module.entry_points.iter().enumerate() {
187 let attributes = match ep.stage {
188 ShaderStage::Vertex | ShaderStage::Fragment => vec![Attribute::Stage(ep.stage)],
189 ShaderStage::Compute => vec![
190 Attribute::Stage(ShaderStage::Compute),
191 Attribute::WorkGroupSize(ep.workgroup_size),
192 ],
193 ShaderStage::Task | ShaderStage::Mesh => unreachable!(),
194 };
195
196 self.write_attributes(&attributes)?;
197 writeln!(self.out)?;
199
200 let func_ctx = back::FunctionCtx {
201 ty: back::FunctionType::EntryPoint(index as u16),
202 info: info.get_entry_point(index),
203 expressions: &ep.function.expressions,
204 named_expressions: &ep.function.named_expressions,
205 };
206 self.write_function(module, &ep.function, &func_ctx)?;
207
208 if index < module.entry_points.len() - 1 {
209 writeln!(self.out)?;
210 }
211 }
212
213 for polyfill in &self.required_polyfills {
215 writeln!(self.out)?;
216 write!(self.out, "{}", polyfill.source)?;
217 writeln!(self.out)?;
218 }
219
220 Ok(())
221 }
222
223 fn write_enable_declarations(&mut self, module: &Module) -> BackendResult {
226 let mut needs_f16 = false;
227 let mut needs_dual_source_blending = false;
228 let mut needs_clip_distances = false;
229
230 for (_, ty) in module.types.iter() {
232 match ty.inner {
233 TypeInner::Scalar(scalar)
234 | TypeInner::Vector { scalar, .. }
235 | TypeInner::Matrix { scalar, .. } => {
236 needs_f16 |= scalar == crate::Scalar::F16;
237 }
238 TypeInner::Struct { ref members, .. } => {
239 for binding in members.iter().filter_map(|m| m.binding.as_ref()) {
240 match *binding {
241 crate::Binding::Location {
242 blend_src: Some(_), ..
243 } => {
244 needs_dual_source_blending = true;
245 }
246 crate::Binding::BuiltIn(crate::BuiltIn::ClipDistance) => {
247 needs_clip_distances = true;
248 }
249 _ => {}
250 }
251 }
252 }
253 _ => {}
254 }
255 }
256
257 let mut any_written = false;
259 if needs_f16 {
260 writeln!(self.out, "enable f16;")?;
261 any_written = true;
262 }
263 if needs_dual_source_blending {
264 writeln!(self.out, "enable dual_source_blending;")?;
265 any_written = true;
266 }
267 if needs_clip_distances {
268 writeln!(self.out, "enable clip_distances;")?;
269 any_written = true;
270 }
271 if any_written {
272 writeln!(self.out)?;
274 }
275
276 Ok(())
277 }
278
279 fn write_function(
285 &mut self,
286 module: &Module,
287 func: &crate::Function,
288 func_ctx: &back::FunctionCtx<'_>,
289 ) -> BackendResult {
290 let func_name = match func_ctx.ty {
291 back::FunctionType::EntryPoint(index) => &self.names[&NameKey::EntryPoint(index)],
292 back::FunctionType::Function(handle) => &self.names[&NameKey::Function(handle)],
293 };
294
295 write!(self.out, "fn {func_name}(")?;
297
298 for (index, arg) in func.arguments.iter().enumerate() {
300 if let Some(ref binding) = arg.binding {
302 self.write_attributes(&map_binding_to_attribute(binding))?;
303 }
304 let argument_name = &self.names[&func_ctx.argument_key(index as u32)];
306
307 write!(self.out, "{argument_name}: ")?;
308 self.write_type(module, arg.ty)?;
310 if index < func.arguments.len() - 1 {
311 write!(self.out, ", ")?;
313 }
314 }
315
316 write!(self.out, ")")?;
317
318 if let Some(ref result) = func.result {
320 write!(self.out, " -> ")?;
321 if let Some(ref binding) = result.binding {
322 self.write_attributes(&map_binding_to_attribute(binding))?;
323 }
324 self.write_type(module, result.ty)?;
325 }
326
327 write!(self.out, " {{")?;
328 writeln!(self.out)?;
329
330 for (handle, local) in func.local_variables.iter() {
332 write!(self.out, "{}", back::INDENT)?;
334
335 write!(self.out, "var {}: ", self.names[&func_ctx.name_key(handle)])?;
338
339 self.write_type(module, local.ty)?;
341
342 if let Some(init) = local.init {
344 write!(self.out, " = ")?;
347
348 self.write_expr(module, init, func_ctx)?;
351 }
352
353 writeln!(self.out, ";")?
355 }
356
357 if !func.local_variables.is_empty() {
358 writeln!(self.out)?;
359 }
360
361 for sta in func.body.iter() {
363 self.write_stmt(module, sta, func_ctx, back::Level(1))?;
365 }
366
367 writeln!(self.out, "}}")?;
368
369 self.named_expressions.clear();
370
371 Ok(())
372 }
373
374 fn write_attributes(&mut self, attributes: &[Attribute]) -> BackendResult {
376 for attribute in attributes {
377 match *attribute {
378 Attribute::Location(id) => write!(self.out, "@location({id}) ")?,
379 Attribute::BlendSrc(blend_src) => write!(self.out, "@blend_src({blend_src}) ")?,
380 Attribute::BuiltIn(builtin_attrib) => {
381 let builtin = builtin_attrib.to_wgsl_if_implemented()?;
382 write!(self.out, "@builtin({builtin}) ")?;
383 }
384 Attribute::Stage(shader_stage) => {
385 let stage_str = match shader_stage {
386 ShaderStage::Vertex => "vertex",
387 ShaderStage::Fragment => "fragment",
388 ShaderStage::Compute => "compute",
389 ShaderStage::Task | ShaderStage::Mesh => unreachable!(),
390 };
391 write!(self.out, "@{stage_str} ")?;
392 }
393 Attribute::WorkGroupSize(size) => {
394 write!(
395 self.out,
396 "@workgroup_size({}, {}, {}) ",
397 size[0], size[1], size[2]
398 )?;
399 }
400 Attribute::Binding(id) => write!(self.out, "@binding({id}) ")?,
401 Attribute::Group(id) => write!(self.out, "@group({id}) ")?,
402 Attribute::Invariant => write!(self.out, "@invariant ")?,
403 Attribute::Interpolate(interpolation, sampling) => {
404 if sampling.is_some() && sampling != Some(crate::Sampling::Center) {
405 let interpolation = interpolation
406 .unwrap_or(crate::Interpolation::Perspective)
407 .to_wgsl();
408 let sampling = sampling.unwrap_or(crate::Sampling::Center).to_wgsl();
409 write!(self.out, "@interpolate({interpolation}, {sampling}) ")?;
410 } else if interpolation.is_some()
411 && interpolation != Some(crate::Interpolation::Perspective)
412 {
413 let interpolation = interpolation
414 .unwrap_or(crate::Interpolation::Perspective)
415 .to_wgsl();
416 write!(self.out, "@interpolate({interpolation}) ")?;
417 }
418 }
419 };
420 }
421 Ok(())
422 }
423
424 fn write_struct(
435 &mut self,
436 module: &Module,
437 handle: Handle<crate::Type>,
438 members: &[crate::StructMember],
439 ) -> BackendResult {
440 write!(self.out, "struct {}", self.names[&NameKey::Type(handle)])?;
441 write!(self.out, " {{")?;
442 writeln!(self.out)?;
443 for (index, member) in members.iter().enumerate() {
444 write!(self.out, "{}", back::INDENT)?;
446 if let Some(ref binding) = member.binding {
447 self.write_attributes(&map_binding_to_attribute(binding))?;
448 }
449 let member_name = &self.names[&NameKey::StructMember(handle, index as u32)];
451 write!(self.out, "{member_name}: ")?;
452 self.write_type(module, member.ty)?;
453 write!(self.out, ",")?;
454 writeln!(self.out)?;
455 }
456
457 writeln!(self.out, "}}")?;
458
459 Ok(())
460 }
461
462 fn write_type(&mut self, module: &Module, ty: Handle<crate::Type>) -> BackendResult {
463 let type_context = WriterTypeContext {
467 module,
468 names: &self.names,
469 };
470 type_context.write_type(ty, &mut self.out)?;
471
472 Ok(())
473 }
474
475 fn write_type_resolution(
476 &mut self,
477 module: &Module,
478 resolution: &proc::TypeResolution,
479 ) -> BackendResult {
480 let type_context = WriterTypeContext {
484 module,
485 names: &self.names,
486 };
487 type_context.write_type_resolution(resolution, &mut self.out)?;
488
489 Ok(())
490 }
491
492 fn write_stmt(
497 &mut self,
498 module: &Module,
499 stmt: &crate::Statement,
500 func_ctx: &back::FunctionCtx<'_>,
501 level: back::Level,
502 ) -> BackendResult {
503 use crate::{Expression, Statement};
504
505 match *stmt {
506 Statement::Emit(ref range) => {
507 for handle in range.clone() {
508 let info = &func_ctx.info[handle];
509 let expr_name = if let Some(name) = func_ctx.named_expressions.get(&handle) {
510 Some(self.namer.call(name))
515 } else {
516 let expr = &func_ctx.expressions[handle];
517 let min_ref_count = expr.bake_ref_count();
518 let required_baking_expr = match *expr {
520 Expression::ImageLoad { .. }
521 | Expression::ImageQuery { .. }
522 | Expression::ImageSample { .. } => true,
523 _ => false,
524 };
525 if min_ref_count <= info.ref_count || required_baking_expr {
526 Some(Baked(handle).to_string())
527 } else {
528 None
529 }
530 };
531
532 if let Some(name) = expr_name {
533 write!(self.out, "{level}")?;
534 self.start_named_expr(module, handle, func_ctx, &name)?;
535 self.write_expr(module, handle, func_ctx)?;
536 self.named_expressions.insert(handle, name);
537 writeln!(self.out, ";")?;
538 }
539 }
540 }
541 Statement::If {
543 condition,
544 ref accept,
545 ref reject,
546 } => {
547 write!(self.out, "{level}")?;
548 write!(self.out, "if ")?;
549 self.write_expr(module, condition, func_ctx)?;
550 writeln!(self.out, " {{")?;
551
552 let l2 = level.next();
553 for sta in accept {
554 self.write_stmt(module, sta, func_ctx, l2)?;
556 }
557
558 if !reject.is_empty() {
561 writeln!(self.out, "{level}}} else {{")?;
562
563 for sta in reject {
564 self.write_stmt(module, sta, func_ctx, l2)?;
566 }
567 }
568
569 writeln!(self.out, "{level}}}")?
570 }
571 Statement::Return { value } => {
572 write!(self.out, "{level}")?;
573 write!(self.out, "return")?;
574 if let Some(return_value) = value {
575 write!(self.out, " ")?;
577 self.write_expr(module, return_value, func_ctx)?;
578 }
579 writeln!(self.out, ";")?;
580 }
581 Statement::Kill => {
583 write!(self.out, "{level}")?;
584 writeln!(self.out, "discard;")?
585 }
586 Statement::Store { pointer, value } => {
587 write!(self.out, "{level}")?;
588
589 let is_atomic_pointer = func_ctx
590 .resolve_type(pointer, &module.types)
591 .is_atomic_pointer(&module.types);
592
593 if is_atomic_pointer {
594 write!(self.out, "atomicStore(")?;
595 self.write_expr(module, pointer, func_ctx)?;
596 write!(self.out, ", ")?;
597 self.write_expr(module, value, func_ctx)?;
598 write!(self.out, ")")?;
599 } else {
600 self.write_expr_with_indirection(
601 module,
602 pointer,
603 func_ctx,
604 Indirection::Reference,
605 )?;
606 write!(self.out, " = ")?;
607 self.write_expr(module, value, func_ctx)?;
608 }
609 writeln!(self.out, ";")?
610 }
611 Statement::Call {
612 function,
613 ref arguments,
614 result,
615 } => {
616 write!(self.out, "{level}")?;
617 if let Some(expr) = result {
618 let name = Baked(expr).to_string();
619 self.start_named_expr(module, expr, func_ctx, &name)?;
620 self.named_expressions.insert(expr, name);
621 }
622 let func_name = &self.names[&NameKey::Function(function)];
623 write!(self.out, "{func_name}(")?;
624 for (index, &argument) in arguments.iter().enumerate() {
625 if index != 0 {
626 write!(self.out, ", ")?;
627 }
628 self.write_expr(module, argument, func_ctx)?;
629 }
630 writeln!(self.out, ");")?
631 }
632 Statement::Atomic {
633 pointer,
634 ref fun,
635 value,
636 result,
637 } => {
638 write!(self.out, "{level}")?;
639 if let Some(result) = result {
640 let res_name = Baked(result).to_string();
641 self.start_named_expr(module, result, func_ctx, &res_name)?;
642 self.named_expressions.insert(result, res_name);
643 }
644
645 let fun_str = fun.to_wgsl();
646 write!(self.out, "atomic{fun_str}(")?;
647 self.write_expr(module, pointer, func_ctx)?;
648 if let crate::AtomicFunction::Exchange { compare: Some(cmp) } = *fun {
649 write!(self.out, ", ")?;
650 self.write_expr(module, cmp, func_ctx)?;
651 }
652 write!(self.out, ", ")?;
653 self.write_expr(module, value, func_ctx)?;
654 writeln!(self.out, ");")?
655 }
656 Statement::ImageAtomic {
657 image,
658 coordinate,
659 array_index,
660 ref fun,
661 value,
662 } => {
663 write!(self.out, "{level}")?;
664 let fun_str = fun.to_wgsl();
665 write!(self.out, "textureAtomic{fun_str}(")?;
666 self.write_expr(module, image, func_ctx)?;
667 write!(self.out, ", ")?;
668 self.write_expr(module, coordinate, func_ctx)?;
669 if let Some(array_index_expr) = array_index {
670 write!(self.out, ", ")?;
671 self.write_expr(module, array_index_expr, func_ctx)?;
672 }
673 write!(self.out, ", ")?;
674 self.write_expr(module, value, func_ctx)?;
675 writeln!(self.out, ");")?;
676 }
677 Statement::WorkGroupUniformLoad { pointer, result } => {
678 write!(self.out, "{level}")?;
679 let res_name = Baked(result).to_string();
681 self.start_named_expr(module, result, func_ctx, &res_name)?;
682 self.named_expressions.insert(result, res_name);
683 write!(self.out, "workgroupUniformLoad(")?;
684 self.write_expr(module, pointer, func_ctx)?;
685 writeln!(self.out, ");")?;
686 }
687 Statement::ImageStore {
688 image,
689 coordinate,
690 array_index,
691 value,
692 } => {
693 write!(self.out, "{level}")?;
694 write!(self.out, "textureStore(")?;
695 self.write_expr(module, image, func_ctx)?;
696 write!(self.out, ", ")?;
697 self.write_expr(module, coordinate, func_ctx)?;
698 if let Some(array_index_expr) = array_index {
699 write!(self.out, ", ")?;
700 self.write_expr(module, array_index_expr, func_ctx)?;
701 }
702 write!(self.out, ", ")?;
703 self.write_expr(module, value, func_ctx)?;
704 writeln!(self.out, ");")?;
705 }
706 Statement::Block(ref block) => {
708 write!(self.out, "{level}")?;
709 writeln!(self.out, "{{")?;
710 for sta in block.iter() {
711 self.write_stmt(module, sta, func_ctx, level.next())?
713 }
714 writeln!(self.out, "{level}}}")?
715 }
716 Statement::Switch {
717 selector,
718 ref cases,
719 } => {
720 write!(self.out, "{level}")?;
722 write!(self.out, "switch ")?;
723 self.write_expr(module, selector, func_ctx)?;
724 writeln!(self.out, " {{")?;
725
726 let l2 = level.next();
727 let mut new_case = true;
728 for case in cases {
729 if case.fall_through && !case.body.is_empty() {
730 return Err(Error::Unimplemented(
732 "fall-through switch case block".into(),
733 ));
734 }
735
736 match case.value {
737 crate::SwitchValue::I32(value) => {
738 if new_case {
739 write!(self.out, "{l2}case ")?;
740 }
741 write!(self.out, "{value}")?;
742 }
743 crate::SwitchValue::U32(value) => {
744 if new_case {
745 write!(self.out, "{l2}case ")?;
746 }
747 write!(self.out, "{value}u")?;
748 }
749 crate::SwitchValue::Default => {
750 if new_case {
751 if case.fall_through {
752 write!(self.out, "{l2}case ")?;
753 } else {
754 write!(self.out, "{l2}")?;
755 }
756 }
757 write!(self.out, "default")?;
758 }
759 }
760
761 new_case = !case.fall_through;
762
763 if case.fall_through {
764 write!(self.out, ", ")?;
765 } else {
766 writeln!(self.out, ": {{")?;
767 }
768
769 for sta in case.body.iter() {
770 self.write_stmt(module, sta, func_ctx, l2.next())?;
771 }
772
773 if !case.fall_through {
774 writeln!(self.out, "{l2}}}")?;
775 }
776 }
777
778 writeln!(self.out, "{level}}}")?
779 }
780 Statement::Loop {
781 ref body,
782 ref continuing,
783 break_if,
784 } => {
785 write!(self.out, "{level}")?;
786 writeln!(self.out, "loop {{")?;
787
788 let l2 = level.next();
789 for sta in body.iter() {
790 self.write_stmt(module, sta, func_ctx, l2)?;
791 }
792
793 if !continuing.is_empty() || break_if.is_some() {
798 writeln!(self.out, "{l2}continuing {{")?;
799 for sta in continuing.iter() {
800 self.write_stmt(module, sta, func_ctx, l2.next())?;
801 }
802
803 if let Some(condition) = break_if {
806 write!(self.out, "{}break if ", l2.next())?;
808 self.write_expr(module, condition, func_ctx)?;
809 writeln!(self.out, ";")?;
811 }
812
813 writeln!(self.out, "{l2}}}")?;
814 }
815
816 writeln!(self.out, "{level}}}")?
817 }
818 Statement::Break => {
819 writeln!(self.out, "{level}break;")?;
820 }
821 Statement::Continue => {
822 writeln!(self.out, "{level}continue;")?;
823 }
824 Statement::ControlBarrier(barrier) | Statement::MemoryBarrier(barrier) => {
825 if barrier.contains(crate::Barrier::STORAGE) {
826 writeln!(self.out, "{level}storageBarrier();")?;
827 }
828
829 if barrier.contains(crate::Barrier::WORK_GROUP) {
830 writeln!(self.out, "{level}workgroupBarrier();")?;
831 }
832
833 if barrier.contains(crate::Barrier::SUB_GROUP) {
834 writeln!(self.out, "{level}subgroupBarrier();")?;
835 }
836
837 if barrier.contains(crate::Barrier::TEXTURE) {
838 writeln!(self.out, "{level}textureBarrier();")?;
839 }
840 }
841 Statement::RayQuery { .. } => unreachable!(),
842 Statement::SubgroupBallot { result, predicate } => {
843 write!(self.out, "{level}")?;
844 let res_name = Baked(result).to_string();
845 self.start_named_expr(module, result, func_ctx, &res_name)?;
846 self.named_expressions.insert(result, res_name);
847
848 write!(self.out, "subgroupBallot(")?;
849 if let Some(predicate) = predicate {
850 self.write_expr(module, predicate, func_ctx)?;
851 }
852 writeln!(self.out, ");")?;
853 }
854 Statement::SubgroupCollectiveOperation {
855 op,
856 collective_op,
857 argument,
858 result,
859 } => {
860 write!(self.out, "{level}")?;
861 let res_name = Baked(result).to_string();
862 self.start_named_expr(module, result, func_ctx, &res_name)?;
863 self.named_expressions.insert(result, res_name);
864
865 match (collective_op, op) {
866 (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::All) => {
867 write!(self.out, "subgroupAll(")?
868 }
869 (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Any) => {
870 write!(self.out, "subgroupAny(")?
871 }
872 (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Add) => {
873 write!(self.out, "subgroupAdd(")?
874 }
875 (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Mul) => {
876 write!(self.out, "subgroupMul(")?
877 }
878 (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Max) => {
879 write!(self.out, "subgroupMax(")?
880 }
881 (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Min) => {
882 write!(self.out, "subgroupMin(")?
883 }
884 (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::And) => {
885 write!(self.out, "subgroupAnd(")?
886 }
887 (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Or) => {
888 write!(self.out, "subgroupOr(")?
889 }
890 (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Xor) => {
891 write!(self.out, "subgroupXor(")?
892 }
893 (crate::CollectiveOperation::ExclusiveScan, crate::SubgroupOperation::Add) => {
894 write!(self.out, "subgroupExclusiveAdd(")?
895 }
896 (crate::CollectiveOperation::ExclusiveScan, crate::SubgroupOperation::Mul) => {
897 write!(self.out, "subgroupExclusiveMul(")?
898 }
899 (crate::CollectiveOperation::InclusiveScan, crate::SubgroupOperation::Add) => {
900 write!(self.out, "subgroupInclusiveAdd(")?
901 }
902 (crate::CollectiveOperation::InclusiveScan, crate::SubgroupOperation::Mul) => {
903 write!(self.out, "subgroupInclusiveMul(")?
904 }
905 _ => unimplemented!(),
906 }
907 self.write_expr(module, argument, func_ctx)?;
908 writeln!(self.out, ");")?;
909 }
910 Statement::SubgroupGather {
911 mode,
912 argument,
913 result,
914 } => {
915 write!(self.out, "{level}")?;
916 let res_name = Baked(result).to_string();
917 self.start_named_expr(module, result, func_ctx, &res_name)?;
918 self.named_expressions.insert(result, res_name);
919
920 match mode {
921 crate::GatherMode::BroadcastFirst => {
922 write!(self.out, "subgroupBroadcastFirst(")?;
923 }
924 crate::GatherMode::Broadcast(_) => {
925 write!(self.out, "subgroupBroadcast(")?;
926 }
927 crate::GatherMode::Shuffle(_) => {
928 write!(self.out, "subgroupShuffle(")?;
929 }
930 crate::GatherMode::ShuffleDown(_) => {
931 write!(self.out, "subgroupShuffleDown(")?;
932 }
933 crate::GatherMode::ShuffleUp(_) => {
934 write!(self.out, "subgroupShuffleUp(")?;
935 }
936 crate::GatherMode::ShuffleXor(_) => {
937 write!(self.out, "subgroupShuffleXor(")?;
938 }
939 crate::GatherMode::QuadBroadcast(_) => {
940 write!(self.out, "quadBroadcast(")?;
941 }
942 crate::GatherMode::QuadSwap(direction) => match direction {
943 crate::Direction::X => {
944 write!(self.out, "quadSwapX(")?;
945 }
946 crate::Direction::Y => {
947 write!(self.out, "quadSwapY(")?;
948 }
949 crate::Direction::Diagonal => {
950 write!(self.out, "quadSwapDiagonal(")?;
951 }
952 },
953 }
954 self.write_expr(module, argument, func_ctx)?;
955 match mode {
956 crate::GatherMode::BroadcastFirst => {}
957 crate::GatherMode::Broadcast(index)
958 | crate::GatherMode::Shuffle(index)
959 | crate::GatherMode::ShuffleDown(index)
960 | crate::GatherMode::ShuffleUp(index)
961 | crate::GatherMode::ShuffleXor(index)
962 | crate::GatherMode::QuadBroadcast(index) => {
963 write!(self.out, ", ")?;
964 self.write_expr(module, index, func_ctx)?;
965 }
966 crate::GatherMode::QuadSwap(_) => {}
967 }
968 writeln!(self.out, ");")?;
969 }
970 }
971
972 Ok(())
973 }
974
975 fn plain_form_indirection(
998 &self,
999 expr: Handle<crate::Expression>,
1000 module: &Module,
1001 func_ctx: &back::FunctionCtx<'_>,
1002 ) -> Indirection {
1003 use crate::Expression as Ex;
1004
1005 if self.named_expressions.contains_key(&expr) {
1009 return Indirection::Ordinary;
1010 }
1011
1012 match func_ctx.expressions[expr] {
1013 Ex::LocalVariable(_) => Indirection::Reference,
1014 Ex::GlobalVariable(handle) => {
1015 let global = &module.global_variables[handle];
1016 match global.space {
1017 crate::AddressSpace::Handle => Indirection::Ordinary,
1018 _ => Indirection::Reference,
1019 }
1020 }
1021 Ex::Access { base, .. } | Ex::AccessIndex { base, .. } => {
1022 let base_ty = func_ctx.resolve_type(base, &module.types);
1023 match *base_ty {
1024 TypeInner::Pointer { .. } | TypeInner::ValuePointer { .. } => {
1025 Indirection::Reference
1026 }
1027 _ => Indirection::Ordinary,
1028 }
1029 }
1030 _ => Indirection::Ordinary,
1031 }
1032 }
1033
1034 fn start_named_expr(
1035 &mut self,
1036 module: &Module,
1037 handle: Handle<crate::Expression>,
1038 func_ctx: &back::FunctionCtx,
1039 name: &str,
1040 ) -> BackendResult {
1041 write!(self.out, "let {name}")?;
1043 if self.flags.contains(WriterFlags::EXPLICIT_TYPES) {
1044 write!(self.out, ": ")?;
1045 self.write_type_resolution(module, &func_ctx.info[handle].ty)?;
1047 }
1048
1049 write!(self.out, " = ")?;
1050 Ok(())
1051 }
1052
1053 fn write_expr(
1057 &mut self,
1058 module: &Module,
1059 expr: Handle<crate::Expression>,
1060 func_ctx: &back::FunctionCtx<'_>,
1061 ) -> BackendResult {
1062 self.write_expr_with_indirection(module, expr, func_ctx, Indirection::Ordinary)
1063 }
1064
1065 fn write_expr_with_indirection(
1078 &mut self,
1079 module: &Module,
1080 expr: Handle<crate::Expression>,
1081 func_ctx: &back::FunctionCtx<'_>,
1082 requested: Indirection,
1083 ) -> BackendResult {
1084 let plain = self.plain_form_indirection(expr, module, func_ctx);
1087 match (requested, plain) {
1088 (Indirection::Ordinary, Indirection::Reference) => {
1089 write!(self.out, "(&")?;
1090 self.write_expr_plain_form(module, expr, func_ctx, plain)?;
1091 write!(self.out, ")")?;
1092 }
1093 (Indirection::Reference, Indirection::Ordinary) => {
1094 write!(self.out, "(*")?;
1095 self.write_expr_plain_form(module, expr, func_ctx, plain)?;
1096 write!(self.out, ")")?;
1097 }
1098 (_, _) => self.write_expr_plain_form(module, expr, func_ctx, plain)?,
1099 }
1100
1101 Ok(())
1102 }
1103
1104 fn write_const_expression(
1105 &mut self,
1106 module: &Module,
1107 expr: Handle<crate::Expression>,
1108 arena: &crate::Arena<crate::Expression>,
1109 ) -> BackendResult {
1110 self.write_possibly_const_expression(module, expr, arena, |writer, expr| {
1111 writer.write_const_expression(module, expr, arena)
1112 })
1113 }
1114
1115 fn write_possibly_const_expression<E>(
1116 &mut self,
1117 module: &Module,
1118 expr: Handle<crate::Expression>,
1119 expressions: &crate::Arena<crate::Expression>,
1120 write_expression: E,
1121 ) -> BackendResult
1122 where
1123 E: Fn(&mut Self, Handle<crate::Expression>) -> BackendResult,
1124 {
1125 use crate::Expression;
1126
1127 match expressions[expr] {
1128 Expression::Literal(literal) => match literal {
1129 crate::Literal::F16(value) => write!(self.out, "{value}h")?,
1130 crate::Literal::F32(value) => write!(self.out, "{value}f")?,
1131 crate::Literal::U32(value) => write!(self.out, "{value}u")?,
1132 crate::Literal::I32(value) => {
1133 if value == i32::MIN {
1137 write!(self.out, "i32({value})")?;
1138 } else {
1139 write!(self.out, "{value}i")?;
1140 }
1141 }
1142 crate::Literal::Bool(value) => write!(self.out, "{value}")?,
1143 crate::Literal::F64(value) => write!(self.out, "{value:?}lf")?,
1144 crate::Literal::I64(value) => {
1145 if value == i64::MIN {
1150 write!(self.out, "i64({} - 1)", value + 1)?;
1151 } else {
1152 write!(self.out, "{value}li")?;
1153 }
1154 }
1155 crate::Literal::U64(value) => write!(self.out, "{value:?}lu")?,
1156 crate::Literal::AbstractInt(_) | crate::Literal::AbstractFloat(_) => {
1157 return Err(Error::Custom(
1158 "Abstract types should not appear in IR presented to backends".into(),
1159 ));
1160 }
1161 },
1162 Expression::Constant(handle) => {
1163 let constant = &module.constants[handle];
1164 if constant.name.is_some() {
1165 write!(self.out, "{}", self.names[&NameKey::Constant(handle)])?;
1166 } else {
1167 self.write_const_expression(module, constant.init, &module.global_expressions)?;
1168 }
1169 }
1170 Expression::ZeroValue(ty) => {
1171 self.write_type(module, ty)?;
1172 write!(self.out, "()")?;
1173 }
1174 Expression::Compose { ty, ref components } => {
1175 self.write_type(module, ty)?;
1176 write!(self.out, "(")?;
1177 for (index, component) in components.iter().enumerate() {
1178 if index != 0 {
1179 write!(self.out, ", ")?;
1180 }
1181 write_expression(self, *component)?;
1182 }
1183 write!(self.out, ")")?
1184 }
1185 Expression::Splat { size, value } => {
1186 let size = common::vector_size_str(size);
1187 write!(self.out, "vec{size}(")?;
1188 write_expression(self, value)?;
1189 write!(self.out, ")")?;
1190 }
1191 _ => unreachable!(),
1192 }
1193
1194 Ok(())
1195 }
1196
1197 fn write_expr_plain_form(
1205 &mut self,
1206 module: &Module,
1207 expr: Handle<crate::Expression>,
1208 func_ctx: &back::FunctionCtx<'_>,
1209 indirection: Indirection,
1210 ) -> BackendResult {
1211 use crate::Expression;
1212
1213 if let Some(name) = self.named_expressions.get(&expr) {
1214 write!(self.out, "{name}")?;
1215 return Ok(());
1216 }
1217
1218 let expression = &func_ctx.expressions[expr];
1219
1220 match *expression {
1229 Expression::Literal(_)
1230 | Expression::Constant(_)
1231 | Expression::ZeroValue(_)
1232 | Expression::Compose { .. }
1233 | Expression::Splat { .. } => {
1234 self.write_possibly_const_expression(
1235 module,
1236 expr,
1237 func_ctx.expressions,
1238 |writer, expr| writer.write_expr(module, expr, func_ctx),
1239 )?;
1240 }
1241 Expression::Override(_) => unreachable!(),
1242 Expression::FunctionArgument(pos) => {
1243 let name_key = func_ctx.argument_key(pos);
1244 let name = &self.names[&name_key];
1245 write!(self.out, "{name}")?;
1246 }
1247 Expression::Binary { op, left, right } => {
1248 write!(self.out, "(")?;
1249 self.write_expr(module, left, func_ctx)?;
1250 write!(self.out, " {} ", back::binary_operation_str(op))?;
1251 self.write_expr(module, right, func_ctx)?;
1252 write!(self.out, ")")?;
1253 }
1254 Expression::Access { base, index } => {
1255 self.write_expr_with_indirection(module, base, func_ctx, indirection)?;
1256 write!(self.out, "[")?;
1257 self.write_expr(module, index, func_ctx)?;
1258 write!(self.out, "]")?
1259 }
1260 Expression::AccessIndex { base, index } => {
1261 let base_ty_res = &func_ctx.info[base].ty;
1262 let mut resolved = base_ty_res.inner_with(&module.types);
1263
1264 self.write_expr_with_indirection(module, base, func_ctx, indirection)?;
1265
1266 let base_ty_handle = match *resolved {
1267 TypeInner::Pointer { base, space: _ } => {
1268 resolved = &module.types[base].inner;
1269 Some(base)
1270 }
1271 _ => base_ty_res.handle(),
1272 };
1273
1274 match *resolved {
1275 TypeInner::Vector { .. } => {
1276 write!(self.out, ".{}", back::COMPONENTS[index as usize])?
1278 }
1279 TypeInner::Matrix { .. }
1280 | TypeInner::Array { .. }
1281 | TypeInner::BindingArray { .. }
1282 | TypeInner::ValuePointer { .. } => write!(self.out, "[{index}]")?,
1283 TypeInner::Struct { .. } => {
1284 let ty = base_ty_handle.unwrap();
1287
1288 write!(
1289 self.out,
1290 ".{}",
1291 &self.names[&NameKey::StructMember(ty, index)]
1292 )?
1293 }
1294 ref other => return Err(Error::Custom(format!("Cannot index {other:?}"))),
1295 }
1296 }
1297 Expression::ImageSample {
1298 image,
1299 sampler,
1300 gather: None,
1301 coordinate,
1302 array_index,
1303 offset,
1304 level,
1305 depth_ref,
1306 clamp_to_edge,
1307 } => {
1308 use crate::SampleLevel as Sl;
1309
1310 let suffix_cmp = match depth_ref {
1311 Some(_) => "Compare",
1312 None => "",
1313 };
1314 let suffix_level = match level {
1315 Sl::Auto => "",
1316 Sl::Zero if clamp_to_edge => "BaseClampToEdge",
1317 Sl::Zero | Sl::Exact(_) => "Level",
1318 Sl::Bias(_) => "Bias",
1319 Sl::Gradient { .. } => "Grad",
1320 };
1321
1322 write!(self.out, "textureSample{suffix_cmp}{suffix_level}(")?;
1323 self.write_expr(module, image, func_ctx)?;
1324 write!(self.out, ", ")?;
1325 self.write_expr(module, sampler, func_ctx)?;
1326 write!(self.out, ", ")?;
1327 self.write_expr(module, coordinate, func_ctx)?;
1328
1329 if let Some(array_index) = array_index {
1330 write!(self.out, ", ")?;
1331 self.write_expr(module, array_index, func_ctx)?;
1332 }
1333
1334 if let Some(depth_ref) = depth_ref {
1335 write!(self.out, ", ")?;
1336 self.write_expr(module, depth_ref, func_ctx)?;
1337 }
1338
1339 match level {
1340 Sl::Auto => {}
1341 Sl::Zero => {
1342 if depth_ref.is_none() && !clamp_to_edge {
1344 write!(self.out, ", 0.0")?;
1345 }
1346 }
1347 Sl::Exact(expr) => {
1348 write!(self.out, ", ")?;
1349 self.write_expr(module, expr, func_ctx)?;
1350 }
1351 Sl::Bias(expr) => {
1352 write!(self.out, ", ")?;
1353 self.write_expr(module, expr, func_ctx)?;
1354 }
1355 Sl::Gradient { x, y } => {
1356 write!(self.out, ", ")?;
1357 self.write_expr(module, x, func_ctx)?;
1358 write!(self.out, ", ")?;
1359 self.write_expr(module, y, func_ctx)?;
1360 }
1361 }
1362
1363 if let Some(offset) = offset {
1364 write!(self.out, ", ")?;
1365 self.write_const_expression(module, offset, func_ctx.expressions)?;
1366 }
1367
1368 write!(self.out, ")")?;
1369 }
1370
1371 Expression::ImageSample {
1372 image,
1373 sampler,
1374 gather: Some(component),
1375 coordinate,
1376 array_index,
1377 offset,
1378 level: _,
1379 depth_ref,
1380 clamp_to_edge: _,
1381 } => {
1382 let suffix_cmp = match depth_ref {
1383 Some(_) => "Compare",
1384 None => "",
1385 };
1386
1387 write!(self.out, "textureGather{suffix_cmp}(")?;
1388 match *func_ctx.resolve_type(image, &module.types) {
1389 TypeInner::Image {
1390 class: crate::ImageClass::Depth { multi: _ },
1391 ..
1392 } => {}
1393 _ => {
1394 write!(self.out, "{}, ", component as u8)?;
1395 }
1396 }
1397 self.write_expr(module, image, func_ctx)?;
1398 write!(self.out, ", ")?;
1399 self.write_expr(module, sampler, func_ctx)?;
1400 write!(self.out, ", ")?;
1401 self.write_expr(module, coordinate, func_ctx)?;
1402
1403 if let Some(array_index) = array_index {
1404 write!(self.out, ", ")?;
1405 self.write_expr(module, array_index, func_ctx)?;
1406 }
1407
1408 if let Some(depth_ref) = depth_ref {
1409 write!(self.out, ", ")?;
1410 self.write_expr(module, depth_ref, func_ctx)?;
1411 }
1412
1413 if let Some(offset) = offset {
1414 write!(self.out, ", ")?;
1415 self.write_const_expression(module, offset, func_ctx.expressions)?;
1416 }
1417
1418 write!(self.out, ")")?;
1419 }
1420 Expression::ImageQuery { image, query } => {
1421 use crate::ImageQuery as Iq;
1422
1423 let texture_function = match query {
1424 Iq::Size { .. } => "textureDimensions",
1425 Iq::NumLevels => "textureNumLevels",
1426 Iq::NumLayers => "textureNumLayers",
1427 Iq::NumSamples => "textureNumSamples",
1428 };
1429
1430 write!(self.out, "{texture_function}(")?;
1431 self.write_expr(module, image, func_ctx)?;
1432 if let Iq::Size { level: Some(level) } = query {
1433 write!(self.out, ", ")?;
1434 self.write_expr(module, level, func_ctx)?;
1435 };
1436 write!(self.out, ")")?;
1437 }
1438
1439 Expression::ImageLoad {
1440 image,
1441 coordinate,
1442 array_index,
1443 sample,
1444 level,
1445 } => {
1446 write!(self.out, "textureLoad(")?;
1447 self.write_expr(module, image, func_ctx)?;
1448 write!(self.out, ", ")?;
1449 self.write_expr(module, coordinate, func_ctx)?;
1450 if let Some(array_index) = array_index {
1451 write!(self.out, ", ")?;
1452 self.write_expr(module, array_index, func_ctx)?;
1453 }
1454 if let Some(index) = sample.or(level) {
1455 write!(self.out, ", ")?;
1456 self.write_expr(module, index, func_ctx)?;
1457 }
1458 write!(self.out, ")")?;
1459 }
1460 Expression::GlobalVariable(handle) => {
1461 let name = &self.names[&NameKey::GlobalVariable(handle)];
1462 write!(self.out, "{name}")?;
1463 }
1464
1465 Expression::As {
1466 expr,
1467 kind,
1468 convert,
1469 } => {
1470 let inner = func_ctx.resolve_type(expr, &module.types);
1471 match *inner {
1472 TypeInner::Matrix {
1473 columns,
1474 rows,
1475 scalar,
1476 } => {
1477 let scalar = crate::Scalar {
1478 kind,
1479 width: convert.unwrap_or(scalar.width),
1480 };
1481 let scalar_kind_str = scalar.to_wgsl_if_implemented()?;
1482 write!(
1483 self.out,
1484 "mat{}x{}<{}>",
1485 common::vector_size_str(columns),
1486 common::vector_size_str(rows),
1487 scalar_kind_str
1488 )?;
1489 }
1490 TypeInner::Vector {
1491 size,
1492 scalar: crate::Scalar { width, .. },
1493 } => {
1494 let scalar = crate::Scalar {
1495 kind,
1496 width: convert.unwrap_or(width),
1497 };
1498 let vector_size_str = common::vector_size_str(size);
1499 let scalar_kind_str = scalar.to_wgsl_if_implemented()?;
1500 if convert.is_some() {
1501 write!(self.out, "vec{vector_size_str}<{scalar_kind_str}>")?;
1502 } else {
1503 write!(self.out, "bitcast<vec{vector_size_str}<{scalar_kind_str}>>")?;
1504 }
1505 }
1506 TypeInner::Scalar(crate::Scalar { width, .. }) => {
1507 let scalar = crate::Scalar {
1508 kind,
1509 width: convert.unwrap_or(width),
1510 };
1511 let scalar_kind_str = scalar.to_wgsl_if_implemented()?;
1512 if convert.is_some() {
1513 write!(self.out, "{scalar_kind_str}")?
1514 } else {
1515 write!(self.out, "bitcast<{scalar_kind_str}>")?
1516 }
1517 }
1518 _ => {
1519 return Err(Error::Unimplemented(format!(
1520 "write_expr expression::as {inner:?}"
1521 )));
1522 }
1523 };
1524 write!(self.out, "(")?;
1525 self.write_expr(module, expr, func_ctx)?;
1526 write!(self.out, ")")?;
1527 }
1528 Expression::Load { pointer } => {
1529 let is_atomic_pointer = func_ctx
1530 .resolve_type(pointer, &module.types)
1531 .is_atomic_pointer(&module.types);
1532
1533 if is_atomic_pointer {
1534 write!(self.out, "atomicLoad(")?;
1535 self.write_expr(module, pointer, func_ctx)?;
1536 write!(self.out, ")")?;
1537 } else {
1538 self.write_expr_with_indirection(
1539 module,
1540 pointer,
1541 func_ctx,
1542 Indirection::Reference,
1543 )?;
1544 }
1545 }
1546 Expression::LocalVariable(handle) => {
1547 write!(self.out, "{}", self.names[&func_ctx.name_key(handle)])?
1548 }
1549 Expression::ArrayLength(expr) => {
1550 write!(self.out, "arrayLength(")?;
1551 self.write_expr(module, expr, func_ctx)?;
1552 write!(self.out, ")")?;
1553 }
1554
1555 Expression::Math {
1556 fun,
1557 arg,
1558 arg1,
1559 arg2,
1560 arg3,
1561 } => {
1562 use crate::MathFunction as Mf;
1563
1564 enum Function {
1565 Regular(&'static str),
1566 InversePolyfill(InversePolyfill),
1567 }
1568
1569 let function = match fun.try_to_wgsl() {
1570 Some(name) => Function::Regular(name),
1571 None => match fun {
1572 Mf::Inverse => {
1573 let ty = func_ctx.resolve_type(arg, &module.types);
1574 let Some(overload) = InversePolyfill::find_overload(ty) else {
1575 return Err(Error::unsupported("math function", fun));
1576 };
1577
1578 Function::InversePolyfill(overload)
1579 }
1580 _ => return Err(Error::unsupported("math function", fun)),
1581 },
1582 };
1583
1584 match function {
1585 Function::Regular(fun_name) => {
1586 write!(self.out, "{fun_name}(")?;
1587 self.write_expr(module, arg, func_ctx)?;
1588 for arg in IntoIterator::into_iter([arg1, arg2, arg3]).flatten() {
1589 write!(self.out, ", ")?;
1590 self.write_expr(module, arg, func_ctx)?;
1591 }
1592 write!(self.out, ")")?
1593 }
1594 Function::InversePolyfill(inverse) => {
1595 write!(self.out, "{}(", inverse.fun_name)?;
1596 self.write_expr(module, arg, func_ctx)?;
1597 write!(self.out, ")")?;
1598 self.required_polyfills.insert(inverse);
1599 }
1600 }
1601 }
1602
1603 Expression::Swizzle {
1604 size,
1605 vector,
1606 pattern,
1607 } => {
1608 self.write_expr(module, vector, func_ctx)?;
1609 write!(self.out, ".")?;
1610 for &sc in pattern[..size as usize].iter() {
1611 self.out.write_char(back::COMPONENTS[sc as usize])?;
1612 }
1613 }
1614 Expression::Unary { op, expr } => {
1615 let unary = match op {
1616 crate::UnaryOperator::Negate => "-",
1617 crate::UnaryOperator::LogicalNot => "!",
1618 crate::UnaryOperator::BitwiseNot => "~",
1619 };
1620
1621 write!(self.out, "{unary}(")?;
1622 self.write_expr(module, expr, func_ctx)?;
1623
1624 write!(self.out, ")")?
1625 }
1626
1627 Expression::Select {
1628 condition,
1629 accept,
1630 reject,
1631 } => {
1632 write!(self.out, "select(")?;
1633 self.write_expr(module, reject, func_ctx)?;
1634 write!(self.out, ", ")?;
1635 self.write_expr(module, accept, func_ctx)?;
1636 write!(self.out, ", ")?;
1637 self.write_expr(module, condition, func_ctx)?;
1638 write!(self.out, ")")?
1639 }
1640 Expression::Derivative { axis, ctrl, expr } => {
1641 use crate::{DerivativeAxis as Axis, DerivativeControl as Ctrl};
1642 let op = match (axis, ctrl) {
1643 (Axis::X, Ctrl::Coarse) => "dpdxCoarse",
1644 (Axis::X, Ctrl::Fine) => "dpdxFine",
1645 (Axis::X, Ctrl::None) => "dpdx",
1646 (Axis::Y, Ctrl::Coarse) => "dpdyCoarse",
1647 (Axis::Y, Ctrl::Fine) => "dpdyFine",
1648 (Axis::Y, Ctrl::None) => "dpdy",
1649 (Axis::Width, Ctrl::Coarse) => "fwidthCoarse",
1650 (Axis::Width, Ctrl::Fine) => "fwidthFine",
1651 (Axis::Width, Ctrl::None) => "fwidth",
1652 };
1653 write!(self.out, "{op}(")?;
1654 self.write_expr(module, expr, func_ctx)?;
1655 write!(self.out, ")")?
1656 }
1657 Expression::Relational { fun, argument } => {
1658 use crate::RelationalFunction as Rf;
1659
1660 let fun_name = match fun {
1661 Rf::All => "all",
1662 Rf::Any => "any",
1663 _ => return Err(Error::UnsupportedRelationalFunction(fun)),
1664 };
1665 write!(self.out, "{fun_name}(")?;
1666
1667 self.write_expr(module, argument, func_ctx)?;
1668
1669 write!(self.out, ")")?
1670 }
1671 Expression::RayQueryGetIntersection { .. }
1673 | Expression::RayQueryVertexPositions { .. } => unreachable!(),
1674 Expression::CallResult(_)
1676 | Expression::AtomicResult { .. }
1677 | Expression::RayQueryProceedResult
1678 | Expression::SubgroupBallotResult
1679 | Expression::SubgroupOperationResult { .. }
1680 | Expression::WorkGroupUniformLoadResult { .. } => {}
1681 }
1682
1683 Ok(())
1684 }
1685
1686 fn write_global(
1690 &mut self,
1691 module: &Module,
1692 global: &crate::GlobalVariable,
1693 handle: Handle<crate::GlobalVariable>,
1694 ) -> BackendResult {
1695 if let Some(ref binding) = global.binding {
1697 self.write_attributes(&[
1698 Attribute::Group(binding.group),
1699 Attribute::Binding(binding.binding),
1700 ])?;
1701 writeln!(self.out)?;
1702 }
1703
1704 write!(self.out, "var")?;
1706 let (address, maybe_access) = address_space_str(global.space);
1707 if let Some(space) = address {
1708 write!(self.out, "<{space}")?;
1709 if let Some(access) = maybe_access {
1710 write!(self.out, ", {access}")?;
1711 }
1712 write!(self.out, ">")?;
1713 }
1714 write!(
1715 self.out,
1716 " {}: ",
1717 &self.names[&NameKey::GlobalVariable(handle)]
1718 )?;
1719
1720 self.write_type(module, global.ty)?;
1722
1723 if let Some(init) = global.init {
1725 write!(self.out, " = ")?;
1726 self.write_const_expression(module, init, &module.global_expressions)?;
1727 }
1728
1729 writeln!(self.out, ";")?;
1731
1732 Ok(())
1733 }
1734
1735 fn write_global_constant(
1740 &mut self,
1741 module: &Module,
1742 handle: Handle<crate::Constant>,
1743 ) -> BackendResult {
1744 let name = &self.names[&NameKey::Constant(handle)];
1745 write!(self.out, "const {name}: ")?;
1747 self.write_type(module, module.constants[handle].ty)?;
1748 write!(self.out, " = ")?;
1749 let init = module.constants[handle].init;
1750 self.write_const_expression(module, init, &module.global_expressions)?;
1751 writeln!(self.out, ";")?;
1752
1753 Ok(())
1754 }
1755
1756 #[allow(clippy::missing_const_for_fn)]
1758 pub fn finish(self) -> W {
1759 self.out
1760 }
1761}
1762
1763struct WriterTypeContext<'m> {
1764 module: &'m Module,
1765 names: &'m crate::FastHashMap<NameKey, String>,
1766}
1767
1768impl TypeContext for WriterTypeContext<'_> {
1769 fn lookup_type(&self, handle: Handle<crate::Type>) -> &crate::Type {
1770 &self.module.types[handle]
1771 }
1772
1773 fn type_name(&self, handle: Handle<crate::Type>) -> &str {
1774 self.names[&NameKey::Type(handle)].as_str()
1775 }
1776
1777 fn write_unnamed_struct<W: Write>(&self, _: &TypeInner, _: &mut W) -> core::fmt::Result {
1778 unreachable!("the WGSL back end should always provide type handles");
1779 }
1780
1781 fn write_override<W: Write>(&self, _: Handle<crate::Override>, _: &mut W) -> core::fmt::Result {
1782 unreachable!("overrides should be validated out");
1783 }
1784
1785 fn write_non_wgsl_inner<W: Write>(&self, _: &TypeInner, _: &mut W) -> core::fmt::Result {
1786 unreachable!("backends should only be passed validated modules");
1787 }
1788
1789 fn write_non_wgsl_scalar<W: Write>(&self, _: crate::Scalar, _: &mut W) -> core::fmt::Result {
1790 unreachable!("backends should only be passed validated modules");
1791 }
1792}
1793
1794fn map_binding_to_attribute(binding: &crate::Binding) -> Vec<Attribute> {
1795 match *binding {
1796 crate::Binding::BuiltIn(built_in) => {
1797 if let crate::BuiltIn::Position { invariant: true } = built_in {
1798 vec![Attribute::BuiltIn(built_in), Attribute::Invariant]
1799 } else {
1800 vec![Attribute::BuiltIn(built_in)]
1801 }
1802 }
1803 crate::Binding::Location {
1804 location,
1805 interpolation,
1806 sampling,
1807 blend_src: None,
1808 } => vec![
1809 Attribute::Location(location),
1810 Attribute::Interpolate(interpolation, sampling),
1811 ],
1812 crate::Binding::Location {
1813 location,
1814 interpolation,
1815 sampling,
1816 blend_src: Some(blend_src),
1817 } => vec![
1818 Attribute::Location(location),
1819 Attribute::BlendSrc(blend_src),
1820 Attribute::Interpolate(interpolation, sampling),
1821 ],
1822 }
1823}