1use alloc::{
2 borrow::ToOwned,
3 boxed::Box,
4 format,
5 string::{String, ToString},
6 vec::Vec,
7};
8use core::num::NonZeroU32;
9
10use crate::front::wgsl::error::{Error, ExpectedToken, InvalidAssignmentType};
11use crate::front::wgsl::index::Index;
12use crate::front::wgsl::parse::number::Number;
13use crate::front::wgsl::parse::{ast, conv};
14use crate::front::wgsl::Result;
15use crate::front::Typifier;
16use crate::{
17 common::wgsl::{TryToWgsl, TypeContext},
18 compact::KeepUnused,
19};
20use crate::{common::ForDebugWithTypes, proc::LayoutErrorInner};
21use crate::{ir, proc};
22use crate::{Arena, FastHashMap, FastIndexMap, Handle, Span};
23
24mod construction;
25mod conversion;
26
27macro_rules! resolve_inner {
41 ($ctx:ident, $expr:expr) => {{
42 $ctx.grow_types($expr)?;
43 $ctx.typifier()[$expr].inner_with(&$ctx.module.types)
44 }};
45}
46pub(super) use resolve_inner;
47
48macro_rules! resolve_inner_binary {
56 ($ctx:ident, $left:expr, $right:expr) => {{
57 $ctx.grow_types($left)?;
58 $ctx.grow_types($right)?;
59 (
60 $ctx.typifier()[$left].inner_with(&$ctx.module.types),
61 $ctx.typifier()[$right].inner_with(&$ctx.module.types),
62 )
63 }};
64}
65
66macro_rules! resolve {
76 ($ctx:ident, $expr:expr) => {{
77 let expr = $expr;
78 $ctx.grow_types(expr)?;
79 &$ctx.typifier()[expr]
80 }};
81}
82pub(super) use resolve;
83
84pub struct GlobalContext<'source, 'temp, 'out> {
86 ast_expressions: &'temp Arena<ast::Expression<'source>>,
88
89 types: &'temp Arena<ast::Type<'source>>,
91
92 globals: &'temp mut FastHashMap<&'source str, LoweredGlobalDecl>,
96
97 module: &'out mut ir::Module,
99
100 const_typifier: &'temp mut Typifier,
101
102 layouter: &'temp mut proc::Layouter,
103
104 global_expression_kind_tracker: &'temp mut proc::ExpressionKindTracker,
105}
106
107impl<'source> GlobalContext<'source, '_, '_> {
108 fn as_const(&mut self) -> ExpressionContext<'source, '_, '_> {
109 ExpressionContext {
110 ast_expressions: self.ast_expressions,
111 globals: self.globals,
112 types: self.types,
113 module: self.module,
114 const_typifier: self.const_typifier,
115 layouter: self.layouter,
116 expr_type: ExpressionContextType::Constant(None),
117 global_expression_kind_tracker: self.global_expression_kind_tracker,
118 }
119 }
120
121 fn as_override(&mut self) -> ExpressionContext<'source, '_, '_> {
122 ExpressionContext {
123 ast_expressions: self.ast_expressions,
124 globals: self.globals,
125 types: self.types,
126 module: self.module,
127 const_typifier: self.const_typifier,
128 layouter: self.layouter,
129 expr_type: ExpressionContextType::Override,
130 global_expression_kind_tracker: self.global_expression_kind_tracker,
131 }
132 }
133
134 fn ensure_type_exists(
135 &mut self,
136 name: Option<String>,
137 inner: ir::TypeInner,
138 ) -> Handle<ir::Type> {
139 self.module
140 .types
141 .insert(ir::Type { inner, name }, Span::UNDEFINED)
142 }
143}
144
145pub struct StatementContext<'source, 'temp, 'out> {
147 ast_expressions: &'temp Arena<ast::Expression<'source>>,
153
154 types: &'temp Arena<ast::Type<'source>>,
159
160 globals: &'temp mut FastHashMap<&'source str, LoweredGlobalDecl>,
164
165 local_table:
180 &'temp mut FastHashMap<Handle<ast::Local>, Declared<Typed<Handle<ir::Expression>>>>,
181
182 const_typifier: &'temp mut Typifier,
183 typifier: &'temp mut Typifier,
184 layouter: &'temp mut proc::Layouter,
185 function: &'out mut ir::Function,
186 named_expressions: &'out mut FastIndexMap<Handle<ir::Expression>, (String, Span)>,
189 module: &'out mut ir::Module,
190
191 local_expression_kind_tracker: &'temp mut proc::ExpressionKindTracker,
201 global_expression_kind_tracker: &'temp mut proc::ExpressionKindTracker,
202}
203
204impl<'a, 'temp> StatementContext<'a, 'temp, '_> {
205 fn as_const<'t>(
206 &'t mut self,
207 block: &'t mut ir::Block,
208 emitter: &'t mut proc::Emitter,
209 ) -> ExpressionContext<'a, 't, 't>
210 where
211 'temp: 't,
212 {
213 ExpressionContext {
214 globals: self.globals,
215 types: self.types,
216 ast_expressions: self.ast_expressions,
217 const_typifier: self.const_typifier,
218 layouter: self.layouter,
219 global_expression_kind_tracker: self.global_expression_kind_tracker,
220 module: self.module,
221 expr_type: ExpressionContextType::Constant(Some(LocalExpressionContext {
222 local_table: self.local_table,
223 function: self.function,
224 block,
225 emitter,
226 typifier: self.typifier,
227 local_expression_kind_tracker: self.local_expression_kind_tracker,
228 })),
229 }
230 }
231
232 fn as_expression<'t>(
233 &'t mut self,
234 block: &'t mut ir::Block,
235 emitter: &'t mut proc::Emitter,
236 ) -> ExpressionContext<'a, 't, 't>
237 where
238 'temp: 't,
239 {
240 ExpressionContext {
241 globals: self.globals,
242 types: self.types,
243 ast_expressions: self.ast_expressions,
244 const_typifier: self.const_typifier,
245 layouter: self.layouter,
246 global_expression_kind_tracker: self.global_expression_kind_tracker,
247 module: self.module,
248 expr_type: ExpressionContextType::Runtime(LocalExpressionContext {
249 local_table: self.local_table,
250 function: self.function,
251 block,
252 emitter,
253 typifier: self.typifier,
254 local_expression_kind_tracker: self.local_expression_kind_tracker,
255 }),
256 }
257 }
258
259 #[allow(dead_code)]
260 fn as_global(&mut self) -> GlobalContext<'a, '_, '_> {
261 GlobalContext {
262 ast_expressions: self.ast_expressions,
263 globals: self.globals,
264 types: self.types,
265 module: self.module,
266 const_typifier: self.const_typifier,
267 layouter: self.layouter,
268 global_expression_kind_tracker: self.global_expression_kind_tracker,
269 }
270 }
271
272 fn invalid_assignment_type(&self, expr: Handle<ir::Expression>) -> InvalidAssignmentType {
273 if let Some(&(_, span)) = self.named_expressions.get(&expr) {
274 InvalidAssignmentType::ImmutableBinding(span)
275 } else {
276 match self.function.expressions[expr] {
277 ir::Expression::Swizzle { .. } => InvalidAssignmentType::Swizzle,
278 ir::Expression::Access { base, .. } => self.invalid_assignment_type(base),
279 ir::Expression::AccessIndex { base, .. } => self.invalid_assignment_type(base),
280 _ => InvalidAssignmentType::Other,
281 }
282 }
283 }
284}
285
286pub struct LocalExpressionContext<'temp, 'out> {
287 local_table: &'temp FastHashMap<Handle<ast::Local>, Declared<Typed<Handle<ir::Expression>>>>,
292
293 function: &'out mut ir::Function,
294 block: &'temp mut ir::Block,
295 emitter: &'temp mut proc::Emitter,
296 typifier: &'temp mut Typifier,
297
298 local_expression_kind_tracker: &'temp mut proc::ExpressionKindTracker,
303}
304
305pub enum ExpressionContextType<'temp, 'out> {
307 Runtime(LocalExpressionContext<'temp, 'out>),
314
315 Constant(Option<LocalExpressionContext<'temp, 'out>>),
322
323 Override,
330}
331
332pub struct ExpressionContext<'source, 'temp, 'out> {
370 ast_expressions: &'temp Arena<ast::Expression<'source>>,
372 types: &'temp Arena<ast::Type<'source>>,
373
374 globals: &'temp mut FastHashMap<&'source str, LoweredGlobalDecl>,
378
379 module: &'out mut ir::Module,
383
384 const_typifier: &'temp mut Typifier,
388 layouter: &'temp mut proc::Layouter,
389 global_expression_kind_tracker: &'temp mut proc::ExpressionKindTracker,
390
391 expr_type: ExpressionContextType<'temp, 'out>,
394}
395
396impl TypeContext for ExpressionContext<'_, '_, '_> {
397 fn lookup_type(&self, handle: Handle<ir::Type>) -> &ir::Type {
398 &self.module.types[handle]
399 }
400
401 fn type_name(&self, handle: Handle<ir::Type>) -> &str {
402 self.module.types[handle]
403 .name
404 .as_deref()
405 .unwrap_or("{anonymous type}")
406 }
407
408 fn write_override<W: core::fmt::Write>(
409 &self,
410 handle: Handle<ir::Override>,
411 out: &mut W,
412 ) -> core::fmt::Result {
413 match self.module.overrides[handle].name {
414 Some(ref name) => out.write_str(name),
415 None => write!(out, "{{anonymous override {handle:?}}}"),
416 }
417 }
418
419 fn write_unnamed_struct<W: core::fmt::Write>(
420 &self,
421 _: &ir::TypeInner,
422 _: &mut W,
423 ) -> core::fmt::Result {
424 unreachable!("the WGSL front end should always know the type name");
425 }
426}
427
428impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
429 #[allow(dead_code)]
430 fn as_const(&mut self) -> ExpressionContext<'source, '_, '_> {
431 ExpressionContext {
432 globals: self.globals,
433 types: self.types,
434 ast_expressions: self.ast_expressions,
435 const_typifier: self.const_typifier,
436 layouter: self.layouter,
437 module: self.module,
438 expr_type: ExpressionContextType::Constant(match self.expr_type {
439 ExpressionContextType::Runtime(ref mut local_expression_context)
440 | ExpressionContextType::Constant(Some(ref mut local_expression_context)) => {
441 Some(LocalExpressionContext {
442 local_table: local_expression_context.local_table,
443 function: local_expression_context.function,
444 block: local_expression_context.block,
445 emitter: local_expression_context.emitter,
446 typifier: local_expression_context.typifier,
447 local_expression_kind_tracker: local_expression_context
448 .local_expression_kind_tracker,
449 })
450 }
451 ExpressionContextType::Constant(None) | ExpressionContextType::Override => None,
452 }),
453 global_expression_kind_tracker: self.global_expression_kind_tracker,
454 }
455 }
456
457 fn as_global(&mut self) -> GlobalContext<'source, '_, '_> {
458 GlobalContext {
459 ast_expressions: self.ast_expressions,
460 globals: self.globals,
461 types: self.types,
462 module: self.module,
463 const_typifier: self.const_typifier,
464 layouter: self.layouter,
465 global_expression_kind_tracker: self.global_expression_kind_tracker,
466 }
467 }
468
469 fn as_const_evaluator(&mut self) -> proc::ConstantEvaluator<'_> {
470 match self.expr_type {
471 ExpressionContextType::Runtime(ref mut rctx) => {
472 proc::ConstantEvaluator::for_wgsl_function(
473 self.module,
474 &mut rctx.function.expressions,
475 rctx.local_expression_kind_tracker,
476 self.layouter,
477 rctx.emitter,
478 rctx.block,
479 false,
480 )
481 }
482 ExpressionContextType::Constant(Some(ref mut rctx)) => {
483 proc::ConstantEvaluator::for_wgsl_function(
484 self.module,
485 &mut rctx.function.expressions,
486 rctx.local_expression_kind_tracker,
487 self.layouter,
488 rctx.emitter,
489 rctx.block,
490 true,
491 )
492 }
493 ExpressionContextType::Constant(None) => proc::ConstantEvaluator::for_wgsl_module(
494 self.module,
495 self.global_expression_kind_tracker,
496 self.layouter,
497 false,
498 ),
499 ExpressionContextType::Override => proc::ConstantEvaluator::for_wgsl_module(
500 self.module,
501 self.global_expression_kind_tracker,
502 self.layouter,
503 true,
504 ),
505 }
506 }
507
508 fn as_diagnostic_display<T>(
514 &self,
515 value: T,
516 ) -> crate::common::DiagnosticDisplay<(T, proc::GlobalCtx<'_>)> {
517 let ctx = self.module.to_ctx();
518 crate::common::DiagnosticDisplay((value, ctx))
519 }
520
521 fn append_expression(
522 &mut self,
523 expr: ir::Expression,
524 span: Span,
525 ) -> Result<'source, Handle<ir::Expression>> {
526 let mut eval = self.as_const_evaluator();
527 eval.try_eval_and_append(expr, span)
528 .map_err(|e| Box::new(Error::ConstantEvaluatorError(e.into(), span)))
529 }
530
531 fn const_eval_expr_to_u32(
532 &self,
533 handle: Handle<ir::Expression>,
534 ) -> core::result::Result<u32, proc::U32EvalError> {
535 match self.expr_type {
536 ExpressionContextType::Runtime(ref ctx) => {
537 if !ctx.local_expression_kind_tracker.is_const(handle) {
538 return Err(proc::U32EvalError::NonConst);
539 }
540
541 self.module
542 .to_ctx()
543 .eval_expr_to_u32_from(handle, &ctx.function.expressions)
544 }
545 ExpressionContextType::Constant(Some(ref ctx)) => {
546 assert!(ctx.local_expression_kind_tracker.is_const(handle));
547 self.module
548 .to_ctx()
549 .eval_expr_to_u32_from(handle, &ctx.function.expressions)
550 }
551 ExpressionContextType::Constant(None) => self.module.to_ctx().eval_expr_to_u32(handle),
552 ExpressionContextType::Override => Err(proc::U32EvalError::NonConst),
553 }
554 }
555
556 fn is_const(&self, handle: Handle<ir::Expression>) -> bool {
558 use ExpressionContextType as Ect;
559 match self.expr_type {
560 Ect::Runtime(ref ctx) | Ect::Constant(Some(ref ctx)) => {
561 ctx.local_expression_kind_tracker.is_const(handle)
562 }
563 Ect::Constant(None) | Ect::Override => {
564 self.global_expression_kind_tracker.is_const(handle)
565 }
566 }
567 }
568
569 fn get_expression_span(&self, handle: Handle<ir::Expression>) -> Span {
570 match self.expr_type {
571 ExpressionContextType::Runtime(ref ctx)
572 | ExpressionContextType::Constant(Some(ref ctx)) => {
573 ctx.function.expressions.get_span(handle)
574 }
575 ExpressionContextType::Constant(None) | ExpressionContextType::Override => {
576 self.module.global_expressions.get_span(handle)
577 }
578 }
579 }
580
581 fn typifier(&self) -> &Typifier {
582 match self.expr_type {
583 ExpressionContextType::Runtime(ref ctx)
584 | ExpressionContextType::Constant(Some(ref ctx)) => ctx.typifier,
585 ExpressionContextType::Constant(None) | ExpressionContextType::Override => {
586 self.const_typifier
587 }
588 }
589 }
590
591 fn local(
592 &mut self,
593 local: &Handle<ast::Local>,
594 span: Span,
595 ) -> Result<'source, Typed<Handle<ir::Expression>>> {
596 match self.expr_type {
597 ExpressionContextType::Runtime(ref ctx) => Ok(ctx.local_table[local].runtime()),
598 ExpressionContextType::Constant(Some(ref ctx)) => ctx.local_table[local]
599 .const_time()
600 .ok_or(Box::new(Error::UnexpectedOperationInConstContext(span))),
601 _ => Err(Box::new(Error::UnexpectedOperationInConstContext(span))),
602 }
603 }
604
605 fn runtime_expression_ctx(
606 &mut self,
607 span: Span,
608 ) -> Result<'source, &mut LocalExpressionContext<'temp, 'out>> {
609 match self.expr_type {
610 ExpressionContextType::Runtime(ref mut ctx) => Ok(ctx),
611 ExpressionContextType::Constant(_) | ExpressionContextType::Override => {
612 Err(Box::new(Error::UnexpectedOperationInConstContext(span)))
613 }
614 }
615 }
616
617 fn gather_component(
618 &mut self,
619 expr: Handle<ir::Expression>,
620 component_span: Span,
621 gather_span: Span,
622 ) -> Result<'source, ir::SwizzleComponent> {
623 match self.expr_type {
624 ExpressionContextType::Runtime(ref rctx) => {
625 if !rctx.local_expression_kind_tracker.is_const(expr) {
626 return Err(Box::new(Error::ExpectedConstExprConcreteIntegerScalar(
627 component_span,
628 )));
629 }
630
631 let index = self
632 .module
633 .to_ctx()
634 .eval_expr_to_u32_from(expr, &rctx.function.expressions)
635 .map_err(|err| match err {
636 proc::U32EvalError::NonConst => {
637 Error::ExpectedConstExprConcreteIntegerScalar(component_span)
638 }
639 proc::U32EvalError::Negative => Error::ExpectedNonNegative(component_span),
640 })?;
641 ir::SwizzleComponent::XYZW
642 .get(index as usize)
643 .copied()
644 .ok_or(Box::new(Error::InvalidGatherComponent(component_span)))
645 }
646 ExpressionContextType::Constant(_) | ExpressionContextType::Override => Err(Box::new(
649 Error::UnexpectedOperationInConstContext(gather_span),
650 )),
651 }
652 }
653
654 fn register_type(
665 &mut self,
666 handle: Handle<ir::Expression>,
667 ) -> Result<'source, Handle<ir::Type>> {
668 self.grow_types(handle)?;
669 let typifier = match self.expr_type {
673 ExpressionContextType::Runtime(ref ctx)
674 | ExpressionContextType::Constant(Some(ref ctx)) => ctx.typifier,
675 ExpressionContextType::Constant(None) | ExpressionContextType::Override => {
676 &*self.const_typifier
677 }
678 };
679 Ok(typifier.register_type(handle, &mut self.module.types))
680 }
681
682 fn grow_types(&mut self, handle: Handle<ir::Expression>) -> Result<'source, &mut Self> {
703 let empty_arena = Arena::new();
704 let resolve_ctx;
705 let typifier;
706 let expressions;
707 match self.expr_type {
708 ExpressionContextType::Runtime(ref mut ctx)
709 | ExpressionContextType::Constant(Some(ref mut ctx)) => {
710 resolve_ctx = proc::ResolveContext::with_locals(
711 self.module,
712 &ctx.function.local_variables,
713 &ctx.function.arguments,
714 );
715 typifier = &mut *ctx.typifier;
716 expressions = &ctx.function.expressions;
717 }
718 ExpressionContextType::Constant(None) | ExpressionContextType::Override => {
719 resolve_ctx = proc::ResolveContext::with_locals(self.module, &empty_arena, &[]);
720 typifier = self.const_typifier;
721 expressions = &self.module.global_expressions;
722 }
723 };
724 typifier
725 .grow(handle, expressions, &resolve_ctx)
726 .map_err(Error::InvalidResolve)?;
727
728 Ok(self)
729 }
730
731 fn image_data(
732 &mut self,
733 image: Handle<ir::Expression>,
734 span: Span,
735 ) -> Result<'source, (ir::ImageClass, bool)> {
736 match *resolve_inner!(self, image) {
737 ir::TypeInner::Image { class, arrayed, .. } => Ok((class, arrayed)),
738 _ => Err(Box::new(Error::BadTexture(span))),
739 }
740 }
741
742 fn prepare_args<'b>(
743 &mut self,
744 args: &'b [Handle<ast::Expression<'source>>],
745 min_args: u32,
746 span: Span,
747 ) -> ArgumentContext<'b, 'source> {
748 ArgumentContext {
749 args: args.iter(),
750 min_args,
751 args_used: 0,
752 total_args: args.len() as u32,
753 span,
754 }
755 }
756
757 fn binary_op_splat(
765 &mut self,
766 op: ir::BinaryOperator,
767 left: &mut Handle<ir::Expression>,
768 right: &mut Handle<ir::Expression>,
769 ) -> Result<'source, ()> {
770 if matches!(
771 op,
772 ir::BinaryOperator::Add
773 | ir::BinaryOperator::Subtract
774 | ir::BinaryOperator::Divide
775 | ir::BinaryOperator::Modulo
776 ) {
777 match resolve_inner_binary!(self, *left, *right) {
778 (&ir::TypeInner::Vector { size, .. }, &ir::TypeInner::Scalar { .. }) => {
779 *right = self.append_expression(
780 ir::Expression::Splat {
781 size,
782 value: *right,
783 },
784 self.get_expression_span(*right),
785 )?;
786 }
787 (&ir::TypeInner::Scalar { .. }, &ir::TypeInner::Vector { size, .. }) => {
788 *left = self.append_expression(
789 ir::Expression::Splat { size, value: *left },
790 self.get_expression_span(*left),
791 )?;
792 }
793 _ => {}
794 }
795 }
796
797 Ok(())
798 }
799
800 fn interrupt_emitter(
805 &mut self,
806 expression: ir::Expression,
807 span: Span,
808 ) -> Result<'source, Handle<ir::Expression>> {
809 match self.expr_type {
810 ExpressionContextType::Runtime(ref mut rctx)
811 | ExpressionContextType::Constant(Some(ref mut rctx)) => {
812 rctx.block
813 .extend(rctx.emitter.finish(&rctx.function.expressions));
814 }
815 ExpressionContextType::Constant(None) | ExpressionContextType::Override => {}
816 }
817 let result = self.append_expression(expression, span);
818 match self.expr_type {
819 ExpressionContextType::Runtime(ref mut rctx)
820 | ExpressionContextType::Constant(Some(ref mut rctx)) => {
821 rctx.emitter.start(&rctx.function.expressions);
822 }
823 ExpressionContextType::Constant(None) | ExpressionContextType::Override => {}
824 }
825 result
826 }
827
828 fn apply_load_rule(
833 &mut self,
834 expr: Typed<Handle<ir::Expression>>,
835 ) -> Result<'source, Handle<ir::Expression>> {
836 match expr {
837 Typed::Reference(pointer) => {
838 let load = ir::Expression::Load { pointer };
839 let span = self.get_expression_span(pointer);
840 self.append_expression(load, span)
841 }
842 Typed::Plain(handle) => Ok(handle),
843 }
844 }
845
846 fn ensure_type_exists(&mut self, inner: ir::TypeInner) -> Handle<ir::Type> {
847 self.as_global().ensure_type_exists(None, inner)
848 }
849}
850
851struct ArgumentContext<'ctx, 'source> {
852 args: core::slice::Iter<'ctx, Handle<ast::Expression<'source>>>,
853 min_args: u32,
854 args_used: u32,
855 total_args: u32,
856 span: Span,
857}
858
859impl<'source> ArgumentContext<'_, 'source> {
860 pub fn finish(self) -> Result<'source, ()> {
861 if self.args.len() == 0 {
862 Ok(())
863 } else {
864 Err(Box::new(Error::WrongArgumentCount {
865 found: self.total_args,
866 expected: self.min_args..self.args_used + 1,
867 span: self.span,
868 }))
869 }
870 }
871
872 pub fn next(&mut self) -> Result<'source, Handle<ast::Expression<'source>>> {
873 match self.args.next().copied() {
874 Some(arg) => {
875 self.args_used += 1;
876 Ok(arg)
877 }
878 None => Err(Box::new(Error::WrongArgumentCount {
879 found: self.total_args,
880 expected: self.min_args..self.args_used + 1,
881 span: self.span,
882 })),
883 }
884 }
885}
886
887#[derive(Debug, Copy, Clone)]
888enum Declared<T> {
889 Const(T),
891
892 Runtime(T),
894}
895
896impl<T> Declared<T> {
897 fn runtime(self) -> T {
898 match self {
899 Declared::Const(t) | Declared::Runtime(t) => t,
900 }
901 }
902
903 fn const_time(self) -> Option<T> {
904 match self {
905 Declared::Const(t) => Some(t),
906 Declared::Runtime(_) => None,
907 }
908 }
909}
910
911#[derive(Debug, Copy, Clone)]
933enum Typed<T> {
934 Reference(T),
936
937 Plain(T),
939}
940
941impl<T> Typed<T> {
942 fn map<U>(self, mut f: impl FnMut(T) -> U) -> Typed<U> {
943 match self {
944 Self::Reference(v) => Typed::Reference(f(v)),
945 Self::Plain(v) => Typed::Plain(f(v)),
946 }
947 }
948
949 fn try_map<U, E>(
950 self,
951 mut f: impl FnMut(T) -> core::result::Result<U, E>,
952 ) -> core::result::Result<Typed<U>, E> {
953 Ok(match self {
954 Self::Reference(expr) => Typed::Reference(f(expr)?),
955 Self::Plain(expr) => Typed::Plain(f(expr)?),
956 })
957 }
958}
959
960enum Components {
966 Single(u32),
967 Swizzle {
968 size: ir::VectorSize,
969 pattern: [ir::SwizzleComponent; 4],
970 },
971}
972
973impl Components {
974 const fn letter_component(letter: char) -> Option<ir::SwizzleComponent> {
975 use ir::SwizzleComponent as Sc;
976 match letter {
977 'x' | 'r' => Some(Sc::X),
978 'y' | 'g' => Some(Sc::Y),
979 'z' | 'b' => Some(Sc::Z),
980 'w' | 'a' => Some(Sc::W),
981 _ => None,
982 }
983 }
984
985 fn single_component(name: &str, name_span: Span) -> Result<'_, u32> {
986 let ch = name.chars().next().ok_or(Error::BadAccessor(name_span))?;
987 match Self::letter_component(ch) {
988 Some(sc) => Ok(sc as u32),
989 None => Err(Box::new(Error::BadAccessor(name_span))),
990 }
991 }
992
993 fn new(name: &str, name_span: Span) -> Result<'_, Self> {
997 let size = match name.len() {
998 1 => return Ok(Components::Single(Self::single_component(name, name_span)?)),
999 2 => ir::VectorSize::Bi,
1000 3 => ir::VectorSize::Tri,
1001 4 => ir::VectorSize::Quad,
1002 _ => return Err(Box::new(Error::BadAccessor(name_span))),
1003 };
1004
1005 let mut pattern = [ir::SwizzleComponent::X; 4];
1006 for (comp, ch) in pattern.iter_mut().zip(name.chars()) {
1007 *comp = Self::letter_component(ch).ok_or(Error::BadAccessor(name_span))?;
1008 }
1009
1010 if name.chars().all(|c| matches!(c, 'x' | 'y' | 'z' | 'w'))
1011 || name.chars().all(|c| matches!(c, 'r' | 'g' | 'b' | 'a'))
1012 {
1013 Ok(Components::Swizzle { size, pattern })
1014 } else {
1015 Err(Box::new(Error::BadAccessor(name_span)))
1016 }
1017 }
1018}
1019
1020enum LoweredGlobalDecl {
1022 Function {
1023 handle: Handle<ir::Function>,
1024 must_use: bool,
1025 },
1026 Var(Handle<ir::GlobalVariable>),
1027 Const(Handle<ir::Constant>),
1028 Override(Handle<ir::Override>),
1029 Type(Handle<ir::Type>),
1030 EntryPoint(usize),
1031}
1032
1033enum Texture {
1034 Gather,
1035 GatherCompare,
1036
1037 Sample,
1038 SampleBias,
1039 SampleCompare,
1040 SampleCompareLevel,
1041 SampleGrad,
1042 SampleLevel,
1043 SampleBaseClampToEdge,
1044}
1045
1046impl Texture {
1047 pub fn map(word: &str) -> Option<Self> {
1048 Some(match word {
1049 "textureGather" => Self::Gather,
1050 "textureGatherCompare" => Self::GatherCompare,
1051
1052 "textureSample" => Self::Sample,
1053 "textureSampleBias" => Self::SampleBias,
1054 "textureSampleCompare" => Self::SampleCompare,
1055 "textureSampleCompareLevel" => Self::SampleCompareLevel,
1056 "textureSampleGrad" => Self::SampleGrad,
1057 "textureSampleLevel" => Self::SampleLevel,
1058 "textureSampleBaseClampToEdge" => Self::SampleBaseClampToEdge,
1059 _ => return None,
1060 })
1061 }
1062
1063 pub const fn min_argument_count(&self) -> u32 {
1064 match *self {
1065 Self::Gather => 3,
1066 Self::GatherCompare => 4,
1067
1068 Self::Sample => 3,
1069 Self::SampleBias => 5,
1070 Self::SampleCompare => 5,
1071 Self::SampleCompareLevel => 5,
1072 Self::SampleGrad => 6,
1073 Self::SampleLevel => 5,
1074 Self::SampleBaseClampToEdge => 3,
1075 }
1076 }
1077}
1078
1079enum SubgroupGather {
1080 BroadcastFirst,
1081 Broadcast,
1082 Shuffle,
1083 ShuffleDown,
1084 ShuffleUp,
1085 ShuffleXor,
1086 QuadBroadcast,
1087}
1088
1089impl SubgroupGather {
1090 pub fn map(word: &str) -> Option<Self> {
1091 Some(match word {
1092 "subgroupBroadcastFirst" => Self::BroadcastFirst,
1093 "subgroupBroadcast" => Self::Broadcast,
1094 "subgroupShuffle" => Self::Shuffle,
1095 "subgroupShuffleDown" => Self::ShuffleDown,
1096 "subgroupShuffleUp" => Self::ShuffleUp,
1097 "subgroupShuffleXor" => Self::ShuffleXor,
1098 "quadBroadcast" => Self::QuadBroadcast,
1099 _ => return None,
1100 })
1101 }
1102}
1103
1104enum AbstractRule {
1106 Concretize,
1108
1109 Allow,
1111}
1112
1113pub struct Lowerer<'source, 'temp> {
1114 index: &'temp Index<'source>,
1115}
1116
1117impl<'source, 'temp> Lowerer<'source, 'temp> {
1118 pub const fn new(index: &'temp Index<'source>) -> Self {
1119 Self { index }
1120 }
1121
1122 pub fn lower(&mut self, tu: ast::TranslationUnit<'source>) -> Result<'source, ir::Module> {
1123 let mut module = ir::Module {
1124 diagnostic_filters: tu.diagnostic_filters,
1125 diagnostic_filter_leaf: tu.diagnostic_filter_leaf,
1126 ..Default::default()
1127 };
1128
1129 let mut ctx = GlobalContext {
1130 ast_expressions: &tu.expressions,
1131 globals: &mut FastHashMap::default(),
1132 types: &tu.types,
1133 module: &mut module,
1134 const_typifier: &mut Typifier::new(),
1135 layouter: &mut proc::Layouter::default(),
1136 global_expression_kind_tracker: &mut proc::ExpressionKindTracker::new(),
1137 };
1138 if !tu.doc_comments.is_empty() {
1139 ctx.module.get_or_insert_default_doc_comments().module =
1140 tu.doc_comments.iter().map(|s| s.to_string()).collect();
1141 }
1142
1143 for decl_handle in self.index.visit_ordered() {
1144 let span = tu.decls.get_span(decl_handle);
1145 let decl = &tu.decls[decl_handle];
1146
1147 match decl.kind {
1148 ast::GlobalDeclKind::Fn(ref f) => {
1149 let lowered_decl = self.function(f, span, &mut ctx)?;
1150 if !f.doc_comments.is_empty() {
1151 match lowered_decl {
1152 LoweredGlobalDecl::Function { handle, .. } => {
1153 ctx.module
1154 .get_or_insert_default_doc_comments()
1155 .functions
1156 .insert(
1157 handle,
1158 f.doc_comments.iter().map(|s| s.to_string()).collect(),
1159 );
1160 }
1161 LoweredGlobalDecl::EntryPoint(index) => {
1162 ctx.module
1163 .get_or_insert_default_doc_comments()
1164 .entry_points
1165 .insert(
1166 index,
1167 f.doc_comments.iter().map(|s| s.to_string()).collect(),
1168 );
1169 }
1170 _ => {}
1171 }
1172 }
1173 ctx.globals.insert(f.name.name, lowered_decl);
1174 }
1175 ast::GlobalDeclKind::Var(ref v) => {
1176 let explicit_ty =
1177 v.ty.map(|ast| self.resolve_ast_type(ast, &mut ctx.as_const()))
1178 .transpose()?;
1179
1180 let (ty, initializer) = self.type_and_init(
1181 v.name,
1182 v.init,
1183 explicit_ty,
1184 AbstractRule::Concretize,
1185 &mut ctx.as_override(),
1186 )?;
1187
1188 let binding = if let Some(ref binding) = v.binding {
1189 Some(ir::ResourceBinding {
1190 group: self.const_u32(binding.group, &mut ctx.as_const())?.0,
1191 binding: self.const_u32(binding.binding, &mut ctx.as_const())?.0,
1192 })
1193 } else {
1194 None
1195 };
1196
1197 let handle = ctx.module.global_variables.append(
1198 ir::GlobalVariable {
1199 name: Some(v.name.name.to_string()),
1200 space: v.space,
1201 binding,
1202 ty,
1203 init: initializer,
1204 },
1205 span,
1206 );
1207
1208 if !v.doc_comments.is_empty() {
1209 ctx.module
1210 .get_or_insert_default_doc_comments()
1211 .global_variables
1212 .insert(
1213 handle,
1214 v.doc_comments.iter().map(|s| s.to_string()).collect(),
1215 );
1216 }
1217 ctx.globals
1218 .insert(v.name.name, LoweredGlobalDecl::Var(handle));
1219 }
1220 ast::GlobalDeclKind::Const(ref c) => {
1221 let mut ectx = ctx.as_const();
1222
1223 let explicit_ty =
1224 c.ty.map(|ast| self.resolve_ast_type(ast, &mut ectx))
1225 .transpose()?;
1226
1227 let (ty, init) = self.type_and_init(
1228 c.name,
1229 Some(c.init),
1230 explicit_ty,
1231 AbstractRule::Allow,
1232 &mut ectx,
1233 )?;
1234 let init = init.expect("Global const must have init");
1235
1236 let handle = ctx.module.constants.append(
1237 ir::Constant {
1238 name: Some(c.name.name.to_string()),
1239 ty,
1240 init,
1241 },
1242 span,
1243 );
1244
1245 ctx.globals
1246 .insert(c.name.name, LoweredGlobalDecl::Const(handle));
1247 if !c.doc_comments.is_empty() {
1248 ctx.module
1249 .get_or_insert_default_doc_comments()
1250 .constants
1251 .insert(
1252 handle,
1253 c.doc_comments.iter().map(|s| s.to_string()).collect(),
1254 );
1255 }
1256 }
1257 ast::GlobalDeclKind::Override(ref o) => {
1258 let explicit_ty =
1259 o.ty.map(|ast| self.resolve_ast_type(ast, &mut ctx.as_const()))
1260 .transpose()?;
1261
1262 let mut ectx = ctx.as_override();
1263
1264 let (ty, init) = self.type_and_init(
1265 o.name,
1266 o.init,
1267 explicit_ty,
1268 AbstractRule::Concretize,
1269 &mut ectx,
1270 )?;
1271
1272 let id =
1273 o.id.map(|id| self.const_u32(id, &mut ctx.as_const()))
1274 .transpose()?;
1275
1276 let id = if let Some((id, id_span)) = id {
1277 Some(
1278 u16::try_from(id)
1279 .map_err(|_| Error::PipelineConstantIDValue(id_span))?,
1280 )
1281 } else {
1282 None
1283 };
1284
1285 let handle = ctx.module.overrides.append(
1286 ir::Override {
1287 name: Some(o.name.name.to_string()),
1288 id,
1289 ty,
1290 init,
1291 },
1292 span,
1293 );
1294
1295 ctx.globals
1296 .insert(o.name.name, LoweredGlobalDecl::Override(handle));
1297 }
1298 ast::GlobalDeclKind::Struct(ref s) => {
1299 let handle = self.r#struct(s, span, &mut ctx)?;
1300 ctx.globals
1301 .insert(s.name.name, LoweredGlobalDecl::Type(handle));
1302 if !s.doc_comments.is_empty() {
1303 ctx.module
1304 .get_or_insert_default_doc_comments()
1305 .types
1306 .insert(
1307 handle,
1308 s.doc_comments.iter().map(|s| s.to_string()).collect(),
1309 );
1310 }
1311 }
1312 ast::GlobalDeclKind::Type(ref alias) => {
1313 let ty = self.resolve_named_ast_type(
1314 alias.ty,
1315 Some(alias.name.name.to_string()),
1316 &mut ctx.as_const(),
1317 )?;
1318 ctx.globals
1319 .insert(alias.name.name, LoweredGlobalDecl::Type(ty));
1320 }
1321 ast::GlobalDeclKind::ConstAssert(condition) => {
1322 let condition = self.expression(condition, &mut ctx.as_const())?;
1323
1324 let span = ctx.module.global_expressions.get_span(condition);
1325 match ctx
1326 .module
1327 .to_ctx()
1328 .eval_expr_to_bool_from(condition, &ctx.module.global_expressions)
1329 {
1330 Some(true) => Ok(()),
1331 Some(false) => Err(Error::ConstAssertFailed(span)),
1332 _ => Err(Error::NotBool(span)),
1333 }?;
1334 }
1335 }
1336 }
1337
1338 crate::compact::compact(&mut module, KeepUnused::Yes);
1342
1343 Ok(module)
1344 }
1345
1346 fn type_and_init(
1348 &mut self,
1349 name: ast::Ident<'source>,
1350 init: Option<Handle<ast::Expression<'source>>>,
1351 explicit_ty: Option<Handle<ir::Type>>,
1352 abstract_rule: AbstractRule,
1353 ectx: &mut ExpressionContext<'source, '_, '_>,
1354 ) -> Result<'source, (Handle<ir::Type>, Option<Handle<ir::Expression>>)> {
1355 let ty;
1356 let initializer;
1357 match (init, explicit_ty) {
1358 (Some(init), Some(explicit_ty)) => {
1359 let init = self.expression_for_abstract(init, ectx)?;
1360 let ty_res = proc::TypeResolution::Handle(explicit_ty);
1361 let init = ectx
1362 .try_automatic_conversions(init, &ty_res, name.span)
1363 .map_err(|error| match *error {
1364 Error::AutoConversion(e) => Box::new(Error::InitializationTypeMismatch {
1365 name: name.span,
1366 expected: e.dest_type,
1367 got: e.source_type,
1368 }),
1369 _ => error,
1370 })?;
1371
1372 let init_ty = ectx.register_type(init)?;
1373 if !ectx.module.compare_types(
1374 &proc::TypeResolution::Handle(explicit_ty),
1375 &proc::TypeResolution::Handle(init_ty),
1376 ) {
1377 return Err(Box::new(Error::InitializationTypeMismatch {
1378 name: name.span,
1379 expected: ectx.type_to_string(explicit_ty),
1380 got: ectx.type_to_string(init_ty),
1381 }));
1382 }
1383 ty = explicit_ty;
1384 initializer = Some(init);
1385 }
1386 (Some(init), None) => {
1387 let mut init = self.expression_for_abstract(init, ectx)?;
1388 if let AbstractRule::Concretize = abstract_rule {
1389 init = ectx.concretize(init)?;
1390 }
1391 ty = ectx.register_type(init)?;
1392 initializer = Some(init);
1393 }
1394 (None, Some(explicit_ty)) => {
1395 ty = explicit_ty;
1396 initializer = None;
1397 }
1398 (None, None) => return Err(Box::new(Error::DeclMissingTypeAndInit(name.span))),
1399 }
1400 Ok((ty, initializer))
1401 }
1402
1403 fn function(
1404 &mut self,
1405 f: &ast::Function<'source>,
1406 span: Span,
1407 ctx: &mut GlobalContext<'source, '_, '_>,
1408 ) -> Result<'source, LoweredGlobalDecl> {
1409 let mut local_table = FastHashMap::default();
1410 let mut expressions = Arena::new();
1411 let mut named_expressions = FastIndexMap::default();
1412 let mut local_expression_kind_tracker = proc::ExpressionKindTracker::new();
1413
1414 let arguments = f
1415 .arguments
1416 .iter()
1417 .enumerate()
1418 .map(|(i, arg)| -> Result<'_, _> {
1419 let ty = self.resolve_ast_type(arg.ty, &mut ctx.as_const())?;
1420 let expr =
1421 expressions.append(ir::Expression::FunctionArgument(i as u32), arg.name.span);
1422 local_table.insert(arg.handle, Declared::Runtime(Typed::Plain(expr)));
1423 named_expressions.insert(expr, (arg.name.name.to_string(), arg.name.span));
1424 local_expression_kind_tracker.insert(expr, proc::ExpressionKind::Runtime);
1425
1426 Ok(ir::FunctionArgument {
1427 name: Some(arg.name.name.to_string()),
1428 ty,
1429 binding: self.binding(&arg.binding, ty, ctx)?,
1430 })
1431 })
1432 .collect::<Result<Vec<_>>>()?;
1433
1434 let result = f
1435 .result
1436 .as_ref()
1437 .map(|res| -> Result<'_, _> {
1438 let ty = self.resolve_ast_type(res.ty, &mut ctx.as_const())?;
1439 Ok(ir::FunctionResult {
1440 ty,
1441 binding: self.binding(&res.binding, ty, ctx)?,
1442 })
1443 })
1444 .transpose()?;
1445
1446 let mut function = ir::Function {
1447 name: Some(f.name.name.to_string()),
1448 arguments,
1449 result,
1450 local_variables: Arena::new(),
1451 expressions,
1452 named_expressions: crate::NamedExpressions::default(),
1453 body: ir::Block::default(),
1454 diagnostic_filter_leaf: f.diagnostic_filter_leaf,
1455 };
1456
1457 let mut typifier = Typifier::default();
1458 let mut stmt_ctx = StatementContext {
1459 local_table: &mut local_table,
1460 globals: ctx.globals,
1461 ast_expressions: ctx.ast_expressions,
1462 const_typifier: ctx.const_typifier,
1463 typifier: &mut typifier,
1464 layouter: ctx.layouter,
1465 function: &mut function,
1466 named_expressions: &mut named_expressions,
1467 types: ctx.types,
1468 module: ctx.module,
1469 local_expression_kind_tracker: &mut local_expression_kind_tracker,
1470 global_expression_kind_tracker: ctx.global_expression_kind_tracker,
1471 };
1472 let mut body = self.block(&f.body, false, &mut stmt_ctx)?;
1473 proc::ensure_block_returns(&mut body);
1474
1475 function.body = body;
1476 function.named_expressions = named_expressions
1477 .into_iter()
1478 .map(|(key, (name, _))| (key, name))
1479 .collect();
1480
1481 if let Some(ref entry) = f.entry_point {
1482 let workgroup_size_info = if let Some(workgroup_size) = entry.workgroup_size {
1483 let mut workgroup_size_out = [1; 3];
1485 let mut workgroup_size_overrides_out = [None; 3];
1486 for (i, size) in workgroup_size.into_iter().enumerate() {
1487 if let Some(size_expr) = size {
1488 match self.const_u32(size_expr, &mut ctx.as_const()) {
1489 Ok(value) => {
1490 workgroup_size_out[i] = value.0;
1491 }
1492 Err(err) => {
1493 if let Error::ConstantEvaluatorError(ref ty, _) = *err {
1494 match **ty {
1495 proc::ConstantEvaluatorError::OverrideExpr => {
1496 workgroup_size_overrides_out[i] =
1497 Some(self.workgroup_size_override(
1498 size_expr,
1499 &mut ctx.as_override(),
1500 )?);
1501 }
1502 _ => {
1503 return Err(err);
1504 }
1505 }
1506 } else {
1507 return Err(err);
1508 }
1509 }
1510 }
1511 }
1512 }
1513 if workgroup_size_overrides_out.iter().all(|x| x.is_none()) {
1514 (workgroup_size_out, None)
1515 } else {
1516 (workgroup_size_out, Some(workgroup_size_overrides_out))
1517 }
1518 } else {
1519 ([0; 3], None)
1520 };
1521
1522 let (workgroup_size, workgroup_size_overrides) = workgroup_size_info;
1523 ctx.module.entry_points.push(ir::EntryPoint {
1524 name: f.name.name.to_string(),
1525 stage: entry.stage,
1526 early_depth_test: entry.early_depth_test,
1527 workgroup_size,
1528 workgroup_size_overrides,
1529 function,
1530 });
1531 Ok(LoweredGlobalDecl::EntryPoint(
1532 ctx.module.entry_points.len() - 1,
1533 ))
1534 } else {
1535 let handle = ctx.module.functions.append(function, span);
1536 Ok(LoweredGlobalDecl::Function {
1537 handle,
1538 must_use: f.result.as_ref().is_some_and(|res| res.must_use),
1539 })
1540 }
1541 }
1542
1543 fn workgroup_size_override(
1544 &mut self,
1545 size_expr: Handle<ast::Expression<'source>>,
1546 ctx: &mut ExpressionContext<'source, '_, '_>,
1547 ) -> Result<'source, Handle<ir::Expression>> {
1548 let span = ctx.ast_expressions.get_span(size_expr);
1549 let expr = self.expression(size_expr, ctx)?;
1550 match resolve_inner!(ctx, expr).scalar_kind().ok_or(0) {
1551 Ok(ir::ScalarKind::Sint) | Ok(ir::ScalarKind::Uint) => Ok(expr),
1552 _ => Err(Box::new(Error::ExpectedConstExprConcreteIntegerScalar(
1553 span,
1554 ))),
1555 }
1556 }
1557
1558 fn block(
1559 &mut self,
1560 b: &ast::Block<'source>,
1561 is_inside_loop: bool,
1562 ctx: &mut StatementContext<'source, '_, '_>,
1563 ) -> Result<'source, ir::Block> {
1564 let mut block = ir::Block::default();
1565
1566 for stmt in b.stmts.iter() {
1567 self.statement(stmt, &mut block, is_inside_loop, ctx)?;
1568 }
1569
1570 Ok(block)
1571 }
1572
1573 fn statement(
1574 &mut self,
1575 stmt: &ast::Statement<'source>,
1576 block: &mut ir::Block,
1577 is_inside_loop: bool,
1578 ctx: &mut StatementContext<'source, '_, '_>,
1579 ) -> Result<'source, ()> {
1580 let out = match stmt.kind {
1581 ast::StatementKind::Block(ref block) => {
1582 let block = self.block(block, is_inside_loop, ctx)?;
1583 ir::Statement::Block(block)
1584 }
1585 ast::StatementKind::LocalDecl(ref decl) => match *decl {
1586 ast::LocalDecl::Let(ref l) => {
1587 let mut emitter = proc::Emitter::default();
1588 emitter.start(&ctx.function.expressions);
1589
1590 let explicit_ty = l
1591 .ty
1592 .map(|ty| self.resolve_ast_type(ty, &mut ctx.as_const(block, &mut emitter)))
1593 .transpose()?;
1594
1595 let mut ectx = ctx.as_expression(block, &mut emitter);
1596
1597 let (_ty, initializer) = self.type_and_init(
1598 l.name,
1599 Some(l.init),
1600 explicit_ty,
1601 AbstractRule::Concretize,
1602 &mut ectx,
1603 )?;
1604
1605 let initializer =
1608 initializer.expect("type_and_init did not return an initializer");
1609
1610 ctx.local_expression_kind_tracker
1616 .force_non_const(initializer);
1617
1618 block.extend(emitter.finish(&ctx.function.expressions));
1619 ctx.local_table
1620 .insert(l.handle, Declared::Runtime(Typed::Plain(initializer)));
1621 ctx.named_expressions
1622 .insert(initializer, (l.name.name.to_string(), l.name.span));
1623
1624 return Ok(());
1625 }
1626 ast::LocalDecl::Var(ref v) => {
1627 let mut emitter = proc::Emitter::default();
1628 emitter.start(&ctx.function.expressions);
1629
1630 let explicit_ty =
1631 v.ty.map(|ast| {
1632 self.resolve_ast_type(ast, &mut ctx.as_const(block, &mut emitter))
1633 })
1634 .transpose()?;
1635
1636 let mut ectx = ctx.as_expression(block, &mut emitter);
1637 let (ty, initializer) = self.type_and_init(
1638 v.name,
1639 v.init,
1640 explicit_ty,
1641 AbstractRule::Concretize,
1642 &mut ectx,
1643 )?;
1644
1645 let (const_initializer, initializer) = {
1646 match initializer {
1647 Some(init) => {
1648 if is_inside_loop
1656 || !ctx.local_expression_kind_tracker.is_const_or_override(init)
1657 {
1658 (None, Some(init))
1659 } else {
1660 (Some(init), None)
1661 }
1662 }
1663 None => (None, None),
1664 }
1665 };
1666
1667 let var = ctx.function.local_variables.append(
1668 ir::LocalVariable {
1669 name: Some(v.name.name.to_string()),
1670 ty,
1671 init: const_initializer,
1672 },
1673 stmt.span,
1674 );
1675
1676 let handle = ctx
1677 .as_expression(block, &mut emitter)
1678 .interrupt_emitter(ir::Expression::LocalVariable(var), Span::UNDEFINED)?;
1679 block.extend(emitter.finish(&ctx.function.expressions));
1680 ctx.local_table
1681 .insert(v.handle, Declared::Runtime(Typed::Reference(handle)));
1682
1683 match initializer {
1684 Some(initializer) => ir::Statement::Store {
1685 pointer: handle,
1686 value: initializer,
1687 },
1688 None => return Ok(()),
1689 }
1690 }
1691 ast::LocalDecl::Const(ref c) => {
1692 let mut emitter = proc::Emitter::default();
1693 emitter.start(&ctx.function.expressions);
1694
1695 let ectx = &mut ctx.as_const(block, &mut emitter);
1696
1697 let explicit_ty =
1698 c.ty.map(|ast| self.resolve_ast_type(ast, &mut ectx.as_const()))
1699 .transpose()?;
1700
1701 let (_ty, init) = self.type_and_init(
1702 c.name,
1703 Some(c.init),
1704 explicit_ty,
1705 AbstractRule::Allow,
1706 &mut ectx.as_const(),
1707 )?;
1708 let init = init.expect("Local const must have init");
1709
1710 block.extend(emitter.finish(&ctx.function.expressions));
1711 ctx.local_table
1712 .insert(c.handle, Declared::Const(Typed::Plain(init)));
1713 return Ok(());
1714 }
1715 },
1716 ast::StatementKind::If {
1717 condition,
1718 ref accept,
1719 ref reject,
1720 } => {
1721 let mut emitter = proc::Emitter::default();
1722 emitter.start(&ctx.function.expressions);
1723
1724 let condition =
1725 self.expression(condition, &mut ctx.as_expression(block, &mut emitter))?;
1726 block.extend(emitter.finish(&ctx.function.expressions));
1727
1728 let accept = self.block(accept, is_inside_loop, ctx)?;
1729 let reject = self.block(reject, is_inside_loop, ctx)?;
1730
1731 ir::Statement::If {
1732 condition,
1733 accept,
1734 reject,
1735 }
1736 }
1737 ast::StatementKind::Switch {
1738 selector,
1739 ref cases,
1740 } => {
1741 let mut emitter = proc::Emitter::default();
1742 emitter.start(&ctx.function.expressions);
1743
1744 let mut ectx = ctx.as_expression(block, &mut emitter);
1745
1746 let (mut exprs, spans) = core::iter::once(selector)
1749 .chain(cases.iter().filter_map(|case| match case.value {
1750 ast::SwitchValue::Expr(expr) => Some(expr),
1751 ast::SwitchValue::Default => None,
1752 }))
1753 .enumerate()
1754 .map(|(i, expr)| {
1755 let span = ectx.ast_expressions.get_span(expr);
1756 let expr = self.expression_for_abstract(expr, &mut ectx)?;
1757 let ty = resolve_inner!(ectx, expr);
1758 match *ty {
1759 ir::TypeInner::Scalar(
1760 ir::Scalar::I32 | ir::Scalar::U32 | ir::Scalar::ABSTRACT_INT,
1761 ) => Ok((expr, span)),
1762 _ => match i {
1763 0 => Err(Box::new(Error::InvalidSwitchSelector { span })),
1764 _ => Err(Box::new(Error::InvalidSwitchCase { span })),
1765 },
1766 }
1767 })
1768 .collect::<Result<(Vec<_>, Vec<_>)>>()?;
1769
1770 let mut consensus =
1771 ectx.automatic_conversion_consensus(&exprs)
1772 .map_err(|span_idx| Error::SwitchCaseTypeMismatch {
1773 span: spans[span_idx],
1774 })?;
1775 if consensus == ir::Scalar::ABSTRACT_INT {
1777 consensus = ir::Scalar::I32;
1778 }
1779 for expr in &mut exprs {
1780 ectx.convert_to_leaf_scalar(expr, consensus)?;
1781 }
1782
1783 block.extend(emitter.finish(&ctx.function.expressions));
1784
1785 let mut exprs = exprs.into_iter();
1786 let selector = exprs
1787 .next()
1788 .expect("First element should be selector expression");
1789
1790 let cases = cases
1791 .iter()
1792 .map(|case| {
1793 Ok(ir::SwitchCase {
1794 value: match case.value {
1795 ast::SwitchValue::Expr(expr) => {
1796 let span = ctx.ast_expressions.get_span(expr);
1797 let expr = exprs.next().expect(
1798 "Should yield expression for each SwitchValue::Expr case",
1799 );
1800 match ctx
1801 .module
1802 .to_ctx()
1803 .eval_expr_to_literal_from(expr, &ctx.function.expressions)
1804 {
1805 Some(ir::Literal::I32(value)) => {
1806 ir::SwitchValue::I32(value)
1807 }
1808 Some(ir::Literal::U32(value)) => {
1809 ir::SwitchValue::U32(value)
1810 }
1811 _ => {
1812 return Err(Box::new(Error::InvalidSwitchCase {
1813 span,
1814 }));
1815 }
1816 }
1817 }
1818 ast::SwitchValue::Default => ir::SwitchValue::Default,
1819 },
1820 body: self.block(&case.body, is_inside_loop, ctx)?,
1821 fall_through: case.fall_through,
1822 })
1823 })
1824 .collect::<Result<_>>()?;
1825
1826 ir::Statement::Switch { selector, cases }
1827 }
1828 ast::StatementKind::Loop {
1829 ref body,
1830 ref continuing,
1831 break_if,
1832 } => {
1833 let body = self.block(body, true, ctx)?;
1834 let mut continuing = self.block(continuing, true, ctx)?;
1835
1836 let mut emitter = proc::Emitter::default();
1837 emitter.start(&ctx.function.expressions);
1838 let break_if = break_if
1839 .map(|expr| {
1840 self.expression(expr, &mut ctx.as_expression(&mut continuing, &mut emitter))
1841 })
1842 .transpose()?;
1843 continuing.extend(emitter.finish(&ctx.function.expressions));
1844
1845 ir::Statement::Loop {
1846 body,
1847 continuing,
1848 break_if,
1849 }
1850 }
1851 ast::StatementKind::Break => ir::Statement::Break,
1852 ast::StatementKind::Continue => ir::Statement::Continue,
1853 ast::StatementKind::Return { value: ast_value } => {
1854 let mut emitter = proc::Emitter::default();
1855 emitter.start(&ctx.function.expressions);
1856
1857 let value;
1858 if let Some(ast_expr) = ast_value {
1859 let result_ty = ctx.function.result.as_ref().map(|r| r.ty);
1860 let mut ectx = ctx.as_expression(block, &mut emitter);
1861 let expr = self.expression_for_abstract(ast_expr, &mut ectx)?;
1862
1863 if let Some(result_ty) = result_ty {
1864 let mut ectx = ctx.as_expression(block, &mut emitter);
1865 let resolution = proc::TypeResolution::Handle(result_ty);
1866 let converted =
1867 ectx.try_automatic_conversions(expr, &resolution, Span::default())?;
1868 value = Some(converted);
1869 } else {
1870 value = Some(expr);
1871 }
1872 } else {
1873 value = None;
1874 }
1875 block.extend(emitter.finish(&ctx.function.expressions));
1876
1877 ir::Statement::Return { value }
1878 }
1879 ast::StatementKind::Kill => ir::Statement::Kill,
1880 ast::StatementKind::Call {
1881 ref function,
1882 ref arguments,
1883 } => {
1884 let mut emitter = proc::Emitter::default();
1885 emitter.start(&ctx.function.expressions);
1886
1887 let _ = self.call(
1888 stmt.span,
1889 function,
1890 arguments,
1891 &mut ctx.as_expression(block, &mut emitter),
1892 true,
1893 )?;
1894 block.extend(emitter.finish(&ctx.function.expressions));
1895 return Ok(());
1896 }
1897 ast::StatementKind::Assign {
1898 target: ast_target,
1899 op,
1900 value,
1901 } => {
1902 let mut emitter = proc::Emitter::default();
1903 emitter.start(&ctx.function.expressions);
1904 let target_span = ctx.ast_expressions.get_span(ast_target);
1905
1906 let mut ectx = ctx.as_expression(block, &mut emitter);
1907 let target = self.expression_for_reference(ast_target, &mut ectx)?;
1908 let target_handle = match target {
1909 Typed::Reference(handle) => handle,
1910 Typed::Plain(handle) => {
1911 let ty = ctx.invalid_assignment_type(handle);
1912 return Err(Box::new(Error::InvalidAssignment {
1913 span: target_span,
1914 ty,
1915 }));
1916 }
1917 };
1918
1919 let target_scalar = match op {
1924 Some(ir::BinaryOperator::ShiftLeft | ir::BinaryOperator::ShiftRight) => {
1925 Some(ir::Scalar::U32)
1926 }
1927 _ => resolve_inner!(ectx, target_handle)
1928 .pointer_automatically_convertible_scalar(&ectx.module.types),
1929 };
1930
1931 let value = self.expression_for_abstract(value, &mut ectx)?;
1932 let mut value = match target_scalar {
1933 Some(target_scalar) => ectx.try_automatic_conversion_for_leaf_scalar(
1934 value,
1935 target_scalar,
1936 target_span,
1937 )?,
1938 None => value,
1939 };
1940
1941 let value = match op {
1942 Some(op) => {
1943 let mut left = ectx.apply_load_rule(target)?;
1944 ectx.binary_op_splat(op, &mut left, &mut value)?;
1945 ectx.append_expression(
1946 ir::Expression::Binary {
1947 op,
1948 left,
1949 right: value,
1950 },
1951 stmt.span,
1952 )?
1953 }
1954 None => value,
1955 };
1956 block.extend(emitter.finish(&ctx.function.expressions));
1957
1958 ir::Statement::Store {
1959 pointer: target_handle,
1960 value,
1961 }
1962 }
1963 ast::StatementKind::Increment(value) | ast::StatementKind::Decrement(value) => {
1964 let mut emitter = proc::Emitter::default();
1965 emitter.start(&ctx.function.expressions);
1966
1967 let op = match stmt.kind {
1968 ast::StatementKind::Increment(_) => ir::BinaryOperator::Add,
1969 ast::StatementKind::Decrement(_) => ir::BinaryOperator::Subtract,
1970 _ => unreachable!(),
1971 };
1972
1973 let value_span = ctx.ast_expressions.get_span(value);
1974 let target = self
1975 .expression_for_reference(value, &mut ctx.as_expression(block, &mut emitter))?;
1976 let target_handle = match target {
1977 Typed::Reference(handle) => handle,
1978 Typed::Plain(_) => {
1979 return Err(Box::new(Error::BadIncrDecrReferenceType(value_span)))
1980 }
1981 };
1982
1983 let mut ectx = ctx.as_expression(block, &mut emitter);
1984 let scalar = match *resolve_inner!(ectx, target_handle) {
1985 ir::TypeInner::ValuePointer {
1986 size: None, scalar, ..
1987 } => scalar,
1988 ir::TypeInner::Pointer { base, .. } => match ectx.module.types[base].inner {
1989 ir::TypeInner::Scalar(scalar) => scalar,
1990 _ => return Err(Box::new(Error::BadIncrDecrReferenceType(value_span))),
1991 },
1992 _ => return Err(Box::new(Error::BadIncrDecrReferenceType(value_span))),
1993 };
1994 let literal = match scalar.kind {
1995 ir::ScalarKind::Sint | ir::ScalarKind::Uint => ir::Literal::one(scalar)
1996 .ok_or(Error::BadIncrDecrReferenceType(value_span))?,
1997 _ => return Err(Box::new(Error::BadIncrDecrReferenceType(value_span))),
1998 };
1999
2000 let right =
2001 ectx.interrupt_emitter(ir::Expression::Literal(literal), Span::UNDEFINED)?;
2002 let rctx = ectx.runtime_expression_ctx(stmt.span)?;
2003 let left = rctx.function.expressions.append(
2004 ir::Expression::Load {
2005 pointer: target_handle,
2006 },
2007 value_span,
2008 );
2009 let value = rctx
2010 .function
2011 .expressions
2012 .append(ir::Expression::Binary { op, left, right }, stmt.span);
2013 rctx.local_expression_kind_tracker
2014 .insert(left, proc::ExpressionKind::Runtime);
2015 rctx.local_expression_kind_tracker
2016 .insert(value, proc::ExpressionKind::Runtime);
2017
2018 block.extend(emitter.finish(&ctx.function.expressions));
2019 ir::Statement::Store {
2020 pointer: target_handle,
2021 value,
2022 }
2023 }
2024 ast::StatementKind::ConstAssert(condition) => {
2025 let mut emitter = proc::Emitter::default();
2026 emitter.start(&ctx.function.expressions);
2027
2028 let condition =
2029 self.expression(condition, &mut ctx.as_const(block, &mut emitter))?;
2030
2031 let span = ctx.function.expressions.get_span(condition);
2032 match ctx
2033 .module
2034 .to_ctx()
2035 .eval_expr_to_bool_from(condition, &ctx.function.expressions)
2036 {
2037 Some(true) => Ok(()),
2038 Some(false) => Err(Error::ConstAssertFailed(span)),
2039 _ => Err(Error::NotBool(span)),
2040 }?;
2041
2042 block.extend(emitter.finish(&ctx.function.expressions));
2043
2044 return Ok(());
2045 }
2046 ast::StatementKind::Phony(expr) => {
2047 let mut emitter = proc::Emitter::default();
2051 emitter.start(&ctx.function.expressions);
2052
2053 let value = self.expression(expr, &mut ctx.as_expression(block, &mut emitter))?;
2054 block.extend(emitter.finish(&ctx.function.expressions));
2055 ctx.named_expressions
2056 .insert(value, ("phony".to_string(), stmt.span));
2057 return Ok(());
2058 }
2059 };
2060
2061 block.push(out, stmt.span);
2062
2063 Ok(())
2064 }
2065
2066 fn expression(
2074 &mut self,
2075 expr: Handle<ast::Expression<'source>>,
2076 ctx: &mut ExpressionContext<'source, '_, '_>,
2077 ) -> Result<'source, Handle<ir::Expression>> {
2078 let expr = self.expression_for_abstract(expr, ctx)?;
2079 ctx.concretize(expr)
2080 }
2081
2082 fn expression_for_abstract(
2083 &mut self,
2084 expr: Handle<ast::Expression<'source>>,
2085 ctx: &mut ExpressionContext<'source, '_, '_>,
2086 ) -> Result<'source, Handle<ir::Expression>> {
2087 let expr = self.expression_for_reference(expr, ctx)?;
2088 ctx.apply_load_rule(expr)
2089 }
2090
2091 fn expression_with_leaf_scalar(
2092 &mut self,
2093 expr: Handle<ast::Expression<'source>>,
2094 scalar: ir::Scalar,
2095 ctx: &mut ExpressionContext<'source, '_, '_>,
2096 ) -> Result<'source, Handle<ir::Expression>> {
2097 let unconverted = self.expression_for_abstract(expr, ctx)?;
2098 ctx.try_automatic_conversion_for_leaf_scalar(unconverted, scalar, Span::default())
2099 }
2100
2101 fn expression_for_reference(
2102 &mut self,
2103 expr: Handle<ast::Expression<'source>>,
2104 ctx: &mut ExpressionContext<'source, '_, '_>,
2105 ) -> Result<'source, Typed<Handle<ir::Expression>>> {
2106 let span = ctx.ast_expressions.get_span(expr);
2107 let expr = &ctx.ast_expressions[expr];
2108
2109 let expr: Typed<ir::Expression> = match *expr {
2110 ast::Expression::Literal(literal) => {
2111 let literal = match literal {
2112 ast::Literal::Number(Number::F16(f)) => ir::Literal::F16(f),
2113 ast::Literal::Number(Number::F32(f)) => ir::Literal::F32(f),
2114 ast::Literal::Number(Number::I32(i)) => ir::Literal::I32(i),
2115 ast::Literal::Number(Number::U32(u)) => ir::Literal::U32(u),
2116 ast::Literal::Number(Number::I64(i)) => ir::Literal::I64(i),
2117 ast::Literal::Number(Number::U64(u)) => ir::Literal::U64(u),
2118 ast::Literal::Number(Number::F64(f)) => ir::Literal::F64(f),
2119 ast::Literal::Number(Number::AbstractInt(i)) => ir::Literal::AbstractInt(i),
2120 ast::Literal::Number(Number::AbstractFloat(f)) => ir::Literal::AbstractFloat(f),
2121 ast::Literal::Bool(b) => ir::Literal::Bool(b),
2122 };
2123 let handle = ctx.interrupt_emitter(ir::Expression::Literal(literal), span)?;
2124 return Ok(Typed::Plain(handle));
2125 }
2126 ast::Expression::Ident(ast::IdentExpr::Local(local)) => {
2127 return ctx.local(&local, span);
2128 }
2129 ast::Expression::Ident(ast::IdentExpr::Unresolved(name)) => {
2130 let global = ctx
2131 .globals
2132 .get(name)
2133 .ok_or(Error::UnknownIdent(span, name))?;
2134 let expr = match *global {
2135 LoweredGlobalDecl::Var(handle) => {
2136 let expr = ir::Expression::GlobalVariable(handle);
2137 match ctx.module.global_variables[handle].space {
2138 ir::AddressSpace::Handle => Typed::Plain(expr),
2139 _ => Typed::Reference(expr),
2140 }
2141 }
2142 LoweredGlobalDecl::Const(handle) => {
2143 Typed::Plain(ir::Expression::Constant(handle))
2144 }
2145 LoweredGlobalDecl::Override(handle) => {
2146 Typed::Plain(ir::Expression::Override(handle))
2147 }
2148 LoweredGlobalDecl::Function { .. }
2149 | LoweredGlobalDecl::Type(_)
2150 | LoweredGlobalDecl::EntryPoint(_) => {
2151 return Err(Box::new(Error::Unexpected(span, ExpectedToken::Variable)));
2152 }
2153 };
2154
2155 return expr.try_map(|handle| ctx.interrupt_emitter(handle, span));
2156 }
2157 ast::Expression::Construct {
2158 ref ty,
2159 ty_span,
2160 ref components,
2161 } => {
2162 let handle = self.construct(span, ty, ty_span, components, ctx)?;
2163 return Ok(Typed::Plain(handle));
2164 }
2165 ast::Expression::Unary { op, expr } => {
2166 let expr = self.expression_for_abstract(expr, ctx)?;
2167 Typed::Plain(ir::Expression::Unary { op, expr })
2168 }
2169 ast::Expression::AddrOf(expr) => {
2170 match self.expression_for_reference(expr, ctx)? {
2173 Typed::Reference(handle) => {
2174 let expr = &ctx.runtime_expression_ctx(span)?.function.expressions[handle];
2175 if let &ir::Expression::Access { base, .. }
2176 | &ir::Expression::AccessIndex { base, .. } = expr
2177 {
2178 if let Some(ty) = resolve_inner!(ctx, base).pointer_base_type() {
2179 if matches!(
2180 *ty.inner_with(&ctx.module.types),
2181 ir::TypeInner::Vector { .. },
2182 ) {
2183 return Err(Box::new(Error::InvalidAddrOfOperand(
2184 ctx.get_expression_span(handle),
2185 )));
2186 }
2187 }
2188 }
2189 return Ok(Typed::Plain(handle));
2191 }
2192 Typed::Plain(_) => {
2193 return Err(Box::new(Error::NotReference(
2194 "the operand of the `&` operator",
2195 span,
2196 )));
2197 }
2198 }
2199 }
2200 ast::Expression::Deref(expr) => {
2201 let pointer = self.expression(expr, ctx)?;
2203
2204 if resolve_inner!(ctx, pointer).pointer_space().is_none() {
2205 return Err(Box::new(Error::NotPointer(span)));
2206 }
2207
2208 return Ok(Typed::Reference(pointer));
2210 }
2211 ast::Expression::Binary { op, left, right } => {
2212 self.binary(op, left, right, span, ctx)?
2213 }
2214 ast::Expression::Call {
2215 ref function,
2216 ref arguments,
2217 } => {
2218 let handle = self
2219 .call(span, function, arguments, ctx, false)?
2220 .ok_or(Error::FunctionReturnsVoid(function.span))?;
2221 return Ok(Typed::Plain(handle));
2222 }
2223 ast::Expression::Index { base, index } => {
2224 let mut lowered_base = self.expression_for_reference(base, ctx)?;
2225 let index = self.expression(index, ctx)?;
2226
2227 if let Typed::Plain(handle) = lowered_base {
2230 if resolve_inner!(ctx, handle).pointer_space().is_some() {
2231 lowered_base = Typed::Reference(handle);
2232 }
2233 }
2234
2235 lowered_base.try_map(|base| match ctx.const_eval_expr_to_u32(index).ok() {
2236 Some(index) => Ok::<_, Box<Error>>(ir::Expression::AccessIndex { base, index }),
2237 None => {
2238 let base = ctx.concretize(base)?;
2244 Ok(ir::Expression::Access { base, index })
2245 }
2246 })?
2247 }
2248 ast::Expression::Member { base, ref field } => {
2249 let mut lowered_base = self.expression_for_reference(base, ctx)?;
2250
2251 if let Typed::Plain(handle) = lowered_base {
2254 if resolve_inner!(ctx, handle).pointer_space().is_some() {
2255 lowered_base = Typed::Reference(handle);
2256 }
2257 }
2258
2259 let temp_ty;
2260 let composite_type: &ir::TypeInner = match lowered_base {
2261 Typed::Reference(handle) => {
2262 temp_ty = resolve_inner!(ctx, handle)
2263 .pointer_base_type()
2264 .expect("In Typed::Reference(handle), handle must be a Naga pointer");
2265 temp_ty.inner_with(&ctx.module.types)
2266 }
2267
2268 Typed::Plain(handle) => {
2269 resolve_inner!(ctx, handle)
2270 }
2271 };
2272
2273 let access = match *composite_type {
2274 ir::TypeInner::Struct { ref members, .. } => {
2275 let index = members
2276 .iter()
2277 .position(|m| m.name.as_deref() == Some(field.name))
2278 .ok_or(Error::BadAccessor(field.span))?
2279 as u32;
2280
2281 lowered_base.map(|base| ir::Expression::AccessIndex { base, index })
2282 }
2283 ir::TypeInner::Vector { .. } => {
2284 match Components::new(field.name, field.span)? {
2285 Components::Swizzle { size, pattern } => {
2286 Typed::Plain(ir::Expression::Swizzle {
2287 size,
2288 vector: ctx.apply_load_rule(lowered_base)?,
2289 pattern,
2290 })
2291 }
2292 Components::Single(index) => {
2293 lowered_base.map(|base| ir::Expression::AccessIndex { base, index })
2294 }
2295 }
2296 }
2297 _ => return Err(Box::new(Error::BadAccessor(field.span))),
2298 };
2299
2300 access
2301 }
2302 ast::Expression::Bitcast { expr, to, ty_span } => {
2303 let expr = self.expression(expr, ctx)?;
2304 let to_resolved = self.resolve_ast_type(to, &mut ctx.as_const())?;
2305
2306 let element_scalar = match ctx.module.types[to_resolved].inner {
2307 ir::TypeInner::Scalar(scalar) => scalar,
2308 ir::TypeInner::Vector { scalar, .. } => scalar,
2309 _ => {
2310 let ty = resolve!(ctx, expr);
2311 return Err(Box::new(Error::BadTypeCast {
2312 from_type: ctx.type_resolution_to_string(ty),
2313 span: ty_span,
2314 to_type: ctx.type_to_string(to_resolved),
2315 }));
2316 }
2317 };
2318
2319 Typed::Plain(ir::Expression::As {
2320 expr,
2321 kind: element_scalar.kind,
2322 convert: None,
2323 })
2324 }
2325 };
2326
2327 expr.try_map(|handle| ctx.append_expression(handle, span))
2328 }
2329
2330 fn binary(
2331 &mut self,
2332 op: ir::BinaryOperator,
2333 left: Handle<ast::Expression<'source>>,
2334 right: Handle<ast::Expression<'source>>,
2335 span: Span,
2336 ctx: &mut ExpressionContext<'source, '_, '_>,
2337 ) -> Result<'source, Typed<ir::Expression>> {
2338 let mut left = self.expression_for_abstract(left, ctx)?;
2340 let mut right = self.expression_for_abstract(right, ctx)?;
2341
2342 ctx.binary_op_splat(op, &mut left, &mut right)?;
2345
2346 match op {
2348 ir::BinaryOperator::ShiftLeft | ir::BinaryOperator::ShiftRight => {
2349 right =
2354 ctx.try_automatic_conversion_for_leaf_scalar(right, ir::Scalar::U32, span)?;
2355
2356 if !ctx.is_const(right) {
2368 left = ctx.concretize(left)?;
2369 }
2370 }
2371
2372 _ => {
2377 ctx.grow_types(left)?;
2378 ctx.grow_types(right)?;
2379 if let Ok(consensus_scalar) =
2380 ctx.automatic_conversion_consensus([left, right].iter())
2381 {
2382 ctx.convert_to_leaf_scalar(&mut left, consensus_scalar)?;
2383 ctx.convert_to_leaf_scalar(&mut right, consensus_scalar)?;
2384 }
2385 }
2386 }
2387
2388 Ok(Typed::Plain(ir::Expression::Binary { op, left, right }))
2389 }
2390
2391 fn call(
2410 &mut self,
2411 span: Span,
2412 function: &ast::Ident<'source>,
2413 arguments: &[Handle<ast::Expression<'source>>],
2414 ctx: &mut ExpressionContext<'source, '_, '_>,
2415 is_statement: bool,
2416 ) -> Result<'source, Option<Handle<ir::Expression>>> {
2417 let function_span = function.span;
2418 match ctx.globals.get(function.name) {
2419 Some(&LoweredGlobalDecl::Type(ty)) => {
2420 let handle = self.construct(
2421 span,
2422 &ast::ConstructorType::Type(ty),
2423 function_span,
2424 arguments,
2425 ctx,
2426 )?;
2427 Ok(Some(handle))
2428 }
2429 Some(
2430 &LoweredGlobalDecl::Const(_)
2431 | &LoweredGlobalDecl::Override(_)
2432 | &LoweredGlobalDecl::Var(_),
2433 ) => Err(Box::new(Error::Unexpected(
2434 function_span,
2435 ExpectedToken::Function,
2436 ))),
2437 Some(&LoweredGlobalDecl::EntryPoint(_)) => {
2438 Err(Box::new(Error::CalledEntryPoint(function_span)))
2439 }
2440 Some(&LoweredGlobalDecl::Function {
2441 handle: function,
2442 must_use,
2443 }) => {
2444 let arguments = arguments
2445 .iter()
2446 .enumerate()
2447 .map(|(i, &arg)| {
2448 let Some(&ir::FunctionArgument {
2450 ty: parameter_ty, ..
2451 }) = ctx.module.functions[function].arguments.get(i)
2452 else {
2453 return self.expression(arg, ctx);
2456 };
2457
2458 let expr = self.expression_for_abstract(arg, ctx)?;
2459 ctx.try_automatic_conversions(
2460 expr,
2461 &proc::TypeResolution::Handle(parameter_ty),
2462 ctx.ast_expressions.get_span(arg),
2463 )
2464 })
2465 .collect::<Result<Vec<_>>>()?;
2466
2467 let has_result = ctx.module.functions[function].result.is_some();
2468
2469 if must_use && is_statement {
2470 return Err(Box::new(Error::FunctionMustUseUnused(function_span)));
2471 }
2472
2473 let rctx = ctx.runtime_expression_ctx(span)?;
2474 rctx.block
2476 .extend(rctx.emitter.finish(&rctx.function.expressions));
2477 let result = has_result.then(|| {
2478 let result = rctx
2479 .function
2480 .expressions
2481 .append(ir::Expression::CallResult(function), span);
2482 rctx.local_expression_kind_tracker
2483 .insert(result, proc::ExpressionKind::Runtime);
2484 result
2485 });
2486 rctx.emitter.start(&rctx.function.expressions);
2487 rctx.block.push(
2488 ir::Statement::Call {
2489 function,
2490 arguments,
2491 result,
2492 },
2493 span,
2494 );
2495
2496 Ok(result)
2497 }
2498 None => {
2499 let span = function_span;
2500 let expr = if let Some(fun) = conv::map_relational_fun(function.name) {
2501 let mut args = ctx.prepare_args(arguments, 1, span);
2502 let argument = self.expression(args.next()?, ctx)?;
2503 args.finish()?;
2504
2505 let argument_unmodified = matches!(
2507 fun,
2508 ir::RelationalFunction::All | ir::RelationalFunction::Any
2509 ) && {
2510 matches!(
2511 resolve_inner!(ctx, argument),
2512 &ir::TypeInner::Scalar(ir::Scalar {
2513 kind: ir::ScalarKind::Bool,
2514 ..
2515 })
2516 )
2517 };
2518
2519 if argument_unmodified {
2520 return Ok(Some(argument));
2521 } else {
2522 ir::Expression::Relational { fun, argument }
2523 }
2524 } else if let Some((axis, ctrl)) = conv::map_derivative(function.name) {
2525 let mut args = ctx.prepare_args(arguments, 1, span);
2526 let expr = self.expression(args.next()?, ctx)?;
2527 args.finish()?;
2528
2529 ir::Expression::Derivative { axis, ctrl, expr }
2530 } else if let Some(fun) = conv::map_standard_fun(function.name) {
2531 self.math_function_helper(span, fun, arguments, ctx)?
2532 } else if let Some(fun) = Texture::map(function.name) {
2533 self.texture_sample_helper(fun, arguments, span, ctx)?
2534 } else if let Some((op, cop)) = conv::map_subgroup_operation(function.name) {
2535 return Ok(Some(
2536 self.subgroup_operation_helper(span, op, cop, arguments, ctx)?,
2537 ));
2538 } else if let Some(mode) = SubgroupGather::map(function.name) {
2539 return Ok(Some(
2540 self.subgroup_gather_helper(span, mode, arguments, ctx)?,
2541 ));
2542 } else if let Some(fun) = ir::AtomicFunction::map(function.name) {
2543 return self.atomic_helper(span, fun, arguments, is_statement, ctx);
2544 } else {
2545 match function.name {
2546 "select" => {
2547 let mut args = ctx.prepare_args(arguments, 3, span);
2548
2549 let reject_orig = args.next()?;
2550 let accept_orig = args.next()?;
2551 let mut values = [
2552 self.expression_for_abstract(reject_orig, ctx)?,
2553 self.expression_for_abstract(accept_orig, ctx)?,
2554 ];
2555 let condition = self.expression(args.next()?, ctx)?;
2556
2557 args.finish()?;
2558
2559 let diagnostic_details =
2560 |ctx: &ExpressionContext<'_, '_, '_>,
2561 ty_res: &proc::TypeResolution,
2562 orig_expr| {
2563 (
2564 ctx.ast_expressions.get_span(orig_expr),
2565 format!("`{}`", ctx.as_diagnostic_display(ty_res)),
2566 )
2567 };
2568 for (&value, orig_value) in
2569 values.iter().zip([reject_orig, accept_orig])
2570 {
2571 let value_ty_res = resolve!(ctx, value);
2572 if value_ty_res
2573 .inner_with(&ctx.module.types)
2574 .vector_size_and_scalar()
2575 .is_none()
2576 {
2577 let (arg_span, arg_type) =
2578 diagnostic_details(ctx, value_ty_res, orig_value);
2579 return Err(Box::new(Error::SelectUnexpectedArgumentType {
2580 arg_span,
2581 arg_type,
2582 }));
2583 }
2584 }
2585 let mut consensus_scalar = ctx
2586 .automatic_conversion_consensus(&values)
2587 .map_err(|_idx| {
2588 let [reject, accept] = values;
2589 let [(reject_span, reject_type), (accept_span, accept_type)] =
2590 [(reject_orig, reject), (accept_orig, accept)].map(
2591 |(orig_expr, expr)| {
2592 let ty_res = &ctx.typifier()[expr];
2593 diagnostic_details(ctx, ty_res, orig_expr)
2594 },
2595 );
2596 Error::SelectRejectAndAcceptHaveNoCommonType {
2597 reject_span,
2598 reject_type,
2599 accept_span,
2600 accept_type,
2601 }
2602 })?;
2603 if !ctx.is_const(condition) {
2604 consensus_scalar = consensus_scalar.concretize();
2605 }
2606
2607 ctx.convert_slice_to_common_leaf_scalar(&mut values, consensus_scalar)?;
2608
2609 let [reject, accept] = values;
2610
2611 ir::Expression::Select {
2612 reject,
2613 accept,
2614 condition,
2615 }
2616 }
2617 "arrayLength" => {
2618 let mut args = ctx.prepare_args(arguments, 1, span);
2619 let expr = self.expression(args.next()?, ctx)?;
2620 args.finish()?;
2621
2622 ir::Expression::ArrayLength(expr)
2623 }
2624 "atomicLoad" => {
2625 let mut args = ctx.prepare_args(arguments, 1, span);
2626 let (pointer, _scalar) = self.atomic_pointer(args.next()?, ctx)?;
2627 args.finish()?;
2628
2629 ir::Expression::Load { pointer }
2630 }
2631 "atomicStore" => {
2632 let mut args = ctx.prepare_args(arguments, 2, span);
2633 let (pointer, scalar) = self.atomic_pointer(args.next()?, ctx)?;
2634 let value =
2635 self.expression_with_leaf_scalar(args.next()?, scalar, ctx)?;
2636 args.finish()?;
2637
2638 let rctx = ctx.runtime_expression_ctx(span)?;
2639 rctx.block
2640 .extend(rctx.emitter.finish(&rctx.function.expressions));
2641 rctx.emitter.start(&rctx.function.expressions);
2642 rctx.block
2643 .push(ir::Statement::Store { pointer, value }, span);
2644 return Ok(None);
2645 }
2646 "atomicCompareExchangeWeak" => {
2647 let mut args = ctx.prepare_args(arguments, 3, span);
2648
2649 let (pointer, scalar) = self.atomic_pointer(args.next()?, ctx)?;
2650
2651 let compare =
2652 self.expression_with_leaf_scalar(args.next()?, scalar, ctx)?;
2653
2654 let value = args.next()?;
2655 let value_span = ctx.ast_expressions.get_span(value);
2656 let value = self.expression_with_leaf_scalar(value, scalar, ctx)?;
2657
2658 args.finish()?;
2659
2660 let expression = match *resolve_inner!(ctx, value) {
2661 ir::TypeInner::Scalar(scalar) => ir::Expression::AtomicResult {
2662 ty: ctx.module.generate_predeclared_type(
2663 ir::PredeclaredType::AtomicCompareExchangeWeakResult(
2664 scalar,
2665 ),
2666 ),
2667 comparison: true,
2668 },
2669 _ => {
2670 return Err(Box::new(Error::InvalidAtomicOperandType(
2671 value_span,
2672 )))
2673 }
2674 };
2675
2676 let result = ctx.interrupt_emitter(expression, span)?;
2677 let rctx = ctx.runtime_expression_ctx(span)?;
2678 rctx.block.push(
2679 ir::Statement::Atomic {
2680 pointer,
2681 fun: ir::AtomicFunction::Exchange {
2682 compare: Some(compare),
2683 },
2684 value,
2685 result: Some(result),
2686 },
2687 span,
2688 );
2689 return Ok(Some(result));
2690 }
2691 "textureAtomicMin" | "textureAtomicMax" | "textureAtomicAdd"
2692 | "textureAtomicAnd" | "textureAtomicOr" | "textureAtomicXor" => {
2693 let mut args = ctx.prepare_args(arguments, 3, span);
2694
2695 let image = args.next()?;
2696 let image_span = ctx.ast_expressions.get_span(image);
2697 let image = self.expression(image, ctx)?;
2698
2699 let coordinate = self.expression(args.next()?, ctx)?;
2700
2701 let (_, arrayed) = ctx.image_data(image, image_span)?;
2702 let array_index = arrayed
2703 .then(|| {
2704 args.min_args += 1;
2705 self.expression(args.next()?, ctx)
2706 })
2707 .transpose()?;
2708
2709 let value = self.expression(args.next()?, ctx)?;
2710
2711 args.finish()?;
2712
2713 let rctx = ctx.runtime_expression_ctx(span)?;
2714 rctx.block
2715 .extend(rctx.emitter.finish(&rctx.function.expressions));
2716 rctx.emitter.start(&rctx.function.expressions);
2717 let stmt = ir::Statement::ImageAtomic {
2718 image,
2719 coordinate,
2720 array_index,
2721 fun: match function.name {
2722 "textureAtomicMin" => ir::AtomicFunction::Min,
2723 "textureAtomicMax" => ir::AtomicFunction::Max,
2724 "textureAtomicAdd" => ir::AtomicFunction::Add,
2725 "textureAtomicAnd" => ir::AtomicFunction::And,
2726 "textureAtomicOr" => ir::AtomicFunction::InclusiveOr,
2727 "textureAtomicXor" => ir::AtomicFunction::ExclusiveOr,
2728 _ => unreachable!(),
2729 },
2730 value,
2731 };
2732 rctx.block.push(stmt, span);
2733 return Ok(None);
2734 }
2735 "storageBarrier" => {
2736 ctx.prepare_args(arguments, 0, span).finish()?;
2737
2738 let rctx = ctx.runtime_expression_ctx(span)?;
2739 rctx.block
2740 .push(ir::Statement::ControlBarrier(ir::Barrier::STORAGE), span);
2741 return Ok(None);
2742 }
2743 "workgroupBarrier" => {
2744 ctx.prepare_args(arguments, 0, span).finish()?;
2745
2746 let rctx = ctx.runtime_expression_ctx(span)?;
2747 rctx.block
2748 .push(ir::Statement::ControlBarrier(ir::Barrier::WORK_GROUP), span);
2749 return Ok(None);
2750 }
2751 "subgroupBarrier" => {
2752 ctx.prepare_args(arguments, 0, span).finish()?;
2753
2754 let rctx = ctx.runtime_expression_ctx(span)?;
2755 rctx.block
2756 .push(ir::Statement::ControlBarrier(ir::Barrier::SUB_GROUP), span);
2757 return Ok(None);
2758 }
2759 "textureBarrier" => {
2760 ctx.prepare_args(arguments, 0, span).finish()?;
2761
2762 let rctx = ctx.runtime_expression_ctx(span)?;
2763 rctx.block
2764 .push(ir::Statement::ControlBarrier(ir::Barrier::TEXTURE), span);
2765 return Ok(None);
2766 }
2767 "workgroupUniformLoad" => {
2768 let mut args = ctx.prepare_args(arguments, 1, span);
2769 let expr = args.next()?;
2770 args.finish()?;
2771
2772 let pointer = self.expression(expr, ctx)?;
2773 let result_ty = match *resolve_inner!(ctx, pointer) {
2774 ir::TypeInner::Pointer {
2775 base,
2776 space: ir::AddressSpace::WorkGroup,
2777 } => base,
2778 ref other => {
2779 log::error!("Type {other:?} passed to workgroupUniformLoad");
2780 let span = ctx.ast_expressions.get_span(expr);
2781 return Err(Box::new(Error::InvalidWorkGroupUniformLoad(span)));
2782 }
2783 };
2784 let result = ctx.interrupt_emitter(
2785 ir::Expression::WorkGroupUniformLoadResult { ty: result_ty },
2786 span,
2787 )?;
2788 let rctx = ctx.runtime_expression_ctx(span)?;
2789 rctx.block.push(
2790 ir::Statement::WorkGroupUniformLoad { pointer, result },
2791 span,
2792 );
2793
2794 return Ok(Some(result));
2795 }
2796 "textureStore" => {
2797 let mut args = ctx.prepare_args(arguments, 3, span);
2798
2799 let image = args.next()?;
2800 let image_span = ctx.ast_expressions.get_span(image);
2801 let image = self.expression(image, ctx)?;
2802
2803 let coordinate = self.expression(args.next()?, ctx)?;
2804
2805 let (class, arrayed) = ctx.image_data(image, image_span)?;
2806 let array_index = arrayed
2807 .then(|| {
2808 args.min_args += 1;
2809 self.expression(args.next()?, ctx)
2810 })
2811 .transpose()?;
2812 let scalar = if let ir::ImageClass::Storage { format, .. } = class {
2813 format.into()
2814 } else {
2815 return Err(Box::new(Error::NotStorageTexture(image_span)));
2816 };
2817
2818 let value =
2819 self.expression_with_leaf_scalar(args.next()?, scalar, ctx)?;
2820
2821 args.finish()?;
2822
2823 let rctx = ctx.runtime_expression_ctx(span)?;
2824 rctx.block
2825 .extend(rctx.emitter.finish(&rctx.function.expressions));
2826 rctx.emitter.start(&rctx.function.expressions);
2827 let stmt = ir::Statement::ImageStore {
2828 image,
2829 coordinate,
2830 array_index,
2831 value,
2832 };
2833 rctx.block.push(stmt, span);
2834 return Ok(None);
2835 }
2836 "textureLoad" => {
2837 let mut args = ctx.prepare_args(arguments, 2, span);
2838
2839 let image = args.next()?;
2840 let image_span = ctx.ast_expressions.get_span(image);
2841 let image = self.expression(image, ctx)?;
2842
2843 let coordinate = self.expression(args.next()?, ctx)?;
2844
2845 let (class, arrayed) = ctx.image_data(image, image_span)?;
2846 let array_index = arrayed
2847 .then(|| {
2848 args.min_args += 1;
2849 self.expression(args.next()?, ctx)
2850 })
2851 .transpose()?;
2852
2853 let level = class
2854 .is_mipmapped()
2855 .then(|| {
2856 args.min_args += 1;
2857 self.expression(args.next()?, ctx)
2858 })
2859 .transpose()?;
2860
2861 let sample = class
2862 .is_multisampled()
2863 .then(|| self.expression(args.next()?, ctx))
2864 .transpose()?;
2865
2866 args.finish()?;
2867
2868 ir::Expression::ImageLoad {
2869 image,
2870 coordinate,
2871 array_index,
2872 level,
2873 sample,
2874 }
2875 }
2876 "textureDimensions" => {
2877 let mut args = ctx.prepare_args(arguments, 1, span);
2878 let image = self.expression(args.next()?, ctx)?;
2879 let level = args
2880 .next()
2881 .map(|arg| self.expression(arg, ctx))
2882 .ok()
2883 .transpose()?;
2884 args.finish()?;
2885
2886 ir::Expression::ImageQuery {
2887 image,
2888 query: ir::ImageQuery::Size { level },
2889 }
2890 }
2891 "textureNumLevels" => {
2892 let mut args = ctx.prepare_args(arguments, 1, span);
2893 let image = self.expression(args.next()?, ctx)?;
2894 args.finish()?;
2895
2896 ir::Expression::ImageQuery {
2897 image,
2898 query: ir::ImageQuery::NumLevels,
2899 }
2900 }
2901 "textureNumLayers" => {
2902 let mut args = ctx.prepare_args(arguments, 1, span);
2903 let image = self.expression(args.next()?, ctx)?;
2904 args.finish()?;
2905
2906 ir::Expression::ImageQuery {
2907 image,
2908 query: ir::ImageQuery::NumLayers,
2909 }
2910 }
2911 "textureNumSamples" => {
2912 let mut args = ctx.prepare_args(arguments, 1, span);
2913 let image = self.expression(args.next()?, ctx)?;
2914 args.finish()?;
2915
2916 ir::Expression::ImageQuery {
2917 image,
2918 query: ir::ImageQuery::NumSamples,
2919 }
2920 }
2921 "rayQueryInitialize" => {
2922 let mut args = ctx.prepare_args(arguments, 3, span);
2923 let query = self.ray_query_pointer(args.next()?, ctx)?;
2924 let acceleration_structure = self.expression(args.next()?, ctx)?;
2925 let descriptor = self.expression(args.next()?, ctx)?;
2926 args.finish()?;
2927
2928 let _ = ctx.module.generate_ray_desc_type();
2929 let fun = ir::RayQueryFunction::Initialize {
2930 acceleration_structure,
2931 descriptor,
2932 };
2933
2934 let rctx = ctx.runtime_expression_ctx(span)?;
2935 rctx.block
2936 .extend(rctx.emitter.finish(&rctx.function.expressions));
2937 rctx.emitter.start(&rctx.function.expressions);
2938 rctx.block
2939 .push(ir::Statement::RayQuery { query, fun }, span);
2940 return Ok(None);
2941 }
2942 "getCommittedHitVertexPositions" => {
2943 let mut args = ctx.prepare_args(arguments, 1, span);
2944 let query = self.ray_query_pointer(args.next()?, ctx)?;
2945 args.finish()?;
2946
2947 let _ = ctx.module.generate_vertex_return_type();
2948
2949 ir::Expression::RayQueryVertexPositions {
2950 query,
2951 committed: true,
2952 }
2953 }
2954 "getCandidateHitVertexPositions" => {
2955 let mut args = ctx.prepare_args(arguments, 1, span);
2956 let query = self.ray_query_pointer(args.next()?, ctx)?;
2957 args.finish()?;
2958
2959 let _ = ctx.module.generate_vertex_return_type();
2960
2961 ir::Expression::RayQueryVertexPositions {
2962 query,
2963 committed: false,
2964 }
2965 }
2966 "rayQueryProceed" => {
2967 let mut args = ctx.prepare_args(arguments, 1, span);
2968 let query = self.ray_query_pointer(args.next()?, ctx)?;
2969 args.finish()?;
2970
2971 let result =
2972 ctx.interrupt_emitter(ir::Expression::RayQueryProceedResult, span)?;
2973 let fun = ir::RayQueryFunction::Proceed { result };
2974 let rctx = ctx.runtime_expression_ctx(span)?;
2975 rctx.block
2976 .push(ir::Statement::RayQuery { query, fun }, span);
2977 return Ok(Some(result));
2978 }
2979 "rayQueryGenerateIntersection" => {
2980 let mut args = ctx.prepare_args(arguments, 2, span);
2981 let query = self.ray_query_pointer(args.next()?, ctx)?;
2982 let hit_t = self.expression(args.next()?, ctx)?;
2983 args.finish()?;
2984
2985 let fun = ir::RayQueryFunction::GenerateIntersection { hit_t };
2986 let rctx = ctx.runtime_expression_ctx(span)?;
2987 rctx.block
2988 .push(ir::Statement::RayQuery { query, fun }, span);
2989 return Ok(None);
2990 }
2991 "rayQueryConfirmIntersection" => {
2992 let mut args = ctx.prepare_args(arguments, 1, span);
2993 let query = self.ray_query_pointer(args.next()?, ctx)?;
2994 args.finish()?;
2995
2996 let fun = ir::RayQueryFunction::ConfirmIntersection;
2997 let rctx = ctx.runtime_expression_ctx(span)?;
2998 rctx.block
2999 .push(ir::Statement::RayQuery { query, fun }, span);
3000 return Ok(None);
3001 }
3002 "rayQueryTerminate" => {
3003 let mut args = ctx.prepare_args(arguments, 1, span);
3004 let query = self.ray_query_pointer(args.next()?, ctx)?;
3005 args.finish()?;
3006
3007 let fun = ir::RayQueryFunction::Terminate;
3008 let rctx = ctx.runtime_expression_ctx(span)?;
3009 rctx.block
3010 .push(ir::Statement::RayQuery { query, fun }, span);
3011 return Ok(None);
3012 }
3013 "rayQueryGetCommittedIntersection" => {
3014 let mut args = ctx.prepare_args(arguments, 1, span);
3015 let query = self.ray_query_pointer(args.next()?, ctx)?;
3016 args.finish()?;
3017
3018 let _ = ctx.module.generate_ray_intersection_type();
3019 ir::Expression::RayQueryGetIntersection {
3020 query,
3021 committed: true,
3022 }
3023 }
3024 "rayQueryGetCandidateIntersection" => {
3025 let mut args = ctx.prepare_args(arguments, 1, span);
3026 let query = self.ray_query_pointer(args.next()?, ctx)?;
3027 args.finish()?;
3028
3029 let _ = ctx.module.generate_ray_intersection_type();
3030 ir::Expression::RayQueryGetIntersection {
3031 query,
3032 committed: false,
3033 }
3034 }
3035 "RayDesc" => {
3036 let ty = ctx.module.generate_ray_desc_type();
3037 let handle = self.construct(
3038 span,
3039 &ast::ConstructorType::Type(ty),
3040 function.span,
3041 arguments,
3042 ctx,
3043 )?;
3044 return Ok(Some(handle));
3045 }
3046 "subgroupBallot" => {
3047 let mut args = ctx.prepare_args(arguments, 0, span);
3048 let predicate = if arguments.len() == 1 {
3049 Some(self.expression(args.next()?, ctx)?)
3050 } else {
3051 None
3052 };
3053 args.finish()?;
3054
3055 let result =
3056 ctx.interrupt_emitter(ir::Expression::SubgroupBallotResult, span)?;
3057 let rctx = ctx.runtime_expression_ctx(span)?;
3058 rctx.block
3059 .push(ir::Statement::SubgroupBallot { result, predicate }, span);
3060 return Ok(Some(result));
3061 }
3062 "quadSwapX" => {
3063 let mut args = ctx.prepare_args(arguments, 1, span);
3064
3065 let argument = self.expression(args.next()?, ctx)?;
3066 args.finish()?;
3067
3068 let ty = ctx.register_type(argument)?;
3069
3070 let result = ctx.interrupt_emitter(
3071 crate::Expression::SubgroupOperationResult { ty },
3072 span,
3073 )?;
3074 let rctx = ctx.runtime_expression_ctx(span)?;
3075 rctx.block.push(
3076 crate::Statement::SubgroupGather {
3077 mode: crate::GatherMode::QuadSwap(crate::Direction::X),
3078 argument,
3079 result,
3080 },
3081 span,
3082 );
3083 return Ok(Some(result));
3084 }
3085
3086 "quadSwapY" => {
3087 let mut args = ctx.prepare_args(arguments, 1, span);
3088
3089 let argument = self.expression(args.next()?, ctx)?;
3090 args.finish()?;
3091
3092 let ty = ctx.register_type(argument)?;
3093
3094 let result = ctx.interrupt_emitter(
3095 crate::Expression::SubgroupOperationResult { ty },
3096 span,
3097 )?;
3098 let rctx = ctx.runtime_expression_ctx(span)?;
3099 rctx.block.push(
3100 crate::Statement::SubgroupGather {
3101 mode: crate::GatherMode::QuadSwap(crate::Direction::Y),
3102 argument,
3103 result,
3104 },
3105 span,
3106 );
3107 return Ok(Some(result));
3108 }
3109
3110 "quadSwapDiagonal" => {
3111 let mut args = ctx.prepare_args(arguments, 1, span);
3112
3113 let argument = self.expression(args.next()?, ctx)?;
3114 args.finish()?;
3115
3116 let ty = ctx.register_type(argument)?;
3117
3118 let result = ctx.interrupt_emitter(
3119 crate::Expression::SubgroupOperationResult { ty },
3120 span,
3121 )?;
3122 let rctx = ctx.runtime_expression_ctx(span)?;
3123 rctx.block.push(
3124 crate::Statement::SubgroupGather {
3125 mode: crate::GatherMode::QuadSwap(crate::Direction::Diagonal),
3126 argument,
3127 result,
3128 },
3129 span,
3130 );
3131 return Ok(Some(result));
3132 }
3133 _ => {
3134 return Err(Box::new(Error::UnknownIdent(function.span, function.name)))
3135 }
3136 }
3137 };
3138
3139 let expr = ctx.append_expression(expr, span)?;
3140 Ok(Some(expr))
3141 }
3142 }
3143 }
3144
3145 fn math_function_helper(
3156 &mut self,
3157 span: Span,
3158 fun: ir::MathFunction,
3159 ast_arguments: &[Handle<ast::Expression<'source>>],
3160 ctx: &mut ExpressionContext<'source, '_, '_>,
3161 ) -> Result<'source, ir::Expression> {
3162 let mut lowered_arguments = Vec::with_capacity(ast_arguments.len());
3163 for &arg in ast_arguments {
3164 let lowered = self.expression_for_abstract(arg, ctx)?;
3165 ctx.grow_types(lowered)?;
3166 lowered_arguments.push(lowered);
3167 }
3168
3169 let fun_overloads = fun.overloads();
3170 let rule = self.resolve_overloads(span, fun, fun_overloads, &lowered_arguments, ctx)?;
3171 self.apply_automatic_conversions_for_call(&rule, &mut lowered_arguments, ctx)?;
3172
3173 if let proc::Conclusion::Predeclared(predeclared) = rule.conclusion {
3177 ctx.module.generate_predeclared_type(predeclared);
3178 }
3179
3180 Ok(ir::Expression::Math {
3181 fun,
3182 arg: lowered_arguments[0],
3183 arg1: lowered_arguments.get(1).cloned(),
3184 arg2: lowered_arguments.get(2).cloned(),
3185 arg3: lowered_arguments.get(3).cloned(),
3186 })
3187 }
3188
3189 fn resolve_overloads<O, F>(
3200 &self,
3201 span: Span,
3202 fun: F,
3203 overloads: O,
3204 arguments: &[Handle<ir::Expression>],
3205 ctx: &ExpressionContext<'source, '_, '_>,
3206 ) -> Result<'source, proc::Rule>
3207 where
3208 O: proc::OverloadSet,
3209 F: TryToWgsl + core::fmt::Debug + Copy,
3210 {
3211 let mut remaining_overloads = overloads.clone();
3212 let min_arguments = remaining_overloads.min_arguments();
3213 let max_arguments = remaining_overloads.max_arguments();
3214 if arguments.len() < min_arguments {
3215 return Err(Box::new(Error::WrongArgumentCount {
3216 span,
3217 expected: min_arguments as u32..max_arguments as u32,
3218 found: arguments.len() as u32,
3219 }));
3220 }
3221 if arguments.len() > max_arguments {
3222 return Err(Box::new(Error::TooManyArguments {
3223 function: fun.to_wgsl_for_diagnostics(),
3224 call_span: span,
3225 arg_span: ctx.get_expression_span(arguments[max_arguments]),
3226 max_arguments: max_arguments as _,
3227 }));
3228 }
3229
3230 log::debug!(
3231 "Initial overloads: {:#?}",
3232 remaining_overloads.for_debug(&ctx.module.types)
3233 );
3234
3235 for (arg_index, &arg) in arguments.iter().enumerate() {
3236 let arg_type_resolution = &ctx.typifier()[arg];
3237 let arg_inner = arg_type_resolution.inner_with(&ctx.module.types);
3238 log::debug!(
3239 "Supplying argument {arg_index} of type {:?}",
3240 arg_type_resolution.for_debug(&ctx.module.types)
3241 );
3242 let next_remaining_overloads =
3243 remaining_overloads.arg(arg_index, arg_inner, &ctx.module.types);
3244
3245 log::debug!(
3253 "Remaining overloads: {:#?}",
3254 next_remaining_overloads.for_debug(&ctx.module.types)
3255 );
3256
3257 if next_remaining_overloads.is_empty() {
3260 let function = fun.to_wgsl_for_diagnostics();
3261 let call_span = span;
3262 let arg_span = ctx.get_expression_span(arg);
3263 let arg_ty = ctx.as_diagnostic_display(arg_type_resolution).to_string();
3264
3265 let only_this_argument = overloads.arg(arg_index, arg_inner, &ctx.module.types);
3268 if only_this_argument.is_empty() {
3269 let allowed: Vec<String> = overloads
3273 .allowed_args(arg_index, &ctx.module.to_ctx())
3274 .iter()
3275 .map(|ty| ctx.type_resolution_to_string(ty))
3276 .collect();
3277
3278 if allowed.is_empty() {
3279 unreachable!("expected all overloads to have the same arity");
3285 }
3286
3287 return Err(Box::new(Error::WrongArgumentType {
3290 function,
3291 call_span,
3292 arg_span,
3293 arg_index: arg_index as u32,
3294 arg_ty,
3295 allowed,
3296 }));
3297 }
3298
3299 let allowed: Vec<String> = remaining_overloads
3307 .allowed_args(arg_index, &ctx.module.to_ctx())
3308 .iter()
3309 .map(|ty| ctx.type_resolution_to_string(ty))
3310 .collect();
3311
3312 let mut remaining_overloads = overloads;
3315 for (prior_index, &prior_expr) in arguments.iter().enumerate() {
3316 let prior_type_resolution = &ctx.typifier()[prior_expr];
3317 let prior_ty = prior_type_resolution.inner_with(&ctx.module.types);
3318 remaining_overloads =
3319 remaining_overloads.arg(prior_index, prior_ty, &ctx.module.types);
3320 if remaining_overloads
3321 .arg(arg_index, arg_inner, &ctx.module.types)
3322 .is_empty()
3323 {
3324 let inconsistent_span = ctx.get_expression_span(arguments[prior_index]);
3326 let inconsistent_ty =
3327 ctx.as_diagnostic_display(prior_type_resolution).to_string();
3328
3329 if allowed.is_empty() {
3330 unreachable!("expected all overloads to have the same arity");
3337 }
3338
3339 return Err(Box::new(Error::InconsistentArgumentType {
3341 function,
3342 call_span,
3343 arg_span,
3344 arg_index: arg_index as u32,
3345 arg_ty,
3346 inconsistent_span,
3347 inconsistent_index: prior_index as u32,
3348 inconsistent_ty,
3349 allowed,
3350 }));
3351 }
3352 }
3353 unreachable!("Failed to eliminate argument type when re-tried");
3354 }
3355 remaining_overloads = next_remaining_overloads;
3356 }
3357
3358 Ok(remaining_overloads.most_preferred())
3361 }
3362
3363 fn apply_automatic_conversions_for_call(
3369 &self,
3370 rule: &proc::Rule,
3371 arguments: &mut [Handle<ir::Expression>],
3372 ctx: &mut ExpressionContext<'source, '_, '_>,
3373 ) -> Result<'source, ()> {
3374 for (i, argument) in arguments.iter_mut().enumerate() {
3375 let goal_inner = rule.arguments[i].inner_with(&ctx.module.types);
3376 let converted = match goal_inner.scalar_for_conversions(&ctx.module.types) {
3377 Some(goal_scalar) => {
3378 let arg_span = ctx.get_expression_span(*argument);
3379 ctx.try_automatic_conversion_for_leaf_scalar(*argument, goal_scalar, arg_span)?
3380 }
3381 None => *argument,
3383 };
3384
3385 *argument = converted;
3386 }
3387
3388 Ok(())
3389 }
3390
3391 fn atomic_pointer(
3392 &mut self,
3393 expr: Handle<ast::Expression<'source>>,
3394 ctx: &mut ExpressionContext<'source, '_, '_>,
3395 ) -> Result<'source, (Handle<ir::Expression>, ir::Scalar)> {
3396 let span = ctx.ast_expressions.get_span(expr);
3397 let pointer = self.expression(expr, ctx)?;
3398
3399 match *resolve_inner!(ctx, pointer) {
3400 ir::TypeInner::Pointer { base, .. } => match ctx.module.types[base].inner {
3401 ir::TypeInner::Atomic(scalar) => Ok((pointer, scalar)),
3402 ref other => {
3403 log::error!("Pointer type to {other:?} passed to atomic op");
3404 Err(Box::new(Error::InvalidAtomicPointer(span)))
3405 }
3406 },
3407 ref other => {
3408 log::error!("Type {other:?} passed to atomic op");
3409 Err(Box::new(Error::InvalidAtomicPointer(span)))
3410 }
3411 }
3412 }
3413
3414 fn atomic_helper(
3415 &mut self,
3416 span: Span,
3417 fun: ir::AtomicFunction,
3418 args: &[Handle<ast::Expression<'source>>],
3419 is_statement: bool,
3420 ctx: &mut ExpressionContext<'source, '_, '_>,
3421 ) -> Result<'source, Option<Handle<ir::Expression>>> {
3422 let mut args = ctx.prepare_args(args, 2, span);
3423
3424 let (pointer, scalar) = self.atomic_pointer(args.next()?, ctx)?;
3425 let value = self.expression_with_leaf_scalar(args.next()?, scalar, ctx)?;
3426 let value_inner = resolve_inner!(ctx, value);
3427 args.finish()?;
3428
3429 let is_64_bit_min_max = matches!(fun, ir::AtomicFunction::Min | ir::AtomicFunction::Max)
3434 && matches!(
3435 *value_inner,
3436 ir::TypeInner::Scalar(ir::Scalar { width: 8, .. })
3437 );
3438 let result = if is_64_bit_min_max && is_statement {
3439 let rctx = ctx.runtime_expression_ctx(span)?;
3440 rctx.block
3441 .extend(rctx.emitter.finish(&rctx.function.expressions));
3442 rctx.emitter.start(&rctx.function.expressions);
3443 None
3444 } else {
3445 let ty = ctx.register_type(value)?;
3446 Some(ctx.interrupt_emitter(
3447 ir::Expression::AtomicResult {
3448 ty,
3449 comparison: false,
3450 },
3451 span,
3452 )?)
3453 };
3454 let rctx = ctx.runtime_expression_ctx(span)?;
3455 rctx.block.push(
3456 ir::Statement::Atomic {
3457 pointer,
3458 fun,
3459 value,
3460 result,
3461 },
3462 span,
3463 );
3464 Ok(result)
3465 }
3466
3467 fn texture_sample_helper(
3468 &mut self,
3469 fun: Texture,
3470 args: &[Handle<ast::Expression<'source>>],
3471 span: Span,
3472 ctx: &mut ExpressionContext<'source, '_, '_>,
3473 ) -> Result<'source, ir::Expression> {
3474 let mut args = ctx.prepare_args(args, fun.min_argument_count(), span);
3475
3476 fn get_image_and_span<'source>(
3477 lowerer: &mut Lowerer<'source, '_>,
3478 args: &mut ArgumentContext<'_, 'source>,
3479 ctx: &mut ExpressionContext<'source, '_, '_>,
3480 ) -> Result<'source, (Handle<ir::Expression>, Span)> {
3481 let image = args.next()?;
3482 let image_span = ctx.ast_expressions.get_span(image);
3483 let image = lowerer.expression_for_abstract(image, ctx)?;
3484 Ok((image, image_span))
3485 }
3486
3487 let image;
3488 let image_span;
3489 let gather;
3490 match fun {
3491 Texture::Gather => {
3492 let image_or_component = args.next()?;
3493 let image_or_component_span = ctx.ast_expressions.get_span(image_or_component);
3494 let lowered_image_or_component = self.expression(image_or_component, ctx)?;
3496
3497 match *resolve_inner!(ctx, lowered_image_or_component) {
3498 ir::TypeInner::Image {
3499 class: ir::ImageClass::Depth { .. },
3500 ..
3501 } => {
3502 image = lowered_image_or_component;
3503 image_span = image_or_component_span;
3504 gather = Some(ir::SwizzleComponent::X);
3505 }
3506 _ => {
3507 (image, image_span) = get_image_and_span(self, &mut args, ctx)?;
3508 gather = Some(ctx.gather_component(
3509 lowered_image_or_component,
3510 image_or_component_span,
3511 span,
3512 )?);
3513 }
3514 }
3515 }
3516 Texture::GatherCompare => {
3517 (image, image_span) = get_image_and_span(self, &mut args, ctx)?;
3518 gather = Some(ir::SwizzleComponent::X);
3519 }
3520
3521 _ => {
3522 (image, image_span) = get_image_and_span(self, &mut args, ctx)?;
3523 gather = None;
3524 }
3525 };
3526
3527 let sampler = self.expression_for_abstract(args.next()?, ctx)?;
3528
3529 let coordinate = self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?;
3530 let clamp_to_edge = matches!(fun, Texture::SampleBaseClampToEdge);
3531
3532 let (class, arrayed) = ctx.image_data(image, image_span)?;
3533 let array_index = arrayed
3534 .then(|| self.expression(args.next()?, ctx))
3535 .transpose()?;
3536
3537 let level;
3538 let depth_ref;
3539 match fun {
3540 Texture::Gather => {
3541 level = ir::SampleLevel::Zero;
3542 depth_ref = None;
3543 }
3544 Texture::GatherCompare => {
3545 let reference =
3546 self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?;
3547 level = ir::SampleLevel::Zero;
3548 depth_ref = Some(reference);
3549 }
3550
3551 Texture::Sample => {
3552 level = ir::SampleLevel::Auto;
3553 depth_ref = None;
3554 }
3555 Texture::SampleBias => {
3556 let bias = self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?;
3557 level = ir::SampleLevel::Bias(bias);
3558 depth_ref = None;
3559 }
3560 Texture::SampleCompare => {
3561 let reference =
3562 self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?;
3563 level = ir::SampleLevel::Auto;
3564 depth_ref = Some(reference);
3565 }
3566 Texture::SampleCompareLevel => {
3567 let reference =
3568 self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?;
3569 level = ir::SampleLevel::Zero;
3570 depth_ref = Some(reference);
3571 }
3572 Texture::SampleGrad => {
3573 let x = self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?;
3574 let y = self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?;
3575 level = ir::SampleLevel::Gradient { x, y };
3576 depth_ref = None;
3577 }
3578 Texture::SampleLevel => {
3579 let exact = match class {
3580 ir::ImageClass::Depth { .. } => self.expression(args.next()?, ctx)?,
3583
3584 ir::ImageClass::Sampled { .. } => {
3587 self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?
3588 }
3589
3590 ir::ImageClass::Storage { .. } | ir::ImageClass::External => {
3594 self.expression(args.next()?, ctx)?
3595 }
3596 };
3597 level = ir::SampleLevel::Exact(exact);
3598 depth_ref = None;
3599 }
3600 Texture::SampleBaseClampToEdge => {
3601 level = crate::SampleLevel::Zero;
3602 depth_ref = None;
3603 }
3604 };
3605
3606 let offset = args
3607 .next()
3608 .map(|arg| self.expression_with_leaf_scalar(arg, ir::Scalar::I32, &mut ctx.as_const()))
3609 .ok()
3610 .transpose()?;
3611
3612 args.finish()?;
3613
3614 Ok(ir::Expression::ImageSample {
3615 image,
3616 sampler,
3617 gather,
3618 coordinate,
3619 array_index,
3620 offset,
3621 level,
3622 depth_ref,
3623 clamp_to_edge,
3624 })
3625 }
3626
3627 fn subgroup_operation_helper(
3628 &mut self,
3629 span: Span,
3630 op: ir::SubgroupOperation,
3631 collective_op: ir::CollectiveOperation,
3632 arguments: &[Handle<ast::Expression<'source>>],
3633 ctx: &mut ExpressionContext<'source, '_, '_>,
3634 ) -> Result<'source, Handle<ir::Expression>> {
3635 let mut args = ctx.prepare_args(arguments, 1, span);
3636
3637 let argument = self.expression(args.next()?, ctx)?;
3638 args.finish()?;
3639
3640 let ty = ctx.register_type(argument)?;
3641
3642 let result = ctx.interrupt_emitter(ir::Expression::SubgroupOperationResult { ty }, span)?;
3643 let rctx = ctx.runtime_expression_ctx(span)?;
3644 rctx.block.push(
3645 ir::Statement::SubgroupCollectiveOperation {
3646 op,
3647 collective_op,
3648 argument,
3649 result,
3650 },
3651 span,
3652 );
3653 Ok(result)
3654 }
3655
3656 fn subgroup_gather_helper(
3657 &mut self,
3658 span: Span,
3659 mode: SubgroupGather,
3660 arguments: &[Handle<ast::Expression<'source>>],
3661 ctx: &mut ExpressionContext<'source, '_, '_>,
3662 ) -> Result<'source, Handle<ir::Expression>> {
3663 let mut args = ctx.prepare_args(arguments, 2, span);
3664
3665 let argument = self.expression(args.next()?, ctx)?;
3666
3667 use SubgroupGather as Sg;
3668 let mode = if let Sg::BroadcastFirst = mode {
3669 ir::GatherMode::BroadcastFirst
3670 } else {
3671 let index = self.expression(args.next()?, ctx)?;
3672 match mode {
3673 Sg::BroadcastFirst => unreachable!(),
3674 Sg::Broadcast => ir::GatherMode::Broadcast(index),
3675 Sg::Shuffle => ir::GatherMode::Shuffle(index),
3676 Sg::ShuffleDown => ir::GatherMode::ShuffleDown(index),
3677 Sg::ShuffleUp => ir::GatherMode::ShuffleUp(index),
3678 Sg::ShuffleXor => ir::GatherMode::ShuffleXor(index),
3679 Sg::QuadBroadcast => ir::GatherMode::QuadBroadcast(index),
3680 }
3681 };
3682
3683 args.finish()?;
3684
3685 let ty = ctx.register_type(argument)?;
3686
3687 let result = ctx.interrupt_emitter(ir::Expression::SubgroupOperationResult { ty }, span)?;
3688 let rctx = ctx.runtime_expression_ctx(span)?;
3689 rctx.block.push(
3690 ir::Statement::SubgroupGather {
3691 mode,
3692 argument,
3693 result,
3694 },
3695 span,
3696 );
3697 Ok(result)
3698 }
3699
3700 fn r#struct(
3701 &mut self,
3702 s: &ast::Struct<'source>,
3703 span: Span,
3704 ctx: &mut GlobalContext<'source, '_, '_>,
3705 ) -> Result<'source, Handle<ir::Type>> {
3706 let mut offset = 0;
3707 let mut struct_alignment = proc::Alignment::ONE;
3708 let mut members = Vec::with_capacity(s.members.len());
3709
3710 let mut doc_comments: Vec<Option<Vec<String>>> = Vec::new();
3711
3712 for member in s.members.iter() {
3713 let ty = self.resolve_ast_type(member.ty, &mut ctx.as_const())?;
3714
3715 ctx.layouter.update(ctx.module.to_ctx()).map_err(|err| {
3716 let LayoutErrorInner::TooLarge = err.inner else {
3717 unreachable!("unexpected layout error: {err:?}");
3718 };
3719 if ty == err.ty {
3725 Box::new(Error::StructMemberTooLarge {
3726 member_name_span: member.name.span,
3727 })
3728 } else {
3729 Box::new(Error::TypeTooLarge {
3732 span: ctx.module.types.get_span(err.ty),
3733 })
3734 }
3735 })?;
3736
3737 let member_min_size = ctx.layouter[ty].size;
3738 let member_min_alignment = ctx.layouter[ty].alignment;
3739
3740 let member_size = if let Some(size_expr) = member.size {
3741 let (size, span) = self.const_u32(size_expr, &mut ctx.as_const())?;
3742 if size < member_min_size {
3743 return Err(Box::new(Error::SizeAttributeTooLow(span, member_min_size)));
3744 } else {
3745 size
3746 }
3747 } else {
3748 member_min_size
3749 };
3750
3751 let member_alignment = if let Some(align_expr) = member.align {
3752 let (align, span) = self.const_u32(align_expr, &mut ctx.as_const())?;
3753 if let Some(alignment) = proc::Alignment::new(align) {
3754 if alignment < member_min_alignment {
3755 return Err(Box::new(Error::AlignAttributeTooLow(
3756 span,
3757 member_min_alignment,
3758 )));
3759 } else {
3760 alignment
3761 }
3762 } else {
3763 return Err(Box::new(Error::NonPowerOfTwoAlignAttribute(span)));
3764 }
3765 } else {
3766 member_min_alignment
3767 };
3768
3769 let binding = self.binding(&member.binding, ty, ctx)?;
3770
3771 offset = member_alignment.round_up(offset);
3772 struct_alignment = struct_alignment.max(member_alignment);
3773
3774 if !member.doc_comments.is_empty() {
3775 doc_comments.push(Some(
3776 member.doc_comments.iter().map(|s| s.to_string()).collect(),
3777 ));
3778 }
3779 members.push(ir::StructMember {
3780 name: Some(member.name.name.to_owned()),
3781 ty,
3782 binding,
3783 offset,
3784 });
3785
3786 offset += member_size;
3787 if offset > crate::valid::MAX_TYPE_SIZE {
3788 return Err(Box::new(Error::TypeTooLarge { span }));
3789 }
3790 }
3791
3792 let size = struct_alignment.round_up(offset);
3793 let inner = ir::TypeInner::Struct {
3794 members,
3795 span: size,
3796 };
3797
3798 let handle = ctx.module.types.insert(
3799 ir::Type {
3800 name: Some(s.name.name.to_string()),
3801 inner,
3802 },
3803 span,
3804 );
3805 for (i, c) in doc_comments.drain(..).enumerate() {
3806 if let Some(comment) = c {
3807 ctx.module
3808 .get_or_insert_default_doc_comments()
3809 .struct_members
3810 .insert((handle, i), comment);
3811 }
3812 }
3813 Ok(handle)
3814 }
3815
3816 fn const_u32(
3817 &mut self,
3818 expr: Handle<ast::Expression<'source>>,
3819 ctx: &mut ExpressionContext<'source, '_, '_>,
3820 ) -> Result<'source, (u32, Span)> {
3821 let span = ctx.ast_expressions.get_span(expr);
3822 let expr = self.expression(expr, ctx)?;
3823 let value = ctx
3824 .module
3825 .to_ctx()
3826 .eval_expr_to_u32(expr)
3827 .map_err(|err| match err {
3828 proc::U32EvalError::NonConst => Error::ExpectedConstExprConcreteIntegerScalar(span),
3829 proc::U32EvalError::Negative => Error::ExpectedNonNegative(span),
3830 })?;
3831 Ok((value, span))
3832 }
3833
3834 fn array_size(
3835 &mut self,
3836 size: ast::ArraySize<'source>,
3837 ctx: &mut ExpressionContext<'source, '_, '_>,
3838 ) -> Result<'source, ir::ArraySize> {
3839 Ok(match size {
3840 ast::ArraySize::Constant(expr) => {
3841 let span = ctx.ast_expressions.get_span(expr);
3842 let const_expr = self.expression(expr, &mut ctx.as_const());
3843 match const_expr {
3844 Ok(value) => {
3845 let len = ctx.const_eval_expr_to_u32(value).map_err(|err| {
3846 Box::new(match err {
3847 proc::U32EvalError::NonConst => {
3848 Error::ExpectedConstExprConcreteIntegerScalar(span)
3849 }
3850 proc::U32EvalError::Negative => {
3851 Error::ExpectedPositiveArrayLength(span)
3852 }
3853 })
3854 })?;
3855 let size =
3856 NonZeroU32::new(len).ok_or(Error::ExpectedPositiveArrayLength(span))?;
3857 ir::ArraySize::Constant(size)
3858 }
3859 Err(err) => {
3860 if let Error::ConstantEvaluatorError(ref ty, _) = *err {
3861 match **ty {
3862 proc::ConstantEvaluatorError::OverrideExpr => {
3863 ir::ArraySize::Pending(self.array_size_override(
3864 expr,
3865 &mut ctx.as_global().as_override(),
3866 span,
3867 )?)
3868 }
3869 _ => {
3870 return Err(err);
3871 }
3872 }
3873 } else {
3874 return Err(err);
3875 }
3876 }
3877 }
3878 }
3879 ast::ArraySize::Dynamic => ir::ArraySize::Dynamic,
3880 })
3881 }
3882
3883 fn array_size_override(
3884 &mut self,
3885 size_expr: Handle<ast::Expression<'source>>,
3886 ctx: &mut ExpressionContext<'source, '_, '_>,
3887 span: Span,
3888 ) -> Result<'source, Handle<ir::Override>> {
3889 let expr = self.expression(size_expr, ctx)?;
3890 match resolve_inner!(ctx, expr).scalar_kind().ok_or(0) {
3891 Ok(ir::ScalarKind::Sint) | Ok(ir::ScalarKind::Uint) => Ok({
3892 if let ir::Expression::Override(handle) = ctx.module.global_expressions[expr] {
3893 handle
3894 } else {
3895 let ty = ctx.register_type(expr)?;
3896 ctx.module.overrides.append(
3897 ir::Override {
3898 name: None,
3899 id: None,
3900 ty,
3901 init: Some(expr),
3902 },
3903 span,
3904 )
3905 }
3906 }),
3907 _ => Err(Box::new(Error::ExpectedConstExprConcreteIntegerScalar(
3908 span,
3909 ))),
3910 }
3911 }
3912
3913 fn resolve_named_ast_type(
3923 &mut self,
3924 handle: Handle<ast::Type<'source>>,
3925 name: Option<String>,
3926 ctx: &mut ExpressionContext<'source, '_, '_>,
3927 ) -> Result<'source, Handle<ir::Type>> {
3928 let inner = match ctx.types[handle] {
3929 ast::Type::Scalar(scalar) => scalar.to_inner_scalar(),
3930 ast::Type::Vector { size, ty, ty_span } => {
3931 let ty = self.resolve_ast_type(ty, ctx)?;
3932 let scalar = match ctx.module.types[ty].inner {
3933 ir::TypeInner::Scalar(sc) => sc,
3934 _ => return Err(Box::new(Error::UnknownScalarType(ty_span))),
3935 };
3936 ir::TypeInner::Vector { size, scalar }
3937 }
3938 ast::Type::Matrix {
3939 rows,
3940 columns,
3941 ty,
3942 ty_span,
3943 } => {
3944 let ty = self.resolve_ast_type(ty, ctx)?;
3945 let scalar = match ctx.module.types[ty].inner {
3946 ir::TypeInner::Scalar(sc) => sc,
3947 _ => return Err(Box::new(Error::UnknownScalarType(ty_span))),
3948 };
3949 match scalar.kind {
3950 ir::ScalarKind::Float => ir::TypeInner::Matrix {
3951 columns,
3952 rows,
3953 scalar,
3954 },
3955 _ => return Err(Box::new(Error::BadMatrixScalarKind(ty_span, scalar))),
3956 }
3957 }
3958 ast::Type::Atomic(scalar) => scalar.to_inner_atomic(),
3959 ast::Type::Pointer { base, space } => {
3960 let base = self.resolve_ast_type(base, ctx)?;
3961 ir::TypeInner::Pointer { base, space }
3962 }
3963 ast::Type::Array { base, size } => {
3964 let base = self.resolve_ast_type(base, &mut ctx.as_const())?;
3965 let size = self.array_size(size, ctx)?;
3966
3967 ctx.layouter.update(ctx.module.to_ctx()).map_err(|err| {
3969 let LayoutErrorInner::TooLarge = err.inner else {
3970 unreachable!("unexpected layout error: {err:?}");
3971 };
3972 Box::new(Error::TypeTooLarge {
3975 span: ctx.module.types.get_span(err.ty),
3976 })
3977 })?;
3978 let stride = ctx.layouter[base].to_stride();
3979
3980 ir::TypeInner::Array { base, size, stride }
3981 }
3982 ast::Type::Image {
3983 dim,
3984 arrayed,
3985 class,
3986 } => ir::TypeInner::Image {
3987 dim,
3988 arrayed,
3989 class,
3990 },
3991 ast::Type::Sampler { comparison } => ir::TypeInner::Sampler { comparison },
3992 ast::Type::AccelerationStructure { vertex_return } => {
3993 ir::TypeInner::AccelerationStructure { vertex_return }
3994 }
3995 ast::Type::RayQuery { vertex_return } => ir::TypeInner::RayQuery { vertex_return },
3996 ast::Type::BindingArray { base, size } => {
3997 let base = self.resolve_ast_type(base, ctx)?;
3998 let size = self.array_size(size, ctx)?;
3999 ir::TypeInner::BindingArray { base, size }
4000 }
4001 ast::Type::RayDesc => {
4002 return Ok(ctx.module.generate_ray_desc_type());
4003 }
4004 ast::Type::RayIntersection => {
4005 return Ok(ctx.module.generate_ray_intersection_type());
4006 }
4007 ast::Type::User(ref ident) => {
4008 return match ctx.globals.get(ident.name) {
4009 Some(&LoweredGlobalDecl::Type(handle)) => Ok(handle),
4010 Some(_) => Err(Box::new(Error::Unexpected(ident.span, ExpectedToken::Type))),
4011 None => Err(Box::new(Error::UnknownType(ident.span))),
4012 }
4013 }
4014 };
4015
4016 Ok(ctx.as_global().ensure_type_exists(name, inner))
4017 }
4018
4019 fn resolve_ast_type(
4021 &mut self,
4022 handle: Handle<ast::Type<'source>>,
4023 ctx: &mut ExpressionContext<'source, '_, '_>,
4024 ) -> Result<'source, Handle<ir::Type>> {
4025 self.resolve_named_ast_type(handle, None, ctx)
4026 }
4027
4028 fn binding(
4029 &mut self,
4030 binding: &Option<ast::Binding<'source>>,
4031 ty: Handle<ir::Type>,
4032 ctx: &mut GlobalContext<'source, '_, '_>,
4033 ) -> Result<'source, Option<ir::Binding>> {
4034 Ok(match *binding {
4035 Some(ast::Binding::BuiltIn(b)) => Some(ir::Binding::BuiltIn(b)),
4036 Some(ast::Binding::Location {
4037 location,
4038 interpolation,
4039 sampling,
4040 blend_src,
4041 }) => {
4042 let blend_src = if let Some(blend_src) = blend_src {
4043 Some(self.const_u32(blend_src, &mut ctx.as_const())?.0)
4044 } else {
4045 None
4046 };
4047
4048 let mut binding = ir::Binding::Location {
4049 location: self.const_u32(location, &mut ctx.as_const())?.0,
4050 interpolation,
4051 sampling,
4052 blend_src,
4053 };
4054 binding.apply_default_interpolation(&ctx.module.types[ty].inner);
4055 Some(binding)
4056 }
4057 None => None,
4058 })
4059 }
4060
4061 fn ray_query_pointer(
4062 &mut self,
4063 expr: Handle<ast::Expression<'source>>,
4064 ctx: &mut ExpressionContext<'source, '_, '_>,
4065 ) -> Result<'source, Handle<ir::Expression>> {
4066 let span = ctx.ast_expressions.get_span(expr);
4067 let pointer = self.expression(expr, ctx)?;
4068
4069 match *resolve_inner!(ctx, pointer) {
4070 ir::TypeInner::Pointer { base, .. } => match ctx.module.types[base].inner {
4071 ir::TypeInner::RayQuery { .. } => Ok(pointer),
4072 ref other => {
4073 log::error!("Pointer type to {other:?} passed to ray query op");
4074 Err(Box::new(Error::InvalidRayQueryPointer(span)))
4075 }
4076 },
4077 ref other => {
4078 log::error!("Type {other:?} passed to ray query op");
4079 Err(Box::new(Error::InvalidRayQueryPointer(span)))
4080 }
4081 }
4082 }
4083}
4084
4085impl ir::AtomicFunction {
4086 pub fn map(word: &str) -> Option<Self> {
4087 Some(match word {
4088 "atomicAdd" => ir::AtomicFunction::Add,
4089 "atomicSub" => ir::AtomicFunction::Subtract,
4090 "atomicAnd" => ir::AtomicFunction::And,
4091 "atomicOr" => ir::AtomicFunction::InclusiveOr,
4092 "atomicXor" => ir::AtomicFunction::ExclusiveOr,
4093 "atomicMin" => ir::AtomicFunction::Min,
4094 "atomicMax" => ir::AtomicFunction::Max,
4095 "atomicExchange" => ir::AtomicFunction::Exchange { compare: None },
4096 _ => return None,
4097 })
4098 }
4099}